第一章：CUDA并行计算基础

GPU架构概览：从图形卡到通用计算

说实话，我第一次接触GPU编程时，脑子里想的是：「这不就是个显卡吗？」后来才发现，GPU和CPU的设计哲学完全不同。

CPU像个全能选手——什么活都能干，但一次只能干几件事。GPU呢？它像个流水线工厂——成千上万个简单工人同时干活。我习惯把GPU想象成一个拥有几千个核心的「计算怪兽」，每个核心虽然跑得慢，但架不住数量多啊。

来看个关键数据：

特性	CPU	GPU
核心数量	4-16个	数千个
缓存大小	大（MB级）	小（KB级）
控制单元	复杂	简单
适合任务	串行、复杂逻辑	并行、数据密集

你想想看，CPU花了大把晶体管做分支预测、乱序执行这些花活。GPU呢？它把晶体管都堆在计算单元上。所以GPU特别适合那种「同样的操作，不同的数据」的场景——比如图像处理、矩阵乘法、物理模拟。

我在项目中遇到过一位同事，非要用CPU做大规模图像滤波。结果呢？一张4K图片处理了3秒。换成GPU后，同样的算法跑了不到10毫秒。这就是架构差异带来的质变。

CUDA编程模型：主机与设备的协作

CUDA把系统分成两部分：主机（Host）和设备（Device）。主机就是CPU那端，设备就是GPU那端。说白了，CPU负责调度，GPU负责干活。

我刚开始写CUDA程序时，总忘记把数据拷贝到设备端。结果GPU算了个寂寞——它拿到的全是垃圾数据。嗯，这里要注意：CPU和GPU有各自独立的内存空间。

典型的CUDA程序流程是这样的：

1. 在主机端分配内存并初始化数据
2. 在设备端分配显存
3. 把数据从主机拷贝到设备
4. 调用核函数（Kernel）在GPU上执行
5. 把结果从设备拷贝回主机
6. 释放内存

来看个最简单的例子：

// 核函数：每个线程加一个数
__global__ void add_kernel(int* a, int* b, int* 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;
    int *d_a, *d_b, *d_c;
    
    // 1. 分配设备内存
    cudaMalloc(&d_a, n * sizeof(int));
    cudaMalloc(&d_b, n * sizeof(int));
    cudaMalloc(&d_c, n * sizeof(int));
    
    // 2. 拷贝数据到设备
    cudaMemcpy(d_a, h_a, n * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b, n * sizeof(int), cudaMemcpyHostToDevice);
    
    // 3. 启动核函数
    add_kernel<<<1, 256>>>(d_a, d_b, d_c, n);
    
    // 4. 拷贝结果回主机
    cudaMemcpy(h_c, d_c, n * sizeof(int), cudaMemcpyDeviceToHost);
    
    // 5. 清理
    cudaFree(d_a); cudaFree(d_b); cudaFree(d_c);
    return 0;
}

线程层次结构：Grid-Block-Thread

CUDA的线程组织方式，我管它叫「三层嵌套」：

• Thread（线程）：最基础的计算单元
• Block（线程块）：一组线程，可以协作
• Grid（网格）：一组线程块，构成整个计算任务

为什么会这样设计？你想想看，GPU有几千个核心，如果让程序员直接管理每个线程，那不得疯掉？所以CUDA搞了个层次结构——你只需要告诉它「我要启动多少个块，每个块里有多少线程」，剩下的调度交给硬件。

我记得第一次写CUDA程序时，设了blockDim.x=1024，结果程序直接崩了。后来查文档才知道，每个block的线程数上限是1024（不同架构略有差异）。

层级	最大数量	说明
Grid维度	2^31-1（x方向）	可以是一维、二维、三维
Block维度	1024（总线程数）	每个维度也有上限（x:1024, y:1024, z:64）
Thread	无上限	由block和grid共同决定

每个线程都有自己唯一的ID，通过内置变量获取：

• threadIdx.x：线程在block内的索引
• blockIdx.x：block在grid内的索引
• blockDim.x：block的维度大小
• gridDim.x：grid的维度大小

计算全局索引的公式：int global_idx = threadIdx.x + blockIdx.x * blockDim.x;

如何选择合适的Grid和Block大小

这个问题，我面试时经常问候选人。答案不是固定的，但有几个原则：

1. block大小最好是32的倍数——因为warp（线程束）是32个线程一组。不是32的倍数会导致warp浪费。
2. block大小建议128-256——太小了SM利用率低，太大了寄存器压力大。
3. grid大小由数据量决定——总线程数要覆盖所有数据。

举个例子，处理1024x1024的图像：

dim3 blockDim(16, 16);  // 每个block 256个线程
dim3 gridDim(64, 64);   // 总共4096个block
// 总线程数：256 * 4096 = 1,048,576，刚好覆盖所有像素

你可能会问：「为什么blockDim用16x16而不是32x32？」嗯，32x32是1024个线程，虽然没超上限，但每个线程能用的寄存器就少了。我一般从128开始试，然后调优。

小结

这一章我们聊了GPU为什么快、CUDA怎么组织线程、以及如何配置Grid和Block。说白了，CUDA并行计算的核心就三件事：数据准备好、线程组织好、核函数写好。

下一章我们会深入原子操作——当多个线程同时读写同一块数据时，怎么保证不出乱子。嗯，那才是真正考验功底的地方。