GPU架构基础：SM、SP、线程束、线程块、网格等核心概念

要理解CUDA占用率调优，首先必须建立对GPU硬件执行模型与软件线程层次结构的清晰认知。CUDA的编程模型之所以能高效运行，是因为它精确地映射了NVIDIA GPU的物理架构。本节将剖析从硬件执行单元（SM、SP）到软件抽象层次（线程、线程块、网格）的核心概念，并阐明它们之间的对应关系。

核心硬件单元：SM与SP

GPU由多个流式多处理器（Streaming Multiprocessor，SM）构成，每个SM内部包含多个流式处理器（Streaming Processor，SP），即CUDA核心。SP是执行浮点/整数运算的最小单元，而SM是独立的调度和执行单位。以NVIDIA Ampere架构的GA102芯片为例，其包含7个GPC（图形处理簇），每个GPC内含多个TPC（纹理处理簇），每个TPC包含2个SM，每个SM包含128个CUDA核心（SP）。

SM内部还集成了共享内存（Shared Memory）、寄存器文件（Register File）、特殊功能单元（SFU）、加载/存储单元（LD/ST）以及线程束调度器（Warp Scheduler）。这些资源共同决定了SM能够同时容纳的线程数量上限，也就是占用率（Occupancy）的物理基础。

软件线程层次：网格、线程块与线程

CUDA的软件线程模型是三层嵌套结构，与硬件执行模型形成映射关系：

• 网格（Grid）：一次内核（Kernel）启动所创建的所有线程的集合。网格由多个线程块组成，对应整个GPU设备（Device）上的所有SM资源池。
• 线程块（Thread Block）：网格中的子集，由多个线程组成。线程块是软件与硬件交互的关键单位——一个线程块会被调度到单个SM上执行，且块内线程可通过共享内存和同步屏障（__syncthreads()）协作。
• 线程（Thread）：执行内核函数的最小单位，每个线程拥有独立的程序计数器、寄存器状态和局部变量。

线程块与SM的绑定关系是占用率调优的核心杠杆。一个SM可以同时驻留多个线程块（受资源限制），但一个线程块不能拆分到多个SM上。这种设计保证了块内通信的低延迟，但也限制了灵活性。

线程束（Warp）：硬件调度的真实单位

尽管编程模型以线程为单位，但GPU硬件实际执行和调度的单位是线程束（Warp）。线程束是32个线程的集合，这32个线程在同一个SM内被锁步（Lockstep）执行——即它们在同一时刻执行同一条指令（SIMT模型，单指令多线程）。

线程束的划分是连续的：threadIdx.x 0~31组成warp 0，32~63组成warp 1，以此类推。当线程块维度不是32的倍数时，最后一个线程束会包含无效线程（仍占用调度槽位），造成资源浪费。因此，线程块大小通常设置为32的倍数（如128、256、512），以最大化线程束利用率。

线程束调度器在每个时钟周期选择一个就绪的线程束发射指令。当某个线程束因访存延迟或同步等待而阻塞时，调度器会切换到其他就绪的线程束，通过线程级并行（TLP）隐藏延迟。这就是为什么高占用率能提升性能——更多的线程束意味着更多的“备选任务”来填补流水线气泡。

SM资源与线程块驻留的约束关系

每个SM能同时驻留多少个线程块，取决于以下四个硬件资源限制：

例如，若内核每个线程使用40个寄存器，则每个SM最多容纳 65536 ÷ 40 = 1638 个线程，但受1536线程上限约束，实际最大驻留1536线程。若每个线程块为256线程，则最多驻留 1536 ÷ 256 = 6 个块。此时占用率为 1536 ÷ 1536 = 100%。

案例：线程块大小对占用率的影响

假设一个内核每个线程使用64个寄存器，SM寄存器总量为65536，最大线程数为1536。我们来对比不同线程块大小的效果：

// 内核声明示例
__global__ void example_kernel(float* data, int N) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < N) {
        // 使用约64个寄存器的计算逻辑
        float a = data[idx] * 2.0f;
        // ... 更多局部变量
    }
}

// 启动配置
int blockSize = 256;  // 可调参数
int gridSize = (N + blockSize - 1) / blockSize;
example_kernel<<<gridSize, blockSize>>>(d_data, N);

计算过程：

• 寄存器限制：65536 ÷ 64 = 1024 线程/SM
• 线程数上限：1536线程/SM
• 实际最大驻留线程数 = min(1024, 1536) = 1024
• 若blockSize=256，则驻留 1024 ÷ 256 = 4 个块，占用率 = 1024 ÷ 1536 ≈ 66.7%
• 若blockSize=128，则驻留 1024 ÷ 128 = 8 个块，占用率仍为66.7%（受寄存器限制）
• 若blockSize=512，则驻留 1024 ÷ 512 = 2 个块，占用率仍为66.7%

可见，当寄存器成为瓶颈时，调整线程块大小无法提升占用率，必须减少每线程寄存器使用量（如使用 __launch_bounds__ 或重构代码）。

误区与复盘：常见概念混淆