GPU架构基础：SM与CUDA Core、寄存器文件结构、寄存器与共享内存的区别
好，咱们直接进入正题。这一章我打算聊聊GPU的硬件骨架——SM、CUDA Core、寄存器文件这些底层玩意儿。说实话，很多做CUDA优化的朋友，写了几年代码，对这些概念还是模模糊糊的。我个人习惯是，先搞清楚硬件长什么样，再谈优化。不然你连数据存在哪儿、怎么流动的都不知道，优化就是瞎蒙。
SM与CUDA Core：别被名字骗了
先说说SM，全称Streaming Multiprocessor，流式多处理器。你可以把它理解成GPU里的“计算小组”。一个GPU芯片里，通常有几十个甚至上百个SM。每个SM内部，才是真正的执行单元。
那CUDA Core呢？嗯，这个名字其实有点误导性。我刚开始接触时也以为它跟CPU的Core是一回事。其实不是。CUDA Core本质上是一个整数/浮点运算单元，它只能执行简单的算术逻辑指令。一个SM里通常有几十到上百个CUDA Core。
举个例子，NVIDIA的A100 GPU，每个SM有64个CUDA Core。但你别以为64个Core就能同时跑64个线程。实际上，SM是以Warp为单位调度的——32个线程一组。一个Warp里的线程，在同一时刻执行同一条指令，只是操作不同的数据。这就是SIMT（单指令多线程）的精髓。
关键点：SM是调度和执行的基本单位，CUDA Core是计算的基本单元。优化时，你要关注的是SM的利用率，而不是CUDA Core的数量。
我在项目中遇到过一个问题：一个同事把线程块（Block）设得特别大，觉得这样能充分利用CUDA Core。结果性能反而下降了。为什么？因为每个SM能同时容纳的线程块数量是有限的。Block太大，SM里能放的Block就少，调度灵活性降低，最终导致部分CUDA Core闲置。嗯，这里要注意，不是线程越多越好，是活跃的Warp越多越好。
寄存器文件结构：最宝贵的“私房钱”
接下来聊聊寄存器文件。每个SM内部都有一块寄存器文件，这是GPU里速度最快的存储资源，没有之一。它的容量是固定的，比如A100每个SM有65536个32位寄存器。
你想想看，这些寄存器要分给SM里所有的线程。每个线程能分到多少？取决于两个因素：
- 每个线程使用的寄存器数量（由编译器决定，或者你用
__launch_bounds__手动限制） - SM里同时活跃的线程总数
举个例子，假设一个SM有65536个寄存器，你每个线程用了32个寄存器。那这个SM最多能同时容纳65536 / 32 = 2048个线程。如果每个线程用64个寄存器，就只能容纳1024个线程。
这里有个常见的坑。我曾经为了减少全局内存访问，在核函数里声明了一大堆局部变量。结果编译器把这些变量都映射到了寄存器上。每个线程用了80多个寄存器！SM能同时跑的线程数直接腰斩，Warp调度器经常“断粮”，性能反而比用共享内存还差。
避坑指南：寄存器是“私有的”，每个线程独享。但它是有限的。如果你发现Occupancy（占用率）上不去，先查查是不是寄存器用太多了。我曾经用--ptxas-options=-v编译选项查看寄存器使用量，发现一个核函数用了96个寄存器，优化后降到32个，性能提升了40%。
寄存器与共享内存的区别：速度与容量
很多初学者搞不清寄存器和共享内存的区别。我简单总结一下：
| 特性 | 寄存器 | 共享内存 |
|---|---|---|
| 作用域 | 线程私有 | 线程块内共享 |
| 速度 | 最快（1个时钟周期） | 较快（约20-30个时钟周期） |
| 容量 | 每个SM固定（如64K） | 每个SM固定（如164KB） |
| 生命周期 | 线程执行期间 | 线程块执行期间 |
| 用途 | 局部变量、中间结果 | 线程间数据共享、缓存 |
说白了，寄存器是“私房钱”，只有你自己能用，但存取极快。共享内存是“家庭基金”，一个Block里的线程都能用，但速度稍慢。如果你只需要线程内部的数据，用寄存器。如果需要线程间协作，用共享内存。
我见过一个典型的错误：有人在核函数里用共享内存来存每个线程的临时计算结果，其实这些结果根本不需要共享。白白浪费了共享内存的带宽，还增加了延迟。你想想看，共享内存虽然快，但跟寄存器比还是差了一个数量级。能用寄存器解决的问题，别用共享内存。
个人经验：我写核函数时，会先估算每个线程需要的寄存器数量。如果超过32个，我会考虑是否能把一些变量放到共享内存里，或者用volatile关键字强制编译器不要过度优化。记住，寄存器溢出（Register Spill）到本地内存（Local Memory）是性能杀手，因为本地内存实际上在全局内存里，速度慢得离谱。
实战建议：如何利用这些知识做优化
好，理论讲完了，咱们来点实际的。你在写核函数时，可以按这个思路来：
- 先看Occupancy：用NVIDIA Nsight Compute或者Visual Profiler，看看你的核函数实际Occupancy是多少。如果低于50%，大概率是寄存器或共享内存用太多了。
- 检查寄存器使用量：编译时加
--ptxas-options=-v，看registers那一行的数值。如果超过32，考虑优化。 - 权衡寄存器与共享内存：如果线程间需要频繁交换数据，用共享内存。如果只是线程内部的计算，尽量用寄存器。
- 手动限制寄存器数量：用
__launch_bounds__告诉编译器你希望每个线程用多少寄存器。比如__launch_bounds__(256, 4)表示每个Block最多256个线程，每个SM至少4个Block。编译器会尽量满足这个约束。
我曾经接手过一个项目，核函数性能很差。一看，每个线程用了64个寄存器，Occupancy只有25%。我手动限制了寄存器数量到32个，Occupancy提升到50%，虽然每个线程的指令数略有增加，但整体吞吐量翻了一倍。这就是寄存器优化的威力。
嗯，这一章就到这里。记住，GPU优化的核心就是让SM尽可能多地同时处理Warp。而寄存器文件，就是决定你能同时处理多少Warp的关键瓶颈之一。下一章咱们聊聊共享内存的bank conflict，那又是一个容易踩坑的地方。