一、SM内部结构全景：GPU计算的核心引擎

大家好，我是你们的老朋友。今天咱们来聊聊SM——Streaming Multiprocessor。说白了，SM就是GPU里真正干活的单元。你想想看，一块GPU有几千个核心，但这些核心不是独立工作的，它们被组织成一个个SM。每个SM就像一个小型计算中心，有自己的调度器、寄存器和内存。

我个人习惯把SM比作一个「微型多核处理器」。为什么这么说？因为它内部包含了完整的计算资源：

  • 计算单元：CUDA核心（也叫SP）、特殊函数单元（SFU）、加载存储单元（LD/ST）
  • 存储资源：寄存器文件、共享内存、一级缓存、常量缓存
  • 控制单元：Warp调度器、指令分发单元

我在项目中遇到过不少同学，一上来就盯着CUDA核心数量看，觉得核心越多性能越好。嗯，这里要注意——SM的内部结构设计，其实决定了你的程序能不能把硬件喂饱。

1.1 SM组成单元：各司其职的团队

一个典型的SM（以NVIDIA Turing架构为例）包含以下关键单元：

单元名称 数量 主要职责
CUDA Core 64个 执行整数和单精度浮点运算
Tensor Core 8个 执行矩阵乘加运算（混合精度）
SFU 4个 执行三角函数、倒数、平方根等特殊函数
LD/ST Unit 16个 处理全局内存和共享内存的加载/存储
Warp Scheduler 4个 管理Warp的调度和指令发射
Register File 65536个 提供最快的寄存器访问
Shared Memory 96KB 线程块内共享数据

这些单元不是独立工作的。举个例子，当你的kernel执行一条浮点乘法指令时，Warp调度器会从寄存器文件读取操作数，送到CUDA Core计算，结果再写回寄存器。整个过程流水线化，每个时钟周期都能发射指令。

关键认知：SM内部的资源是有限的。一个SM最多能同时驻留多少个线程？这取决于寄存器文件和共享内存的容量。我曾经调试过一个程序，发现occupancy只有25%，查了半天——原来是每个线程用了太多寄存器，导致SM能容纳的线程块数量骤减。

1.2 寄存器文件：最快的存储，也是最稀缺的资源

寄存器文件是SM内部最宝贵的资源。为什么？因为它离计算单元最近，访问延迟只有1个时钟周期。相比之下，全局内存的延迟是几百个时钟周期。

每个SM的寄存器文件大小是固定的。以Volta架构为例，每个SM有65536个32位寄存器。这些寄存器被分配给驻留在SM上的所有线程。假设你的kernel每个线程用了32个寄存器，那么一个SM最多能容纳65536 / 32 = 2048个线程。

避坑指南：我曾经见过一个项目，团队为了优化计算精度，在kernel里用了大量double类型变量。结果每个线程的寄存器使用量从32飙升到64，occupancy直接腰斩。性能不但没提升，反而下降了30%。记住：寄存器不是免费的，用多了会挤占线程并行度。

怎么查看寄存器使用量？用nvcc编译时加--ptxas-options=-v参数：

nvcc --ptxas-options=-v mykernel.cu

// 输出示例：
// ptxas info: Used 32 registers, 8192 bytes smem, 48 bytes cmem[0]

我个人习惯在优化kernel时，先看寄存器使用量。如果超过32，我会考虑：

  • 能不能把一些局部变量改成共享内存？
  • 能不能拆分成多个小kernel？
  • 能不能用__launch_bounds__限制每个线程块的最大线程数？

1.3 共享内存：可编程的片上缓存

共享内存是SM内部的一块可编程SRAM。它和L1缓存共享同一块物理存储，你可以通过编译选项配置两者的比例。比如96KB共享内存 + 32KB L1，或者64KB + 64KB。

共享内存的访问延迟大约是20-30个时钟周期，比寄存器慢，但比全局内存快得多。它的最大优势是：同一个线程块内的所有线程都能访问同一块共享内存。

来看一个典型的共享内存使用场景——矩阵转置：

