Open
feat(triton): add bwd_dadt_cumsum_fused_kernel operator for Mamba backward pass #90
feat(triton): add bwd_dadt_cumsum_fused_kernel operator for Mamba backward pass #90
Open
TsukiAdd on Jul 2
Tsuki
TsukiCollaborators
Jul 2

新增 bwd_dadt_cumsum_fused_kernel 算子,用于 Mamba 模型反向传播中 ddt 的融合计算。
ut:
e56a34d5-8234-441b-80bc-755faf4e1732.png
atk:
042403b0-e0eb-45f6-acc7-402e883dfe2b.png

三、910B实测性能

328918b0-2659-4c76-a658-b76cd22aa403.png

四、A100等效性能折算

f738ada5-ecf7-44b1-83b4-56902e3d8b75.png

结论

  1. 功能正确性 :UT测试全部通过,精度误差 < 1e-6
  2. 精度达标 :ATK精度评估 Pass,30个用例通过率100%
  3. 性能表现 :在910B上实现高效的cumulative sum计算
  4. A100等效性能 :折算后性能满足CV类算子要求(Cube算力bound场景下达到A100的0.8倍性能)

使用示例

from mindspeed_ops.api.triton.bwd_dadt_cumsum_fused import bwd_dadt_cumsum_fused

# 输入张量: [B, H, S]
ddt_out = bwd_dadt_cumsum_fused(ddA_cs, ddA_cs_rev, dA_cs, dA_cs_rev, ddt_out, chunk_size=64)
likedislike
Merge blocked
ascend-robot
ascend-robotCollaboratorsJul 2Start code review1
mindspeed_ops/arch32/triton/bwd_dadt_cumsum_fused.py
Outdated
@@ -0,0 +1,2 @@
1+# Copyright (c) 2026, HUAWEI CORPORATION. All rights reserved.
2+# pylint: disable=import-error
3+import triton
ascend-robot
ascend-robotJul 2Comments:

This code comment section is from+1To+3

【openlibing.ci】识别到代码检查告警抑制注释,匹配工具:pylint,请Committer检视其合理性。

likedislike
ascend-robot
ascend-robotCollaboratorsJul 2Start code review1
mindspeed_ops/arch32/triton/bwd_dadt_cumsum_fused.py
@@ -0,0 +3,4 @@
3+import triton
4+import triton.language as tl # pylint: disable=E0611
5+ 
6+ 
7+@triton.jit # pylint: disable=E1101
ascend-robot
ascend-robotJul 2Comments:

This code comment section is from+3To+8

【openlibing.ci】识别到代码检查告警抑制注释,匹配工具:pylint,请Committer检视其合理性。

likedislike
ascend-robot
ascend-robotCollaboratorsJul 2Start code review1
mindspeed_ops/arch32/triton/bwd_dadt_cumsum_fused.py
@@ -0,0 +4,4 @@
4+import triton.language as tl # pylint: disable=E0611
5+ 
6+ 
7+@triton.jit # pylint: disable=E1101
8+def bwd_dadt_cumsum_fused_kernel(
ascend-robot
ascend-robotJul 2Comments:

This code comment section is from+4To+9

【openlibing.ci】识别到代码检查告警抑制注释,匹配工具:pylint,请Committer检视其合理性。

likedislike
ascend-robotascend-robotCollaborators
Jul 2 Added label: ascend-cla/yes
ascend-robot
ascend-robotCollaborators
Jul 2 Comments:

Thanks for your pull-request.
The full list of commands accepted by me can be found at here.
You can get sig-info at here.
You can self-configure the PR merge rules for this repository. For more details, please refer to Here.


PR Approval Progress

⚠️ This PR does not yet meet the following requirements:lgtm (requires ≥ 2 person(s) per module)、approve (requires ≥ 1 person(s) per module)

Module Approval Details

module lgtm status approve status
repo-Ascend/MindSpeed-Ops ❌ (0/2)(You can also ask: guihaowen666, tichang, 周蓓蓉, 雷珍珍, HANHU1CHEN) ❌ (0/1)(You can also ask: 刘荣, 华郁秀, 周蓓蓉, 雷珍珍, 郑加利)

💡 Tip:

  • Committer can comment /approve or /lgtm
  • Commenting /approve implies both code review (lgtm) and intent to merge (approve)

CLA Signature Pass

xiaomaomao666, thanks for your pull request. All authors of the commits have signed the CLA. 👍

likedislike
114 messages are collapsed here。 See more
AtlasAccountAtlasAccountCollaborators
Sep 11 Removed label: ci-pipeline-running
AtlasAccountAtlasAccountCollaborators
Sep 11 Added label: ci-pipeline-passed
AtlasAccount
AtlasAccountCollaborators
Sep 11 Comments:
流水线 PR-pipeline_MindSpeed-Ops#717 [ commitID:7464a793 ] 已完成
>>>代码风格自动修复执行成功(无修复内容)
阶段 任务名 状态 详情
恶意代码检查 Antipoison ✅ COMPLETED >>>
编码安全与规范检查 Only_doc_commit ✅ COMPLETED >>>
pre-commit ✅ COMPLETED >>>
开源片段检查 SCA ✅ COMPLETED >>>
开发者测试 UT_MindSpeed_1 ✅ COMPLETED >>>
UT_MindSpeed_2 ✅ COMPLETED >>>
UT_MindSpeed_3 ✅ COMPLETED >>>
UT_MindSpeed_4 ✅ COMPLETED >>>
流水线 PR-pipeline_MindSpeed-Ops ✅ COMPLETED >>>
此流水线已支持下列评论快捷指令,仅PR创建者和白名单成员[wujinyuan1, liuzhexu]评论有效
  • compile : 运行流水线
  • retry : 重试流水线所有失败子任务
  • retry <任务名> : 仅重试指定失败子任务
  • stop : 停止流水线
likedislike
LinShuaCollaboratorsSep 15Start code review1
mindspeed_ops/api/triton/bwd_dadt_cumsum_fused.py
@@ -0,0 +1,2 @@
1+# Copyright (c) 2023, Tri Dao, Albert Gu
2+# Copyright (c) 2026, Huawei Technologies Co., Ltd. All rights reserved.
3+ 
LinShuaSep 15Comments:

需要补充GPU上pytorch小算子与GPU上开源triton算子的精度对齐截图(以确保pytorch小算子标杆的准确性)

likedislike
SeaAsters
SeaAstersCollaborators
4 days ago Comments:

开发者你好,参照下面的标准再补充对应的信息:
(1)ATK的精度和性能通过截图
(2)UT测试用例通过的截图
(3)triton算子在NPU和GPU上的性能对比
(4)triton算子在NPU和GPU上的精度对比(同数据输入场景)或GPU上验证pytorch小算子与GPU 开源triton算子的精度对比(以验证pytorch小算子标杆准确性)

likedislike