3、核函数（Kernel）基础：核函数定义与声明、<<<>>>执行配置语法、核函数的同步与异步执行

好，咱们正式开始聊核函数。这是CUDA编程里最核心的东西，没有之一。你想想看，GPU上跑的所有并行计算，最终都要落到核函数上。我刚开始学CUDA那会儿，觉得这玩意儿不就是个加了修饰符的函数嘛，后来踩的坑多了才明白——这里面的门道，真不少。

3.1 核函数定义与声明

核函数，说白了就是在GPU上执行的函数。它跟普通的C++函数长得有点像，但有几个关键区别。

先看一个最简单的例子：

__global__ void vectorAdd(float *A, float *B, float *C, int N) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

注意那个 __global__ 修饰符。这是CUDA C的扩展关键字，告诉编译器：这个函数要在GPU上跑，而且可以从CPU端调用。我个人习惯把 __global__ 放在返回类型前面，这样一眼就能看出来这是个核函数。

核函数有几个硬性规定，我列一下：

返回类型必须是 void。嗯，不能返回任何值。
参数必须能通过寄存器或共享内存传递。结构体可以传，但别太大。
不能使用可变参数（像 printf 那种不行）。
不能是类的成员函数（静态成员函数可以）。
不能使用函数指针。

我在项目中遇到过有人试图在核函数里用 std::vector，结果编译报错一脸懵。记住，核函数里只能用CUDA内置的类型和简单的C/C++结构。

重要：核函数里声明的局部变量，默认存在每个线程的寄存器或本地内存中。每个线程都有自己的副本，互不干扰。这就是GPU并行计算的基础。

3.2 <<<>>> 执行配置语法

调用核函数的时候，要用到那个看起来很奇怪的尖括号语法：

vectorAdd<<<gridDim, blockDim, sharedMem, stream>>>(d_A, d_B, d_C, N);

这四个参数，我来拆开讲：

参数	含义	类型	默认值
gridDim	网格维度，即有多少个线程块	dim3	无默认值，必须指定
blockDim	线程块维度，即每个块有多少个线程	dim3	无默认值，必须指定
sharedMem	每个块动态分配的共享内存字节数	size_t	0
stream	CUDA流，用于异步执行	cudaStream_t	0（默认流）

举个例子，我要启动1024个线程块，每个块256个线程：

int blockSize = 256;
int gridSize = (N + blockSize - 1) / blockSize;
vectorAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);

这里有个小技巧。我习惯用 (N + blockSize - 1) / blockSize 这个公式来计算网格大小，它能保证覆盖所有数据元素，不会漏掉尾巴。你想想看，如果N=1000，blockSize=256，那gridSize就是4，最后一个块只有232个线程干活，但至少不会漏数据。

经验之谈：blockSize一般选32的倍数，最好是128或256。为什么？因为GPU的warp大小是32，选32的倍数能避免warp内线程闲置。我试过用37这种奇怪的数，性能直接掉了一半。

3.3 核函数的同步与异步执行

核函数调用是异步的。什么意思？就是CPU调用核函数后，不会等它执行完，而是立刻返回继续执行后面的代码。这跟普通函数调用完全不一样。

看这段代码：

vectorAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, N);
// CPU不会在这里等待，而是直接往下执行
printf("核函数已经启动\n");
cudaDeviceSynchronize();  // 手动等待GPU完成
printf("核函数执行完毕\n");

如果不加 cudaDeviceSynchronize()，那 "核函数执行完毕" 可能会在GPU还没算完的时候就打印出来。我曾经在调试一个bug时，就是因为忘了同步，结果CPU端拷贝数据时GPU还在算，拷出来的全是垃圾数据。

同步方式主要有三种：

cudaDeviceSynchronize()：阻塞CPU，直到所有GPU操作完成。最常用，但别滥用。
cudaStreamSynchronize(stream)：只等待指定流里的操作完成。
cudaEventSynchronize(event)：等待某个事件触发。

注意：频繁调用 cudaDeviceSynchronize() 会严重拖慢性能。因为它强制CPU等待GPU，破坏了异步执行的优势。我建议只在需要拷贝数据回CPU，或者调试时使用。

异步执行的好处是什么？CPU可以在GPU干活的同时，做自己的事情。比如准备下一批数据，或者处理之前的结果。这就是所谓的「CPU-GPU并发执行」。

我画了一张图，帮你理解核函数的执行流程：

从图上你能看到，CPU调用核函数后，GPU开始干活，CPU也没闲着。直到CPU调了 cudaDeviceSynchronize()，才停下来等GPU。这种设计让CPU和GPU能同时工作，效率高很多。

不过要注意，核函数内部的线程之间，默认是不同步的。如果你想在核函数里让所有线程同步，得用 __syncthreads()。这个我们后面会细讲。

避坑指南：我曾经在核函数里用了一个全局变量做累加，结果发现每次跑出来的结果都不一样。为什么？因为多个线程同时写同一个全局变量，产生了数据竞争。记住，核函数里多个线程访问同一块数据时，一定要考虑同步问题。

最后总结一下核函数的基础要点：

用 __global__ 声明，返回类型必须是 void
用 <<<grid, block>>> 指定执行配置
调用是异步的，需要手动同步
每个线程有自己的局部变量，共享数据需要同步

嗯，核函数的基础就这些。看起来简单，但真正用好它，需要理解GPU的硬件架构和执行模型。下一节我们会深入线程层次结构，到时候你就知道为什么blockSize要选32的倍数了。

📚 课程目录

3、核函数（Kernel）基础：核函数定义与声明、<<<>>>执行配置语法、核函数的同步与异步执行

3.1 核函数定义与声明

3.2 <<<>>> 执行配置语法

3.3 核函数的同步与异步执行