反向 Kernel A/B 优化通用方法论 —— ncu 指标解析 + 共同优化手段总结

日期：2026-07-22 GPU：H20-3e（sm_90a，Hopper）。本文是 Kernel A（串行 dstate 反向 scan） 和 Kernel B（并行 per-token 梯度，split 成 4 个小 kernel） 两条优化战役的横向提炼：把两边反复用到的同一批 ncu 指标讲透，把两边共享的优化套路总结成可复用的 playbook。

姊妹篇（逐阶段细节）：bwd-kernelA-beginner-guide-full.md、bwd-kernelB-beginner-guide-full.md。
术语从零扫盲：见 Kernel B 小白版第 0 部分（GPU 层级、三种内存、occupancy、bank conflict、 coalescing、cp.async、Tensor Core/MMA/cute/swizzle/LDSM、ncu/roofline/stall）。本文不重复扫盲，只做跨 kernel 的对照与提炼。

一句话读法：先看第 1 部分学会读 ncu 那几个数字，再看第 2 部分的"诊断流程图"知道该往哪修，第 3 部分是两边都验证过的具体手段，第 4 部分是两边都踩过的坑。

第 1 部分：两条战役共用的 ncu 指标解析

优化不是瞎试，是先用 ncu 定位墙、再对症下药。A/B 两边从头到尾就盯着下面这几个数字。逐个讲清 它是什么、数值往哪个方向算好、看到它该做什么。

1.1 occupancy（占用率）—— 第一个要看的数

定义：一个 SM 上实际同时活跃的 warp 数 ÷ 硬件上限。H20 每 SM 上限 64 warp，占用率 12.5% = 只有 8 个活跃 warp。
好坏方向：越高越好。高 = SM 上随时有别的 warp 顶上 → 一个 warp 卡住时 SM 不空转 → 延迟被藏住。
被谁限制：每个 SM 的 smem 和寄存器是固定的一池。一个 block 吃的 smem 越多 / 寄存器越多，一个 SM 能同时住的 block 就越少 → 活跃 warp 越少 → 占用率越低。smem 通常是主凶。
两边的同一病根：
- Kernel A 基线：smem/block = 124.48KB → 放 2 块要 249KB > 227KB 上限 → 1 block/SM，占用率 12.5%。
- Kernel B 基线（未拆的大融合 kernel）：smem = 176KB > 233KB/2 → 同样 1 block/SM，占用率 12.5%。
看到它该做什么：占用率是"smem 阶梯函数"——只有把 smem 砍过某个台阶（能多塞一个 block）才跳一档。 A 靠寄存器融合把 124→50KB（跨到 4 block/SM）；B 靠拆 kernel 把 176→70KB（跨到 3 block/SM）。

1.2 stall reason（停顿原因）—— warp 卡住时在等什么

ncu 会统计"warp 没发指令时，卡在什么原因上"，占比最高的就是当前的墙。A/B 反复出现这四个：

stall reason	含义	典型触发	归到哪类墙
Short Scoreboard	等一个短延迟片上访问（几乎总是 smem 读 LDS/LDSM）的结果	bank conflict 让 LDS 变慢/变多	smem 管线
Long Scoreboard	等一个长延迟 gmem 读 LDG 的结果	gmem 访问没被 cp.async/流水藏住	gmem 延迟
MIO Throttle	MIO（所有 smem 指令 + 类型转换走的管线）排满，发不进新指令	LDS/STS 太多、或 bf16 打包 F2FP 太多	smem 管线拥塞
Barrier	等 `__syncthreads` 全员到齐	跨 warp 数据交接太多、warp 太少干等	同步/占用率

好坏方向：这些都是"浪费的周期占比"，越低越好；但关键不是绝对值，是谁排第一——第一名就是这一轮该打的墙。
两边的用法：
- Kernel A 基线首要 stall = MIO throttle 31.4%，DRAM 仅 4% → 不是带宽受限，是 smem 管线受限 → 去修 bank conflict，不去动 gmem。
- Kernel A s3a（换了 ldmatrix 但没 swizzle）：Short Scoreboard 40.3% + MIO Throttle 32.1% = 72% 全在 smem → 确认瓶颈是 bank conflict 不是指令条数 → 逼出 swizzle。
- Kernel B split-B 路径：主 stall 是 gmem 的 Long Scoreboard（dS 流式读）→ 判定为 latency-bound，只有提占用率有用。

