1. 为什么我劝你先啃透 warp 调度与延迟隐藏先说一个我这两年里反复遇到的场景一个看着很朴素的带宽型 kernel比如向量加、归约、或者转置用 ncu 一看 occupancy 已经 100% 了Achieved Occupancy 跑到 64 warps/SM 满值但 DRAM Throughput 只有峰值的 40% 出头。这时候很多人第一反应是是不是 block 开少了是不是 grid 不够大于是把 grid 从 108 拉到 10800结果数字纹丝不动。真正的原因藏在smsp__average_warps_issue_stalled_long_scoreboard_per_issue_active这个指标里——它在告诉你绝大多数 warp 其实都卡在等全局内存返回调度器每一轮都在面对一堆没资格发射的 warp。CUDA里所有性能话题最终都会落到两个地方SM里的warp 调度器在每个周期挑谁发射以及当被挑中的 warp 需要等 500 个周期才能拿到数据时怎么让别的 warp 把这段空档填满也就是访存延迟隐藏。这两件事是同一枚硬币的两面调度机制决定谁能上延迟隐藏决定上去之后能不能把流水线喂饱。搞不清楚这两层你做算子优化基本就是瞎调参数。这篇文章我打算按自己的理解路径来写先把 SM 的物理结构拆开看看有哪些部件、各自的预算上限是多少再钻进调度器的状态机看它每个周期到底在做什么决策然后用 Littles Law 把我到底需要多少 warp 才能藏住延迟算成一道小学算术题最后拿一个真实的带宽型 kernel 做实录一边改代码一边看指标把 42% 拉到 88%。中间会穿插 Nsight Compute 的指标速查表和我这些年踩过的坑。适合已经能写能跑 CUDA kernel、但性能始终卡在天花板下的朋友也适合刚接触 GPU 性能分析、想建立一套系统判断框架的人。不需要你懂电路但需要你会写for循环。1.1 一个反直觉的现象occupancy 拉满带宽只有四成我拿归约 kernel 举过太多次例子了。假设你写了一个最朴素的版本每个线程用网格跨步循环累加最后atomicAdd到全局。这个 kernel 的 occupancy 通常能轻松打满因为每线程只用了十几个寄存器共享内存为零块大小设成 256 的话一个 SM 能塞下 8 个块、64 个 warp正好顶到硬件上限。但它的性能往往很难看。原因不在线程不够多而在每个线程只携带了极少的内存并行度。具体说线程执行s x[i]这条语句时会先发起一次 32-bit 的全局加载然后立刻需要这个值来做加法。加载指令在 LSU 里排队、穿过 L1、L2、到 HBM再原路返回来这一趟在 A100 上大约要 400 到 600 个周期。返回之前这条线程所在的 warp 处于long scoreboard阻塞状态调度器不会选它。如果所有 64 个 warp 都处于这种发一条、等回来的节奏那么整个 SM 在每个瞬间真正在飞的访存请求数其实非常有限。粗算一下一个 warp 发一条LDG.E.32指令请求 128 字节32 线程 × 4 字节。64 个 warp 全部各有一条在飞也就 8KB 的数据在路上。而 A100 满带宽时每 SM 每周期需要的字节吞吐是十几个字节500 周期的延迟意味着你需要大约 6 到 7KB 的数据持续在飞才能把管道填满——看起来 8KB 够了问题是现实中没有这么理想地址计算、循环变量更新、请求返回后的消费阶段都会占掉时间片实际同时悬挂的请求数往往只有理论值的一半甚至三分之一。这就是为什么 occupancy 100% 却只能跑出四成带宽。1.2 这篇文章会带你走完的完整链路整条链路的逻辑顺序是这样的SM 被切成四个子分区每个子分区有自己的 warp 调度器和寄存器堆调度器每周期从就绪的 warp 里挑一个发射被选中的前提是该 warp 的下一条指令的操作数已经准备好如果没准备好warp 就带着一个 stall 原因挂起等资源就绪后重新进入候选池而所谓隐藏延迟本质就是让调度器手上永远有足够多的备用 warp 可选这样即使一批 warp 在等内存另一批仍然能把发射端口占满。把这套逻辑想通之后你会发现几乎所有 CUDA 优化手段都能被归到三类里增加并发 warp 数提高 TLP靠 occupancy、增加每个 warp 的在飞请求数提高 MLP靠展开和向量化、降低单次访问的延迟或提高带宽利用率靠合并访问、共享内存、缓存策略。剩下的都是这三类的变体。后面每一章我都会落到具体的公式、代码或者指标上不停留在概念。2. 拆开 SM调度器到底在什么样的物理环境里工作很多人学 CUDA 时把 SM 当成一个黑盒只知道线程块扔进去、结果吐出来。但要理解调度必须先知道调度器手上有多少资源可分配、每个周期能花多少钱。这一章我们把 SM 的图纸画出来顺便把几个关键的数字钉死后面算延迟隐藏的时候要用。2.1 SM 分四个 SMSP每个 SMSP 是独立小工厂从 Volta 开始NVIDIA 的 SM 结构就固定成了四分之一分区的形态。一个 SM 内部被切成 4 个处理块文档里叫 processing block 或者 sub-partition我习惯直接叫 SMSP。每个 SMSP 内部包含一个 warp 调度器、一个指令分发单元dispatch unit、一份独立的寄存器堆、一个 L0 指令缓存以及一组执行单元。执行单元这一块从 Volta 到 Hopper 大致是这样演进的Volta 每个 SMSP 有 16 个 FP32 单元加 16 个 INT32 单元Turing 把 FP32 和 INT32 合并成 16 个可双用的单元再加 16 个 FP32Ampere 的 GA10x 每 SM 有 128 个 FP32分摊到 SMSP 就是 32 个Hopper 也维持 128 FP32/SM 的规模。除此之外每个 SMSP 还带一组 LSU加载存储单元、SFU特殊函数单元、以及 Tensor Core。关键的一点是这四个 SMSP 之间是高度独立的。它们各自调度各自的 warp各自发射各自的指令互不干涉。一个 SMSP 因为所有 warp 都在等内存而空转另外三个 SMSP 完全不受影响。这个特性在分析性能时有实际意义如果你发现 4 个 SMSP 的利用率严重不均那通常是 warp 分配或者分支分歧导致的问题而不是带宽问题。寄存器堆也是按 SMSP 划分的。A100 每个 SM 有 65536 个 32-bit 寄存器摊到 4 个 SMSP 就是每个 16384 个。这个数字直接决定了你能塞进多少 warp——后面算 occupancy 的时候会用到。2.2 warp 为什么必须是 32 个线程这个问题我在很多次分享里被问到过。答案不复杂GPU 的执行单元是按通道组织的一个 SMSP 里的 32 个 FP32 单元排成一条流水线一条指令进来32 个操作数同时分发到 32 个通道上并行执行。这就是 SIMT 的本意。所以一个 warp 自然就是 32 个线程它对应一次指令发射所能覆盖的最大并行宽度。换句话说warp 不是一个抽象概念它是硬件发射粒度的直接映射。调度器每次做决策选的是一个 warp而不是单个线程指令缓存里存的也是 warp 级的指令流寄存器分配的物理单位同样是按 warp 组织的。这也解释了为什么当 warp 内部出现分支分歧时代价那么大——32 个通道里有 16 个走 if 分支、16 个走 else硬件没有部分发射这种能力只能把两条路径串行执行两遍各屏蔽掉一半通道。从 Volta 开始NVIDIA 引入了独立线程调度Independent Thread Scheduling每个线程有了自己的 PC 和调用栈理论上可以让分歧的两半真正并行推进。但这不等于没有代价重汇聚点变得不确定__syncwarp()从基本没用变成了必须显式加。如果你在 Volta 之后的架构上写依赖 warp 同步的代码比如 warp 级归约却没有加__syncwarp()是有概率踩到正确性问题的。2.3 每个周期的发射预算与资源上限调度器的预算这个词很贴切每个 SMSP 的调度器每周期最多只能发射1 条指令。注意是 1 条不是 2 条。早期 Fermi 时代某些 SM 有双发射能力但那个设计早被砍掉了从 Kepler 之后主流架构都是每调度器每周期 1 条。所以一个 SM 每周期最多发射 4 条指令4 个 SMSP 各一条。这个上限意味着什么意味着指令发射吞吐本身就是一种稀缺资源。如果你的 kernel 每条指令能干的活太少比如全是标量 4 字节加载那即使内存带宽没跑满发射端口也可能先被打满。这是smsp__issue_active和smsp__inst_executed这类指标的意义所在。我把几个关键上限整理成一张表方便后面查阅。不同架构的具体数值会有差异下表以 Ampere GA100A100为主其他架构我会在括号里标注资源项每 SM 上限每 SMSP 上限说明常驻 warp 数6416自 Volta 起固定Hopper 相同常驻线程数2048512与 warp 上限一致常驻线程块数32—块数过多会导致尾部分配碎片32-bit 寄存器6553616384分配粒度通常为每 warp 256 个共享内存 / L1192KBHopper 256KB—可配置为共享内存的比例每周期发射指令41硬上限无法突破FP32 单元12832Hopper 相同注意寄存器分配有粒度限制。即使你的 kernel 平均每线程只用 20 个寄存器编译器也会按 warp 粒度通常是 256 个寄存器对齐分配所以算出来的 occupancy 往往比手算值低一档。这一点在做 occupancy 精算时特别容易忽略。举个具体的数如果你用__launch_bounds__(256)不限制寄存器编译器给了 40 个寄存器/线程。一个块 256 线程 8 个 warp每 warp 分配 40×32 1280 个寄存器向上取整到 1280已经是 256 的倍数5 个 256 单位。8 个 warp 就是 10240 个寄存器。65536 / 10240 6.4所以只能放 6 个块 48 个 warpoccupancy 是 48/64 75%。如果你把寄存器压到 32 个每块 8192 个能放 8 个块 64 warp正好 100%。这中间差的就是那 8 个寄存器。3. warp 调度器每周期只做一件小事的决策者现在进入核心。调度器每个周期都在重复同一套动作扫描候选 warp、判断是否就绪、挑一个、发射一条指令。听起来简单但里面的策略细节决定了整个 GPU 的行为特征。这一章我把状态机、调度策略和分支处理讲清楚。3.1 warp 的四种状态resident / eligible / stalled / selected任何时刻一个 SMSP 里的 16 个 warp 槽位中每个 warp 都处于以下某种状态。理解这四个状态是所有性能分析的基础。Resident常驻指的是这个 warp 已经被分配到这个 SM 上占着寄存器和槽位。常驻不等于能跑它可能正在等任何东西。Eligible就绪指的是这个 warp 的下一条指令所需的所有操作数都已经准备好了理论上可以立刻发射。调度器只会在 eligible 集合里做选择。Stalled阻塞指的是 warp 在等某个资源典型的有等全局内存数据回来long scoreboard、等共享内存或寄存器依赖short scoreboard、等固定延迟的运算结果wait、等 barrier__syncthreads、等指令缓存填充no instruction。每个 stall 都会被硬件计数这就是 ncu 里那些 stall 指标的来源。Selected被选中指的是这一周期调度器挑了它指令被送进分发单元。选中的下一周期这个 warp 大概率会进入 stalled 状态如果它下一条指令依赖本次结果也可能继续 eligible如果下一条是独立的。这里有个容易混淆的点stall_not_selected这个指标代表warp 已经就绪但这一周期没被选中。它高不是坏事恰恰说明你的并行度很充裕调度器挑不过来。真正要警惕的是 long scoreboard 和 short scoreboard 占比高。状态之间的迁移是硬件自动完成的程序员无法直接干预但可以通过改变代码结构去间接影响比如把一条访问拆成多路独立的访问就能让 warp 在等第一路数据的同时第二路请求已经发出去了从而缩短 stalled 的时间占比。3.2 贪心-最老优先GTO为什么不用简单轮询调度策略这件事NVIDIA 从来没有官方公开过细节但学术界用微基准测试做过大量逆向分析结论比较一致实际行为非常接近贪心-最老优先策略也就是 Greedy-Then-Oldest。具体怎么理解假设当前周期有 3 个 warp 处于 eligible 状态。如果上一周期被选中的那个 warp 这一周期仍然 eligible调度器会继续选它——这就是贪心的部分。这个设计的动机很好理解连续发射同一个 warp 的指令可以复用上一周期已经译码的结果省掉一次指令译码的开销也避免在多个 warp 之间来回切换上下文。那什么时候换人当当前 warp 因为依赖阻塞而变成 stalled 时调度器就把它从候选里踢出去转向下一个。这时候的选择依据是优先级而优先级的核心因素是这个 warp 有多久没发射过指令了。等待时间最长的那个 warp 优先获得发射机会这就是最老的部分。为什么要加最老这一条如果只贪心不轮转理论上可能出现某个活跃 warp 长期霸占发射端口其他 warp 迟迟拿不到机会极端情况下会导致部分线程饿死。加入年龄因素后硬件保证了公平性下限。这个策略对写代码有两个直接的启示。第一长的依赖链对调度器是友好的一个 warp 连续执行十几条独立的计算指令会被连续发射效率很高。第二不要让一个 warp 的指令流频繁地在不同资源之间跳跃比如算一条、访一次存、再算一条、再访一次这样每次访存都会强制上下文切换白白损失发射效率。这实际上是把访存批量攒起来再一起发这种优化手法的理论依据。3.3 双发射的兴衰与 Volta 之后的独立线程调度我在网上看到过不少人还在讲SM 每周期可以发射 2 条指令这个说法在 Fermi 的某些型号上有过对应的硬件设计但早就不是主流了。从 Kepler 开始主流 SM 设计就是每个调度器每周期发射 1 条指令。所以任何基于双发射的性能模型推导出来的结论在现代架构上都不成立。Volta 带来的真正重要的变化是独立线程调度。在这之前一个 warp 里的 32 个线程共享一个 PC遇到分歧只能串行化Volta 之后每个线程有自己的 PC 和调用栈warp 在硬件层面被进一步细分成了 4 组、每组 8 个线程来管理调度器维护的是所谓的线程束内组级别的状态。这带来的实际影响有两个。其一是分歧的代价模型变了以前是两条路径串行执行、时间相加现在是两组线程可以交替推进硬件在每个周期选择一组发射总时间仍然会长于无分歧情况但不完全是简单相加了。其二是同步语义变了__syncwarp()从可选变成了必需。我在一个 warp 级 shuffle 归约的代码里就踩过这个坑在 Kepler 上跑得好好的换到 Ampere 上偶发结果错误加了__syncwarp()之后问题消失。3.4 warp 分歧、重汇聚与 ITL 的实际代价关于分歧我想补充一个容易被忽略的细节分歧的代价不仅取决于分支的条数还取决于分歧的持续时间。如果 if/else 两个分支各自只有两三条指令硬件处理起来很快如果其中一个分支里有一整个循环那这个循环的每一轮迭代都要在屏蔽状态下执行代价就很大了。常见的处理手法有这么几类。第一类是把分歧改成无分支计算用算术或位运算替代条件跳转比如用x flag ? a : b这种三元表达式编译器通常会生成SEL指令而不是分支避免了分歧。第二类是按数据重排把一个块内的线程按照某种属性排序让同一个 warp 里的 32 个线程尽量走同一条路径这在处理不规则数据比如稀疏矩阵、粒子模拟时很有效。第三类是warp 级协作把原本需要分支的负载均衡问题改成所有线程合作为所有数据分流典型代表是 GPU 上的基数排序。从 Volta 开始的硬件在分支处理上还有一个新特性硬件会维护一个主动线程掩码active mask栈自动处理嵌套分歧的重汇聚。但收敛点不再保证一定在分支结束处所以如果你的代码里有依赖分支结束后所有线程自动同步这个假设的逻辑必须显式加__syncwarp()。4. 访存延迟隐藏的数学用 Littles Law 算出你到底要多少 warp这一章是全文的核心。前面所有关于调度器的讨论最终都要回答一个问题要多少并行度才能让内存带宽跑满。这个问题有精确的数学答案不需要凭感觉调参。4.1 各级存储的延迟量级先把数据准备好。下表是我在 A100 和 RTX 系列上实测得到的量级参考单位是 SM 时钟周期。不同架构、不同频率下会有浮动但数量级是稳的访问类型典型延迟周期备注寄存器访问0同周期直接接在流水线上FP32 加法依赖4 左右固定延迟可被其它 warp 掩盖共享内存 / L1 命中20 ~ 30走 MIO 管道short scoreboardL2 命中180 ~ 250跨 SM 共享HBM 访问未命中 L2400 ~ 800最常见的 long scoreboard 来源看这张表要抓住一个关键对比HBM 延迟是 L1 命中的 20 倍以上。这就解释了为什么很多 kernel 的性能瓶颈看起来是内存带宽实际上是延迟没藏住、带宽被浪费——如果每个 warp 发完请求就傻等LSU 和内存控制器之间会出现大量空档带宽利用率自然上不去。另外一个值得关注的点是延迟和频率的关系。HBM 延迟的绝对时间纳秒基本由内存控制器和 DRAM 时序决定但换算成周期数时SM 频率越高周期数就越大。所以同一张卡在标称频率和降频状态下能藏住延迟所需的 warp 数是不一样的。这也是为什么笔记本 GPU 和桌面版同型号的性能差异有时会比纸面参数更大。实测小技巧想量准你自己机器上的内存延迟可以写一个指针追逐pointer chasing微基准让一个 warp 沿随机地址链跳转用clock64()计时跳若干次后除以跳数。这个方法比查文档可靠得多因为实际延迟受显存配置、ECC、温度影响。4.2 把带宽需求折算成在飞字节数现在上公式。Littles Law 在排队论里的标准形式是系统中的平均并发数 到达率 × 平均停留时间。搬到 GPU 的场景里翻译一下就是需要的在飞字节数 目标带宽 × 访存延迟这个式子非常好用因为它把我需要多少并发变成了两个已知量的乘积。让我用 A100 的实际数字算一遍。A100 SXM 版本的 HBM2e 带宽峰值是 2039 GB/sSM 数量 108 个加速频率约 1.41 GHz。先算每个 SM 每个周期需要搬运多少字节每 SM 每周期字节数 2039e9 / (108 × 1.41e9) ≈ 13.4 bytes/cycle也就是说为了让 HBM 带宽满载每个 SM 平均每个周期要向内存系统提交 13.4 字节的请求。现在乘以延迟取 500 周期这个中间值在飞字节数 13.4 × 500 ≈ 6700 bytes结论每个 SM 需要大约 6.7KB 的数据持续悬停在访存管道里才能把 HBM 带宽吃满。现在把这个数字折算成 warp 级请求。一条 128-bit 的 warp 访存指令LDG.E.128请求的字节数是32 线程 × 16 bytes 512 bytes所以需要的并发请求数是6700 / 512 ≈ 13 条 warp 级访存指令同时在路上分摊到 4 个 SMSP每个 SMSP 大约需要3.3 条访存指令同时在飞。这个数字看起来很小但注意它的前提是每条在飞指令必须携带 512 字节。如果你用的是 32-bit 标量加载一条指令只带 128 字节那需要的并发指令数就要乘 4变成 13 条×4 52 条每 SM也就是 13 条每 SMSP。而每个 SMSP 最多只有 16 个 warp 槽位这意味着你必须让几乎每一个 warp 都时刻挂着一条访存请求才能勉强达到带宽峰值。没有任何容错余量。这个推导直接给出了两个优化方向要么增加每个 warp 携带的字节数向量化让一条指令带 512 字节而不是 128 字节要么增加每个 warp 同时悬挂的请求数展开循环造 MLP。两个方向本质是同一件事的两种做法。4.3 TLP 与 MLP两条腿走路上一节算出来的 6.7KB 在飞数据可以通过两种方式凑出来。这两条路在文献里分别叫 TLPThread Level Parallelism线程级并行和 MLPMemory Level Parallelism内存级并行我更喜欢把它们叫做多找几个人排队和让每个人多拿几件东西。TLP 路线就是提高 occupancy也就是增加常驻 warp 数。每个 warp 发一条访存指令然后等靠 warp 数量堆出并发请求。这条路的天花板很清楚每个 SMSP 最多 16 个 warp每 warp 一条在飞指令总共 16 条如果每条带 128 字节就是 2KB——远低于 6.7KB 的需求。这解释了为什么纯靠 occupancy 打不满带宽。而且现实更糟因为常驻 warp 里总有一部分在处理非访存指令地址计算、循环判断、结果消费真正挂着访存请求的比例更低。MLP 路线是让每个 warp 一次发出多条独立的访存请求然后等所有数据都回来再一起消费。做法有两种循环展开unroll和向量化vectorized load。循环展开的效果最直观。原本的循环是发请求 → 等 → 用 → 下一轮展开 4 路之后变成发请求 1、2、3、4 → 等 → 四个都用。四个请求是背靠背发出的不需要等前一个返回于是同一个 warp 的在飞字节数从 128 涨到了 512。如果每个 SMSP 有 8 个 warp 处于这种状态就是 4KB 在飞配合另外 8 个 warp 的其他活动基本能顶到目标值。向量化是另一条更便宜的路径。把 4 次 32-bit 加载合并成 1 次 128-bit 加载指令数少了 4 倍但在飞字节数不变——因为 512 字节还是一次请求。等等这里要澄清一个常见误解向量化的主要收益不是增加在飞字节数而是减少指令数和提高访存效率。同样的 512 字节用一条LDG.128发出只占一个 LSU 流水线槽位、一次地址计算用四条LDG.32发出占四个槽位、四次地址计算还会消耗四倍的发射带宽。所以在发射端口成为瓶颈的场景下向量化收益巨大在纯延迟受限的场景下向量化配合展开才是最优解。我用一个具体的代码片段说明展开的写法。这是一个累加 kernel 的两种实现// 版本 A无展开每个 warp 每次只有一条在飞请求 __global__ void sum_naive(const float* __restrict__ x, float* __restrict__ out, int n) { int i blockIdx.x * blockDim.x threadIdx.x; int stride gridDim.x * blockDim.x; float s 0.0f; for (; i n; i stride) { s x[i]; // 加载后立即依赖warp 立刻 stalled } atomicAdd(out, s); } // 版本 B4 路展开一个 warp 同时悬挂 4 条独立请求 __global__ void sum_unroll4(const float* __restrict__ x, float* __restrict__ out, int n) { int base blockIdx.x * blockDim.x * 4 threadIdx.x; int stride gridDim.x * blockDim.x * 4; float s0 0.f, s1 0.f, s2 0.f, s3 0.f; for (int i base; i 3 * blockDim.x n; i stride) { s0 x[i]; s1 x[i blockDim.x]; s2 x[i 2 * blockDim.x]; s3 x[i 3 * blockDim.x]; // 四条加载互相独立硬件可以背靠背发射全部在飞 } atomicAdd(out, (s0 s1) (s2 s3)); }版本 B 的关键点在于四个累加器s0到s3是独立的。编译器看到s0 x[i]和s1 x[i blockDim.x]之间没有依赖关系就会把四条加载指令连续排布全部发出去之后再等结果。硬件层面四条请求同时在 LSU 和内存系统里穿行有效延迟被摊薄了。如果你写成s x[i] x[i1] x[i2] x[i3]那就退化成串行依赖了因为有单条依赖链。这里有个细节要提醒四个累加器不能让编译器优化掉或者合并成一个所以必须显式写成四个独立变量且在最后汇总。用#pragma unroll让编译器自动展开循环也能达到类似效果但编译器展开后的调度顺序不一定符合你的预期关键路径上我还是倾向于手写。4.4 occupancy 是入场券不是成绩单讲到这里我想强调一个反直觉但很重要的观点occupancy 只在它限制你的并发请求数时才有意义。当你的 kernel 处于带宽受限、MLP 不足的状态时提高 occupancy 有帮助因为更多 warp 意味着更多并发请求。但当你的 kernel 已经通过展开和向量化获得了足够的 MLP继续提高 occupancy 收益就很小了甚至可能有害——因为寄存器压力会让编译器减少寄存器用量导致指令调度空间变小、出现溢出访问local memory spill反而拖慢速度。我在做矩阵乘法的时候深有体会。一个经典的 128×128 tile 分块实现用__launch_bounds__(256, 2)把每个 SM 限制成 2 个块寄存器给到 128 个occupancy 只有 25%16 个 warp / 64。但它的性能远好于一个 occupancy 100% 的朴素版本因为它每个线程扛着 8×8 的输出累加器计算访存比极高共享内存访问模式也经过了精心设计。这时候 occupancy 低反而说明寄存器用得好。所以正确的思路是先用 ncu 判断瓶颈类型带宽受限 / 延迟受限 / 计算受限 / 发射受限再决定往哪个方向使劲。如果smsp__issue_active已经接近 100%说明发射端口满了提高 occupancy 没用得减少指令数。如果dram__throughput只有一半而 long scoreboard 很高那才是 MLP/TLP 的问题。5. 实操把一个带宽跑不满的 kernel 从 42% 拉到 88%前面都是理论这一章上真家伙。我拿一个实际手头做过的归约 kernel 做样本把每一步改动、每一次指标变化都记录下来。测试环境是一张 A100 40GB PCIe 卡CUDA 12.1驱动 530 系列数据量 256M 个 float1GB 数据。5.1 先搭测量基线别急着改代码第一步永远是量化。我用 ncu 跑一遍完整指标集重点关注三类数据吞吐、occupancy、stall 分布。# 编译时务必带上行号信息否则 ncu 无法把指标映射回源码 nvcc -O3 -lineinfo -archsm_80 reduce_bench.cu -o reduce_bench # 抓两组指标全局吞吐 per-SMSP 的调度与 stall 分布 ncu --kernel-name regex:sum_naive \ --metrics dram__throughput.avg.pct_of_peak_sustained_elapsed,\ sm__warps_active.avg.pct_of_peak_sustained_active,\ smsp__average_warps_issue_stalled_long_scoreboard_per_issue_active.ratio,\ smsp__average_warps_issue_stalled_not_selected_per_issue_active.ratio,\ smsp__issue_active.avg.pct_of_peak_sustained_active \ ./reduce_bench关于-lineinfo这个选项会略微增大二进制体积、对性能有极小影响但它让你在 ncu 里能把每一行源码和指标对应起来性价比远高于那一点点开销。生产构建里可以去掉。基线数据是这样的为了可读性我把指标名简化了指标基线值我的解读DRAM 吞吐占峰值42.3%带宽严重未跑满Achieved Occupancy99.6%说明不是线程数不够发射活跃度51.2%发射端口有一半时间闲着long scoreboard / issue6.8平均每条发射指令背着 6.8 个等内存的 warpnot selected / issue0.4就绪但没被选中的很少说明并行度不富余这组数据把问题定性得非常清楚occupancy 满、发射半空、stall 集中在 long scoreboard、并且没有多余的就绪 warp 可以调度。典型的 MLP 不足 TLP 无余量的组合。这时候继续加 block 数量是无效的因为 occupancy 已经到顶了得从每个 warp 的在飞请求数下手。5.2 第一刀展开循环把 MLP 从 1 提到 4按 4.3 节的思路我把 kernel 改成 4 路展开四个独立累加器。同时为了防止编译器把展开后的加载顺序打乱我用__restrict__明确告诉编译器x和out不重叠给它更多重排自由。改完之后立刻复测。指标变化是这样的指标基线4 路展开后变化DRAM 吞吐42.3%71.5%29.2 个百分点发射活跃度51.2%74.8%大幅提升long scoreboard / issue6.83.1降低约 55%not selected / issue0.42.2出现富余就绪 warplong scoreboard 降低是因为同一个 warp 的等待时间被多条请求平摊了——原本等一次 500 周期的数据才做 1 次加法现在等一次能填满 4 次加法单位有效工作对应的等待比例下降。not selected 上升是好消息说明调度器现在有得挑。但 71.5% 还没到位。继续往下挖。5.3 第二刀向量化 参数约束下一刀是两个动作叠加。第一把 4 次 32-bit 加载换成 1 次 128-bit 加载减少指令数、降低发射压力。第二加__launch_bounds__(256)约束块大小让编译器把寄存器控制在合理范围。向量化这部分有个必须注意的细节地址必须 16 字节对齐。我用了__builtin_assume_aligned和显式类型转换来保证这一点同时因为数据量是 256M 个 float整除 4所以尾部不用特殊处理。__global__ void __launch_bounds__(256) sum_vec4(const float4* __restrict__ x4, float* __restrict__ out, int n4) { int base blockIdx.x * blockDim.x * 4 threadIdx.x; int stride gridDim.x * blockDim.x * 4; float s0 0.f, s1 0.f, s2 0.f, s3 0.f; // 每个线程一次拿 4 个 float4即 16 个 float凑出 64 字节在飞 for (int i base; i 3 * blockDim.x n4; i stride) { float4 v0 x4[i]; float4 v1 x4[i blockDim.x]; float4 v2 x4[i 2 * blockDim.x]; float4 v3 x4[i 3 * blockDim.x]; s0 (v0.x v0.y) (v0.z v0.w); s1 (v1.x v1.y) (v1.z v1.w); s2 (v2.x v2.y) (v2.z v2.w); s3 (v3.x v3.y) (v3.z v3.w); } atomicAdd(out, (s0 s1) (s2 s3)); }这个版本每个 warp 每次循环悬挂的字节数是32 线程 × 4 条 float4 × 16 字节 2048 字节。对比基线的 128 字节提升了 16 倍。按 4.2 节算的需求每 SM 6.7KB只要有三四个 warp 处于这个状态就能满足。实测结果指标基线展开4向量化 展开4DRAM 吞吐42.3%71.5%88.4%发射活跃度51.2%74.8%62.1%long scoreboard / issue6.83.12.4指令总数相对值100%103%38%注意发射活跃度从 74.8% 掉到了 62.1%但带宽反而涨了。这完全符合预期向量化把指令数砍掉了 62%同样的工作量需要的发射次数大幅减少所以发射端口可以更闲而每个周期搬运的字节数更多了。这是一个从发射受限转向带宽受限的典型信号。5.4 第三刀网格配置与尾部效应88.4% 已经很接近了剩下的 11.6% 去哪儿了我抓了一次 timeline 视图发现两个问题。第一个是网格跨步循环的尾部效应。当 1GB 数据被 grid-stride 循环切分时最后一轮迭代中只有一部分线程还在工作其余都退出了。这时候 SM 上的活跃 warp 数骤降带宽自然掉下来。解决办法是让网格大小正好整除数据量或者让每个线程处理的元素数固定通过多次 kernel launch 覆盖剩余部分。我用的是第二种算出每个 SM 需要多少块、每块处理多少元素凑成一个不产生尾部的配置。第二个是块大小对 warp 调度的影响。我把块从 256 调到 512 再调到 128 各测了一遍得到的数据挺有意思块大小寄存器/线程Achieved OccupancyDRAM 吞吐12832100%86.1%25632100%88.4%51232100%87.2%10243275%79.6%128 略低是因为块数多、调度开销略大512 差距不大1024 明显掉下来因为每 SM 只能塞 2 个块块数太少导致调度器在不同阶段难以找到足够的就绪 warp而且寄存器压力让 occupancy 掉到了 75%。结论是 256 到 512 是这类内存密集型 kernel 的甜点区和 CUDA 官方 best practices 里的建议一致。最终定格在 88.4%剩下那 11.6% 主要是 DRAM 刷新开销、ECC 校验、以及 PCIe 版本的带宽上限PCIe 卡的实际可达带宽通常比 SXM 版本低一些。一个容易被忽略的点dram__throughput这个指标的分母是理论峰值而实际可用的持续带宽通常只有理论值的 90% 到 93%。所以当你的指标跑到 88% 时很可能已经接近物理极限了没必要再折腾。判断方法是对比一次cudaMemcpy的带宽——纯拷贝能达到的上限就是你的 kernel 能达到的上限。5.5 把优化前后的数据整理成对照为了让你有个整体印象我把四个版本的完整对比列在这里。所有数值都是同一张卡、同一份数据、连跑五次的稳定值版本DRAM 吞吐Occupancylong sb/issue相对耗时v0 朴素网格跨步42.3%99.6%6.81.00v1 4 路展开71.5%99.6%3.10.59v2 向量化 展开88.4%99.6%2.40.48v3 v2 网格调优88.4%99.6%2.30.47从 1.00 到 0.47一倍多的提速全程没有换算法、没有换精度纯粹是把并发请求数从 128 字节提到了 2048 字节。这就是延迟隐藏在实践里的分量。6. Nsight Compute 指标怎么读stall 原因速查表理论讲完了实操也做了但我发现真正卡住大多数人的不是不懂原理而是打开 ncu 看到一屏指标不知道怎么下手。这一章把最常用的 stall 指标整理成一张速查表再补充一些使用上的注意事项。6.1 stall 原因对照每个数字背后是什么ncu 会给出每个 SMSP 的平均 stall 分布这些指标名都很长我在下面统一简写。所有 stall 指标都是每条发射指令平均挂着多少个这种状态的 warp数值越大说明这种等待越普遍指标简写全名关键词含义典型对策long sbstalled_long_scoreboard等全局/本地内存返回展开造 MLP、向量化、提高 L1 命中、调整访问模式short sbstalled_short_scoreboard等共享内存、MIO 管道结果消除 bank conflict、减少 shared 访问、改用 shufflewaitstalled_wait等固定延迟指令FMA、SFU通常可被其它 warp 掩盖过高说明 occupancy 或 ILP 不足not selectedstalled_not_selected就绪但未被选中正常现象若远高于其它项说明并行度过剩math throttlestalled_math_pipe_throttle运算单元排队该用 Tensor Core 或降低精度了mio throttlestalled_mio_throttleMIO 指令队列满减少共享内存/特殊函数指令密度lg throttlestalled_lg_throttleLSU 队列满减少访存指令数向量化有效barrierstalled_barrier等__syncthreads检查负载均衡、减少同步次数imc missstalled_imc_miss常量缓存未命中常量内存访问模式要广播化别让不同线程访问不同地址no instructionstalled_no_instruction指令缓存未命中缩小 kernel 体积、减少展开规模、提高 i-cache 局部性dispatch stallstalled_dispatch_stall分发端口冲突少见于单纯原因一般是组合症状用这张表的时候有一个重要技巧先看占比最高的那一项而不是看绝对值。如果 long scoreboard 占全部 stall 的 60%那它就是主因如果它只占 15%而 math throttle 占 50%那你该去优化算术而不是访存。新手最容易犯的错误是看到 long scoreboard 数字大就去改访存结果发现另一项才是真正的瓶颈。另外ncu 里还有一组smsp__average_warps_issue_stalled_*_per_issue_active的指标和上面这些smsp__average_warps_issue_stalled_*_per_issue_active.ratio是同一族命名规则在不同版本里有变化但含义一致都是平均每条发射指令背后有多少个 warp 因为某个原因在等待。6.2 采样、replay 与测量误差ncu 的工作原理是内核重放它会拦住你的 kernel 执行采集一部分指标然后把 kernel 重新跑一遍再采集另一组。一个 kernel 可能被重放几十次。这带来几个实际影响必须知道。第一计时结果不代表真实耗时。ncu 报告的耗时可能比实际运行高几倍因为它要停下来做采集。想测真实时间必须用cudaEvent或者多次 launch 取平均别信 ncu 的 duration。第二重放会改变缓存状态。如果指标采集需要多次重放那第一次重放时 L2 是冷的后续几次可能命中了上一次留下的数据。ncu 通常会尝试在每次重放前清空缓存用--cache-control all但这本身也有开销。所以在看 L2 命中率这类指标时要留意这个因素。第三--set full非常慢。一个 kernel 全量采集可能被重放上百次跑几分钟很正常。我的习惯是先跑--set speedoflight或者只抓几个关键指标定位方向之后再针对性地做深度采集。这样迭代速度快很多。# 快速定位瓶颈类型只要 speedoflight 视图里的三行 ncu --set speedoflight -k regex:reduce ./reduce_bench # 定向采集只看访存延迟相关 ncu -k regex:reduce \ --metrics dram__bytes.sum,lts__t_sector_hit_rate.pct,\ smsp__average_warps_issue_stalled_long_scoreboard_per_issue_active.ratio \ ./reduce_bench # 采集并导出报告方便反复看 ncu -o report_01 --set detailed ./reduce_bench ncu-ui report_01.ncu-rep关于 ncu 版本不同 CUDA 版本附带的 ncu 支持的指标名会变。如果你从教程里抄来的指标名报错先跑一次ncu --query-metrics | grep scoreboard看看实际支持的名称。这个坑我踩过不止一次尤其是从 CUDA 11 的教程抄到 12 上跑。6.3 我常用的一套排查顺序经过这几年的使用我形成了固定的排查流程写下来供你参考。第一步看 Speed of Light。这个视图给了三个百分比Compute、Memory、DRAM。谁的百分比最高谁就是瓶颈候选。如果三个都低于 60%那多半是延迟问题或者同步问题。第二步看 Achieved Occupancy。如果它远低于 Theoretical Occupancy说明你在启动配置上限制了并行度块大小、共享内存、寄存器。如果它和理论值一样高但性能还是差那不是 occupancy 的问题。第三步看 warp state 统计。也就是前面那张 stall 表。找到占比最高的两项它们大概率能解释 80% 的性能损失。第四步看访存效率。l1tex__t_sectors_per_request这个比值能告诉你访存是否合并。理想值是 4一次 warp 的 128 字节请求正好覆盖 4 个 32 字节 sector。如果显著大于 4说明有跨行访问或者随机访问。第五步看源码级指标。有了-lineinfo之后ncu 会把指标映射到具体代码行。切换到 Source 视图按 stall 数排序直接能看到是哪几行的访存拖了后腿。这一步往往是最高效的因为它跳过了所有中间推断直接告诉你问题在哪一行。7. 踩坑记录与高频疑问这一章写点不那么教科书的东西。前面几章讲的是方法论但真正让我掉头发的是那些文档里不会写的细节。7.1 几个我实际踩过的坑坑一--use_fast_math让访存优化失效。有一次我为了提升浮点性能打开了这个开关结果发现带宽反而降了。排查半天才明白-use_fast_math会启用一系列激进的优化包括把float累加重排成允许更早收缩的形式编译器借此把原本独立的四个累加器重新串成了一条依赖链MLP 直接归零。这个教训是数学优化开关和性能优化目标不一定同向打开之前先测一次。坑二PTX 版本不匹配导致的性能漂移。在一次跨机器部署时我发现同一个二进制在两台配置相同的机器上性能差了 30%。原因是我编译时只指定了-archsm_80而没有指定 PTX 版本导致在一台驱动较早的机器上走了 JIT 编译路径。JIT 出来的 SASS 和 AOT 编译的 SASS 在指令调度上可能有差异进而影响延迟隐藏效果。正确做法是同时指定 SASS 和 PTXnvcc -gencode archcompute_80,codesm_80 \ -gencode archcompute_80,codecompute_80 \ -O3 kernel.cu -o kernel这样既有针对 sm_80 的原生二进制又保留了 compute_80 的 PTX 供更高版本架构 JIT兼顾了兼容性和性能。如果你的卡是 4060TiAdasm_89或者 4090编译目标至少要包含compute_89否则会走 PTX JIT 路径。坑三块大小改一下就从 88% 掉到 79%。前面表格里那个 1024 块大小的数据就是真实的教训。当时我天真地以为块越大越好调度开销越小结果 occupancy 从 100% 掉到 75%带宽掉了一大截。在内存密集型 kernel 上块大小的甜点区通常比计算密集型更窄因为你需要足够多的块来给调度器留出调度余量。坑四环境层面的坑不全是环境问题。有时候你会遇到cuda .run gzip: stdin: invalid compressed>