GPU架构基础：CUDA核心与SM、显存层次结构
各位同学，今天我们来聊聊GPU的“骨架”和“血管”。说白了，你要做并发推理调度，不懂GPU的硬件结构，就像开车不看仪表盘——迟早要翻车。我个人习惯，拿到一块新GPU，第一件事不是跑benchmark，而是先看它的SM数量和显存带宽。
CUDA核心与SM：真正的“工人”与“车间”
很多人以为CUDA核心就是一个个独立的计算单元，其实没那么简单。我刚开始做GPU优化时也犯过这个错。
SM（Streaming Multiprocessor）才是真正的调度单位。你可以把SM想象成一个车间，CUDA核心就是车间里的工人。一个SM里通常有几十到上百个CUDA核心，但它们共享指令缓存、共享内存和调度器。
关键点：CUDA核心不是独立工作的，它们以“warp”（32个线程为一组）为单位执行。同一个warp里的线程执行相同的指令，只是处理不同的数据。这就是SIMT（单指令多线程）的精髓。
我在项目中遇到过一个问题：某个kernel的occupancy（占用率）一直上不去。查了半天，发现是每个线程使用的寄存器太多，导致SM能同时容纳的warp数量减少。嗯，这里要注意——寄存器是SM的共享资源，不是越多越好。
| GPU型号 | SM数量 | 每个SM的CUDA核心数 | 最大warp数/SM |
|---|---|---|---|
| A100 | 108 | 64 | 64 |
| V100 | 80 | 64 | 64 |
| RTX 3090 | 82 | 128 | 48 |
你想想看，为什么A100的SM数量比RTX 3090多，但每个SM的CUDA核心反而少？因为A100的设计更注重计算密度和能效比，而不是单纯堆核心数。我建议你在选型时，不要只看“多少TFlops”，要看SM的架构设计。
显存层次结构：HBM、L2、L1
GPU的显存层次，说白了就是“速度越快的越贵，容量越小的越近”。我经常跟团队说：优化GPU程序，本质上就是在管理数据在各级存储间的移动。
HBM（高带宽显存）
HBM是GPU的主存，容量大（几十GB），但延迟高（几百纳秒）。A100用的是HBM2e，带宽高达2TB/s。听起来很吓人对吧？但你要知道，如果所有数据都从HBM读，再快的带宽也扛不住。
我的经验：在做推理调度时，尽量把模型参数常驻在显存中，不要频繁换入换出。我曾经见过一个团队，每次推理都从CPU内存拷贝模型权重到GPU，结果带宽全浪费在数据传输上了。
L2缓存
L2缓存是所有SM共享的，容量通常在几十MB。它的延迟比HBM低一个数量级（几十纳秒）。我习惯把L2看作“短期记忆”——频繁访问的数据应该尽量留在L2里。
为什么会这样？因为L2的带宽是HBM的几倍。举个例子，A100的L2带宽约4TB/s，而HBM只有2TB/s。如果你的数据访问模式有局部性，L2能帮你省下大量时间。
L1缓存/共享内存
L1是每个SM私有的，容量最小（通常几十KB到几百KB），但速度最快（几个时钟周期）。这里有个坑：L1和共享内存其实是同一块硬件，你可以通过配置来分配它们的比例。
我曾经踩过的坑：在写一个矩阵乘法的kernel时，我把所有数据都塞进共享内存，结果因为容量不够，只能处理很小的矩阵。后来改成“分块+流水线”的方式，用L1缓存做数据重用，性能提升了3倍。记住：共享内存不是越大越好，要跟你的算法匹配。
内存带宽与延迟：GPU优化的核心矛盾
带宽和延迟，是GPU优化的两个核心指标。但很多人只盯着带宽看，忽略了延迟的影响。
- 带宽：单位时间内能传输的数据量。HBM的带宽很高，但如果你访问模式是随机的，实际有效带宽可能只有理论值的10%。
- 延迟：从发出请求到收到数据的时间。HBM的延迟约400-800个时钟周期，L2约100-200个，L1约20-30个。
我建议你记住一个公式：有效带宽 = 理论带宽 × 利用率。利用率取决于你的访问模式是否连续、是否对齐、是否有数据重用。
举个例子，我优化过一个BERT推理的kernel。原始版本直接从HBM读取attention矩阵，带宽利用率只有15%。后来我把Q、K、V矩阵分块加载到共享内存，利用率提升到了70%，推理速度翻了4倍。
// 优化前：直接从HBM读取
float q = Q[row * N + col]; // 每次访问都走HBM，延迟高
// 优化后：分块加载到共享内存
__shared__ float Q_shared[BLOCK_SIZE][BLOCK_SIZE];
Q_shared[threadIdx.y][threadIdx.x] = Q[row * N + col];
__syncthreads();
// 后续从共享内存读取，延迟降低一个数量级核心原则：计算是免费的，数据移动是昂贵的。GPU的算力远超内存带宽，所以瓶颈往往在数据搬运上。你做并发推理调度时，要优先考虑如何减少数据移动，而不是如何提高计算速度。
知识体系总览
下面这张图是我自己画的，把GPU架构的核心概念串起来了。你仔细看，会发现所有优化技巧都围绕“数据在哪、怎么搬、搬多快”这三个问题。
这张图里，从上到下是“整体→组件→指标→优化目标”的递进关系。你每次做优化时，都可以对照这张图，看看自己卡在了哪个环节。
我的建议：刚开始学GPU优化时，不要贪多。先把SM和显存层次搞透，再去看具体的优化技巧。我见过太多人一上来就学cuDNN、TensorRT，结果连warp divergence是什么都不知道，最后调优全靠试错。
好了，这一章的内容就到这里。记住：GPU架构不是死记硬背的参数，而是你做并发推理调度时的“地图”。知道数据在哪、怎么搬、搬多快，你就能写出高效的代码。
想要更多资料，联系可可