1.3 Warp Cycles per Instruction Issued —— 每条指令有多堵

定义：平均每发射一条指令，一个 warp 要耗多少周期。
好坏方向：越小越好（越接近"发一条走一条"）。变大 = 每条指令更堵，多半在等 smem/gmem。
经典反直觉案例（Kernel A s3a）：换 ldmatrix 后指令总数下降，但 Warp Cycles 32→46（更差） → 证明 kernel 不是被指令数量限制，而是被 bank conflict 限制 → 砍指令数治不了病。这是"别只看指令数" 的铁证。

1.4 L1/TEX、L1/shared 吞吐利用率（SoL）—— 找"真墙"

定义：某个硬件单元的吞吐用到理论峰值的百分比（Speed-of-Light）。ncu 里 L1/shared 85.8%、 Compute 42.6%、L1/TEX 97.5% 都是这类。
好坏方向：利用率最高的那个单元 = 墙。想提速就得减少打到那面墙的流量，或把负载挪到没到顶的单元。
陷阱：利用率高不一定等于在干有用的活。Kernel A s3a 的 L1/TEX 钉死 97.5%，但它忙的是反复串行重试 bank conflict（忙于重试，不是搬数据）→ 高利用率反而是"病"的信号。要结合 stall reason 一起判。
两边用法：Kernel B 用 L1/shared 85.8% vs Compute 42.6% 判定"墙是 smem 吞吐、不是算力" → 优化方向定为"砍 smem 访存流量"，而不是重写 GEMM。

1.5 bank-conflict 计数：`conf_ld` / `conf_st`、`LDS.U16` / `op_st`

定义：ncu 统计的共享内存冲突事件数。conf_ld=load 冲突、conf_st=store 冲突；text-mode （--page source）能进一步拆到 SASS 指令级：LDS.U16（bf16 标量读）、LDS（32 位读）、 LDGSTS（cp.async）、LDSM/STSM（矩阵读写）各贡献多少冲突 wavefront。
好坏方向：越低越好；但必须先确认它在关键路径上（见 1.2 的 stall 排序），否则修了不涨。
两边最有价值的一课（源码级归因）：
- Kernel A s11：把 1.99B 冲突按 SASS 拆开，发现 79% 来自 R1 没被 swizzle 碰过 的 token-strided LDS.U16，而所有 LDSM/STSM（MMA 操作数）= 0 冲突 → 修法从"给 MMA 操作数加 swizzle"改成 “pad R1 raw buffer”。这是 atom-guessing 与 source-attribution 的分水岭。
- 数值证据链：s11 后 conf_ld 1.89B→1.16B（−39%）、short_scoreboard 2.53→1.20（−53%）。

1.6 用哪些 ncu 命令拿到这些数（共用命令行）

# 1) 总览：占用率 + 各单元 SoL + 首要 stall（先跑这个定墙）
ncu --set full -k <kernel_name_regex> -c 1 ./bench

# 2) 只要 stall / 吞吐 / 占用率的精简 section（迭代时快）
ncu --section SpeedOfLight --section Occupancy \
    --section WarpStateStats -k <regex> -c 1 ./bench

# 3) bank-conflict 计数（load/store 冲突事件）
ncu --metrics \
  l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_ld.sum,\
  l1tex__data_bank_conflicts_pipe_lsu_mem_shared_op_st.sum \
  -k <regex> -c 1 ./bench

# 4) ★源码级归因：把冲突 wavefront 拆到 SASS 指令（s11 靠它翻案）
ncu --page source --import <report>.ncu-rep     # text-mode，逐 SASS 看 conflict 归属
#   或 GUI 里打开 .ncu-rep → Source page → 按 "L1 Wavefronts Shared Excessive" 排序

迭代纪律（两边共用）：永远同一 session、单变量改一处、重测一次。Kernel A 曾差点把 caching-allocator 的抖动误判成优化收益；A/B 测试不隔离就会追到噪声。

