Star318
78
代码
代码
Issues76
Pull Requests55
Actions
项目讨论
Wiki
分析
Star318
78
已合并
knowledge(triton) 新增 A2 SparseFlashAttentionGrad triton 算子优化知识 #174
guomc创建于 8 天前
knowledge(triton) 新增 A2 SparseFlashAttentionGrad triton 算子优化知识 #174
已合并
guomc创建于 8 天前
从guomc9/cannbot-knowledgesparse-attn-grad-runbook-v2合入到cann/cannbot-knowledgemaster
共 6 个文件变更+271-138
Aknowledge/ops/triton/runbooks/performance/block_sparse_attention_grad_optimization.md+122-0
@@ -0,0 +1,122 @@
1+---
2+type: Runbook
3+title: BlockSparseAttentionGrad 算子优化
4+description: BlockSparseAttentionGrad 的 Triton-Ascend 优化：在 device 侧转置块掩码并构建 kv 主序选择表，避免 strided gather；dk/dv 使用连续写回与 0 轴归约，tile 按 UB 和 L0C 容量独立选取
5+tags: [task.optimization, topic.attention, topic.implementation.masking, topic.tiling, topic.memory, topic.precision]
6+status: draft
7+sources: []
8+created_at: '2026-10-27T00:00:00Z'
9+updated_at: '2026-10-28T00:00:00Z'
10+verified:
11+- by: process:opencode-deepseek-flash-tt36
12+ at: '2026-10-21T00:00:00Z'
13+platforms: [a2]
14+---
15+ 
16+# BlockSparseAttentionGrad 算子优化
17+ 
18+> 类别约束见 [attention_writing_classification.md](../../guides/constraints/attention_writing_classification.md)，
19+> 全局约束见 [global_constraints.md](../../guides/constraints/global_constraints.md)。
20+ 
21+## 任务特征
22+ 
23+- BNSD 布局、MHA、D=128，输入 dO、O 和 softmax 统计量，输出 dQ、dK、dV
24+- 块掩码为 uint8，形状是 [B,N,ceilQ,ceilKV]。实测 bx 为 16、32、48、64、128，by 为 16、32、64、80、128，两方向可以不同
25+- 评测共 17 case，覆盖 causal、滑窗、bigbird、随机和稠密掩码，包含非整除边界
26+ 
27+## 优化 1：在 device 侧构建选择表
28+ 
29+host 只做 shape、alloc 和 view。掩码转置与前缀计数都在 device 侧完成，不要物化 token 级稠密 mask 再传入。
30+ 
31+类别约束默认使用无 dot 的建表 kernel。下面的 2D-dot 是 Ascend 910B3、triton_ascend 3.6.0 和 CANN 9.1.0 下的实测例外，仅对应本 17 case 的选择表，前缀矩阵边长 T 不超过 128，不作为类别通用替代规则。
32+ 
33+```python
34+# ❌ host 侧物化 token 级稠密 mask
35+token_mask = block_mask.repeat_interleave(bx, 2).repeat_interleave(by, 3)
36+ 
37+# 本卡受限场景：比较矩阵配合一次 cube matmul 计算前缀位置
38+U = (cc[:, None] < cc[None, :]).to(tl.float16)
39+pos = tl.dot(keep.to(tl.float16), U).to(tl.int32)
zhanghz1
zhanghz1zhanghz18 天前

新 Runbook 推荐“比较矩阵配合一次 cube matmul”建表，并给出 tl.dot 示例；但它开头引用的类别约束第 74 行 (line 74)要求选择表“必须用无 dot 的 Triton kernel 建”。

同时消费两张卡时，会得到相反指令。

建议：在新 Runbook 中明确这是哪个环境、选择表规模下的实测例外，以及哪些旧约束仍适用；将 dot、比较归约和 cumsum 的对照条件写清楚。 若要升级为类别通用规则，再同步对应类别卡。

likedislike
guomc
guomc
7 天前 评论：
40+tl.store(sel_ptr + row * PAD + pos, col_ids, mask=keep)
41+```
42+ 
43+对照使用相同输入并产生相同选择表：tl.cumsum 为 0.13ms，3D broadcast-sum 为 0.17ms，2D-dot 前缀为 0.095ms。终版建表 kernel 约 78µs，其中约 30µs 是小 dot 的 cube 固定时延。表规模或编译环境变化后需要重新比较，tl.atomic_add 建表在该后端直接报 aivec error，不使用。
44+ 
45+掩码转置仍使用无 dot、无 atomic 的 kernel。device 侧建表、host 不物化稠密 mask、选择表与 kernel 粒度对齐等类别约束不变，program 粒度对齐掩码块粒度，不逐 token 发 program。
46+ 
47+## 优化 2：kv 主序选择表避开 strided 列 gather
48+ 
49+```python
50+# ❌ ub overflow 401KB，需要 3211264 bits，只有 1572864 bits
51+keep = tl.load(bmask_ptr + (bnidx * NQBLK + qts) * NKBLK + kb, mask=qts < NQB, other=0)
52+# 触发条件：gather 长度达到 128，或 gather 外面套了任意外层循环
53+ 
54+# ✅ 先用 device 侧 kernel 转置掩码，之后只读连续行
55+v = tl.load(bmask_ptr + base + j[None, :] * NKBLK + i[:, None], mask=mk, other=0)
56+tl.store(bmt_ptr + pid * NKBLK * NQBLK + i[:, None] * NQBLK + j[None, :], v, mask=mk)
57+```
58+ 
59+同样的紧致化流程，q 主序的连续 load 在 128 长度加 4 段循环下都能编过，只有 strided gather 爆 UB。转置是 uint8 小 tile 一次完成，实测结果与 mask.permute 逐位相等。
60+ 
61+sel_k 的列是 q tile 粒度，转置掩码的列是掩码块粒度。bx 大于 BMK 时必须做映射，否则 dk/dv 全错：
62+ 
63+```python
64+mqs = qts // (bx // BMK)
65+```
66+ 
67+## 优化 3：连续写回并沿 0 轴归约
68+ 
69+- dk/dv 写回走连续布局。实测连续写为 5.97ms，gather-store 为 107.6ms，atomic_add 为 156.2ms
70+- 归约沿 0 轴逐 lane 累加，快轴归约有 8~10 倍惩罚
71+ 
72+## 优化 4：dk/dv 独立选择 tile
73+ 
74+- BMK 必须整除 bx，从 64、48、32、16 中依次选择首个满足条件的值。实测 bx=48 时取 48，bx=128 时取 64；没有候选值时不套用本规则，需要重新做编译和正确性验证，不能直接使用 min(bx, 64)
75+- BT 取整除 by 的最大值，by 不超过 32 时取 by，否则依次试 32、48、16
76+- 同一套 BM=BN=64、BT=64 在 dq 能编过，在 dkdv 需要 224KB，直接越界
77+- L0C 住 cube 累加器，910B 为 128KB。两个 [128,128] fp32 累加器会触发 Cc overflow，dk/dv 双输出的 BT=128 不可用，需要降到 64 以下
78+- tl.dot 三参形态能把累加器放上 cube 复用，但双输出仍各占一个累加器，降 BT 才是生效的修复
79+ 
80+## 精度要点
81+ 
82+- q/k/v/dout 用原生 dtype 进 tl.dot，累加器用 fp32，不要把操作数升 fp32
83+- p 和 ds 是真 fp32 量，喂 dot 前做 hi+lo 二段拆分
84+- 边界列用 tl.where(valid, s, -3e38)，越界行地址先钳制再 mask
85+ 
86+## 正确性验收
87+ 
88+- 验证入口为 `.agents/skills/triton-op-verifier/scripts/verify.py --op_name BlockSparseAttentionGrad --verify_dir output/opt2_iter_2/verify`，运行环境为 Ascend 910B3、CANN 9.1.0、triton_ascend 3.6.0、torch 2.10.0+cpu、torch_npu 2.10.0.post7
89+- 本次 verify 使用 eco_open_source 判据，dQ、dK、dV 均要求 matched_ratio 不低于 0.99，并同时满足以下误差限制：
90+ 
91+| dtype | atol 与 rtol | 逐元素误差上限 |
92+|---|---|---|
93+| fp16 | 2^-9，约 1.953e-3 | max(0.1, 32·ULP) |
94+| bf16 | 2^-6，约 1.562e-2 | max(1.0, 32·ULP) |
95+ 
96+- 最终 `verify_result.json` 记录 17/17 通过。`cmp_case.py` 只用于 CANN 与 Triton 的核时对比，不做精度判定，正确性结论只来自上述 verify
97+ 
98+## 调试方法
99+ 
100+- aicore timeout 用消融定位，逐个禁用 kernel 跑问题 case，禁掉谁不 hang 谁就是元凶
101+- 输出 buffer 用 torch.empty 且没被 kernel 全覆盖写时，精度失败先查覆盖率再查数值
102+- tile 参数做成环境变量可调，扫描时触发 hang 的组合自动回退默认值
103+ 
104+### 总结
105+ 
106+1. 在 device 侧转置掩码并构建 kv 主序选择表，不用 strided gather
107+2. dk/dv 连续写回，归约沿 0 轴
108+3. dq 与 dk/dv 分别按 UB 和 L0C 容量选择 tile
109+ 
110+## 实测观测
111+ 
112+- **现象**：初版正确性通过，但性能远低于竞品
113+- **负载**：BNSD/MHA/D=128 块稀疏反向，17 case
114+- **解决层**：tiling + memory
115+- **严重度**：high
116+- **置信度**：multi_run
117+- **复现次数**：优化 10 轮，多轮复测
118+- **信号**：
119+ - 全 17 case 的 Triton 总核时从 5934.7ms 降到 21.67ms，提升 273.8 倍
120+ - 对竞品 geomean 0.9283，最低 0.6815，17/17 超过 0.6；4 个 case 反超，最高 3.41 倍
121+- **实测最优延迟**：全 17 case 总核时 21.67ms
122+- **证据**：910B3 实测，triton_ascend 3.6.0 + CANN 9.1.0；竞品为 ops-transformer 的 aclnnBlockSparseAttentionGrad，逐 case 独立进程加 profiler kernel_details 核时口径；官方 verify 17/17 通过
Mknowledge/ops/triton/runbooks/performance/index.md+2-1
@@ -16,4 +16,5 @@
16- [大规模 Sum 归约优化](reduction_sum_large.md)：大规模归约（sum）非reduce轴很大优化：计算重组减少归约次数，在优先占满UB前提下为reduce轴分配较大切分尺寸（BLOCK_SIZE_N=1024最优685.65us），适用于非reduce轴非常大（6万+）、reduce轴中等（千级）的2D归约场景16- [大规模 Sum 归约优化](reduction_sum_large.md)：大规模归约（sum）非reduce轴很大优化：计算重组减少归约次数，在优先占满UB前提下为reduce轴分配较大切分尺寸（BLOCK_SIZE_N=1024最优685.65us），适用于非reduce轴非常大（6万+）、reduce轴中等（千级）的2D归约场景
17- [Weighted SwiGLU Backward 融合算子优化](reduction_weighted_swiglu.md)：3D融合算子（Weighted SwiGLU Backward）优化：Reshape降维将前两维合并简化并行策略，行二次切分避免超UB，在优先占满UB前提下为reduce轴分配较大切分尺寸，grid数较大时可能性能更优，适用于3D张量逐元素+reduce融合的场景17- [Weighted SwiGLU Backward 融合算子优化](reduction_weighted_swiglu.md)：3D融合算子（Weighted SwiGLU Backward）优化：Reshape降维将前两维合并简化并行策略，行二次切分避免超UB，在优先占满UB前提下为reduce轴分配较大切分尺寸，grid数较大时可能性能更优，适用于3D张量逐元素+reduce融合的场景
18- [ChunkKdaFwd MFU 优化实测（分块 KDA 前向，1.995% → 5.152%）](chunk_kda_fwd_mfu.md)：12_ChunkKdaFwd 在 910B2C 上的 MFU 优化轨迹：起点 11.610 ms/1.995% 到 4.497 ms/5.152%，含逐项收益、kernel 拆解、求逆截断与守卫快速路径的实测数据与验证信号。18- [ChunkKdaFwd MFU 优化实测（分块 KDA 前向，1.995% → 5.152%）](chunk_kda_fwd_mfu.md)：12_ChunkKdaFwd 在 910B2C 上的 MFU 优化轨迹：起点 11.610 ms/1.995% 到 4.497 ms/5.152%，含逐项收益、kernel 拆解、求逆截断与守卫快速路径的实测数据与验证信号。
19-- [BlockSparseAttentionGrad 算子（block-sparse-attention-bwd） 实现手册](sparse_attention_grad_optimization.md)：BlockSparseAttentionGrad（BNSD/MHA/D=128、bx/by 16~128 可非对称、17 case）实测：kv 主序表若用 strided 列 gather 建表会 ub overflow（401KB，带循环或无循环 P2=128 均复现），改用 device 侧 uint8 掩码转置 kernel + 连续行 load 后一次编过；dk/dv kernel 必须独立选 tile（BMK=min(bx,64)；BT 取整除 by 的最大值：by≤32 取 by，否则依次试 32/48/16）否则 224KB 越界；17/17 精度通过，基线 geomean 0.034x（Phase 3 未优化）。19+- [BlockSparseAttentionGrad 算子优化](block_sparse_attention_grad_optimization.md)：BlockSparseAttentionGrad 的 Triton-Ascend 优化：在 device 侧转置块掩码并构建 kv 主序选择表，避免 strided gather；dk/dv 使用连续写回与 0 轴归约，tile 按 UB 和 L0C 容量独立选取
20+- [SparseFlashAttentionGrad 算子优化](sparse_flash_attention_grad_optimization.md)：SparseFlashAttentionGrad 的 Triton-Ascend 优化：先用微基准区分 gather 带宽与 dot 形状瓶颈，再用重数矩阵把梯度展开为大 GEMM；中间量按连续布局存放，tile 兼顾 L0C 容量与核间并行
Dknowledge/ops/triton/runbooks/performance/sparse_attention_grad_optimization.md+0-137
@@ -1,137 +0,0 @@
1----
2-type: Runbook
3-title: BlockSparseAttentionGrad 算子（block-sparse-attention-bwd） 实现手册
4-description: BlockSparseAttentionGrad（BNSD/MHA/D=128、bx/by 16~128 可非对称、17 case）实测：kv 主序表若用 strided 列 gather 建表会 ub overflow（401KB，带循环或无循环 P2=128 均复现），改用 device 侧 uint8 掩码转置 kernel + 连续行 load 后一次编过；dk/dv kernel 必须独立选 tile（BMK=min(bx,64)；BT 取整除 by 的最大值：by≤32 取 by，否则依次试 32/48/16）否则 224KB 越界；17/17 精度通过，基线 geomean 0.034x（Phase 3 未优化）。
5-tags: [task.optimization, topic.attention, topic.implementation.masking, topic.tiling, topic.memory, topic.precision]
6-status: draft
7-sources: []
8-created_at: '2026-10-21T00:00:00Z'
9-updated_at: '2026-10-22T00:00:00Z'
10-verified:
11-- by: process:opencode-deepseek-flash-tt36
12- at: '2026-10-21T00:00:00Z'
13-platforms: [a2]
14----
15- 
16-# BlockSparseAttentionGrad 算子（block-sparse-attention-bwd） 实现手册
17- 
18-> 类别约束见 [block_sparse_attention_constraints.md](../../guides/constraints/block_sparse_attention_constraints.md)，
19-> 全局约束见 [global_constraints.md](../../guides/constraints/global_constraints.md)。
20- 
21-## 任务特征
22- 
23-- 形态：三·2 块跳过 + `sp-bwd`（入参 dO/O/LSE，出 dQ/dK/dV），MHA（无 GQA），BNSD、D=128。
24-- 掩码：`blockMask[B,N,ceil(S/bx),ceil(S/by)] uint8`，bx/by ∈ {16,32,48,64,80,128}（可非对称），golden 用 token 级 `where(mask)` 表达。
25-- 17 case：S=1024~8192，fp16×16 + bf16×1，pattern causal/window/bigbird/random/dense，含 S=2000/4032 非整除边界。
26- 
27-## 骨架（5+1 kernel，全部 device 侧建表，host 只做 shape/alloc/view）
28- 
29-```
30-mask_transpose -> sel_q(q 主序) / sel_k(kv 主序) -> delta -> dq -> dkdv
31-```
32- 
33-- `sel_q[B,N,NRQ,NKB]`：行 = q tile（BM 行），列域 = kv 子块（BN 行）。
34-- `sel_k[B,N,NRK,NQB]`：行 = kv 子块（BT 行），列域 = q tile（BMK 行）。
35-- 两表都用**比较矩阵前缀计数**紧致化（`pos = sum(where(loc[:,None] > loc[None,:], keep, 0), 1)`），禁 `tl.cumsum`；计数值换算 fp16 可再省一半 UB。
36- 
37-## 坑 1（头号，必踩）：kv 主序表不要用 strided 列 gather
38- 
39-```python
40-# ❌ ub overflow 401KB（requires 3211264 bits while 1572864 available）
41-keep = tl.load(bmask_ptr + (bnidx * NQBLK + qts) * NKBLK + kb, mask=qts < NQB, other=0)
42-# 触发条件实测：gather 长度 >= 128；或 gather + 任意外层循环（即便 NCH=1）
43-```
44- 
45-最小复现结论：同样的紧致化流程，**q 主序的连续 load 在 P2=128/带 4 段 chunk 循环下都能编过**，只有「strided gather + 循环」或「strided gather 长度 128」爆 UB。
46- 
47-```python
48-# ✅ 先做一个 device 侧 uint8 掩码转置 kernel（纯布局，无 dot/atomic，tile = next_pow2(max(nqb,nkb))）
49-v = tl.load(bmask_ptr + base + j[None, :] * NKBLK + i[:, None], mask=mk, other=0) # [nkb, nqb]
50-tl.store(bmt_ptr + pid * NKBLK * NQBLK + i[:, None] * NQBLK + j[None, :], v, mask=mk)
51-# 之后 sel_k 只读连续行：tl.load(bmt_ptr + rbase + mqs, mask=qts < NQB)
52-```
53- 
54-转置 kernel 的 stried 访问是**一次**完成且 tile 很小（uint8，≤16KB），与「每个 (b,kt) 行都 gather 一遍」量级不同——实测转置结果与 `mask.permute(0,1,3,2)` 逐位相等。
55- 
56-## 坑 2：kv 主序表的列粒度映射
57- 
58-`sel_k` 的列是 **q tile（BMK 粒度）**，而转置掩码的列是**掩码块（bx 粒度）**。`bx > BMK` 时必须映射，否则索引越界且 dk/dv 全错（首版 maxabs=2.8、meanrel=12）：
59- 
60-```python
61-mqs = qts // (bx // BMK) # q tile -> mask block column，再 tl.load(bmt + rbase + mqs)
62-```
63- 
64-## 坑 3：dk/dv kernel 的 tile 必须独立于 dq
65- 
66-同一个 `BM=BN=64`、`BT=64`、D=128 在 dkdv 上需要 224KB（两个 `[BT,D]` fp32 累加器 + 4 个 `[BT,BM]` fp32 中间量 + q/dout 各两份）→ ub overflow。实测可编梯度：
67- 
68-| kernel | 配置 | 结果 |
69-|---|---|---|
70-| dq | BM=64, BN=64 / BN=80 | 可编 |
71-| dkdv | BMK=64, BT=32 | 可编（最大 case S=8192 亦通过） |
72-| dkdv | BMK=BM=64, BT=64 | ub overflow |
73- 
74-规则：`BMK = min(bx, 64)`（整除 bx）、`BT = by if by<=32 else 32 if by%32==0 else (48 if by%48==0 else 16)`（整除 by，保证 kv 子块唯一归属一个掩码块）。
75- 
76-## 精度配方（fp32 参考、fp16/bf16 输出，eco 口径 17/17 通过）
77- 
78-- `q/k/v/dout` 原生 dtype 进 `tl.dot`，fp32 累加器（禁升 fp32 operand）。
79-- `p = exp(s - lse)` 与 `ds = p*(dp - delta)` 是真 fp32 量：喂 dot 前 **hi+lo 二段拆分**（`ph = p.to(DTYPE); pl = (p - ph.to(f32)).to(DTYPE)`），dq/dk/dv 三路都做。
80-- 边界列 `tl.where(valid, s, -3e38)`（constexpr），越界行地址先 `min(idx, S-1)` 再 mask。
81-- 输出 lse/delta 全为**输入给定/fp32**，反向无除法，故无 NaN 风险；`cnt==0` 行自然得 0。
82- 
83-## 实测（Ascend 910B3 20 AIC / 40 AIV，triton_ascend 3.6.0（tt36）+ CANN 9.1.0 + torch 2.10.0 / torch_npu 2.10.0.post7）
84- 
85-| 项 | 值 |
86-|---|---|
87-| 精度 | 17/17（eco：fp16 atol=rtol=2^-9，bf16 2^-6，matched_ratio≥0.99） |
88-| 基线 geomean（Phase 3，未优化） | **0.0340x**（impl 几何均值 185.87ms / fw 6.31ms） |
89-| kernel 热点 | `_dkdv_kernel` 181.6ms、`_dq_kernel` 146.9ms |
90-| 最差 case | bx48/by80 causal（0.0059）、dense S2048（0.0149） |
91- 
92-> 该基线是「正确性优先」形态（数据依赖 `range(cnt)` 循环、hi+lo 双份 dot、dkdv 里 q/dout 各 load 两次、无流水/多缓冲）。
93- 
94-## 最终优化结果（同环境，6 轮有效 + round2 追加 2 轮）
95- 
96-| 项 | 值 |
97-|---|---|
98-| 全 17 case triton 总核时 | 5934.7 ms → **21.67 ms（273.8x）** |
99-| 对 AscendC 标杆（aclnnBlockSparseAttentionGrad，ops-transformer） | geomean **0.9283**，min 0.6815，**17/17 > 0.6** |
100-| 反超标杆 case | 4 个（最高 3.41x） |
101- 
102-## 后续优化方向（实测结论）
103- 
104-原三个方向的验证/否决结果：
105- 
106-- ① 取消 q/dout 二次 load —— **已验证**：单遍 5-dot 融合（dq/dkdv 合并为每 pair 1 次迭代）+ dk/dv partial workspace + reduce，geomean 0.47 → 0.672 档。
107-- ② 重扫 `BN/BT` —— **已验证**：`tl.dot(a,b,acc)` 3 参把累加器放上 cube，dq tile 上界 128×128、dk/dv BMK=128；`BT=128`（dk/dv 输出行）不可用——两个 [128,128] fp32 累加器触发 `Cc overflow`。
108-- ③ 数据依赖循环流水 —— **否决/中性**：slot 外提、fp16 transpose、reduce 双子块一轮改动全部在噪声内，拒绝合入。
109- 
110-round2 追加（case7 专攻，0.5954 → 0.6815）实测结论：
111- 
112-- `_delta_kernel` 改 (bn, chunk) 仿射网格 + ROB 128→64 + 去 padding 行 `tl.where`（0.092 → 0.054 ms；ROB=128 会 ub overflow）。
113-- 搬运型 kernel（`_reduce_kernel` / `_fused_kernel`）program 数 `nc = 2*n_aic`：128MB copy 微基准 NC=20 → 1.18 TB/s、NC=40 → 2.29 TB/s（1.94x）；cube 密集 kernel 无此收益。
114-- 建表 78µs 中约 30µs 为小 dot 的 cube 固定时延，不可无损消除；`tl.cumsum`（0.13 ms）/ 3D broadcast-sum（0.17 ms）变体更慢，选 2D-dot 前缀（0.095 ms）；`tl.atomic_add` 建表在该后端直接 aivec error，禁用。
115-- 仍未实施：partial 降 fp8 / k-v 驻 L2 重算、手写 CV 流水（PIPE_STAGES=2，天花板约再 1.5-2x）。
116- 
117-## 小算子时延下限与调试判据（910B3 实测）
118- 
119-**小 dot 的 cube 固定时延**：块选择表构建 kernel 78µs 中约 30µs 是小 dot 的 cube 固定时延，不可无损消除（triton_ascend 3.6.0 / 910B3 实测）。同功能前缀计数实现对比：`tl.cumsum` 0.13 ms、3D broadcast-sum 0.17 ms、2D-dot 前缀 0.095 ms——选 2D-dot。`tl.atomic_add` 建表在该后端直接 aivec error，禁用。实测案例见 [BlockSparseAttentionGrad 实现手册](../../runbooks/performance/sparse_attention_grad_optimization.md)。
120- 
121-**调试判据**：输出 buffer 用 `torch.empty` 分配且未被 kernel 全覆盖写时，精度失败应先怀疑覆盖率（未写区域是脏数据）而非数值精度——先核 kernel 写出索引是否覆盖全部输出元素，再排查数值路径。
122- 
123-## Program 数 = DMA 并行流数（910B3 实测）
124- 
125-搬运型 kernel 的 GM 带宽随 program 数（grid）增长，超过物理核数仍有显著收益。同一 128MB copy kernel（[64,128] fp16 tile，triton_ascend 3.6.0 / Ascend 910B3 实测）：
126- 
127-| program 数 NC | 实测带宽 |
128-|---|---|
129-| 20（= n_aic） | 1.18 TB/s |
130-| 40（= 2 × n_aic） | 2.29 TB/s（1.94x） |
131- 
132-结论：大块连续搬运型 kernel（reduce/partial、delta、elementwise 归约）以 `nc = 2 * n_aic` 起步再扫；cube 密集 kernel 无此收益（瓶颈在 cube↔vector 同步时延，不吃 DMA 吞吐）。
133- 
134-## L0C（Cc）容量边界（cube 累加器，910B3 实测）
135- 
136-Cube 累加器驻留 L0C（910B 为 128KB）。两个 [128,128] fp32 累加器（2 × 64KB）即触发 `Cc overflow`：dk/dv 类双累加器 kernel 的输出行 `BT=128` 不可用，需降到 BT ≤ 64；`tl.dot(a, b, acc)` 3 参形态可把累加器放上 cube 复用，但 dk/dv 双输出仍各占一个累加器，BT=128 时两个 [128,128] fp32 累加器仍触发 Cc overflow，实际生效的修复是降 BT。
137- 
Aknowledge/ops/triton/runbooks/performance/sparse_flash_attention_grad_optimization.md+102-0
@@ -0,0 +1,102 @@
1+---
2+type: Runbook
3+title: SparseFlashAttentionGrad 算子优化
4+description: SparseFlashAttentionGrad 的 Triton-Ascend 优化：先用微基准区分 gather 带宽与 dot 形状瓶颈，再用重数矩阵把梯度展开为大 GEMM；中间量按连续布局存放，tile 兼顾 L0C 容量与核间并行
5+tags: [task.optimization, topic.attention, topic.implementation.masking, topic.tiling, topic.memory, topic.precision]
6+status: draft
7+sources: []
8+created_at: '2026-10-27T00:00:00Z'
9+updated_at: '2026-10-28T00:00:00Z'
10+verified:
11+- by: process:opencode-deepseek-flash-tt36
12+ at: '2026-10-22T00:00:00Z'
13+platforms: [a2]
14+---
15+ 
16+# SparseFlashAttentionGrad 算子优化
17+ 
18+> 类别约束见 [attention_writing_classification.md](../../guides/constraints/attention_writing_classification.md)，
19+> 全局约束见 [global_constraints.md](../../guides/constraints/global_constraints.md)。
20+ 
21+## 任务特征
22+ 
23+- BSND 布局、MQA，N2=1、D=512，输入 dO、O 和 softmax 统计量，输出 dQ、dK、dV
24+- 每个 q token 带 K 个选中索引，K 从 1024 到 8192
25+- 评测共 18 case，K 分四档，S1 与 S2 可不等长，覆盖 fp16 与 bf16
26+ 
27+## 优化 1：先用微基准定位瓶颈
28+ 
29+旧架构按 N1 维铺小 dot，一度认定瓶颈是 gather 带宽。微基准直接证伪：
30+ 
31+- 裸 gather 452~555 GB/s，纯连续读 1538 GB/s，而 prep kernel 有效带宽只有 46 GB/s
32+- 裸 matmul 128×128×64 达到 109 TFLOPS，而 N1=4 的 prep 只有 0.28 TFLOPS
33+ 
34+瓶颈是小 N1 的 dot 形状喂不饱 cube，不是访存。先量出硬件能力和当前 kernel 吞吐的数量级差，再决定是否重写算法。
35+ 
36+## 优化 2：dense-k 重构，梯度全部大 GEMM 化
37+ 
38+选择结构和重数矩阵都在 device 侧构建，host 只做 shape、alloc 和 view。把每个 q token 组的选中集合在稠密 k 域展开，显式计算重数矩阵 cnt[bt, kk]：
39+ 
40+- 同一个 kk 被多个 j 选中时，它的 s、p、dp、ds 完全相同，所以 Σ_{idx=kk} ds = cnt[kk] · ds_dense[kk]
41+- dq、dk、dv 全部变成 M、N、K 不小于 64 的大 GEMM。低密度的稠密化浪费最多 8 倍，被 GEMM 约 400 倍的效率差覆盖
42+- 中间量 R 按 [b][k-tile][m][j] 分块，qT/doT 按 [b][m/BKM][d][j] 分块，让 dk 的两个操作数都连续，dk 从 20~55ms 降到 2.5~6.5ms
43+ 
44+program 粒度对齐 q token 组粒度，不要逐 token 发 program。
45+ 
46+## 优化 3：连续写回并沿 0 轴归约
47+ 
48+- dk/dv 写回走连续布局。实测连续写为 5.97ms，gather-store 为 107.6ms，atomic_add 为 156.2ms
49+- 归约沿 0 轴逐 lane 累加，快轴归约有 8~10 倍惩罚
50+ 
51+## 优化 4：tile 兼顾 L0C 容量与并行度
52+ 
53+- dk/dv 沿 D 每多切一刀就要把 R 和 qT 重读一遍。ND 取最小切分，使 fp32 累加 tile 不超 64KB，且 program 数不少于核数
54+- bk 在 MM 不超 128 时取 128。case11 实测 bk 从 64 升到 128、ND 从 8 降到 4，总时间从 8.64ms 降到 6.73ms
55+- 搬运型 kernel 的 program 数按两倍核数起步。同一 128MB copy kernel，20 个 program 实测 1.18 TB/s，40 个为 2.29 TB/s；cube 密集 kernel 没有这项收益
56+ 
57+## 精度要点
58+ 
59+- q/k/v/dout 用原生 dtype 进 tl.dot，累加器用 fp32，不要把操作数升 fp32
60+- cnt 用 uint8 会在 mode3 溢出，改用 int16
61+- 转置 kernel 的末块必须补 mask，避免越界
62+- qT 分块与 dk 的 BKM 必须对齐，强制 BKTQ 不超 BKM，且 BKM 整除 BKT
63+- offs_m 整除 N1 在 N1=32 会被后端错编译，该形状走展开布局
64+ 
65+## 正确性验收
66+ 
67+- 验证入口为 `.agents/skills/triton-op-verifier/scripts/verify.py --op_name SparseFlashAttentionGrad --verify_dir output/opt4_iter_1/verify`，运行环境为 Ascend 910B3、CANN 9.1.0、triton_ascend 3.6.0、torch 2.10.0+cpu、torch_npu 2.10.0.post7
68+- 本次 verify 使用 eco_open_source 判据，dQ、dK、dV 均要求 matched_ratio 不低于 0.99，并同时满足以下误差限制：
69+ 
70+| dtype | atol 与 rtol | 逐元素误差上限 |
71+|---|---|---|
72+| fp16 | 2^-9，约 1.953e-3 | max(0.1, 32·ULP) |
73+| bf16 | 2^-6，约 1.562e-2 | max(1.0, 32·ULP) |
74+ 
75+- 最终 `verify_result.json` 记录 18/18 通过。性能脚本只做核时对比；`cmp_opt/bench_case.py` 使用已验证实现的缓存 golden 和 eco_limits 做快速回归，不能替代从参考实现重新生成 golden 的完整 verify
76+ 
77+## 调试方法
78+ 
79+- aicore timeout 用消融定位，逐个禁用 kernel 跑问题 case，禁掉谁不 hang 谁就是元凶
80+- 输出 buffer 用 torch.empty 且没被 kernel 全覆盖写时，精度失败先查覆盖率再查数值
81+- tile 参数做成环境变量可调，扫描时触发 hang 的组合自动回退默认值
82+ 
83+### 总结
84+ 
85+1. 先用微基准区分访存瓶颈与算力瓶颈，本算子的瓶颈是 dot 形状
86+2. 用重数矩阵把稀疏梯度展开为大 GEMM，中间量按消费顺序连续存放
87+3. 写回连续化，归约沿 0 轴，tile 同时满足 L0C 容量与核间并行要求
88+ 
89+## 实测观测
90+ 
91+- **现象**：初版正确性通过，但性能远低于竞品
92+- **负载**：BSND/MQA/D=512 topK 索引反向，18 case
93+- **解决层**：tiling + memory + 算法重构
94+- **严重度**：high
95+- **置信度**：multi_run
96+- **复现次数**：优化 4 轮，多轮复测
97+- **信号**：
98+ - geomean 轨迹为 0.2432 → 0.3493 → 2.1549 → 2.5626 → 2.6177
99+ - 全 18 case 总核时从 6764.3ms 降到 199.5ms，提升 33.9 倍
100+ - 18/18 超过 0.6，最低 case11 为 0.7424，最高 case13 为 6.96 倍
101+- **实测最优延迟**：全 18 case 总核时 199.5ms
102+- **证据**：910B3 实测，triton_ascend 3.6.0 + CANN 9.1.0；竞品为 ops-transformer 的 npu_sparse_flash_attention_grad，逐 case 独立进程加 profiler kernel_details 核时口径；官方 verify 18/18 通过
Mlogs/2026-10-27.md+33-0
@@ -1,5 +1,16 @@
1# 2026-10-271# 2026-10-27
2 2 
3+## [20:23] knowledge | 稀疏反向 runbook 统一接入 attention 写法分类
4+ 
5+### Summary
6+ 
7+按评审意见移除新增的 sparse_flash_attention 设计期约束卡，两张稀疏注意力反向 runbook 统一链接 attention 家族写法分类。
8+ 
9+### Details
10+ 
11+- 删除约束卡及 constraints 索引条目，保留优化实测内容在对应 runbook 中。
12+- 既有 block_sparse_attention 约束卡不改动。
13+ 
3## [20:16] knowledge | PR158 候选整合与验证14## [20:16] knowledge | PR158 候选整合与验证
4 15 
5### Summary16### Summary
@@ -37,6 +48,17 @@
37- 优化清单第 11–14 项统一继承 §3.5 的前提；Vector 搬运说明限定为原测实现的观测，目标工具链另行确认。48- 优化清单第 11–14 项统一继承 §3.5 的前提；Vector 搬运说明限定为原测实现的观测，目标工具链另行确认。
38- 固定旧模板来源使用 primary，移除旧 PR 的 supporting 项；保留原代码块、历史性能和精度记录，本轮未执行新的 NPU 实验。49- 固定旧模板来源使用 primary，移除旧 PR 的 supporting 项；保留原代码块、历史性能和精度记录，本轮未执行新的 NPU 实验。
39 50 
51+## [17:02] knowledge | 稀疏注意力反向经验拆分为两张算子卡
52+ 
53+### Summary
54+ 
55+将合并的稀疏注意力反向 runbook 拆为 BlockSparseAttentionGrad 与 SparseFlashAttentionGrad 两张算子卡，并补充 sparse_flash_attention 类算子设计约束。
56+ 
57+### Details
58+ 
59+- 处理决定：split；块掩码与 topK 索引形态分别维护，原有实测与验证信号归入对应算子卡。
60+- 同步 performance 与 constraints 目录索引，新增约束卡链接全局约束和对应 runbook。
61+ 
40## [11:12] governance | 日志时间倒序检查降为 P1 并修复主线乱序条目62## [11:12] governance | 日志时间倒序检查降为 P1 并修复主线乱序条目
41 63 
42### Summary64### Summary
@@ -50,3 +72,14 @@
50- 同步 Lint 日志检查调度与检查项说明，保留原有 `log.order` 问题码。72- 同步 Lint 日志检查调度与检查项说明，保留原有 `log.order` 问题码。
51- 26 项 Lint 行为和 FormatCheck 摘要测试通过，覆盖时间倒序、同时间并列、P1 不阻断以及非法时间仍为 P0；修改文件的 Ruff 和差异空白检查通过。73- 26 项 Lint 行为和 FormatCheck 摘要测试通过，覆盖时间倒序、同时间并列、P1 不阻断以及非法时间仍为 P0；修改文件的 Ruff 和差异空白检查通过。
52- 全库严格 Lint 检查 9052 个知识 Markdown 页面，P0/P1/P2 均为 0，日志时间顺序问题已清除。74- 全库严格 Lint 检查 9052 个知识 Markdown 页面，P0/P1/P2 均为 0，日志时间顺序问题已清除。
75+ 
76+## [10:55] knowledge | 稀疏注意力反向卡扩为两形态并按目录体例重整
77+ 
78+### Summary
79+ 
80+`runbooks/performance/sparse_attention_grad_optimization.md` 由 BlockSparseAttentionGrad 单算子手册扩写为稀疏注意力反向算子优化卡，并入 SparseFlashAttentionGrad 的 dense-k 重构等实测结论，结构与表述按 runbook 目录既有体例重整。
81+ 
82+### Details
83+ 
84+- 处理决定：update，不新建卡片。
85+- 验证：`knowledge_lint` P0=P1=P2=0；`sync-indexes --check` 无过期。
Mlogs/2026-10-28.md+12-0
@@ -14,6 +14,18 @@
14- 平台口径：`950PR/950DT` 的编译参数为 `--npu-arch 3510`，内部 target 为 `c310`（`base_dsl/arch.py` 的 `_NPU_ARCH_TO_TARGET` 与 README 支持矩阵），三者同族，`platforms: ['950']` 成立。14- 平台口径：`950PR/950DT` 的编译参数为 `--npu-arch 3510`，内部 target 为 `c310`（`base_dsl/arch.py` 的 `_NPU_ARCH_TO_TARGET` 与 README 支持矩阵），三者同族，`platforms: ['950']` 成立。
15- 验证：`knowledge_lint --strict-navigation` 全库 P0/P1/P2 均为 0；`sync-indexes --check` stale 0；Query 索引 build/full verify 通过；检索质量门禁 recall@k=0.560、MRR=0.466（72 用例）通过；`test_markdown_structure` 的上游存量失败经逐文件过滤后本目录命中 0 项。15- 验证：`knowledge_lint --strict-navigation` 全库 P0/P1/P2 均为 0；`sync-indexes --check` stale 0；Query 索引 build/full verify 通过；检索质量门禁 recall@k=0.560、MRR=0.466（72 用例）通过；`test_markdown_structure` 的上游存量失败经逐文件过滤后本目录命中 0 项。
16 16 
17+## [09:21] knowledge | 稀疏反向 runbook 补齐适用范围与验收依据
18+ 
19+### Summary
20+ 
21+按评审意见修正 BlockSparseAttentionGrad 的 BMK 整除规则和块粒度适用范围，明确 dot 建表的实测例外条件，并为两张稀疏注意力反向 runbook 补充正确性验收依据。
22+ 
23+### Details
24+ 
25+- BMK 改为从已验证候选中选择整除 bx 的值，不再把整个 16～128 区间写成已支持。
26+- dot 建表限定到实际编译环境和选择表规模，保留类别约束中的无 dot 默认规则。
27+- 两张卡分别记录 verify 入口、运行环境、fp16/bf16 判据与通过结果，区分性能脚本和正式精度验收。
28+ 
17## [03:26] update | 修订 Reg 矢量理论吞吐与单双发口径29## [03:26] update | 修订 Reg 矢量理论吞吐与单双发口径
18 30 
19### Summary31### Summary