2.1 核函数

内核函数是CUDA 每个线程 执行的函数。CUDA使用扩展的C语言编写内核函数，关键字为global。内核函数返回值只能是void。

下面是一段简单的内核函数，用于求两个数组的和：

// Kernel definition
__global__ void VecAdd(float* A, float* B, float* C)
{
    int i = threadIdx.x;
    C[i] = A[i] + B[i];
}

int main()
{
    ...
    // Kernel invocation with N threads
    VecAdd<<<1, N>>>(A, B, C);
    ...
}

threadIdx.x是线程在所处线程块中的X方向的ID。由于本例中是定义的1维排布，因此X方向ID即为线程的ID。

由于GPU中的每个线程都会执行相同的VecAdd函数，因此不同的线程需要使用自己独有的ID来区分彼此，来获取不同的数据。这就是SIMT的概念，即“相同指令，不同线程”。

在main()函数中，我们注意到，VecAdd函数的调用使用了<<<blockPerGrid, threadsPerBlock>>>关键字。这是调用内核函数所独有的。程序员通过该关键字，制定网格中线程块和线程的排布方式。排布方式与数据息息相关。

下面举一个2维排布的例子，用于做矩阵加法：

// Kernel definition
__global__ void MatAdd(float A[N][N], float B[N][N],
float C[N][N])
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    int j = blockIdx.y * blockDim.y + threadIdx.y;
    if (i < N && j < N)
        C[j][i] = A[j][i] + B[j][i];
}

int main()
{
    ...
    // Kernel invocation
    dim3 threadsPerBlock(16, 16);
    dim3 numBlocks(N / threadsPerBlock.x, N / threadsPerBlock.y);
    MatAdd<<<numBlocks, threadsPerBlock>>>(A, B, C);
...
}

首先看主函数，当排布不使用1维时，需要使用dim3数据类型。该程序每个线程块中线程为16x16排布，而线程块的排布依赖于数据的多少。

在内核函数中，i代表x方向上的ID，j代表y方向上的ID。blockDim代表当前线程块的尺寸。从程序中可以看到，x方向为行方向，y方向为列方向。（注意，这里官方文档里面写的有些错误）

每个线程读取自己ID对应的数据A[j][i]和B[j][i]，并将结果写回C[j][i]。其中A、B、C都存储在GPU的全局内存上（后面会提及）