1. CUDA编程模型回顾

做GPU优化这么多年，我见过太多人一上来就盯着代码调优，结果连最基本的线程模型都没搞明白。说实话，CUDA编程模型就像盖房子的地基——地基不稳，后面再怎么装修也是白搭。今天咱们就把这块地基彻底夯实了。

1.1 Grid-Block-Thread层级结构

先看这张图，它概括了CUDA最核心的层级关系：

Grid（网格） Block (0,0) Block (1,0) Block (2,0) T(0,0) T(1,0) T(2,0) T(0,1) T(1,1) T(2,1) ... ... ... 一个Grid包含多个Block，每个Block包含多个Thread

这张图我画了很多遍，每次给团队新人讲课时都会拿出来。你看，最外层是Grid，里面是Block，最里面是Thread。这个三层结构，说白了就是GPU用来组织并行任务的骨架。

Grid 是整个任务的总称。你调用一个kernel函数，就启动了一个Grid。Grid可以是一维、二维或三维的，但实际项目中我99%的情况只用一维或二维——三维用得少，除非处理体数据。

Block 是中间层，也是优化的关键。每个Block内的线程可以协作——通过共享内存和同步操作。我记得刚入行时犯过一个错误：把Block设得太大，结果共享内存不够用，程序直接崩了。后来养成了习惯，先算好共享内存预算再定Block大小。

Thread 是最小执行单元。每个线程执行相同的kernel代码，但处理不同的数据——这就是SIMT的精髓。你想想看，成千上万个线程同时跑，每个线程拿自己的索引去取数据，效率能不高吗？

关键约束：

  • 每个Block最多1024个线程（现代GPU，老架构是512或768）
  • Grid维度最大为2^31-1（一维）或65535×65535（二维）
  • Block内的线程通过 __syncthreads() 同步

1.2 内存层次结构——谁快谁慢，心里要有数

做GPU优化，说白了就是在跟内存打交道。我见过太多人把数据一股脑扔进全局内存，然后抱怨GPU慢。其实GPU的内存层次很丰富，用对了地方，性能能翻好几倍。

内存类型 位置 访问延迟 生命周期 作用域
Register SM内部 ~1 cycle 线程内 单个线程
Shared Memory SM内部 ~5 cycles Block内 Block内所有线程
Local Memory 显存（DRAM） ~100 cycles 线程内 单个线程（寄存器溢出时）
Global Memory 显存（DRAM） ~400-800 cycles 整个程序 所有线程 + Host
Constant Memory 显存（带缓存） ~100 cycles（缓存命中时更快） 整个程序 所有线程（只读）
Texture Memory 显存（带缓存） ~100 cycles 整个程序 所有线程（只读，支持插值）

这张表我建议你打印出来贴在工位上。每次写kernel之前看一眼，问自己三个问题：

  1. 这个数据能不能放寄存器？能就别放共享内存。
  2. 这个数据能不能放共享内存？能就别放全局内存。
  3. 这个数据是不是只读且所有线程都访问？是就考虑常量内存。

我的经验： 寄存器是最快的，但数量有限。每个线程最多255个寄存器（现代GPU）。如果你写的代码用了太多局部变量，编译器会把多余的变量"溢出"到Local Memory——这玩意儿其实在显存里，慢得很。我曾经优化过一个矩阵乘法的kernel，就是靠减少寄存器使用量，把性能从60%提升到了85%。

1.3 Shared Memory——Block内的"小团队协作区"

共享内存是CUDA优化的核心武器。它位于SM内部，速度比全局内存快两个数量级。但它的容量很小——通常每个SM只有48KB到164KB（取决于架构）。

共享内存的典型用法：

  • 数据复用： 多个线程需要访问同一块数据时，先加载到共享内存
  • 合并访问： 把非合并的全局内存访问模式，转成合并的共享内存访问
  • 线程间通信： Block内的线程通过共享内存交换数据
