📚 课程目录第3章 CUDA执行模型

CUDA执行模型：线程层次结构、线程调度与线程束分化

核心原理：从线程到线程束的硬件映射

CUDA的线程组织并非扁平结构，而是遵循“网格（Grid）→ 线程块（Block）→ 线程（Thread）”的三级层次。这种设计直接服务于GPU的硬件架构：线程块是独立调度到流多处理器（SM）上的最小逻辑单元，而SM内部则以线程束（Warp）为真正的硬件执行单位。理解这一映射关系，是分析占用率与性能的起点。

线程块（Block）：由程序员指定维度（1D/2D/3D）与大小，块内线程可通过共享内存和同步屏障协作。一个块只会被分配到一个SM上，不会跨SM拆分。
线程束（Warp）：硬件将块内线程按连续线程ID每32个分为一组，形成warp。warp是SM中线程调度器（Scheduler）的调度单位，也是SIMT（单指令多线程）执行的物理载体。
网格（Grid）：由多个块组成，块在SM间的分配顺序由硬件工作分配器决定，对程序员透明，但影响缓存局部性。

一个典型映射示例：若块大小为256线程，则每个块被拆分为8个warp（256 ÷ 32 = 8）。SM中的调度器每个时钟周期选择一个就绪warp发射指令，而非单独调度线程。这意味着，warp数量直接决定了隐藏延迟的能力，而块大小则影响SM内warp的分配粒度。

线程调度机制：延迟隐藏与占用率的关系

GPU通过大量线程束的快速切换来隐藏访存延迟。当一个warp因等待全局内存数据而阻塞时，调度器立即切换到另一个就绪warp执行计算，无需保存/恢复上下文（寄存器与状态均物理驻留）。这种机制使得计算与访存重叠，是GPU高吞吐量的核心来源。

调度器在每个SM中通常有4个（取决于架构），每个调度器每周期可发射1条指令给其管理的warp。例如，在Volta架构中，每个SM有4个调度器，每个调度器管理16个warp，共64个warp槽位。要达到理想的延迟隐藏，通常需要满足：

所需warp数 ≈ 访存延迟（周期） ÷ 指令发射间隔（周期）
例如：访存延迟600周期，每周期发射1条指令 → 至少600个活动warp才能完全隐藏延迟
但SM最多64 warp → 实际只能隐藏约10%的延迟，剩余依赖指令级并行（ILP）

因此，占用率（Occupancy）并非越高越好。当warp数超过隐藏延迟所需阈值后，继续增加块只会消耗寄存器与共享内存资源，甚至导致局部性下降。调优时应先计算最小所需warp数，再反推块大小与块数量。

线程束分化：性能杀手与规避策略

线程束分化（Warp Divergence）发生在同一warp内线程执行不同分支路径时。由于SIMT架构要求warp内所有线程共享程序计数器（PC），硬件会串行执行每个分支路径，并屏蔽不活跃的线程。例如：

if (tid % 2 == 0) {
    // 偶数线程执行：16个线程活跃
    result = computeA(tid);
} else {
    // 奇数线程执行：16个线程活跃
    result = computeB(tid);
}
// 总执行周期 = computeA周期 + computeB周期（而非并行）

这意味着分化导致最多2倍性能损失（两个分支各执行一遍）。更严重的是嵌套分支，可能呈指数级放大。规避策略包括：

数据重排：将同一分支的线程聚集到同一warp（如按条件排序后分段处理）。
使用三元运算符：当分支体简单时，编译器可能生成预测指令（Predication），避免真正分支。
循环展开：减少循环内分支判断次数。
利用warp聚合：让整个warp执行相同操作，用掩码（__ballot_sync）收集结果。

案例：矩阵乘法中的块大小与分化分析

以经典的SGEMM（单精度矩阵乘法）为例，假设矩阵尺寸为1024×1024，采用朴素平铺（Tiling）算法。我们对比不同块大小对占用率和分化程度的影响：

块大小	每块warp数	SM中块数（上限）	实际占用率	分化风险
64（8×8）	2	32	64/64 = 100%	低（每warp内线程访问连续行）
256（16×16）	8	8	64/64 = 100%	低（边界处理时分化）
512（32×16）	16	4	64/64 = 100%	中（列方向非连续访问）
1024（32×32）	32	2	64/64 = 100%	高（共享内存bank冲突+分化）

实际测试中，256线程块（16×16平铺）通常表现最佳，因为：

每块8个warp，足以让4个调度器各分配2个warp，负载均衡。
共享内存平铺尺寸16×16×4字节 = 1KB，可容纳多块同时驻留。
边界条件（矩阵非倍数）导致的分化仅影响边缘warp，占比小。

而512或1024线程块虽然占用率相同，但块内同步开销增大，且共享内存分配粒度导致SM只能驻留少量块，调度灵活性下降。

误区与复盘：常见调优陷阱

误区1：盲目追求100%占用率。当kernel为计算密集型（算术指令多）时，50%占用率可能已完全隐藏延迟，继续增加warp只会增加寄存器溢出（Spill）到本地内存，反而降低性能。应通过cudaOccupancyMaxPotentialBlockSize获取理论上限，再结合实测调整。

误区2：忽略warp分化在边界条件中的影响。例如处理非对齐数据时，if (globalIdx < N)的判断在最后一个warp中可能只有部分线程活跃。若N不是32的倍数，最后一个warp的分化会造成约50%的效率损失。解决方案：填充数据到32的倍数，或使用__builtin_assume_aligned帮助编译器优化。

复盘技巧：使用NVIDIA Nsight Compute的Warp State Statistics页面，查看Stalled Warp原因分布。若Long Scoreboard占比高，说明访存延迟未隐藏，应增加块数或提高ILP；若Branch Resolving占比高，则需处理分化问题。

另一个常见误区是混淆线程块大小与warp大小。块大小不必是32的倍数，但非倍数会导致最后一个warp不满员，造成资源浪费。例如块大小100，实际占用4个warp（128线程槽位），但只有100个活跃线程，浪费28个槽位。建议块大小设为32的倍数（如128、256），同时确保块内线程数能整除总数据量，避免剩余线程空转。

最后，不要忽略寄存器分配对占用率的间接影响。编译器可能为每个线程分配较多寄存器以优化ILP，但寄存器总量固定（如每SM 65536个），若每线程使用40个寄存器，则最多驻留1638线程（约51%占用率）。此时应通过__launch_bounds__(256, 8)限制每块线程数并提示编译器控制寄存器用量，换取更高块驻留数。