已开启
feat(triton): add bwd_segsum_dadt operator for Mamba backward pass #89
feat(triton): add bwd_segsum_dadt operator for Mamba backward pass #89
已开启
Tsuki创建于 7月2日
Tsuki
Tsuki成员
7月2日

This PR adds the bwd_segsum_dadt operator which computes the backward segment sum for the Mamba architecture's dA_dt computation.
ut测试:
ee233a1f-c11f-4be1-9f9b-b43ce0a8d5e4.png
atk测试:
825528fc-d365-4796-b798-545fef535d20.png
Files added:

  • mindspeed_ops/arch32/triton/bwd_segsum_dadt.py: Triton kernel implementation
  • mindspeed_ops/api/triton/bwd_segsum_dadt.py: Python API wrapper
  • tests/unit_tests/triton/test_bwd_segsum_dadt.py: Unit tests with CPU golden reference

The operator performs:

  1. Transposed loading of dSSdA matrix
  2. Element-wise multiplication with exp(SSdA)
  3. Reverse cumulative sum along axis=0
  4. Causal masking (zeroing upper triangular)
  5. Row-wise summation and accumulation into ddt_out

Supported data types: float32, float16, bfloat16

All 18 test cases passed.

likedislike
合并受阻
ascend-robot
ascend-robot成员7月2日进行代码检视1
mindspeed_ops/arch32/triton/bwd_segsum_dadt.py
@@ -0,0 +1,2 @@
1+# Copyright (c) 2026, HUAWEI CORPORATION. All rights reserved.
2+# pylint: disable=import-error
3+import triton
ascend-robot
ascend-robot7月2日评论:

此条代码评论区间+1至+4

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

likedislike
ascend-robot
ascend-robot成员7月2日进行代码检视1
mindspeed_ops/arch32/triton/bwd_segsum_dadt.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-robot7月2日评论:

此条代码评论区间+3至+7

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

likedislike
ascend-robot
ascend-robot成员7月2日进行代码检视1
mindspeed_ops/arch32/triton/bwd_segsum_dadt.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_segsum_dadt_kernel(
ascend-robot
ascend-robot7月2日评论:

此条代码评论区间+4至+8

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

likedislike
ascend-robotascend-robot成员
7月2日 添加了label:ascend-cla/no
ascend-robot
ascend-robot成员
7月2日 评论:

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


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: 周蓓蓉, iansheng, bigdog1206, 郑加利, 孙银磊) ❌ (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
此处折叠了116条消息 查看更多
ascend-robotascend-robot成员
8月18日 删除了label:ci-pipeline-running
ascend-robot
ascend-robot成员
8月18日 评论:
流水线 PR-pipeline_MindSpeed-Ops#604 [ commitID:41d4e5d2 ] 已终止运行
>>>代码风格自动修复执行成功(无修复内容)
阶段 任务名 状态 详情
恶意代码检查 Antipoison 🟨 >>>
编码安全与规范检查 Only_doc_commit ✅ >>>
pre-commit ✅ >>>
开源片段检查 SCA ✅ >>>
开发者测试 UT_MindSpeed_1 🕚 >>>
UT_MindSpeed_2 🕚 >>>
UT_MindSpeed_3 🕚 >>>
UT_MindSpeed_4 🕚 >>>
流水线 PR-pipeline_MindSpeed-Ops 🟨 >>>
此流水线已支持下列评论快捷指令,仅PR创建者和白名单成员[wujinyuan1, liuzhexu]评论有效
  • compile : 运行流水线
  • retry : 重试流水线所有失败子任务
  • retry <任务名> : 仅重试指定失败子任务
  • stop : 停止流水线
likedislike
ascend-robotascend-robot成员
8月18日 添加了label:ci-pipeline-passed
ascend-robot
ascend-robot成员
8月18日 评论:
流水线 PR-pipeline_MindSpeed-Ops#605 [ commitID:41d4e5d2 ] 已完成
>>>代码风格自动修复执行成功(无修复内容)
阶段 任务名 状态 详情
恶意代码检查 Antipoison ✅ >>>
编码安全与规范检查 Only_doc_commit ✅ >>>
pre-commit ✅ >>>
开源片段检查 SCA ✅ >>>
开发者测试 UT_MindSpeed_1 ✅ >>>
UT_MindSpeed_2 ✅ >>>
UT_MindSpeed_3 ✅ >>>
UT_MindSpeed_4 ✅ >>>
流水线 PR-pipeline_MindSpeed-Ops ✅ >>>
此流水线已支持下列评论快捷指令,仅PR创建者和白名单成员[wujinyuan1, liuzhexu]评论有效
  • compile : 运行流水线
  • retry : 重试流水线所有失败子任务
  • retry <任务名> : 仅重试指定失败子任务
  • stop : 停止流水线
likedislike
SeaAstersSeaAsters成员
9月10日 关联了issue:【Feature】triton算子(bwd_segsum_dadt_kernel)迁移mindspeed-ops仓