GPU架构基础：核心架构与内存层次
大家好，我是你们的老朋友。今天咱们聊聊GPU的骨架——它的核心架构和内存系统。说实话，我刚开始接触GPU编程时，觉得这东西就是个黑盒子。后来踩了不少坑，才慢慢摸清楚里面的门道。
你想想看，CPU和GPU最大的区别在哪？CPU是「精打细算型」——几个强壮的核，每个都能处理复杂任务。GPU呢？它是「人海战术型」——成千上万个简单核，专为并行而生。我习惯把GPU想象成一个超级工厂，里面全是流水线工人。
SM：GPU的基本计算单元
SM，全称Streaming Multiprocessor，流式多处理器。它是GPU里最核心的模块。每个SM就像一个小型计算中心，里面包含了CUDA Core、Tensor Core、共享内存、寄存器文件等资源。
我在项目中遇到过一个问题：同样的算法，在A卡上跑得飞快，换到B卡就慢得像蜗牛。后来一查，原来是SM数量不同导致的。嗯，这里要注意——SM的数量直接决定了GPU的并行能力。
关键点：SM是独立执行指令的单元。每个SM有自己的指令调度器、执行单元和缓存。GPU的并行度，说白了就是SM数量 × 每个SM能同时处理的线程数。
CUDA Core：最基础的执行单元
CUDA Core，也叫流处理器。它是SM里最基础的算术逻辑单元。每个CUDA Core可以执行整数运算和浮点运算。我刚开始写CUDA程序时，总以为CUDA Core越多越好。其实不然——利用率才是关键。
举个例子：一块RTX 3090有82个SM，每个SM有128个CUDA Core。总共10496个CUDA Core。听起来很吓人对吧？但如果你写的程序让这些核大部分时间都在空转，那再多的核也没用。
我的习惯：写代码前，先算算你的算法需要多少线程。别一股脑全扔给GPU。我曾经写过一段代码，开了100万个线程，结果每个线程只干了一点点活。调度开销比计算本身还大，得不偿失。
Tensor Core：AI时代的加速器
Tensor Core是NVIDIA从Volta架构开始引入的专用单元。它专门为矩阵乘法设计。你想想看，深度学习里最耗时的操作是什么？没错，就是矩阵乘法。Tensor Core可以在一个时钟周期内完成4×4矩阵的乘加运算。
我记得第一次用Tensor Core时，那感觉就像开了挂。同样的矩阵乘法，用CUDA Core要跑10毫秒，换成Tensor Core只要1毫秒。但要注意——Tensor Core对数据精度有要求。它主要支持FP16、BF16、INT8等低精度格式。
避坑指南：我曾经在项目里直接用了FP32数据去调用Tensor Core，结果发现根本没加速。后来才意识到，Tensor Core需要特定的数据格式和内存对齐。使用前一定要检查你的数据精度和布局。
内存层次结构：数据流动的命脉
GPU的内存系统，说白了就是「越靠近计算单元，速度越快，容量越小」。我习惯把它想象成三级火箭：
| 内存类型 | 速度 | 容量 | 作用域 |
|---|---|---|---|
| 寄存器 | 最快（1个时钟周期） | 极小（每个线程几十到几百字节） | 单个线程 |
| 共享内存 | 快（几个时钟周期） | 小（每个SM几十KB） | 同一个线程块 |
| 全局内存 | 慢（几百个时钟周期） | 大（几GB到几十GB） | 所有线程 |
寄存器：线程的私有空间
寄存器是GPU里最快的存储。每个线程都有自己的寄存器文件。我刚开始写CUDA时，总喜欢定义很多局部变量。结果呢？寄存器不够用，数据被「溢出」到本地内存，速度直接掉了一个数量级。
为什么会这样？因为每个SM的寄存器总数是有限的。比如一个SM有65536个寄存器，如果你每个线程用64个，那这个SM最多能同时处理1024个线程。如果你每个线程用128个，那只能处理512个了。嗯，这里要权衡——寄存器用多了，并行度就下降了。
我的建议：写kernel时，尽量控制局部变量的数量。能用int就别用double。我一般会先用__launch_bounds__指定线程块大小，然后看编译器报告，再调整寄存器使用量。
共享内存：线程块的协作空间
共享内存是SM内部的一块高速缓存。同一个线程块里的所有线程都能访问它。它的速度比全局内存快得多，但容量有限——通常每个SM只有几十KB到一百多KB。
我在项目中遇到过这样一个场景：做矩阵转置时，直接读写全局内存，带宽利用率只有20%。后来改成用共享内存做中转，利用率直接飙到80%以上。说白了，共享内存就是用来做数据复用的。
关键技巧：使用共享内存时，一定要注意bank conflict。共享内存被分成32个bank，如果多个线程同时访问同一个bank，就会产生冲突，导致串行化。我习惯在声明共享内存数组时，加一个padding来避免这种情况。
全局内存：大容量但慢速
全局内存是GPU里最大的存储空间。所有线程都能访问它，但速度也是最慢的。访问一次全局内存，大概需要几百个时钟周期。你想想看，如果每个线程都频繁访问全局内存，那性能肯定好不了。
我记得有一次优化一个图像处理算法，发现大部分时间都花在全局内存访问上。后来用了合并访问的技巧——让连续的线程访问连续的内存地址——带宽利用率从30%提升到了90%。
避坑指南：我曾经犯过一个错误——在循环里反复访问全局内存。每次循环都去读一次，结果性能惨不忍睹。正确的做法是：先把数据从全局内存读到共享内存或寄存器，处理完再写回去。记住：减少全局内存访问次数，比优化单次访问速度更重要。
Warp与线程束调度
Warp是GPU执行指令的基本单位。在NVIDIA的架构里，32个线程组成一个Warp。SM以Warp为单位来调度和执行指令。你想想看，如果Warp里的32个线程都执行同样的指令，那效率最高。但如果它们分支了——有的走if，有的走else——那就麻烦了。
为什么会这样？因为SM的指令调度器一次只能给一个Warp发一条指令。如果Warp里的线程执行不同路径，那这些路径就得串行执行。这就是所谓的Warp divergence（线程束分化）。
核心原则：尽量让同一个Warp里的线程走相同的分支。我习惯在写条件判断时，先想想这个条件会不会导致Warp分化。如果会，就尽量用数据预处理或重新组织线程的方式来避免。
举个例子：假设你要处理一个数组，只处理其中大于0的元素。如果数据是随机分布的，那每个Warp里可能既有大于0的，也有小于等于0的。这时候就会产生分化。我一般会先对数据做一次排序或分组，让同类数据集中在一起。
我的经验：Warp调度还有一个容易被忽略的点——占用率。占用率是指每个SM上活跃的Warp数量。占用率越高，SM越能通过切换Warp来隐藏内存访问延迟。但占用率不是越高越好，太高了可能会挤占寄存器或共享内存。我一般会先用profiler跑一下，看看占用率和性能的关系曲线。
好了，这一章的内容就到这里。GPU架构这块，说白了就是理解「并行」和「层次」这两个词。下一章我们会聊聊如何用这些知识来写高效的CUDA代码。到时候我会分享一些实战中的优化技巧，保证让你少走弯路。
本章小结：
- SM是GPU的基本计算单元，SM数量决定并行能力
- CUDA Core执行基本运算，Tensor Core专为矩阵乘法加速
- 内存层次：寄存器最快，共享内存次之，全局内存最慢
- Warp是调度单位，避免Warp分化是性能关键
- 占用率要适度，过高或过低都会影响性能
公众号:蓝海资料掘金营，微信deep3321