// 共享内存的典型用法：矩阵分块
__global__ void matmul_shared(float *A, float *B, float *C, int N) {
    __shared__ float As[TILE_SIZE][TILE_SIZE];
    __shared__ float Bs[TILE_SIZE][TILE_SIZE];
    
    int bx = blockIdx.x, by = blockIdx.y;
    int tx = threadIdx.x, ty = threadIdx.y;
    
    int row = by * TILE_SIZE + ty;
    int col = bx * TILE_SIZE + tx;
    
    float sum = 0.0f;
    
    for (int t = 0; t < N / TILE_SIZE; t++) {
        // 协作加载数据到共享内存
        As[ty][tx] = A[row * N + t * TILE_SIZE + tx];
        Bs[ty][tx] = B[(t * TILE_SIZE + ty) * N + col];
        __syncthreads();
        
        // 从共享内存计算
        for (int k = 0; k < TILE_SIZE; k++) {
            sum += As[ty][k] * Bs[k][tx];
        }
        __syncthreads();
    }
    
    C[row * N + col] = sum;
}

注意： 共享内存使用不当会导致 bank conflict。共享内存被分成32个bank（每个bank 4字节宽），如果多个线程同时访问同一个bank的不同地址，就会发生冲突，导致串行化。我踩过这个坑——一个简单的转置kernel，因为没处理好bank conflict，性能直接腰斩。

1.4 Warp执行模型与SIMT架构

这是CUDA最核心、也最容易让人困惑的概念。说白了，GPU并不是真的同时执行所有线程——它把32个线程打包成一个Warp，以Warp为单位调度执行。

为什么是32？ 这是NVIDIA的设计选择。32个线程共享一个指令单元，执行相同的指令。如果线程之间分支不同（比如if-else），就会发生Warp Divergence——一部分线程执行，另一部分等待，性能直接打折扣。

Warp Divergence 示例：

// 不好的写法——会导致Warp Divergence
if (threadIdx.x % 2 == 0) {
    // 偶数线程执行这里
    result = a + b;
} else {
    // 奇数线程执行这里
    result = a * b;
}

// 好的写法——避免分支
int mask = threadIdx.x & 1;  // 0或1
result = mask ? (a * b) : (a + b);

我曾经在一个图像处理项目里，因为一个简单的条件判断，导致Warp内一半线程闲置，性能掉了40%。后来改成用位运算和三元表达式，问题就解决了。你想想看，32个线程里如果有16个走分支A、16个走分支B，那这16个等待的时候，计算单元就空转了。

SIMT vs SIMD：

  • SIMD（单指令多数据）： 所有处理单元必须执行相同指令，数据宽度固定（如SSE的128位）
  • SIMT（单指令多线程）： 每个线程有自己的程序计数器，可以独立分支（但分支会导致性能损失）

SIMT更灵活，但代价就是Warp Divergence。所以写CUDA代码时，要尽量让同一个Warp内的线程走相同的执行路径。

1.5 实战中的避坑指南

说了这么多理论，最后分享几个我在项目中踩过的坑：

避坑1： 我曾经以为Block越大越好，结果把Block设成了1024，但每个线程用了太多寄存器，导致Occupancy（占用率）只有25%。后来把Block改成256，Occupancy提升到75%，性能反而翻倍了。记住：不是线程越多越好，是活跃的Warp越多越好。

避坑2： 共享内存的bank conflict不是总能避免的。如果实在避不开，可以尝试padding——在共享内存数组的每一行末尾加几个空元素，改变地址对齐方式。我有个同事靠这招把性能提升了30%。

避坑3： 全局内存的合并访问。Warp访问全局内存时，如果32个线程访问的地址是连续的（且对齐到128字节），硬件会合并成一次或几次内存事务。否则，每个线程单独访问，效率极低。我见过有人写 array[threadIdx.x * 32] 这种模式，结果性能惨不忍睹。

好了，这一章的内容就到这里。记住这些基础概念，后面我们讲具体优化技巧时，你会发现自己能更快地理解问题出在哪里。

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