第 2 部分：共用的诊断流程（该往哪修）

两条战役收敛出同一张决策图。拿到一个 kernel，按这个顺序问：

①  占用率够吗？（ncu Occupancy）
    └─低（≤12.5%，1 block/SM）── 先修占用率：砍 smem 跨阶梯（见 3.1）。别的先别碰。
    └─够 ─────────────────────► 进 ②

②  首要 stall 是谁？（ncu WarpStateStats）
    ├─ Long Scoreboard（等 gmem）── gmem 是墙：cp.async 分组重叠 / 减 gmem 流量（见 3.4）。
    │                               ★此时去修 smem bank conflict 是白费（s7 教训）。
    ├─ Short Scoreboard / MIO（等 smem）── smem 是墙：进 ③
    └─ Barrier（等同步）── 跨 warp 交接太多 / warp 太少：register-fuse 消 barrier（3.1）或补并行度（3.5）。

③  smem 墙 —— 冲突到底在哪条指令？（ncu --page source 源码级归因）
    ├─ 冲突在 MMA 操作数（LDSM/STSM）── swizzle 读写成对（见 3.2、3.3）。
    ├─ 冲突在 raw 标量读（LDS.U16 token-strided）── pad 行 stride 破整除（见 3.2 的 padding）。
    └─ 已到 4-way 硬地板（16B-coalescing 锁死）── 换更小写粒度 + 互质 stride（见 3.2 末）。

④  smem 冲突耗尽、但某配置（尤其 B1/低 N）仍不动？
    └─ 是并行度饥饿，不是冲突：补 grid 并行（split-V / grid remap / context-parallel，见 3.5）。

两条铁律（贯穿全图）：

先占用率、后冲突：占用率没修好时，bank conflict 被延迟盖住看不出来；顺序反了会误判。
先定 stall、后动手：不看"绑定 stall 是谁"就去追某个 ncu 计数，十有八九追到不在关键路径上的红鲱鱼。

第 3 部分：两边共享的优化手段（可复用套路）

下面每条都是 A 和 B 都实际用过并验证有效的手段。列出"原理 / A 侧落点 / B 侧落点 / 收益量级"。

3.1 用"更省 smem 的写法"跨占用率阶梯（两边最大的赢）

原理：占用率被 smem 卡住时，唯一有效的动作是把 smem/block 砍到能多塞一个 block 的台阶。有两条路：
- register-fusion（寄存器融合）：把 fp32 中间缓冲从 smem 搬进寄存器（每线程私有、不占 smem 池）。
- kernel-splitting（拆 kernel）：把大融合 kernel 拆成几个小 kernel，每个只留自己那步的 smem，中间结果写回 gmem 传递。用"多写几趟 gmem"换"占用率翻倍"。
Kernel A 落点：register-fuse 5 个 GEMM 的中间量（dU/dvcorr/bdv/T1），干掉 64KB 的"大象"T1_smem → 124.48→50.75KB → 1→4 block/SM → 占用率 12.5%→50%。收益 −38/−37.5/−37%。
Kernel B 落点：把大融合 kernel 拆成 4 个（B1/B2/B0/B3），每个 smem ~70KB → 1→3 block/SM → 占用率 12.5%→37.5%。
共同前提：两者都在 latency/occupancy-bound 场景才净赚——延迟被占用率藏住，多出来的 gmem 写/寄存器压力被更高并行度盖过。若本就带宽受限，拆 kernel 反而亏。

3.2 swizzle / padding 消 bank conflict（smem 墙的主武器）

