刚上手 CUDA 的人几乎都会撞上同一个困惑明明按文档把 block 开到了 256 线程、grid 铺满了整卡跑出来的算力却只有理论峰值的三成。更让人抓狂的是把线程数再翻一倍速度纹丝不动把寄存器限制一下、把 occupancy 拉满反而更慢了。这些现象背后指向同一个东西——SM 里的 warp 调度器和它赖以生存的访存延迟隐藏机制。这两个概念是 CUDA 性能模型的地基理解了它们你再看 Nsight Compute 里那一堆smsp__warp_issue_stalled_*指标才不会发懵调优也才有方向而不是靠多开线程加__restrict__这种碰运气的做法。下面我把这套机制从硬件形态一直讲到实测数据顺带把我在真实项目里踩过的几个反直觉的坑一并交代清楚无论你是刚开始写 kernel 还是已经能看懂 profiling 报告应该都能从中拿到能直接用的东西。1. SM 到底在藏什么延迟先算清这笔账1.1 一条访存指令的延迟有多离谱要理解调度器为什么存在先得知道它要对付的敌人有多大。现代 GPU 上一条普通的全局内存加载指令从发射到数据真正回到寄存器中间要穿过 L1、L2、内存控制器、DRAM 颗粒再原路返回。这条链路在不同层级上的耗时量级大致是这样的数值随架构和频率浮动这里给的是量级不是精确值访问目标典型延迟SM 周期说明寄存器操作数0操作数直通不算延迟FP32 ALU 结果依赖4~6背靠背依赖链的最小间隔共享内存 load20~30走 MIO 队列short scoreboard 管辖L1 命中30~40全局访存但命中片上缓存L2 命中约 200跨 SM 共享走片上网状互连设备内存 DRAM400~800真正去颗粒取数注意最后一行DRAM 的延迟大约是 ALU 依赖延迟的一百倍。这意味着如果 SM 傻乎乎地发射一条访存指令然后原地等数据回来再发射下一条那它 99% 的时间都在发呆。CPU 的做法是堆大缓存 乱序执行 深流水线把这种等待尽量抹平GPU 的做法完全相反——它不消除延迟它用海量并发把延迟填满。这就是延迟隐藏latency hiding四个字的全部含义SM 上常驻几十个 warp某个 warp 卡在等内存的时候调度器立刻切到另一个已经就绪的 warp 去发射指令。流水线永远不空代价是每个 warp 自己看起来都很慢。1.2 用吞吐反推需要多少在飞行的数据光说要用并发填是空话我们可以用最朴素的利特尔法则Littles Law把它算成一个具体数字这样后面判断到底缺不缺 warp就有了标尺。假设目标是打满 A100 的显存带宽约 1555 GB/sDRAM 往返延迟按 600 ns 算那么任意时刻必须处于已发出、未返回状态的数据量是在飞字节数 带宽 × 延迟 ≈ 1555e9 B/s × 600e-9 s ≈ 933,000 B ≈ 0.93 MBA100 有 108 个 SM平均分摊下来每个 SM 需要大约8.6 KB的数据在飞。一个 warp 执行一条 32 位 load32 个 lane × 4 字节 128 字节正好 4 个 32 字节 sector。所以每个 SM 至少要有需要的在飞 load 指令数 8.6 KB / 128 B ≈ 67 条而 A100 每个 SM 最多常驻 64 个 warp。也就是说哪怕每个 warp 都恰好有一条 load 在飞也才勉强够线。一旦考虑到指令发射本身有间隔、大部分 warp 处在计算或等待状态、尾部的 block 还没铺满这个数字就明显不够了。结论很直接只靠一块 warp 一条 load是喂不饱显存的必须让单个 warp 同时挂起多条互不依赖的 load这就是后面要讲的 ILP 和 MLP。有了这个 67 的基准数字你在 profiling 时看到long_scoreboard占比极高就知道该往哪个方向使劲了。1.3 为什么 GPU 不做乱序执行有人会问既然延迟这么要命为什么不像 CPU 那样做乱序执行答案在面积和功耗上。乱序执行需要重排序缓冲、寄存器重命名、复杂的发射仲裁逻辑这套东西在 CPU 上要吃掉大块芯片面积而 GPU 靠的是把同样的面积铺成几百个简单的顺序执行单元。SM 的设计哲学是单线程性能用规模换每个 warp 内部严格顺序执行编译期由编译器做静态调度跨 warp 之间由硬件调度器做动态调度。这个分工决定了 CUDA 编程的一条铁律——想提高性能要么给调度器更多可选的 warp提高占用率要么给编译器更多可重排的独立指令提高 ILP。所有优化手段本质上都落在这两条路上。2. warp 调度器的物理形态四个分区各管一摊2.1 SM 被切成四个 SMSPwarp 是绑定的从 Volta 开始Turing、Ampere、Hopper 一路延续一个 SM 内部被划成四个几乎对称的处理块NVIDIA 文档里叫 SM Sub-Partition简称SMSP。每个 SMSP 拥有自己的一组执行单元FP32/INT32 单元、部分 SFU 等、自己的指令缓冲、自己的寄存器文件切片以及——本篇的主角——一个 warp 调度器。关键在于warp 到 SMSP 的分配是静态的。当一个 CTA 被 GigaThread 引擎派发到某个 SM 上它内部的 warp 会按 warp id 轮流round-robin分配到这 4 个分区上分配完之后这个 warp 的整个生命周期都待在同一个分区里不会迁移。这个设计是为了让寄存器文件和指令缓存控制在可接受的规模不必做全局的仲裁。它带来一个很实际的推论如果你的 block 只有 1 个 warp32 线程那这个 block 只会占用 SM 的 1/4 个分区。虽然 SM 上允许同时驻留多个 block 来把这 4 个分区填满但 Ampere 世代每个 SM 的常驻 block 上限是 32 个GA10x也就是说极端情况下 32 个单 warp block 只能凑出 32 个 warp而分区能容纳的容量是 64 个 warp占用率天然被卡在 50%。所以 blockDim 一般不建议小于 1284 个 warp这样至少能保证每个分区都分到一个 warp往下再靠多 block 叠加把分区喂满。2.2 一个周期只发射一条从双发射到单发射这是很多老资料里的坑。Kepler 和 Maxwell 时代的 SM 支持双发射dual issue一个调度器在一个周期内可以向两条不同的流水线各发一条指令论文和教材里讲的每个调度器维护两个 warp slot、贪心地先看 slot0 再看 slot1就是从那个时代来的。到了 Volta、Turing、Ampere、Hopper双发射被取消了每个调度器每周期只发射一条 warp 指令规则变成简单的贪心 轮转在属于自己分区的就绪 warp 里挑一个发指令如果这个 warp 下一条指令还没就绪就换下一个。这个变化把发射带宽变成了硬约束。以 GA10x 这类消费级芯片为例每个分区有 32 个 FP32 单元一条 warp 级 FP32 指令刚好覆盖 32 个 lane。也就是说分区 FP32 峰值吞吐 1 条 warp 指令/周期 ⇔ 32 lane/周期 ⇔ 128 FLOP/周期FMA 计 2 次要让这个 32 lane 的单元每个周期都有活干调度器就必须每个周期都至少有一个就绪 warp 可发。四个分区同时满足这个条件才叫跑满 FP32 峰值。而实际 kernel 里smsp__issue_active能到 60% 就算相当健康了很多访存密集型 kernel 只有 30%~40%。你在 profiling 报告里看到的算力利用率上不去绝大多数时候不是 ALU 不够快而是调度器经常手上没有可发的 warp。2.3 warp 的四种状态谁才有资格被选中要搞清楚调度器为什么没得选得先理解硬件的状态机。一个 warp 在任一时刻处于以下状态之一Active已经分配了资源、驻扎在 SM 上但可能还没开始执行或已执行完。Eligible就绪下一条指令的所有操作数都齐了随时可以被发射。这才是调度器的可选项。Selected被选中调度器已经把它挑出来准备在这个周期发指令。Stalled停顿因为某种原因发不出去比如操作数还没回来、要走的内存队列满了、在执行__syncthreads()等。调度器每一拍的逻辑就是从自己分区的 warp 池里找 Eligible 的挑一个变成 Selected 并发射。如果一个都找不到这一拍就空转记作 idle执行单元闲置。所以提升性能的核心问题可以一句话概括如何让调度器每一拍都能找到 Eligible warp。这里有个容易被忽略的细节——停顿原因里的not_selected。它表示这个 warp 是就绪的但这一拍没轮到它。如果你在报告里看到not_selected占比很高说明可选的 warp 已经足够多了问题不在占用率上这时候再去堆线程是白费力气得从指令效率、访存合并、减少冗余计算这些角度下手。反过来如果long_scoreboard占比高那才是真的缺 warp 或缺 MLP。区分这两者是读 profiling 报告的第一步。3. 停顿原因逐条拆解调度器究竟在等什么3.1 两级记分牌long scoreboard 与 short scoreboard硬件里维护着一组记分牌scoreboard用来追踪每条结果还没回来的指令防止后续依赖它的指令提前发射。现代架构上大致分成两类Long scoreboard管的是长延迟操作——全局内存 load、本地内存 load、纹理取数等。这些操作的结果可能要几百个周期才回来对应的记分牌槽位少、持有时间长。long_scoreboard停顿占主导意味着 SM 在等显存/L2 的数据是最常见也最需要处理的瓶颈。Short scoreboard管的是相对短延迟的片上操作——共享内存访问、部分特殊功能单元的结果。它走的是 MIOMemory Input/Output队列。short_scoreboard高通常是共享内存访问太密集或者 bank 冲突严重导致 MIO 队列排长队。注意它虽然单次延迟只有 20~30 周期但因为量大积少成多照样能把调度器拖住。两者的处理思路完全不同前者要靠增加在飞访存提高占用率或 ILP、用异步拷贝绕过寄存器后者要靠优化共享内存的访问模式、减少 bank conflict、考虑把一部分数据换到寄存器里。混为一谈会让调优彻底跑偏。3.2 那些容易被忽略的停顿barrier、throttle 与 drain除了两个 scoreboard还有一些停顿原因经常出现在报告里但很多人不知道它们对应什么停顿原因触发场景处理方向wait固定延迟依赖比如等前一条 ALU 指令的结果展开循环、制造独立指令流barrier卡在__syncthreads()上等同一个 block 内的其他 warp减少同步次数检查负载均衡membar内存栅栏等待访存顺序保证减少不必要的__threadfencelg_throttleLSU加载存储单元的指令队列满了减少访存指令密度用向量化访问mio_throttleMIO 队列满共享内存/特殊指令太多降低 shared 访问频率math_pipe_throttle数学流水线排队说明计算已经很饱和好事但也可能是长依赖链卡住drainwarp 退出前等待已发出的内存操作排空减少退出前的未完成访存imc_miss常量缓存未命中把常量改成__constant__或参数传值no_instruction指令缓存未命中或 warp 刚启动减小代码体积避免巨型 kernelnot_selected就绪但没被选中占用率已够转攻指令效率math_pipe_throttle这个指标特别值得说它出现说明 SM 的计算流水线已经排满了属于幸福的烦恼。如果你的报告里这个指标很高同时算力利用率也高那基本不用调了。但如果它高而利用率低那就是流水线里有长依赖链——比如一串 FMA 全部互相依赖每个都要等前一个的 4 个周期虽然流水线被占住实际吞吐却上不去。这种情况要靠展开循环制造并行的累加器来解决也就是常说的 ILP。3.3 把停顿原因读出来几行命令的事用 Nsight Compute 抓这些指标其实很简单关键是抓对 section# 全量采集看整体画像较慢适合定位阶段 ncu --set full -o profile_base ./my_app # 只抓调度器与 warp 状态速度快很多 ncu --section SchedulerStats --section WarpStateStats ./my_app # 直接盯几个关键指标 ncu --metrics \ sm__warps_active.avg.pct_of_peak_sustained_active,\ smsp__issue_active.avg.pct_of_peak_sustained_active,\ smsp__average_warps_issue_stalled_long_scoreboard_per_issue_active.ratio,\ smsp__average_warps_issue_stalled_short_scoreboard_per_issue_active.ratio \ ./my_app读报告的顺序我一般是这样先看SpeedOfLight里的计算利用率和显存利用率判断是计算瓶颈还是访存瓶颈再看Occupancy的理论值与实测值差多少差值大说明有别的因素在限制最后看WarpStateStats里排前三的停顿原因它们直接告诉你调度器在等什么。这三步走下来绝大多数 kernel 的瓶颈都能定位剩下的就是针对性改代码。提示smsp__average_warps_issue_stalled_*_per_issue_active.ratio这类指标的含义是平均每次发射时有多少个 warp 卡在这个原因上不是百分比数值越大约严重。刚看的人很容易误解一定要注意分母。4. 延迟隐藏的定量做法占用率、ILP 和 MLP4.1 占用率的真实天花板怎么算占用率occupancy的定义是SM 上实际常驻的 warp 数 / 硬件支持的最大 warp 数。A100 和 GA10x 的上限都是 64 warp2048 线程。理论占用率受三个因素共同压制谁先到顶谁说了算寄存器A100 每个 SM 有 65536 个 32 位寄存器每个线程用的寄存器数 × 32 就是每个 warp 的占用且要按分配粒度通常是 256 个寄存器/warp向上取整。共享内存A100 每个 SM 可配置最多 164 KB 的共享内存每个 block 的分配也有粒度。块数上限GA10x 每个 SM 最多 32 个 block这条限制经常被忽视。举个具体的例子。假设 kernel 用了 40 个寄存器/线程blockDim 是 256每个 warp 的寄存器 40 × 32 1280已是 256 的整数倍 每 SM 可容纳 warp 数 65536 / 1280 51.2 → 51 个 每 block 的 warp 数 256 / 32 8 每 SM 可容纳 block 数 51 / 8 6.375 → 6 个 实际常驻 warp 6 × 8 48 理论占用率 48 / 64 75%注意 51 个 warp 只能凑出 6 个完整的 block剩下 3 个 warp 的位置白白浪费了——这就是分配粒度带来的损耗。如果改成 blockDim 128每 block 4 个 warp51/4 12 个 block受 32 个 block 上限约束不成立12 × 4 48 个 warp还是 75%但如果把寄存器压到 32 个每个 warp 1024 个寄存器65536/1024 64 个 warp此时 blockDim 128 就能拿到 16 个 block × 4 64 warp占用率 100%而 blockDim 256 只能拿 8 个 block 64 warp 也是 100%。可见寄存器数量对占用率的影响是阶跃式的少一个寄存器有时候就能多塞好几个 warp。代码里可以用__launch_bounds__告诉编译器你的目标让它主动控制寄存器用量// 告诉编译器每个 block 最多 256 线程我希望每个 SM 至少能跑 6 个 block // 编译器会据此把寄存器压到 65536 / (256 * 6 / 32) 附近 __global__ void __launch_bounds__(256, 6) saxpy_kernel(const float* __restrict__ x, const float* __restrict__ y, float* __restrict__ out, float a, int n) { int i blockIdx.x * blockDim.x threadIdx.x; if (i n) out[i] a * x[i] y[i]; }但这里有个大坑后面第 5 节会专门讲盲目追求高占用率是新手最容易犯的错。4.2 ILP 与 MLP让单个 warp 挂起更多访存占用率是横向铺开ILP 是纵向挖深。两者是互补的而且在寄存器受限的场景下ILP 往往比堆占用率更划算。**ILPInstruction Level Parallelism**指一个线程内部有多少条互不依赖的指令。最典型的改造就是手写展开累加// 低 ILP每次迭代的 load 结果立刻被使用形成长依赖链 __global__ void sum_low_ilp(const float* __restrict__ in, float* __restrict__ out, int n) { int i blockIdx.x * blockDim.x threadIdx.x; float acc 0.f; for (; i n; i gridDim.x * blockDim.x) acc in[i]; // 下一条 load 的地址依赖 acc不依赖但使用点卡住了 out[i] acc; }上面这段其实不算最糟因为地址计算是独立的编译器能提前发出若干条 load。真正的问题在于acc形成了一个串行的加法依赖链每条 FADD 要等上一条。改用四个独立累加器就能让这四条链并行推进// 高 ILP四个独立累加器 四路展开同时挂起更多 load __global__ void sum_ilp4(const float* __restrict__ in, float* __restrict__ out, int n) { int tid blockIdx.x * blockDim.x threadIdx.x; int stride gridDim.x * blockDim.x; float a0 0.f, a1 0.f, a2 0.f, a3 0.f; int i tid; // 主循环每轮处理 4 个跨步元素四个累加器互不依赖 for (; i 3 * stride n; i 4 * stride) { a0 in[i]; a1 in[i stride]; a2 in[i 2 * stride]; a3 in[i 3 * stride]; } // 尾巴补齐剩余元素 for (; i n; i stride) a0 in[i]; out[tid] (a0 a1) (a2 a3); }这段代码的关键不在少循环了几次而在于四条 load 的地址互不依赖硬件可以在一个 warp 上同时挂起 4 条在飞访存正好把第 1.2 节算出来的 67 条缺口补上一部分。四个累加器最后的归约顺序(a0a1)(a2a3)也是刻意的比顺序相加少了两次依赖延迟。**MLPMemory Level Parallelism**更直接就是让单线程同时发出多条访存。向量化加载是最省事的做法——用float4一次读 16 字节硬件层面就是一条指令带来 4 个 sector 的数据// 前提指针按 16 字节对齐且 n 是 4 的倍数或者做好边界处理 __global__ void add_vec4(const float4* __restrict__ a, const float4* __restrict__ b, float4* __restrict__ c, int n4) { int i blockIdx.x * blockDim.x threadIdx.x; int stride gridDim.x * blockDim.x; for (; i n4; i stride) { float4 va a[i]; float4 vb b[i]; float4 vc; vc.x va.x vb.x; vc.y va.y vb.y; vc.z va.z vb.z; vc.w va.w vb.w; c[i] vc; } }向量化有个前提必须注意对齐。用reinterpret_castfloat4*强转指针时如果原始指针不是 16 字节对齐会触发未对齐访问错误或者性能断崖。稳妥做法是在 host 端用cudaMalloc分配它天然返回 256 字节对齐的地址或者用__align__(16)修饰结构体。如果是处理数组中间的一段宁可老老实实走标量路径也别赌对齐。4.3 Ampere 之后的异步拷贝绕开记分牌从 Ampere 开始硬件提供了cp.async系列指令可以在不占用寄存器的情况下把数据从全局内存直接拷到共享内存。它的价值在于传统路径global → register → shared需要经过寄存器那就要占用记分牌槽位、制造 long scoreboard 停顿异步拷贝走的是旁路通道发射之后直接用cp.async.wait_group统一等待中间这段时间 warp 可以去做别的事。Hopper 上进一步演化为 TMATensor Memory Accelerator由专门的硬件单元按描述符搬运整块数据连地址计算都省了。不过异步拷贝并不是万灵药。它最适合的形态是软件流水——预取下一块数据的同时计算当前块。如果 kernel 本身没有可重叠的计算用了异步拷贝也不会有收益反而增加了代码复杂度。我一般建议先用传统方式把 kernel 写对、测出基线确认瓶颈确实是 long scoreboard再考虑上异步拷贝不要一上来就套模板。5. 调优路上几个反直觉的实测结论5.1 占用率不是越高越好这条值得单独拎出来讲因为它几乎每个新人都会踩。占用率只是手段不是目标。它真正的作用是给调度器提供足够多的备选 warp。一旦调度器已经有了足够多的 Eligible warp表现为not_selected开始显著再往上堆占用率就毫无收益还要付出代价为了压低寄存器用量编译器不得不把变量溢出到本地内存local memory而本地内存本质上是走 L1/L2 的全局内存于是访存指令暴增long scoreboard 反而恶化性能直接倒退。真实场景里 60%~75% 的占用率往往就是甜点区。我做过的一个 stencil kernel寄存器从 40 压到 32 之后占用率从 75% 涨到 100%但long_scoreboard占比从 22% 涨到 41%整体耗时反而增加 18%。后来用__launch_bounds__(256, 5)把寄存器放开回 44 个占用率降到 62%性能回到了最优。判断依据永远是指标不是占用率数字好不好看。5.2 访存合并与调度器是同一件事的两面很多人把访存合并coalescing当成一个纯访存话题其实它和调度器是一回事。一个 warp 的 32 个 lane 如果访问的地址跨越了很多个 32 字节 sector硬件就得拆成多次内存事务每条 load 在 long scoreboard 上挂的时间成倍增长而一个记分牌槽位对应的事务越多停顿时长就越夸张。对比一下两种访问方式对同一个 warp 的影响访问模式单条 warp load 产生的事务数有效带宽scoreboard 占用连续地址lane i 访问 basei4 个 sector接近峰值短跨步访问lane i 访问 basei*32最多 32 个 sector1/8 峰值极长广播所有 lane 同地址1 个 sector看缓存短但竞争跨步 32 的情况相当于每个 lane 单独取一个 sector一次 warp load 要拆成 32 笔事务带宽利用率只有 1/8同时把 LSU 和内存队列全部堵死。这时候你看到的就是lg_throttle和long_scoreboard一起飙高。解法通常是先改数据布局——把结构体数组改成数组结构体AoS→SoA——或者用共享内存做一次转置把不规则访问变成规则访问。这一步做完的效果往往比调占用率、调 blockDim 加起来都大。5.3 warp divergence 被过度关注了最后说一个我认为被高估的点。warp divergence分支发散确实会造成串行化——一个 warp 内 32 个 lane 走了不同分支硬件只能把两条路径都执行一遍用掩码屏蔽不参与的 lane。但它的代价是多执行了分支路径而不是调度器空转。在很多 kernel 里发散带来的额外执行量只要几毫秒而 long scoreboard 带来的空转可能是几十毫秒。先看停顿原因再谈发散顺序反了就会把时间浪费在优化一个不是瓶颈的地方。真正值得警惕的发散是那种分支体内有大量访存的情况——因为两条路径都执行时内存事务也翻倍了这才是真伤。纯粹的计算分支发散在 GPU 这种吞吐导向的架构上没那么可怕。5.4 一个我常用的排查顺序把这么多东西串起来我给自己的 kernel 定了一套固定检查流程基本能覆盖八成情况先用ncu --set full抓一次看 SpeedOfLight确认是 compute bound 还是 memory bound。看 Occupancy实测值远低于理论值时查是不是寄存器、共享内存或 block 上限卡住的。看 WarpStateStats 前三名。long_scoreboard→ 加 ILP/向量化/异步拷贝short_scoreboard→ 查共享内存 bank conflictnot_selected→ 停止堆占用率改看指令效率wait→ 展开循环打破依赖链lg_throttle/mio_throttle→ 检查访存模式和 shared 访问密度。改一处、测一次只改一个变量。性能优化最忌讳一次改五处最后不知道哪个起了作用。记录下来。我习惯在代码注释里留一行改了什么、指标怎么变隔几个月回头看能省下大量重复试错的时间。还有个容易被忽略的小技巧固定 GPU 时钟频率再对比性能。GPU 有动态调频同一份代码两次跑的耗时可能差 10% 以上调优时如果不锁频很容易被噪声带到沟里去。用nvidia-smi -lgc把频率锁到一个固定值跑完记得-rgc复位测量数据才有可比性。踩过几次这种坑之后我现在做任何性能对比都会先锁频这算是用时间换来的一条经验。
刚上手 CUDA 的人几乎都会撞上同一个困惑明明按文档把 block 开到了 256 线程、grid 铺满了整卡跑出来的算力却只有理论峰值的三成。更让人抓狂的是把线程数再翻一倍速度纹丝不动把寄存器限制一下、把 occupancy 拉满反而更慢了。这些现象背后指向同一个东西——SM 里的 warp 调度器和它赖以生存的访存延迟隐藏机制。这两个概念是 CUDA 性能模型的地基理解了它们你再看 Nsight Compute 里那一堆smsp__warp_issue_stalled_*指标才不会发懵调优也才有方向而不是靠多开线程加__restrict__这种碰运气的做法。下面我把这套机制从硬件形态一直讲到实测数据顺带把我在真实项目里踩过的几个反直觉的坑一并交代清楚无论你是刚开始写 kernel 还是已经能看懂 profiling 报告应该都能从中拿到能直接用的东西。1. SM 到底在藏什么延迟先算清这笔账1.1 一条访存指令的延迟有多离谱要理解调度器为什么存在先得知道它要对付的敌人有多大。现代 GPU 上一条普通的全局内存加载指令从发射到数据真正回到寄存器中间要穿过 L1、L2、内存控制器、DRAM 颗粒再原路返回。这条链路在不同层级上的耗时量级大致是这样的数值随架构和频率浮动这里给的是量级不是精确值访问目标典型延迟SM 周期说明寄存器操作数0操作数直通不算延迟FP32 ALU 结果依赖4~6背靠背依赖链的最小间隔共享内存 load20~30走 MIO 队列short scoreboard 管辖L1 命中30~40全局访存但命中片上缓存L2 命中约 200跨 SM 共享走片上网状互连设备内存 DRAM400~800真正去颗粒取数注意最后一行DRAM 的延迟大约是 ALU 依赖延迟的一百倍。这意味着如果 SM 傻乎乎地发射一条访存指令然后原地等数据回来再发射下一条那它 99% 的时间都在发呆。CPU 的做法是堆大缓存 乱序执行 深流水线把这种等待尽量抹平GPU 的做法完全相反——它不消除延迟它用海量并发把延迟填满。这就是延迟隐藏latency hiding四个字的全部含义SM 上常驻几十个 warp某个 warp 卡在等内存的时候调度器立刻切到另一个已经就绪的 warp 去发射指令。流水线永远不空代价是每个 warp 自己看起来都很慢。1.2 用吞吐反推需要多少在飞行的数据光说要用并发填是空话我们可以用最朴素的利特尔法则Littles Law把它算成一个具体数字这样后面判断到底缺不缺 warp就有了标尺。假设目标是打满 A100 的显存带宽约 1555 GB/sDRAM 往返延迟按 600 ns 算那么任意时刻必须处于已发出、未返回状态的数据量是在飞字节数 带宽 × 延迟 ≈ 1555e9 B/s × 600e-9 s ≈ 933,000 B ≈ 0.93 MBA100 有 108 个 SM平均分摊下来每个 SM 需要大约8.6 KB的数据在飞。一个 warp 执行一条 32 位 load32 个 lane × 4 字节 128 字节正好 4 个 32 字节 sector。所以每个 SM 至少要有需要的在飞 load 指令数 8.6 KB / 128 B ≈ 67 条而 A100 每个 SM 最多常驻 64 个 warp。也就是说哪怕每个 warp 都恰好有一条 load 在飞也才勉强够线。一旦考虑到指令发射本身有间隔、大部分 warp 处在计算或等待状态、尾部的 block 还没铺满这个数字就明显不够了。结论很直接只靠一块 warp 一条 load是喂不饱显存的必须让单个 warp 同时挂起多条互不依赖的 load这就是后面要讲的 ILP 和 MLP。有了这个 67 的基准数字你在 profiling 时看到long_scoreboard占比极高就知道该往哪个方向使劲了。1.3 为什么 GPU 不做乱序执行有人会问既然延迟这么要命为什么不像 CPU 那样做乱序执行答案在面积和功耗上。乱序执行需要重排序缓冲、寄存器重命名、复杂的发射仲裁逻辑这套东西在 CPU 上要吃掉大块芯片面积而 GPU 靠的是把同样的面积铺成几百个简单的顺序执行单元。SM 的设计哲学是单线程性能用规模换每个 warp 内部严格顺序执行编译期由编译器做静态调度跨 warp 之间由硬件调度器做动态调度。这个分工决定了 CUDA 编程的一条铁律——想提高性能要么给调度器更多可选的 warp提高占用率要么给编译器更多可重排的独立指令提高 ILP。所有优化手段本质上都落在这两条路上。2. warp 调度器的物理形态四个分区各管一摊2.1 SM 被切成四个 SMSPwarp 是绑定的从 Volta 开始Turing、Ampere、Hopper 一路延续一个 SM 内部被划成四个几乎对称的处理块NVIDIA 文档里叫 SM Sub-Partition简称SMSP。每个 SMSP 拥有自己的一组执行单元FP32/INT32 单元、部分 SFU 等、自己的指令缓冲、自己的寄存器文件切片以及——本篇的主角——一个 warp 调度器。关键在于warp 到 SMSP 的分配是静态的。当一个 CTA 被 GigaThread 引擎派发到某个 SM 上它内部的 warp 会按 warp id 轮流round-robin分配到这 4 个分区上分配完之后这个 warp 的整个生命周期都待在同一个分区里不会迁移。这个设计是为了让寄存器文件和指令缓存控制在可接受的规模不必做全局的仲裁。它带来一个很实际的推论如果你的 block 只有 1 个 warp32 线程那这个 block 只会占用 SM 的 1/4 个分区。虽然 SM 上允许同时驻留多个 block 来把这 4 个分区填满但 Ampere 世代每个 SM 的常驻 block 上限是 32 个GA10x也就是说极端情况下 32 个单 warp block 只能凑出 32 个 warp而分区能容纳的容量是 64 个 warp占用率天然被卡在 50%。所以 blockDim 一般不建议小于 1284 个 warp这样至少能保证每个分区都分到一个 warp往下再靠多 block 叠加把分区喂满。2.2 一个周期只发射一条从双发射到单发射这是很多老资料里的坑。Kepler 和 Maxwell 时代的 SM 支持双发射dual issue一个调度器在一个周期内可以向两条不同的流水线各发一条指令论文和教材里讲的每个调度器维护两个 warp slot、贪心地先看 slot0 再看 slot1就是从那个时代来的。到了 Volta、Turing、Ampere、Hopper双发射被取消了每个调度器每周期只发射一条 warp 指令规则变成简单的贪心 轮转在属于自己分区的就绪 warp 里挑一个发指令如果这个 warp 下一条指令还没就绪就换下一个。这个变化把发射带宽变成了硬约束。以 GA10x 这类消费级芯片为例每个分区有 32 个 FP32 单元一条 warp 级 FP32 指令刚好覆盖 32 个 lane。也就是说分区 FP32 峰值吞吐 1 条 warp 指令/周期 ⇔ 32 lane/周期 ⇔ 128 FLOP/周期FMA 计 2 次要让这个 32 lane 的单元每个周期都有活干调度器就必须每个周期都至少有一个就绪 warp 可发。四个分区同时满足这个条件才叫跑满 FP32 峰值。而实际 kernel 里smsp__issue_active能到 60% 就算相当健康了很多访存密集型 kernel 只有 30%~40%。你在 profiling 报告里看到的算力利用率上不去绝大多数时候不是 ALU 不够快而是调度器经常手上没有可发的 warp。2.3 warp 的四种状态谁才有资格被选中要搞清楚调度器为什么没得选得先理解硬件的状态机。一个 warp 在任一时刻处于以下状态之一Active已经分配了资源、驻扎在 SM 上但可能还没开始执行或已执行完。Eligible就绪下一条指令的所有操作数都齐了随时可以被发射。这才是调度器的可选项。Selected被选中调度器已经把它挑出来准备在这个周期发指令。Stalled停顿因为某种原因发不出去比如操作数还没回来、要走的内存队列满了、在执行__syncthreads()等。调度器每一拍的逻辑就是从自己分区的 warp 池里找 Eligible 的挑一个变成 Selected 并发射。如果一个都找不到这一拍就空转记作 idle执行单元闲置。所以提升性能的核心问题可以一句话概括如何让调度器每一拍都能找到 Eligible warp。这里有个容易被忽略的细节——停顿原因里的not_selected。它表示这个 warp 是就绪的但这一拍没轮到它。如果你在报告里看到not_selected占比很高说明可选的 warp 已经足够多了问题不在占用率上这时候再去堆线程是白费力气得从指令效率、访存合并、减少冗余计算这些角度下手。反过来如果long_scoreboard占比高那才是真的缺 warp 或缺 MLP。区分这两者是读 profiling 报告的第一步。3. 停顿原因逐条拆解调度器究竟在等什么3.1 两级记分牌long scoreboard 与 short scoreboard硬件里维护着一组记分牌scoreboard用来追踪每条结果还没回来的指令防止后续依赖它的指令提前发射。现代架构上大致分成两类Long scoreboard管的是长延迟操作——全局内存 load、本地内存 load、纹理取数等。这些操作的结果可能要几百个周期才回来对应的记分牌槽位少、持有时间长。long_scoreboard停顿占主导意味着 SM 在等显存/L2 的数据是最常见也最需要处理的瓶颈。Short scoreboard管的是相对短延迟的片上操作——共享内存访问、部分特殊功能单元的结果。它走的是 MIOMemory Input/Output队列。short_scoreboard高通常是共享内存访问太密集或者 bank 冲突严重导致 MIO 队列排长队。注意它虽然单次延迟只有 20~30 周期但因为量大积少成多照样能把调度器拖住。两者的处理思路完全不同前者要靠增加在飞访存提高占用率或 ILP、用异步拷贝绕过寄存器后者要靠优化共享内存的访问模式、减少 bank conflict、考虑把一部分数据换到寄存器里。混为一谈会让调优彻底跑偏。3.2 那些容易被忽略的停顿barrier、throttle 与 drain除了两个 scoreboard还有一些停顿原因经常出现在报告里但很多人不知道它们对应什么停顿原因触发场景处理方向wait固定延迟依赖比如等前一条 ALU 指令的结果展开循环、制造独立指令流barrier卡在__syncthreads()上等同一个 block 内的其他 warp减少同步次数检查负载均衡membar内存栅栏等待访存顺序保证减少不必要的__threadfencelg_throttleLSU加载存储单元的指令队列满了减少访存指令密度用向量化访问mio_throttleMIO 队列满共享内存/特殊指令太多降低 shared 访问频率math_pipe_throttle数学流水线排队说明计算已经很饱和好事但也可能是长依赖链卡住drainwarp 退出前等待已发出的内存操作排空减少退出前的未完成访存imc_miss常量缓存未命中把常量改成__constant__或参数传值no_instruction指令缓存未命中或 warp 刚启动减小代码体积避免巨型 kernelnot_selected就绪但没被选中占用率已够转攻指令效率math_pipe_throttle这个指标特别值得说它出现说明 SM 的计算流水线已经排满了属于幸福的烦恼。如果你的报告里这个指标很高同时算力利用率也高那基本不用调了。但如果它高而利用率低那就是流水线里有长依赖链——比如一串 FMA 全部互相依赖每个都要等前一个的 4 个周期虽然流水线被占住实际吞吐却上不去。这种情况要靠展开循环制造并行的累加器来解决也就是常说的 ILP。3.3 把停顿原因读出来几行命令的事用 Nsight Compute 抓这些指标其实很简单关键是抓对 section# 全量采集看整体画像较慢适合定位阶段 ncu --set full -o profile_base ./my_app # 只抓调度器与 warp 状态速度快很多 ncu --section SchedulerStats --section WarpStateStats ./my_app # 直接盯几个关键指标 ncu --metrics \ sm__warps_active.avg.pct_of_peak_sustained_active,\ smsp__issue_active.avg.pct_of_peak_sustained_active,\ smsp__average_warps_issue_stalled_long_scoreboard_per_issue_active.ratio,\ smsp__average_warps_issue_stalled_short_scoreboard_per_issue_active.ratio \ ./my_app读报告的顺序我一般是这样先看SpeedOfLight里的计算利用率和显存利用率判断是计算瓶颈还是访存瓶颈再看Occupancy的理论值与实测值差多少差值大说明有别的因素在限制最后看WarpStateStats里排前三的停顿原因它们直接告诉你调度器在等什么。这三步走下来绝大多数 kernel 的瓶颈都能定位剩下的就是针对性改代码。提示smsp__average_warps_issue_stalled_*_per_issue_active.ratio这类指标的含义是平均每次发射时有多少个 warp 卡在这个原因上不是百分比数值越大约严重。刚看的人很容易误解一定要注意分母。4. 延迟隐藏的定量做法占用率、ILP 和 MLP4.1 占用率的真实天花板怎么算占用率occupancy的定义是SM 上实际常驻的 warp 数 / 硬件支持的最大 warp 数。A100 和 GA10x 的上限都是 64 warp2048 线程。理论占用率受三个因素共同压制谁先到顶谁说了算寄存器A100 每个 SM 有 65536 个 32 位寄存器每个线程用的寄存器数 × 32 就是每个 warp 的占用且要按分配粒度通常是 256 个寄存器/warp向上取整。共享内存A100 每个 SM 可配置最多 164 KB 的共享内存每个 block 的分配也有粒度。块数上限GA10x 每个 SM 最多 32 个 block这条限制经常被忽视。举个具体的例子。假设 kernel 用了 40 个寄存器/线程blockDim 是 256每个 warp 的寄存器 40 × 32 1280已是 256 的整数倍 每 SM 可容纳 warp 数 65536 / 1280 51.2 → 51 个 每 block 的 warp 数 256 / 32 8 每 SM 可容纳 block 数 51 / 8 6.375 → 6 个 实际常驻 warp 6 × 8 48 理论占用率 48 / 64 75%注意 51 个 warp 只能凑出 6 个完整的 block剩下 3 个 warp 的位置白白浪费了——这就是分配粒度带来的损耗。如果改成 blockDim 128每 block 4 个 warp51/4 12 个 block受 32 个 block 上限约束不成立12 × 4 48 个 warp还是 75%但如果把寄存器压到 32 个每个 warp 1024 个寄存器65536/1024 64 个 warp此时 blockDim 128 就能拿到 16 个 block × 4 64 warp占用率 100%而 blockDim 256 只能拿 8 个 block 64 warp 也是 100%。可见寄存器数量对占用率的影响是阶跃式的少一个寄存器有时候就能多塞好几个 warp。代码里可以用__launch_bounds__告诉编译器你的目标让它主动控制寄存器用量// 告诉编译器每个 block 最多 256 线程我希望每个 SM 至少能跑 6 个 block // 编译器会据此把寄存器压到 65536 / (256 * 6 / 32) 附近 __global__ void __launch_bounds__(256, 6) saxpy_kernel(const float* __restrict__ x, const float* __restrict__ y, float* __restrict__ out, float a, int n) { int i blockIdx.x * blockDim.x threadIdx.x; if (i n) out[i] a * x[i] y[i]; }但这里有个大坑后面第 5 节会专门讲盲目追求高占用率是新手最容易犯的错。4.2 ILP 与 MLP让单个 warp 挂起更多访存占用率是横向铺开ILP 是纵向挖深。两者是互补的而且在寄存器受限的场景下ILP 往往比堆占用率更划算。**ILPInstruction Level Parallelism**指一个线程内部有多少条互不依赖的指令。最典型的改造就是手写展开累加// 低 ILP每次迭代的 load 结果立刻被使用形成长依赖链 __global__ void sum_low_ilp(const float* __restrict__ in, float* __restrict__ out, int n) { int i blockIdx.x * blockDim.x threadIdx.x; float acc 0.f; for (; i n; i gridDim.x * blockDim.x) acc in[i]; // 下一条 load 的地址依赖 acc不依赖但使用点卡住了 out[i] acc; }上面这段其实不算最糟因为地址计算是独立的编译器能提前发出若干条 load。真正的问题在于acc形成了一个串行的加法依赖链每条 FADD 要等上一条。改用四个独立累加器就能让这四条链并行推进// 高 ILP四个独立累加器 四路展开同时挂起更多 load __global__ void sum_ilp4(const float* __restrict__ in, float* __restrict__ out, int n) { int tid blockIdx.x * blockDim.x threadIdx.x; int stride gridDim.x * blockDim.x; float a0 0.f, a1 0.f, a2 0.f, a3 0.f; int i tid; // 主循环每轮处理 4 个跨步元素四个累加器互不依赖 for (; i 3 * stride n; i 4 * stride) { a0 in[i]; a1 in[i stride]; a2 in[i 2 * stride]; a3 in[i 3 * stride]; } // 尾巴补齐剩余元素 for (; i n; i stride) a0 in[i]; out[tid] (a0 a1) (a2 a3); }这段代码的关键不在少循环了几次而在于四条 load 的地址互不依赖硬件可以在一个 warp 上同时挂起 4 条在飞访存正好把第 1.2 节算出来的 67 条缺口补上一部分。四个累加器最后的归约顺序(a0a1)(a2a3)也是刻意的比顺序相加少了两次依赖延迟。**MLPMemory Level Parallelism**更直接就是让单线程同时发出多条访存。向量化加载是最省事的做法——用float4一次读 16 字节硬件层面就是一条指令带来 4 个 sector 的数据// 前提指针按 16 字节对齐且 n 是 4 的倍数或者做好边界处理 __global__ void add_vec4(const float4* __restrict__ a, const float4* __restrict__ b, float4* __restrict__ c, int n4) { int i blockIdx.x * blockDim.x threadIdx.x; int stride gridDim.x * blockDim.x; for (; i n4; i stride) { float4 va a[i]; float4 vb b[i]; float4 vc; vc.x va.x vb.x; vc.y va.y vb.y; vc.z va.z vb.z; vc.w va.w vb.w; c[i] vc; } }向量化有个前提必须注意对齐。用reinterpret_castfloat4*强转指针时如果原始指针不是 16 字节对齐会触发未对齐访问错误或者性能断崖。稳妥做法是在 host 端用cudaMalloc分配它天然返回 256 字节对齐的地址或者用__align__(16)修饰结构体。如果是处理数组中间的一段宁可老老实实走标量路径也别赌对齐。4.3 Ampere 之后的异步拷贝绕开记分牌从 Ampere 开始硬件提供了cp.async系列指令可以在不占用寄存器的情况下把数据从全局内存直接拷到共享内存。它的价值在于传统路径global → register → shared需要经过寄存器那就要占用记分牌槽位、制造 long scoreboard 停顿异步拷贝走的是旁路通道发射之后直接用cp.async.wait_group统一等待中间这段时间 warp 可以去做别的事。Hopper 上进一步演化为 TMATensor Memory Accelerator由专门的硬件单元按描述符搬运整块数据连地址计算都省了。不过异步拷贝并不是万灵药。它最适合的形态是软件流水——预取下一块数据的同时计算当前块。如果 kernel 本身没有可重叠的计算用了异步拷贝也不会有收益反而增加了代码复杂度。我一般建议先用传统方式把 kernel 写对、测出基线确认瓶颈确实是 long scoreboard再考虑上异步拷贝不要一上来就套模板。5. 调优路上几个反直觉的实测结论5.1 占用率不是越高越好这条值得单独拎出来讲因为它几乎每个新人都会踩。占用率只是手段不是目标。它真正的作用是给调度器提供足够多的备选 warp。一旦调度器已经有了足够多的 Eligible warp表现为not_selected开始显著再往上堆占用率就毫无收益还要付出代价为了压低寄存器用量编译器不得不把变量溢出到本地内存local memory而本地内存本质上是走 L1/L2 的全局内存于是访存指令暴增long scoreboard 反而恶化性能直接倒退。真实场景里 60%~75% 的占用率往往就是甜点区。我做过的一个 stencil kernel寄存器从 40 压到 32 之后占用率从 75% 涨到 100%但long_scoreboard占比从 22% 涨到 41%整体耗时反而增加 18%。后来用__launch_bounds__(256, 5)把寄存器放开回 44 个占用率降到 62%性能回到了最优。判断依据永远是指标不是占用率数字好不好看。5.2 访存合并与调度器是同一件事的两面很多人把访存合并coalescing当成一个纯访存话题其实它和调度器是一回事。一个 warp 的 32 个 lane 如果访问的地址跨越了很多个 32 字节 sector硬件就得拆成多次内存事务每条 load 在 long scoreboard 上挂的时间成倍增长而一个记分牌槽位对应的事务越多停顿时长就越夸张。对比一下两种访问方式对同一个 warp 的影响访问模式单条 warp load 产生的事务数有效带宽scoreboard 占用连续地址lane i 访问 basei4 个 sector接近峰值短跨步访问lane i 访问 basei*32最多 32 个 sector1/8 峰值极长广播所有 lane 同地址1 个 sector看缓存短但竞争跨步 32 的情况相当于每个 lane 单独取一个 sector一次 warp load 要拆成 32 笔事务带宽利用率只有 1/8同时把 LSU 和内存队列全部堵死。这时候你看到的就是lg_throttle和long_scoreboard一起飙高。解法通常是先改数据布局——把结构体数组改成数组结构体AoS→SoA——或者用共享内存做一次转置把不规则访问变成规则访问。这一步做完的效果往往比调占用率、调 blockDim 加起来都大。5.3 warp divergence 被过度关注了最后说一个我认为被高估的点。warp divergence分支发散确实会造成串行化——一个 warp 内 32 个 lane 走了不同分支硬件只能把两条路径都执行一遍用掩码屏蔽不参与的 lane。但它的代价是多执行了分支路径而不是调度器空转。在很多 kernel 里发散带来的额外执行量只要几毫秒而 long scoreboard 带来的空转可能是几十毫秒。先看停顿原因再谈发散顺序反了就会把时间浪费在优化一个不是瓶颈的地方。真正值得警惕的发散是那种分支体内有大量访存的情况——因为两条路径都执行时内存事务也翻倍了这才是真伤。纯粹的计算分支发散在 GPU 这种吞吐导向的架构上没那么可怕。5.4 一个我常用的排查顺序把这么多东西串起来我给自己的 kernel 定了一套固定检查流程基本能覆盖八成情况先用ncu --set full抓一次看 SpeedOfLight确认是 compute bound 还是 memory bound。看 Occupancy实测值远低于理论值时查是不是寄存器、共享内存或 block 上限卡住的。看 WarpStateStats 前三名。long_scoreboard→ 加 ILP/向量化/异步拷贝short_scoreboard→ 查共享内存 bank conflictnot_selected→ 停止堆占用率改看指令效率wait→ 展开循环打破依赖链lg_throttle/mio_throttle→ 检查访存模式和 shared 访问密度。改一处、测一次只改一个变量。性能优化最忌讳一次改五处最后不知道哪个起了作用。记录下来。我习惯在代码注释里留一行改了什么、指标怎么变隔几个月回头看能省下大量重复试错的时间。还有个容易被忽略的小技巧固定 GPU 时钟频率再对比性能。GPU 有动态调频同一份代码两次跑的耗时可能差 10% 以上调优时如果不锁频很容易被噪声带到沟里去。用nvidia-smi -lgc把频率锁到一个固定值跑完记得-rgc复位测量数据才有可比性。踩过几次这种坑之后我现在做任何性能对比都会先锁频这算是用时间换来的一条经验。
