1. CUDA编程模型回顾：Grid-Block-Thread层级结构、Kernel Launch语法、主机与设备内存交互

好，咱们正式开始。第一章，我打算先带大家快速回顾一下CUDA编程模型的核心骨架。这些东西你可能看过很多遍了，但相信我，很多坑都藏在这些基础概念里。我自己带团队做GPU加速项目时，发现80%的性能问题，根源都是对这几个基本概念理解得不够透彻。

1.1 Grid-Block-Thread：GPU的“三体”结构

先问大家一个问题：CPU和GPU最大的区别在哪？

CPU是“大力士”，核心少但每个核心都很强，适合处理复杂的串行任务。GPU是“蚂蚁军团”，核心多但每个核心相对简单，适合处理大量并行的简单任务。

CUDA为了让咱们能高效地指挥这支“蚂蚁军团”，设计了一套三层组织架构：Grid（网格）→ Block（线程块）→ Thread（线程）。

  • Thread（线程）：最基础的执行单元。每个线程执行相同的代码（即kernel函数），但处理不同的数据。
  • Block（线程块）：一组线程的集合。同一个Block内的线程可以共享一块内存（Shared Memory），并且可以通过同步指令（__syncthreads()）来协调执行顺序。
  • Grid（网格）：一组Block的集合。一个Grid对应一个Kernel函数的启动。

嗯，这里要注意：Block内的线程数量是有限制的。我见过不少新手一上来就设1024个线程在一个Block里，结果程序跑不起来。为什么？因为每个Block能用的资源（寄存器、共享内存）是有限的。具体限制可以查你的GPU计算能力（Compute Capability）对应的文档。

核心记忆点：

  • Grid 是任务的总称
  • Block 是任务的分组
  • Thread 是干活的最小单位

为了让你更直观地理解这个结构，我画了一张图：

Grid（网格） Block (0,0) Block (1,0) Block (2,0) T0 T1 T2 T3 T4 T5 ... ... ... Thread

1.2 Kernel Launch 语法：怎么把任务扔给GPU

搞清楚了结构，接下来就是怎么启动它。在CUDA里，启动一个GPU函数（叫Kernel）的语法长这样：

// 定义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];
    }
}

// 启动Kernel
int main() {
    int N = 1 << 20;  // 约100万个元素
    int threadsPerBlock = 256;
    int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;

    vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);
    return 0;
}

这里有几个关键点，我当年刚学时也踩过坑：

  • __global__：这个修饰符告诉编译器，这个函数是在GPU上运行，但可以从CPU调用。说白了，它就是GPU程序的入口。
  • <<<grid, block>>>：这是CUDA特有的启动语法。第一个参数是Grid里有多少个Block，第二个参数是每个Block里有多少个Thread。
  • threadIdx.x, blockIdx.x, blockDim.x：这三个内置变量帮你计算当前线程的全局ID。我个人习惯用 idx = threadIdx.x + blockIdx.x * blockDim.x 这个公式，几乎成了肌肉记忆。

小技巧：启动Kernel时，Block数量最好设成SM（流多处理器）数量的整数倍。比如你的GPU有20个SM，那Block数设成40、80、160效果会更好。为什么？因为这样能让所有SM都忙起来，不会有的SM闲着看热闹。

1.3 主机与设备内存交互：数据怎么搬家

GPU不能直接访问CPU的内存（主机内存），CPU也不能直接访问GPU的内存（设备内存）。所以数据需要“搬家”。这个搬家过程，说白了就是三步：

  1. 分配设备内存：用 cudaMalloc() 在GPU上划一块地
  2. 数据拷贝：用 cudaMemcpy() 把数据从CPU搬到GPU，或者从GPU搬回CPU
  3. 释放内存：用 cudaFree() 把GPU上的地还回去

来看个完整的例子：

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

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

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

// 4. 数据从主机拷贝到设备
cudaMemcpy(d_A, h_A, N * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_B, h_B, N * sizeof(float), cudaMemcpyHostToDevice);

// 5. 启动Kernel
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

// 6. 结果从设备拷贝回主机
cudaMemcpy(h_C, d_C, N * sizeof(float), cudaMemcpyDeviceToHost);

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

避坑指南：我曾经犯过一个低级错误——在Kernel还没执行完就调用了 cudaMemcpy 把结果拷回来。结果拷回来的全是垃圾数据。记住：cudaMemcpy 默认是同步操作，它会等Kernel执行完再拷贝。但如果你用了异步流（Stream），那就得手动加 cudaDeviceSynchronize() 了。

1.4 内存拷贝的代价：别让数据搬家拖慢你

你想想看，PCIe 3.0 x16的理论带宽大约是16GB/s，而GPU显存带宽动辄几百GB/s。数据拷贝的速度远低于GPU计算的速度。所以，尽量减少主机和设备之间的数据拷贝次数，这是优化的第一原则。

操作 带宽（约） 延迟（约）
CPU内存 → GPU显存（PCIe 3.0） 12 GB/s 几微秒
GPU显存内部 900 GB/s（H100） 几百纳秒
GPU共享内存 极高（片上） 几十纳秒

看到差距了吧？显存内部带宽是PCIe的几十倍。所以我的建议是：尽量在GPU上把数据算完，再一次性拷回CPU。别算一步拷一步，那样性能会惨不忍睹。

本章核心要点：

  • Grid-Block-Thread 是CUDA的并行组织架构，理解它才能用好GPU
  • Kernel启动用 <<<grid, block>>> 语法，注意Block大小不要超过硬件限制
  • 主机和设备内存交互用 cudaMemcpy，但尽量少用，因为PCIe带宽是瓶颈
  • 记得检查每个CUDA API的返回值，用 cudaGetLastError() 捕获错误

好了，第一章就到这里。这些基础概念虽然简单，但它们是后面所有优化的基石。下一章我们会深入聊聊线程束（Warp）和线程调度，那才是真正开始“榨干”GPU性能的地方。

可可精品资料库

← 上一章 📖 返回目录 下一章 →