第3章：GPU编程基础（CUDA）——从零开始写你的第一个Kernel

好，我们进入CUDA编程的正题了。说实话，很多初学者一上来就被线程层次、内存模型这些东西搞懵了。我当年也是，第一次看到Grid、Block、Thread这些概念时，脑子里全是问号。但别急，咱们一步步来拆解。

3.1 CUDA编程模型——CPU+GPU的协作模式

CUDA的编程模型，说白了就是「CPU负责调度，GPU负责干活」。CPU我们叫它Host，GPU叫Device。你想想看，CPU擅长逻辑控制，GPU擅长并行计算，各司其职。

典型的流程是这样的：

1. CPU把数据从内存拷贝到GPU显存
2. CPU告诉GPU：「嘿，启动这个Kernel，用1024个线程跑」
3. GPU执行完，把结果拷回CPU内存

我在项目中遇到过一个问题：有人把数据拷贝和Kernel执行混在一起写，结果性能惨不忍睹。记住，数据传输是瓶颈，尽量减少Host和Device之间的数据搬运。

3.2 线程层次结构——Grid/Block/Thread

这是CUDA最核心的概念。我习惯用一个比喻来理解：

• Thread（线程）：最基层的工人，每个工人干一小块活
• Block（线程块）：一个小组，组内工人可以互相聊天（共享内存）
• Grid（网格）：整个工厂，包含所有小组

为什么会设计成三层？因为GPU硬件就是这样组织的。一个SM（流多处理器）可以同时运行多个Block，每个Block里的Thread可以协作。嗯，这里要注意：同一个Block里的Thread可以通过共享内存通信，不同Block的Thread不能直接通信。

看一个简单的配置例子：

// 定义Grid和Block的维度
dim3 gridDim(2, 2);      // 2x2 = 4个Block
dim3 blockDim(16, 16);   // 16x16 = 256个Thread/Block

// 启动Kernel
myKernel<<<gridDim, blockDim>>>(d_data);

这个配置下，总共有 4 × 256 = 1024 个线程在跑。每个线程可以通过内置变量知道自己是谁：

• threadIdx.x：线程在Block内的x索引（0~15）
• blockIdx.x：Block在Grid内的x索引（0~1）
• blockDim.x：Block的x维度大小（16）
• gridDim.x：Grid的x维度大小（2）

我建议你刚开始时，先用一维的Grid和Block，等熟悉了再玩多维。我曾经在项目里为了炫技用了三维Grid，结果调试了三天才发现索引算错了……

3.3 内存层次结构——Global/Shared/Local

GPU的内存模型，说白了就是「越靠近计算单元，越快但越小」。我整理了一张表，你一看就明白：

内存类型	位置	访问速度	生命周期	作用域
Global Memory	显存	慢（~400周期）	程序运行期间	所有线程
Shared Memory	SM内部	快（~5周期）	Block生命周期	同一Block内
Local Memory	寄存器/显存	快/慢	线程生命周期	单个线程
Constant Memory	显存（带缓存）	快（只读）	程序运行期间	所有线程

这里有个坑：Local Memory并不总是快的。当寄存器不够用时，编译器会把变量「溢出」到显存，这时候Local Memory就变成了Global Memory的速度。我曾经在写一个矩阵乘法时，因为寄存器压力太大，性能直接腰斩。后来通过减少每个线程的临时变量，才把性能救回来。

3.4 第一个Kernel——向量加法

好了，理论说完了，咱们写个真正的Kernel。向量加法是CUDA的「Hello World」，简单但能说明所有核心概念。

先看CPU版本的代码：

// CPU版本：串行加法
void vecAddCPU(float* a, float* b, float* c, int n) {
    for (int i = 0; i < n; i++) {
        c[i] = a[i] + b[i];
    }
}

再看GPU版本：

// GPU版本：并行加法Kernel
__global__ void vecAddGPU(float* a, float* b, float* c, int n) {
    // 计算当前线程的全局ID
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    
    // 边界检查：防止越界
    if (tid < n) {
        c[tid] = a[tid] + b[tid];
    }
}

// 主函数中的调用
int main() {
    int n = 1024;
    size_t bytes = n * sizeof(float);
    
    // 1. 分配Host内存
    float *h_a, *h_b, *h_c;
    h_a = (float*)malloc(bytes);
    h_b = (float*)malloc(bytes);
    h_c = (float*)malloc(bytes);
    
    // 初始化数据...
    
    // 2. 分配Device显存
    float *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);
    
    // 3. 数据从Host拷贝到Device
    cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, bytes, cudaMemcpyHostToDevice);
    
    // 4. 启动Kernel
    int threadsPerBlock = 256;
    int blocksPerGrid = (n + threadsPerBlock - 1) / threadsPerBlock;
    vecAddGPU<<<blocksPerGrid, threadsPerBlock>>>(d_a, d_b, d_c, n);
    
    // 5. 结果拷回Host
    cudaMemcpy(h_c, d_c, bytes, cudaMemcpyDeviceToHost);
    
    // 6. 释放内存
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    free(h_a);
    free(h_b);
    free(h_c);
    
    return 0;
}

注意看几个关键点：

• __global__ 修饰符：告诉编译器这是个Kernel函数，由CPU调用、GPU执行
• cudaMalloc：在显存上分配空间，用法和malloc一样
• cudaMemcpy：数据传输，最后一个参数指定方向
• 边界检查：if (tid < n) 这行不能少。因为数据量不一定能被线程数整除，多余的线程什么都不做

3.5 调试与验证——别信你的第一个Kernel

写完了Kernel，怎么知道它算对了？我建议你这样做：

1. 先用小数据量（比如n=16）测试
2. CPU跑一遍串行版本，GPU跑一遍并行版本
3. 对比结果，看误差是否在可接受范围内

这里有个小技巧：用 cudaDeviceSynchronize() 确保Kernel执行完毕，再用 cudaGetLastError() 检查是否有错误。我曾经因为忘了加同步，结果拿到的全是垃圾数据，排查了半天……

下一章我们会深入Shared Memory的优化技巧，到时候你会看到，用好Shared Memory能让你的Kernel性能翻倍。嗯，先消化今天的内容吧。