已开启
feat(triton): add bwd_dtrap_ddt_kernel from mamba to mindspeed ops #91
feat(triton): add bwd_dtrap_ddt_kernel from mamba to mindspeed ops #91
已开启
Tsuki创建于 7月2日
Tsuki
Tsuki成员
7月2日

ut测试结果:
f022e151-1a21-441c-9fc9-81e9ec6119f5.png
atk测试结果:
23c176a2-5dbc-4b2b-aef7-e93f463a680f.png
3307076d-1fea-4546-ab4c-56bca6cc85aa.png

UT 单元测试

  • 命令: pytest tests/unit_tests/triton/test_bwd_dtrap_ddt.py -v
  • 72/72 全部通过 ,耗时 7.39s
  • 覆盖 6 shapes × 4 chunk_sizes × 3 dtypes

ATK 精度测试

  • 命令: atk node --backend triton --devices 0 node --backend cpu --devices 0 task -c result/.../all_bwd_dtrap_ddt.json --task accuracy -p triton_bwd_dtrap_ddt.py
  • 60/60 全部通过 , acc_pass_result:Pass
  • 覆盖 fp32 / fp16 / bf16 各 20 用例

910B2 原生测试数据

Shape (B,H,S) dtype VEC 耗时 (µs) Triton 耗时 (µs) 原生加速比
(1,1,8) fp32 316.6 77.1 4.11x
(2,4,64) fp32 306.3 78.5 3.90x
(4,32,512) fp32 270.2 115.4 2.34x
平均 - 323.6 84.7 4.22x

A100 折算验证

指标 数值
平均性能比 3.722
最差 Case 性能比 2.068
验收阈值 0.795
达标判定 PASS

结论:算子在 910B2 上的表现折算至 A100 后,远超 VV 类 0.9x 验收标准。

创建的文件

  1. 内核实现 - bwd_dtrap_ddt.py
  2. API包装器 - bwd_dtrap_ddt.py
  3. 单元测试 - test_bwd_dtrap_ddt.py

测试结果

所有 72 个测试用例全部通过,覆盖 float32/bfloat16/float16 三种数据类型和多种形状。

likedislike
合并受阻
ascend-robot
ascend-robot成员7月2日进行代码检视1
mindspeed_ops/arch32/triton/bwd_dtrap_ddt.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+7

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

likedislike
ascend-robot
ascend-robot成员7月2日进行代码检视2
mindspeed_ops/arch32/triton/bwd_dtrap_ddt.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
System
系统消息系统
7月11日 评论:

changed this line on fb07503e view diff detail

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

此条代码评论区间+4+12

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

likedislike
ascend-robotascend-robot成员
7月2日 添加了label:ascend-cla/yes
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: bigdog1206, 孙银磊, 雷珍珍, 郭志华, cjy840282) ❌ (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
此处折叠了105条消息 查看更多
ascend-robot
ascend-robot成员
8月18日 评论:
流水线 PR-pipeline_MindSpeed-Ops#607 [ commitID:b4f92f01 ] 已完成
>>>代码风格自动修复执行成功(无修复内容)
阶段 任务名 状态 详情
恶意代码检查 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成员
12 天前 关联了issue:【Feature】triton算子(bwd_dtrap_ddt_kernel)迁移mindspeed-ops仓
SeaAstersSeaAsters成员
12 天前 删除了关联的issue:【Feature】triton算子(bwd_dtrap_ddt_kernel)迁移mindspeed-ops仓
SeaAstersSeaAsters成员
12 天前 关联了issue:【Feature】triton算子(bwd_dtrap_ddt_kernel)迁移mindspeed-ops仓
LinShua成员7 天前进行代码检视1
mindspeed_ops/api/triton/bwd_dtrap_ddt.py
@@ -0,0 +1,2 @@
1+# Copyright (c) 2023, Tri Dao, Albert Gu
2+# Copyright (c) 2026, Huawei Technologies Co., Ltd.
3+#
LinShua7 天前评论:

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

likedislike