我們學任何一門語言和框架箱玷，都要從hello world學起怨规。cuda的hello world可以是：

#include <cstdio>
#include <algorithm>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>

__global__ void add(float *dx, float *dy)
{
    int id = blockIdx.x * blockDim.x + threadIdx.x;
    dy[id] += dx[id];
}

int main()
{
    int N = 1 << 20;

    float *hx, *hy, *dx, *dy;

    hx = new float[N];
    hy = new float[N];
    cudaMalloc(&dx, N * sizeof(float));
    cudaMalloc(&dy, N * sizeof(float));

    for (int i = 0; i < N; i++)
    {
        hx[i] = 1.0f;
        hy[i] = 2.0f;
    }

    cudaMemcpy(dx, hx, N * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(dy, hy, N * sizeof(float), cudaMemcpyHostToDevice);

    int threadNums = 256;
    int blockNums = (N + threadNums - 1) / threadNums;

    add << <blockNums, threadNums >> >(N, dx, dy);

    cudaDeviceSynchronize();

    cudaMemcpy(hy, dy, N * sizeof(float), cudaMemcpyDeviceToHost);

    float maxError = 0.0f;
    for (int i = 0; i < N; i++)
        maxError = std::max(maxError, std::abs(hy[i] - 3.0f));

    printf("Max error: %.6f", maxError);

    delete[] hx;
    delete[] hy;
    cudaFree(dx);
    cudaFree(dy);

    return 0;
}

首先要注意，這段代碼文件的后綴名是.cu锡足。cuda有兩種文件后綴名波丰，.cu和.cuh，相當于c++中.cpp和.h的關系舶得。

程序要做的就是兩個向量x掰烟、y的相加：y += x。向量長度為2^20沐批。下面我將逐行分析這段代碼：

1.引入c++頭文件

#include <cstdio>
#include <algorithm>

2.引入cuda頭文件

#include <cuda_runtime.h>
#include <device_launch_parameters.h>

3.向量相加核函數(shù)（kernel）

__global__ void add(float *dx, float *dy)
{
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    dy[idx] += dx[idx];
}

核函數(shù)是運行在gpu上的函數(shù)纫骑，能在很多線程上并行運行。

n是向量的長度九孩，dx和dy分別是指向向量的指針先馆。blockIdx和threadIdx分別是當前block和thread的下標，gridDim是block的數(shù)量捻撑，blockDim是一個block里thread的數(shù)量磨隘。blockIdx、threadIdx顾患、gridDim番捂、blockDim的關系如下圖所示（gridDim沒有畫出，grid和block的關系其實就相當于block和thread的關系）：

block/thread關系.png

id定位到具體的thread下標江解，正好也等于數(shù)組的下標（每個線程對應一個數(shù)組的元素）设预。

4.main函數(shù)入口

int main()

5.初始化向量數(shù)據(jù)

int N = 1 << 20;

float *hx, *hy, *dx, *dy;

hx = new float[N];
hy = new float[N];
cudaMalloc(&dx, N * sizeof(float));
cudaMalloc(&dy, N * sizeof(float));

for (int i = 0; i < N; i++)
{
    hx[i] = 1.0f;
    hy[i] = 2.0f;
}

cudaMemcpy(dx, hx, N * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(dy, hy, N * sizeof(float), cudaMemcpyHostToDevice);

N為向量長度。hx犁河、hy為host鳖枕，即cpu端的數(shù)據(jù)魄梯；dx、dy為device宾符，即gpu端的數(shù)據(jù)酿秸。host數(shù)據(jù)使用new分配，device數(shù)據(jù)使用cudaMalloc分配魏烫。device數(shù)據(jù)分配完畢之后辣苏，需要使用cudaMemcpy將host數(shù)據(jù)拷貝到device數(shù)據(jù)上。

int threadNums = 256;
int blockNums = (N + threadNums - 1) / threadNums;

add << <blockNums, threadNums >> >(N, dx, dy);

threadNums指的是一個block中有多少thread哄褒，即上面add函數(shù)中的blockDim稀蟋；blockNums指的是block的數(shù)量，對應gridDim呐赡。

add << <blockNums, threadNums >> >(N, dx, dy);一句調(diào)用了add核函數(shù)退客，通過<<<>>>符號傳入核函數(shù)運行配置參數(shù)。核函數(shù)的運行配置參數(shù)的完整形式為<<<gridDim, blockDim, Ns, S>>>链嘀，其中Ns和S一般使用默認值萌狂。

gridDim和blockDim都是Dim類型，可以是Dim1怀泊、Dim2和Dim3粥脚，即可以是一維、二維和三維包个。gridDim每個維度最大可以是65535，blockDim的值最大可以是1024冤留。所以理論上我們可以配置的最大線程數(shù)量為6553565535655351024 = 2.910^{17個碧囊。需要注意的是，這里的線程指的是軟件的線程纤怒，硬件上真正能并行的線程數(shù)要遠遠小于這個天文數(shù)字糯而，一般來說只有10}4的數(shù)量級。

6.同步gpu線程

cudaDeviceSynchronize();

其實沒有必要泊窘，詳見stackoverflow的解釋熄驼。

7.計算并打印出最大的誤差值

float maxError = 0.0f;
for (int i = 0; i < N; i++)
    maxError = std::max(maxError, std::abs(hy[i] - 3.0f));

printf("Max error: %.6f", maxError);

8.釋放資源

delete[] hx;
delete[] hy;
cudaFree(dx);
cudaFree(dy);

最后打印的結(jié)果當然不出所料，是0烘豹。

我們可以寫一段對應的c++程序：

#include <cstdio>
#include <algorithm>

int main()
{
    int N = 1 << 20;
    float *x, *y;
    x = new float[N];
    y = new float[N];

    for (int i = 0; i < N; i++)
    {
        x[i] = 1.0f;
        y[i] = 2.0f;
    }

    for (int i = 0; i < N; i++)
    {
        y[i] += x[i];
    }

    float maxError = 0.0f;
    for (int i = 0; i < N; i++)
        maxError = std::max(maxError, std::abs(y[i] - 3.0f));

    printf("Max error: %.6f", maxError);

    delete[] x;
    delete[] y;

    return 0;
}

最后可以做一個簡單的profile（使用c++11的chrono計時器）：

**cuda：

std::chrono::time_point<std::chrono::system_clock> begin = std::chrono::system_clock::now();

add << <blockNums, threadNums >> > (N, dx, dy);

std::chrono::time_point<std::chrono::system_clock> end = std::chrono::system_clock::now();
std::chrono::duration<double> elapsedTime = end - begin;
printf("Time: %.6lfs\n", elapsedTime.count());

結(jié)果：

cuda運行結(jié)果.png

c++：

std::chrono::time_point<std::chrono::system_clock> begin = std::chrono::system_clock::now();

for (int i = 0; i < N; i++)
{
    y[i] += x[i];
}

std::chrono::time_point<std::chrono::system_clock> end = std::chrono::system_clock::now();
std::chrono::duration<double> elapsedTime = end - begin;
printf("Time: %.6lfs\n", elapsedTime.count());

結(jié)果：

c++運行結(jié)果.png

feel the power of cuda!