/label add good-first-issue


/label add triage-review


接取这个任务


接取这个任务
欢迎认领任务,请参考https://gitcode.com/Ascend/MindSpeed-Ops/issues/3 社区任务池明确该任务的:
- 完成的截止日期
- 开发进展反馈
- 微信答疑群
- 任务交付注意事项
等信息。如果您同时认领了多项任务,但无法都能进行投入,可以在部分任务中回复退出.
麻烦您加入到对应微信群,群备注名修改为"社区任务+您的gitcode账号", 后续有相关消息和问题都可以在微信群咨询答疑。 等您加入到微信群后,我这边会在社区任务池里面登记任务责任人。


认领这个任务






正在设计优化方案


已经提交pr


认领这个任务


认领该任务


chunk_gla_fwd_o_gk:Triton 算子 NPU 迁移实践与交付(issue 提交)
交付 MR:Ascend/MindSpeed-Ops !92 — head 17b843e968ab05586b95db62a2c87562a0e52836,base master;当前 open,等待评审与合入。
交付内容:生产源码(mindspeed_ops/api/triton/、mindspeed_ops/arch32/triton/)+ 仓内 UT(tests/unit_tests/triton/)+ ATK 三件套(tests/atk_tests/triton/)+ 文档(docs/triton/chunk_gla_fwd_o_gk.md)。
验收证据(四项)
针对 review comment #192727253:
需要补充以下信息:
(1)ATK通过截图
(2)UT通过截图
(3)NPU与GPU性能对比达标数据
(4)NPU上的triton算子与GPU上的triton算子的精度对比情况或者GPU上该pytorch算子与triton算子的精度对齐情况
已在远端 head 17b843e968ab05586b95db62a2c87562a0e52836(本 MR 当前 head,与复验工作树逐字节一致)上逐项复验,证据如下。
(1)ATK 通过截图

ATK 26.9.24,物理 NPU 6(Ascend910_9382 physical=6/card=3/chip=0),40 例:
| 检查 | 结果 |
|---|---|
| accuracy | 40/40 success,failed 0,acc_pass_result: Pass |
| performance_device | 40/40 success,device 通过率 100.0%,平均 device 性能比 44.5103x,is_device_perf_pass: Pass |
| memory_device | 40/40 success,内存通过率 30.0%,is_device_memory_pass: Failed |
说明:memory_device 已定量归因为框架启动 workspace —— Triton-Ascend 驱动按 workspace_size * grid 每次 launch 分配,编译 workspace_size 22016–81920B;候选/参考内存比随 grid 单调上升(grid≤2 为 0.85–1.02、grid=4 为 1.03–1.11、grid 6..64 为 1.15–1.44),不是候选可移除缓冲;ATK 三项均不在合并流水线门禁内(流水线只跑仓库原生 UT)。
(2)UT 通过截图

