GPU架构基础：SM、Warp与内存层次结构

在解读Nsight Compute的性能指标之前，必须先理解这些指标所描述的硬件对象。GPU并非一个“巨大的并行核”，而是一组由流式多处理器（Streaming Multiprocessor，SM）构成的阵列。每个SM内部又包含多个计算单元、寄存器堆与共享内存，并按照Warp为单位调度指令。性能指标中的占用率（Occupancy）、吞吐量（Throughput）、延迟（Latency）都与这三层结构直接相关。本节将拆解SM的内部组织、Warp的执行模型以及完整的内存层次，为后续指标解读建立硬件坐标系。

2.1 SM：GPU的基本计算单元

一个GPU由多个SM组成，每个SM是独立的计算与调度实体。以NVIDIA Ampere架构的GA100为例，每个SM包含4个处理块（Processing Block），每个处理块内有32个FP32 CUDA核心、16个FP64核心、16个INT32核心以及4个Tensor Core。这些核心共享一套调度单元与寄存器堆。SM内部的关键部件如下：

• CUDA核心：执行浮点与整数运算，是吞吐量指标（如FP32 FLOPS）的物理基础。
• Tensor Core：执行矩阵乘加运算，在混合精度计算中显著提升AI工作负载的吞吐。
• 寄存器堆（Register File）：每个SM拥有64K个32位寄存器，是线程私有的最快存储。
• 共享内存（Shared Memory）：每个SM最多可配置228KB（Ampere），供块内线程共享数据。
• 调度器（Warp Scheduler）：每个处理块拥有1个调度器，负责将Warp派发到执行单元。

Nsight Compute中的sm__throughput.avg.pct_of_peak_sustained_elapsed指标，衡量的就是SM内所有执行单元（包括CUDA核心与Tensor Core）的利用率。若该值接近100%，说明SM的计算资源已被充分压满；若长期低于60%，则可能存在访存瓶颈或调度开销。

2.2 Warp：SIMT执行的基本单位

Warp是GPU调度与执行的最小单位，由32个线程组成。这32个线程在硬件上以单指令多线程（SIMT）模式运行：同一时刻，一个Warp中的所有线程执行同一条指令，但可以操作不同的数据。若线程间发生分支分歧（即部分线程走if路径，部分走else路径），Warp会串行执行两个分支，导致部分执行单元空闲。

Nsight Compute中的smsp__warp_issue_stalled_*系列指标，正是统计Warp因各种原因（如等待数据返回、等待执行单元空闲、同步屏障等）而无法发射指令的周期数。理解Warp的调度机制，有助于判断瓶颈是计算密集还是延迟密集：

指标名称	含义	典型瓶颈
smsp__average_warp_latency_per_inst	Warp执行一条指令的平均延迟（周期数）	高延迟可能来自访存或依赖链
smsp__warp_issue_stalled_long_scoreboard	因等待全局内存数据返回而停顿的周期	全局访存延迟未隐藏
smsp__warp_issue_stalled_short_scoreboard	因等待共享内存或L1命中数据而停顿	共享内存访问冲突或延迟
smsp__warp_issue_stalled_wait	因固定延迟指令（如整数除法）等待	特殊函数或高延迟算术指令
smsp__warp_issue_stalled_barrier	因块内同步屏障等待其他Warp	负载不均或同步过频

当SM中活跃Warp数量足够多时，一个Warp因访存停顿，调度器可以立即切换到另一个就绪Warp，从而隐藏延迟。这就是延迟隐藏机制。Nsight Compute中的sm__warps_active指标反映了当前活跃Warp数，而launch__occupancy_limit_warps则给出了理论最大Warp数。两者之比即为占用率。

2.3 内存层次结构：从寄存器到全局内存

GPU的内存层次从最快、最小的寄存器到最慢、最大的全局内存，形成了金字塔结构。Nsight Compute的访存指标（如l1tex__throughput、dram__throughput）分别对应不同层级。理解各层级的容量与延迟，是定位访存瓶颈的前提。

1. 寄存器（Register）：每个线程私有，访问延迟约1个周期。容量有限（每个SM 64K个32位寄存器），若每个线程使用过多寄存器，会降低占用率。
2. 共享内存（Shared Memory）：每个SM内块间共享，延迟约20-30周期。容量可配置（Ampere最大228KB），适合块内数据复用与协作。
3. L1缓存/纹理缓存：每个SM私有，用于缓存全局内存访问。与共享内存共用物理存储（可配置比例）。
4. L2缓存：GPU全局共享，容量约40MB（A100）。所有SM的全局内存访问都经过L2，是跨SM数据共享的枢纽。
5. 全局内存（DRAM）：显存，容量最大（可达80GB），延迟约400-800周期。通过PCIe或NVLink与主机通信。

Nsight Compute的memory__dram_read_bytes与l1tex__data_bank_reads等指标，分别统计了DRAM与L1层级的实际流量。若DRAM吞吐接近峰值而L1命中率低，说明数据复用性差；若L1吞吐高但DRAM吞吐低，则可能数据已在缓存中命中。

2.4 案例：从指标反推硬件行为

假设Nsight Compute报告以下关键指标：

sm__throughput.avg.pct_of_peak_sustained_elapsed : 45%
sm__warps_active.avg.pct_of_peak_sustained_active : 62%
smsp__warp_issue_stalled_long_scoreboard : 38%
smsp__warp_issue_stalled_short_scoreboard : 12%
smsp__warp_issue_stalled_barrier : 25%
l1tex__data_bank_reads.sum : 2.1e9
dram__bytes_read.sum : 8.3e8

解读步骤：

• SM吞吐仅45%，说明计算资源未饱和。
• 活跃Warp为62%，占用率中等，尚有提升空间。
• 停顿原因中，long_scoreboard占38%，表明全局内存延迟是主要停顿源；barrier占25%，说明块内同步等待也较显著。
• L1读取量（2.1e9）远大于DRAM读取量（8.3e8），说明存在一定的数据复用，但仍有大量全局访问未命中L1。

综合判断：该内核属于全局访存延迟敏感型，且存在同步负载不均。优化方向：

1. 增加每个线程的独立工作，减少同步频率（降低barrier停顿）。
2. 使用__ldg()或只读缓存修饰符，提高L1命中率。
3. 若数据复用模式允许，将热点数据显式放入共享内存，减少long_scoreboard。

2.5 误区与复盘：硬件认知偏差