CUDA核函数

启动一个核函数

一个CUDA函数的调用

kernel_name <<<grid, block>>>(argument list);

<<<>>>中间的部分，是核函数的运行配置。执行配置的第一个值是网格维度，也就是启动块的数目；第二个值是块维度，也就是每个块中线程的数目。

同一个块中的线程可以相互协作，不同块内的线程不能协作。

由于数据在全局内存中是线性存储的，因此可以用变量blockIdx.x和threadIdx.x来进行以下操作。

• 在网格中标识一个唯一的线程
• 建立线程和数据元素之间的映射关系

核函数的调用与主机线程是异步的。核函数调用结束后，控制权立刻返回给主机端。但是可以通过调用以下函数来强行令主机端等待所有的核函数执行结束：

cudaError_t cudaDeviceSynchronize(void);

举个栗子

我们在调用核函数之后，控制权立刻返回给主机端，所以先打印了hello cpu，然后才打印了GPU的核函数的输出。

强制让主机端等待后，先输出了核函数的打印，才打印了主机端的CPU。

另外一些CUDA运行时API在主机端与设别端是隐式同步的，比如cudaMemcpy函数，在之前所有的核函数调用完成后，开始拷贝数据。直到数据拷贝完成后，将控制权返回给主机端。

编写核函数

核函数的定义

__global__ void kerel_name(argument list){
    //...
}

核函数由__global__声明，返回类型必须是void

下面说明CUDA中函数类型限定符

__global__	在设备端执行	可以从主机端调用 也可以从计算能力为3的设备中调用
__device__	在设备端执行	仅能从设备端调用
__host__	在主机端执行	仅能从主机端调用（可以省略）

device限定符和host限定符可以一起使用，这样函数可以同时在主机端和设备端进行编译。

CUDA核函数的限制

• 只能访问设备内存
• 必选返回void类型
• 不支持可变数量的参数
• 不支持静态变量
• 显示异步行为

举个栗子

将两个大小为N的向量A,B相加，如果在主机端完成，代码如下

void sumArray(int* A, int* B, int* C, const int N){
    for(int i=0; i<N; i++)
        c[i] = A[i] + B[i];
}

我们对比一下在GPU核函数上进行相同的操作

__global__ void addArrayOnGPU(const int* a, const int* b, int* c) {
    int i = threadIdx.x;
    c[i] = a[i] + b[i];
}

int main()
{
    using namespace std;
    
    int a[] = { 1,2,3,4 };
    int b[] = { 1,2,3,4 };
    int c[4];

    int* d_a, * d_b, * d_c;
    cudaMalloc((int**)&d_a, sizeof(a));
    cudaMalloc((int**)&d_b, sizeof(b));
    cudaMalloc((int**)&d_c, sizeof(c));

    cudaMemcpy(d_a, a, sizeof(a), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, sizeof(b), cudaMemcpyHostToDevice);
    addArrayOnGPU << <1, 4 >> > (d_a, d_b, d_c);
    cudaMemcpy(c, d_c, sizeof(a), cudaMemcpyDeviceToHost);
    
    cudaDeviceReset();

    for (int i = 0; i < 4; i++)
        cout << c[i] << " ";

    return 0;
}

我们注意到，循环体消失了，用内置的线程坐标变量替换了数组索引。因为我们在主机端调用核函数时，启动了4个线程，每个线程都从GPU全局内存中读取数据并且进行了修改。