Merged
docs: 摄入 cannbot skill 下面的reference 知识并补全 API 合同 #33
docs: 摄入 cannbot skill 下面的reference 知识并补全 API 合同 #33
Merged
汪超Add on Sep 14
Total: 701 files changed+25765-108
@@ -93,3 +93,6 @@ assert(assertFlag == 12);
93[ASSERT]` assertFlag == 12 ' at /home/.../test.cpp:3693[ASSERT]` assertFlag == 12 ' at /home/.../test.cpp:36
94```94```
95 95 
96+ 
97+ 
98+- [assert(AI CPU API):固定 reference 调用约束](../../guides/diagnostics/aicpu_assert_reference_contract.md)
@@ -110,3 +110,6 @@ fmt string 291
110fmt string 3.140000110fmt string 3.140000
111TEST 0xdfffd6fddd1c111TEST 0xdfffd6fddd1c
112```112```
113+ 
114+ 
115+- [printf(AI CPU API):固定 reference 调用约束](../../guides/diagnostics/aicpu_printf_reference_contract.md)
@@ -109,3 +109,6 @@ vector_bool mask = asc_create_mask_b16(PAT_ALL);
109asc_loadalign(src, src_addr); // src_addr是外部输入的Unified Buffer(UB)内存空间地址。109asc_loadalign(src, src_addr); // src_addr是外部输入的Unified Buffer(UB)内存空间地址。
110asc_bfloat162e1m2x2_rn(dst, src, mask);110asc_bfloat162e1m2x2_rn(dst, src, mask);
111```111```
112+ 
113+ 
114+- [asc_bfloat162e1m2x2 (废弃):固定 reference 调用约束](../../../../guides/regbase/asc_bfloat162e1m2x2_deprecated_reference_contract.md)
@@ -73,3 +73,6 @@ PIPE_V
73## 调用示例73## 调用示例
74 74 
75见[地址寄存器调用示例](../defs/type/data_type_definition.md#example4)。75见[地址寄存器调用示例](../defs/type/data_type_definition.md#example4)。
76+ 
77+ 
78+- [asc_create_iter_reg (废弃):固定 reference 调用约束](../../../../guides/regbase/asc_create_iter_reg_deprecated_reference_contract.md)
@@ -76,4 +76,6 @@ vector_uint8_t dst;
76vector_uint16_t src;76vector_uint16_t src;
77asc_pack(dst, src); // 将src的低8位写入dst的低半部分77asc_pack(dst, src); // 将src的低8位写入dst的低半部分
78asc_pack_v2(dst, src); // 将src的低8位写入dst的高半部分78asc_pack_v2(dst, src); // 将src的低8位写入dst的高半部分
79-```79+```
80+ 
81+- [asc_pack (废弃):固定 reference 调用约束](../../../../guides/regbase/asc_pack_deprecated_reference_contract.md)
@@ -80,3 +80,6 @@ asc_loadalign(src, src_addr); // src_addr是外部输入的Unified Buffer(UB
80vector_int8_t dst;80vector_int8_t dst;
81asc_squeeze_v2(dst, src, mask);81asc_squeeze_v2(dst, src, mask);
82```82```
83+ 
84+ 
85+- [AscSqueezeV2:固定 reference 调用约束](../../../../guides/regbase/asc_squeeze_deprecated_reference_contract.md)
@@ -168,3 +168,6 @@ int main()
168 return passed ? 0 : 1;168 return passed ? 0 : 1;
169}169}
170```170```
171+ 
172+ 
173+- [AscLoad:固定 reference 调用约束](../../../../../guides/regbase/asc_load_reference_contract.md)
@@ -217,3 +217,6 @@ int main()
217 return passed ? 0 : 1;217 return passed ? 0 : 1;
218}218}
219```219```
220+ 
221+ 
222+- [AscDiv:固定 reference 调用约束](../../../../../guides/regbase/asc_div_reference_contract.md)
@@ -189,3 +189,6 @@ int main()
189 return passed ? 0 : 1;189 return passed ? 0 : 1;
190}190}
191```191```
192+ 
193+ 
194+- [AscCeil:固定 reference 调用约束](../../../../../guides/regbase/asc_ceil_reference_contract.md)
@@ -219,3 +219,6 @@ int main()
219 return passed ? 0 : 1;219 return passed ? 0 : 1;
220}220}
221```221```
222+ 
223+ 
224+- [AscFloat2Half:固定 reference 调用约束](../../../../../guides/regbase/asc_float2half_reference_contract.md)
@@ -189,3 +189,6 @@ int main()
189 return passed ? 0 : 1;189 return passed ? 0 : 1;
190}190}
191```191```
192+ 
193+ 
194+- [AscFloor:固定 reference 调用约束](../../../../../guides/regbase/asc_floor_reference_contract.md)
@@ -189,3 +189,6 @@ int main()
189 return passed ? 0 : 1;189 return passed ? 0 : 1;
190}190}
191```191```
192+ 
193+ 
194+- [AscRint:固定 reference 调用约束](../../../../../guides/regbase/asc_rint_reference_contract.md)
@@ -189,3 +189,6 @@ int main()
189 return passed ? 0 : 1;189 return passed ? 0 : 1;
190}190}
191```191```
192+ 
193+ 
194+- [AscRound:固定 reference 调用约束](../../../../../guides/regbase/asc_round_reference_contract.md)
@@ -195,3 +195,6 @@ int main()
195 return passed ? 0 : 1;195 return passed ? 0 : 1;
196}196}
197```197```
198+ 
199+ 
200+- [AscTrunc:固定 reference 调用约束](../../../../../guides/regbase/asc_trunc_reference_contract.md)
@@ -233,3 +233,6 @@ int main()
233 return passed ? 0 : 1;233 return passed ? 0 : 1;
234}234}
235```235```
236+ 
237+ 
238+- [AscCopy:固定 reference 调用约束](../../../../../guides/regbase/asc_copy_reference_contract.md)
@@ -205,3 +205,6 @@ int main()
205 return passed ? 0 : 1;205 return passed ? 0 : 1;
206}206}
207```207```
208+ 
209+ 
210+- [AscGather:固定 reference 调用约束](../../../../../guides/regbase/asc_gather_reference_contract.md)
@@ -215,3 +215,6 @@ int main()
215 return 0;215 return 0;
216}216}
217```217```
218+ 
219+ 
220+- [AscArange:固定 reference 调用约束](../../../../../guides/regbase/asc_arange_reference_contract.md)
@@ -173,3 +173,6 @@ int main()
173 return passed ? 0 : 1;173 return passed ? 0 : 1;
174}174}
175```175```
176+ 
177+ 
178+- [AscCreateMask:固定 reference 调用约束](../../../../../guides/regbase/asc_create_mask_reference_contract.md)
@@ -188,3 +188,6 @@ int main()
188 return passed ? 0 : 1;188 return passed ? 0 : 1;
189}189}
190```190```
191+ 
192+ 
193+- [AscUpdateMask:固定 reference 调用约束](../../../../../guides/regbase/asc_update_mask_reference_contract.md)
@@ -252,3 +252,6 @@ int main()
252 return passed ? 0 : 1;252 return passed ? 0 : 1;
253}253}
254```254```
255+ 
256+ 
257+- [AscDeintlv:固定 reference 调用约束](../../../../../guides/regbase/asc_deintlv_reference_contract.md)
@@ -225,3 +225,6 @@ int main()
225 return passed ? 0 : 1;225 return passed ? 0 : 1;
226}226}
227```227```
228+ 
229+ 
230+- [AscIntlv:固定 reference 调用约束](../../../../../guides/regbase/asc_intlv_reference_contract.md)
@@ -197,3 +197,7 @@ int main()
197}197}
198```198```
199 199 
200+ 
201+ 
202+- [AscSelect:固定 reference 调用约束](../../../../../guides/regbase/asc_select_reference_contract.md)
203+- [比较与选择:固定 reference 调用约束](../../../../../guides/regbase/compare_and_select_reference_contract.md)
@@ -188,3 +188,7 @@ int main()
188 return passed ? 0 : 1;188 return passed ? 0 : 1;
189}189}
190```190```
191+ 
192+ 
193+- [AscSqueeze:固定 reference 调用约束](../../../../../guides/regbase/asc_squeeze_reference_contract.md)
194+- [比较与选择:固定 reference 调用约束](../../../../../guides/regbase/compare_and_select_reference_contract.md)
@@ -192,3 +192,6 @@ int main()
192 return passed ? 0 : 1;192 return passed ? 0 : 1;
193}193}
194```194```
195+ 
196+ 
197+- [AscUnsqueeze:固定 reference 调用约束](../../../../../guides/regbase/asc_unsqueeze_reference_contract.md)
@@ -167,3 +167,6 @@ int main()
167 return passed ? 0 : 1;167 return passed ? 0 : 1;
168}168}
169```169```
170+ 
171+ 
172+- [AscReduceMax:固定 reference 调用约束](../../../../../guides/regbase/asc_reduce_max_reference_contract.md)
@@ -167,3 +167,6 @@ int main()
167 return passed ? 0 : 1;167 return passed ? 0 : 1;
168}168}
169```169```
170+ 
171+ 
172+- [AscReduceMin:固定 reference 调用约束](../../../../../guides/regbase/asc_reduce_min_reference_contract.md)
@@ -197,3 +197,6 @@ int main()
197 return passed ? 0 : 1;197 return passed ? 0 : 1;
198}198}
199```199```
200+ 
201+ 
202+- [AscReduceSum:固定 reference 调用约束](../../../../../guides/regbase/asc_reduce_sum_reference_contract.md)
@@ -235,3 +235,6 @@ int main()
235 return passed ? 0 : 1;235 return passed ? 0 : 1;
236}236}
237```237```
238+ 
239+ 
240+- [AscStore:固定 reference 调用约束](../../../../../guides/regbase/asc_store_reference_contract.md)
@@ -167,3 +167,6 @@ int main()
167 return passed ? 0 : 1;167 return passed ? 0 : 1;
168}168}
169```169```
170+ 
171+ 
172+- [AscStoreUnalign:固定 reference 调用约束](../../../../../guides/regbase/asc_storeunalign_reference_contract.md)
@@ -227,3 +227,6 @@ int main()
227 return passed ? 0 : 1;227 return passed ? 0 : 1;
228}228}
229```229```
230+ 
231+ 
232+- [AscGather:固定 reference 调用约束](../../../../../guides/regbase/asc_gather_reference_contract.md)
@@ -115,3 +115,6 @@ __aicore__ inline void mmad_gemv_example(__gm__ half* a_gm, __gm__ half* b_gm, _
115 mmad(atom, l0c, l0a_tensor, l0b);115 mmad(atom, l0c, l0a_tensor, l0b);
116}116}
117```117```
118+ 
119+ 
120+- [基础APIGEMV:固定 reference 调用约束](../../../../../guides/blaze/tensor_api_gemv_reference_contract.md)
@@ -106,3 +106,6 @@ copy(fixpipe_atom, gm, l0c);
106| 128 | 256 | 512 | Disable | 8 | 4156 | 4225 |106| 128 | 256 | 512 | Disable | 8 | 4156 | 4225 |
107| 128 | 128 | 512 | Enable | 8 | 2108 | 2172 |107| 128 | 128 | 512 | Enable | 8 | 2108 | 2172 |
108| 128 | 128 | 512 | Disable | 8 | 2108 | 2172 |108| 128 | 128 | 512 | Disable | 8 | 2108 | 2172 |
109+ 
110+ 
111+- [UnitFlag:固定 reference 调用约束](../../../../../guides/blaze/tensor_api_unit_flag_reference_contract.md)
@@ -225,3 +225,6 @@ simt vf: d: (0, 1, 0), f: 3.140000, s: pass
225simt vf: d: (1, 1, 0), f: 3.140000, s: pass225simt vf: d: (1, 1, 0), f: 3.140000, s: pass
226......226......
227```227```
228+ 
229+ 
230+- [Printf:固定 reference 调用约束](../../../guides/diagnostics/ascendc_printf_reference_contract.md)
@@ -0,0 +1,31 @@
1+---
2+type: Guide
3+title: Attention 在线输出归一化与最终除法必须一致
4+description: 纠正设计中的二次除法,保存mask、分核和内存尺度条件。
5+tags:
6+- task.optimization
7+- topic.precision
8+- topic.implementation.tiling
9+- topic.memory
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-performance-best-practices/references/reduce/templates/fused_attention_online_softmax_design.md
13+ role: primary
14+created_at: '2026-09-11T07:01:45Z'
15+updated_at: '2026-09-11T07:01:45Z'
16+platforms: []
17+---
18+ 
19+# Attention 在线输出归一化与最终除法必须一致
20+ 
21+在线Attention可维护未归一化a:mNew=max(mOld,tileMax),alpha=exp(mOld-mNew),p=exp(score-mNew),lNew=lOld*alpha+sum(p),aNew=aOld*alpha+pV,最后a/l。也可每步维护已归一化O:ONew=OOld*(lOld*alpha/lNew)+(pV)/lNew,此时最终直接输出O。
22+ 
23+来源更新伪代码采用第二种已归一化形式,但外层结束又ElemwiseDiv(oUb,lUb),会二次除法。取两个score都0、V为2/4,则l2、正确O3,源码组合最后变1.5。必须选定并保持一种表示;章节称O_acc并不能代替数学合同。若输出L接口实际期望logsumexp则应m+log(l),不能把示例写L=lFinal直接当通用Attention接口。
24+ 
25+causal/window判断要统一含尾或开区间:以query/key含尾[q0,q1]/[k0,k1]、合法k≤q时,k0>q1整块跳过,k1≤q0整块有效,其余逐元素mask。滑窗q-W<k≤q时,k1≤q0-W可跳过过旧块,边界等号与真实接口preTokens/nextTokens定义对照。source理论近50%只适用等分方阵因果块计数,不同S实测百分比无环境证据不推广。
26+ 
27+m/l/a可由同一核在同次launch跨S2保留,GM workspace不是在线算法必需;跨核SplitKV、跨调用或状态放不下时才需要显式暂存/归并。MTE2加载UB/L1各需对应管线,PV既可Cube也可Vector实现,不能说910B必由Vector计算PV或固定只有8核。decode也不必绝对memory-bound,取决于batch、并行度、KV布局及计算。
28+ 
29+内存要按当前query tile Sq、KV tile Sk和D计量:score Sq*Sk、输出状态Sq*D、m/l各Sq,以及Q/K/V、float转换、workspace和buffer份数;不物化S²中间矩阵不等于所有配置都固定65KB。来源单Q行D64/Sk128表逐项合计约68224B,D128约134400B,未含未列buffer/对齐/队列,不能凭总计和192KB断言任意平台安全。全局Q/K/V/O内存仍随S*D增长。
30+ 
31+向量Mul+Reduce和标量PV只是可表达的基线,不是生产高性能证明。收益表明确预计,单launch省调度不能自动保证1.2–5倍。全掩码首块、首尾query/key、精度、跨tile最大值上涨与真实端到端性能需单独验证。
@@ -0,0 +1,32 @@
1+---
2+type: Guide
3+title: 反向 Cube 融合需保留全部梯度路径
4+description: 区分 dS 与概率矩阵,补齐累加器、转置输入和输出写回的寿命。
5+tags:
6+- task.optimization
7+- topic.attention
8+- topic.memory
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/cards/P79_grad_cube23_fusion.md
12+ role: primary
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/playbooks/P79_grad_cube23_fusion.md
14+ role: supporting
15+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/preconditions/P79.yaml
16+ role: supporting
17+created_at: '2026-09-11T02:13:08Z'
18+updated_at: '2026-09-11T02:13:08Z'
19+platforms: []
20+---
21+ 
22+# 反向 Cube 融合需保留全部梯度路径
23+ 
24+Flash Attention反向中,若dS代表分数梯度、P代表概率矩阵,相应矩阵链包括dQ=dS*K、dK=dS^T*Q、dV=P^T*dY;scale、dropout、mask和head归并仍按目标梯度定义处理。来源用S同时指代部分矩阵,不能把dS与P视为同一数据而错误共享。
25+ 
26+来源headDim=64且sparseMode=0时调度Cube23,将dS与其转置分别载入L0A ping/pong,将K/Q载入L0B ping/pong,再维护dQ/dK各自结果。两种方向可以共享原始L1数据,但所示骨架仍各调用一次正常LoadData和转置LoadData,并未证明从两次搬运降为一次;实际收益可能来自复用、调度与重叠,需要指令/流量证据。
27+ 
28+该片段只写出dQ,未展开dK写回及dV计算,不能当作完整Backward实现。所有必需输出都应保留对应任务坐标、有效M/N、归约维度和最后写回。每个L0C累加链首次初始化,其余贡献累加;交替L0A/B并不要求或允许把同一结果的部分和随意分散到不同L0C后只取一份。
29+ 
30+来源四个L0C buffer、合计128KiB是特定块尺寸/份数预算,不能仅因d=64就认定必然需要或容得下。应计算dQ/dK/dV的真实累加dtype、矩阵布局、同时活跃结果和Fixpipe消费寿命,支持其他headDim的实现也可以采用另行证明的融合方案;来源条件只约束这条特化分支。
31+ 
32+原子dQ写回还需要目的初始化、正确输出类型、累加范围与退出原子状态,不能由SetAtomicType一行推导完整正确性。验证三条梯度、K/query尾部、GQA共享head归并、多个partial与循环排空,比较误差及总流量/时间;源码省略处未被本批编译或设备验证。
@@ -0,0 +1,36 @@
1+---
2+type: Guide
3+title: 因果窗口、GQA 与整块跳过的证明
4+description: 补全闭区间边界、query 坐标和空行语义,限定 Attention 性能案例。
5+tags:
6+- task.optimization
7+- topic.compute.reduction
8+- topic.implementation.tiling
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-knowledge/references/a3/algorithm_insights/attention_family.md
12+ role: primary
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-knowledge/references/a3/optimization_patterns/causal_block_skip.md
14+ role: supporting
15+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-knowledge/references/a3/proven_solutions/INDEX.md
16+ role: supporting
17+created_at: '2026-09-11T03:58:39Z'
18+updated_at: '2026-09-11T03:58:39Z'
19+platforms: []
20+---
21+ 
22+# 因果窗口、GQA 与整块跳过的证明
23+ 
24+MHA令Q/KV头数相等,MQA令KV头数为1,GQA在整除且组内共享KV的约定下令G=Hq/Hkv、kv_head=floor(q_head/G)。直接索引可避免repeat_interleave实体扩展,但头分组、实际布局与序列坐标必须对应。BSND、BNSD和BSH的轴/stride不同;query的绝对token位置包括本序列的块起点、偏移或缓存位置,不能用head偏移替代,也不能把块内行号单独当所有块的全局token。
25+ 
26+对标准j≤i的因果关系,Q闭区间[q0,q1]、K闭区间[k0,k1]:k0>q1时整块无效;k1≤q0时整块有效;其余逐元素mask。对角线有效,所以跳过条件不能改成≥。尾块q1需min(实际末行,块末行),有效KV块数还要限制在真实KV范围内。来源proven_solutions称无效块在“below diagonal”错误,应为该坐标约定的上三角。
27+ 
28+滑窗i-W<j≤i还排除过旧key:当k1≤q0-W时可证明整个块过旧。来源用严格<会少跳过等号边界但通常不漏合法值;区间若改成半开应重新推导。preTokens/nextTokens等接口参数的包含性、非等长Q/K位置对齐、稀疏模式与padding要按具体合同,不能仅套标准三角。
29+ 
30+若n个等大Q/K块且完整等长序列,有效块数n(n+1)/2,占全部n²的(n+1)/(2n),可跳过比例(n-1)/(2n)趋近1/2。n是块数而非无条件token数。不同tile、负载分配、尾块与流水固定开销会改变真实时间收益。来源S=512/2048/8192/32768报告约22/37/43/45.5%降低及1.003x→1.83x,是其910B3案例,未在本批复现,也不能把降低45.5%写成仅加速45.5%。
31+ 
32+只在被跳块对最终输出无贡献且无其他必须副作用时省掉加载、QK、softmax更新与PV;Cube/Vector的任务数、信号和workspace槽必须一起调整,避免一侧仍等待不存在的结果。exp(-inf)=0不意味着全为-inf的softmax必然输出0,整行无有效key需目标合同,见[在线状态与归并](online_softmax_state_and_partition_merge.md)。
33+ 
34+在线Attention同时重缩放max、指数和与加权输出;存已归一化O时应乘旧l*exp(m_old-m_new)/l_new,未归一化a则采用对应更新,不能混用。状态能常驻时无需每个KV块写GM,是否溢出到workspace由容量与调度决定。QK和PV都含矩阵乘,来源把softmax+PV全部指定给Vector不是通用Cube/Vector映射。
35+ 
36+Decode常为每请求单query但仍有batch/head并行与Split-KV选项;不能由Q长1断言完全无需Q侧调度或总是带宽瓶颈。Prefill处理一段query与完整有效KV,KV-cache和分页管理减少重算/碎片的作用不同。来源FlashAttention-2/3及PagedAttention简介仅为外部候选方向,GPU的warp专用化、FP8支持和A100 2x成绩不迁为Ascend已验证能力。
@@ -0,0 +1,34 @@
1+---
2+type: Guide
3+title: CV 交错流水的任务槽与数据槽
4+description: 明确 MM1、Vec1、MM2、Vec2 的三期排程、索引范围和跨核生产消费及回收关系。
5+tags:
6+- task.design
7+- topic.parallelism
8+- topic.implementation.synchronization
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/playbooks/P14_cv_pipeline_preload.md
12+ role: primary
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/preconditions/P14.yaml
14+ role: supporting
15+created_at: '2026-09-11T01:38:31Z'
16+updated_at: '2026-09-11T01:38:31Z'
17+platforms: []
18+---
19+ 
20+# CV 交错流水的任务槽与数据槽
21+ 
22+对任务 j 的依赖链 `MM1(j)→Vec1(j)→MM2(j)→Vec2(j)`,可以考虑把多个独立任务错期开工。来源的轮次安排为:轮 i 发起MM1(i),处理Vec1(i-1)及MM2(i-1),处理Vec2(i-2)。Vec1和MM2对同一任务仍有生产消费关系,不能因为列在同一轮就同时无等待地执行。
23+ 
24+若任务总数为 L,轮次0到L+1可排空这一特定排程,但每个阶段应显式满足任务索引在[0,L):MM1用i<L,Vec1/MM2用i>=1且i-1<L,Vec2用i>=2且i-2<L。L=0时不读任何任务信息。来源仅以isValid辅助判断,数组却未完整初始化,短任务及排空轮可能读到未写状态;所有元数据初态必须定义,数值范围守卫也不能省略。
25+ 
26+三份RunInfo分别保存i、i-1、i-2,可以用 i%3、(i+2)%3、(i+1)%3定位;这些是任务描述槽。workspace的两份数据槽是另一维度,不能因为有三份描述就证明两份数据足够。MM1、Vec1和MM2的结果应有明确区域、跨度和最后消费者;若MM1(i)与尚待Vec2消费的MM2(i-2)共用物理区,同为i%2会发生覆盖。
27+ 
28+每条边都要建立真实完成与可见性,例如Cube完成的结果发布给Vector、Vector产生的数据发布给Cube,以及最后消费者完成后归还槽位的反向通知。数据槽复用不能只依赖元数据isValid=false。参与AIC/AIV、子核映射、flag编号、pipe与API版本必须匹配;来源“先不加跨核同步,精度失败再补”没有正确性依据。
29+ 
30+空间按各阶段结果的实际字节数、数据槽数及需要私有副本的核组数量求和。源文`原workspace*2*coreNum`只有当“原”严格指每核单槽且所有区域都翻倍时才成立;全局基数不能重复乘coreNum。元数据缓存与输出running state也需计入各自存储域。缩为单数据槽后必须重新安排等待/复用,不是只改数字即可保持排程。
31+ 
32+Vec2包含浮点累计或在线Softmax状态时,还要保证任务顺序与统计基准一致。使用DMA原子模式的局部区域应按实际接口在完成后恢复普通模式,避免污染后续写入,但不能据此要求所有输出更新都用原子。完整流水需验证0/1/2/3/多任务、每槽首次使用、尾部两轮、空工作核与异步搬运结束。
33+ 
34+这是静态排程和数据生命期的设计知识,未声明该骨架已经在设备上编译运行,也不承诺30–50%收益。是否值得交错取决于阶段时长、数据共享、资源竞争和任务数量。
@@ -0,0 +1,28 @@
1+---
2+type: Guide
3+title: Decode 拼接搬运与 online tile 归约
4+description: 融合 nope/rope 的物理暂存,并保存两级点积、有效域与在线 softmax 状态。
5+tags:
6+- task.optimization
7+- topic.compute.composite
8+- topic.data_movement.layout_transform
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/discovered/disc_X4.md
12+ role: primary
13+created_at: '2026-09-11T02:29:01Z'
14+updated_at: '2026-09-11T02:29:01Z'
15+platforms: []
16+---
17+ 
18+# Decode 拼接搬运与 online tile 归约
19+ 
20+当每个key的nope与rope分量分别存储时,可用合法的两组带stride搬运,将对应tile写入同一UB行的两段,形成[T,Dqk]逻辑拼接,供后续点积读取。必须分别定义源行距、目的行距、字节单位、Dnope/Drope有效长度与padding,避免覆写相邻行。拼接后的ND布局并不会自动成为Cube需要的NZ,两次相邻调用也不等于一条物理DMA。
21+ 
22+一个KV tile搬入后,各key行与相同query做Mul,再通过块级部分和与整行合并形成score[r]。两级归约必须覆盖全部Dqk,包括非满块,并按真实物理跨度定位部分和;Mul→归约→下一Vector/Scalar消费者都有数据依赖。按tile批量搬入可将逐key搬运/同步的调用数从O(Skv)降到O(ceil(Skv/T))量级,前提是实现确实批量加载;计算部分仍可能逐行展开,不能只改循环名字即宣称完成。
23+ 
24+跨KV tile维护FP32的max m、分母l与未归一化输出o。合入分数s时,m'=max(m,max(s))、alpha=exp(m-m')、p=exp(s-m'),l'=alpha*l+sum(p),o'=alpha*o+pV;最终输出o/l。初始m=-inf,l=0,o=0,但空或全屏蔽tile需避免-inf-(-inf)并按语义处理。更新max后只重标l而漏掉o会改变结果,mask、scale、sink与V列顺序也必须保持。
25+ 
26+Vector的一lane Exp替代Scalar libm只是实现候选,仍有mask设置、Vector→Scalar同步和调用成本。V预取必须计入额外缓冲及释放时刻,不能为了减少score阶段时延覆盖后续PV所需数据。算法等价不意味着浮点逐位一致,需检查溢出、特殊值及长期在线误差。
27+ 
28+来源T=64约150KiB/192KiB、T=80更慢以及0.11–0.27到0.25–0.79的归一化性能都是特定配置,未在本批复现。验证D与Skv双尾、不同拼接长度、全masked与多tile状态,并计布局准备及最终归一化的完整代价。
@@ -0,0 +1,34 @@
1+---
2+type: Guide
3+title: 双 AIV 行范围与 Workspace 份数
4+description: 保存对齐分工公式、空消费者与物理槽容量,避免混淆核比例和核总数。
5+tags:
6+- task.design
7+- topic.parallelism
8+- topic.tiling
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/playbooks/P73_cv_parallel_aic_aiv_ratio.md
12+ role: primary
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/playbooks/P75_dual_aiv_workload_split.md
14+ role: supporting
15+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/preconditions/P73.yaml
16+ role: supporting
17+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/preconditions/P75.yaml
18+ role: supporting
19+created_at: '2026-09-11T02:07:16Z'
20+updated_at: '2026-09-11T02:07:16Z'
21+platforms: []
22+---
23+ 
24+# 双 AIV 行范围与 Workspace 份数
25+ 
26+Cube+Vector融合中可按实际执行模式让两个AIV消费一个Cube任务的不同M/S1行。均匀元素分割可用start=floor(i*W/V)、end=floor((i+1)*W/V)(V>0);乘法需足够宽,且W的单位须是可独立处理的工作项,不能拆断需要共同统计的一行而不增加正确归并。
27+ 
28+来源按16行对齐切分,令m<=16时v=m,否则 `v=ceil(ceil(m/16)/2)*16`,两份为[0,v)和[v,m)。第二份可以为0,如m=17时v=16而不是平均8/9;m<=16第二个AIV没有行工作,仍需遵守其同步参与角色。16是该路径布局候选,不是所有Vector的统一分裂粒度,M/S1是否同一轴应从实际布局确认。
29+ 
30+GetBlockIdx()%2只有在已验证全局核编号和成组关系时才能作为AIV组内编号;通常应依据具体执行模型的subblock接口与Host/编译配置。仅赋aicAivRatio字段不保证已启动所需Kernel类型,1:2也不必然减少芯片上的AIC物理核数。真正的吞吐取舍来自任务分配、资源及Cube/Vector相对工作量。
31+ 
32+两个消费者必须写互不冲突的逻辑范围,并各自有正确的workspace/UB地址。四份workspace是来源一项流水候选;从两份变四份是这部分容量翻倍,不是总GM必然变四倍。总量以每槽真实shape/dtype、对齐、并发任务及读写寿命计算。槽序号、跨核发布/回收和Vector double buffer需一起调整,四份本身不保证消除气泡。
33+ 
34+验证W/m为0、1、16、17、31、32、33,空消费者信号、首轮填充/末轮排空以及核组身份。测每阶段等待和完整时延,不能用出现后处理代码来断言Vector已是瓶颈。
@@ -0,0 +1,35 @@
1+---
2+type: Guide
3+title: A2 FA Stage 资源与组件适用范围
4+description: 保存四段流水和轮转容量,修正平台与管线泛化。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.implementation.synchronization
9+- topic.performance.analysis
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/a2-a3-flash-attention-stage-design.md
13+ role: primary
14+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/flash-attention.md
15+ role: supporting
16+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-handcraft-pitfalls.md
17+ role: supporting
18+created_at: '2026-09-11T05:21:50Z'
19+updated_at: '2026-09-11T05:21:50Z'
20+platforms: []
21+---
22+ 
23+# A2 FA Stage 资源与组件适用范围
24+ 
25+来源A2/A3组件方案由AIC QK、AIV online softmax、AIC PV、AIV rescale组成,中间S/P/O经独立GM workspace交接。单块依赖顺序不应违反,但不同块可流水重叠;不能把“禁止改变顺序”理解为四段全局串行。生产完成后ready、消费前wait及最后消费者结束后的release共同保护slot;只在stage入口/出口握手是一种粒度选择,内部更细tile流水也可合法。
26+ 
27+preLaunch2对应3个轮转槽,元素offset候选(core*3+slot)*WORKSPACE_BLOCK_SIZE_DB,常量131072必须明确是元素还是字节。来源S/Otemp/Oupdate每核131072*12B、P每核131072*6B,对应三槽FP32或FP16;总计每核5505024B=5.25MiB。按实际row/stack、dtype、核数和阶段深度可重推,不能把该大容量公式当所有FA最小需要。
28+ 
29+源资源表UB192KiB/L1512KiB/L0A/B64KiB为所述配置候选,L0C一处写256KiB,与同批其他A2资料128KiB冲突,未核验前不能合并成硬件事实。mask[128,1024]FP16单独已256KiB,大于该UB192KiB,必须分片或其他合法设计;源码示意预算表并非可直接分配的配置。超容量可能编译/静态检查或运行时失败,不是一律编译干净。
30+ 
31+M_MTE1八个、FIX_M两个、MTE1_MTE2八个预置信用及出口wait是来源特定Block的事件生命周期,不能给任何FA kernel无条件加18对;组件自身也可能初始化。MTE1不是L0→UB、MTE2不是通用UB运算,需按真实搬运方向建依赖。硬同步基址只用于依赖该协议的launch,系统workspace函数也不能仅因FA名字统一禁用。
32+ 
33+L1/L0<128,128,128>、qS128、KV128、qN按分组推导是候选实例,静态K、动态D及尾处理须按组件定义,不能同时宣称静态K128又要求任意D≤128都与K相等。PAGED=false历史设备异常仅约束所测组件版本,恒等页表可作其连续KV适配,不证明所有A2非分页attention不可实现。
34+ 
35+FAIQK旧模板与FAQK新模板、参数6/9项及Tail特化差异需固定实际依赖版本核验;GitCode HEAD不保证长期兼容,空DispatchPolicy声明也不证明存在实现。“手搓默认/复用默认/不许读外部源码”等冲突是Skill流程,本卡不将它们作为能力限制。历史cube_util96%、1159/1166us或手写慢3–9倍不能推成必然性能。
@@ -0,0 +1,37 @@
1+---
2+type: Guide
3+title: FA 性能输入身份与覆盖分布
4+description: 保存结构张量生成、长度分桶及多batch故障探测。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.implementation.synchronization
9+- topic.performance.analysis
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-perf-gate-evaluation.md
13+ role: primary
14+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-kernel-compound-features.md
15+ role: supporting
16+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-handcraft-pitfalls.md
17+ role: supporting
18+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/flash-attention.md
19+ role: supporting
20+created_at: '2026-09-11T05:21:50Z'
21+updated_at: '2026-09-11T05:21:50Z'
22+platforms: []
23+---
24+ 
25+# FA 性能输入身份与覆盖分布
26+ 
27+性能与精度必须使用同一实际配置和合法输入域;环境旋钮、tiling、已编译实例、binary身份应能对应。source aog driver优先get_input_groups而非model.json、某版本忽略tensor value随机生成负lens,均说明声明shape/value不足以证明真实执行输入。lens/cu_seqlens/blockTable属于结构数据,随机负数或固定SEGS=1只跑block0会造成假快,必须验证有效范围和覆盖任务。
28+ 
29+B>1且各batch Q/KV长度不同可揭示前缀累计遗漏,多个core/slot/任务可揭示workspace别名;B1矩阵通过不等于这些维度覆盖。源码与实际CMake源列表、quoted include解析位置不一致也会测到另一个tiling版本。旧输入/输出与同名目录误删重建问题应通过运行身份和制品对应排除,不能只凭文件存在。
30+ 
31+来源从按时长2/4/8/6选20case改为按KV长度分桶,因为优化和测量修复会使时长桶漂移。长度是更稳定输入轴,但仍需dtype、B/H/GQA、D、mask、tail和内存边界覆盖;B*KV总token放大与单序列KV拉长是不同实验。多算子各自采样与混合总比例也不能混为一种覆盖。
32+ 
33+来源FA长度桶≤16k、16–64k、64–200k、>200k各5例,MLA≤2k、2–16k、16–100k、100–204.8k各5例,是该历史套件选择。原生成器上限16384之外运行成功不自动扩展官方支持合同,也不能把458752或204800写成全平台可支持上限。
34+ 
35+KV数据基本预算B*Sk*Hkv*(Dk+Dv)*dtypeBytes,另计Q/O、padding、量化scale和workspace;source简式kv*D*2B*2漏掉一般B/Hkv及异构D,不能通用。CPU黄金S矩阵还可能O(Sq*Sk),设备能放下KV不保证Host黄金不OOM。长度桶与老时长命中表不能同时声称是当前同一口径。
36+ 
37+来源sink H>16、H16同进程随机161001、varlen需TND等是具体CANN9.1/torch_npu2.6路径经验,回退分解基线需显式说明语义和性能对象已变。首次调用、争用和工具修复后应使用同轮数据分桶并计算ratio,不能挑不同轮的最好分母/分子。
@@ -0,0 +1,35 @@
1+---
2+type: Guide
3+title: FA 复合 Causal Varlen 与 Sink 数学
4+description: 修正局部列坐标、有限Mask与Sink分子缩放。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.implementation.synchronization
9+- topic.performance.analysis
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-kernel-compound-features.md
13+ role: primary
14+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/flash-attention.md
15+ role: supporting
16+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/a2-a3-flash-attention-stage-design.md
17+ role: supporting
18+created_at: '2026-09-11T05:21:50Z'
19+updated_at: '2026-09-11T05:21:50Z'
20+platforms: []
21+---
22+ 
23+# FA 复合 Causal Varlen 与 Sink 数学
24+ 
25+等长左上causal为c0+j≤r0+i,因此局部最后可见列lim=r0+i-c0。来源vis0=r0-c0却写lim=i-vis0,符号反了;vis0≤0的注释也与c0≤r0矛盾。整块全可见需最大key≤最小query,整块全不可见需最小key>最大query,其余都需行mask。不能仅用某个对角chunk不等式覆盖边界等号或跨多个chunk的query块。
26+ 
27+右下causal还需实际sk-sq偏移,varlen用各batch长度。因果可省约半矩阵仅对近似方形长序列且忽略块/启动开销,不能要求时间恰好减半。source NO_MASK0/MASK_SPEC1/MASK_CAUSUAL2是内部枚举,不等同外部sparse_mode2/3的全部语义。
28+ 
29+无效K/V补0仍使score0对应exp非0;必须在softmax有效集合上屏蔽。-3e38是有限sentinel,对全部无效行会相减得0并产生概率,极大合法score/负scale也可破坏假设。若mask在Muls前写入,scale0会抹掉mask,负scale会翻转;按已缩放score施加真实mask并定义全mask/非有限策略更明确。
30+ 
31+前缀varlen mask只限制KV,Q的实际长度及输出域仍需单独处理。预建每batch K长mask可减少Scalar构造,但需B*K*dtypeBytes容量及数据更新失效条件;逐batchlaunch、批界截断和连续TND是不同组织,不能无条件宣称完全等价。
32+ 
33+learnable sink若定义为只进分母的虚拟logit z,无V贡献:mF=max(m,z),alpha=exp(m-mF),lF=alpha*l+exp(z-mF),未归一化oF=alpha*o。来源向量链更新m/l却未展示o缩放,必须在RescaleO中确认仅一次对应缩放,否则z>m时结果错;也可对已归一化旧输出按等价权重校正。sink处于哪个scale单位由接口决定,不能把原始logit与已缩放logit混用。
34+ 
35+aux smax/ssum来源按每行8FP32物理槽输出,实际有效槽/复制布局需匹配消费者;Brcb+slot0掩码写法与同文Brcb不可靠断言冲突,按具体重载验证。可选null输出应由分支避免写或分配真实内部缓冲,不能把null直接传给仍写aux的Kernel。
@@ -0,0 +1,35 @@
1+---
2+type: Guide
3+title: 手写 FA 缓冲与状态故障模式
4+description: 保留多任务覆盖和事件问题,修正共享UB与最终Max公式。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.implementation.synchronization
9+- topic.performance.analysis
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-handcraft-pitfalls.md
13+ role: primary
14+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-kernel-compound-features.md
15+ role: supporting
16+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/flash-attention.md
17+ role: supporting
18+created_at: '2026-09-11T05:21:50Z'
19+updated_at: '2026-09-11T05:21:50Z'
20+platforms: []
21+---
22+ 
23+# 手写 FA 缓冲与状态故障模式
24+ 
25+FAQK两个实例若共享事件/资源管理且构造析构都操作同一信用,第二实例即使未调用也可能破坏配平;来源建议单实例分趟是该实现修复,不是所有CATLASS同类对象不可共存。每事件生命周期需平衡,不能简化为所有id每tile必恰好1Set1Wait,预置、借用、多消费者和排空会改变局部计数。
26+ 
27+mode0x2来源双AIV→AIC使用双方set和一个逻辑wait,反向单set可供双AIV;参与者不可单边return。其真实硬件/封装语义需匹配核型和版本,不能替换为950mode4表。纯Vector merge没有混合subblock语义,需按其真实block数覆盖全部行。
28+ 
29+来源“两个subblock共享物理UB”不能从覆写症状推出;物理核资源、LocalTensor地址空间、手工allocator和GM共享工作区应分别核验。需要给实际共享区按sub/core分段,不应无证据把所有AIV私有UB再强制减半。ubMax/ubSum512B对齐只在其广播/写出布局合同下成立,不是所有归约API通用起点下限。
30+ 
31+oTmp单槽在多task复用时,上一任务最后块也必须等待消费结束后释放;只保护非末块会跨task覆盖。GM双槽基址应含depth,例如(core*2+slot)*slotSize,不能core*slotSize+slot*slotSize使相邻核重叠。B1通过不能覆盖这些多任务映射,Q/O和mask前缀也需随batch递增。
32+ 
33+两launch方案第一阶段保存每块P=exp(S-m_j)、m_j和最终l,第二阶段需exp(m_j-m_final)缩放PV结果再归一化。m_final为最后有效running max或各独立分区max的最大值,来源“末块m+sum(m_j)”错误。P采用固定物理pitch时尾块必须写正确中性值;逻辑valid长度和pitch不同,不能混用。
34+ 
35+将计算和搬运分离或物理分开scratch可以消除某些复用冲突,但PIPE_ALL后仍错不证明该API无法同步,需检查跨核和真正最后消费者。设备printf、int→float、lambda均有编译模式/API范围,来源个例不能变成Ascend C全局禁令。同步标记、binary指纹和输入/输出身份比单看mtime更能排除stale观测。
@@ -0,0 +1,35 @@
1+---
2+type: Guide
3+title: FA 布局与 SeqLen 接口变体
4+description: 区分BNSD、TND、分页格式和两类长度数组。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.implementation.synchronization
9+- topic.performance.analysis
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/flash-attention.md
13+ role: primary
14+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-kernel-compound-features.md
15+ role: supporting
16+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-varlen-handcraft-recipe.md
17+ role: supporting
18+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/a2-a3-flash-attention-stage-design.md
19+ role: supporting
20+created_at: '2026-09-11T05:21:50Z'
21+updated_at: '2026-09-11T05:21:50Z'
22+platforms: []
23+---
24+ 
25+# FA 布局与 SeqLen 接口变体
26+ 
27+来源一种公开接口为连续FP16 BNSD Q[B,H,Sq,D]、K/V[B,H,Sk,D],内部Q/O为BSND[B,S,H,D]、KV为[numBlocks,128,H,D]。Q和O需相应转置,KV逐batch分块并补齐;原始K文件B*H*Sk*D*2B与设备块分配numBlocks*128*H*D*2B分开。接口选择不排除BF16、GQA、TND或设备布局转换的其他实现。
28+ 
29+BNSD可在MHA且映射一致时把B*H重解释成batch、内部head1,Q/O无转置;KV也需具有相同每head连续面及正确分页身份。GQA/MQA的KV头少于Q头,不能简单把同一BH数套给KV造成重复索引或越界。连续页表只改变寻址方式,不自动把非连续布局变连续。
30+ 
31+TND Q/O[sum Sq,H,D]可按prefixQ*H*D直接偏移;KV可选择连续TND直读,或每batch单独pad到128后转块。两条路线的Host长度合同不同:来源FAInfer适配用B项原始长度,Host/kernel自行累加;varlen s1s2 recipe用B项累加终点,b0取首项、其余差分。它们不能互换,也不同于长度B+1且首项0的常见cu_seqlens形式。
32+ 
33+实际合同需检查数组长度、单调性、非负差分、最后终点与真实token总数、Q和KV各自前缀。TND直读无需为了这个recipe自动转块;paged路线也不能只mask尾部却允许DMA读入下一batch。只有布局/长度/输入内容不变时可复用Host tiling、页表和H2D副本,指针相同不证明内容未改。
34+ 
35+BNSD固定stride的batch间距与紧密TND实际长度前缀不同,若BNSD中保留padding却按实际长度累加会错位。布局转换属于完整方案的成本,应在端到端口径中计入;不是所有转换都必须在Host,也不是只做kernel计时就能说明整体更快。
@@ -0,0 +1,32 @@
1+---
2+type: Guide
3+title: FA 优化归因与内存证据
4+description: 纠正Codegen唯一根因和大输入容量估算。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.implementation.synchronization
9+- topic.precision
10+- topic.implementation.tiling
11+status: draft
12+sources:
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-kernel-recipes.md
14+ role: primary
15+created_at: '2026-09-11T05:27:29Z'
16+updated_at: '2026-09-11T05:27:29Z'
17+platforms: []
18+---
19+ 
20+# FA 优化归因与内存证据
21+ 
22+数学算法等价且tiling字段符合某公式,不证明剩余性能差只能来自工具链。地址布局、同步、指令数量、寄存器/缓存、调度和编译选项仍可能不同;要归因codegen需要实际编译产物与控制变量证据。相邻baseline漂移6%、候选±2–7%只说明该实验收益不清楚,不能证明编译器已消除所有Barrier或源码无优化空间。
23+ 
24+TaskDuration减最长核耗时的611us对43us可以量化时间差,但92.9%差值不唯一等于“同步头开销”;可能含启动、调度、不同统计范围和测量聚合。aiv_scalar_ratio高也不独自证明Scalar计算,API发射次数是可进一步核验的候选。Tile32/64、任务数12–32及N-pack减少任务是可调参数,不是所有shape固定最优。
25+ 
26+历史11case几何1.68倍偏向H1小样本,另批大H/N得到0.31–0.52倍,体现测试分布会改变汇总,不能把前者作为全族收益。序列轴M/K和并行度BH/N应分开扫描;sumkerneldevice与Host墙钟基线不可交叉算ratio。profiler失败期旧健康窗口基线与新event数据应明示不可保证同环境。
27+ 
28+FP16 Q加O基本字节数2*B*H*Sq*D*2。B1,H16,Sq458752,D128为3.5GiB;B16,H64,Sq200000,D128为104857600000B=97.65625GiB,来源9.8GiB少10倍,尚未计KV、baseline副本和workspace,不能以64GB卡无OOM推广。固定B/H/D/Sk的Sq扫描若中途缩BH,就不再是单变量实验。
29+ 
30+短序列S1和长KV流式适用性受有效shape/特性与完整资源影响;S1所谓“长序列458752”实际是Sq长而Sk128,不能据此说一次softmax覆盖458752列。源平均0.922既非展示行的简单比值也未给全部20case,只能保留为来源汇总,不能重算成已验证统计。
31+ 
32+错误码507015、全零、profiling空目录和timeout都有多种可能,重试成功不排除初始化竞态,换case恢复也不证明原kernel无错误。可保存故障原始证据与受限实验条件,不能由一次7shape/5次压力/30case通过宣称所有FA变体生成完备。
@@ -0,0 +1,33 @@
1+---
2+type: Guide
3+title: FA 分页与量化操作数合同
4+description: 补齐物理Head偏移、尾页和反量化尺度。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.implementation.synchronization
9+- topic.performance.analysis
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-kernel-compound-features.md
13+ role: primary
14+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/flash-attention.md
15+ role: supporting
16+created_at: '2026-09-11T05:21:50Z'
17+updated_at: '2026-09-11T05:21:50Z'
18+platforms: []
19+---
20+ 
21+# FA 分页与量化操作数合同
22+ 
23+分页先按逻辑token确定table项,再由物理块号与页内位置、KV head、D得到K/V地址。来源table[B*H,maxBlocks]与物理KV[blocks,128,H,D]同时出现时,必须明确定义table是每head独立还是每batch共享;示意phys*128*H*D还缺当前head*D,除非kBlocksPlane已是相应视图。K/V必须使用同一逻辑页对应各自布局,物理页顺序不能替代causal位置。
24+ 
25+chunk拆128子块时cols//128会漏非整页尾部,应ceil并传实际有效列。每块S写入packed chunk需要正确行距/列偏移;如果softmax仍按完整chunk归一化,不能同时声称已经每128块更新状态而不修改PV缩放。改变粒度要保持P对应的max、sum和oacc缩放,不能只改bmm1指针。
26+ 
27+blockTable身份、范围、尾页valid和共享页规则按实际接口检查。恒等页表可与dense做数学对照,但不同分块/浮点次序不保证位级相同。重复物理页是否是合法逻辑重复输入由合同决定,不宜一概去重。
28+ 
29+q8来源量化定义q=round(x*s),反量化为q/s;若输入scale本来表示dequant乘系数,则应q*scale。必须区分两种命名,s0、溢出饱和、舍入和zero-point也需规定。Golden需依据实际量化整数重建,不能只把原始x除s。
30+ 
31+dav_c220来源采用INT8→half再Muls,缺INT8→float某直转重载不意味着所有BF16/FP32输出方案都不可能。MTE2加载后Cast需真实依赖,PIPE_ALL只是较粗候选;AIV反量化发布后AIC才能读取。逐tile生产却只在整batch一次ready,必须保证全部消费数据已写完,否则通知粒度不匹配。
32+ 
33+完整GM staging增加INT8读、half写、后续half读及计算,流量按具体缓存/复用计算,不能一律“每元素6B纯带宽”。on-load转换可减少往返但需组件支持和数值合同。q8仅FP16、varlen/sink几乎免费及0.3%/4.3%收益属于来源实现和历史样本,不是功能叠加的性能保证。
@@ -0,0 +1,39 @@
1+---
2+type: Guide
3+title: FA 性能采集范围与汇总口径
4+description: 保存跨框架计时陷阱,修正公平下界与唯一测量方法的泛化。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.implementation.synchronization
9+- topic.performance.analysis
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-perf-gate-evaluation.md
13+ role: primary
14+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-kernel-compound-features.md
15+ role: supporting
16+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/flash-attention.md
17+ role: supporting
18+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-handcraft-pitfalls.md
19+ role: supporting
20+created_at: '2026-09-11T05:21:50Z'
21+updated_at: '2026-09-11T05:21:50Z'
22+platforms: []
23+---
24+ 
25+# FA 性能采集范围与汇总口径
26+ 
27+性能比较需同公式、输入、dtype、layout、设备/争用窗口和计时范围。device kernel执行、同stream事件区间及包含Host调度/布局转换的端到端时长是不同指标。Host初始化销毁和executor重建可显著扩大后两者,不能据纯device结果解释为Host框架开销带来的加速。
28+ 
29+来源torch进程ctypes launch逐次event取min出现约1.6–1.7us恒值和391倍假加速,是事件未覆盖实际外来stream/操作等问题的线索。可用实际launch stream上的ACL事件组回放并与profiler核对;不能推广为全部torch event永远无法量外来kernel。分组3*10取min/均值/median都要标明聚合顺序,某一最小值尤其易放大漏测。
30+ 
31+ATB executor重建例event4061us/device76us说明区间可能含Host造成的GPU/NPU空闲,选择profiler纯device可隔离该部分;端到端吞吐仍可合法用同步包围的多次回放,不能称墙钟一律无效。NNAL ATB环境脚本和torch_npu导出解析触发属于来源安装/版本要求,不是所有安装路径固定或仅export_chrome_trace一种导出。
32+ 
33+profiler过滤只留AI_VECTOR_CORE/AI_CORE会漏MIX_AIC融合kernel;类型字段及名称按导出版本核验。磁盘满、trace截断、进程OOM和解析未完成都可造成异常短时长,不能只凭15us对72–78us差异认定工具盲区。独立进程逐case可降低累积内存影响,也会改变缓存/初始化条件。
34+ 
35+多kernel链逐项device duration之和忽略Host间隙,但若核间/stream间重叠,它也不是必然的端到端公平下界;需区分工作量总和与时间线span。分解链与融合基线应显式命名,source33.5倍墙钟对1.16倍device不能当同一加速指标。
36+ 
37+gate可用exp(mean(log(tRef/tCustom))),要求各时长正且有限、逐case配对及权重明示;失败/缺测/0/0不能算通过或静默排除。算术均值对大ratio敏感,几何均值也不替代逐case回退/尾部风险。来源1.025、3.047、0.107–0.364等不同工程历史值不能作新实现门槛。
38+ 
39+争用时将memset和kernel分事件区间可定位排队成本,但队列等待不保证只发生在首操作;若算法每launch确实需要清workspace,端到端仍应计入。kfc只清64KiB取决于实际系统区合同。两侧同窗口交替、多轮、同配置可降低环境偏差,不能保证完全相同争用。torch.cuda.Event并非NPU事件替代;同步墙钟均摊仍含Host部分,应与device口径分列。
@@ -0,0 +1,36 @@
1+---
2+type: Guide
3+title: FA Q Head 打包与 Block 生命周期
4+description: 保存GQA行映射和组件内部槽位,限制硬事件照抄。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.implementation.synchronization
9+- topic.precision
10+- topic.implementation.tiling
11+status: draft
12+sources:
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-kernel-handcraft.md
14+ role: primary
15+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-kernel-recipes.md
16+ role: supporting
17+created_at: '2026-09-11T05:27:29Z'
18+updated_at: '2026-09-11T05:27:29Z'
19+platforms: []
20+---
21+ 
22+# FA Q Head 打包与 Block 生命周期
23+ 
24+FAInfer候选每task为(batch,qSBlock,qNBlock),qS128,qN=min(max(floor(128/qLen)//2*2,1),groupSize),groupSize=Hq/Hkv须为正整数。每组qN块数ceil(group/qN),task数逐batch为qSBlocks*qNBlocksPerGroup*Hkv。该启发式在qLen0除零,且取min到奇数group或尾head后并不保证偶数;两AIV分头必须处理余头,不能都只取floor(qN/2)。
25+ 
26+rowNum=validQRows*validQHeads,把BSND的[S,head,D]片段按真实stride载入为二维Q。kvHead=qNBlock/qNBlocksPerGroup,qHead=kvHead*group+(qNBlock%qNBlocksPerGroup)*qNTile;Q/O和K/V使用各自平面基址。qN1可按query行分两sub,qN>1可按head分,写回必须逆映射且覆盖奇数/尾行。
27+ 
28+AIC/AIV对同一task使用一致的逻辑core编号、步长和batch累计。源伪码curTotalTaskNum=0后while立即++curBatch会跳过batch0,只有另有正确首batch初始化才可用;空batch连续跳过还需boIdx边界。AIV GetBlockIdx除subBlockNum与GetBlockNum语义必须匹配launch,不能两者都机械除2。
29+ 
30+选定FAIQK的Q经GM→L1一次装载并在任务内复用,K/V用L1双缓冲再送L0,L0A/B与L0C生命周期分开。四个128KV块叠512列,L0C可能两块成对累加/写出;gCOffset=(stackIdx//2*2)*128需与输出stride和有效列一致。尾K、最后奇数KV块必须执行正确initMmad/endItr,不能只靠整4块流程。
31+ 
32+computeQK伪码列GM→L1、L1→L0、Mmad、Fixpipe事件,但省略部分反向复用和consumer释放,不能当完整可执行协议。M_MTE1保护Mmad读完后MTE1覆盖L0,MTE1_MTE2保护MTE1读完后MTE2覆盖L1,M_FIX/FIX_M交接L0C计算与写出;源把MTE1说成L0→UB是错误解释。
33+ 
34+来源AIV预置表2+4+4+1实际11项,自查却写10;AIC18与AIV11也只适用于那些成员组件和event池。换FAQK若构造已管理6类/槽信用,重复手动预置会破坏协议;析构行为也需核对。“多余Set导致缺token”的描述不构成完整因果,必须推演真实计数。
35+ 
36+OnlineSoftmax把S/P、local/global max/sum、diff和scratch按向量/行广播布局分配,8192元素S槽和固定字节偏移只是特化。preLaunch延迟的RescaleO要读对应块的alpha和最终l,不能拿当前最新版覆盖旧消费者。常量表、14参数入口和模板特化语法仍需与真实类声明/Host ABI核对。
@@ -0,0 +1,32 @@
1+---
2+type: Guide
3+title: FAQK 交接与 Vector Count 约束
4+description: 保存Token复用和批量布局知识,限制器件故障泛化。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.implementation.synchronization
9+- topic.precision
10+- topic.implementation.tiling
11+status: draft
12+sources:
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-kernel-recipes.md
14+ role: primary
15+created_at: '2026-09-11T05:27:29Z'
16+updated_at: '2026-09-11T05:27:29Z'
17+platforms: []
18+---
19+ 
20+# FAQK 交接与 Vector Count 约束
21+ 
22+来源FAQK可配置紧密LayoutC,而FAIQK选定特化带512stride与KV堆叠;换组件需同时更新调用签名、Q装载、page能力、资源分配和构造析构事件。FAQK示意接受可变pingpongFlag和isFirst,后者表示该调用真实A重新装载,不是任何调用永远true。
23+ 
24+连续调用复用L1A时,下一GM→L1写必须等待前次MTE1读完。来源Wait(MTE1_MTE2,prev)再Set恢复token可在特定组件内部后续仍需消费该credit时成立;需核验实际事件所有者及调用步数,不能机械插入或用“短于32用全Barrier、长128用dance”当通用阈值。跨阶段、0/1调用和出口残留均需账目一致。
25+ 
26+count重载Add/Muls/Exp/Cast/Adds表达元素数,不是repeat次数;repeat参数重载则另有mask/stride/次数语义。ReduceMax/Sum“每64只出一个/128静默错”可能来自具体重载、block/whole归约或workspace误用,不是所有高阶Reduce固定只支持64。分段64再合并是可选保守实现,末尾不足64要有效mask,scratch256float也需实际API而非固定经验。
27+ 
28+批量16行可减少发射与搬运,前提是S0等确实紧密;若UB行pitch=KSEG,bulk load也必须按同一物理pitch,不能一边紧排一边以KSEG逐行计算。source330次*100ns约33us,与其约25us估计不同,计数/单次成本是粗模型不能当精确因果。
29+ 
30+来源2D store仅首块、局部小stride load混行、Duplicate后SetValue全零等观测不足以禁止所有2D搬运或Scalar/Vector混合。应检查stride单位、blockLen、地址、元素位宽和S↔V/MTE依赖;逐行1D及物理独立scratch是候选修复,PIPE_ALL通过也不证明唯一根因。
31+ 
32+UB偏移93188与93184相差4B说明源实例存在32B对齐条件,但全零正常返回不是该条件唯一症状。“器件文档不可信”“dump慢就不是race”等绝对结论不入库;模板注释、目标头文件、最小探针和真实运行证据应一起核验。
@@ -0,0 +1,40 @@
1+---
2+type: Guide
3+title: Attention 特性的可见集合与量化交互
4+description: 抽取 causal/window/sink、稀疏索引、RoPE 和后量化的组合边界。
5+tags:
6+- task.optimization
7+- topic.implementation.tiling
8+- topic.performance.analysis
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-tiling-design/references/flashattention/foundation/composition.md
12+ role: primary
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-tiling-design/references/flashattention/foundation/extension_points.md
14+ role: supporting
15+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-tiling-design/references/flashattention/specialization/feature_flags.md
16+ role: supporting
17+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-tiling-design/references/flashattention/specialization/sparse_design.md
18+ role: supporting
19+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-tiling-design/references/flashattention/subfamilies/quantization_design.md
20+ role: supporting
21+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-tiling-design/references/flashattention/patterns.md
22+ role: supporting
23+created_at: '2026-09-11T04:27:49Z'
24+updated_at: '2026-09-11T04:27:49Z'
25+platforms: []
26+---
27+ 
28+# Attention 特性的可见集合与量化交互
29+ 
30+causal与window可以组合,合法KV集合通常按各自规则取交集;sink扩展的集合还需按接口处理因果边界,不能简单说二者互斥或前N项永远无mask。sink可能是额外KV,也可能是分母logit等其他算子语义,不能从同名字段推断。prefix需绑定物理KV布局、分页/长度和全局位置;chunk query偏移按全序列位置,不能把前一query块的输出当当前query的累积状态。
31+ 
32+来源将PSE扩写为Partial Sum Extension并解释为部分输出累加,没有算子合同证明。应按真实pseType、realShift/score偏置及scale应用次序解释该参数,不能由缩写发明新的阶段。RoPE对Q/K的位置、配对维和cos/sin取值有明确数学作用;与量化/反量化的顺序会改变数值,也不能仅声称可在L1→L0间插入Vector计算而不证明执行位置支持。
33+ 
34+稀疏设计来源明确尚未验证,但可保留可见集合约束:块索引需合法、处理重复/空列表并保证没有遗漏有效token;动态Host索引生成需要数据确实可在Host取得。部分mask块的实际Cube发射FLOP可能大于有效FLOP。gather/pack有效token后仍可能用Cube,不能断言稀疏无法Cube计算。
35+ 
36+在线rescale由最大值变化决定,不由已处理块数决定。跳过块应保持已有状态,所有块都跳过时仍要完成空输出/事件;最后有效块不一定是循环末尾。量化数据与scale需按同一绝对KV坐标取值,不能跳过数据却独立推进scale计数。稀疏选择可能打散原32元素scale分组,需要保留原分组映射或重新量化。
37+ 
38+后量化应对最终归一化后的O应用合同规定的scale、舍入和饱和,不要求输入也使用MX。amax=max(abs(O)),不能用max(O)替代;全负输出即为反例。amax=0、NaN/Inf及overflow需明确定义,E8M0表达正的幂次scale时不能直接把任意amax/qmax当可精确表示的编码。RINT等转换对异常值的行为受真实API约束,不作统一猜测。
39+ 
40+量化并不取消mask/padding和统计FP32预算;1-bit mask只在实现真正采用该编码时才节省存储。特性表中的勾选与“无相互影响”不是已测试的组合能力,本次保存可组合的数学条件及来源未验证限制,不把路线图提升为支持承诺。
@@ -0,0 +1,36 @@
1+---
2+type: Guide
3+title: FFA 正交 Batch 映射与流式状态
4+description: 保存双KV对角模型的四路GEMM和完整平面偏移。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.implementation.synchronization
9+- topic.precision
10+- topic.implementation.tiling
11+status: draft
12+sources:
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-kernel-recipes.md
14+ role: primary
15+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-kernel-handcraft.md
16+ role: supporting
17+created_at: '2026-09-11T05:27:29Z'
18+updated_at: '2026-09-11T05:27:29Z'
19+platforms: []
20+---
21+ 
22+# FFA 正交 Batch 映射与流式状态
23+ 
24+来源双KV对角FFA使用Q[N,M,D]、K1/V1按N分组、K2/V2按M分组;同一(n,m)输出的score按k为scale*(dot(Q[n,m],K1[n,k])+dot(Q[n,m],K2[m,k])),输出按该同一概率加权V1[n,k]+V2[m,k]。它不同于将两套KV沿k拼接并扩大softmax域。
25+ 
26+可用Q0[N,M,D]及Q1[M,N,D]两个视图/转换,QK1任务(n,mTile)、QK2任务(m,nTile),PV同理。S0[N,M,K]、S1[M,N,K]各自连续k,Vector用各自地址相加;P也需两种布局供两路PV。正交分组令共同连续排列困难,但并非数学证明所有其他布局/融合均不高效。
27+ 
28+PV若复用ColumnMajor B组件,物理V^T[D,K]的leading dimension为K,应按对应layoutVB(K,D)表达,块j起点+j*KB而非+j*KB*D。运行问题K轴长度KB与原始整平面strideK分开,D==K或K==KB会掩盖错误。新增BH后V1平面大小N*D*K、V2为M*D*K,S/P/O与mask各有独立广播/平面规则。
29+ 
30+终态候选tile32、KB128:A0/A1分别形成S0[i,r,k]/S1[r,i,k],Vector16行批量加scale/mask并online更新,C1/C2生成O1/O2,Oacc=alpha*old+O1+O2,末块除l。首步直接初始化输出可避免未初始化NaN乘0;nK需ceil且处理真实尾,来源nK=K/128只支持其整除合同。
31+ 
32+per-core S0/S1/P0/P1均按tile²*KB,O1/O2及Oacc按tile²*D并计各自双缓冲,工作区随核数而非全K线性增长。O2[r,i,D]不能对i/r误作连续块;S1/P1的转置访问同理。全K物化可以减少逐块握手但增加空间,两路都可能合法,“大K只有流式可行”取决于可用内存和实际shape。
33+ 
34+来源aux[B,H,N,M,8]仅槽0有效且其余0是特定输出合同,Host清零和Kernel只写槽0须配套;不同API可能复制8槽。mask[B,1,N,1,K]可在同n多m行共享,不能跨n误用;强制k0不mask只覆盖非空行,不能代替合法全mask测试。scale默认1.0与1/sqrt(D)不同,必须保留接口默认。
35+ 
36+四路GEMM可以用5条分别ready或先合并两QK/两PV后3条ready,但只在所有相关生产者均完成后通知;不是四个GEMM天然必须5个flag,也不是原模板所有事件/slot无需改动。闲核是否可退出取决于pair协议还是全核barrier,source各段1:1、2Set:1Wait及“全核参与”不能混用。
@@ -0,0 +1,33 @@
1+---
2+type: Guide
3+title: FFA TensorList 前缀与组件合同
4+description: 区分拼接KV和双KV对角模型,保存FD LSE与量化入口语义。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.implementation.synchronization
9+- topic.performance.analysis
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/ffa-handcraft-recipe.md
13+ role: primary
14+created_at: '2026-09-11T05:21:50Z'
15+updated_at: '2026-09-11T05:21:50Z'
16+platforms: []
17+---
18+ 
19+# FFA TensorList 前缀与组件合同
20+ 
21+FIA来源接口族包含query、key/value TensorList、共享key/valuePrefix与实际prefix长度、query/keyRope及rope antiquant scale、learnableSink、deq/quant scales与offset、key/value antiquant scales/offset、Q/KV长度和padding、blocktable、attentionOut与可选softmaxLse。称37参不替代真实签名,参数可选性、dtype/shape、pse/mask与scale顺序需逐项对齐。
22+ 
23+sharedPrefix与每batch local KV沿序列轴拼接时,逻辑长度是sharedLen+localLen_b,两段进入同一softmax归一化。TensorList各成员可能不同存储池,必须有能表达pool/base与页号的寻址,不能只令blocktable物理号连续就自动访问多个不连续tensor。无需Host全量拼接是可实现方案,但不是任意已有kernel都天然支持。
24+ 
25+来源另称“双KV对角FFA”的S=Q*K1^T加特定对角Q*K2贡献,与prefix/local的序列拼接是不同算子合同,不能用同名FFA互相替代。sink-in-epilogue也需区分虚拟logit与额外KV;rope旁路维度与scale顺序同样不能只套一个加号。
26+ 
27+组件路线包含FAIQK/FAIPV及Tail、resetBlockStart(kvStart,pagedBlockSize)、ND或K nZ/V zN布局、MAX_KV_STACK_LEN约512、PRE_LAUNCH延迟PV、CrossCore mode0x2和GM轮转。nZ/zN不是更改layout名字即可转换物理数据;INT4 view-only转换也需原始packed字节、stride与元素数量已经符合。
28+ 
29+IS_FD让每核处理指定s2分段,仍需正确的partial max/sum/o归并;空KV任务由epilogueInitOut定义输出,LSE按[totalQTokens,qHeads]真实布局。mask压缩视图、fullMask、sink参数须传到所有相关分支,不是只有主循环一个调用点。
30+ 
31+registry tilingKey可编码dtype、LSE、TND、paged、KVNZ、mask、split-fuse等组合,direct入口可以用其他经验证dispatch,不必照搬key。循环从kvStart非0开始时,延迟PV条件应基于已生产步数,不能只写绝对kvSIdx>=PRE_LAUNCH就读取未生产槽。末块按当前分段而非全局kvLoop决定,slot保留到消费者结束。
32+ 
33+smax/ssum与基线位级一致不是所有实现必需,应按明确数值合同验证;source几何2.19倍只对应其历史双KV对角实现,不能转借给sharedPrefix组件路径。
@@ -0,0 +1,43 @@
1+---
2+type: Guide
3+title: G 轴分工与变长 KV 线段切分
4+description: 补齐核号映射、零维、贪心超额分核及独立输出和 KV 归约的区别。
5+tags:
6+- task.design
7+- topic.attention
8+- topic.tiling
9+- topic.parallelism
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/cards/P57_flashdecode_g_partition.md
13+ role: primary
14+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/cards/P58_tnd_load_balance_partition.md
15+ role: supporting
16+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/playbooks/P57_flashdecode_g_partition.md
17+ role: supporting
18+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/playbooks/P58_tnd_load_balance_partition.md
19+ role: supporting
20+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/preconditions/P57.yaml
21+ role: supporting
22+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/preconditions/P58.yaml
23+ role: supporting
24+created_at: '2026-09-11T01:56:34Z'
25+updated_at: '2026-09-11T01:56:34Z'
26+platforms: []
27+---
28+ 
29+# G 轴分工与变长 KV 线段切分
30+ 
31+当B*Hkv的独立任务不足时,可把每组G个query head/输出行进一步分给多个执行者。G在具体算子中的含义必须先确认:将独立query输出分开,不需要把这些不同G输出相加;如果分的是同一query的KV贡献,才需要正确的FlashDecode partial归并。不能由策略名“G分核归约”推断所有分片都要sum。
32+ 
33+来源先令P=min(floor(C/(B*Hkv)),ceil(G/16)),再取S=ceil(G/P),修正P为min(P,ceil(G/S)),尾长为G-(P-1)*S。必须先处理B*Hkv=0、G=0、C=0及B*Hkv>C导致P=0的情况,且乘法/向上取整避免溢出。来源固定48核、B*Hkv<=24、最小粒度16是案例策略;“用ceil(G/16)限制P”并不保证每片>=16:G=17、P=2时得到9与8。
34+ 
35+每核还需从核号得到所属(b,hkv)任务与局部分片p,例如在连续映射下 group_id=core/P、p=core%P,然后start=p*S,valid=min(S,G-start)。来源所有核gSizeStart=0只是未展开helper的占位,若直接使用会重复工作。超过有效任务的核可以跳过计算,但若仍是屏障参与者则不能一律return。
36+ 
37+变长TND使用真实累计边界,T一般为各batch有效token数之和,不是固定B*S。来源按U=512把每个KV长度变为w_b=ceil(KV_b/U),以ceil(sum(w)/C)作目标;这只近似一个固定query/head范围的代价。多个query tile、G或head、causal/mask与尾块成本不同,还需相应计入,QS*G<=128也是案例tile约束。
38+ 
39+源贪心方案“放不下当前整段就换核”不保证核数上界:段长[2,2,2]、C=2、平均上取整3时,会依次分给核0、1、2。不可拆分的大段也可能超过平均。应选择有明确C上界的完整任务分配,允许不均衡、重新平衡,或沿KV拆出可合并partial;不能只丢弃超额核任务。全零长度、空batch和单段过大均需定义行为,定长输入也不禁止使用可证明合法的相同分配框架。
40+ 
41+跨核切分同一query的KV线段时,metadata要保存batch/query/head、KV区间、分片数量与workspace位置;状态至少包含实际需要的max、sum及未归一化输出。workspace容量不能仅写batchSize*sizeof(splitInfo),还要乘真实分片/行数、输出D及dtype。归并按max重缩放,不能直接相加已归一化输出,见[online归并](online_softmax_state_and_partition_merge.md)。
42+ 
43+两种分工组合后,要重新证明总核数、任务覆盖与输出所有权,而非机械规定先P58再P57就正确。验证0维、G=17、非整分片、变长[2,2,2]单位段、跨多个核的同query、超额空核和tail,测每核与整体时间并计Host分析和归并成本。上述两个反例仅由源整数公式复现。
@@ -0,0 +1,36 @@
1+---
2+type: Guide
3+title: GQA 任务几何与 Split-KV 并行条件
4+description: 区分 head 分组、查询行、物理 stride 和可用核数。
5+tags:
6+- task.optimization
7+- topic.implementation.tiling
8+- topic.performance.analysis
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-tiling-design/references/flashattention/design/shape.md
12+ role: primary
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-tiling-design/references/flashattention/foundation/composition.md
14+ role: supporting
15+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-tiling-design/references/flashattention/patterns.md
16+ role: supporting
17+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-tiling-design/references/flashattention/subfamilies/base_design.md
18+ role: supporting
19+created_at: '2026-09-11T04:27:49Z'
20+updated_at: '2026-09-11T04:27:49Z'
21+platforms: []
22+---
23+ 
24+# GQA 任务几何与 Split-KV 并行条件
25+ 
26+GQA中G=Hq/Hkv要求Hkv>0且整除;实际计算行数mEff=curBq*curG,容量规划通常使用Bq*gMax,不必一次包含全部G。分组融合可改善K/V复用,但不是数学正确性的必需条件;curG=1或分组切分仍可计算相同输出,只是资源/性能不同。
27+ 
28+一个明确任务网格为B*Hkv*ceil(Sq/Bq)*ceil(G/gMax),每任务保留query/head起点和curBq/curG尾数。来源r=g*Bq+i是head外层、query内层映射,不能因文字称Bq-major就改变公式。行间使用名义Bq还是curBq作stride,必须在Q加载、mask、P/V与输出解码中一致。BSND输出元素偏移为(((b*Sq+q)*Hq+h)*Dv+d),query行步长Hq*Dv;gOff须明确按head、行还是元素计量。
29+ 
30+DataCopyPad参数长度若接口单位为字节,D需乘dtype字节;示例{1,D}不能跨重载复制。Dqk决定QK归约,Dv决定PV输出列,两者无需相同。MHA是G=1的数学特例,不要求源码与GQA完全同一分支。
31+ 
32+Split-KV按K/V区间生成局部最大值、指数和与加权V,再用全局最大值重标度合并;局部区间内部仍可使用Cube和在线softmax,不是Vector专属。空/全mask区间需中性状态,不能计算-inf-(-inf)后把NaN传播给合并。
33+ 
34+是否缺少并行应先比较原始任务数与物理可用AIC数。若usedCoreNum=min(totalTasks,aicNum),条件totalTasks<usedCoreNum永远不成立,不能作为启用Split-KV的判据。Sq=1或Sk>2048是场景启发,已有足够任务时也可能因其他瓶颈受益,仍需计额外workspace、归约、尾片和launch开销。
35+ 
36+Vector FA快/慢10–100倍、固定__mix__(1,2)以及统一增加AIV均无通用证据。增加AIV可能争用UB/带宽或改变同步参与者;更多AIV对应AIC:AIV数值比降低,不能把比例名称与方向混淆。
@@ -0,0 +1,32 @@
1+---
2+type: Guide
3+title: 共享 KV head 的 Decode GEMV 合并
4+description: 以同组 query heads 为矩阵 M 维,保留分组、数值稳定性、任务并行和分块累加。
5+tags:
6+- task.optimization
7+- topic.compute.composite
8+- topic.parallelism
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/discovered/disc_X5.md
12+ role: primary
13+- resource: https://gitcode.com/qq_41667743/cannbot-knowledge/blob/a70025d581ae45da8a93b4938b5ba19da6fa6f92/knowledge/ops/ascendc/operators/aclnn/transformer/attention/incre_flash_attention.md
14+ role: constraint
15+created_at: '2026-09-11T02:29:01Z'
16+updated_at: '2026-09-11T02:29:01Z'
17+platforms: []
18+---
19+ 
20+# 共享 KV head 的 Decode GEMV 合并
21+ 
22+Decode/MTP即使query序列S只有1或2,也可将共享同一KV head的多个query head排为矩阵行。对Nkv=1,Q为[Nq,Dqk],QKᵀ为[Nq,Skv],P·V为[Nq,Dv],将原来Nq个GEMV合成GEMM。Nkv>1且Nq能被整除时,可逐KV组使用M=G=Nq/Nkv;原资料笼统排除所有Nkv>1过严,但不能把读取不同K/V的所有组直接合成共享同一右矩阵的GEMM。MHA的G=1没有相同的组内行扩展收益。
23+ 
24+每个(b,s,kvGroup)任务必须取对应Q行、K/V缓存和输出head位置。Nkv=1时只有B*S个任务,B*S很小仍可能闲核;可进一步按合法维度分块并处理合并成本,不能声称这种网格天然占满设备。不同head可有不同mask/PSE,只有确实相同的部分才共享预处理。S=2的causal位置也不能沿用S=1的可见列。
25+ 
26+AIV可负责布局、mask和softmax,AIC负责QKᵀ及PV;交接的GM/片上地址、dtype、发布/消费/回收必须闭合。源码MARSHAL/CS/VP一次性flag是特定串行阶段协议,每任务需重置/消耗正确,不能直接假定事件0–7无冲突;1:2时两AIV行域、到达顺序及空任务必须一致。串行阶段会付同步成本,不自动实现CV重叠。
27+ 
28+分数“通常O(1)”不足以跳过max-shift。需按输入域证明或保留稳定softmax与所有KV tile的在线归一化。若PV按Dv分块,每个N tile都可能重新读取P;QK的NC分块也可能重载Q。统计真实重复流量,不能同时宣称Q/K/P/V都只读一次。Fixpipe直接写最终y只在该输出的全K累计和softmax缩放完成后成立,不能让后续tile覆盖部分和。
29+ 
30+示例NC=256、KC=256、NV=128,FP32 L0C保存[Nq,128]时Nq=64占32KiB、Nq=128占64KiB;还需计算A/B装载、双缓冲、L1和UB。源L0C128KiB不是所有平台规格。实际NZ布局、尾部、Mmad初始化/续累加和Fixpipe行距按接口合同建立。
31+ 
32+已有[IncreFlashAttention](../../operators/aclnn/transformer/attention/incre_flash_attention.md)也采用M=G的矩阵视图,但其公开V4的S1=1、dtype、page及route guards不等于本来源S=2专用实现的已验证支持。分别验证Nkv1/多组、不同heads、Skv/Dv尾块、多tile状态与全masked行。来源decode11–33倍、geomean0.5745→3.527未在本批设备复现。
@@ -0,0 +1,40 @@
1+---
2+type: Guide
3+title: 线性 Attention 的分组 head 缓存与梯度归属
4+description: 按真实数据依赖限定 HK 缓存,保持跨 stage 的 workspace 映射,并区分共享输入与需累加输出。
5+tags:
6+- task.design
7+- topic.memory
8+- topic.parallelism
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/a2-a3-linear-attention-stage-design.md
12+ role: primary
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/linear-attention.md
14+ role: supporting
15+created_at: '2026-09-10T14:50:50Z'
16+updated_at: '2026-09-10T14:50:50Z'
17+platforms: []
18+---
19+ 
20+# 线性 Attention 的分组 head 缓存与梯度归属
21+ 
22+在多个 value head 共享一个 key head 的分组合同下,若 `HV % HK == 0` 且各组连续均分,映射为 `hk = floor(hv / (HV/HK))`。这描述的是一种明确的分组约定;不把它推广到任意 Attention 接口或非均匀映射。
23+ 
24+## 缓存键包含全部数学依赖
25+ 
26+只依赖 `(batch, chunk, hk)` 的 K-side 中间量,可以在对应 value head 组内共享。若中间量还包含按 hv 变化的 gate、beta、scale 或状态,就不能只按 hk 缓存。缓存复用范围应由计算依赖决定,而不是由输入张量名字决定。
27+ 
28+来源的 Catlass 多 stage 设计使用稳定 workspace slot 连接生产者与消费者。同一逻辑数据在所有 stage 中必须使用一致的 slot 公式;GM 跨 stage slot 与 UB ping/pong 是不同的生命周期。窗口轮转前需要确认旧消费者结束,新数据发布也需要满足生产完成依赖。固定四槽公式或两个 head 的调度只是特定实现选择,不足以自行证明复用安全。
29+ 
30+## 共享读取与梯度累加不同
31+ 
32+多个 hv 读取同一份 K-side 缓存不需要对读取做归约;若多个 hv 产生同一 hk 输出的梯度贡献,则需要累加归属。把整组贡献交给唯一执行者,可以在组内按确定顺序初始化、累加和写回;若把同组拆到多个核,必须另行设计原子或分阶段归约,不能继续使用无保护的读改写。
33+ 
34+按 batch/chunk 分核并在核内遍历 head 组,可以保持归属,但 batch/chunk 任务过少时会降低并行利用率。增加 head 维并行度时,需同时修改输出归属及 workspace 协议。
35+ 
36+## 分块边界
37+ 
38+物理 chunk/tile 尺寸与实际 validRows 分开管理,varlen 任务从真实序列和 chunk 索引得到基址。不能仅依据总 T 推断每个序列尾块。L1 resident、scratch 和跨 stage workspace 分别核算存活范围;大 V 不天然等价于 GEMM 的 K 分割,是否累加以及何时最终写回必须按每个矩阵乘的数学轴判断。
39+ 
40+验证应覆盖 HK=HV、HV>HK、尾 chunk、不同 gate/状态,以及同组贡献的多核分配。来源提到的固定 tile、分钟级长跑与性能收益缺少本批原始执行证据,不在这里发布为平台规律。
@@ -0,0 +1,78 @@
1+# attention
2+ 
3+- [分块 Attention 的 online softmax 状态与归并](online_softmax_state_and_partition_merge.md):保持 max、指数和与未归一化加权和一致,正确处理跨 KV 分块缩放、Split-KV 和分块 softmax 输出。
4+ 
5+- [Attention 的逻辑 token、head 与物理缓存地址](logical_tokens_heads_and_cache_addresses.md):分开处理 TND batch 边界、GQA head 映射、稀疏 token 选择与 page table,防止布局组合后混用坐标。
6+- [Attention 打分维度与 Value 输出维度](score_dimensions_and_value_dimensions.md):在 MLA 等分段输入路径中分别管理 QK 打分维度、Value 宽度及 latent 数据消费范围。
7+- [Attention 的 score 可见性与额外 KV 来源](score_visibility_and_additional_kv_sources.md):区分 causal/显式 mask 与额外 sink KV,保证有效列集合和 online softmax 的统一归一化。
8+ 
9+- [线性 Attention 的分组 head 缓存与梯度归属](grouped_head_cache_and_gradient_ownership.md):按真实数据依赖限定 HK 缓存,保持跨 stage 的 workspace 映射,并区分共享输入与需累加输出。
10+ 
11+- [CV 交错流水的任务槽与数据槽](cv_pipeline_task_and_data_slots.md):明确 MM1、Vec1、MM2、Vec2 的三期排程、索引范围和跨核生产消费及回收关系。
12+ 
13+- [Online Softmax 状态槽的接续与最终选择](online_state_slots_and_final_selection.md):明确 max、sum、缩放因子和输出累加器的任务身份,避免双槽递推丢失或重复计数。
14+- [Workspace 槽地址与延迟梯度 Scatter](workspace_slots_and_deferred_gradient_scatter.md):区分阶段槽、梯度聚合域和最终刷新时机,保留完整 dK/dV 贡献。
15+ 
16+- [G 轴分工与变长 KV 线段切分](g_axis_partition_and_variable_length_segments.md):补齐核号映射、零维、贪心超额分核及独立输出和 KV 归约的区别。
17+ 
18+- [双 AIV 行范围与 Workspace 份数](dual_aiv_row_ranges_and_workspace_multiplicity.md):保存对齐分工公式、空消费者与物理槽容量,避免混淆核比例和核总数。
19+- [稀疏与稠密分发需要有效集合等价](sparse_dense_dispatch_requires_set_equivalence.md):区分块数量预算和实际 token 覆盖,保留分发后的掩码与资源合同。
20+ 
21+- [反向 Cube 融合需保留全部梯度路径](backward_cube_fusion_preserves_all_gradients.md):区分 dS 与概率矩阵,补齐累加器、转置输入和输出写回的寿命。
22+- [Sink 任务与跳块后的流水状态](sink_tasks_and_skipped_block_state.md):补齐 sink 与 KV 的集合关系、任务字段及 online 状态跨任务归并。
23+- [Softmax 反向行点积的预计算](softmax_backward_row_dot_precomputation.md):保存梯度公式、每 query 行结果容量及 dY/O 版本一致性。
24+ 
25+- [Decode 拼接搬运与 online tile 归约](decode_concat_dma_and_online_tiles.md):融合 nope/rope 的物理暂存,并保存两级点积、有效域与在线 softmax 状态。
26+- [共享 KV head 的 Decode GEMV 合并](grouped_decode_gemv_to_gemm.md):以同组 query heads 为矩阵 M 维,保留分组、数值稳定性、任务并行和分块累加。
27+ 
28+- [因果窗口、GQA 与整块跳过的证明](causal_windows_gqa_and_block_skip_proofs.md):补全闭区间边界、query 坐标和空行语义,限定 Attention 性能案例。
29+ 
30+- [Attention 特性的可见集合与量化交互](features_sparse_visibility_and_quantization_interactions.md):抽取 causal/window/sink、稀疏索引、RoPE 和后量化的组合边界。
31+- [GQA 任务几何与 Split-KV 并行条件](gqa_task_geometry_and_split_kv_capacity.md):区分 head 分组、查询行、物理 stride 和可用核数。
32+- [MLA 扩展假设与 Attention 特化范围](mla_extension_hypotheses_and_specialization_scope.md):保留扩展点而不把未验证路线写成平台能力。
33+- [Attention 在线 Softmax 的空行与量化阶段合同](softmax_empty_rows_and_quantized_stage_contracts.md):修正有限 mask、初始 rescale 及 QK/PV 归约轴混淆。
34+- [Attention Tiling 的完整内存预算与流式边界](tiling_memory_budget_and_streaming_limits.md):按物理布局和同时存活对象核算 UB、L1、L0 与 workspace。
35+ 
36+- [Attention 模式组合保持逻辑 Token 身份](pattern_composition_preserves_logical_token_identity.md):从导航提炼变长、分组、分页、稀疏及可见性的组合条件。
37+ 
38+- [Sage 缓冲布局与容量冲突检查](sage_buffers_layouts_and_capacity_conflicts.md):保存GM和片上地址合同,揭示源布局表中的重叠与扩容缺口。
39+- [Sage 编译 Helper 与 API 适用范围](sage_compiler_helpers_and_api_scope.md):保留地址空间、寄存器和向量辅助实现的条件。
40+- [Sage 游标 Split-KV 与归并边界](sage_cursor_split_merge_and_mask_boundaries.md):保存任务首步、完整状态槽与因果位置合同。
41+- [Sage 分层取证与证据边界](sage_diagnostics_and_evidence_limits.md):保留拟合和中间量探针,避免余弦相似度与症状唯一归因。
42+- [Sage 四阶段流水与版本事件所有权](sage_pipeline_flags_and_version_ownership.md):保留mode4槽位协议,限定Bootstrap与双重等待结论。
43+- [Sage 量化与 K 平滑补偿合同](sage_quantization_and_smoothing_contract.md):保存INT8 QK与FP16 PV路径,修正精确性和类型转换推断。
44+ 
45+- [A2 FA Stage 资源与组件适用范围](fa_a2_stage_resources_and_component_scope.md):保存四段流水和轮转容量,修正平台与管线泛化。
46+- [FA 性能输入身份与覆盖分布](fa_benchmark_inputs_and_coverage_distribution.md):保存结构张量生成、长度分桶及多batch故障探测。
47+- [FA 复合 Causal Varlen 与 Sink 数学](fa_compound_causal_varlen_and_sink_math.md):修正局部列坐标、有限Mask与Sink分子缩放。
48+- [手写 FA 缓冲与状态故障模式](fa_handcrafted_buffer_and_state_failures.md):保留多任务覆盖和事件问题,修正共享UB与最终Max公式。
49+- [FA 布局与 SeqLen 接口变体](fa_layout_and_seqlen_interface_variants.md):区分BNSD、TND、分页格式和两类长度数组。
50+- [FA 分页与量化操作数合同](fa_paged_and_quantized_operand_contract.md):补齐物理Head偏移、尾页和反量化尺度。
51+- [FA 性能采集范围与汇总口径](fa_performance_capture_scope_and_aggregation.md):保存跨框架计时陷阱,修正公平下界与唯一测量方法的泛化。
52+- [FFA TensorList 前缀与组件合同](ffa_tensorlist_prefix_and_component_contract.md):区分拼接KV和双KV对角模型,保存FD LSE与量化入口语义。
53+- [Varlen L1 Carry 与 SameAB 约束](varlen_l1_carry_and_sameab_constraints.md):保存驻留优化条件,修正跨batch无条件复用与容量常量。
54+- [Varlen 任务前缀与 Sparse 边界](varlen_task_prefixes_and_sparse_bounds.md):保存四类累计量、TND偏移与尾块双重约束。
55+ 
56+- [FA 优化归因与内存证据](fa_optimization_attribution_and_memory_evidence.md):纠正Codegen唯一根因和大输入容量估算。
57+- [FA Q Head 打包与 Block 生命周期](fa_q_head_packing_and_block_lifetimes.md):保存GQA行映射和组件内部槽位,限制硬事件照抄。
58+- [FAQK 交接与 Vector Count 约束](faqk_handoffs_vector_counts_and_claim_limits.md):保存Token复用和批量布局知识,限制器件故障泛化。
59+- [FFA 正交 Batch 映射与流式状态](ffa_orthogonal_batch_mapping_and_streaming.md):保存双KV对角模型的四路GEMM和完整平面偏移。
60+- [KFC Chunk 原子累加与尾部合同](kfc_chunk_atomic_rescale_and_tail_contract.md):保留Prescale顺序和双缓冲状态,限定尾块性能案例。
61+- [KFC Launch ABI 与 Matmul 异步范围](kfc_launch_abi_and_matmul_async_scope.md):保留系统区和用户区分离,修正参数数目和硬编码地址窗口。
62+- [S1 Sink Softmax API 与缓冲复用](s1_sink_softmax_api_and_buffer_reuse.md):保存播种法和状态写回,纠正Tiling配置及行列循环混淆。
63+- [S1 三槽流水与排空身份](s1_three_slot_pipeline_and_drain_identity.md):修正模3时间标签和伪任务索引,保留L1Reuse配对。
64+- [S1 Workspace 单位与 Tiling 候选](s1_workspace_units_and_tiling_candidates.md):保存三段双缓冲公式,限定资源覆盖和缓存Shape条件。
65+ 
66+- [MLA 延迟 PV 槽位与跨任务状态](mla_delayed_pv_slots_and_cross_task_state.md):保存全局和局部PingPong、前取与事件所有权。
67+- [MLA FD Partial 表示与容量](mla_fd_partial_representation_and_capacity.md):区分归一化Partial和指数和,补齐全核发布及数据类型。
68+- [MLA Latent 组件与特化合同](mla_latent_component_and_specialization_contract.md):保存两段Score、V复用与模板差异,限定历史适用域。
69+- [MLA Split 规划与任务一一映射](mla_split_planning_and_task_bijection.md):保存贪心切分与双解码,补齐页容量及蛇形一致性。
70+- [MLA Tiling Word 偏移与字段冲突](mla_tiling_words_offsets_and_field_conflicts.md):保存400加17布局,纠正Float数组和绝对地址混淆。
71+- [Paged FA 组件布局与 ABI 修正](paged_fa_component_layout_and_abi_repairs.md):保存三项关键修复,补齐GQA基址和声明宏问题。
72+- [Paged FA 延迟 Finalize 与 Mask 限制](paged_fa_delayed_finalize_and_mask_limits.md):保存排空和Tail调用,纠正任意Mask套Causal及无符号边界。
73+ 
74+- [Softmax 在线重标定与有限 Mask 限制](softmax_online_rescale_and_finite_mask_limits.md):保存分块状态和压缩偏移,限定高阶API精度声明。
75+ 
76+- [Attention 在线输出归一化与最终除法必须一致](attention_normalized_output_and_final_division_must_agree.md):纠正设计中的二次除法,保存mask、分核和内存尺度条件。
77+ 
78+- [推理 RunInfo 初始化与 Query/KV 坐标读取](inference_runinfo_initialization_and_coordinate_reads.md):从 FA 模板的条件初始化与消费路径识别潜伏未初始化字段,按推理语义建立默认值及逐实例验证。
@@ -0,0 +1,38 @@
1+---
2+type: Guide
3+title: 推理 RunInfo 初始化与 Query/KV 坐标读取
4+description: 从 FA 模板的条件初始化与消费路径识别潜伏未初始化字段,按推理语义建立默认值及逐实例验证。
5+tags:
6+- task.debugging
7+- topic.attention
8+- topic.memory
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ascendc-port-orchestrator/kb/target/ascendc/ERROR_CORRECTIONS.md
12+ role: primary
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ascendc-port-orchestrator/kb/target/ascendc/patterns/domains/fa_class/templates/op_kernel/wholeport/wp_util_regbase.h
14+ role: header
15+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ascendc-port-orchestrator/kb/target/ascendc/patterns/domains/fa_class/templates/op_kernel/wholeport/wp_block_cube.h
16+ role: implementation
17+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ascendc-port-orchestrator/kb/target/ascendc/patterns/domains/fa_class/templates/op_kernel/wholeport/wp_kernel_train.h
18+ role: implementation
19+created_at: '2026-09-14T04:24:48Z'
20+updated_at: '2026-09-14T04:24:48Z'
21+platforms: []
22+---
23+ 
24+# 推理 RunInfo 初始化与 Query/KV 坐标读取
25+ 
26+共享 FA 模板中,训练实例可正常运行,而推理特化首次启用时暴露未初始化状态。固定来源 EC-86 报告过 query/kv 行偏移随干净重建变化的案例;本轮读取其固定代码可确认字段声明和部分消费路径,未运行该模板,也不把历史“+1 行”当作全部未初始化读取的固定症状。
27+ 
28+`wp_util_regbase.h` 的 COMMON_RUN_INFO 声明 preTokensPerBatch、nextTokensPerBatch 时没有成员默认值;RunInfo<true> 另有无默认值的 queryLeftPaddingSize、kvLeftPaddingSize。另一个类型 RunParamStr<true> 虽为 pre/next 字段设置了 MAX_PRE_NEXT_TOKENS,并不自动初始化 RunInfo 的同名字段。结构体名相近与内存独立性不能混淆。
29+ 
30+固定 `wp_block_cube.h` 中,CalcS1Coord 在 isInfer 分支读取 queryLeftPaddingSize,并对负 nextTokensPerBatch 加相反数;CalcS2Coord 在 isInfer 分支读取 kvLeftPaddingSize。它们不再以 hasAtten 保护这些读取。EC-86 将四个字段都归于 CalcS1Coord 的说法过宽:当前查看的函数并未读取 preTokensPerBatch,KV padding 由 CalcS2Coord 读取;preTokens 的其他消费者需单独追踪。
31+ 
32+`wp_kernel_train.h::InitUniqueRunInfo` 只在 hasAtten 且执行于 AIV 时为 runInfo.preTokensPerBatch/nextTokensPerBatch 赋值。这项赋值不能单独证明 AIC 对象、hasAtten=false 或新推理实例已得到同值;也没有在该函数里初始化两个 padding 字段。是否最终发生未初始化读取,还需检查具体对象的构造方式、值初始化、所有实际调用路径和写入覆盖,不能只由声明缺默认值就断言每次运行必错。
33+ 
34+为状态字段建立符合语义的初始化不变量,可以使用声明默认值或覆盖全部构造路径的显式初始化。无左 padding 的合同允许默认零;真实带 padding 输入则必须写入有效尺寸。pre/next 表示注意力可见窗口,默认值应符合无 mask 或相应 mask 模式,不能把四个字段机械清零而改变可见范围。每批、每任务重复使用对象时,也需要更新会变化的字段。
35+ 
36+仅在某个消费点加 `if constexpr(hasAtten)` 可能改变原本合法的推理坐标语义,并遗漏其他读点。先列出字段在 isInfer/hasAtten、AIC/AIV 及布局组合下的来源与读者,再选择一致的初始化方式。声明处修复减少漏改机会,但是否影响生成代码、布局或性能仍需实际构建验证,不能宣称天然零开销。
37+ 
38+验证应实际实例化 isInfer=true,覆盖无/有 mask、无/有左 padding、非零窗口及两类核的消费路径。使用不同 query/kv 行的可识别值检查坐标,比较重复执行与独立干净构建;后一项可暴露构建敏感的未初始化状态,但两次相等也不构成没有未定义行为的证明。原记录中未合并算子的历史状态不代表当前所有用户均触发此问题。
@@ -0,0 +1,34 @@
1+---
2+type: Guide
3+title: KFC Chunk 原子累加与尾部合同
4+description: 保留Prescale顺序和双缓冲状态,限定尾块性能案例。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.implementation.synchronization
9+- topic.precision
10+- topic.implementation.tiling
11+status: draft
12+sources:
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-kernel-recipes.md
14+ role: primary
15+created_at: '2026-09-11T05:27:29Z'
16+updated_at: '2026-09-11T05:27:29Z'
17+platforms: []
18+---
19+ 
20+# KFC Chunk 原子累加与尾部合同
21+ 
22+来源D128 FP16 BNSD路径每query块128行、KV chunk1024列,S/P两槽,FP32 max/sum按rows*8状态,P16降精度后PV。每chunk QK完成→softmax产生P和alpha;上一chunk PV完成后才能Oacc*=alpha,再提交本chunk AtomicAdd PV。初次Oacc必须初始化0,最后PV完成后才可除最终l和写输出。
23+ 
24+原子PV与非原子prescale同时读改写相同Oacc是竞争,即使所有PV自身使用atomic也不能保护Vector写。来源16列组NaN是其表现,不能当所有架构原子RMW粒度事实。每个chunk WaitIterateAll/End、prescale DMA完成及下次PV之间需完整依赖;若使用o2独立槽后显式合并,也需相同数学与生命周期。
25+ 
26+alpha=exp(mOld-mNew)要保留旧max再更新,按对应chunk槽存储,不能把下一chunk的alpha提前用于当前PV;若无sink且初m无效/全mask,需专门首步逻辑避免-inf差。16行softmax配两个8行Cast是该容量策略,不是count<8的所有Vector API都会trap。
27+ 
28+S按packed cols、P也按同样物理pitch才能供PV,原矩阵K/V的leading dimension与chunk有效宽度分开。Tail需要所有shape字段每次正确设置,不能先短尾再下个task满块沿用旧尾而只在tail条件时调用SetTail。source把非128倍数GEMM一概判坏与API一般tail能力不符,仅作为其路径未支持范围。
29+ 
30+pad到128并补K/V合法地址、softmax屏蔽无效score是一种处理候选,但source最终仍声明S%1024≠0的原子路径竞态未解决;不能因前文“标准解法”便声称任意长度已支持。Sq尾与KV尾是独立维,B/H/D和dtype范围也要与测过路径区分。
31+ 
32+source总UB约100KiB与F0列S64+P32+tmp32+state17+O16KiB加aux至少161KiB互相矛盾,需按真实同存活区核算。per-worker约1.2MiB GM也需含S/P双槽和Oacc及padding,不直接作为最小配置。状态若经int32逐字拷tiling,须保证POD size对4可整除或补尾字节、对齐及合法别名。
33+ 
34+来源流程在F0第7步后结束,并非包含所称11步完整实现;省略段不能补造成已验证代码。历史geomean0.693、458752自注意力22.7s及S1固定Sk128同Sq约5.7ms不是同一工作量,应分别记录接口和输入轴。
@@ -0,0 +1,48 @@
1+---
2+type: Guide
3+title: KFC Launch ABI 与 Matmul 异步范围
4+description: 保留系统区和用户区分离,修正参数数目和硬编码地址窗口。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.implementation.synchronization
9+- topic.precision
10+- topic.implementation.tiling
11+status: draft
12+sources:
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-kernel-recipes.md
14+ role: primary
15+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-sink-handcraft-recipe.md
16+ role: supporting
17+- resource: https://gitcode.com/qq_41667743/cannbot-knowledge/blob/a70025d581ae45da8a93b4938b5ba19da6fa6f92/knowledge/ops/ascendc/apis/simd_api/adv_api/cube_compute/matmul_kernel/iterate_all.md
18+ role: constraint
19+- resource: https://gitcode.com/qq_41667743/cannbot-knowledge/blob/a70025d581ae45da8a93b4938b5ba19da6fa6f92/knowledge/ops/ascendc/apis/simd_api/adv_api/cube_compute/matmul_kernel/wait_iterate_all.md
20+ role: constraint
21+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ascendc-port-orchestrator/kb/target/ascendc/ERROR_CORRECTIONS.md
22+ role: supporting
23+created_at: '2026-09-11T05:27:29Z'
24+updated_at: '2026-09-14T04:24:48Z'
25+platforms: []
26+---
27+ 
28+# KFC Launch ABI 与 Matmul 异步范围
29+ 
30+KFC混合入口需要按实际编译/runtime约定注册Matmul对象、系统workspace和tiling;AIC只执行注册分支、AIV调度是该入口模型,不是所有Matmul都按此写。多个框架Matmul对象可合法注册,不应混同底层FAQK实例共享资源问题。
31+ 
32+来源workspace属性位于第8参、地址观测差4GiB、系统保留16MiB、消息区清64KiB及launch中间槽nullptr为特定CANN9.1/dav2201案例。它们不能推广为地址恒等式或任何入口固定参数位置;尤其原文称14参却只列9个名字,不能凭标题构造ABI。Host和Device真实签名、属性及指针应逐字段匹配。
33+ 
34+将用户workspace作为独立普通GM_ADDR可避免从特殊系统参数错误派生基址,但仍需真实系统区容量与Host allocation边界。不能把GetSysWorkSpacePtr()与入参差值直接手工加减4GiB,也不能仅因一次测试通过就只清固定64KiB。workspace初始化若为每launch必需应计入完整执行成本。
35+ 
36+MatmulApiTiling可推导类型、形状、bias及buffer设置,GetTiling状态必须检查;之后手改shareL1/L0C/depth可能破坏计算器结果,不能据“手填TCubeTiling挂死”宣称所有明确合法手工tiling都禁止。类型级transB与运行时SetTensorB的支持组合按实际实现验证,source类型级转置错值仅约束其配置。
37+ 
38+IterateAll<false>为异步,消费结果前使用匹配WaitIterateAll及End;已有API分别有GM和Local输出重载,不能因GM版本正确就直接替UB输出。单次输出容量至少singleCoreM*singleCoreN。Iterate/GetTensorC也有合法同步模型,来源慢4倍或数值错的尝试不能形成“唯一又快又对的调用形态”。
39+ 
40+同一对象先提交c+1再等待c需确认API允许并发请求和可区分结果;双GM槽不自动建立请求队列。AIV worker数应从真实核型/资源与task规划推导,固定min(40,tasks)不是跨平台值,闲worker也要遵循注册/同步参与协议。
41+ 
42+## 编译分支与系统 Workspace 初始化的版本条件
43+ 
44+插件错误汇总 EC-57 记录,某 CANN 9.0.0 DYNAMIC_MODE 构建只设置全局 ASCENDC_MATMUL_AICORE,未使 AIC/AIV 分别选择 KFC 的服务端/客户端分支;来源用分 pass 的 SPLIT_CORE_CUBE、SPLIT_CORE_VEC 修正。其后更正又称 CANN 9.1.0 可从 __DAV_CUBE__/__DAV_VEC__ 派生这两个宏。因此可复用要求是核对真实编译命令、预处理分支与注册协议,不是所有 SDK 必须手动加同一组宏。不能把两类宏同时全局定义来代替分 pass 配置,也不能把任何 507015 都归因于这些宏。
45+ 
46+EC-58 的构建中 matmul_intf.h 由显式 include 或相应 pass 的 interface library 传播;kernel_operator.h 的存在不单独证明 MatmulImpl/REGIST_MATMUL_OBJ 已可见。检查实际头文件和 include 链,不能将别的高级 API 的传递包含方式统一套用。
47+ 
48+EC-68 在 CANN 9.1.T500 的独立 ACLRT_LAUNCH 案例里记录了先设置 SetSysWorkspaceForce(workspaceGM)、再初始化 TPipe/Matmul 及派生用户区的修复。这是该入口依赖的系统基址合同;其 16MiB 偏移不应固化到所有版本。Host 分配必须覆盖真实系统保留区及全部用户区,单纯把普通入参命名为 workspace 不证明运行时使用了这个基址。上述版本记录来自固定说明,本轮未验证对应 SDK 分支。
@@ -0,0 +1,38 @@
1+---
2+type: Guide
3+title: Attention 的逻辑 token、head 与物理缓存地址
4+description: 分开处理 TND batch 边界、GQA head 映射、稀疏 token 选择与 page table,防止布局组合后混用坐标。
5+tags:
6+- task.design
7+- topic.compute.indexing
8+- topic.data_movement.layout_transform
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/tilelang2ascendc-ops-generator/skills/tilelang2ascend-translator/references/attention-patterns/references/page_attention_pattern.md
12+ role: primary
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/tilelang2ascendc-ops-generator/skills/tilelang2ascend-translator/references/attention-patterns/references/tnd_pattern.md
14+ role: supporting
15+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/tilelang2ascendc-ops-generator/skills/tilelang2ascend-translator/references/attention-patterns/references/head_sharing_pattern.md
16+ role: supporting
17+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/tilelang2ascendc-ops-generator/skills/tilelang2ascend-translator/references/attention-patterns/references/topk_kv_sparse_pattern.md
18+ role: supporting
19+created_at: '2026-09-10T14:32:33Z'
20+updated_at: '2026-09-10T14:32:33Z'
21+platforms: []
22+---
23+ 
24+# Attention 的逻辑 token、head 与物理缓存地址
25+ 
26+Ascend C Attention Kernel 组合变长、共享 head、分页和稀疏访问时,应分开建立逻辑任务坐标与物理地址,再把它们接到同一 QK/softmax/PV 数据流。它们不要求统一采用某一条固定流水或强制增加 Gather stage。
27+ 
28+对无前导零的累计长度数组 L,batch b 的起点为 `b==0 ? 0 : L[b-1]`,长度为 `L[b]-start`。Q 与 KV 分别恢复边界;若接口提供带前导零的数组或非累计长度,公式随之改变。任务到 batch 的映射不能在变长时使用固定每 batch 块数,尾块也按该 batch 有效长度计算。
29+ 
30+对按连续 head 组共享 KV 的 GQA,令 `G=Hq/Hkv`,需要正整数 G,`kv_head=floor(q_head/G)`;MQA 对应 Hkv=1。K 与 V 必须使用一致的映射,输出仍按 query head 排列。Host 显式展开可以作为参考计算方式,但实际 Kernel 无需为了 head 共享而物化完整重复 K/V。
31+ 
32+当索引合同明确表示逻辑 token 时,先检查 sentinel、长度和可见性,再按 token 寻址。若索引表示 block、page 或已经编码的物理 offset,则应使用该合同的转换,不能重复套 token 公式。例如原库 [NsaSelectedAttentionInfer](../../operators/aclnn/transformer/attention/nsa_selected_attention_infer.md) 的 selected index 表示区间块,不能当作本卡的逐 token 索引。
33+ 
34+对 page 大小 P,逻辑 token t 映射为 `logical_page=t/P`、`row=t%P`、`physical_page=table[b,logical_page]`。在连续 `[num_pages,P,Hkv,D]` 布局下,元素 offset 为 `((physical_page*P+row)*Hkv+kv_head)*D+d`;若实际布局是 `[num_pages,Hkv,P,D]`,head 和 row 的乘法顺序不同。必须分别验证 page table 范围、物理 page 范围及 K/V 各自跨度。
35+ 
36+连续逻辑 tile 只有在实际物理段可由当前搬运接口表达时才能直接加载。跨 page、离散选择或复杂布局可拆分搬运,或先 gather 成密集 workspace;并非所有分页都需要 gather,也不能只凭“逻辑连续”假定物理连续。
37+ 
38+Gather 的 K/V 行顺序与 score 列顺序必须一致;V 暂存要保留到 PV 消费完成。无效行的零填充保护搬运和矩阵输入,但还需要 score 可见性处理。验证使用不规则 page table、非等长 batch、不同 head 数据、负 sentinel、尾块和 ring 复用,并检查每个逻辑 token 的来源。
@@ -0,0 +1,33 @@
1+---
2+type: Guide
3+title: MLA 延迟 PV 槽位与跨任务状态
4+description: 保存全局和局部PingPong、前取与事件所有权。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.implementation.synchronization
9+- topic.implementation.tiling
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-mla-paged.md
13+ role: primary
14+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-mla-handcraft-recipe.md
15+ role: supporting
16+created_at: '2026-09-11T05:34:07Z'
17+updated_at: '2026-09-11T05:34:07Z'
18+platforms: []
19+---
20+ 
21+# MLA 延迟 PV 槽位与跨任务状态
22+ 
23+来源通用MLA先prologue QK0→softmax0而无PV,随后n1..nLoop生产QKn/softmaxn并消费PV(n-1)/rescale(n-1),最后可预取下一有效task的首块以延续流水。该调度匹配所选组件;单块锁步也只有重新设计完整事件和资源协议后才能替换,不是数学attention只能延迟PV。
24+ 
25+全局pp跨task保留,S元素offset=core*32768+(pp%2)*16384,P=core*65536+(pp%2)*32768,OTmp=core*131072+((pp-1)%2)*65536,OUpdate=core*65536。生产与消费对同一slot达成一致,pp0之前不能取pp-1;locPingPongIdx是组件内部另一个状态,taskPingPongFlag保护跨task统计,不能全部统一为一个奇偶变量。
26+ 
27+lookahead要找到下一真正有效process并保存其batch/head/split/首page身份;跳过空任务不能遗漏首块,已预取的task进入主循环也不能再prologue一次。下一task softmax可能早于上一task最后rescale,m/l及alpha必须按task槽隔离,glFlag[2]真实初始化与组件消费协议一致;非null数组不是通用替代所有状态初始化。
28+ 
29+QK从blocktable[realBatch*maxBlocks+startKV/128+n]解析latent和Rope各自stride。PV来源调用不显式传V GM,依赖专用组件保留的latent/K资源关系;换通用PV必须补齐数据来源,不能假定省掉指针后仍自动读到正确V。相邻task输入不同不得无条件重用旧L1。
30+ 
31+来源AIC预置8+2+8+6+1=25、AIV3+7+5=15,实际预置/析构和排空顺序需匹配组件;“反序”文字与所列V_MTE2 drain4,0,1,2,3不是简单逆序,应按依赖核验。裸CrossCoreWait与Set<0x2,PIPE>配对可能由封装默认语义实现,不能简单比较模板文字数字判不配对或适用于所有版本。
32+ 
33+资源对象只构造一次、先读取tor再构造Epilogue;源片段在声明tor前已使用它,说明仍是有省略的示例。AIV logical core=blockIdx/subBlockNum是该混合任务映射,原文“裸编号0与20都映0”并非这条公式本身推出,应核对真实编号范围。宏分流、同步基址和Host核型也需与真实编译结果匹配。
@@ -0,0 +1,32 @@
1+---
2+type: Guide
3+title: MLA 扩展假设与 Attention 特化范围
4+description: 保留扩展点而不把未验证路线写成平台能力。
5+tags:
6+- task.optimization
7+- topic.implementation.tiling
8+- topic.performance.analysis
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-tiling-design/references/flashattention/foundation/composition.md
12+ role: primary
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-tiling-design/references/flashattention/foundation/extension_points.md
14+ role: supporting
15+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-tiling-design/references/flashattention/patterns.md
16+ role: supporting
17+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-tiling-design/references/flashattention/design/_governance.md
18+ role: supporting
19+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-tiling-design/references/flashattention/specialization/feature_flags.md
20+ role: supporting
21+created_at: '2026-09-11T04:27:49Z'
22+updated_at: '2026-09-11T04:27:49Z'
23+platforms: []
24+---
25+ 
26+# MLA 扩展假设与 Attention 特化范围
27+ 
28+MLA涉及潜在KV表示、投影及可能的矩阵吸收,不能仅把KV head设为1,再在V1后加入两个逐元素操作就视为完成。共享latent cache也不意味着所有外部接口必须Hkv=1;QK score维度Dqk、value维度Dv及投影shape应由实际方程确定,投影FLOP和流量可能显著。
29+ 
30+来源提出用int32中间“修复FP32精度”等说法但无范围、scale和溢出证明,不能采为正确实现。应保存为未建立的假设,并保留真正可复用的扩展位置:Q/K预处理、score偏置/可见性、P格式转换、V/输出投影及最终后量化。每个位置都需要输入输出单位、布局、精度和依赖合同。
31+ 
32+特化可按shape、dtype、mask及硬件能力选不同Bq/Bk、G分片、流水深度和执行比例,但模板分支不自动证明所有组合可编译/可执行。容量不等式、映射覆盖、数值与同步证据应分别建立;来源的治理步骤、强制审批/Agent角色和声称最优的流程不作为知识库指令。
@@ -0,0 +1,35 @@
1+---
2+type: Guide
3+title: MLA FD Partial 表示与容量
4+description: 区分归一化Partial和指数和,补齐全核发布及数据类型。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.implementation.synchronization
9+- topic.implementation.tiling
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-mla-paged.md
13+ role: primary
14+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-mla-handcraft-recipe.md
15+ role: supporting
16+created_at: '2026-09-11T05:34:07Z'
17+updated_at: '2026-09-11T05:34:07Z'
18+platforms: []
19+---
20+ 
21+# MLA FD Partial 表示与容量
22+ 
23+FD归并需要先确定partial表示。若保存未归一化a_j、m_j、l_j,则m=max(m_j)、w_j=exp(m_j-m),O=sum(w_j*a_j)/sum(w_j*l_j);若保存已归一化o_j=a_j/l_j,则分子为sum(w_j*l_j*o_j)。不能给未归一化a再次乘l,或给已归一化o漏乘l。
24+ 
25+也可保存o_j与logZ_j=m_j+log(l_j),按exp(logZ_j-maxLogZ)加权归并。这解释了为什么某实现只需每row/split一个统计float,但source仅列gl一float同时用m/l公式,未说明其真实编码,本批不能推断gl就是普通sum或已含max。必须核对实际MLARescaleO写端和FDRescaleO读端后绑定格式。
26+ 
27+maxSplit来自tiling后再分配:oFd至少Dc*H*totalTokens*maxSplit*sizeof(partialType),l按该实现实际编码及同split布局。来源较新表partial和OUpdate为FP32,旧表globalo写FP16会少配一半;GemmType输出半精度不决定中间Update类型。源“最小S/P尺寸”比所列slot公式更宽裕,不能称已证明最小。
28+ 
29+S公式每core32768FP32、P每core65536元素、OTmp每core131072FP32、Update65536FP32只是地址跨度下界,还要检查最大row/col physical tile是否装入每半槽。资源分配和实际传入Element类型一致,不能同一offset有时字节有时元素。
30+ 
31+所有partial写完后才进入全AIV归并,局部HardEvent drain不能单独保证其他核完成。CrossCoreBarrier模式、PIPE_MTE3及参与者按真实launch核型确定,闲任务核不可在函数头return躲开barrier。SetAtomicNone、SetMaskNorm和mask复位需符合后续操作与调用者状态约定。
32+ 
33+归并阶段按物理AIV编号分发,不是主循环logical AIC组编号。ComputeEleNum6144、Dc512得到headsPerLoop=min(12,16)=12,H16为12+4尾头;若Dc>6144会得0,须守卫而非继续ceil除零。每batch每qtoken与head段只写一次,空KV/全mask/非有限统计要显式处理,maxSplit比实际大时未写槽不得参与。
34+ 
35+来源同篇开头声称20case全过,后文又说只验证未split的4点、FD仍待实现;不能以总标题断言当前recipe FD已实测通过。gate3.047/3.116/约3.19和历史33.5倍分解墙钟属于不同批次/口径,不能拼成新的性能证据。
@@ -0,0 +1,31 @@
1+---
2+type: Guide
3+title: MLA Latent 组件与特化合同
4+description: 保存两段Score、V复用与模板差异,限定历史适用域。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.implementation.synchronization
9+- topic.implementation.tiling
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-mla-paged.md
13+ role: primary
14+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-mla-handcraft-recipe.md
15+ role: supporting
16+created_at: '2026-09-11T05:34:07Z'
17+updated_at: '2026-09-11T05:34:07Z'
18+platforms: []
19+---
20+ 
21+# MLA Latent 组件与特化合同
22+ 
23+该MLA接口Q[tokens,H,512]与QRope[tokens,H,64]分开,paged K latent与KRope分别宽512/64,score=scale*(Q*Klatent^T+QRope*KRope^T),V为同一latent512维,输出[tokens,H,512]。默认scale1/sqrt576只适用于此明确维度与默认合同,不能误用sqrt512,也不能对两段score独立softmax。输入RoPE部分若已经编码,不应再次应用位置变换。
24+ 
25+来源q1–4 decode/MTP、KV≥128且尾页可不满、H16/32/64通用、H128 TP1/AMLA、FP16/BF16是其声明范围,不是所有MLA定义。q>4仅打印告警不会实际执行接口拒绝;历史ATB只支持q1与KV16384以上未经其原验证的区间须版本限定。黄金应计算该latent域语义,不能擅自解压V再比较。
26+ 
27+通用组件候选MmadAtlasA2MLAQK/MLAPV是该版本普通struct,FAIQK则可能是模板;BlockMmad实际GemmType参数与某伪码裸类型不同。Softmax、RescaleO、FDRescaleO参数个数也不能照另一版本类比。TP1候选MLAQKTp1Spec/MLAPVTp1Spec、AMLAPVTp1Spec及对应MLATP1/AMLATP1 epilogue只提供导航,换名字并不完成全部task布局特化。
28+ 
29+源H128一处称K常驻L1,一处称4-block且无K常驻,不能同时采为已定架构。L1/L0静态Shape<128,128,576>不证明完整tile一次驻留每级:单份128*576*2=147456B已大于典型64KiB L0A/B,需要组件真实内部分块。L1.Q/K/P分区和stage份数也要逐项计算,K576不自动证明跨task复用。
30+ 
31+预处理__DAV_CUBE__/__DAV_VEC__分支、FA模板特化及KFC __mix__是不同编译/调度模型。来源某组合死锁不能推出__mix__属性自身必错;匹配实际二进制核型、组件事件协议及entry才是可复用约束。单kernel、小GM轮转、专用两段QK、KV复用、split-KV是性能候选,不能宣称缺任一项就必慢一个数量级。
@@ -0,0 +1,33 @@
1+---
2+type: Guide
3+title: MLA Split 规划与任务一一映射
4+description: 保存贪心切分与双解码,补齐页容量及蛇形一致性。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.implementation.synchronization
9+- topic.implementation.tiling
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-mla-paged.md
13+ role: primary
14+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-mla-handcraft-recipe.md
15+ role: supporting
16+created_at: '2026-09-11T05:34:07Z'
17+updated_at: '2026-09-11T05:34:07Z'
18+platforms: []
19+---
20+ 
21+# MLA Split 规划与任务一一映射
22+ 
23+来源split启用候选maxKV≥blockDim*2*128、batch≤0.8*blockDim、maxQ==1;blockDim20时阈值5120,历史测试4096→6144跳档不证明精确边界6144。MAX_SPLIT8与组件可能允许KV_SPLIT_MAX64是两层上限,不能混用;planner还需每段至少512token的启发式。
24+ 
25+贪心从各batch1段开始,总分配<blockDim时选择alignedKV/split最大且split<8、alignedKV/(split+1)≥512者递增,无候选即停。实际page段长ceil(ceil(KV/128)/assignedSplit)*128,实际段数ceil(KV/段长)可能少于assignedSplit;globalMax若取分配数可作容量上界,但任务前缀必须用实际段数,merge只读真实写出分区。
26+ 
27+连续前缀路径process定位batch,再curN=process-prefixBatch且qHeadSplitIdx0;蛇形路径按headSplit×maxSplit解码。若用realProcess变换batch,却用原process解head/split,可组合出不对应同一task的坐标,尤其波次尾不整齐时需证明完整双射。所有轴应由同一逻辑task身份解码,或给出严格等价映射;不能只以每核process递增证明无重漏。
28+ 
29+KV空batch、curN≥该batch实际split均不应生产数据,但同一pair两侧必须作相同跳过,后续全AIV barrier还需全员到达。末段长度KV-curN*perCore在合法范围算,nLoop=ceil(curKV/128),Q pack行数qLen*actualHeadSplit并round16,尾head/尾KV同时约束。
30+ 
31+源sum(KV长度)≤numBlocks*128不足以证明独立分页容量:两个129token序列需要4个独占page,258≤3*128仍会错误接受3页。应核验每个有效table项范围、page内valid和共享规则;若允许共享page,sum长度检查也不是必要条件。mask前缀与实际压缩/广播布局分开。
32+ 
33+QHead候选列表128/64/32/16/8/1按ceil(log2(qLen))选择并cap,qLen0不得求log。H128按每token重组的spec不能沿用通用batch条目数/地址表。排序或task预取只在实际数据身份不变时有cache收益,不能把多batch称“免费性能”。
@@ -0,0 +1,35 @@
1+---
2+type: Guide
3+title: MLA Tiling Word 偏移与字段冲突
4+description: 保存400加17布局,纠正Float数组和绝对地址混淆。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.implementation.synchronization
9+- topic.implementation.tiling
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-mla-paged.md
13+ role: primary
14+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-mla-handcraft-recipe.md
15+ role: supporting
16+created_at: '2026-09-11T05:34:07Z'
17+updated_at: '2026-09-11T05:34:07Z'
18+platforms: []
19+---
20+ 
21+# MLA Tiling Word 偏移与字段冲突
22+ 
23+MLA来源使用32位word表,通用分配400+17*batch个uint32,TP1按token重组时需其实际条目数。不是把所有字段数值转为FP32;整数长度/偏移按uint32读,tor仅将float位模式写同一word再用float视图读。gTilingFp64变量实际GlobalTensor<float>,名字不代表double。
24+ 
25+recipe头区为0batch、1H、2Dc、3numBlocks、4blockSize、5maxBlocks、6tor、7KVheads、8headSize、9paraSize、10headSplitSize、11headSplitNum、12mask、13Dr、14maxKV、15split、16maxSplit、17maxQ、18totalTokens、19formerTasks、20tailTasks、21processNum。旧paged说明把processNum放19,与recipe21冲突,必须由真实Host writer和Device reader同一版本决定,不能混抄。
26+ 
27+每batch段起点header[8]+header[9]*b,来源默认400+17*b:0qLen、1kvLen、2排序前batch号、3pageSize、4/5Q偏移高低字、6/7QRope、8/9mask、11/12l、13/14oFd、15KV每核段长、16该batch实际split数。第10槽未定义不能擅自赋义;40位以上偏移重建需先uint64(high)<<32再或low。
28+ 
29+来源多处称这些数是GM绝对地址,但Host公式计算的是相对各张量的元素offset,后续GlobalTensor[offset]也再次加基址;应明确是元素偏移还是绝对byte地址,不能两者兼用。Q和O同shape可复用相同元素offset,却仍是不同base;不意味着把Q指针当O指针或允许原地别名。
30+ 
31+l偏移累计H*qLen*globalMaxSplit,oFd基础offset累计H*qLen*Dc,使用时再乘globalMaxSplit;已乘split的量不得再乘。FP16输入、FP32partial及word表各自元素宽度不同。per-batch排序后必须保留原Q/QRope/O/页表对应关系,不能把sorted索引直接当原batch。
32+ 
33+CUTASK前缀在word25起,声明最多130项则batch+1≤130,还要不与其他header字段重叠。前缀终值与processNum、实际每batch split总数一致,0代表另一解码路径时不得同时写非空split表却让Device走蛇形。分配前检查溢出和完整表长度。
34+ 
35+tilingKey来源通用0/1、AMLA4/5、TP1候选7/8,H128分支及numTokens%blockDim≤10、batch≤40是特定策略;早文写%20是假设20核,不能搬到其他blockDim。Host须检查GetMLATilingParam状态,不能先用未填tiling尺寸分配工作区。
@@ -0,0 +1,34 @@
1+---
2+type: Guide
3+title: 分块 Attention 的 online softmax 状态与归并
4+description: 保持 max、指数和与未归一化加权和一致,正确处理跨 KV 分块缩放、Split-KV 和分块 softmax 输出。
5+tags:
6+- task.design
7+- topic.compute.reduction
8+- topic.memory
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-tiling-design/references/flashattention/foundation/fundamentals.md
12+ role: primary
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-blaze-best-practice/references/scenarios/cube-matmul-then-vector-softmax-reduce/online-softmax-reduce-epilogue-design.md
14+ role: supporting
15+created_at: '2026-09-10T14:01:59Z'
16+updated_at: '2026-09-10T14:01:59Z'
17+platforms: []
18+---
19+ 
20+# 分块 Attention 的 online softmax 状态与归并
21+ 
22+对一个 query 行,分块 Attention 可以保留当前最大 score m、指数和 l,以及未归一化加权和 a,逐个 KV 块更新,而不在 GM 中物化完整 attention score 矩阵。score 应包含该接口所需的 scale、mask 等变换;max、sum 和加权和必须对应同一组有效 key。
23+ 
24+对新块的有限有效 score s 和对应 value V,令 m_new=max(m,max(s))、alpha=exp(m-m_new)、p=exp(s-m_new),更新 l_new=alpha*l+sum(p),a_new=alpha*a+p·V。处理完全部 KV 后再求 a/l。旧指数和和旧加权和必须一起重缩放;只更新 max/sum 而遗漏 a 的缩放会改变结果。
25+ 
26+每个并行 KV 分区可以独立形成 (m_i,l_i,a_i)。归并时令 m=max_i(m_i),再计算 l=sum_i(exp(m_i-m)*l_i)、a=sum_i(exp(m_i-m)*a_i),最后输出 a/l。由此可见,Split-KV 需要显式归并,不能直接相加各分区已归一化的输出。该等式针对同一 query 的不重叠、完整 key 分区;重复或遗漏 key 会改变 attention。
27+ 
28+空分区和全掩码行需要显式语义。直接对初始化的 m=-inf 做 -inf-(-inf) 会产生 NaN,不能只把 l 和 a 初始化为零就假定所有空路径正确。可按接口定义跳过无贡献分区,并为整行无有效 key、NaN/Inf score 等情况规定行为。上述有限有效 score 的公式本身不替接口决定这些特殊值语义。
29+ 
30+对于需要输出完整 softmax 概率的 Matmul+softmax 场景,若某 tile 保存的是基于历史最大值 m_t 的 exp(s-m_t),最终归一化需乘 exp(m_t-m_final)/l_final。保存了指数块却丢失其对应 m_t,会使后续跨核归一化用错缩放;这与 Attention 只保留加权和的存储方案不同。
31+ 
32+状态应按 query/head/task 隔离,并与流水槽的生命周期一致。大 value 维度可能使 a 无法常驻 UB,可采用 GM 暂存,但每次读改写仍需遵守相同数学状态更新和同步条件。GQA 合并 query heads 是减少重复 KV 搬运的一种调度选择,不是 attention 正确性的必需条件。
33+ 
34+Ascend C 的具体 Softmax 接口还有独立的 dtype、形状和临时空间约束,不能由数学公式推出任意调用均受支持。验证可用小规模明确计算的 attention 对照,改变 KV 分块和分区数,覆盖尾块、mask、空分区、不同 query/value 维度以及多个流水槽;浮点归并顺序变化需要按目标数值合同评估。
@@ -0,0 +1,37 @@
1+---
2+type: Guide
3+title: Online Softmax 状态槽的接续与最终选择
4+description: 明确 max、sum、缩放因子和输出累加器的任务身份,避免双槽递推丢失或重复计数。
5+tags:
6+- task.design
7+- topic.attention
8+- topic.memory
9+- topic.compute.reduction
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/playbooks/P38_softmax_state_resident.md
13+ role: primary
14+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/preconditions/P38.yaml
15+ role: supporting
16+created_at: '2026-09-11T01:49:44Z'
17+updated_at: '2026-09-11T01:49:44Z'
18+platforms: []
19+---
20+ 
21+# Online Softmax 状态槽的接续与最终选择
22+ 
23+online softmax 的状态沿同一 query/head 的 KV 分块接续。分配两个 max/sum/exp 物理槽,只能使不同版本有独立存储;第二块仍必须读取第一块更新后的状态,不能让两个槽独立从初始值轮流计算。来源模板取 `s2%PRELOAD_NUM`,却没有展示 inSumTensor/inMaxTensor 如何指向真正的前一状态,因此不足以证明跨块递推正确。
24+ 
25+对于串行递推版本,第 k 块的新状态已经包含此前全部有效块,最终使用最后一个有效版本;把最后两个槽直接相加会重复累计。只有明确采用互不重叠 KV 子集的独立 partial,才可使用[带 max 重缩放的归并](online_softmax_state_and_partition_merge.md),并一并处理未归一化输出累加器。softmaxExp 若保存前后 max 的缩放因子,还要保留到对应输出 rescale 的消费者结束。
26+ 
27+max/sum 的数学状态通常按 query 行保存,exp 及具体 API 临时区可能采用复制、广播或其他布局,三者并不必然等大。状态容量不能仅由 KV 循环次数 s2Loops 推出;较长 KV 增加迭代数,不必线性增加每行状态。Mmad 的输入输出存储层次也不能由这份 UB 状态模板直接推定合法。
28+ 
29+如果确实三份状态每槽各 S 字节、深度 p,则基础占用3*S*p;一般应逐缓冲求和,另计输出累加器、输入、scratch 和对齐。`slot*(S/sizeof(T))` 转元素偏移要求 S 可整除且每个槽满足地址对齐,使用的 T 必须是实际视图类型。L0A/L0B/L0C 是不同存储资源,不能一并当作 UB 字节扣除。
30+ 
31+初始化需满足实际 Softmax 接口的首轮模式或显式旧状态。数学空状态的 m=-Inf、l=0 仍需处理全空/全mask的 -Inf-(-Inf);-FLT_MAX/-1e30 不是与 -Inf 对所有输入等价的替代。exp 是否需初值以及初值为0/1,由其真正含义和首次读写顺序决定,不能统一设0。
32+ 
33+在初始化阶段建立长期缓冲有利于维持分配和状态生命期;不能把循环内反复 InitBuffer 简化为固定的函数调用开销,实际还需遵守资源管理接口是否允许重初始化及如何释放。单槽状态可以与其他独立流的搬运并行;双槽也不保证更快。p>2 或运行时选择是否可用,应由资源/API合同决定,而非来源的建议数字。
34+ 
35+状态 spill 到 GM 时,须保持 query/head、状态分量和版本的完整字节布局,在 reload 前完成发布,在下一次覆写前等消费者结束。不能只复制 max 而丢失 sum/输出状态。最终输出与所有延迟消费者完成后才可释放槽。
36+ 
37+验证0/1/2/3个KV块、多次轮转、最后一块被mask、切换query/head、不同value维度以及全mask行。用不分块计算对照,并检查每个中间版本覆盖的 token 集合。
@@ -0,0 +1,31 @@
1+---
2+type: Guide
3+title: Paged FA 组件布局与 ABI 修正
4+description: 保存三项关键修复,补齐GQA基址和声明宏问题。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.implementation.synchronization
9+- topic.implementation.tiling
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-paged-handcraft-recipe.md
13+ role: primary
14+created_at: '2026-09-11T05:34:07Z'
15+updated_at: '2026-09-11T05:34:07Z'
16+platforms: []
17+---
18+ 
19+# Paged FA 组件布局与 ABI 修正
20+ 
21+该FA实例Q/K/V/P/O为half、S/OTmp/Update为float,K ColumnMajor、V RowMajor;MLA的V消费视图不同,不能直接套用。FAIQK/FAIPV<true,false>、EpilogueOnlineSoftmax与RescaleO三GemmType参数均为所选版本,真实Tail名来源明确FAITailQK/FAITailPV,不能把导航别名FAIQKTail当已存在符号。
22+ 
23+QK在L1分配Q一份、K两份,128*128*2*(1+2)=98304B=96KiB;PV构造从L1_QK_SIZE之后分配可避免源实例重叠,实际资源还需PV/Tail/额外区总预算。不同组件共用相同区只有生命周期互斥才合法,不能只因构造传了offset就认定整体无冲突。
24+ 
25+AIV softmax以P物理rowStride512写,PV必须用同stride读取即LayoutP(rowNum,validCols,512);two-arg只在其默认stride恰好一致时等价。OTmp行距为alignUp(D,16),O完整行距为Hq*D。S/P/OTmp虽然slot元素offset形式相同,但buffer基址与dtype不同。
26+ 
27+GQA任务已算kvHeadIdx/qHeadIdx,Q/O显式加head偏移;源NO_MASK调用却直接传gK/gV基址且无kvHead入参。除非组件另有可证明的head视图来源,必须给K/V正确kvHeadIdx*D偏移,不能只传strideKV=Hkv*D就自动选中KV头。masked片段出现gmKOffset/gmVOffset也要真实定义。Hkv1测试会掩盖此遗漏。
28+ 
29+若GM_ADDR是宏展开为__gm__uint8_t*,写GM_ADDR q,k会让后者不带星号,独立声明或真正指针类型alias才能一致;这不是C++禁止所有逗号声明。source Step2要求逐声明、Step9却再次给逗号列表,不能照抄后者。类内显式模板特化也需目标编译器支持或改为合法特化组织。
30+ 
31+FA Params13个GM指针加entry hardwareSync共14参,Host顺序、POD FATilingData的uint64字段对齐及float scale位模式须一致。Host先GetFATilingParam再按结果分配四段,来源步骤先malloc尺寸后填tiling顺序颠倒。四段可在一块充分对齐allocation中划分并传独立指针,不必四次malloc;禁合并分配没有数学或ABI依据。
@@ -0,0 +1,35 @@
1+---
2+type: Guide
3+title: Paged FA 延迟 Finalize 与 Mask 限制
4+description: 保存排空和Tail调用,纠正任意Mask套Causal及无符号边界。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.implementation.synchronization
9+- topic.implementation.tiling
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-paged-handcraft-recipe.md
13+ role: primary
14+created_at: '2026-09-11T05:34:07Z'
15+updated_at: '2026-09-11T05:34:07Z'
16+platforms: []
17+---
18+ 
19+# Paged FA 延迟 Finalize 与 Mask 限制
20+ 
21+普通路径4个128KV页组成最多512列stack,preLaunch2给3槽,QK/softmax当前stack生产P,PV/rescale消费两stack前。收尾需消费剩余真实stack;primary rescale不finalize,secondary仅当消费序号==totalStackSeqNum-1才除l写最终O,不能所有drain步isLast恒1。
22+ 
23+一到多个stack、KV0、未满128/512尾、多个task复用须独立验证。若按无符号kvLoopNum-1判尾,0循环数会下溢;loop推进的slot与实际生产序号应对应,不能等待未生产P。源码AIV预置标10项实际11,AIC25项与骨架旧表18又不同,需目标组件全生命周期核验。
24+ 
25+Mask0/1/2必须区别。源对所有maskType!=0套右下causal noSkip公式,却把MASK_SPEC解释为任意逐位mask;任意mask不能据query位置跳掉本可能可见的key,也不能假定只有对角块需要加载mask。只有明确压缩causal/band/prefix合同及正确索引才可结构性裁剪。
26+ 
27+令qStart为query块起点、qEnd为有效行exclusive终点、delta=Sk-Sq。右下causal候选最大可见key exclusive上界为max(0,min(Sk,qEnd+delta)),逐行边界仍是k≤q+delta。delta和noSkip-qRows须在宽有符号域计算并夹紧;Sq>Sk时uint32差值与后续减法可回绕,让本应空可见区错误变成全KV或巨长区。
28+ 
29+来源noMaskKvS/noSkipKvS正确区间应是[0,noMask)全可见候选、[noMask,noSkip)需mask、[noSkip,Sk)跳过,原文把最后区间说成对角块有误。全可见前缀可保守收缩,但不能漏有效key;masked宽度也不总等于qRows,尤其早期全/部分空行。totalStack=ceil(noMask/512)+1只有确有一个masked stack时成立。
30+ 
31+noMaskTailS=prefix%128、maskedStart回退页、noMaskTailInteStackNum及延迟preLaunchStackNum的变步长是特定压缩布局索引。整数步长0可能需要靠stackSeqCount改变下一分支,必须证明整体循环进展和数据覆盖,不能只看kvIdx。Tail QK/PV多两个尾参、最后PVTail选择也须与这些枚举任务一一对应。
32+ 
33+NO_MASK外层显式wait qkReady,masked epilogue可能内部wait;若外层也wait同一token会双消费,需核验实际签名和所有权。1024² half固定模板通常2MiB GM,不是任意大mask可直接无压缩复制到UB;1表示屏蔽实际上构造的是上三角指示,source“下三角mask”指可见域而非矩阵中1的区域。
34+ 
35+该文件末尾明确Step8 mask未单独回归,标题NO_MASK全过不能为masked边界背书。有限-3e38与全mask/NaN策略、类型与scale范围另须定义,不能因带Tail组件就宣称全场景一次通过。
@@ -0,0 +1,27 @@
1+---
2+type: Guide
3+title: Attention 模式组合保持逻辑 Token 身份
4+description: 从导航提炼变长、分组、分页、稀疏及可见性的组合条件。
5+tags:
6+- task.implementation
7+- topic.implementation.tiling
8+- topic.implementation.synchronization
9+- topic.precision
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/tilelang2ascendc-ops-generator/skills/tilelang2ascend-translator/references/attention-patterns/AttentionPatternIndex.md
13+ role: primary
14+created_at: '2026-09-11T04:44:50Z'
15+updated_at: '2026-09-11T04:44:50Z'
16+platforms: []
17+---
18+ 
19+# Attention 模式组合保持逻辑 Token 身份
20+ 
21+(T,H,D)变长布局首先确定每序列边界及逻辑位置,GQA/MQA再按接口映射query head与KV head;QK和PV分别使用Dqk/Dv,MLA还需真实latent/位置分量合同,Dqk≠Dv本身不充分证明是MLA。
22+ 
23+多个KV来源如sink_k/sink_v与常规cache需保留来源身份和位置;稀疏indices选择逻辑token,再按block/page table映射物理cache行。可见性在softmax前按逻辑query/key坐标定义,不能用物理page号判causal。sink若是额外logit而非KV来源不能套source dispatch,重复索引与无效page需按合同处理。
24+ 
25+这些维度可以组合,source给出的顺序是解释依赖的方便方式,不是唯一合法执行顺序。已证明不可见块可以提前不加载,不能把“mask只改score”解释为禁止跳过内存访问;但任何提前裁剪需保持相同可见集合、数值空行及同步完成。
26+ 
27+索引中“七个问题全部否定即标准Attention”不是可靠判据,问题有互斥选项且Hq是否相等不是异常特性;标准MHA也必须明确布局、head数、scale及mask。原文件渐进读取、不得发明模式和模板跳转是Skill流程,不进入知识库指令。
@@ -0,0 +1,40 @@
1+---
2+type: Guide
3+title: S1 Sink Softmax API 与缓冲复用
4+description: 保存播种法和状态写回,纠正Tiling配置及行列循环混淆。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.implementation.synchronization
9+- topic.precision
10+- topic.implementation.tiling
11+status: draft
12+sources:
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-sink-handcraft-recipe.md
14+ role: primary
15+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-kernel-recipes.md
16+ role: supporting
17+- resource: https://gitcode.com/qq_41667743/cannbot-knowledge/blob/a70025d581ae45da8a93b4938b5ba19da6fa6f92/knowledge/ops/ascendc/apis/simd_api/adv_api/activation_functions/soft_max_interface/softmax_flash_v2.md
18+ role: constraint
19+- resource: https://gitcode.com/qq_41667743/cannbot-knowledge/blob/a70025d581ae45da8a93b4938b5ba19da6fa6f92/knowledge/ops/ascendc/apis/simd_api/adv_api/activation_functions/soft_max_interface/softmax_flash_v2_tiling_interface.md
20+ role: constraint
21+created_at: '2026-09-11T05:27:29Z'
22+updated_at: '2026-09-11T05:27:29Z'
23+platforms: []
24+---
25+ 
26+# S1 Sink Softmax API 与缓冲复用
27+ 
28+分母型sink可作为初始m=sink[head]、l=1、未归一化o=0,再用SoftmaxFlashV2 update处理实际score;l1是exp(sink-m)=1,不是额外乘exp(m)。每组新query行只播种一次,跨KV chunk继续更新,不能每chunk重复加入sink。无sink可首轮使用非update或正确空状态,不能令m取任意sink占位值且忽略全mask行。
29+ 
30+来源调用支持src/dst、max/inMax、sum/inSum同址的候选,但不等于所有tensor/scratch都可互相别名。已有API卡约束isReuseSource为预留默认false,状态末轴通常32B;source启用isReuse=true并把expMax按rows*64分配不能推成通用必需布局。应按实际API/config的广播模式和最小scratch查询确定。
31+ 
32+来源反复强调Tiling isUpdate=false而Kernel模板true;与已有Tiling API卡要求isUpdate和isBasicBlock两侧一致冲突。不能将此经验当正确固定配方,本批未读外部归档实现,应使用相匹配的接口配置或取得该旧版本特例证据。ShapeInfo需同时表达逻辑shape、物理pitch和有效size,不是任意二维声明即可修复数值。
33+ 
34+ProcessVec1文字称realSplitN=ceil(s2Real/列预算),代码却以loopIdx*vecS1Base切query行,并计算s1Real-loopIdx*vecS1Base;这实际要求按行分批的loop数。若真按列拆,则需同一query的online更新及alpha接续,不能沿用这些行偏移。vecS1Tail、S读址、P写址及max/sum槽必须同一解释。
35+ 
36+maskPse分别加载且遵循PSE外乘/内加合同,scale与mask先后不能任意换。尾列显式mask,compressed-mask快路径也要确保所有padding列已无贡献。hasSink时maskPing被复用为exp输出,需在mask最后消费者结束后才能覆写;相应V_MTE2释放移后是该共享生命周期的一部分。
37+ 
38+来源UB命名区11+11+16+32+32+16+8+8+8+32=174KiB,仍需实际分配对齐及其他框架区。max/sum各8KiB攒多批行后写GM,copyOutLimit必须正且不越容量;Vec2从相同batch/head/query偏移读回l,不能依赖已被复用的UB旧状态。
39+ 
40+exp输出中alpha位置应按真实合同读取,若自算exp(mOld-mNew),旧max在update覆写前保存并保留至延迟消费者。P降精度而sum来自FP32指数是一种近似路径,不能据sum相同就证明O bitwise一致。来源“手写softmax慢430倍”仅为历史失败实现,不是API以外算法均不可用。
@@ -0,0 +1,34 @@
1+---
2+type: Guide
3+title: S1 三槽流水与排空身份
4+description: 修正模3时间标签和伪任务索引,保留L1Reuse配对。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.implementation.synchronization
9+- topic.precision
10+- topic.implementation.tiling
11+status: draft
12+sources:
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-sink-handcraft-recipe.md
14+ role: primary
15+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-kernel-recipes.md
16+ role: supporting
17+created_at: '2026-09-11T05:27:29Z'
18+updated_at: '2026-09-11T05:27:29Z'
19+platforms: []
20+---
21+ 
22+# S1 三槽流水与排空身份
23+ 
24+S1选定路径每task覆盖一块query并一次处理该任务全部有效KV;s1s2则有KV内层。三槽extraInfo为调度元数据,S/P/O两槽为数据存储,两种深度不能混用。若当前taskId=t,(t+2)%3对应t-1,(t+1)%3对应t-2;来源将它们分别叫task-2和task-1颠倒。
25+ 
26+按所列代码,提交QK(t),消费QK(t-1)做softmax并提交PV(t-1),等待/消费PV(t-2)做Vec2。每次复用同一Matmul对象、GM槽或metadata前应满足对应完成条件;只能保存三条extraInfo不证明允许无限未完成Matmul请求。WaitBmm1/2、End及异步模式需按目标API合同,而非仅以序号解释。
27+ 
28+非L1Reuse多2个空迭代用于排空生产链;这些伪task只能消费已有工作,不得先ComputeAxisIdx(idx)和GetS2LoopRange读取超过totalSize的batch/head/长度。空核、单task、尾task以及奇数task都需独立推演,不能因notLast抑制了Iterate就认为前面的元数据访问安全。
29+ 
30+L1Reuse来源以blockIdx/2给一对核分共同段、奇偶各自处理对应任务,tiling splitFactor和limit额外排空/补偶要整体匹配。单独乘2会重复或遗漏,单独开模板也会错;但D64–128不是所有FA必须开L1Reuse的数学条件。S1示例禁用、s1s2候选启用是其策略差异。
31+ 
32+连续任务段按ceil(total/core)切可能有闲核;稀疏sparseStartIdx表按估算工作量分配,不可只看任务数。QK/softmax/PV在不同硬件单元执行,即使同一global入口也不表示AIC/AIV变为同一物理核或只需本核事件,框架消息协议仍存在。
33+ 
34+来源s1s2章节称mm1Res和stage1Res全在UB,S1具体GlobalTensor和workspace公式却明确S/P落GM。应按实际实例的地址空间选择资源和同步;不能以“单核融合”标题证明无GM往返,也不能凭该标题归因2–3倍收益。
@@ -0,0 +1,40 @@
1+---
2+type: Guide
3+title: S1 Workspace 单位与 Tiling 候选
4+description: 保存三段双缓冲公式,限定资源覆盖和缓存Shape条件。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.implementation.synchronization
9+- topic.precision
10+- topic.implementation.tiling
11+status: draft
12+sources:
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-sink-handcraft-recipe.md
14+ role: primary
15+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-kernel-recipes.md
16+ role: supporting
17+- resource: https://gitcode.com/qq_41667743/cannbot-knowledge/blob/a70025d581ae45da8a93b4938b5ba19da6fa6f92/knowledge/ops/ascendc/apis/simd_api/adv_api/cube_compute/matmul_kernel/set_org_shape.md
18+ role: constraint
19+- resource: https://gitcode.com/qq_41667743/cannbot-knowledge/blob/a70025d581ae45da8a93b4938b5ba19da6fa6f92/knowledge/ops/ascendc/apis/simd_api/adv_api/cube_compute/matmul_kernel/iterate_all.md
20+ role: constraint
21+created_at: '2026-09-11T05:27:29Z'
22+updated_at: '2026-09-11T05:27:29Z'
23+platforms: []
24+---
25+ 
26+# S1 Workspace 单位与 Tiling 候选
27+ 
28+令aN=alignUp(S2,16)、A=alignUp(s1Base*aN,256)*sizeof(T)、C=alignUp(s1Base*alignUp(D,16),256)*sizeof(T)。来源两份mm1Res占2A,两份P总占Bstage,两份mm2Res占2C,再加按512B对齐的PSE区;每核基址core*(2A+Bstage+2C+PSE)。这里256是元素对齐再乘dtype,不是简单256B。
29+ 
30+T=FP32且INPUT_T=FP16时Bstage=A,两份P每份A/2;INPUT_T=FP32时来源Bstage=A*sizeof(INPUT_T)/2=2A,两份P各A。该等式依T/INPUT_T具体选择,不能推广任意类型。dropout bit mask若占workspace前缀,应先跳其512B对齐容量,再按核分区;userWorkspace与系统prefix也需分开。
31+ 
32+commonTBuf32KiB若用[0]/[8192]作为两个槽,8192的元素宽度必须明确:half偏移16KiB可分半,float偏移32KiB已到末端。GM字节指针与GlobalTensor元素索引同样不能混用。所有mm1/mm2 SetTail不改变实际分配容量,不得超singleCore输出空间。
33+ 
34+来源s1BB=min(128,alignUp(S1,16))、s2BB及dBB类似;nRatio默认4,alignedS2<128候选1,NZ/ND特殊条件含S2%64、S264、S2≥704且B*H*S1>20480、D≤256,均为特定策略启发式。s1Ratio grow/shrink试图平衡并行与workspace,容量应使用实际alignedS2而非固定s2BB;S21024与s2BB128可差8倍。
35+ 
36+s1Base=s1BB*ratio、s1Outer=ceil(S1/s1Base)、total=B*H*s1Outer、core=min(total,aivNum)、split=ceil(total/core),0任务/0资源先处理。L0C计算fixsplitN=floorAligned(L0C/(s1BB*4),16)仍需真实dtype、双缓冲和其他占用;GetTiling后强行shareL1=1MiB/shareL0C256KiB/depth2不能当合法覆盖,目标资源不足会失效。
37+ 
38+SetOrgShape表达原始矩阵完整尺寸及A/B各自K维,不是通用任意stride setter;来源把第4参orgKb解释成V行stride D、第5参叫N2过度简化。需按实际布局/转置和API解释,保存真实leading dimension。只在s1Real变化时重设Shape的缓存还须保证N/K/stride/布局未变;varlen batch同Q长不同KV长会破坏单字段缓存。
39+ 
40+d%16≠0时该实例采用NZ及NdToNz/NzToNd完整链,其他合法padding/布局实现仍可存在。选Skv≤1024的S1或小BNG*S1*S2容量B模板需实际resource/feature评估,1024不是所有平台最优硬阈值。
@@ -0,0 +1,39 @@
1+---
2+type: Guide
3+title: Sage 缓冲布局与容量冲突检查
4+description: 保存GM和片上地址合同,揭示源布局表中的重叠与扩容缺口。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.precision
9+- topic.implementation.synchronization
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-sageattention-architecture.md
13+ role: primary
14+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-sageattention-rules.md
15+ role: supporting
16+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-sageattention-troubleshooting.md
17+ role: supporting
18+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-sageattention-recipe.md
19+ role: supporting
20+created_at: '2026-09-11T05:16:48Z'
21+updated_at: '2026-09-11T05:16:48Z'
22+platforms: []
23+---
24+ 
25+# Sage 缓冲布局与容量冲突检查
26+ 
27+量化Qhat为BSND[B*sqPad,H*D] INT8,Khat为[B*kvBlocks*128,Hkv*D] INT8;V为[B*skPad,Hkv*D] FP16,O写回HND。Q加载GQA下GM行间跨其他head,来源srcStride=(H-1)*D对INT8可同时按元素/字节得到同数,换dtype需重新检查API单位。K索引同时含batch、KV head和逻辑块,不能只换head数量。
28+ 
29+SQ/SK以[blocks,8]FP32复制,索引必须block*8;mu为[B*Hkv*kvBlocks,D*8]FP32,来源读取每块首D值、写8份向量,不能擅自解释成每个元素交错复制8次。复制是该实现的物理ABI,不是量化数学要求。po[totalWork,128,128]FP32保留D有效列,pml[totalWork,512]FP32含sub0在0的m64/l64和sub1在256的m64/l64及padding。
30+ 
31+来源dav-3510示例资源为L0A/B各64KiB、L0C256KiB、L1512KiB、每AIV UB248KiB,核数例28AIC/56AIV但需实际查询;这些值未在本批设备核验。FAIQK静态Shape读(M,N,K),FAIPV读(M,K,N),需按真实版本确认。QK INT8 <128,128,256>每A/B槽32KiB,与PV HALF <128,128,128>每槽32KiB匹配;运行D128不等于静态K也必须128。共用从0起的地址与事件时,槽覆盖范围和生命周期必须兼容,不同槽距并非所有独立设计都禁止。
32+ 
33+来源L0C QK/PV各两份64KiB共256KiB,已占满示例容量;cursor顺序分配不自动防溢出。L1四份P各32KiB共128KiB,加Q/K各两份32KiB及V两份32KiB,总320KiB,另计实际组件其他区。
34+ 
35+UB表S两份共64KiB,O两份64KiB,oacc32KiB,O16 16KiB,P两份32KiB,累计208KiB。QF从208KiB开始却列INT8 8KiB+half8KiB+FP32 16KiB=32KiB,跨度到240KiB;mu又从232KiB开始,状态233KiB、scratch236KiB,按表同时存活会重叠。来源声称合计约240KiB需由实际chunk分配/合法别名证明,不能直接照表当无冲突布局。
36+ 
37+加深S/O从2到4仅这两者就额外128KiB,若P也加深再加32KiB为160KiB;原240KiB即使删32KiB QF也不能放进248KiB。KV tile加到256虽然可将QK单份L0C128KiB+PV128KiB装入256KiB,却也扩大S/P和量化块合同,不能只核算L0C就断言可实现。Fixpipe DEQF16省流量候选还会改变int32 score的数值范围与scale路径。
38+ 
39+D64时来源PV向UB写紧密[64,64],RS不能固定行距128;po GM则仍可能以128为物理行距。Cast与输出长度按有效SH*D,合法1D连续写可避免某2D配置错误,但不是所有128B blockLen搬运都非法。
@@ -0,0 +1,31 @@
1+---
2+type: Guide
3+title: Sage 编译 Helper 与 API 适用范围
4+description: 保留地址空间、寄存器和向量辅助实现的条件。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.precision
9+- topic.implementation.synchronization
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-sageattention-troubleshooting.md
13+ role: primary
14+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-sageattention-architecture.md
15+ role: supporting
16+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-sageattention-recipe.md
17+ role: supporting
18+created_at: '2026-09-11T05:16:48Z'
19+updated_at: '2026-09-11T05:16:48Z'
20+platforms: []
21+---
22+ 
23+# Sage 编译 Helper 与 API 适用范围
24+ 
25+来源dav-3510某count API用tensor[off]生成LocalTensor切片,tensor+off未匹配目标重载;StoreUnAlign地址参数可能要求可修改左值,可先声明__ubuf__指针再调用。不能从一次报错断言所有LocalTensor加法或RegTensor赋值均不支持;RegTensor t0=f0失败可通过真实copy原语或合法分支组织解决,需具体寄存器API和编译器版本。
26+ 
27+__simd_vf__可接收的地址空间与普通Device helper不同,GM指针不能任意传入只接受UB的接口。CATLASS_DEVICE普通函数也需遵守设备语言支持。Host aclrtMalloc得到void*、入口GM_ADDR为__gm__uint8_t*时需合法Host桥接和显式转换,来源uint8_t*加(void**)&是其工程形式,不能因此抹掉Device地址空间限制。
28+ 
29+某device Scalar float除法无法lower时,SageDivScr用三个8FP32 scratch区、Duplicate、Vector Div、完成后GetValue作为候选。8元素和禁止原地属于此helper的保守配置,实际API count/别名限制需看重载;不能扩展成全部Ascend C标量除法都不可用。除数0替换1只适合已定义的零输入分支,不是通用除法语义。
30+ 
31+CATLASS_ARCH未设可能使默认2201分支跳过3510的ScaleGranularity/FP8声明,从而产生连锁模板报错;应核对编译宏、目标npu架构和实际include版本,不是遇到任何ScaleGranularity错误都添加同一宏。源码reference中的“最新CATLASS”没有固定依赖commit,不能当可重现版本标识。
@@ -0,0 +1,37 @@
1+---
2+type: Guide
3+title: Sage 游标 Split-KV 与归并边界
4+description: 保存任务首步、完整状态槽与因果位置合同。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.precision
9+- topic.implementation.synchronization
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-sageattention-architecture.md
13+ role: primary
14+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-sageattention-troubleshooting.md
15+ role: supporting
16+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-sageattention-precision.md
17+ role: supporting
18+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-sageattention-recipe.md
19+ role: supporting
20+created_at: '2026-09-11T05:16:48Z'
21+updated_at: '2026-09-11T05:16:48Z'
22+platforms: []
23+---
24+ 
25+# Sage 游标 Split-KV 与归并边界
26+ 
27+work g对应task=g//SK和shard=g%SK,jStart=floor(jEndF*shard/SK)、jEnd=floor(jEndF*(shard+1)/SK)。区间连续且大小差≤1,但jEndF<SK会有空片;来源“无空片”与后文空片紧凑化矛盾。Host cap splitK到全局kvBlocks仍不能消除causal早期query只见少量块的空片。totalTasks0也不能先算ceil(aicNum/totalTasks)。
28+ 
29+Cursor Next先保存jCur再j++,返回后使用jCur;isFirst为jCur==jStart,不能统一jCur==0。首步状态重新初始化与全局step槽轮转不同:step s读(s-1)%N的同一work旧状态,写s%N,新work首步不可读取前一query的残留。Q借用L1仅在同一work真实Q不变且最后消费者结束后可轮转。
30+ 
31+每步mNew=max(mOld,max(S)),alpha=exp(mOld-mNew),l=alpha*l+sum(exp(S-mNew)),oacc=alpha*oacc+P*V;max、sum、alpha保留到对应延迟RS消费者。分区写未归一化oacc及m/l,归并mStar=max(mk)、wk=exp(mk-mStar)、lStar=sum(wk*lk)、O=sum(wk*ok)/lStar。MulDstAdd不能凭名字当acc+=a*b;按目标API确认其乘的是dst还是两个source。
32+ 
33+pml紧凑化需要保存两AIV所有有效m/l;来源整slot512float含sub1偏移256,copy128只搬sub0会丢半行。可以copy完整512或分别搬两段有效值,后者须保持正确同步/stride。k==nValid时可免copy但仍要计数,不能continue跳过计数。若po没同步紧凑化,必须保留有效原shard编号,不能拿移动后的m/l配错po。
34+ 
35+右下causal候选key≤query+(sk-sq),属于该接口而非普遍默认。块上界由最后有效query行和kvBlocks夹紧,负上界、Sq>Sk及空可见行要处理,不能依C++负整数截断自动得到数学floor。阈值lim-1与列索引比较方向不能反;padding后的零K会产生score0,必须显式mask使其不进入分母。
36+ 
37+分区全空、整行全mask及非有限score不能无条件计算-inf-(-inf);输出与LSE按接口特殊值策略。GQA head映射需H可被Hkv整除且正确映到各batch的KV头,不能只凭8:1等已测比例宣称通用支持。
@@ -0,0 +1,39 @@
1+---
2+type: Guide
3+title: Sage 分层取证与证据边界
4+description: 保留拟合和中间量探针,避免余弦相似度与症状唯一归因。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.precision
9+- topic.implementation.synchronization
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-sageattention-precision.md
13+ role: primary
14+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-sageattention-troubleshooting.md
15+ role: supporting
16+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-sageattention-rules.md
17+ role: supporting
18+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-sageattention-recipe.md
19+ role: supporting
20+created_at: '2026-09-11T05:16:48Z'
21+updated_at: '2026-09-11T05:16:48Z'
22+platforms: []
23+---
24+ 
25+# Sage 分层取证与证据边界
26+ 
27+可分原始输入独立FP32 golden、表达量化/补偿/online/split语义的emulator、设备实现三层比较。但cosine0.9999不证明设计普遍正确,kernel与emu差也未必都是kernel bug,emu和dump仍可能错。若emu直接使用Kernel A产出的Qhat/Khat/mu/scale,它只能局部检查Kernel B,必须另用原始输入独立验证量化产物。
28+ 
29+余弦不检测纯正倍缩放:O和2O的cosine为1;还需目标合同要求的绝对/相对误差、特殊值、shape和必要逐阶段统计。来源22case、cos0.999592–0.999909、128²约31us/1k²约73us/4k causal约505us/200k约719ms为历史声称,未在本批重跑;不同文段还给不同区间,不能合并为新的实测结论。只抽样max-planes防golden OOM时应明示未检查平面范围。
30+ 
31+比较Qhat*Khat整数S、comp逐行、m/l/alpha每step、P*V及最终merge可缩小故障位置。dScale拟合轮流代入各块scale,merge最小二乘拟合po列组合,可发现候选错序/错系数;矩阵秩不足、多个假设相同或误差可拟合时不能唯一归因。cols0–63、rows32–63或随块数恶化是线索,不是某个L0/子核/scale bug的证明。
32+ 
33+来源512KiB qdbg将S保留区、P(+128KiB)、O(+144/+336KiB)、状态400–416KiB、comp416–424KiB、AIC进度250KiB、AIV452–460KiB、merge464–496KiB用于不同探针。多个范围可能复用或重叠,P还有双sub竞写;只能在明确相位、slot、step、写者及范围内解读,不能当无冲突通用dump布局。长序列step快照必须限制容量或轮转并存step身份。
34+ 
35+相位alive/smDone/rsWait/rsDone和迭代号可定位最后进展,AIC/AIV分别记录;Scalar GM写及小DataCopyPad不落盘在来源为特定环境观测,“32B必失败/≥1KiB必可靠/AIV标量store全失败”没有普遍依据。扩大拷贝可能改变时序,仍需合法长度、双方同步、Host回读完成;不能把8倍复制当API全局要求。
36+ 
37+srun返回码、输出存在性与mtime能排除许多旧文件假通过,但秒级mtime、并发同目录和旧binary仍可漏判。失败/timeout产物可留作取证,不能当成功运行精度证据;每次运行以唯一身份、制品和完整输出对应更可靠。npu-smi Health正常不证明所有执行状态无异常,“挂两次换卡”是来源运行策略,不迁移为知识规则。
38+ 
39+进程总时长需拆setDevice、内核、同步、reset/finalize。来源1.8s/4.6s初始化销毁与约4ms内核只说明那次测量,持久进程能摊薄而非保证全消失。header未重编、JSON true被atoll读0、批量replace未命中或括号破坏也可导致旧制品/错误配置,应以真实编译和输入解析证据区别设备故障。
@@ -0,0 +1,37 @@
1+---
2+type: Guide
3+title: Sage 四阶段流水与版本事件所有权
4+description: 保留mode4槽位协议,限定Bootstrap与双重等待结论。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.precision
9+- topic.implementation.synchronization
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-sageattention-architecture.md
13+ role: primary
14+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-sageattention-rules.md
15+ role: supporting
16+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-sageattention-recipe.md
17+ role: supporting
18+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-sageattention-troubleshooting.md
19+ role: supporting
20+created_at: '2026-09-11T05:16:48Z'
21+updated_at: '2026-09-11T05:16:48Z'
22+platforms: []
23+---
24+ 
25+# Sage 四阶段流水与版本事件所有权
26+ 
27+来源混合launch为__mix__(1,2),每AIC配两AIV且各处理64行;AIV逻辑AIC编号GetBlockIdx()>>1、子块GetSubBlockIdx()仅在该启动合同下解释。独立__vector__ merge不能依赖混合子块索引,候选sub=blockIdx&1、task=blockIdx>>1、步长blockNum>>1要求成对核数且步长非0。
28+ 
29+四相位流水按it推进QK(it)、SM(it-1)、PV(it-2)、RS(it-3),循环到P+3排空;单步数据依赖QK→SM→PV→RS不允许违反,但不同step可重叠。QK/S、PV/O各2槽,P到L1用4槽。来源mode4候选C1_V1={0,1}、V1_C2={2,3,4,11}、C2_V2={5,6}、MM2={7,8}、MM1={9,10}、PV_CONSUMED={12,13,14,15};AIC通知双AIV时set id及16+id,AIV set单id由该模式路由,AIC等双方id。此表仅适用对应组件/模式,不能挪到A2或其他事件分配。
30+ 
31+MM1表示AIV已消费sUb,MM2表示已消费oUb;每步只释放本步槽,两个槽都set会使未来等待提前通过。P在UB经MTE3送L1后,下一Cast覆写pUb需等MTE3完成;AIC读完pL1后还需PV_CONSUMED,两个缓冲的保护不能互相替代。首次填充可以跳过不存在消费者的等待,复用才wait,按槽真实使用次数判定,不能不分任务边界机械用it>4。
32+ 
33+来源2026-08所谓“新版”BlockMmadTla在内部wait MM1/MM2;若外层再wait同一token会死锁。需读取实际安装组件确定等待所有权,不能一概删除所有版本外层wait。来源冷PIPE_FIX预置不传播和AIC PIPE_ALL必挂死未附本批底层验证;保留为指定环境故障线索。通用要求是核型支持的pipe集合、实际producer完成及bootstrap/steady/drain计数正确,不能推广成所有AIC永远禁止任何Barrier。
34+ 
35+Set/Wait<V_MTE3>表达V生产后MTE3可消费,不是DMA完成后V可覆写;后者需要MTE3→V或实际等价完成协议。来源“kernel退出不保证MTE3排空”不能脱离runtime/接口强制所有出口插PIPE_ALL;需要确认已提交写出与下个kernel/Host读的完成合同。跨launch残留token是单发对、双发错的一个候选,不是唯一根因。
36+ 
37+新增tiling字段要保持Host序列化与Device逐字段拷贝一致,漏字段会未初始化,P0只是可能表现。CMake依赖追踪是否缺失依ASC工具链,给add_executable列头文件或touch TU不是所有版本根本修复;应验证真正重编及加载制品指纹。
@@ -0,0 +1,35 @@
1+---
2+type: Guide
3+title: Sage 量化与 K 平滑补偿合同
4+description: 保存INT8 QK与FP16 PV路径,修正精确性和类型转换推断。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.precision
9+- topic.implementation.synchronization
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-sageattention-architecture.md
13+ role: primary
14+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-sageattention-recipe.md
15+ role: supporting
16+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-sageattention-precision.md
17+ role: supporting
18+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-sageattention-troubleshooting.md
19+ role: supporting
20+created_at: '2026-09-11T05:16:48Z'
21+updated_at: '2026-09-11T05:16:48Z'
22+platforms: []
23+---
24+ 
25+# Sage 量化与 K 平滑补偿合同
26+ 
27+来源v1为HND输入Q[B,H,Sq,D]、K/V[B,Hkv,Sk,D],FP16输出;每128有效K行求mu_j,令Ks=K-mu,sK=127/max(amax|Ks|,amax|mu|),Khat=round(Ks*sK)转INT8;Q块用sQ=127/amax|Q|和Qhat=round(Q*sQ)。K尾块均值的分母是nValid而非128,padding不应污染均值/scale;树归约破坏原K后需从保存的half重建再减mu。
28+ 
29+重建score为S_int*sm/(sQ*sK_j)+comp_j,S_int=Qhat*Khat^T,comp_j=(Qhat*mu_j)*sm/sQ。sK只在主项中,补偿里已消去;若写Q而非Qhat,则应使用Q的实际近似值而不是额外再除sQ。source recipe一处dot(Q,mu)*(sm/sQ)与其余定义不一致,不能同时采用。mu是FP32,Qhat*mu的dot不是整数dot,不能据结果小于2M就称FP32精确或与另一增广GEMM逐位相同。
30+ 
31+对有限有效score按online max/sum更新,P在[0,1]后降FP16,由FP16 Cube PV产生FP32未归一化输出,最终除sum再转FP16。量化与P降精度都有误差,不是原始全精度attention的代数完全等价。all-zero块的amax0需约定scale和量化零;来源helper把除数0换1是一种实现选择,NaN/Inf与饱和/舍入仍需显式合同。
32+ 
33+来源两级Cast K/Q经FP32→FP16→INT8,RINT时点及饱和必须与目标支持一致,不能凭函数名断言无双重舍入。bf16位模式不能当half读取;先数值转half可修正位解释,但会引入新的舍入和FP16范围限制,不等于完整保留BF16接口值域。生成BF16用FP32高16位截断也不同于最近偶数转换,生成器需明确实际舍入。
34+ 
35+quantScheme=1的FP8 MX PV在来源未完成。MmadFAIPVMx可能要求L1已是指定FP8类型及P/V scale,不能把half缓存交给DataCopy期待隐式量化;FP8格式、scale粒度、布局及转换路径需由实际组件确认。LSE、varlen、NHD和D96是来源v1未补齐的能力,22个历史case通过不能覆盖这些缺口。
@@ -0,0 +1,26 @@
1+---
2+type: Guide
3+title: Attention 打分维度与 Value 输出维度
4+description: 在 MLA 等分段输入路径中分别管理 QK 打分维度、Value 宽度及 latent 数据消费范围。
5+tags:
6+- task.design
7+- topic.compute.composite
8+- topic.data_movement.layout_transform
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/tilelang2ascendc-ops-generator/skills/tilelang2ascend-translator/references/attention-patterns/references/mla_pattern.md
12+ role: primary
13+created_at: '2026-09-10T14:32:33Z'
14+updated_at: '2026-09-10T14:32:33Z'
15+platforms: []
16+---
17+ 
18+# Attention 打分维度与 Value 输出维度
19+ 
20+Attention 的 QK 归约宽度 Dqk 与 PV 输出宽度 Dv 可以不同。Kernel 的 Q/K 搬运、score 计算、V 搬运、输出跨度及资源预算应分别使用对应维度,不能用一个 headDim 替代全部含义。
21+ 
22+在合同将 Q/K 分为 nope 与 rope 两段时,score 的两段点积先合成同一 S,再施加所需 scale、mask 和 softmax。不能对两段独立归一化。默认 scale 若按 Dqk 定义,就不能误用 Dv;显式 scale、位置编码和模型特有缩放仍以实际接口为准。
23+ 
24+某些已经完成投影或吸收变换的 latent cache 布局允许 K 路读完整 Dqk、V 路仅读同一行前 Dv 列。这是特定表示的消费合同,不能推广到所有 MLA。`Dqk!=Dv` 本身也不能识别全部 MLA,更不能证明 weight absorption 已完成。
25+ 
26+若缓存采用分页或离散选择,K/V 必须指向同一逻辑 token 的相应数据段。即使 workspace 为 Dqk 宽,PV 和最终写回也只覆盖约定 Dv;物理 pitch 可包含额外补齐。验证使用两段独立可识别数据、Dqk/Dv 不同取值、非零 scale 及分页尾块,分别核对合成 score 和输出。
@@ -0,0 +1,30 @@
1+---
2+type: Guide
3+title: Attention 的 score 可见性与额外 KV 来源
4+description: 区分 causal/显式 mask 与额外 sink KV,保证有效列集合和 online softmax 的统一归一化。
5+tags:
6+- task.design
7+- topic.compute.reduction
8+- topic.compute.composite
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/tilelang2ascendc-ops-generator/skills/tilelang2ascend-translator/references/attention-patterns/references/mask_causal_pattern.md
12+ role: primary
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/tilelang2ascendc-ops-generator/skills/tilelang2ascend-translator/references/attention-patterns/references/sink_pattern.md
14+ role: supporting
15+created_at: '2026-09-10T14:32:33Z'
16+updated_at: '2026-09-10T14:32:33Z'
17+platforms: []
18+---
19+ 
20+# Attention 的 score 可见性与额外 KV 来源
21+ 
22+Mask 约束 score 的可见列。无效 K/V 的零填充不能替代 softmax 前的 mask,因为零点积仍可能得到非零概率。显式 mask 的极性、dtype、广播规则及加性或布尔含义由接口规定,不能固定为某条 uint8→half→float 转换链。
23+ 
24+Causal 判断应使用同一个 batch 内的逻辑 Q/K 位置,不能使用物理 page 顺序或全局拼接后的 T 行号。左上对齐与右下对齐是不同合同;后者常写为 `k<=q+(Lkv-Lq)`,该偏移不是所有 causal 接口的默认值。整 tile 跳过也必须对所有有效 query 行成立,并保持流水参与者与通知配对。
25+ 
26+当接口额外提供独立 sink_k/sink_v 时,它们是另一组 KV 来源。若数学合同规定它们与 local KV 一起参加 attention,sink 与 local 的 score 必须在同一归一化域中合并;不能将两次独立 softmax 的输出直接相加。可沿一套 online 状态逐块处理,或按等价的分区状态归并。
27+ 
28+本卡的 sink 专指额外 KV 张量;其他接口可能把同名字段定义为只进入分母的 logit,需要另读接口合同。各来源有各自 layout、有效长度和 tail;独立连续 sink 不必复用 local 的 page table。K 的来源选择还必须延续到对应 V 的加载,不能只在 QK 阶段接入新来源。
29+ 
30+验证包括全掩码行、零有效 key、长度不等的 causal、sink/local 不同尾块以及改变 KV 分块后的结果。全掩码状态不能直接套有限 score 更新而产生 `-inf-(-inf)`。状态与归并公式见 [online softmax](online_softmax_state_and_partition_merge.md);物理地址见[逻辑 token 与缓存地址](logical_tokens_heads_and_cache_addresses.md)。
@@ -0,0 +1,30 @@
1+---
2+type: Guide
3+title: Sink 任务与跳块后的流水状态
4+description: 补齐 sink 与 KV 的集合关系、任务字段及 online 状态跨任务归并。
5+tags:
6+- task.design
7+- topic.attention
8+- topic.implementation.synchronization
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/playbooks/P80_sink_task_pipeline_injection.md
12+ role: primary
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/preconditions/P80.yaml
14+ role: supporting
15+created_at: '2026-09-11T02:13:08Z'
16+updated_at: '2026-09-11T02:13:08Z'
17+platforms: []
18+---
19+ 
20+# Sink 任务与跳块后的流水状态
21+ 
22+Sink KV可作为独立任务进入与普通稀疏KV相同的计算流水,为每个query/S1G块贡献相应Attention状态。isSinkTensor决定输入来自keySink/valueSink还是常规cache,但每个任务还需完整初始化query/head/batch、S2范围、有效长度、槽位及mask/位置字段;来源仅设置少数字段的局部TaskInfo不足以作为运行模板。
23+ 
24+普通循环从sinkNumber开始只在sink与普通序列的逻辑定义确实对应“前缀已由sink处理”时成立。独立sink Tensor若是额外集合,直接跳过cache前sinkNumber个token会漏项;若cache本已包含同一sink而又重复加入,则可能重复贡献。keySinkNumber与sinkNumber也需有明确单位与关系,K/V数量、D及地址不能混用。
25+ 
26+Sink和每个有效稀疏task必须更新同一query的online max、sum和未归一化加权和,最后统一归一化;不能分别归一化后直接相加。注入新task会影响首轮初始化、最终输出task识别、workspace槽轮转和排空长度。sinkNumber超过单task容量时应切成多个合法任务,并非每个query始终恰好增加一项。
27+ 
28+IsSkipCal只能跳过全无效的S2区,部分有效区需保留正确mask。跳块后按实际已发布任务推进流水,避免对未生产的槽等待或把“最后枚举块”误当最后有效任务。空稀疏但有sink、无sink、全mask行均须定义结果与同步行为。
29+ 
30+来源sInnerSize=256是效率候选;缩小块可能增加调度和减少无效计算,也可能降低有效块计算效率。既有sink存储可直接复用,不必因改变调度再分配一份GM;策略适用性取决于实际sink语义,不只取决于Streaming LLM标签。验证集合覆盖和逐query状态,再比较含任务增加的总时间。
@@ -0,0 +1,31 @@
1+---
2+type: Guide
3+title: Softmax 反向行点积的预计算
4+description: 保存梯度公式、每 query 行结果容量及 dY/O 版本一致性。
5+tags:
6+- task.optimization
7+- topic.attention
8+- topic.compute.reduction
9+- topic.precision
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/playbooks/P91_softmax_grad_precompute.md
13+ role: primary
14+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/preconditions/P91.yaml
15+ role: supporting
16+created_at: '2026-09-11T02:13:08Z'
17+updated_at: '2026-09-11T02:13:08Z'
18+platforms: []
19+---
20+ 
21+# Softmax 反向行点积的预计算
22+ 
23+对于相应Attention反向公式,行统计delta[q]=sum_d(dY[q,d]*O[q,d])不随该query的KV/S2块变化时,可预先计算并供主循环反复使用。随后分数梯度包含P[q,k]*(dY[q]·V[k]-delta[q])及目标定义的scale/mask等;来源只写Sub,没有展开乘P,不能因此删除这一步。dropout等变体需确认使用的O与梯度公式确实对应。
24+ 
25+预计算沿真实value/output通道D归约,按batch、query head和有效query行保存一个结果。headNum*s1Loops*sizeof(float)只有s1Loops恰好表示所有实际行数且其他维已计入时成立;若它是tile次数,还需每tile有效行数,定长需B,变长需实际累计query边界。消费端的offset、广播行和尾块必须匹配同一布局。
26+ 
27+来源独立VecSfmg阶段示意24KiB×2输入、48KiB×2扩宽区,并未完整区分dY与O两份输入/转换及乘积、输出、归约scratch。应按同时活跃的真实字节数核算,不能把每个阶段≤UB的80%当成总容量证明;两个TPipe也可能管理同一物理空间。TBuf并不因InitBuffer传一个2就自动获得合法ping/pong使用。
28+ 
29+预计算结果写GM完成并发布后,Main才能读取;同核局部交接与跨核发布采用各自协议,不能以一个SyncAll名字证明覆盖全部范围。最后一块写出、空核参与和阶段资源回收都需保留。FP16→FP32等扩宽按实际支持的Cast模式执行,不能仅由CAST_RINT与CAST_NONE标签断言某次精度问题根因。
30+ 
31+能省去多少重复计算取决于S2复用次数,而非s1Loops<2就必然无收益。若dY/O在原循环中发生有效更新,前移会读取错误版本。用多个query/head、D尾部、多个S2块和空输入验证delta与完整dQ/dK/dV,评估新增workspace写读、初始化和同步后的总时间。
@@ -0,0 +1,38 @@
1+---
2+type: Guide
3+title: Attention 在线 Softmax 的空行与量化阶段合同
4+description: 修正有限 mask、初始 rescale 及 QK/PV 归约轴混淆。
5+tags:
6+- task.optimization
7+- topic.implementation.tiling
8+- topic.performance.analysis
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-tiling-design/references/flashattention/subfamilies/base_design.md
12+ role: primary
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-tiling-design/references/flashattention/subfamilies/quantization_design.md
14+ role: supporting
15+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-tiling-design/references/flashattention/design/execution.md
16+ role: supporting
17+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-tiling-design/references/flashattention/implementation_ref.md
18+ role: supporting
19+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-tiling-design/references/flashattention/specialization/feature_flags.md
20+ role: supporting
21+created_at: '2026-09-11T04:27:49Z'
22+updated_at: '2026-09-11T04:27:49Z'
23+platforms: []
24+---
25+ 
26+# Attention 在线 Softmax 的空行与量化阶段合同
27+ 
28+每KV块在线更新为mNew=max(mOld,maxS),alpha=exp(mOld-mNew),P=exp(S-mNew),lNew=alpha*lOld+sum(P),ONew=alpha*OOld+P*V,末尾O/l。首个有限有效块mOld=-inf时alpha数学值为0,不是1;若跳过旧状态读取并直接写首结果,可以避免初始化旧GM,但不能通过0*未初始化NaN达到同等效果。
29+ 
30+空或全mask行需显式定义中性状态和最终输出。将所有S填-1e30后减本行max会得到全0,exp全部为1,产生均匀权重,不是零权重;禁止-inf也不能修复此问题。可以使用有效性mask和对空行的分支,采用-inf时避免-inf-(-inf)。在至少一个有限有效score的相对max表示下l≥1;任意加epsilon会改变算子数学,不能当默认修复。Inf/NaN输入和完全不可见行需遵循接口合同。
31+ 
32+GQA中mEff包含不同head,同一query在多个head重复;因果边界最大query为qOff+curBq-1,不能用qOff+mEff-1。SoftMaxShapeInfo逻辑行数、mPad物理行数及isBasicBlock需对应真实布局,padding行不得写成有效输出。
33+ 
34+S和PV累加通常保留FP32,P可能转换为FP16/BF16/MX格式;V输入保持BF16不表示整个Vector阶段无需Cast。来源把量化BMM1结果写成1字节与其FP32累加路径矛盾,不能据此缩减S缓冲。Cast、Fixpipe转换及位置支持应查实际平台/重载,不能把某路径仅能UB转换扩大为所有转换指令限制。
35+ 
36+MX中Q/K的归约轴是Dqk,P/V的归约轴是Sk,因此V scale的分组轴沿Sk,不能沿Dv替代。来源32元素分组、E8M0与特定存储是其路径合同;scale的head stride、padding及打包格式需独立于数据位宽核算。P及对应scale应在每个生成块交给C2前共同发布,并保持共同生命周期,不能只在最后KV迭代更新P scale。
37+ 
38+来源c1v1Loop将BMM1 K写为S2是轴混淆:QK归约K=Dqk,S2是score的N列;PV归约K才是S2。ceil(actS2/256)可能描述N分块,不能当QK reduction循环。K_BASE、yStep和末块长度需有明确单位。输入量化不会让FP32 S/PV缓冲减半,也不自动使单缓冲调度安全;INT8反量化、MX与其他accumulator路径不能互换合同。
@@ -0,0 +1,26 @@
1+---
2+type: Guide
3+title: Softmax 在线重标定与有限 Mask 限制
4+description: 保存分块状态和压缩偏移,限定高阶API精度声明。
5+tags:
6+- task.implementation
7+- topic.precision
8+- topic.implementation.tiling
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/collaborative-agent-kernel-evolution/skills/cake-evo/references/meta_prompts/strategies/acc_05_softmax_stability.md
12+ role: primary
13+created_at: '2026-09-11T05:47:04Z'
14+updated_at: '2026-09-11T05:47:04Z'
15+platforms: []
16+---
17+ 
18+# Softmax 在线重标定与有限 Mask 限制
19+ 
20+标准有限输入softmax先减全行max再exp即可限制每个指数≤1;完整score矩阵的O(Sq*Sk)存储压力不等于exp本身必溢出。在线softmax节省中间矩阵并维护分块running max/sum/output,不应把空间收益说成只有在线算法才能数值稳定。
21+ 
22+若块max为m_b、按m_b算l_b与a_b,合并m=max(m_old,m_b),alpha=exp(m_old-m)、beta=exp(m_b-m),l=alpha*l_old+beta*l_b,a=alpha*a_old+beta*a_b。若当前块直接按新m算exp,可省显示beta;来源伪码只写running_sum*alpha+block_sum,成立前提是block_sum已在新m域,不能把局部m_b域统计直接相加。最终输出a/l,旧output重标定应对应未归一化表示。
23+ 
24+有限mask值-10000并不在任意score/scale下保证无效概率0,特别是有效score更小、全mask或scale作用顺序改变时;应定义全mask、空row、Inf/NaN和输出策略。Causal/Band/Prefix压缩偏移须与所选mask物理形状对应,ComputeOffsetForCausal的delta≤0取min(-delta,s1Base)、否则min(delta,s2Base)*maskS2只是该压缩模板的地址算法,宽有符号边界和布局仍需检查。
25+ 
26+M对齐8/K对齐64及SoftmaxConfig{false}是来源组件配置,不是所有SoftMax接口的常量。高阶SoftMax<float,false,false>仍需正确scratch、tiling与有效shape;不能由API名称证明其内部必有更高精度、Kahan、额外特殊值处理或一定优于手写版本。保持输入scale/mask、padding和输出Cast合同后,才比较有限数误差及特殊值。
@@ -0,0 +1,24 @@
1+---
2+type: Guide
3+title: 稀疏与稠密分发需要有效集合等价
4+description: 区分块数量预算和实际 token 覆盖,保留分发后的掩码与资源合同。
5+tags:
6+- task.design
7+- topic.attention
8+- topic.compute.indexing
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/cards/P62_sparse_dense_dual_path_dispatch_dra.md
12+ role: primary
13+created_at: '2026-09-11T02:07:16Z'
14+updated_at: '2026-09-11T02:07:16Z'
15+platforms: []
16+---
17+ 
18+# 稀疏与稠密分发需要有效集合等价
19+ 
20+Sparse Attention 可以在有效稀疏选择覆盖完整可达区域时走稠密路径,避免无收益的索引和Gather。来源判断 `curMaxS2 <= actualSelectedBlockCount*selectedBlockSize` 只比较长度与容量;它不是覆盖证明。重复块、索引孔洞、非零起点、页边界、最后块有效长度,以及不同query的causal范围,都可能使乘积足够但实际集合仍缺项。
21+ 
22+合法分发应证明稀疏路径与稠密路径处理相同的逻辑token集合及相同权重语义;若原语义允许重复索引贡献,去重后的集合相同仍不充分。sink、mask、位置编码、head映射及有效序列长度必须一致。空集合需要定义输出,不应以0<=0机械调用不支持空输入的路径。
23+ 
24+Host条件、Kernel分支、输入起点、workspace大小与同步事件需匹配各路径。稠密路径免除稀疏Gather,不意味着可复用尚未发布的数据或取消其他阶段依赖。用全覆盖、有孔洞、重复项、跨页和causal尾部对照两条路径的输出;再评价分发判断及资源增加的总成本。该来源虽标为draft/不建议直接应用,其覆盖判定机制仍可提炼;示意条件不被认证为完备实现。
@@ -0,0 +1,38 @@
1+---
2+type: Guide
3+title: Attention Tiling 的完整内存预算与流式边界
4+description: 按物理布局和同时存活对象核算 UB、L1、L0 与 workspace。
5+tags:
6+- task.optimization
7+- topic.implementation.tiling
8+- topic.performance.analysis
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-tiling-design/references/flashattention/design/resources.md
12+ role: primary
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-tiling-design/references/flashattention/design/shape.md
14+ role: supporting
15+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-tiling-design/references/flashattention/design/execution.md
16+ role: supporting
17+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-tiling-design/references/flashattention/implementation_ref.md
18+ role: supporting
19+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-tiling-design/references/flashattention/foundation/extension_points.md
20+ role: supporting
21+created_at: '2026-09-11T04:27:49Z'
22+updated_at: '2026-09-11T04:27:49Z'
23+platforms: []
24+---
25+ 
26+# Attention Tiling 的完整内存预算与流式边界
27+ 
28+UB需逐项列出V1输入/输出、max/sum/rescale统计量、V2输入/输出、更新scratch、mask、量化辅助区及真实buffer份数,并逐项满足对齐。max、sum、rescale是不同状态;若每行按32B存储,slots*m*32只计一类,不能遗漏另外两类或其他副本。queue深度是生产消费容量,未必等于已分配buffer数量,应以InitBuffer和同时存活量为准。
29+ 
30+按固定字节容量分行时chunkRows=floor(bytes/(Dpad*computeBytes))可能为0;若单行已超容量,需要列分块及可合并的行统计/更新,或明确不支持,不能宣称固定buffer覆盖任意D。仅切M仍保留D因子,未完全解除大D限制。整块m*D驻留在能容纳时是合法选择;V1/V2独立按各自物理维度预算,所谓“对称分块”不要求相同chunk数。
31+ 
32+L1中的A1/B1是API逻辑位置,不能无证据视为完全独立物理容量或绝对端口隔离。Q/P/K/V、scale及双缓冲按真实共享资源与生命周期合计;来源复用区需要前一消费者完成。LoadData多M-fractal读取必须对应完整连续NZ布局,只准备一个fractal却以mStep>1读取会越过有效布局。尾M/N/K的有效区和padding需分别传递。
33+ 
34+L0A预算m*k*ABytes*buffers,L0B为k*n*BBytes*buffers;L0C对S和PV取max只在两者生命周期不重叠且Fixpipe完成后复用时成立,否则应相加或采用独立分配。物理fractal与对齐后的形状才是容量依据。
35+ 
36+workspace按核数乘slot大小仅适合前一任务所有消费者完成后可复用的区。Split-KV局部结果通常需按query/split身份保留,稀疏metadata也可能按任务索引;不能一概禁止per-task空间。运行期实际shape分配与最大shape保留空间各有取舍,graph capture等场景可能需要稳定容量。
37+ 
38+Host预算必须与Kernel所有实际InitBuffer、局部张量/外部workspace偏移逐项相符;只复制文档中的不完整公式不能证明安全。模板已编码的模式未必需要冗余Tiling字段,但Host dispatch必须确保选中匹配变体。所有字节乘积、对齐、slot索引都需防溢出并处理0容量。
@@ -0,0 +1,29 @@
1+---
2+type: Guide
3+title: Varlen L1 Carry 与 SameAB 约束
4+description: 保存驻留优化条件,修正跨batch无条件复用与容量常量。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.implementation.synchronization
9+- topic.performance.analysis
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-varlen-handcraft-recipe.md
13+ role: primary
14+created_at: '2026-09-11T05:21:50Z'
15+updated_at: '2026-09-11T05:21:50Z'
16+platforms: []
17+---
18+ 
19+# Varlen L1 Carry 与 SameAB 约束
20+ 
21+来源needL1Carry候选条件为非SameAB、非FP32、无Rope、S1基本块和D/Dv在实现上限内,并切到UnpackInit/ProcessL1Carry及PairGetS1LoopRange。shareL1Size设0表示改变组件间资源分配,不是物理L1变0或驻留数据无需内存。
22+ 
23+KV整batch驻留需要真实KV字节加P/Q、阶段和对齐满足目标L1;只看S1/D阈值若未同时限制KV长度仍不足。可以跨同batch的多个Q任务复用相同KV,跨batch通常数据不同需重新加载,除非证明输入内容/身份确实共享。来源“跨task/batch复用”不能理解成保留上个batch的KV继续算。
24+ 
25+source首次不开Carry、shareL1Size正常1MB是其工程建议与配置描述,不等于A2物理L1容量1MiB;必须按实际API字段语义和资源查询核算。禁用L1Reuse奇偶对而改needL1Carry也属于两种组织差异,不是varlen数学要求。
26+ 
27+SameAB指Q/K/V为同一GM张量且shape/stride/语义兼容,Sq==Sk仅为必要关系之一;一般self-attention也可能三次投影产生不同Q/K/V,不能只因self-attention标签自动共享加载。Q_EVENT0=2/KV_EVENT0=4/P_EVENT0=6是该类自管候选事件分配,不能与其他组件重复占用。
28+ 
29+aux max/sum[sum Sq,Hq,8]和TND输出应与每batch有效query总数一致。直接entry的REGIST、混合核型和userWorkspace合同来自选定实现,不因增加qList/kvList两个指针就自动完成全部长度适配。
@@ -0,0 +1,31 @@
1+---
2+type: Guide
3+title: Varlen 任务前缀与 Sparse 边界
4+description: 保存四类累计量、TND偏移与尾块双重约束。
5+tags:
6+- task.implementation
7+- topic.attention
8+- topic.implementation.synchronization
9+- topic.performance.analysis
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/catlass-op-develop/references/patterns/fa-varlen-handcraft-recipe.md
13+ role: primary
14+created_at: '2026-09-11T05:21:50Z'
15+updated_at: '2026-09-11T05:21:50Z'
16+platforms: []
17+---
18+ 
19+# Varlen 任务前缀与 Sparse 边界
20+ 
21+该s1s2路径的长度输入为B项累加终点,差分得到qLen_b/kvLen_b。任务数sum_b ceil(qLen_b/s1Base)*n2G,其中n2G代表KV头与GQA组的组合数量;按实际长度切任务避免给短序列分配padding任务。整数范围、0长度batch和全部空输入需在ceil除法/索引前处理。
22+ 
23+定位任务起点和跨batch推进时同时维护任务前缀s1OuterSizeAcc、Q token前缀s1SizeAcc、KV token前缀s2SizeAcc、显式逐batch mask元素前缀attenB1SSOffset=sum(qLen*kvLen)。跨多个空/短batch要while且守住boIdx<B,不能单次if。若mask是广播或压缩存储,最后那个乘积前缀应改用其真实布局,不能无条件使用。
24+ 
25+局部task经qBlocks=ceil(qLen/s1Base)解码n2o=(idx-prefix)//qBlocks//g、go=(idx-prefix)//qBlocks%g、s1o=(idx-prefix)%qBlocks。Q地址为prefixQ*(Hq*D)+s1o*s1Base*(Hq*D)+(n2o*g+go)*D;K为prefixKV*(Hkv*D)+(s2Start+s2Loop*s2Base)*(Hkv*D)+n2o*D,V需用独立Dv与stride。不能对K/V使用Q前缀或b*固定长度。
26+ 
27+纯varlen的s2范围[0,kvLen);左上causal块上界min((s1o+1)*s1Base,kvLen),右下另加kvLen-qLen,并夹到合法非负范围。BAND取max(qBlockStart-pre,0)与min(qBlockEnd+next,kvLen),这些是块候选范围,行内仍要mask。来源空band退回全扫只可作为带完整mask的计算fallback,不能让原本无效key获得概率。
28+ 
29+每次Matmul SetTail约束Q/KV/归约轴的实际尾部,softmax还要将无效score列置-inf或等价mask;SetTail不独自定义softmax分母。s2LoopLimit=ceil((end-start)/base)-1时空范围不进入循环,但finalize必须处理该任务空语义,不能读上个task状态。
30+ 
31+extraInfo三槽沿全局已发布流水step递增,跨s2内层不能随意重置;每task首s2块才切其softmax状态,保留task身份。needNz2Nd、bmm1Nz与普通bmm1分发由dtype/尾部/布局真实条件决定,不能仅模板名字推断所有FP32/BF16尾都必走NZ。
@@ -0,0 +1,41 @@
1+---
2+type: Guide
3+title: Workspace 槽地址与延迟梯度 Scatter
4+description: 区分阶段槽、梯度聚合域和最终刷新时机,保留完整 dK/dV 贡献。
5+tags:
6+- task.design
7+- topic.attention
8+- topic.memory
9+- topic.implementation.synchronization
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/cards/P43_grad_pingpong_deferred_scatter.md
13+ role: primary
14+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/playbooks/P40_workspace_double_buffer.md
15+ role: supporting
16+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/playbooks/P43_grad_pingpong_deferred_scatter.md
17+ role: supporting
18+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/preconditions/P40.yaml
19+ role: supporting
20+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/plugins-community/ops-perf-evolution/skills/evolution-strategies/references/preconditions/P43.yaml
21+ role: supporting
22+created_at: '2026-09-11T01:49:44Z'
23+updated_at: '2026-09-11T01:49:44Z'
24+platforms: []
25+---
26+ 
27+# Workspace 槽地址与延迟梯度 Scatter
28+ 
29+每个核的多阶段 GM workspace 应以字节为单位分区。若阶段 j 的每槽需要 B_j 字节、具有 p_j 个物理槽,则每核至少 sum_j(p_j*B_j),总容量再按实际独立核组数量和对齐计算。对于已经包含核数/槽数的上层 size,不要重复乘。SetGlobalBuffer 前的 workspace 指针加法必须确认底层指针元素宽度,否则名为 offset 的字节数可能再次按类型缩放。
30+ 
31+来源 `aiCoreIdx*dbWorkspaceRatio*mmResUbSize*sizeof(MM1_OUT_T)` 将 mmResUbSize 当作元素数;而后续 `GlobalTensor<T>[slot*bmm2ResUbSize]` 也按元素寻址。MM1/MM2/vec2 区可能 dtype和长度不同,不能共享未经核对的大小。dbWorkspaceRatio 与 preLoadNum 至少要保证所有被选择槽都有真实容量,preLoadNum>0,实际计算长度不能超槽。
32+ 
33+`(loop-1)%preLoadNum` 在 loop=0 之前必须经过任务有效性判断,无符号回绕后取模可能恰好落到合法槽,但数据仍未被生产。阶段游标不一定与外层循环同速,ready 和所有消费者完成后的 release 都要绑定真实任务;双缓冲扩容本身不使 Cube 写与 Vector 读自动并行。
34+ 
35+反向 Attention 的 dK/dV 可先在选定梯度 workspace 累积,再在一个明确的 S1 聚合范围完成后执行 scatter-add,减少细粒度外写。四槽索引 `idx=(idx+1)%4` 只是存储策略,与 Matmul 的二槽 mmPingPongIdx 是不同状态。源 selectedKGmOffset 在增加 idx 前计算,后续消费者必须保留这一已选槽,不能误读增加后的索引。
36+ 
37+changeS1 应发布刚完成的聚合组,且初组不能刷新未初始化数据。最后一组通常没有下一次 changeS1,必须在结束时显式刷新;空组、稀疏跳块和尾部也要有确切行为。每槽在重新使用前需完成本组全部累计、scatter消费和释放,四槽不能无限保存尚未刷新的组。对重复索引、跨核相同目标和 dK/dV 两路,必须保留全部贡献并遵守目标数值顺序要求。
38+ 
39+来源只展示 CrossCoreWaitFlag,没有相应的生产端和回收端,且具体 mode/pipe依赖平台与Kernel类型;不能据此确立完整协议。集合同步、局部事件和原子累计各自承担不同职责。私有partial之后固定顺序归并的边界另见[确定性workspace](../reductions/deterministic_workspace_merge.md)。
40+ 
41+若每个S1组原来按其S2子任务逐次scatter,延迟后调用次数可从子任务总数降到非空聚合组数;这不等于无条件由 O(S2) 降为 O(S1),两轴大小、循环嵌套和每次scatter处理量都应明示。测量总流量、workspace峰值、等待和整体时延,验证首组/最后组、跨四槽、多核重复目标和不同真实序列长度。
@@ -0,0 +1,37 @@
1+---
2+type: Guide
3+title: Blaze 硬件基形与运行期尾块范围的区别
4+description: 将 Block/Copy 的合法 base shape 与 Scheduler 的当前逻辑范围分开,按实际消费者校验 tiling、workspace
5+ 和融合后处理。
6+tags:
7+- task.design
8+- topic.tiling
9+- topic.alignment
10+status: draft
11+sources:
12+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-blaze-best-practice/references/kernel-design/tiling-selection.md
13+ role: primary
14+created_at: '2026-09-10T12:37:22Z'
15+updated_at: '2026-09-10T12:37:22Z'
16+platforms: []
17+---
18+ 
19+# Blaze 硬件基形与运行期尾块范围的区别
20+ 
21+Blaze tiling 字段的含义来自消费它的 Scheduler、Policy、Block 和 Copy 特化。同名字段或相似样例不足以证明两个方案可共用 tiling。
22+ 
23+用于片上资源或搬运原语实例化的 baseM/baseN/baseK 是硬件基形,需要满足该特化的粒度、对齐、容量和交叉约束;curM/curN/curK 或等价当前块范围表达本轮有效逻辑数据。将基形设置成 `min(逻辑尺寸, 首选基形)` 可能产生不合法的小块。
24+ 
25+一个可执行尾块方案需要同时具备合法 base shape、当前有效范围以及所选 Block/Copy 对该范围的处理能力。若该特化不支持某种尾块,应选择具有对应支持的实现,不能仅靠缩小 base 字段宣称兼容。
26+ 
27+## Tiling 与融合的联合约束
28+ 
29+Host 生成的字段应逐项对应 Device 消费者,明确类型、单位、合法范围、核数、workspace 与输出偏移。Grouped 输入还要核对组元数据,量化路径还要核对 scale 和 packing。
30+ 
31+在 MatMul 后融合 Elementwise/Broadcast Epilogue 时,额外输入、广播映射、UB 暂存、slot 和输出应纳入同一资源及生命周期预算。SplitK/StreamK 产生的 partial 必须先按算法完成合并,再交给要求完整结果的后处理;普通逐 tile Epilogue 不能自动获得 final 语义。
32+ 
33+## 验证
34+ 
35+对有粒度约束的轴分别覆盖小于粒度、等于粒度、base+1 和多 tile,Grouped 还需覆盖空组。不同失败用例若共享同一个小 M,不能仅依据用例名叫 K-tail 或 N-tail 就定位到对应轴。
36+ 
37+本卡提供字段与边界的推导原则,没有给出某个 Blaze 版本的固定合法参数组;具体支持域依赖选定源码特化,故不声明统一平台和固定默认值。
@@ -0,0 +1,36 @@
1+---
2+type: Guide
3+title: 权重反量化的 Cast 顺序与 V→C 释放关系
4+description: 补全来源 INT8/FP8 两步转换、Bias 元信息、前后融合边界及诊断限制。
5+tags:
6+- task.design
7+- topic.compute.quantization
8+- topic.implementation.synchronization
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-blaze-best-practice/references/scenarios/weight-dequant-prologue-fusion/layout-variant-boundary.md
12+ role: primary
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-blaze-best-practice/references/scenarios/weight-dequant-prologue-fusion/vc-sync-and-l1-layout-design.md
14+ role: supporting
15+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-blaze-best-practice/references/scenarios/weight-dequant-prologue-fusion/pitfalls-and-diagnosis.md
16+ role: supporting
17+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-blaze-best-practice/references/scenarios/weight-dequant-prologue-fusion/weight-quant-tiling-contract.md
18+ role: supporting
19+created_at: '2026-09-11T03:07:47Z'
20+updated_at: '2026-09-11T03:07:47Z'
21+platforms: []
22+---
23+ 
24+# 权重反量化的 Cast 顺序与 V→C 释放关系
25+ 
26+来源逐通道反量化先Cast(B),再按N广播加offset、乘scale,然后Matmul。INT8例为int8→half→BF16,FP8例为FP8→FP32→BF16,再在BF16进行加法与乘法,最后A/B反量化值转FP32累加并转目标输出。每步舍入/特殊值应与实际设备合同一致;把所有运算提升到FP32后最后只cast一次通常改变结果。
27+ 
28+Golden的关键是模拟实际输入与逐步表示,不能按来源断言“任何C++ FP32模拟都不允许”:若在每个规定位置正确执行BF16舍入,也可形成相应参考。PyTorch名称本身不证明其计算路径等于目标设备。来源重复列出E4M3而在支持表还称E5M2已验证,不能据此提升未具体取证的格式为本批支持。
29+ 
30+Cast mask控制的lane域与Add/Mul/Store的操作数域可能不同,升宽后需按实际API和寄存器重排核验,不机械复用源mask。CastTrait作为模板实参的存储期/声明形式取决于具体模板参数规则;“必须static constexpr”是来源故障线索,不是任意编译期常量的通用法则。
31+ 
32+AIV将B_dequant写L1,AIC从同一L1范围加载L0B,原AIC GM→L1 B路径不能再覆盖它。AIV写完发布ready,AIC最后读完发布可复用;V+C仍有反向释放依赖。分离AIC/AIV是不同执行核,不能因同一logical block而仅用核内HardEvent替代跨核协议。首轮许可、两槽、真实参与AIV及末轮drain与布局共同闭合。
33+ 
34+Perchannel scale/offset只沿N变化,pergroup还需明确K/group维度及额外加载循环;两者不能仅切换一个标签。原生量化Matmul的scale路径与先生成浮点权重再乘的位置和舍入可能不同。前置反量化和后置Vector Epilogue可以在同一数学需求共存;来源称二者不可能同时出现只是其场景路由限制。实现V+C+V需新的生命周期/资源组合,不能由两个单独通过的叶子证明。
35+ 
36+Dequant-only、使用预反量化B的A-only Matmul与Full分别隔离数值转换、Cube计算和协作,但都要保持相同表示和偏移。某组合失败时检查transpose切分、metadata、Bias类型与UB/L1容量;不靠放宽精度或固定错误码判根因。本批没有复现来源历史验证。
@@ -0,0 +1,42 @@
1+---
2+type: Guide
3+title: Blaze 逐元素 Epilogue 的广播映射与加载边界
4+description: 补全 MemBase/RegBase 的 operand 映射、独立 slot 预算与 mask 前整向量读取。
5+tags:
6+- task.design
7+- topic.memory
8+- topic.precision
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-blaze-best-practice/references/scenarios/elementwise-broadcast-epilogue-fusion/elementwise-broadcast-epilogue-fusion-design.md
12+ role: primary
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-blaze-best-practice/references/scenarios/elementwise-broadcast-epilogue-fusion/elementwise-broadcast-epilogue-fusion-development.md
14+ role: supporting
15+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-blaze-best-practice/references/scenarios/elementwise-broadcast-epilogue-fusion/block-l0c2ub-extension.md
16+ role: supporting
17+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-blaze-best-practice/references/scenarios/elementwise-broadcast-epilogue-fusion/fused-kernel-development.md
18+ role: supporting
19+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-blaze-best-practice/references/scenarios/elementwise-broadcast-epilogue-fusion/epilogue-membase-design.md
20+ role: supporting
21+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-blaze-best-practice/references/scenarios/elementwise-broadcast-epilogue-fusion/epilogue-regbase-design.md
22+ role: supporting
23+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-blaze-best-practice/references/scenarios/elementwise-broadcast-epilogue-fusion/precision-diagnosis.md
24+ role: supporting
25+created_at: '2026-09-11T03:07:47Z'
26+updated_at: '2026-09-11T03:07:47Z'
27+platforms: []
28+---
29+ 
30+# Blaze 逐元素 Epilogue 的广播映射与加载边界
31+ 
32+Matmul最终T[m,n]之后的逐元素/broadcast DAG,应使每个Y[m,n]仅依赖同一位置T及明确映射的operand。常见scalar索引为0、row输入取m、column取n、同形输入取(m,n),实际物理地址再由各自stride与dtype得到。quant/dequant若scale已给定且仅逐元素应用可满足此条件;现场求行scale、Softmax、LayerNorm、scan、gather或邻域运算则增加跨元素依赖,需要不同状态/调度。
33+ 
34+L0C有效范围、实际Copy extent、目的UB pitch与最终GM有效列是不同量。每个slot分别预算C、每路operand staging、DAG中间值、输出和必要guard;不能给每slot都分“全部UB”。GetTensor(slot)返回字节偏移、元素偏移还是Tensor视图须按接口处理一次。SplitM已把本地C放slot起点时不再叠sub偏移,GM全量operand/输出则仍要正确全局行列。
35+ 
36+MemBase通常以LocalTensor表达中间运算,RegBase可在VF内保留寄存器中间值。操作数量不能独自决定优劣:API可用性、broadcast实现、实际load/store、寄存器活跃量、UB staging、事件与数值顺序都影响选择。较长链可能减少UB往返,也可能增加寄存器压力或不支持的调用。
37+ 
38+active mask只约束具体API允许的运算/存储行为。如果某load在mask应用之前读取完整向量,即使只计算尾部几个lane,也必须有合法地址和完整物理容量,必要时增加按真实读取跨度计算的guard。guard保护访问范围,不改变逻辑N,也不能让padding参与结果。各operand的窄dtype/不同pitch单独计算,不能套C的容量。
39+ 
40+输入ready、输出store完成及下一轮覆盖是不同依赖;cross-core wait要覆盖Epilogue真正首消费者,然后按需要连接本地MTE/V。全常量输入可能掩盖行列地址错,用行编码、列编码、one-hot及非零known-C检查非零offset、tail、slot复用和空sub。tanh/GELU的指数写法在大值可能先Inf再Inf/Inf,需按原合同选择稳定实现并验证中间值,不能仅凭最终absmax忽略NaN。
41+ 
42+来源五模式和单变量对照可支持定位,但固定拒绝所有非有限值、强制某条工程路线、固定验证次数不是所有算子的数学合同。本文不移入PLAN/资产操作规则;设备收益与支持范围仍需实际实例验证。
@@ -0,0 +1,38 @@
1+---
2+type: Guide
3+title: Blaze 执行引擎、布局标签与特化
4+description: 纠正框架总览的流水方向与核数查询,保留抽象层和数据 Trait 的区别。
5+tags:
6+- task.design
7+- topic.programming_model
8+- topic.data_movement.layout_transform
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-blaze-best-practice/references/fundamentals/blaze-framework-overview.md
12+ role: primary
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-blaze-best-practice/references/fundamentals/blaze-sync-patterns.md
14+ role: supporting
15+- resource: https://gitcode.com/qq_41667743/cannbot-knowledge/blob/a70025d581ae45da8a93b4938b5ba19da6fa6f92/knowledge/ops/ascendc/apis/utils_api/platform_info/platform_ascend_c/get_core_num_aic.md
16+ role: constraint
17+- resource: https://gitcode.com/qq_41667743/cannbot-knowledge/blob/a70025d581ae45da8a93b4938b5ba19da6fa6f92/knowledge/ops/ascendc/apis/utils_api/platform_info/platform_ascend_c/get_core_num_aiv.md
18+ role: constraint
19+- resource: https://gitcode.com/qq_41667743/cannbot-knowledge/blob/a70025d581ae45da8a93b4938b5ba19da6fa6f92/knowledge/ops/ascendc/apis/utils_api/platform_info/platform_ascend_c/get_core_mem_size.md
20+ role: constraint
21+created_at: '2026-09-11T02:57:41Z'
22+updated_at: '2026-09-11T02:57:41Z'
23+platforms: []
24+---
25+ 
26+# Blaze 执行引擎、布局标签与特化
27+ 
28+在来源所述分离AIC/AIV模型中,常见AIC数据流为GM→MTE2→L1→MTE1→L0A/L0B→Mmad→L0C→Fixpipe→UB/GM;AIV为GM→MTE2→UB→Vector→UB→MTE3→GM。总览图把MTE1/MTE2方向写反,并把MTE3与Fixpipe混同,应按真实数据流修正。不能据总览“AIC包含Vector”推导在分离架构的AIC实例中任意Vector调用均有效。
29+ 
30+GetCoreNumAic查询Cube核数,GetCoreNumAiv查询Vector核数;总览将后者赋给aicNum并注释AIC不正确。耦合架构的返回含义与分离架构不同。L1/L0A/L0B/L0C/UB容量应来自目标平台和实际可用资源,来源512/64/64/256/256KiB表只是特定上下文的候选值,不能跨版本硬编码;GetCoreMemSize的真实原型、存储枚举和返回方式也应核对,不能照抄GetCoreMemSize(0/1)的示意。
31+ 
32+Blaze在选定Kernel/Block/Policy中封装tile循环、K迭代、缓冲和同步;Tensor API提供张量描述与Copy/Mmad路由;Ascend C原语执行实际指令。封装范围随组件而变,不代表所有组合均有完整同步,也不能因header-only/编译期标签便宣称整个运算零运行时开销。
33+ 
34+Tensor API的Shape、Stride、Coord可为嵌套tuple,Layout还携带Pattern/Trait。来源NDExt/DNExt用于外层ND/转置表示,ND/DN用于相应内部表示,NZ/ZN为分形标签,ZZ/NN及ScaleA/B的ND/DN标签用于相应Scale路由。标签不是实际数据变换;源/目标位置、arch、Pattern和dtype/Trait共同决定可用Copy特化,Ext与非Ext不能只按名称互换。
35+ 
36+32B数据载体对FP16/FP32/8bit/FP4分别为16/8/32/64逻辑元素,但Scale Trait可能使用C0=2,L0C分形边长也不能套32/sizeof规则。具体Trait决定布局,见[物理存储](matmul_logical_shape_and_physical_storage.md)。MakeFrameLayout构造描述,MakeCopy/MakeMmad组装调用对象,with注入参数;描述有效不证明真实消费者支持非连续stride或尾部。
37+ 
38+Policy的ScheduleType可驱动模板偏特化,具体Kernel还依赖ProblemShape、BlockMmad、Epilogue、Scheduler等类型。公开定义存在不证明某组类型已经构成可执行实例;Batch、Grouped、量化、MX、融合是可能叠加的需求轴,各自组合仍须满足同一入口的真实绑定。本文不继承来源固定版本支持全集,也未进行设备验证。
@@ -0,0 +1,42 @@
1+---
2+type: Guide
3+title: APACE FragmentTensor 的离散视图与搬运消费者
4+description: 以地址表描述离散GM段的逻辑拼接,区分虚拟切片、实际gather/scatter、填充长度和异步生命周期。
5+tags:
6+- task.design
7+- topic.data_movement.layout_transform
8+- topic.memory
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-mc2-best-practice/references/foundations/apace/fundamentals/compute.md
12+ role: primary
13+created_at: '2026-09-10T14:42:48Z'
14+updated_at: '2026-09-10T14:42:48Z'
15+platforms: []
16+---
17+ 
18+# APACE FragmentTensor 的离散视图与搬运消费者
19+ 
20+来源描述的 APACE FragmentTensor 用地址表把多个 GM 段沿一条逻辑轴拼接成张量视图。构造视图或调用 Slice 本身不把各段搬到连续内存;实际搬运由支持 fragment 遍历的消费者执行。
21+ 
22+## 逻辑跨度与有效跨度
23+ 
24+来源中的 FragmentParam 分别记录 assembledShape、assembleAxis、fragmentSize、realFragmentSize 和 fragmentCnt。其模型用 fragmentSize 描述每段的逻辑跨度,用 realFragmentSize 限制本段有效范围;拼接轴跨度与 fragmentSize*fragmentCnt 对应。尾块填充必须同时反映在视图和实际 copy 的截尾处理里,不能只扩大 shape 而继续读取无效行。
25+ 
26+地址表按本次逻辑拼接的顺序排列。A、ScaleA 和输出段可能使用各自的地址表或映射,重排一个表后必须逐项确认下游坐标仍指向同一份数学数据。地址表项数上限和对象的固定数组容量属于具体版本实现约束,应分别核对,不能由 rank 数或通信轮次数相互推断。
27+ 
28+## 消费者需要真实支持离散段
29+ 
30+来源里的 FragmentSliceCopy<false> 遍历离散段执行 gather,FragmentSliceCopy<true> 执行 scatter,并在相关路径按有效跨度截尾。其 Fragment 版 BlockMmad 接入相应搬运路径;只接受 GM_ADDR 的普通输出参数不能直接接收多段 FragmentTensor。
31+ 
32+一个可表示跨 fragment 的 Slice,不代表任何 Tensor API Copy 或 Matmul 特化都能消费它。支持范围需沿调用实例确认;参考文档中的 ForEachFragment 仅是伪代码名称。精确模板签名、容量常量及架构门控以实际 APACE/Blaze 版本为准。
33+ 
34+## 地址表和数据的生命周期
35+ 
36+来源说明 MakeFragmentTensor 保留地址表指针,因此表本身必须活到所有视图使用结束;局部临时数组退出作用域后,视图会失去有效描述。数据缓冲区还必须活到异步搬运完成。
37+ 
38+UpdateAddrList 或复用底层段之前,应确认旧视图及其消费者已完成;新一轮通信就绪和旧一轮读取完成是两个条件。对于来自通信窗口的段,需要先满足相应生产完成依赖;本地输入段仍需要满足其自身的生产和复用条件。
39+ 
40+## 验证
41+ 
42+用每段不同的值和非零 Slice 起点检查段内与跨段地址,覆盖最后一个 fragment、非整齐尾块和多轮地址更新。对照显式拼接后的数学结果,同时检查真实访问不越过有效段。若用它减少中间重排,应由搬运轨迹确认收益,视图的零搬运构造不等于整个算子零拷贝。相关边界见[布局与Copy特化](tensor_layout_and_copy_specialization.md)。
@@ -0,0 +1,38 @@
1+---
2+type: Guide
3+title: Blaze 融合输出的整行量化与阶段边界
4+description: 在 Matmul 或 GLU 后按完整逻辑行求 scale,区分 workspace 行距、分块宽度、量化语义和异步完成范围。
5+tags:
6+- task.design
7+- topic.compute.quantization
8+- topic.precision
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-blaze-best-practice/references/kernel-design/per-token-quant-development.md
12+ role: primary
13+created_at: '2026-09-10T14:12:04Z'
14+updated_at: '2026-09-10T14:12:04Z'
15+platforms: []
16+---
17+ 
18+# Blaze 融合输出的整行量化与阶段边界
19+ 
20+设融合计算产生逻辑 workspace `W[R,Q]`,每个输出 token 对应一整行。per-token scale 的统计域是该行全部 Q 个有效元素;单个 Matmul N tile 或一个量化 chunk 不能独立计算 scale 后冒充整行量化。物理行距可以大于 Q,但 padding 不属于逻辑统计域。
21+ 
22+## 数值合同
23+ 
24+对约定 `a=max_j(abs(W[r,j]))`、`s=max(a,scaleMin)/quantMax` 的对称量化,不能改为 `s=max(a/quantMax,scaleMin)`。这两种 clamp 的位置和单位不同。量化整数范围、舍入规则、零行以及非有限数的处理需要来自实际算子定义;不能从 INT8 dtype 推出统一的 `[-127,127]` 或 `[-128,127]` 合同。
25+ 
26+整行装入 UB 时,应同时计入原值、绝对值/归约临时区、转换结果、输出和异步在途槽位。整行不能常驻时,可第一遍按 chunk 求全行 absmax,得到最终 scale 后第二遍重读原值,所有 chunk 使用同一 scale。两遍读取要求输入内容在该期间保持不变。已有 [DynamicQuant A2/A3 实现](../../operators/aclnn/nn/quant/dynamic_quant.md)记录了完整行与两遍扫描的具体实现,本卡补充融合生产者的完成范围。
27+ 
28+## 融合与完成范围
29+ 
30+Matmul/Elementwise/GLU 的各输出 tile 可能由不同任务生成。量化某行之前,必须保证该行所有有效列均已完成最终写入;进入下一阶段不能只等待当前 tile。若存在 SplitK,完整 K 部分和还必须先合并。
31+ 
32+scale 写出完成与整行量化结果写出完成是不同事件。UB 输出槽位要等搬出结束后才能复用;workspace 要等所有读取者完成后才能覆写。只有证明两个阶段不重叠并完成必要资源重初始化,才可按两个阶段预算的最大值分配,而非把声明在不同函数当作内存已经释放。
33+ 
34+## 验证与边界
35+ 
36+覆盖 Q 跨越多个 N tile/chunk、物理行距大于 Q、零值行、尾行、正负数及舍入半整数附近。scale 相同只说明行最大值相同,不能排除其他列的 workspace 值出错。应分别比较失败列的设备中间值、设备 scale 和相同舍入链;浮点 ULP 差异在量化阈值附近可能改变整数结果。
37+ 
38+本卡来自固定版本技术说明,未执行设备验证,不给出统一的 Blaze 模板参数、UB 容量或误差放宽阈值。
@@ -0,0 +1,44 @@
1+---
2+type: Guide
3+title: GLU 与逐元素输出接整行量化的联合约束
4+description: 保留 BMM/Grouped 行映射、paired metadata、单两遍选择与执行拓扑边界。
5+tags:
6+- task.design
7+- topic.compute.quantization
8+- topic.memory
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-blaze-best-practice/references/scenarios/cube-matmul-then-vector-glu/cube-matmul-then-vector-glu-design.md
12+ role: primary
13+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-blaze-best-practice/references/scenarios/cube-matmul-then-vector-glu/cube-matmul-then-vector-glu-development.md
14+ role: supporting
15+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-blaze-best-practice/references/scenarios/cube-matmul-then-vector-glu-then-vector-pertoken-quant/cube-matmul-then-vector-glu-then-vector-pertoken-quant-design.md
16+ role: supporting
17+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-blaze-best-practice/references/scenarios/cube-matmul-then-vector-glu-then-vector-pertoken-quant/cube-matmul-then-vector-glu-then-vector-pertoken-quant-development.md
18+ role: supporting
19+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-blaze-best-practice/references/scenarios/cube-matmul-then-vector-pertoken-quant/cube-matmul-then-vector-pertoken-quant-design.md
20+ role: supporting
21+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-blaze-best-practice/references/scenarios/cube-matmul-then-vector-pertoken-quant/cube-matmul-then-vector-pertoken-quant-development.md
22+ role: supporting
23+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-blaze-best-practice/references/scenarios/cube-matmul-then-vector-elementwise-then-vector-pertoken-quant/cube-matmul-then-vector-elementwise-then-vector-pertoken-quant-design.md
24+ role: supporting
25+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-blaze-best-practice/references/scenarios/cube-matmul-then-vector-elementwise-then-vector-pertoken-quant/cube-matmul-then-vector-elementwise-then-vector-pertoken-quant-development.md
26+ role: supporting
27+created_at: '2026-09-11T03:01:48Z'
28+updated_at: '2026-09-11T03:01:48Z'
29+platforms: []
30+---
31+ 
32+# GLU 与逐元素输出接整行量化的联合约束
33+ 
34+Matmul直接量化、Matmul→GLU→量化、Matmul→逐元素DAG→量化的量化输入不同。直接路径消费最终FP Matmul结果;GLU消费完整K的两路投影并产[R,Q],只有两路等宽拼成2Q时Q=N/2;同位置逐元素DAG保持[R,N]。Softmax/LayerNorm含跨元素归约,不能当作普通unary后处理。bias、scale、dequant及其顺序只在原数学合同要求时出现。
35+ 
36+GLU的N轴weight/bias/scale按同一qOffset/curQ成对切分,M/K轴metadata不随两分支自动复制。小Q时两路独立padding,gate起点由物理act跨度决定;外部逻辑shape、mask和输出仍保留Q。BMM继续使用真实batch解码和offset,Epilogue可消费batch-local view或展开row base;GMM需要自己的group合同,不能把GMM资产强用于MM/BMM。
37+ 
38+R是实际逻辑行数:普通MM为M,等形BMM可为B*M,Grouped需由各组有效行和输出布局确定。前段列tile可以跨任务产生,同一量化行必须等全部逻辑列final写入后再分配给行消费者。GLU量化域是Q,其他两路为N;workspace dtype、物理pitch、容量及所有输出y/yScale必须一致,padding不参与absmax。
39+ 
40+能直接以Fixpipe生成最终FP workspace时,不必为了形式接量化再加identity Vector。生产者与消费者可以采用已证明的单MIX阶段协议,或在允许的实现合同中采用同stream独立Kernel顺序。GLU/elementwise组合来源指定单entry是其场景实现约束,不能推广为数学上禁止分Kernel;任何拓扑都须证明参与者、初始化、完成、数据可见性与复用,并以同设备测量评估开销。
41+ 
42+来源通用CV1/V2组合层用tuple承载两个Epilogue,前段拥有Matmul/GLU/group遍历,后段按完整行切slice;V2不能继承局部N tile当完整行。AIV-only SyncAll仅在AIC到AIV最终交接已闭合时足够,不能替换为本核PipeBarrier。静态single-pass/two-pass量化entry可由Host共享UB布局公式选择,选择器与实际分配一致,不能保留single-pass却删掉更宽行能力。
43+ 
44+宽行两遍量化第一遍全行absmax、第二遍重读用同一scale;0行、tiny、饱和、半整数舍入、absmax在首个向量/首tile之后以及容量边界前/等于/后均需覆盖。若MMAD、elementwise和最终INT8有差异,分别比较workspace、yScale和y;相同scale不证明每列原值一致。half-tie分类应使用生成参考整数输出的同一个normalized中间值,不能重新计算后误分边界。两阶段单独通过不能替代联合回归。
@@ -0,0 +1,38 @@
1+---
2+type: Guide
3+title: Blaze GLU 双分支的布局与元数据配对
4+description: 保持 act/gate 的列配对、padding 和量化元数据一致,并在完整 K 累加后执行非线性融合。
5+tags:
6+- task.design
7+- topic.data_movement.layout_transform
8+- topic.compute.quantization
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-blaze-best-practice/references/kernel-design/glu-development.md
12+ role: primary
13+created_at: '2026-09-10T14:12:04Z'
14+updated_at: '2026-09-10T14:12:04Z'
15+platforms: []
16+---
17+ 
18+# Blaze GLU 双分支的布局与元数据配对
19+ 
20+对 `act=A×B_act`、`gate=A×B_gate`、`Y=activation(act)×gate` 的 GLU 融合,两个分支在输出列 q 上有一一对应关系。只有输入合同明确把两个 Q 列分支拼成 N=2Q 时,才能用 Q=N/2;独立权重、连续拼接与交错存储需要不同地址映射。
21+ 
22+权重、bias 和沿输出 N/Q 轴变化的量化 scale 应一起按 act/gate 配对。沿 M 或 K 变化的元数据仍按各自坐标索引,不能因为 N 轴拆成两支就一律复制或偏移全部 scale。
23+ 
24+## 逻辑列与物理存储
25+ 
26+若内部存储为 `[act有效列, act填充, gate有效列, gate填充]`,gate 起点由 act 的物理跨度决定,不能使用逻辑 Q。内部 packed 宽度、外部 GM 行距、Scheduler 当前有效宽度和 Copy 特化所需基形需要分别计算。
27+ 
28+小逻辑 Q 不代表小物理资源开销:两支对齐后的 L0C/UB 暂存、元数据、slot 和搬运粒度都计入预算。MX 路径还需对齐数据与 ScaleB 的分支坐标,以及逻辑 K 与物理 K 的 padding;填充数据和 scale 的编码要以实际格式为准。
29+ 
30+## 计算与搬运边界
31+ 
32+非线性 activation 必须消费已经完成全部 K 累加的 act。对各 K 分段先激活再求和通常不等价,因此 SplitK/StreamK 的部分结果不能直接送入最终 GLU Epilogue。
33+ 
34+L0C 到 UB 的布局与传输粒度应根据实际累加结果类型和 Copy 特化确认,不能只按 A/B 的输入 dtype 推算;即使两个类型字节数相同,也不足以证明同一特化支持它们。经 GM 绕转得到正确结果,可以帮助隔离 Cube 计算与 L0C→UB 链路,但不能证明直接链路已经正确。
35+ 
36+验证应让 act/gate 权重、bias 和列 scale 具有可区分值,覆盖小 Q、非整块 Q、多 N tile、K 分段及两支物理 padding。分别比较 act、gate 与融合结果,避免只观察最终误差而把布局问题归因于 activation 近似。
37+ 
38+本卡保留双分支合同与验证方法,不把来源中的具体模板名、固定 tile、局部近零近似或修正 kernel 当作已验证的通用实现。
@@ -0,0 +1,30 @@
1+---
2+type: Guide
3+title: Grouped Matmul 的行编码与逐组权重地址
4+description: 补齐 cumsum/count、空组、完整 Golden 写回及独立 group/batch 物理映射。
5+tags:
6+- task.design
7+- topic.data_movement.layout_transform
8+- topic.implementation.operator_registration
9+status: draft
10+sources:
11+- resource: https://gitcode.com/cann/cannbot-skills/blob/502e3ed33f1440aa01bdc5a5bd396e4bf539c9ea/ops/ascendc-blaze-best-practice/references/kernel-design/group-matmul-delta.md
12+ role: primary
13+created_at: '2026-09-11T03:07:47Z'
14+updated_at: '2026-09-11T03:07:47Z'
15+platforms: []
16+---
17+ 
18+# Grouped Matmul 的行编码与逐组权重地址
19+ 
20+来源GMM以group i的半开行区间[start_i,end_i)读取共同A的不同行,使用自己的B_i,写回C对应行。其groupListType=0表示累计端点:start取前一端点、end取当前值、totalM取最后端点;1表示count:start是前缀和、end=start+count、totalM为总和。这是参考ABI枚举,不自动适用所有Grouped算子。
21+ 
22+正维执行需group数和数组长度一致,cumsum非负非递减且末值=M,count非负且求和=M,前缀和/地址算术不溢出。重复端点或零count表示空组,跳过计算但保留后续prefix与专家索引;不能静默排序错误端点,也不能对cumsum再次求和。
23+ 
24+M=0/E=0仅在算子合同允许时走单独空路径,不解引用不存在的group buffer;仍要校验该空路径规定的metadata、shape和输出。合法空组不意味着任意E=0都能进入正常tiler。全零数值行仍是实际工作,不能按空Tensor跳过。
25+ 
26+连续ND基线中A组基址为A+start_i*lda,B组基址为B+i*weightGroupStride,C组基址为C+start_i*ldc,偏移按对应Tensor元素解释;字节接口只转换一次。转置、NZ、独立Tensor列表、低比特packed或paired权重各自推导物理地址,group编码本身不定义stride。metadata改变后必须同步Copy真正消费的pitch。
27+ 
28+Golden要覆盖完整M行,可先分配完整mm[M,N]并逐非空组填对应slice,再做后处理;也可对行独立后处理逐组写最终对应slice。不能循环覆盖一个局部mm后在循环外只量化最后一组。GLU用每组两路权重并写对应M区间,未参与数学合同的x_scale/weight_scale不能因旧签名存在而加入。
29+ 
30+测试首/中/末及连续空组、不均匀group、专家各不相同、前缀非零、多M tile与非法编码拒绝;分别比较每组C及融合结果,避免所有组误用同一B却被常量数据掩盖。来源设备507015是具体故障,不能作为逻辑K误写baseK的唯一识别信号。