2 GPU架构基础：SM、Warp、Grid-Block-Thread层次、内存层次结构
搞GPU编程，你不理解硬件长什么样，写出来的代码大概率跑不满性能。我最早接触CUDA时，觉得把任务扔给GPU就完事了，结果发现一个简单的向量加法，带宽利用率不到20%。后来老老实实啃了一遍架构手册，才明白问题出在哪。
这一节我们就拆开GPU看看。你不需要成为芯片设计师，但得知道那些玩意儿叫什么、怎么配合。
2.1 SM与Warp——GPU的微架构心脏
先从一个核心概念说起：SM（Streaming Multiprocessor）。你可以把它理解成GPU里的“计算单元小集群”。一块GPU有多少个SM，直接决定了它的并行能力。每个SM里有一堆CUDA Core、共享内存、寄存器文件、调度器等等。
那Warp呢？Warp是SM调度的基本单位，一个Warp包含32个线程。为什么是32？这是NVIDIA设计时的权衡——32个线程一起执行，可以最大化指令发射效率。我曾在项目里写过一些分支很重的kernel，结果一个Warp里32个线程走不同的分支，性能掉了差不多一半。说白了，同一个Warp里的线程，执行路径越一致，性能越好。
| 概念 | 说明 | 我踩过的坑 |
|---|---|---|
| SM | 物理计算单元，包含CUDA Core、寄存器、共享内存等 | 我曾经以为SM数越多越好，却没看每个SM的寄存器总量，导致线程数过多触发寄存器溢出 |
| Warp | 32个线程组成的调度单元，同一个Warp执行相同的指令 | 记住：一个分支让Warp串行化，32个线程变16+16… 慢得你想哭 |
个人技巧：调优时先看Occupancy（占用率）报告。NVCC的--ptxas-options=-v会告诉你每个SM用了多少寄存器和共享内存，据此调整block大小和线程数。
2.2 Grid-Block-Thread——你的任务怎么拆
写CUDA程序时，你需要把一个大任务拆成三层：Grid → Block → Thread。这对应硬件上的：GPU整体（Grid）→ SM（Block）→ Warp & Thread。
// 典型的调用方式
dim3 gridSize(16, 16); // 16x16 = 256个Block
dim3 blockSize(32, 4); // 每个Block有128个线程（32*4）
myKernel<<<gridSize, blockSize>>>();这里要注意几点：
- 一个Block只会被调度到一个SM上，但一个SM可以同时跑多个Block。
- Block里的线程通过共享内存通信，Block之间不能直接共享数据（除非用全局内存再同步，但效率低）。
- 每个Block里的线程数最好是32的倍数（因为Warp是32），我习惯用128或256，既保证Occupancy又不会让寄存器压力太大。
2.3 内存层次结构——数据该放哪儿
GPU有多个内存层级，搬数据的速度天差地别。我画个图你就明白了（后面有SVG）。先列个表格让你一目了然：
| 内存类型 | 位置 | 容量 | 延迟（大约） | 作用域 |
|---|---|---|---|---|
| 寄存器 | 每个线程私有 | 通常256KB/SM | ~1 cycle | 线程 |
| 共享内存 | SM内部 | 48-96KB/SM | ~5-20 cycles | Block内线程 |
| L1/L2缓存 | SM/芯片 | 几十KB~几MB | ~30-200 cycles | 自动管理 |
| 全局内存 | 显存 | 几GB~几十GB | ~400-800 cycles | 所有线程 |
| 常量内存 | 显存+缓存 | 64KB | ~400 cycles + 广播 | 所有线程只读 |
| 纹理内存 | 显存+特殊缓存 | 取决于显存 | ~400 cycles + 空间局部性优化 | 全局只读 |
我个人经验：全局内存访问尽量合并（coalescing）。什么意思？相邻线程访问相邻地址，数据会被打包成128字节的cache line一次性拿到。如果你乱序访问——比如线程0读地址0，线程1读地址1000——那就得两次访存，带宽直接砍半。我曾经在图像处理代码里写了个out[y * width + x]，看起来没问题，但线程ID和内存布局不匹配，导致带宽利用率只有40%。改成out[threadIdx.x + blockDim.x * blockIdx.x]后，带宽飙到95%。
核心原则：能用共享内存解决的数据局部性，就别去碰全局内存。共享内存延迟低一个数量级，而且可以手动tile（分块）。
2.4 一张图看懂GPU架构层级
下面我用SVG画了整个内存与执行单元的依赖关系。你看着它，再回看前面的文字，会更清晰。
从图上你能看到：整个任务（Grid）拆分成多个Block，每个Block被分配给一个SM。SM内部有共享内存和寄存器，Block内的线程通过共享内存交换数据。线程又组成Warp，32个一组，指令同步执行。最下面就是全局内存——数据的大本营，但访问它代价高昂。
实践经验：如果你在写kernel时发现性能卡在访存上，先检查是否可以利用共享内存。比如矩阵乘法，tile size设为16x16或32x32，把子块载入共享内存，计算时反复从共享内存读取，能快好几倍。
2.5 实战小贴士：如何查看硬件参数
想知道你的GPU有多少SM、每个SM多大共享内存？直接用CUDA API：
#include <cuda_runtime.h>
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0); // 查询设备0
printf("SM数量：%d\n", prop.multiProcessorCount);
printf("每个SM的共享内存大小：%zu bytes\n", prop.sharedMemPerMultiprocessor);
printf("每个Block最大线程数：%d\n", prop.maxThreadsPerBlock);
printf("显存总量：%.2f GB\n", prop.totalGlobalMem / 1024.0 / 1024.0 / 1024.0);我一般在写新项目前都会跑一下这个代码，确认硬件的具体数值。毕竟架构不同（比如Volta和Ampere的SM设计就有差别），调参策略也得跟着变。
小结
GPU架构不复杂，但细节很多。记住几个关键数字：Warp = 32线程，Block内线程用共享内存，全局内存访问要合并。后面章节我们会把这些概念用到实际代码里，你会看到同样的算法，改一下内存布局，性能差距有多大。