一、GPU架构与CUDA编程模型：从硬件到实战

各位同学，欢迎来到《CUDA编程与NCCL联合调优实战》的第一章。

说实话，我见过太多人一上来就写CUDA代码，结果性能惨不忍睹。为什么？因为没搞懂GPU到底是怎么工作的。今天我们就从硬件架构讲起，把GPU的底裤扒干净。

1.1 GPU硬件架构：SM、Warp与Memory Hierarchy

先看一张图，这是我个人习惯——学任何东西之前，先看整体框架。

1.2 SM（流式多处理器）—— GPU的"心脏"

SM是GPU最核心的计算单元。你可以把它想象成一个"迷你CPU"，但它跟CPU完全不同。

SM内部长什么样？

CUDA Cores：每个SM里有几十到上百个。比如A100有128个，H100有256个。这些才是真正干活的"工人"。
Warp Scheduler：负责调度线程束。每个SM有4个，可以同时调度4个warp。
Register File：寄存器文件，容量巨大。A100每个SM有65536个32位寄存器。
Shared Memory：共享内存，可编程的L1缓存。A100是192KB，可以配置成L1或Shared。
Tensor Cores：专门做矩阵乘法的硬件单元。搞AI训练的同学，这个你们最熟。

我踩过的坑： 有一次我写了一个kernel，每个线程用了太多寄存器（超过128个），结果occupancy从100%掉到了25%。性能直接腰斩。后来我强制编译器用local memory，虽然慢了，但occupancy上来了，整体反而更快。嗯，这就是个trade-off。

1.3 Warp——GPU调度的基本单位

Warp是32个线程组成的集合。这是GPU执行指令的最小单位。

关键点：

一个warp里的32个线程执行同一条指令（SIMT模型）
如果线程走不同的分支（if-else），就会发生warp divergence
warp divergence会导致部分线程被"屏蔽"，性能下降

实战建议： 我写代码时，会刻意让同一个warp里的线程处理相邻的数据。这样不仅能避免分支发散，还能利用内存合并访问。你想想看，32个线程读连续的内存地址，一次就能搞定，多爽。

1.4 Memory Hierarchy——GPU的"存储金字塔"

GPU的内存层次，说白了就是"越快的越贵，越贵的越小"。

存储层次	大小	带宽	延迟	作用域
寄存器	256KB/SM	~20TB/s	1 cycle	线程私有
共享内存	48-164KB/SM	~10TB/s	~20 cycles	Block内共享
L1缓存	128KB/SM	~5TB/s	~30 cycles	SM内
L2缓存	6-40MB	~2TB/s	~200 cycles	芯片全局
全局内存	16-80GB	~1.5TB/s	~400 cycles	所有线程

注意： 全局内存的延迟是寄存器的400倍！所以，尽量减少全局内存访问。我见过有人把循环里的变量声明在全局内存里，结果性能惨不忍睹。记住：能用寄存器就别用共享内存，能用共享内存就别用全局内存。

1.5 CUDA线程层次结构：Grid、Block、Thread

CUDA的线程模型是分层的。说白了就是：

Thread：最底层的工作单元，每个线程执行一个任务
Block：一组线程，共享共享内存，可以同步
Grid：一组Block，构成一个完整的kernel

为什么会这样设计？因为GPU的硬件就是这样映射的：

一个Block映射到一个SM
Block里的线程被分成warp执行
多个Block可以同时运行在同一个SM上（取决于资源）

// 典型的CUDA kernel启动
dim3 gridDim(16, 16);    // 16x16 = 256个Block
dim3 blockDim(32, 32);   // 32x32 = 1024个线程/Block

myKernel<<<gridDim, blockDim>>>(d_data, N);

// 在kernel内部获取索引
__global__ void myKernel(float* data, int N) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    // 或者用2D索引
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
}

我的经验： Block大小选256或512通常比较稳妥。太小了SM利用率低，太大了寄存器不够用。我一般从256开始试，然后看occupancy，再调整。

1.6 Kernel Launch机制——GPU到底怎么执行你的代码？

Kernel Launch看起来就是一行代码，但背后发生的事情可不少：

CPU端准备：把数据从CPU拷贝到GPU（cudaMemcpy）
Kernel提交：CPU把kernel函数和参数提交到GPU的命令队列
GPU调度：GPU的GigaThread引擎把Block分配到SM上
Warp调度：SM上的warp scheduler开始调度warp执行
指令发射：每个cycle，warp scheduler选择一个ready的warp发射指令
结果写回：执行完成后，数据写回全局内存
同步：CPU等待GPU完成（cudaDeviceSynchronize）

避坑指南： 我曾经犯过一个错误——在循环里反复launch小kernel。每次launch都有几百微秒的开销，结果性能还不如在CPU上跑。后来我把多个小kernel合并成一个，性能提升了10倍。记住：kernel launch不是免费的！

1.7 实战：如何用occupancy API优化配置

CUDA提供了occupancy API，可以帮你计算最优的block大小和shared memory用量。

// 计算最优block大小
int minGridSize, blockSize;
cudaOccupancyMaxPotentialBlockSize(
    &minGridSize, 
    &blockSize, 
    myKernel, 
    0,  // dynamic shared memory size
    0   // block size limit (0 = no limit)
);

// 或者指定shared memory用量
size_t dynamicSmem = 48 * 1024;  // 48KB
cudaOccupancyMaxPotentialBlockSize(
    &minGridSize, 
    &blockSize, 
    myKernel, 
    dynamicSmem, 
    0
);

printf("Recommended block size: %d\n", blockSize);
printf("Min grid size: %d\n", minGridSize);

核心要点： occupancy不是越高越好。有时候occupancy 50%但每个线程用更多寄存器，反而比occupancy 100%但寄存器溢出到local memory要快。我一般会跑几个配置对比一下，选最快的那个。

1.8 本章小结

好了，这一章我们讲了GPU的硬件架构和CUDA编程模型。说白了就是三件事：

硬件：SM是核心，warp是调度单位，内存有层次
线程模型：Grid→Block→Thread，映射到硬件上
Kernel Launch：不是免费的，要合理配置

下一章我们会深入讲CUDA内存管理，包括各种内存的用法和优化技巧。到时候我会分享一些实际项目中的调优案例，保证让你收获满满。

想要更多资料，联系可可

📚 课程目录

一、GPU架构与CUDA编程模型：从硬件到实战

1.1 GPU硬件架构：SM、Warp与Memory Hierarchy

1.2 SM（流式多处理器）—— GPU的"心脏"

1.3 Warp——GPU调度的基本单位

1.4 Memory Hierarchy——GPU的"存储金字塔"

1.5 CUDA线程层次结构：Grid、Block、Thread

1.6 Kernel Launch机制——GPU到底怎么执行你的代码？

1.7 实战：如何用occupancy API优化配置

1.8 本章小结