菜单
triton-ascend 大 K 归约精度修复：Kahan 补偿求和替代简单累加，消除 NPU Cube 引擎 FP32 仿真路径与 Triton 顺序累加路径的精度差异
用 Codex 或 Claude 帮你安装 复制这段 Prompt，粘贴到 Codex、Claude 或其他助手里，让它检查 Skill 页面并帮你完成安装。
基于 SOC 职业分类
| name | triton-ascend-kahan-precision-fix |
| description | triton-ascend 大 K 归约精度修复：Kahan 补偿求和替代简单累加，消除 NPU Cube 引擎 FP32 仿真路径与 Triton 顺序累加路径的精度差异 |
| category | fix |
| version | 1.0.0 |
| metadata | {"case_type":"fix","backend":"ascend","dsl":"triton_ascend","hardware":"Atlas A2, Atlas A3, Atlas A5"} |
# 错误：简单累加，K 大时误差 O(K × eps)
acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)
for k in range(0, K, BLOCK_K):
a = tl.load(...)
b = tl.load(...)
acc += tl.dot(a, b)
# 修复：Kahan 补偿，误差降为 O(eps)
acc = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)
comp = tl.zeros((BLOCK_M, BLOCK_N), dtype=tl.float32)
for k in range(0, K, BLOCK_K):
a = tl.load(...)
b = tl.load(...)
partial = tl.dot(a, b)
y = partial - comp
t = acc + y
comp = (t - acc) - y
acc = t
浮点加法不满足结合律：(a + b) + c ≠ a + (b + c)。当 acc 很大而 partial 很小时，acc + partial 会丢失 partial 的低位精度。Kahan 算法通过一个补偿变量 comp 追踪每次加法丢失的精度，下次累加时补回。
逐步拆解：
partial = tl.dot(a, b) # 本次 dot 结果
y = partial - comp # 减去上次丢失的精度（补偿）
t = acc + y # 累加
comp = (t - acc) - y # 捕获本次丢失的精度
acc = t # 更新累加器
y = partial - comp：把上次丢失的部分补回来t = acc + y：执行实际累加comp = (t - acc) - y：(t - acc) 是实际加入 acc 的值，减去 y 得到本次丢失的低位comp 恒为零，但浮点运算中它捕获了舍入误差| 场景 | 是否适用 |
|---|---|
| matmul K ≥ 4096 | ✅ |
| reduction（sum/mean）大维度 | ✅ |
| matmul K < 4096 | ❌ 简单累加足够 |
| elementwise / 无归约 | ❌ 无累加误差 |
AscendC direct-invoke 崩溃/挂起修复索引：Kernel timeout、hang、Segmentation Fault、aic error、buffer 死锁、plog/memcheck 调试。
AscendC direct-invoke 精度失败修复索引：输出全 0/随机值、DataCopy 对齐、EnQue/DeQue 同步、FP16/FP32 精度、Cast RoundMode、DumpTensor 分段定位。
AscendC direct-invoke 工程契约：WA 使用 kernel.py + ascendc_op/，ModelNew 调用 torch.ops.npu.*，adapter 负责复制工程、CMake 构建和 npu-arch patch。适用于 dsl=ascendc 的 autoresearch 任务。
把注册式 AscendC 算子迁移为 direct-invoke 工程的保真原则：kernel 算法和 tiling 公式不乱改，只替换注册框架胶水、入口 ABI、host launch 与 PyTorch extension。
CATLASS TileShape 与 on-chip 缓存容量约束：L1/L0A/L0B/L0C 预算公式、fp16/fp32 Pingpong 双缓冲、512B 对齐与排布对 Tile 选型的影响。调参前必读。
CATLASS Gemm 性能调优：DispatchPolicy、Tile 与分核负载均衡、Swizzle、何时用 padding/Split-K/Preload。面向 AR 修改 catlass_kernel.asc 中的类型别名。