4.2 SM结构与占用率

4.2 SM（Streaming Multiprocessor）结构与资源分配（CUDA Core、Tensor Core、RT Core）

4.2 SM（Streaming Multiprocessor）结构与资源分配（CUDA Core、Tensor Core、RT Core）

在GPU计算的宏大图景中，Streaming Multiprocessor（SM）无疑是其最核心的执行单元。如果说GPU是一支由成千上万士兵组成的军团，那么SM便是其中一个个精锐作战小队；如果说CUDA编程模型是构建高性能并行算法的蓝图，那么SM则是这张蓝图得以落地的物理载体。深入理解SM的内部结构及其资源分配机制，不仅是掌握CUDA性能优化的关键，更是洞察现代异构计算架构演进逻辑的窗口。

自2006年CUDA问世以来，NVIDIA对SM架构进行了持续迭代，从早期的Tesla架构到如今的Hopper和Blackwell架构，SM不仅在数量上大幅增加，在功能单元的构成上也经历了革命性变化——CUDA Core作为通用浮点/整数运算的核心始终存在，而Tensor Core与RT Core则分别于Volta（2017）和Turing（2018）架构中被引入，标志着GPU从纯粹的通用并行处理器向AI加速器与光线追踪引擎的融合体演进。这种“三位一体”的计算单元布局，构成了当代SM最鲜明的技术特征。

SM的基本组成与调度机制

每个SM是一个高度集成的微处理器模块，内部包含多个功能子单元：寄存器文件（Register File）、共享内存（Shared Memory）、指令缓存、调度器（Warp Scheduler）、分发单元（Dispatch Unit），以及最重要的执行资源——CUDA Core、Tensor Core和RT Core（视具体架构而定）。以Ampere架构（如A100 GPU）为例，一个SM包含：

• 128个FP32 CUDA Core

• 64个INT32 Core（可与FP32交替执行）

• 4个Tensor Core（支持TF32、FP16、INT8等混合精度矩阵运算）

• 无RT Core（RT Core始于Turing，在Ampere中增强，但在数据中心卡如A100中通常不启用或不存在）

SM的执行单位并非单个线程，而是Warp——一组32个连续线程的集合。Warp是SM调度和执行的基本粒度。每个SM通常配备4个Warp调度器，每个调度器可在每个时钟周期向执行单元发射一条指令（前提是该Warp处于就绪状态且无数据依赖）。这种多调度器设计使得SM能够在不同Warp之间快速切换，有效掩盖内存访问延迟，这是GPU高吞吐量的关键所在。

值得注意的是，SM内部的资源是有限且共享的。例如，寄存器总量固定（如A100中每个SM有65536个32位寄存器），共享内存容量也有限（如164 KB，可配置为L1缓存+共享内存的组合）。当一个线程块（Thread Block）被分配到某个SM上执行时，它会占用一定数量的寄存器和共享内存。若资源不足，则无法启动更多线程块，从而限制了SM的occupancy（占用率）——即活跃Warp数量与理论最大值的比值。高occupancy虽非性能的充分条件，但通常是实现高吞吐的重要前提。

图注：Ampere架构SM内部资源与调度关系示意图。Warp调度器管理多个Warp的指令发射，执行单元按类型分工协作。

CUDA Core：通用并行计算的基石

CUDA Core是SM中最基础、最广泛的计算单元。尽管名称中带有“Core”，但它并非传统意义上的完整CPU核心，而更接近于一个标量ALU（算术逻辑单元）。在Fermi架构之后，NVIDIA将FP32和INT32运算单元分离，使得在一个时钟周期内可以同时执行一个浮点运算和一个整数运算（前提是来自不同Warp）。

在Ampere架构中，每个SM拥有128个FP32 CUDA Core，这意味着理论上每个周期可完成128次单精度浮点运算。若配合FP64单元（比例为2:1，即64个FP64 Core），则双精度性能约为单精度的1/2（消费级显卡通常更低，如1/32或1/64）。这种设计反映了GPU对高吞吐、低精度计算的偏好——这正是深度学习和图形渲染等负载的典型特征。

CUDA Core的执行遵循SIMT（Single Instruction, Multiple Thread）模型：同一Warp中的32个线程执行相同的指令，但操作各自的数据。当Warp内出现分支分歧（divergence）时，例如部分线程进入if分支而其余进入else分支，SM必须串行执行所有路径，导致性能下降。因此，良好的CUDA程序应尽量避免Warp内的控制流分歧。

从编程角度看，开发者通过编写kernel函数定义线程行为，编译器（nvcc）将其转化为PTX（Parallel Thread Execution）中间码，再经由驱动进一步编译为SASS（Surface Assembly）机器码，最终在CUDA Core上执行。这一过程对开发者透明，但理解底层执行机制有助于写出更高效的代码——例如，合理安排线程块大小以匹配SM资源，或通过coalesced memory access提升带宽利用率。

Tensor Core：面向AI的矩阵加速引擎

如果说CUDA Core是“通才”，那么Tensor Core就是专为深度学习而生的“专才”。自Volta架构起，Tensor Core被引入以加速矩阵乘加运算（Matrix Multiply-Accumulate, MMA），这是神经网络训练与推理中最耗时的操作。其基本运算形式为：

其中，A 和 B 是输入矩阵，C 是累加矩阵，D 是输出矩阵。在Ampere架构中，Tensor Core支持多种数据类型组合，例如：

• 输入为TF32（19位浮点，含8位指数、10位尾数、1位符号），累加为FP32

• 输入为FP16，累加为FP32

• 输入为INT8，累加为INT32

