GPU架构与内核执行基础：SM架构、线程层次与内存层次

在进入NVIDIA Nsight Compute（以下简称Ncu）的性能分析之前，必须首先建立对GPU底层执行模型的清晰认知。Ncu的所有采样指标、吞吐率计算、瓶颈判定逻辑，本质上都是围绕“SM（Streaming Multiprocessor，流式多处理器）如何调度线程、线程如何访问内存”这两个核心问题展开的。如果对硬件执行模型理解不透彻，分析报告中的“SM Busy”“Warp Stall”“DRAM Throughput”等指标就只是一堆无意义的数字。本章将系统拆解GPU的物理执行单元（SM）、逻辑线程层次（Grid/Block/Warp）以及完整的内存层次结构，并说明它们如何共同决定内核的实际运行效率。

2.1 SM架构：GPU的并行计算核心

SM（Streaming Multiprocessor）是GPU中最基本的物理计算单元。以NVIDIA Ampere架构的A100为例，每个A100 GPU包含108个SM，而每个SM内部又划分为4个处理块（SM子分区）。每个处理块拥有32个FP32 CUDA核心、16个FP64单元、16个INT32单元以及4个Tensor Core（第三代）。这种设计使得一个SM能够同时执行来自不同Warp的指令，实现指令级并行与线程级并行的叠加。

理解SM架构的关键在于“Warp调度”机制。每个SM子分区拥有独立的Warp调度器和指令分发单元，每个周期可以发射一条指令到对应的执行单元。当某个Warp因等待内存访问结果而暂停时，调度器会立即切换到另一个就绪的Warp，这种机制称为“零开销上下文切换”。正是这种设计，使得GPU能够用“大量线程并行”来掩盖“单线程延迟”，从而在高延迟的内存访问场景下依然保持较高的吞吐率。

在Ncu的Source页中，Executed Ipc Active指标反映了SM子分区每个周期的指令发射数。如果该值长期低于理论峰值（如Ampere的4），说明Warp调度器存在空闲周期，可能是由于Warp阻塞（等待内存、同步）或指令依赖导致。此时应结合Warp Stall详情页判断具体阻塞原因。

2.2 线程层次：Grid、Block与Warp的协作逻辑

CUDA的线程层次是程序员直接控制的逻辑结构，它决定了数据如何被划分、线程如何协作。内核启动时，程序员指定Grid维度（即Block的排列方式）和Block维度（即线程的排列方式）。例如，kernel<<<gridDim, blockDim>>>()中，gridDim和blockDim可以是三维向量，但实际使用中通常只使用x维度。

Block是线程调度的基本单位。GPU的GigaThread引擎会将Block分配到各个SM上，一个SM可以同时驻留多个Block（取决于资源限制）。Block内的线程通过共享内存和__syncthreads()进行协作。Warp是硬件调度的最小单位，一个Warp包含32个连续的线程ID。当Block大小为128时，它会被拆分为4个Warp，每个Warp独立调度执行。

理解线程层次对性能优化至关重要，因为它直接决定了两个关键指标：Occupancy（占用率）和Block调度效率。Occupancy定义为“驻留在SM上的Warp数 / SM最大支持Warp数”。较高的Occupancy有助于隐藏内存延迟，但过高的Occupancy也可能导致寄存器溢出或缓存抖动。

在Ncu的Launch Statistics页面中，可以查看Block Limit（受限于寄存器、共享内存或Warp槽位）和Occupancy。如果发现Block Limit Registers为限制因素，说明每个线程使用的寄存器过多，可以通过__launch_bounds__或maxrregcount编译选项限制寄存器使用量。

// 示例：使用 __launch_bounds__ 限制寄存器使用
__global__ void __launch_bounds__(256, 8) myKernel(float* data) {
    // 最多256线程/Block，至少8个Block/SM
    // 编译器将据此限制寄存器分配，确保Occupancy
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    // ... 内核逻辑
}

2.3 内存层次：从寄存器到全局内存的完整路径

GPU的内存层次设计遵循“容量小、速度快”与“容量大、速度慢”的权衡原则。从线程视角看，内存分为以下几个层级：

• 寄存器文件（Register File）：每个线程私有的最快存储，访问延迟约1个周期。SM的寄存器文件总容量固定（如A100为256KB），被所有驻留线程瓜分。若每个线程使用超过64个寄存器，将显著降低Occupancy。
• 共享内存（Shared Memory）：Block内所有线程共享，访问延迟约20-30周期。共享内存与L1缓存共享物理存储（统一架构），可通过cudaFuncSetAttribute动态调整比例。共享内存支持Bank并行访问，但存在Bank Conflict问题。
• 局部内存（Local Memory）：当寄存器溢出或访问动态索引的局部数组时，数据被放置在局部内存中。局部内存物理上位于DRAM，但会经过L1/L2缓存，延迟远高于寄存器。
• L2缓存：所有SM共享的片上缓存，容量约40-50MB（A100）。L2是全局内存访问的必经之路，也是跨SM数据共享的枢纽。
• 全局内存（Global Memory）：即显存DRAM（HBM2e或GDDR6），容量大但延迟高达400-800周期。全局内存的访问效率高度依赖“合并访问”（Coalescing）模式。

在Ncu的Memory Workload Analysis页面中，可以查看Memory Throughput、L1 Hit Rate、L2 Hit Rate以及DRAM Throughput。这些指标帮助判断内核是“计算密集型”还是“内存密集型”。如果DRAM吞吐率接近峰值（如A100的2TB/s），说明内核受内存带宽限制；如果L2命中率低且DRAM吞吐率不高，则可能存在访问模式问题。

2.4 执行模型：Warp调度与延迟隐藏

GPU的执行模型可以概括为“大量Warp并行 + 零开销切换”。每个SM子分区拥有一个Warp调度器，每个周期从就绪Warp中选择一个发射指令。所谓“就绪”，是指该Warp的指令操作数已准备就绪（寄存器已写入、内存访问已完成）。当Warp因等待内存数据而阻塞时，调度器切换到其他Warp，从而隐藏内存延迟。

延迟隐藏的效率取决于两个因素：Occupancy和指令级并行（ILP）。高Occupancy意味着有更多Warp可供调度，但每个Warp的指令级并行也能减少调度压力。例如，一个Warp内连续多条独立指令可以连续发射，无需等待。Ncu的Warp State页面会显示Warp处于Active、Stalled、Waiting等状态的比例，帮助定位阻塞原因。

常见的Warp阻塞原因包括：Long Scoreboard（等待全局内存）、Short Scoreboard（等待共享内存）、Barrier（等待同步）、MIO Throttle（内存输入输出队列满）、Not Selected（调度器选择了其他Warp）。在Ncu的Warp Stall详情页中，可以按阻塞原因排序，找出最突出的瓶颈。

在实际优化中，需要根据Ncu的指标动态调整策略。如果MIO Throttle占比高，说明共享内存或L1的访问队列已满，应减少共享内存访问频率或优化访问模式；如果Barrier占比高，说明Block内同步过于频繁，应减少__syncthreads()调用或调整Block大小。