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
