1. 启动CUDA内核

1.1启动CUDA内核函数格式

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

• kernel在device上执行时实际上是启动很多线程，一个kernel所启动的所有线程称为一个网格（grid)。
• 同一block内的线程可以轻松地相互通信，属于不同block的线程无法协作。
• 对于给定的问题，可以使用不同的grid和block布局来组织线程。例如，假设有32个数据元素用于计算。可以将8个元素分组到每个块中，并按以下方式启动4个块：

kernel_name<<<4, 8>>>(argument list);

thread编号.png

1.2 CUDA内核启动

• 与C函数调用不同，所有CUDA内核启动都是异步的。 CUDA内核被调用后，控制权立即返回给CPU。
• 可以调用以下函数来强制主机应用程序等待所有内核完成。

cudaError_t cudaDeviceSynchronize(void);

• 一些CUDA运行时API在主机和设备之间执行隐式同步。使用cudaMemcpy在主机和设备之间复制数据时，将在主机端执行隐式同步，并且主机应用程序必须等待数据复制完成。在所有先前的内核调用完成后，它将开始复制。复制完成后，控制权立即返回主机。

cudaError_t cudaMemcpy(void* dst, const void* src, size_t count, cudaMemcpyKind kind);

2. 写自己的GPU kernel

关键词 global会告诉编译器这个函数在CPU上调用，在GPU上执行。

__global__ void helloFromGPU(void) {
    printf(“Hello World from GPU!\n”); 
}

• kernel的返回类型必须是void.

3. 示例

• 矩阵A和矩阵B的简单相加。
  CPU代码

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

GPU代码：

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

内置的thread坐标变量用来替换array坐标i，使用启动GPU kernel来启动N个threads。

sumArraysOnGPU<<<1,32>>>(float *A, float *B, float *C);

4. 验证Kernel正确性

• 在Fermi及后面的版本中，可以使用printf函数打印
• 设置执行参数为 <<<1,1>>>, 强制 kernel运行一个block和一个thread