__global__ void transpose_shared(float *in, float *out, int width) {
    __shared__ float tile[TILE_SIZE][TILE_SIZE];
    
    int x = blockIdx.x * TILE_SIZE + threadIdx.x;
    int y = blockIdx.y * TILE_SIZE + threadIdx.y;
    
    // 协作加载到共享内存
    tile[threadIdx.y][threadIdx.x] = in[y * width + x];
    __syncthreads();
    
    // 从共享内存转置写出
    out[x * width + y] = tile[threadIdx.x][threadIdx.y];
}

这里有个细节：__syncthreads()是必须的。为什么？因为要保证所有线程都写完了共享内存，才能开始读。我曾经犯过一个错误——忘了加同步，结果读到的全是脏数据，调试了整整一个下午。

经验之谈：共享内存虽然快，但要注意bank conflict。共享内存被分成32个bank，如果多个线程同时访问同一个bank的不同地址，就会发生冲突，导致串行化。解决办法？数据填充（padding）或者重新组织访问模式。

1.4 调度单元：Warp调度器的工作原理

SM的调度单元是Warp调度器。一个Warp包含32个线程，它们执行相同的指令（SIMT模型）。Warp调度器负责从多个Warp中选择一个，发射下一条指令。

为什么需要多个Warp调度器？因为单个调度器可能无法喂饱所有的计算单元。以Turing架构为例，每个SM有4个Warp调度器，每个调度器每个时钟周期可以发射一条指令。这意味着每个时钟周期最多能发射4条指令。

调度器的核心策略是「零开销上下文切换」。当一个Warp因为访存延迟而等待时，调度器会立即切换到另一个就绪的Warp。这个过程不需要保存和恢复寄存器——因为每个Warp有自己的寄存器文件。

核心概念：隐藏延迟。GPU通过大量线程的并行执行来隐藏内存访问延迟。只要SM上有足够多的Warp，调度器就能在等待访存的同时执行其他Warp的指令。这就是为什么occupancy很重要——更多的Warp意味着更好的延迟隐藏能力。

我建议你在优化时关注两个指标：

  • Occupancy：活跃Warp数量 / 最大Warp数量。一般建议达到50%以上。
  • Warp Stall原因：用NVIDIA Nsight Compute分析，看看Warp是因为什么原因停滞的。是访存？是计算依赖？还是同步等待？

举个例子，我曾经优化一个流体力学模拟程序。初始occupancy只有33%，Warp大部分时间在等待全局内存。我把一些频繁访问的数据搬到了共享内存，occupancy提升到67%，性能翻了一倍。

1.5 知识体系总览

下面这张图展示了SM内部的核心结构和工作流程：

Streaming Multiprocessor (SM) Warp Scheduler 0 指令发射 Warp Scheduler 1 指令发射 Warp Scheduler 2 指令发射 Warp Scheduler 3 指令发射 计算单元 (Compute Units) CUDA Core x64 整数/浮点运算 Tensor Core x8 矩阵乘加 SFU x4 特殊函数 LD/ST Unit x16 加载/存储 存储资源 (Memory Resources) 寄存器文件 65536个32位寄存器 延迟：1 cycle 共享内存 96KB（可配置） 延迟：20-30 cycles L1/常量缓存 与共享内存共享 延迟：30-40 cycles 调度单元 计算单元 存储资源

从这张图可以看出，SM内部是一个高度并行的流水线结构。Warp调度器从寄存器文件读取指令和操作数，送到计算单元执行，结果再写回寄存器或共享内存。整个过程环环相扣，任何一个环节成为瓶颈，都会拖慢整体性能。

我的建议：刚开始学习SM结构时，不要试图一次性记住所有细节。先理解三个核心概念：Warp调度、寄存器文件、共享内存。这三个是GPU性能优化的基石。等你熟练了，再深入Tensor Core、SFU这些高级特性。

好了，这一章的内容就到这里。记住：SM是GPU的心脏，理解它的内部结构，你才能写出真正高效的CUDA程序。下一章我们会深入Warp调度和指令发射的细节，到时候见。

想要更多资料，联系可可

← 上一章 📖 返回目录 下一章 →