| 文件 | 最后提交记录 | 最后更新时间 |
|---|---|---|
[Feat] Add tilelang op sinkhorn Co-authored-by: zhuweichen<calvin_zhu0210@outlook.com> # message auto-generated for no-merge-commit merge: !73 merge sinkhorn into master [Feat] Add tilelang op sinkhorn Created-by: zhuweichen Commit-by: zhuweichen Merged-by: ascend-robot Description: ## Related Issue https://gitcode.com/Ascend/MindSpeed-Ops/issues/50 ## What this PR does / why we need it? This PR adds a forward-only TileLang implementation of the HC-Split Sinkhorn operator for Ascend arch32. Main changes: - Add mindspeed_ops.api.tilelang.sinkhorn.sinkhorn as the public Python API. - Add the arch32 TileLang kernel implementation for HC-Split Sinkhorn. - Support hc_mult=4. - Support torch.float32 and torch.bfloat16 inputs/outputs. For BF16, the kernel keeps BF16 IO and uses FP32 UB buffers for internal sigmoid, exp, reduce, and Sinkhorn normalization. - Add input validation for shape, dtype, hc_mult, sinkhorn_iters, and unsupported arch35. - Add unit tests and ATK accuracy cases for TileLang Sinkhorn. - Add operator documentation at docs/tilelang/sinkhorn.md. ## Does this PR introduce any user-facing change? Yes. This PR introduces a new TileLang Sinkhorn API: python from mindspeed_ops.api.tilelang.sinkhorn import sinkhorn Supported usage: - mixes: shape [batch_size, seq_len, (2 + hc_mult) * hc_mult] - hc_scale: shape [3] - hc_base: shape [(2 + hc_mult) * hc_mult] - hc_mult: 4 - dtype: torch.float32 or torch.bfloat16 - platform: Ascend arch32 only The operator is forward-only and does not provide autograd backward. Usage details are documented in docs/tilelang/sinkhorn.md. ## How was this patch tested? Verified on Ascend NPU environment: Ascend910_9382, CANN 9.0.0, Python 3.11, TileLang-Ascend installed. 1. Unit test bash pytest -q tests/unit_tests/tilelang/test_sinkhorn.py Result: 42 passed 2. ATK accuracy test sh cd tests/atk_tests/tilelang/sinkhorn export PYTHONPATH=/home/z30046196/MindSpeed-Ops-pr73:$PYTHONPATH atk node --backend triton --devices 0 node --backend npu --devices 0 task -c result/sinkhorn/json/all_sinkhorn.json --task accuracy -tup ./ -p tilelang_sinkhorn.py ATK accuracy log: sh +-------+----------+------------------+------------------+--------------+----------------------+--------+--------------+ | 名称 | 总用例数 | 执行成功用例个数 | 执行失败用例个数 | 通过用例个数 | 错误信息匹配用例个数 | 通过率 | 精度是否达标 | +-------+----------+------------------+------------------+--------------+----------------------+--------+--------------+ | npu_0 | 36 | 36 | 0 | 36 | 0 | 100.0 | Pass | +-------+----------+------------------+------------------+--------------+----------------------+--------+--------------+ Total Task: 36, success 36, failed 0 Summary info: acc_pass_result:Pass is_acc_dc_pass:None is_e2e_perf_pass:None is_benchmark_perf_pass:None is_device_perf_pass:None is_device_memory_pass:None is_memory_check_pass:None . atk task success! atk task finish. cost time: 9 s 3. ATK device performance test sh cd tests/atk_tests/tilelang/sinkhorn export PYTHONPATH=/home/z30046196/MindSpeed-Ops-pr73:$PYTHONPATH atk node --backend triton --devices 0 node --backend npu --devices 0 task -c result/sinkhorn/json/all_sinkhorn.json --task performance_device -tup ./ -p tilelang_sinkhorn.py ATK performance log: sh +-------+----------+------------------+------------------+--------------------+ | 名称 | 总用例数 | device性能通过率 | 平均device性能比 | device性能是否达标 | +-------+----------+------------------+------------------+--------------------+ | npu_0 | 36 | 94.4444 | 7.2842 | Pass | +-------+----------+------------------+------------------+--------------------+ Total Task: 36, success 36, failed 0 Summary info: acc_pass_result:None is_acc_dc_pass:None is_e2e_perf_pass:None is_benchmark_perf_pass:None is_device_perf_pass:Pass is_device_memory_pass:None is_memory_check_pass:None . save result excel file: /home/z30046196/MindSpeed-Ops-pr73/tests/atk_tests/tilelang/sinkhorn/atk_output/all_sinkhorn_2026-07-09-09-54-25-397923/report/ all_sinkhorn_reports_2026-07-09-09-54-26.xlsx atk task success! atk task finish. cost time: 537 s See merge request: Ascend/MindSpeed-Ops!73 | 6 天前 | |
[Bugfix] Fix chunk_kda_bwd precision Co-authored-by: zhuweichen<calvin_zhu0210@outlook.com> # message auto-generated for no-merge-commit merge: !105 merge kda_fix into master [Bugfix] Fix chunk_kda_bwd precision Created-by: zhuweichen Commit-by: zhuweichen Merged-by: ascend-robot Description: ## What this PR does / why we need it? Fix chunk_kda_bwd precision. ## Does this PR introduce any user-facing change? No. ## How was this patch tested? ATK & UT   See merge request: Ascend/MindSpeed-Ops!105 | 2 天前 | |
feat: add gdn ops Co-authored-by: sinle4cat<wangqingyu9@huawei.com> # message auto-generated for no-merge-commit merge: !62 merge gdn into master feat: add gdn ops Created-by: sinle4cat Commit-by: sinle4cat Merged-by: ascend-robot Description: ## What this PR does / why we need it? This PR adds arch35 Triton support and related test coverage for several GDN-related operators, and fixes a regression introduced during arch35 enablement. ## Does this PR introduce any user-facing change? No breaking user-facing API change is introduced. - Existing Python operator entry points remain unchanged. - chunk_scaled_dot_kkt now works on arch35 while keeping the existing arch32 behavior. - For variable-length inputs, the varlen path is only taken on arch35 when cu_seqlens is explicitly provided. - In ATK, an optional cu_seqlens input is added so varlen cases can be described explicitly. No documentation file was updated in this PR. ## How was this patch tested?       See merge request: Ascend/MindSpeed-Ops!62 | 23 天前 | |
feat: Add add_rms_norm_bias compiletion Co-authored-by: Liccol<740821011@qq.com> # message auto-generated for no-merge-commit merge: !54 merge aclnn-case into master feat: Add add_rms_norm_bias compiletion Created-by: Liccol Commit-by: Liccol Merged-by: ascend-robot Description: ## What this PR does / why we need it? 新增 add_rms_norm_bias 自定义 ACLNN 算子，将逐元素加法（Add）与 RMS 归一化合并为单个 AscendC Kernel 执行，用于 Transformer 模型推理中残差连接 + RMSNorm 的融合优化。 主要变更： - 新增 add_rms_norm_bias 算子源码 - 新增 cann/ 下的 ACLNN 绑定及 torch_binding 算子注册 - 新增 Python 调用接口（api/aclnn/）及 JIT 构建器（op_builder/） - 引入 cmake/common/scripts/utils 等编译工具链及公共组件 - 新增算子使用文档及开发指南（docs/aclnn/） ## Does this PR introduce any user-facing change? 是，新增 npu_add_rms_norm_bias 接口 ## How was this patch tested? 通过单元测试与 CPU 参考实现对比验证： - 对比分离执行 Add + RMSNorm 与融合算子的数值精度 - 覆盖 float32 / float16 / bfloat16 三种数据类型 - 覆盖有 beta / 无 beta 两种场景  See merge request: Ascend/MindSpeed-Ops!54 | 26 天前 | |
[Feature] Add Aclnn ops compilation and dependencies Co-authored-by: YuhanBai<baiyuhan1@h-partners.com> # message auto-generated for no-merge-commit merge: !52 merge aclnn-adapter into master [Feature] Add aclnn ops compiletion and dependencies Created-by: yuhanBai Commit-by: YuhanBai Merged-by: ascend-robot Description: ## What this PR does / why we need it? 本PR参考vllm-ascend中的算子工程，添加了aclnn算子的编译依赖 RFC: https://gitcode.com/Ascend/MindSpeed-Ops/issues/33 Co-Auther with: [@Liccol](https://gitcode.com/Liccol) [@qq_45570708](https://gitcode.com/qq_45570708) ## Does this PR introduce any user-facing change? Please describe whether the PR will result in any user-facing usage changes. If there is related documentation, please specify its path. ## How was this patch tested? Please explain how to verify the correctness and effectiveness of this feature, as well as its usage constraints and limitations. See merge request: Ascend/MindSpeed-Ops!52 | 1 个月前 | |
feat: prepare_wy_repr ops Co-authored-by: liuxi_<liuxi75@huawei.com> # message auto-generated for no-merge-commit merge: !20 merge prepare_wy_repr into master feat: prepare_wy_repr ops Created-by: liuxi_ Commit-by: liuxi_ Merged-by: ascend-robot Description: ## What this PR does / why we need it? prepare_wy_repr ops [#1](https://gitcode.com/Ascend/MindSpeed-Ops/issues/1) ## Does this PR introduce any user-facing change? 不涉及 ## How was this patch tested? TA 3.2.1 ut  atk   See merge request: Ascend/MindSpeed-Ops!20 | 2 个月前 |