仓库原生 tests/unit_tests/triton/test_chunk_gla_fwd_o_gk.py:19 passed, 1 warning in 3.71s(物理 NPU 6,torch_npu 先于 triton 导入;即仓库 CI stage3 的合并门禁)。
(3)NPU 与 GPU 性能对比达标数据
口径:同一评测记录、同一窗口背靠背;任务书“相对原 Triton 算子提升 20%”的判据分母是 NPU 侧 triton_ascend_reference。
- 窗口:候选
eval_20260929_080625,参考刷新baseline_20260929_075143(warmup 3 / repeat 5)。 - geomean(原 triton(NPU) / 候选) = 1.4754,min = 1.0648,≥1.2 覆盖 18/20。
- 与 GPU(H100
triton_cuda_reference) 的绝对对照:geomean(候选 / H100) = 1.8164(候选比 H100 慢约 1.8×,为 NPU/H100 硬件代差,不是任务书判据)。
| case | 候选 NPU (us) | 原 triton (NPU, us) | 比值(原/候选) | ≥1.2 | H100 triton (us) | 候选/H100 |
|---|---|---|---|---|---|---|
| 1 | 3.24 | 5.40 | 1.6667 | yes | 2.17 | 1.4931 |
| 2 | 3.88 | 6.00 | 1.5464 | yes | 2.30 | 1.6870 |
| 3 | 4.84 | 6.66 | 1.3760 | yes | 2.18 | 2.2202 |
| 4 | 4.06 | 7.32 | 1.8030 | yes | 3.30 | 1.2303 |
| 5 | 4.50 | 7.42 | 1.6489 | yes | 3.30 | 1.3636 |
| 6 | 4.84 | 7.90 | 1.6322 | yes | 3.36 | 1.4405 |
| 7 | 3.70 | 6.80 | 1.8378 | yes | 2.85 | 1.2982 |
| 8 | 4.34 | 6.42 | 1.4793 | yes | 3.52 | 1.2330 |
| 9 | 3.82 | 5.74 | 1.5026 | yes | 2.24 | 1.7054 |
| 10 | 4.58 | 7.70 | 1.6812 | yes | 3.46 | 1.3237 |
| 11 | 6.50 | 8.28 | 1.2738 | yes | 4.19 | 1.5513 |
| 12 | 4.28 | 6.16 | 1.4393 | yes | 3.33 | 1.2853 |
| 13 | 9.88 | 10.52 | 1.0648 | no | 3.55 | 2.7831 |
| 14 | 7.08 | 9.40 | 1.3277 | yes | 4.35 | 1.6276 |
| 15 | 4.62 | 7.70 | 1.6667 | yes | 3.30 | 1.4000 |
| 16 | 20.94 | 27.22 | 1.2999 | yes | 3.62 | 5.7845 |
| 17 | 4.46 | 5.94 | 1.3318 | yes | 2.62 | 1.7023 |
| 18 | 15.40 | 16.94 | 1.1000 | no | 3.46 | 4.4509 |
| 19 | 12.24 | 17.42 | 1.4232 | yes | 3.71 | 3.2992 |
| 20 | 4.98 | 8.60 | 1.7269 | yes | 2.91 | 1.7113 |
(4)精度对齐
以 fp64 CPU golden(torch) 为共同高精度基准(独立于 NPU 与 GPU 两条执行路径):
- NPU 候选(本 MR 算子):20/20 通过 domain-aware 精度检查。
- GPU(H100) triton 参考:15/20 通过;未通过的 5 例为 checker 的 special-domain(相消)判定,
normal_mismatch_count=0,无正常值失配。 - 逐 case 相对该 golden 的
max_diff/mere两侧一致(下表):NPU 候选与 GPU triton 对同一高精度基准的误差同量级。 - 另:ATK
accuracy(NPU 候选 vs 独立 NPU torch 小算子参考)40/40 Pass。
| case | NPU 候选 max_diff | NPU mere | GPU(H100) triton max_diff | GPU mere | NPU 判定 | GPU checker |
|---|---|---|---|---|---|---|
| 1 | 0.000244140625 | 2.113e-04 | 0.000244140625 | 2.113e-04 | domain_aware_pass | Pass |
| 2 | 0.000244140625 | 2.327e-04 | 0.000244140625 | 2.329e-04 | domain_aware_pass | Pass |
| 3 | 0.0001220703125 | 1.766e-04 | 0.0001220703125 | 1.766e-04 | domain_aware_pass | Pass |
| 4 | 0.00048828125 | 2.427e-04 | 0.00048828125 | 2.426e-04 | domain_aware_pass | Pass |
| 5 | 0.0001220703125 | 1.468e-04 | 0.0001220703125 | 1.469e-04 | domain_aware_pass | checker-flag |
| 6 | 0.0001220703125 | 1.723e-04 | 0.0001220703125 | 1.723e-04 | domain_aware_pass | checker-flag |
| 7 | 0.0009765625 | 5.802e-04 | 0.0009765625 | 5.802e-04 | domain_aware_pass | Pass |
| 8 | 0.0009765625 | 1.024e-03 | 0.0009765625 | 1.024e-03 | domain_aware_pass | Pass |
| 9 | 0.0001220703125 | 1.820e-04 | 0.0001220703125 | 1.813e-04 | domain_aware_pass | Pass |
| 10 | 0.0001220703125 | 1.523e-04 | 0.0001220703125 | 1.524e-04 | domain_aware_pass | Pass |
| 11 | 0.0001220703125 | 1.443e-04 | 0.0001220703125 | 1.439e-04 | domain_aware_pass | Pass |
| 12 | 0.00048828125 | 0.000e+00 | 0.00048828125 | 0.000e+00 | domain_aware_pass | Pass |
| 13 | 0.0001220703125 | 1.550e-04 | 0.0001220703125 | 1.550e-04 | domain_aware_pass | checker-flag |
| 14 | 0.000244140625 | 1.541e-04 | 0.000244140625 | 1.540e-04 | domain_aware_pass | Pass |
| 15 | 0.0009765625 | 5.936e-04 | 0.0009765625 | 5.936e-04 | domain_aware_pass | Pass |
| 16 | 0.0009765625 | 5.960e-04 | 0.0009765625 | 5.960e-04 | domain_aware_pass | Pass |
| 17 | 6.103515625e-05 | 1.448e-04 | 6.103515625e-05 | 1.447e-04 | domain_aware_pass | Pass |
| 18 | 0.0001220703125 | 1.565e-04 | 0.0001220703125 | 1.565e-04 | domain_aware_pass | checker-flag |
| 19 | 0.001953125 | 5.422e-04 | 0.001953125 | 5.422e-04 | domain_aware_pass | Pass |
| 20 | 0.0001220703125 | 1.323e-04 | 0.0001220703125 | 1.323e-04 | domain_aware_pass | checker-flag |
证据:截图存放在提交者 fork 的独立 evidence/pr92-chunk-gla-atk 分支(commit e15f671a8a2390681741e72c0992470092f4d179),不属于本 PR 的源码分支与交付文件 diff;ATK/UT 生成物与日志未提交到算子目录。
实践文档(PR 内 docs/triton/chunk_gla_fwd_o_gk.md 全文)
chunk_gla_fwd_o_gk 算子
概述
chunk_gla_fwd_o_gk 是 GLA(Gated Linear Attention)chunk forward 的输出阶段。
算子将 recurrent state 项与 chunk 内因果 attention 项相加,输出
o[B, T, H, V]:
o = scale * (q * exp_or_exp2(g)) @ h + tril(A) @ v
当前实现面向 arch32 NPU,使用 Triton-Ascend kernel。arch35 暂未提供
实现,调用时会明确抛出 NotImplementedError。
模型接口
def chunk_gla_fwd_o_gk(
q: torch.Tensor,
v: torch.Tensor,
g: torch.Tensor,
A: torch.Tensor,
h: torch.Tensor,
scale: float = 0.125,
chunk_size: int = 64,
use_exp2: bool = False,
) -> torch.Tensor:
"""Run the dense GLA chunk forward output stage."""
输入说明
| 参数 | 类型 | shape | 数据类型 | 说明 |
|---|---|---|---|---|
q |
torch.Tensor |
[B, T, H, K] |
float16、bfloat16 |
query |
v |
torch.Tensor |
[B, T, H, V] |
与 q 相同 |
value |
g |
torch.Tensor |
[B, T, H, K] |
float32 |
cumulative gate |
A |
torch.Tensor |
[B, T, H, C] |
float32 |
chunk 内 attention,C=chunk_size |
h |
torch.Tensor |
[B, ceil(T/C), H, K, V] |
float32 |
recurrent state |
scale |
float |
标量 | float |
recurrent state 项缩放系数,默认 0.125 |
chunk_size |
int |
标量 | int |
支持 32 或 64,默认 64 |
use_exp2 |
bool |
标量 | bool |
使用 exp2(g);否则使用 exp(g) |
输入必须位于 NPU,q 与 v 的 dtype 必须一致。当前公开接口只支持 dense、
non-varlen 路径;varlen metadata 不属于本次交付范围。
输出说明
| 输出 | shape | 数据类型 | 说明 |
|---|---|---|---|
o |
[B, T, H, V] |
与 v 相同 |
recurrent state 项与 chunk 内因果项之和 |
实现原理
每个 Triton program 负责一个 V tile、一个 sequence chunk 和一个 batch/head
组合,按以下步骤计算:
- 分 K tile 读取
q、g和h,计算(q * exp_or_exp2(g)) @ h,其中scale由
按算子声明的 scale 与输入 dtype 三选一的载载路径引入(见「迁移要点」第 1 条)。 - 读取当前 chunk 的
A和v,通过下三角 mask 计算tril(A) @ v。 - 将两部分结果在 fp32 accumulator 中相加,再转换为输出 dtype 写回。
- sequence、K 和 V 尾块均使用边界检查或显式 mask。
生产 dispatch 只依赖运行时 tensor metadata:
| 条件 | kernel | 目的 |
|---|---|---|
B == 1, C == 64, K == 64, V <= 64, H >= 4 |
chunk_gla_fwd_kernel_o_b1_blockptr_local_scale_qg |
用 block pointer 优化局部 A/V/o 访问 |
H == 1 且 C == 64 或输入仅一个 chunk |
chunk_gla_fwd_kernel_o_h1_scale_qg |
简化单 head 的 grid 和地址计算 |
| 其它合法 dense 输入 | chunk_gla_fwd_kernel_o_scale_qg |
覆盖通用 shape、dtype 和尾块 |
迁移要点
原始 Triton 实现使用一个通用 chunk_gla_fwd_kernel_o,在同一 kernel 中完成
recurrent state 和 chunk-local 两部分计算,并通过 Triton autotune 选择 K/V tile。
本 PR 保持相同数学语义,针对 Ascend 做了以下调整:
scale不再在 half-precision 上预先融合(那样会把已缩放的qg操作数舍入两次后送进
cube)。改为按算子声明的scale与输入 dtype三选一:|scale| > 1走累加后置乘
(避免折叠后的操作数越出半精度范围);float16 输入用qg的 hi/lo 半精度对驱动两次
tl.dot(22 位有效位);bfloat16 输入把缩放后的qg留在 fp32 并用 fp32 cube 操作数
(8 位尾数在两项半精度拆分里只剩 16 位)。- 根据 shape metadata 增加 H=1 和 B=1/K=64/V<=64 两个实现分支。
- 使用确定性的
BK、BV和 multibuffer 配置,不依赖运行时 autotune 缓存。BK/BV
只按运行时 K/V 范围选择:K<=32用 32-lane、K为 128 的倍数时用 128-lane(把
reduction 收进单次迭代)、其余 64-lane;V轴在V > 64 且 32|V 且 128∤V时用
32-lane(该 tile 无 padding),否则V > 64用 128-lane、其余 64-lane。 - 保留原始来源版权声明;UT 使用独立 Torch 小算子 reference 验证数学语义。
生产 API 只启动本 PR 自有 Triton-Ascend kernel,不导入测试 reference,不调用
Torch、CPU、供应商整算子或其它后端 fallback。
使用示例
import math
import torch
from mindspeed_ops.api.triton.chunk_gla_fwd_o_gk import chunk_gla_fwd_o_gk
B, T, H, K, V, C = 1, 128, 4, 64, 64, 64
q = torch.randn((B, T, H, K), device="npu", dtype=torch.float16)
v = torch.randn((B, T, H, V), device="npu", dtype=torch.float16)
g = torch.randn((B, T, H, K), device="npu", dtype=torch.float32) * 0.08
A = torch.randn((B, T, H, C), device="npu", dtype=torch.float32) * 0.08
h = torch.randn((B, math.ceil(T / C), H, K, V), device="npu", dtype=torch.float32) * 0.08
o = chunk_gla_fwd_o_gk(q, v, g, A, h, scale=K**-0.5, chunk_size=C)
print(o.shape)
适用场景
适用于 GLA dense、non-varlen forward 中,按固定 chunk 合并 recurrent state
贡献和 chunk 内因果 attention 贡献的输出阶段。
注意事项
- 当前仅支持
arch32NPU;arch35明确报不支持。 q/v仅支持float16或bfloat16,g/A/h必须为float32。chunk_size只接受非布尔整数32或64,use_exp2只接受布尔值。- UT/ATK reference 只用于测试,生产路径不存在 Torch、CPU、reference、
供应商整算子或其它后端 fallback。