原理：让一个 warp 的 32 线程访问落到 32 个不同 bank。两种做法：
- swizzle（异或置换）：对 smem 地址做固定 XOR 变换，把本会撞同 bank 的元素散开。适合写地址可自由重排的操作数。
- padding（填充行 stride）：给每行尾部塞废元素，把行 stride 改成不是 32 的倍数，破掉"竖读一列时行行 bank 重合"。适合不能改写指令、只能改布局的 raw buffer。
Kernel A 落点：qd/kd/kr/dS 用 swizzle（3b/3c/3d，累计 −61%）；R1 的 raw q_sm/k_sm/gc_sm 用 padding（s11，stride 136/132，−7.7~10.6%）。
Kernel B 落点：dS 用 swizzle 转置对偶（见 3.3）；某处 fp32 buffer 行 stride 从 D 改 D+4=132 （%32=4，每行错开 4 bank，8 行落 8 bank）消 ~75% 冲突。
两个共同约束（选 pad 值时必守）：(1) 错开量 ≠ 0（真的破 %32 整除）；(2) 每行仍是 16 字节 cp.async 最小搬运块的整数倍（A 的 136 bf16=272B=16×17；B 的 132 fp32=528B=16×33）。
4-way 硬地板（两边都会撞的天花板）：bf16 token-strided 读 在 16B-coalescing 约束下，STRIDE/2 被迫是 4 的倍数 → 最多 8 bank = 4-way 冲突地板，swizzle/pad 都破不了。唯一出路 = 换更小写粒度： Kernel A s12 用 STRIDE=130（65 与 32 互质）+ 4B cp.async.ca（放弃 16B）把读散满 32 bank，代价是 4× LDGSTS，只在延迟绑定的 B1 净赚（−4.2%）。

3.3 transpose-dual：一份物理字节喂两个方向（省一次转置）

原理：一个矩阵常被"正着乘"又被"转置乘"。用互为转置的 K_INTER（正读 LDSM_N）/ MN_INTER （转置读 LDSM_T） swizzle atom，让两个方向命中同一份物理字节 → 存一次，正读+转置读都无冲突、零 padding。
Kernel A 落点：dS 的 DSb（K_INTER）/DSc（MN_INTER），s3d，−22.8/−17.9%。
Kernel B 落点：dS 存一次、B1（dvcorr/dkr）正读、B2（dgT-dot）转置读；sin 的 SwSinLayout/ SwSinLayoutT 双方向读同一份（B2 的 dqd/dkd 收缩 sin 的行用 LDSM_T，kdS 收缩 sin 的列用 LDSM_N）。
共同点：省的是"写 smem→同步→按另一方向读回"整套动作，且正确性由构造保证（store/load 引用同一 layout 对象 = 一个置换 + 它的逆 → 数据恒等，无需重新验证数值）。

3.4 cp.async 分组重叠：把 gmem 延迟藏到 R1 计算背后

原理：普通 gmem→smem 要经寄存器且线程必须等。cp.async 异步搬运：先把一批操作数 commit 不等，马上去算别的（R1 重建），等真正要用时再 wait。用 FIFO 分组控制"先算的先到"。
Kernel A 落点：Group A（q/k/gc）先 commit、wait<1> 后跑 R1 前半；Group B（inv/mqk/do）后 commit、 wait<0> 在 stage-2 前 → 把两批 gmem 读的延迟叠在 R1 计算背后（s8/s9，B1 −10.5%）。
Kernel B 落点：先把所有 bf16 GEMM 操作数（dS/U/INV/Mqk/do）cp.async 发出 + commit 不等，马上跑 R1，用完再 wait（B3 是双 group A/B）。
★关键分诊（两边一致）：只有当首要 stall 是 gmem Long Scoreboard 时，cp.async 才有用。若已是 occupancy-bound（延迟本就被占用率藏住），加 cp.async 流水纯是开销 → ~0% 或负（Kernel B 0719§4.1 的 l9 就是负收益；Kernel A s7 在 gmem 未 crush 前去修 shared 也白费）。

3.5 补并行度：split-V / grid remap（治"低 N/B1 冲突优化不动"）

原理：conflict 优化修的是"每个 warp 跑多快"；当 grid 里 block 太少（B=1、N 小）填不满 SM 时，病在"有没有足够多 warp"，与 bank conflict 正交，swizzle 帮不上。得切出更多 block。
Kernel A 落点：split-V——把状态列轴 v 切给多个 CTA（blockIdx.z），grid (N,H)→(N,H,VSplit)， block 数 ×VSplit。代价是每个 v-band 重算一遍 R1。B1 基线大赢 −33%，饱和的 B32/B8 中性。
Kernel B 落点：拆成 4 个 (总tile数, H) 的 kernel，tile 间无依赖 → 天然打满并行；进一步的杠杆是 context-parallel（沿时间维切段给 B=1 补并行度，与 kernel 内优化正交）。
共同判据：“某配置在整个 conflict 战役里一直持平” = 它缺的是并行度不是速度，该上这类正交杠杆。

