3. CUDA核函数编写基础：从修饰符到线程索引
好，咱们今天聊聊CUDA核函数的基础。说实话，这是整个CUDA编程的基石。你想想看，没有核函数，GPU再多的计算单元也只是摆设。我刚开始学CUDA的时候，就是被这些修饰符和启动配置搞得有点晕，后来在项目中踩过几次坑，才算真正搞明白。
3.1 修饰符三兄弟：__global__、__device__、__host__
这三个修饰符，说白了就是告诉编译器：这段代码跑在哪儿、谁来调用它。
- __global__：这是核函数的标志。它修饰的函数跑在GPU上，但由CPU调用。嗯，你可以把它理解成「GPU的入口函数」。
- __device__：这个修饰的函数也是跑在GPU上，但它只能被其他GPU函数调用（比如从__global__函数里调用）。它不能直接被CPU调用。
- __host__：这是默认的，就是普通的CPU函数。其实你平时写的函数不加修饰符，默认就是__host__。
关键点：__global__函数必须返回void，而且不能是类的成员函数（除非是静态成员）。这个限制我一开始没注意，编译报错找了半天原因。
举个例子你就明白了：
// 核函数：由CPU调用，在GPU上执行
__global__ void vecAdd(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];
}
}
// 设备函数：只能被GPU上的其他函数调用
__device__ float square(float x) {
return x * x;
}
// 主机函数：普通的CPU函数
__host__ void initData(float* arr, int N) {
for (int i = 0; i < N; i++) {
arr[i] = (float)rand() / RAND_MAX;
}
}我的习惯：我一般把__device__函数放在核函数前面，这样阅读代码时能先看到工具函数，再看到主逻辑。当然，这只是个人偏好，没有硬性规定。
3.2 核函数启动配置：<<>>
这个语法看起来有点奇怪，但用多了就习惯了。它的格式是：
kernel_name<<<gridDim, blockDim, sharedMem, stream>>>(args...);其中前两个参数是必须的，后两个可选。咱们先聚焦前两个：
- gridDim：网格的维度，也就是有多少个线程块。类型是dim3，可以是一维、二维或三维。
- blockDim：线程块的维度，也就是每个块里有多少个线程。也是dim3类型。
为什么会设计成网格和块两层结构？说白了，这是为了兼顾灵活性和硬件限制。每个线程块里的线程可以共享内存、同步协作，但块的数量可以很多，覆盖大规模数据。
我曾经踩过的坑：有一次我处理一个100万元素的数组，直接用了100万个线程块，每个块1个线程。结果性能惨不忍睹。后来才意识到，线程块太小会导致硬件利用率极低。一般建议每个块至少128或256个线程。
来看一个典型的启动配置：
int N = 1 << 20; // 约100万
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);这里用了一个经典技巧：(N + threadsPerBlock - 1) / threadsPerBlock 向上取整，确保覆盖所有元素。然后核函数内部用 if (idx < N) 做边界检查。
3.3 线程索引计算：找到你的位置
每个线程在GPU上执行时，都需要知道自己处理的是哪部分数据。这就需要计算线程的全局索引。
先记住几个内置变量：
| 变量 | 含义 | 类型 |
|---|---|---|
| threadIdx | 线程在块内的索引 | dim3 |
| blockIdx | 块在网格内的索引 | dim3 |
| blockDim | 块的维度（每块线程数） | dim3 |
| gridDim | 网格的维度（每网格块数） | dim3 |
一维情况最简单：
int idx = threadIdx.x + blockIdx.x * blockDim.x;二维情况稍微复杂一点，但逻辑是一样的：
int col = threadIdx.x + blockIdx.x * blockDim.x;
int row = threadIdx.y + blockIdx.y * blockDim.y;
int idx = row * width + col; // 展平为一维索引三维情况以此类推。我个人习惯在处理图像或矩阵时用二维网格，处理向量时用一维网格。这样代码更直观。
核心公式：全局线程ID = threadIdx + blockIdx * blockDim。记住这个，90%的核函数都靠它。
3.4 知识体系总览
下面这张图帮你理清本章的核心逻辑：
3.5 实战中的避坑指南
嗯，这里分享几个我实际项目中遇到的坑：
- 线程块大小不是越大越好：每个线程块最多1024个线程（取决于GPU架构），但并不是越大越好。我试过用1024，结果寄存器压力太大，反而变慢了。一般256或512是比较稳妥的选择。
- 别忘了边界检查：当数据量不是线程块大小的整数倍时，多余的线程会访问越界内存。我刚开始写的时候经常忘加if判断，结果数据对不上，排查了半天。
- dim3的默认值：如果你只给一个维度赋值，比如dim3 block(256)，那么block.y和block.z默认是1。这个特性用好了可以简化代码。
我的小技巧：调试核函数时，我习惯先用很小的数据量（比如N=16）和单线程块（<<<1, N>>>）跑一遍，确认逻辑正确后再放大规模。这样能快速定位问题。
好了，这一章的内容就这些。修饰符、启动配置、线程索引，这三个东西搞明白了，你就掌握了CUDA核函数的基本功。下一章咱们会深入聊聊内存模型，那才是真正影响性能的关键。
想要更多资料，联系可可