4、CUDA编程模型入门：主机与设备、核函数(Kernel)概念、线程层次结构

好，咱们正式开始接触CUDA编程的核心了。

说实话，我当年刚接触CUDA时，最困惑的就是这个“主机”和“设备”到底怎么分工。你想想看，一个程序里有两套内存，两套执行流程，稍不留神就搞混了。这一章，咱们就把这些基础概念彻底捋清楚。

4.1 主机(Host)与设备(Device)：CPU和GPU的分工

在CUDA的世界里，主机指的是CPU及其内存（系统内存），设备指的是GPU及其显存（设备内存）。

说白了，CPU是“总指挥”，负责逻辑控制、串行任务、数据预处理。GPU是“打工人”，专门干那些可以并行计算的脏活累活。

我习惯把CPU比作一个项目经理，GPU就是手底下的一千个工人。项目经理把任务拆好、分下去，工人们同时开干，最后项目经理再把结果收回来。

我在项目中遇到过一个问题：有个同事把数据指针直接传给了GPU，以为GPU能直接访问CPU内存。结果程序跑起来要么报错，要么出乱数据。嗯，这里要注意——GPU不能直接访问主机内存（除非你用统一内存，那是后话了）。

4.2 核函数(Kernel)：在GPU上运行的函数

核函数，就是你在GPU上执行的函数。写法上有个明显的标志：__global__ 修饰符。

// 一个简单的核函数：给数组每个元素加1
__global__ void add_one(float *d_data, int n) {
    int idx = threadIdx.x + blockIdx.x * blockDim.x;
    if (idx < n) {
        d_data[idx] += 1.0f;
    }
}

你看，这个函数前面加了 __global__，表示它是在设备上执行、但从主机端调用的。调用方式也很特别：

// 主机端调用核函数
int n = 1024;
float *d_data;
cudaMalloc(&d_data, n * sizeof(float));
// ... 拷贝数据到设备 ...

// 启动核函数：<<>>
add_one<<<1, 256>>>(d_data, n);

cudaDeviceSynchronize(); // 等待GPU完成

这里 <<<1, 256>>> 就是启动配置，告诉GPU：用1个Block，每个Block里256个线程。总共256个线程并行执行这个函数。

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

这是CUDA最核心的概念，也是初学者最容易晕的地方。我当年花了整整一周才彻底搞明白。

CUDA的线程组织分三层：

• Thread（线程）：最小的执行单元，每个线程执行相同的核函数，但处理不同的数据
• Block（线程块）：一组线程的集合，Block内的线程可以共享内存、同步协作
• Grid（网格）：一组Block的集合，一个Grid对应一个核函数的启动

你可以想象成：Grid是一个工厂，Block是车间，Thread是工人。每个工人干一样的活，但负责不同的零件。

计算全局唯一线程ID的公式：

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

举个例子：假设你启动 <<<4, 256>>>，那么总共有 4×256 = 1024 个线程。第0个Block里的线程ID是0~255，第1个Block里的线程ID是256~511，以此类推。

4.4 一维、二维、三维：灵活的数据映射

CUDA支持一维、二维、三维的线程组织。为什么需要多维？因为你处理的数据可能是图像（二维）、体素（三维）。

二维Block的例子：

// 处理一张 32x32 的图像
dim3 blockDim(16, 16);  // 每个Block 16x16 = 256个线程
dim3 gridDim(2, 2);     // 2x2 = 4个Block

kernel<<<gridDim, blockDim>>>(d_image);

// 核函数内部
__global__ void kernel(float *d_image) {
    int x = threadIdx.x + blockIdx.x * blockDim.x;
    int y = threadIdx.y + blockIdx.y * blockDim.y;
    int idx = y * width + x;  // 二维转一维
    // 处理像素...
}

我个人习惯：处理图像时用二维Block，处理一维数组时用一维Block。别为了炫技硬上三维，除非你真的需要。

4.5 一个完整的入门示例

咱们写一个完整的向量加法，把上面所有概念串起来：

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

// 核函数：向量加法
__global__ void vec_add(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];
    }
}

int main() {
    int n = 1024;
    size_t bytes = n * sizeof(float);

    // 1. 分配主机内存
    float *h_a = (float*)malloc(bytes);
    float *h_b = (float*)malloc(bytes);
    float *h_c = (float*)malloc(bytes);

    // 初始化数据...
    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, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);

    // 3. 拷贝数据到设备
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);

    // 4. 启动核函数
    int threads_per_block = 256;
    int blocks_per_grid = (n + threads_per_block - 1) / threads_per_block;
    vec_add<<<blocks_per_grid, threads_per_block>>>(d_a, d_b, d_c, n);

    // 5. 等待完成，拷贝结果回主机
    cudaDeviceSynchronize();
    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);

    // 6. 验证结果
    for (int i = 0; i < 10; i++) {
        printf("h_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;
}

这个例子虽然简单，但包含了CUDA编程的完整流程：分配内存 → 拷贝数据 → 启动核函数 → 同步 → 拷回结果 → 清理。每一步都不能少。

4.6 本章小结

这一章咱们把CUDA编程模型的三块基石讲清楚了：

• 主机与设备：CPU负责调度，GPU负责计算，数据要显式搬运
• 核函数：用 __global__ 声明，用 <<<grid, block>>> 启动
• 线程层次：Thread → Block → Grid，用 threadIdx、blockIdx、blockDim 定位

下一章，咱们会深入内存管理，聊聊全局内存、共享内存、常量内存这些“坑”和“宝”。到时候我会分享一个我在Jetson上优化图像处理时踩过的内存带宽的坑，保证让你少走弯路。