2. CUDA编程模型：理解Host与Device、Kernel函数、线程层次结构

好，咱们直接进入正题。CUDA编程模型，说白了就是搞清楚「谁在干活」和「活怎么分」。

我刚接触CUDA那会儿，最困惑的就是CPU和GPU到底怎么配合。后来踩了不少坑才明白——CPU是老板，GPU是工人。老板负责发号施令、准备材料，工人只管埋头苦干。

2.1 Host与Device：老板和工人的分工

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

我习惯把Host比作「项目经理」，Device比作「施工队」。项目经理不亲自搬砖，但他得把图纸（代码）、材料（数据）准备好，然后喊一声「开工！」——这就是Kernel启动。

嗯，这里要注意：Host和Device有各自独立的内存空间。你不能在GPU上直接访问CPU的内存，反过来也不行。我在项目中遇到过新手直接传指针给Kernel，结果程序崩溃——因为那个指针指向的是系统内存，GPU根本访问不到。

正确的做法是三步走：

1. cudaMalloc 在Device上分配显存
2. cudaMemcpy 把数据从Host拷到Device
3. 启动Kernel，算完再拷回来

// 典型的数据搬运流程
float *d_data;
cudaMalloc(&d_data, N * sizeof(float));          // 1. 在GPU上分配
cudaMemcpy(d_data, h_data, N * sizeof(float), 
           cudaMemcpyHostToDevice);              // 2. 从CPU拷到GPU
myKernel<<<grid, block>>>(d_data);            // 3. 启动Kernel
cudaMemcpy(h_result, d_data, N * sizeof(float),
           cudaMemcpyDeviceToHost);              // 4. 结果拷回CPU
cudaFree(d_data);                                // 5. 释放显存

2.2 Kernel函数：GPU上执行的代码

Kernel函数就是跑在GPU上的函数。你想想看，CPU上跑的是普通函数，GPU上跑的就是Kernel。

定义Kernel时，要用 __global__ 修饰符。调用时用尖括号语法 <<<grid, block>>> 指定线程组织方式。

// 一个简单的向量加法Kernel
__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];
    }
}

// 调用方式
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

我个人习惯在Kernel里总是加边界检查——就是那个 if (idx < N)。为什么？因为线程总数可能超过数据量，多出来的线程如果不检查，就会访问越界内存。我曾经调试过一个图像处理程序，就是少了这行检查，结果图像边缘出现奇怪的条纹，查了两天才找到原因。

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

这是CUDA最核心的概念，也是初学者最容易晕的地方。我尽量用大白话讲清楚。

想象一下：你要给一栋大楼的每个房间刷墙。

• Thread（线程） = 一个工人，负责刷一个房间
• Block（线程块） = 一个施工队，负责刷一层楼
• Grid（网格） = 整个项目，负责整栋楼

每个Thread执行相同的Kernel代码，但处理不同的数据。它们通过内置变量来区分自己是谁：

内置变量	含义	取值范围
threadIdx.x	线程在Block内的索引	0 到 blockDim.x - 1
blockIdx.x	Block在Grid内的索引	0 到 gridDim.x - 1
blockDim.x	Block内线程的数量	由 <<<grid, block>>> 指定
gridDim.x	Grid内Block的数量	由 <<<grid, block>>> 指定

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

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

嗯，这里要注意：Block内的线程可以协作（通过共享内存和同步），但不同Block之间的线程是独立的，不能直接通信。我在项目中遇到过有人试图用全局内存做跨Block的同步，结果死锁了——因为不同Block的执行顺序是不确定的。

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

CUDA支持一维、二维、三维的线程组织。说白了，就是你的数据是什么形状，线程就排成什么形状。

• 一维：处理数组、向量
• 二维：处理图像、矩阵
• 三维：处理体数据、3D网格

// 二维Kernel示例：处理一张图片
__global__ void blurImage(unsigned char *in, unsigned char *out, 
                          int width, int height) {
    int col = threadIdx.x + blockIdx.x * blockDim.x;  // 列
    int row = threadIdx.y + blockIdx.y * blockDim.y;  // 行
    
    if (col < width && row < height) {
        int idx = row * width + col;
        // 执行模糊操作...
        out[idx] = in[idx];
    }
}

// 调用方式
dim3 blockDim(16, 16);  // 每个Block 16x16=256个线程
dim3 gridDim((width + 15) / 16, (height + 15) / 16);
blurImage<<<gridDim, blockDim>>>(d_in, d_out, width, height);

我曾经在做一个3D地震数据处理时，把Block的z维度设成了128，结果Kernel启动失败。查了半天文档才发现z维度最大只有64。从那以后，我每次写Kernel前都会先查一下目标设备的计算能力。

2.5 避坑指南：我踩过的那些坑

好了，这一章的内容就到这里。记住：Host是大脑，Device是肌肉；Kernel是动作，线程是细胞。理解了这个层次关系，后面的优化技巧才能派上用场。