任务描述
基于fla-org/flash-linear-attention开源仓库的triton源码,迁移优化triton算子在NPU上跑通及性能优化。
本期任务的具体信息如下:
(1)算子名称:chunk_gla_fwd_o_gk
(2) fla-org/flash-linear-attention开放仓库的算子源码链接:
https://github.com/fla-org/flash-linear-attention/blob/v0.4.1/fla/ops/gla/chunk.py
(3) 迁移指导文档:https://gitcode.com/Ascend/triton-ascend/blob/main/docs/zh/programming_guide.md
参考 skill 地址:https://gitcode.com/Ascend/agent-skills/tree/master/skills/simple-vector-triton-gpu-to-npu
(4) 开发平台:Atlas 800T A2或者A3
验收标准
一、任务交付件
本期任务为基于fla-org/flash-linear-attention开放仓库代码进行功能扩展,请合入开发代码。主要开发点如下:
(1) triton算子针对NPU的适配修改优化后的代码
(2) 针对该triton算子的测试用例UT
二、验收标准:
1)精度/性能要求
精度:
对比GPU相同输入算子精度误差小于0.1%
若无GPU标杆,与CPU小算子对齐,精度误差小于0.1%
性能:
补充说明:具体性能对比数据呈现,按 CV类算子达到0.7x竞品, VV类算子达到0.9x竞品
以下以竞品A100为例说明计算具体的性能对比要求:
910B2 VS A100
计算式子:以CV类算子的Cube算力bound场景为例:1.13 * 0.7 =0.8
2)实践文档:
triton算子介绍1篇
3)任务完成标准
本次任务完成标准为:
精度/性能(根据实际要求)达标,PR完成合入,实践文档提交到仓库issue。
PR 合入
本地完成测试验证后,向MindSpeed-Ops的master分支发起PR。
对接人
LinSHua
欢迎加入社区,感谢您对社区的贡献 🎉!