3、CUDA编程模型：主机与设备、Kernel函数概念、线程层次结构

好，咱们今天聊聊CUDA的编程模型。说实话，这是整个CUDA学习的基石。你如果把这个搞明白了，后面写Kernel就跟写串行代码一样顺手。

我第一次接触CUDA时，最困惑的就是：这代码到底在CPU上跑，还是在GPU上跑？ 后来才明白，CUDA程序其实是“两拨人”在干活——CPU和GPU各司其职。

3.1 主机与设备：CPU和GPU的分工

在CUDA的世界里，主机（Host） 指的是CPU及其内存系统，设备（Device） 指的是GPU及其显存。说白了，CPU负责逻辑控制、串行任务，GPU负责大规模并行计算。

我习惯把CPU比作“项目经理”，GPU比作“施工队”。项目经理定方案、分任务，施工队闷头干活。项目经理一次只能处理一两件事，但施工队可以几百人同时砌墙。

你想想看，如果让CPU去算几百万个像素点，它得累死。但GPU天生就是干这个的——它有成百上千个核心，每个核心虽然跑得慢，但架不住人多啊。

3.2 Kernel函数：在GPU上跑的函数

Kernel函数，就是你在GPU上执行的函数。它的定义方式很简单——在函数前面加个 __global__ 修饰符。

// 一个简单的Kernel函数
__global__ void vecAdd(float* A, float* B, float* C, int n) {
    int i = threadIdx.x;  // 获取当前线程的ID
    if (i < n) {
        C[i] = A[i] + B[i];
    }
}

注意几个关键点：

• __global__ 表示这个函数在设备上执行，但从主机端调用
• Kernel函数返回类型必须是 void
• Kernel函数内部不能调用标准C++的 printf（调试时可以用 printf，但正式代码别用）

嗯，这里要注意：Kernel函数是异步的。你调用它之后，CPU不会等它执行完就继续往下跑。如果你需要等结果，得手动加 cudaDeviceSynchronize()。

3.3 线程层次结构：Grid、Block、Thread

这是CUDA最核心的概念。你想想看，GPU有几千个核心，你怎么组织这些线程？CUDA给出了一个三层结构：

• Thread（线程）：最小的执行单元，每个线程执行Kernel函数的一份拷贝
• Block（线程块）：一组线程的集合，它们可以共享内存、同步执行
• Grid（网格）：一组Block的集合，一个Kernel对应一个Grid

我画了个图，帮你理解这个层次关系：

每个线程都有一个唯一的ID，通过 threadIdx、blockIdx、blockDim 和 gridDim 来定位。比如：

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

这个公式我建议你背下来。它把三维的线程索引展平成一维，方便你访问数组。

3.4 <<<>>> 语法：启动Kernel的“开关”

启动Kernel的语法很特别，用三括号 <<>> 来指定线程配置：

// 启动Kernel：1个Grid，每个Grid有1个Block，每个Block有256个线程
vecAdd<<<1, 256>>>(d_A, d_B, d_C, N);

第一个参数是Grid的维度（多少个Block），第二个参数是Block的维度（多少个Thread）。你还可以传第三个参数（共享内存大小）和第四个参数（流）。

我个人习惯这样计算Grid大小：

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

这个 (N + blockSize - 1) / blockSize 是向上取整的技巧，保证覆盖所有数据。你想想看，如果N=1000，blockSize=256，那么gridSize=4，最后一个Block只有232个线程干活，剩下的24个线程闲着。这没问题，只要在Kernel里加个边界判断就行。

3.5 完整的CUDA程序示例

把上面这些串起来，就是一个完整的CUDA程序：

#include <cuda_runtime.h>
#include <stdio.h>

// Kernel函数：向量加法
__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;
    size_t size = n * sizeof(float);

    // 1. 分配主机内存
    float *h_A, *h_B, *h_C;
    h_A = (float*)malloc(size);
    h_B = (float*)malloc(size);
    h_C = (float*)malloc(size);

    // 初始化数据
    for (int i = 0; i < n; i++) {
        h_A[i] = i * 1.0f;
        h_B[i] = i * 2.0f;
    }

    // 2. 分配设备内存
    float *d_A, *d_B, *d_C;
    cudaMalloc(&d_A, size);
    cudaMalloc(&d_B, size);
    cudaMalloc(&d_C, size);

    // 3. 拷贝数据到设备
    cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
    cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

    // 4. 配置并启动Kernel
    int blockSize = 256;
    int gridSize = (n + blockSize - 1) / blockSize;
    vecAdd<<<gridSize, blockSize>>>(d_A, d_B, d_C, n);

    // 5. 拷贝结果回主机
    cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

    // 6. 验证结果
    for (int i = 0; i < 10; i++) {
        printf("C[%d] = %f\n", i, h_C[i]);
    }

    // 7. 释放内存
    cudaFree(d_A);
    cudaFree(d_B);
    cudaFree(d_C);
    free(h_A);
    free(h_B);
    free(h_C);

    return 0;
}

这个流程你记住：分配内存 → 拷贝数据 → 启动Kernel → 拷贝结果 → 释放内存。少了哪一步都不行。

好了，这一章的内容就这些。主机与设备的分工、Kernel函数的写法、线程层次结构、<<<>>> 语法，这四个点你吃透了，CUDA编程的大门就算打开了。