2、CUDA编程模型回顾：线程层次结构、内存层次结构、核函数基础

好，咱们正式开始第二讲。说实话，很多朋友一上来就急着学向量化指令，结果连线程怎么组织的都没搞明白，最后代码跑出来比CPU还慢。嗯，这我见得太多了。所以今天咱们花点时间，把CUDA编程模型的三个核心概念彻底捋一遍：线程层次结构、内存层次结构、还有核函数怎么写才靠谱。

2.1 线程层次结构：别把GPU当CPU用

我第一次写CUDA程序时，犯过一个特别蠢的错误——我直接把一个for循环拆成1000个线程，心想这不就并行了吗？结果跑出来比单线程还慢。为什么？因为我不懂线程是怎么组织的。

GPU的线程不是平铺的，它是分层的。你想想看，一个GPU里有成百上千个核心，如果每个线程都独立调度，那调度器得累死。所以NVIDIA搞了一套层级结构：线程（Thread）→ 线程块（Block）→ 网格（Grid）。

  • 线程（Thread）：最基础的执行单元。每个线程执行相同的核函数，但处理不同的数据。
  • 线程块（Block）：一组线程的集合。同一个Block里的线程可以共享数据（通过共享内存），也能同步（通过__syncthreads()）。
  • 网格（Grid）：所有Block的集合。Grid里的Block之间是独立的，没法直接通信。

这里有个关键点：Block里的线程数不能太多。我记得有一次项目里，我贪心地把Block设成1024个线程，结果发现SM（流多处理器）里资源不够，实际并行度反而下降了。一般来说，每个Block 256或512个线程是比较稳妥的选择。

核心公式：
总线程数 = GridDim.x × BlockDim.x（一维情况）
每个线程的全局索引 = blockIdx.x × blockDim.x + threadIdx.x

我习惯用一维索引来处理数组，简单直观。但如果你处理的是图像或矩阵，二维索引会更方便。举个例子：

// 一维索引，适合处理数组
int tid = blockIdx.x * blockDim.x + threadIdx.x;

// 二维索引，适合处理图像
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
int idx = y * width + x;

我的小技巧：调试时先用一维索引把逻辑跑通，再优化成二维。别一上来就搞复杂索引，容易把自己绕晕。

2.2 内存层次结构：速度与容量的博弈

说到内存，我得先吐槽一句：很多人以为GPU显存和CPU内存一样，随便读写就行。大错特错！GPU的内存是有层次的，每一层的速度、容量、作用都不一样。我当年在做一个流体模拟项目时，就是因为没用好共享内存，性能死活上不去。后来一改，速度翻了5倍。

咱们来看看这张表：

内存类型 位置 速度 容量 作用域
全局内存 显存 慢（几百周期） 大（GB级） 所有线程
共享内存 SM内部 快（几个周期） 小（几十KB） 同一个Block
寄存器 SM内部 极快（1周期） 极小（每个线程几十个） 单个线程
常量内存 显存（带缓存） 快（有缓存） 小（64KB） 所有线程（只读）
纹理内存 显存（带缓存） 快（有缓存） 大 所有线程（只读）

你看，全局内存虽然大，但慢得要命。共享内存虽然小，但快得飞起。所以优化的核心思路就是：尽量把数据从全局内存搬到共享内存里，让线程们反复用。

注意：共享内存是有限的！我曾经在一个项目里，每个Block用了48KB共享内存，结果发现SM只能同时运行很少的Block，并行度大打折扣。后来我压缩到32KB，性能反而提升了。记住：共享内存不是越大越好，要留点空间给寄存器和其他资源。

另外，内存合并访问是个老生常谈的话题。说白了，就是让相邻的线程访问相邻的内存地址。比如线程0访问地址0，线程1访问地址1，这样GPU可以一次把一整段数据读进来。如果线程0访问地址0，线程1访问地址100，那就得读两次，效率直接腰斩。

2.3 核函数基础：怎么写才不翻车

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

  • 用__global__修饰符声明
  • 返回类型必须是void
  • 不能使用可变参数（比如printf那种）
  • 不能是类的成员函数（静态成员函数可以）

我刚开始写核函数时，总喜欢在里面放很多if-else分支。结果发现性能很差。为什么？因为GPU是SIMT架构，同一个warp（32个线程）里的线程必须执行相同的指令。如果有分支，有些线程走if，有些走else，那没走的线程就得等着，这叫warp divergence。

避坑指南：尽量减少核函数里的分支。如果实在避免不了，让同一个warp里的线程走同一个分支。比如用if (threadIdx.x < 16)，这样前16个线程走if，后16个走else，不会产生warp divergence。

来看一个最简单的核函数示例：

__global__ void vec_add(float *A, float *B, float *C, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        C[idx] = A[idx] + B[idx];
    }
}

调用方式也很简单：

int N = 1024;
int threads_per_block = 256;
int blocks_per_grid = (N + threads_per_block - 1) / threads_per_block;
vec_add<<<blocks_per_grid, threads_per_block>>>(d_A, d_B, d_C, N);

这里有个细节：blocks_per_grid的计算用了向上取整。因为N不一定能被threads_per_block整除，多出来的线程会在核函数里通过if (idx < N)过滤掉。嗯，这个写法很常见，但别忘了加边界检查，否则会访问越界。

我的习惯：在核函数开头先算好全局索引，然后做一次边界检查。这样后面写代码就不用担心越界了，省心。

2.4 总结：记住这三句话

好了，这一章的内容就这些。说实话，CUDA编程模型说复杂也复杂，说简单也简单。我个人觉得，只要记住三句话就够了：

  1. 线程是分层的：Grid → Block → Thread，别搞混了索引。
  2. 内存是有层次的：全局内存慢但大，共享内存快但小，用共享内存做缓存是性能优化的关键。
  3. 核函数要简洁：少分支，多合并访问，别忘了边界检查。

下一章咱们就要进入正题了——向量化指令集。说实话，向量化指令才是真正让GPU飞起来的东西。但如果你连线程和内存都没搞明白，向量化指令用起来也是白搭。所以，这一章的内容一定要消化透。

嗯，今天就到这儿。有什么问题，咱们下节课再聊。

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