master

ascend-robotfeat: add gdn ops

1a738ef5创建于 25 天前历史提交

文件	最后提交记录	最后更新时间
tilelang	feat/fix/docs/style/refactor/adaptor/chore/test(backend): refactor Table of Contents Co-authored-by: shiyuan680<yangcheng104@huawei.com>	2 个月前
triton	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	25 天前
__init__.py	adapt causal_conv1d and layer_norm_gated_bwd kernels into arch35 & add tests Co-authored-by: wanggangguo<wanggangguo@huawei.com> # message auto-generated for no-merge-commit merge: !63 merge op into master adapt causal_conv1d and layer_norm_gated_bwd kernels into arch35 & add tests Created-by: isfrapples Commit-by: wanggangguo Merged-by: ascend-robot Description: What this PR does / why we need it? Ports and optimizes causal_conv1d and rms_norm_gated Triton kernels from FLA to Ascend NPU arch35. The initial port had severe performance issues caused by @triton.autotune expanding kernel launch counts to ~300x on NPU. This PR removes autotune, adopts a single-config approach with num_warps=8, and fixes BD tile sizing, bringing the GPU/NPU performance ratio from 0.09x to 0.39x. Does this PR introduce any user-facing change? No. The public APIs (causal_conv1d, rms_norm_gated) remain unchanged. All changes are internal to the arch35 kernel implementations and test files. Unit tests for both test_causal_conv1d and test_rmsnormgated are now routed to arch35 implementations when running on arch35 hardware. How was this patch tested? Tested via performance benchmarking (benchmark_tool.py with msprof profiling) on 10 cases covering B×T×D combinations from (1×1024×64) to (32×8192×128) and (2×4096×1200). Average GPU/NPU ratio reached 0.39x, exceeding the 0.32x target. Accuracy validated through existing ATK test suites for conv1d, causal_conv1d_bwd, rmsnormgated, and layer_norm_gated_bwd. See merge request: Ascend/MindSpeed-Ops!63	26 天前