第一章：CUDA编程模型回顾——从线程到内存，把GPU算力榨干

各位同学，欢迎来到《CUDA算子开发进阶》。

我是这门课的主讲。做了十年GPU计算，踩过的坑比写过的代码还多。今天咱们先不急着写算子，把地基打牢。

第一章，说白了就是一次“回锅”。把CUDA编程模型里最核心的三个东西——线程层次、内存层次、核函数写法——重新捋一遍。我保证，哪怕你觉得自己已经懂了，这一章也能挖出点新东西。

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

先问个问题：GPU为什么快？

因为它能同时跑成千上万个线程。但这么多线程，总得有个组织方式吧？

CUDA的组织方式很简单：Grid → Block → Thread，三层结构。

Thread：最基础的执行单元。每个线程执行相同的核函数，但处理不同的数据。
Block：一组线程的集合。同一个Block里的线程可以协作——共享内存、同步、通信。
Grid：所有Block的集合。一个Grid对应一次核函数调用。

我刚开始学的时候，总觉得这玩意儿就是个概念。直到有一次写一个矩阵乘法，Block大小设成32×32，结果性能死活上不去。后来发现，Block里的线程数太少，SM（流多处理器）的资源没喂饱。嗯，这里要注意：Block大小不是随便选的。

核心原则：

Block的线程数最好是32的倍数（warp大小）
一般建议128~512个线程/Block
Grid的维度取决于数据量，但别让Block太少，否则SM空闲

举个例子，你要处理一个1024×1024的矩阵，每个线程处理一个元素：

dim3 blockDim(16, 16);   // 每个Block 256个线程
dim3 gridDim(64, 64);    // 总共4096个Block
kernel<<<gridDim, blockDim>>>(d_data);

为什么Block用16×16？因为256个线程刚好能喂饱一个SM上的资源，又不至于让寄存器压力太大。我个人的习惯是，先试32×32（1024线程），如果寄存器溢出就降下来。

小技巧：用 cudaOccupancyMaxPotentialBlockSize 这个API，它能帮你算出当前硬件上最优的Block大小。省得你一个个试。

1.2 内存层次结构：Global、Shared、Local、Register

GPU的内存，说白了就是“越近越快，越远越慢”。但很多人只记住了这句话，没搞懂怎么用。

内存类型	位置	速度	作用域	生命周期
Register	SM内部	最快（~1 cycle）	单个线程	核函数执行期间
Shared Memory	SM内部	快（~5 cycle）	单个Block	Block生命周期
Local Memory	显存中	慢（~100 cycle）	单个线程	核函数执行期间
Global Memory	显存中	最慢（~400 cycle）	所有线程	程序生命周期

你看这张表，Register最快，但数量有限。每个线程大概能用255个（具体看架构）。Shared Memory也快，但一个Block通常只有48KB~96KB。Global Memory最慢，但容量最大。

我曾经犯过一个低级错误：把一个频繁访问的变量忘了加 __shared__，结果每次循环都去读Global Memory。性能直接掉了3倍。你想想看，一个循环里读100次Global Memory，跟读1次Shared Memory再反复用，差距有多大？

避坑指南：

别把大数组声明成局部变量——它会自动落到Local Memory里，速度跟Global一样慢
Shared Memory用多了会减少每个SM能同时运行的Block数，这叫“占用率下降”
Register溢出到Local Memory时，编译器不会报错，但性能会暴跌。用 --ptxas-options=-v 编译选项可以查看寄存器使用量

1.3 核函数编写基础

核函数，就是跑在GPU上的函数。写法跟C函数差不多，但有几个硬规矩：

用 __global__ 修饰
返回类型必须是 void
不能有可变参数（比如 ...）
不能是类的成员函数（静态成员函数可以）

来，看个最简单的例子——向量加法：

__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 N = 1 << 20;
int threadsPerBlock = 256;
int blocksPerGrid = (N + threadsPerBlock - 1) / threadsPerBlock;
vecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

这里有个细节：if (idx < N) 这个边界检查不能省。因为数据量不一定能被Block大小整除，最后一个Block可能有多余的线程。我见过有人偷懒不写，结果数组越界，数据写到了别人的内存里——调试了一整天。

核函数编写要点：

用 threadIdx、blockIdx、blockDim 计算全局索引
总是做边界检查
尽量减少Global Memory访问，多用Shared Memory
避免线程发散（同一个warp里的线程走不同分支）

说到线程发散，我再多嘴一句。warp是GPU调度的基本单位，32个线程一起执行。如果这32个线程里，有的走if分支，有的走else分支，那它们就得串行执行。性能直接砍半。所以写核函数时，尽量让同一个warp里的线程走相同的路径。

个人经验：我写核函数时，习惯先把数据从Global Memory搬到Shared Memory里，然后在Shared Memory里做计算，最后再写回Global Memory。这样能减少90%的Global Memory访问。代价是多写几行代码，但性能提升是实打实的。

本章知识体系

下面这张图，把本章的核心逻辑串起来了。你可以把它当成一张“地图”，以后写算子时对照着看。

这张图从左到右，从上到下，就是CUDA编程模型的完整脉络。你写任何算子，都逃不出这个框架。

好了，第一章就到这里。内容不多，但都是硬通货。下一章我们会深入Shared Memory的优化技巧——那才是真正拉开性能差距的地方。

想要更多资料，联系可可

📚 课程目录

第一章：CUDA编程模型回顾——从线程到内存，把GPU算力榨干

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

1.2 内存层次结构：Global、Shared、Local、Register

1.3 核函数编写基础

本章知识体系