最关键的是，Tensor Core以4×4×4的瓦片（tile）为基本操作单元。一次MMA指令可完成两个4×4矩阵相乘，并将结果累加到另一个4×4矩阵上。由于硬件并行性极高，一个Tensor Core在一个时钟周期内可完成64次乘加运算（4×4×4=64）。Ampere SM中有4个Tensor Core，因此单SM每周期可执行256次乘加操作。

然而，要充分利用Tensor Core，开发者需满足严格的内存布局和对齐要求。传统cuBLAS库已自动支持Tensor Core加速，但对于自定义kernel，需使用WMMA（Warp Matrix Multiply Accumulate）API或更底层的mma.sync PTX指令。WMMA API要求将矩阵分块加载到Warp级别的寄存器片段（fragment）中，再调用mma_sync进行计算。这虽然增加了编程复杂度，但换来了数量级的性能提升。

以A100为例，其FP16 Tensor Core峰值性能可达312 TFLOPS，而FP32 CUDA Core仅为19.5 TFLOPS——相差16倍。这种巨大的性能鸿沟解释了为何现代AI框架（如PyTorch、TensorFlow）纷纷默认启用混合精度训练（Mixed-Precision Training）：用FP16进行前向/反向传播，用FP32维护主权重副本，既节省显存又大幅提升速度。

RT Core：实时光线追踪的硬件加速器

与Tensor Core服务于AI不同，RT Core专为实时光线追踪（Ray Tracing）而设计，首次亮相于2018年的Turing架构。其核心任务是加速光线-几何体求交测试（Ray-Triangle Intersection Test），这是光线追踪算法中最耗时的步骤。

传统上，这类计算由CUDA Core通过软件实现，效率低下。RT Core则内置专用硬件单元，可高效执行包围体层次结构（Bounding Volume Hierarchy, BVH）遍历和射线-三角形求交。具体而言，RT Core接收来自SM的射线数据和BVH节点信息，快速判断射线是否与某三角形相交，并返回最近交点或命中信息。

在SM层面，RT Core与CUDA Core、Tensor Core协同工作：CUDA Core负责生成射线、着色计算等通用任务；RT Core处理几何求交；Tensor Core可用于降噪（如DLSS技术中的AI超分辨率）。这种异构协同使得现代GPU能在游戏和专业可视化中实现实时光线追踪效果。

值得注意的是，RT Core主要存在于GeForce和Quadro/RTX系列消费级与工作站GPU中，而在数据中心GPU（如A100、H100）中通常不包含或未启用。这反映了产品定位的差异：前者强调图形与AI混合负载，后者聚焦纯计算。

资源竞争与性能权衡

SM内部的多种计算单元虽能并行工作，但并非完全独立。它们共享前端调度器、寄存器文件、共享内存和内存带宽。当一个kernel同时使用CUDA Core和Tensor Core时，两者可能竞争寄存器资源或内存带宽。例如，Tensor Core操作需要大量寄存器存储矩阵片段，若寄存器不足，可能导致寄存器溢出（spilling）到本地内存，严重拖慢性能。

此外，不同类型指令的发射也受调度器限制。尽管Warp调度器可交错发射FP32、INT32和Tensor Core指令，但若某一类指令密集且依赖性强，仍可能造成其他单元空闲。因此，负载均衡和资源感知编程成为高级CUDA优化的关键。

一个典型的例子是混合精度矩阵乘法：理想情况下，数据加载、Tensor Core计算、结果写回应形成流水线。若内存带宽成为瓶颈（如处理大矩阵时），即使Tensor Core空闲，整体性能也无法提升。此时，需通过分块（tiling）、异步数据传输（cudaMemcpyAsync）或使用统一内存（Unified Memory）来缓解。

最新进展：Hopper与Blackwell架构的演进

2022年发布的Hopper架构（如H100 GPU）对SM进行了重大升级。其SM引入了第四代Tensor Core，支持FP8数据类型，并新增Transformer Engine，可动态调整FP8与FP16的精度以优化大语言模型训练。同时，Hopper SM的寄存器文件扩大至131072个32位寄存器，共享内存增至256 KB，显著提升了大模型训练的occupancy。

更引人注目的是DPX（Dual Precision eXecution）指令的支持，允许在单条指令中同时执行FP32和FP64运算，兼顾科学计算与AI负载。此外，Hopper的SM支持异步事务屏障（Asynchronous Transaction Barrier），进一步优化了大规模分布式训练中的同步开销。

展望2024年发布的Blackwell架构，SM继续进化：每个SM包含更多Tensor Core（据传达8个），并支持稀疏计算（Sparsity） 的硬件加速——自动跳过权重矩阵中的零值，实现2倍理论性能提升。同时，Blackwell引入第五代NVLink和更大容量的共享内存，使得单SM可处理更复杂的计算图。

这些演进表明，SM正从“通用并行处理器”向“领域专用加速器集群”转变。未来的SM或许不再强调“Core”的数量，而是关注计算密度、能效比与特定负载的适配性。

结语：SM作为异构计算的微观战场

回望SM的发展历程，我们看到的不仅是一组硬件单元的堆叠，而是一种计算范式的迁移。从纯粹的SIMT执行，到融合AI、图形、科学计算的多模态加速，SM已成为现代异构计算体系中最富活力的微观战场。在这里，每一纳秒的指令调度、每一块寄存器的分配、每一次内存访问的优化，都直接影响着万亿次浮点运算的成败。

对于研究人员而言，理解SM不仅是调优代码的工具，更是洞察计算未来趋势的透镜。当摩尔定律渐趋终结，架构创新成为性能提升的主引擎，SM的每一次迭代都在回答同一个问题：如何在有限的硅面积与功耗预算下，最大化特定负载的计算效率？

答案或许没有终点，但探索永不停歇。