3. CUDA 编程模型：核函数、<<<>>>语法与线程索引计算

好，我们直接进入正题。

CUDA 编程模型，说白了就是一套规则。它规定了你怎么让 CPU 去指挥 GPU 干活。CPU 是主机（Host），GPU 是设备（Device）。你写的普通 C 代码在 CPU 上跑，但想榨干 GPU 的性能，就得靠一个叫“核函数”（Kernel）的东西。

3.1 核函数：GPU 的入口点

核函数是什么？
它是你写的一段 C 函数，但前面加了个 __global__ 声明。这告诉编译器：这个函数要跑在 GPU 上，由 CPU 调用。

我个人习惯，写核函数时注意三点：

• 返回类型必须是 void。没办法，CUDA 就这规矩。
• 不能有可变参数（比如 ...）。
• 里面可以调 __device__ 函数，但不能调 __host__ 函数。

3.2 <<<>>> 语法：启动配置

调用核函数时，你会看到一种特殊的语法：

kernel_name<<<gridDim, blockDim, sharedMem, stream>>>(args);

这里最关键的是前两个参数：

举个例子。你想要 512 个线程，可以这样：

int N = 512;
myKernel<<<1, N>>>(d_a, d_b, d_c);  // 网格 1 个块，块里 512 个线程

但如果你数据量超过 1024（每块最大线程数），就必须分成多个块：

int N = 2048;
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;  // 向上取整
myKernel<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c);

3.3 线程索引计算：找到自己的位置

每个线程都有一个唯一的 ID。怎么算？CUDA 提供了几个内置变量：

• threadIdx：当前线程在线程块里的索引（维度 x、y、z）。
• blockIdx：当前线程块在网格里的索引（维度 x、y、z）。
• blockDim：线程块的维度（每块多少个线程）。
• gridDim：网格的维度（每网格多少个块）。

你想想看，一个一维场景，要拿到全局线程 ID：

int tid = blockIdx.x * blockDim.x + threadIdx.x;

二维场景稍微复杂点：

int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int idx = y * width + x;  // 线性化

为什么需要这个？
因为你要操作的数据是扁平的数组。核函数里，每个线程必须知道它负责哪个数据元素。用这个公式，你可以把线程的 逻辑位置 映射到 物理内存地址。

全局线程ID = 块索引 * 块大小 + 线程在块内的索引

3.4 实战：向量加法

来个最简单的例子，两个向量相加：

__global__ void vecAdd(float *a, float *b, float *c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {  // 边界检查
        c[i] = a[i] + b[i];
    }
}

int main() {
    int n = 1000000;
    float *h_a = (float*)malloc(n * sizeof(float));
    float *h_b = (float*)malloc(n * sizeof(float));
    float *h_c = (float*)malloc(n * sizeof(float));

    // ... 初始化 h_a, h_b ...

    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, n * sizeof(float));
    cudaMalloc(&d_b, n * sizeof(float));
    cudaMalloc(&d_c, n * sizeof(float));

    cudaMemcpy(d_a, h_a, n * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, n * sizeof(float), cudaMemcpyHostToDevice);

    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, n);

    cudaMemcpy(h_c, d_c, n * sizeof(float), cudaMemcpyDeviceToHost);

    // ... 检查结果 ...

    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    free(h_a); free(h_b); free(h_c);
    return 0;
}

代码很简单，但有几个坑：

• 别忘了边界检查。因为 blocksPerGrid * threadsPerBlock 可能超过 n。
• 核函数里用 if 保护数组访问，防止越界。
• 我记得第一次写这个时，忘了 cudaMemcpy 回宿主，结果打印全是 0。嗯，这种错犯一次就够了。

3.5 知识体系：编程模型全景图

下面这张图，把整个 CUDA 线程层次结构和索引计算关系画清楚了：

通过这张图，你能看到完整的模型：

• 顶层的 Grid 由多个 Block 组成。
• 每个 Block 包含多个 Thread。
• 索引公式 blockIdx * blockDim + threadIdx 把所有线程映射到连续的数组索引上。

3.6 常见避坑：我踩过的坑

讲到这里，我想分享一下实战中容易犯的错误：

1. 忘记向上取整——数据量不是块大小的整数倍时，最后一波数据被忽略。我吃过一次亏，图像处理时边缘像素全没了。
2. 索引越界——多算出来的线程不处理，所以一定要加 if (i < n) 保护。
3. 忘了 cudaMemcpy——核函数算完了，结果还在 GPU 显存里，CPU 侧打印全是垃圾值。嗯，这个错我记忆犹新。