第 4 部分：两边都踩过的坑（反面教训）

别追不在关键路径上的 ncu 指标（A 的 s7/s10、B 的 0719§4.4）。先看绑定 stall 是谁；gmem 未 crush 前追 shared conflict、latency-bound 时追 issue-rate，都是红鲱鱼。source-level 归因（–page source）胜过按 atom 瞎猜（A s11 靠它发现 79% 冲突在 raw 读而非 MMA 操作数）。
swizzle atom 别贪大（A s10：dS INTER→SW32 全线回归 +11~12%）。INTER（zero-xor）对 16 宽子 tile 已最优；更大 atom 每次访问多花 XOR 地址运算 → 寄存器压力 → spill → 连不受冲突限制的 B1 都退（铁证不是冲突权衡）。
swizzle 必须读写成对，且只对"可自由重排写地址"的操作数做（A s6/s7）。R1 逐元素产出的可 swizzle （LDSM+STSM 成对）；cp.async 16B 载入的（mqk/inv/do）一 swizzle 就打碎 16B 向量化 → 要么保持普通、要么把 swizzle 白嫖进 cp.async scatter，绝不加独立的生产者 pass（A s7 plan-b 因此 +14.5%）。
向量化时别打乱"同一次乘法里 co-read 的 swizzled 操作数"的 bank pattern（B 0719§4.4 的 l13：把 dS 的 per-lane 分配改连续，顺手打乱了 sin_bf 的 swizzle bank pattern → 引新冲突 → +5~8%）。向量化收益要减去它给邻居带来的新冲突。
占用率是阶梯函数，砍 smem 没跨台阶 = 0 收益。只有把 smem 砍过"能多塞一个 block"那条线才跳档；砍了一半但没过台阶，占用率纹丝不动。
构建陷阱（必记）：s*/l* 都是被 launcher #include 的 header。改 header 不 touch 对应的 bwd_launch.cu 就不重编，stale .so 给假结果。改任何 header 后必须 touch launcher 再 build，且 A/B 测试同 session、单变量。

附：A/B 对照速查表

维度	Kernel A（串行 dstate scan）	Kernel B（并行 per-token 梯度）
grid	(N, H[, VSplit])	(总tile数, H)，拆成 4 个 kernel
基线占用率病根	smem 124.48KB → 1 block/SM 12.5%	smem 176KB → 1 block/SM 12.5%
跨占用率台阶的手段	register-fusion（消 T1_smem 大象）→ 50KB/4 block	kernel-split（4 个 ~70KB）→ 3 block
首要 stall（基线）	MIO throttle 31.4%（smem 墙）	Long Scoreboard（gmem 流 dS，latency-bound）
swizzle 主战场	qd/kd/kr/dS + R1 raw buffer padding	dS 转置对偶 + sin 双方向 + fp32 buffer padding
transpose-dual	dS 的 DSb/DSc	dS 存一次两读 + sin 双方向
cp.async	Group A/B FIFO 重叠 R1	全操作数发出不等、用完再 wait
补并行度杠杆	split-V（切 v 列轴）+ context-parallel	4-kernel 天然并行 + context-parallel
撞到的硬地板	bf16 token-strided 4-way（s12 用 stride-130+4B 破）	occupancy-bound（只有提占用率有用）
累计收益	−71.8/−69.4/−55.4%（B32/B8/B1）	见 kernelB 指南（拆分 + l 系列占用率实验）

一句话总纲：两条战役是同一套方法论的两个实例——先用 ncu 定墙（占用率→stall→源码级冲突归因），先修占用率再修冲突，swizzle/pad 读写成对且只碰可重排的操作数，cp.async 只在 gmem-bound 时上，低 N 配置的病在并行度要用 split-V/CP 这类正交杠杆治。

bwd-ncu-and-optimization-playbook