已开启
【Feature】triton算子(chunk_gla_fwd_o_gk)迁移mindspeed-ops仓 #7
LinShua创建于  5月18日
LinShua成员
5月18日 创建

任务描述

基于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为例说明计算具体的性能对比要求:

  • 1.先计算当前运行NPU设备与对应竞品GPU的算力换算比例
硬件规格 910B2 A100 910B2/A100
Cube算力 354 312 1.13x
Vector算力 11.06 19.5 0.56x
带宽 1800 2039 0.88x
  • 2.根据算力换算比例计算对应算子类在不同bound场景的实际性能要求

910B2 VS A100

bound场景 CV类算子(0.7X) VV类算子(0.9x)
Cube算力bound场景 0.8x 1.0X
Vector算力bound场景 0.4x 0.5X
Memory bound场景(HBM) 0.6x 0.8x

计算式子:以CV类算子的Cube算力bound场景为例:1.13 * 0.7 =0.8
2)实践文档:
triton算子介绍1篇
3)任务完成标准
本次任务完成标准为:
精度/性能(根据实际要求)达标,PR完成合入,实践文档提交到仓库issue。

PR 合入

本地完成测试验证后,向MindSpeed-Ops的master分支发起PR。

对接人

LinSHua

欢迎加入社区,感谢您对社区的贡献 🎉!

likedislike
LLinShua成员
5月18日 修改标题为 “【Feature】triton算子(chunk_gla_fwd_o_gk)迁移mindspeed-ops仓”,原标题为“【Feature】triton算子迁移mindspeed-ops仓”
LLinShua成员
6月4日 添加了label:feature
Ronald1995成员
6月11日 评论:

/label add good-first-issue

likedislike
ascend-robotascend-robot成员
6月11日 添加了label:good-first-issue
wangx700
wangx700成员
6月12日 评论:

/label add triage-review

likedislike
ascend-robotascend-robot成员
6月12日 添加了label:triage-review
gcw_Ii4nWh2d
6月25日 评论:

接取这个任务

likedislike
Ronald1995成员
6月25日 评论:

接取这个任务

@gcw_Ii4nWh2d

欢迎认领任务,请参考https://gitcode.com/Ascend/MindSpeed-Ops/issues/3 社区任务池明确该任务的:

  • 完成的截止日期
  • 开发进展反馈
  • 微信答疑群
  • 任务交付注意事项

等信息。如果您同时认领了多项任务,但无法都能进行投入,可以在部分任务中回复退出.

麻烦您加入到对应微信群,群备注名修改为"社区任务+您的gitcode账号", 后续有相关消息和问题都可以在微信群咨询答疑。 等您加入到微信群后,我这边会在社区任务池里面登记任务责任人。

likedislike
gcw_MTeIQk9o
6月29日 评论:

认领这个任务

likedislike
yuhanBai成员
6月30日 评论:

接取这个任务

@gcw_Ii4nWh2d

你好,感谢您对昇腾平台的支持,请及时在本issue评论区更新进展,超过一周无进展更新自动视为放弃该任务。:)

likedislike
yuhanBai成员
6月30日 评论:

认领这个任务

@gcw_MTeIQk9o

你好,感谢您对昇腾平台的支持,本任务已有责任人认领,请关注仓上其他开放任务。

likedislike
gcw_Ii4nWh2d
7月1日 评论:

正在设计优化方案

likedislike
zeshengzongzeshengzong成员
7月2日 关联了看板:@zeshengzong的看板 20260702
Ggcw_Ii4nWh2d
7月2日 关联了pull request:feat(triton): add chunk gla fwd o gk kernel
ascend-robotascend-robot成员
7月3日 关联了看板:MindStudio ISSUE管理
gcw_Ii4nWh2d
7月6日 评论:

已经提交pr

likedislike
Lliuzhexu成员
7月11日 关联了里程碑:MindSpeed 26.2.0
gcw_K5CvmS79成员
7月23日 评论:

认领这个任务

likedislike
bbl0812
8月4日 评论:

认领这个任务

likedislike
LLinShua成员
8月7日 修改了issue 的描述
梵高的呐喊
梵高的呐喊
9月17日 评论:

认领该任务

likedislike
Xxmz成员
9月23日 关联了里程碑:MindSpeed 26.3.0
gcw_Ii4nWh2d
18 天前 评论:

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 通过截图

chunk_gla_fwd_o_gk ATK validation

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 通过截图

chunk_gla_fwd_o_gk UT validation

仓库原生 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
组合,按以下步骤计算:

  1. 分 K tile 读取 q、g 和 h,计算 (q * exp_or_exp2(g)) @ h,其中 scale 由
    按算子声明的 scale 与输入 dtype 三选一的载载路径引入(见「迁移要点」第 1 条)。
  2. 读取当前 chunk 的 A 和 v,通过下三角 mask 计算 tril(A) @ v。
  3. 将两部分结果在 fp32 accumulator 中相加,再转换为输出 dtype 写回。
  4. 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 做了以下调整:

  1. scale 不再在 half-precision 上预先融合(那样会把已缩放的 qg 操作数舍入两次后送进
    cube)。改为按算子声明的 scale 与输入 dtype三选一:|scale| > 1 走累加后置乘
    (避免折叠后的操作数越出半精度范围);float16 输入用 qg 的 hi/lo 半精度对驱动两次
    tl.dot(22 位有效位);bfloat16 输入把缩放后的 qg 留在 fp32 并用 fp32 cube 操作数
    (8 位尾数在两项半精度拆分里只剩 16 位)。
  2. 根据 shape metadata 增加 H=1 和 B=1/K=64/V<=64 两个实现分支。
  3. 使用确定性的 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。
  4. 保留原始来源版权声明;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 贡献的输出阶段。

注意事项

  1. 当前仅支持 arch32 NPU;arch35 明确报不支持。
  2. q/v 仅支持 float16 或 bfloat16,g/A/h 必须为 float32。
  3. chunk_size 只接受非布尔整数 32 或 64,use_exp2 只接受布尔值。
  4. UT/ATK reference 只用于测试,生产路径不存在 Torch、CPU、reference、
    供应商整算子或其它后端 fallback。
likedislike