feature(cmake): update to new Ascend CMake infra - #215
Merged
Conversation
added 5 commits
July 14, 2026 06:25
Contributor
There was a problem hiding this comment.
Pull request overview
This PR updates the Ascend kernel build/launch integration to support the newer CANN 9.x “ASC language” CMake infrastructure, replacing reliance on legacy toolchain–generated aclrtlaunch_<kernel>.h headers with explicit pto_launch_* shims compiled alongside the kernels.
Changes:
- Add host-callable
pto_launch_*launch shims in kernel sources to centralize<<< >>>launches in kernel-compiled translation units. - Switch host-side
EXEC_KERNEL_CMDto callpto_launch_<kernel>instead ofACLRT_LAUNCH_KERNEL(<kernel>), and update host headers to declare the new launch shims. - Update CMake to optionally use
find_package(ASC)/enable_language(ASC)(default ON), including SOC→arch mapping for--npu-arch.
Reviewed changes
Copilot reviewed 42 out of 42 changed files in this pull request and generated 4 comments.
Show a summary per file
| File | Description |
|---|---|
| csrc/kernel/kernel_tri_inv_trick.cpp | Adds pto_launch_tri_inv_trick_fp16 shim. |
| csrc/kernel/kernel_tri_inv_rec_unroll.cpp | Adds pto_launch_tri_inv_rec_unroll_{bf16,fp16} shims. |
| csrc/kernel/kernel_tri_inv_ns.cpp | Adds pto_launch_tri_inv_ns_fp16 shim. |
| csrc/kernel/kernel_tri_inv_col_sweep.cpp | Adds pto_launch_triv_inv_col_sweep_{fp16,fp32} shims. |
| csrc/kernel/kernel_swiglu.cpp | Adds pto_launch_swiglu_fp16 shim. |
| csrc/kernel/kernel_simple_matmul.cpp | Adds pto_launch_simple_matmul_{bf16,fp16,fp32} shims. |
| csrc/kernel/kernel_scan_ul1.cpp | Adds pto_launch_scan_ul1_{fp16,fp32} shims. |
| csrc/kernel/kernel_kda_wy.cpp | Adds pto_launch_kda_wy shim. |
| csrc/kernel/kernel_kda_kkt.cpp | Adds pto_launch_kda_kkt shim. |
| csrc/kernel/kernel_kda_gate_cumsum.cpp | Adds pto_launch_kda_gate_cumsum shim. |
| csrc/kernel/kernel_kda_chunk_o.cpp | Adds pto_launch_kda_chunk_o shim. |
| csrc/kernel/kernel_kda_chunk_h.cpp | Adds pto_launch_kda_chunk_h shim. |
| csrc/kernel/kernel_gdn_wy_fast.cpp | Adds pto_launch_gdn_wy_fast shim. |
| csrc/kernel/kernel_gdn_scaled_dot_kkt.cpp | Adds pto_launch_gdn_scaled_dot_kkt shim. |
| csrc/kernel/kernel_gdn_chunk_o.cpp | Adds pto_launch_gdn_chunk_o shim. |
| csrc/kernel/kernel_gdn_chunk_h.cpp | Adds pto_launch_gdn_chunk_h shim. |
| csrc/kernel/kernel_gdn_chunk_cumsum.cpp | Adds pto_launch_gdn_chunk_cumsum_fp32 shim. |
| csrc/kernel/kernel_csr_gather.cpp | Adds pto_launch_csr_gather_{fp16,fp32} shims. |
| csrc/kernel/kernel_batch_matrix_square.cpp | Adds pto_launch_batch_matrix_square_{fp16,fp32} shims. |
| csrc/kernel/kernel_abs.cpp | Adds pto_launch_vabs_{fp16,fp32} shims. |
| csrc/host/utils.h | Adds <cstdint> and routes EXEC_KERNEL_CMD through pto_launch_<kernel>. |
| csrc/host/torch_tri_inv.h | Replaces generated aclrtlaunch_* includes with pto_launch_* declarations. |
| csrc/host/torch_tri_inv_trick.h | Replaces generated aclrtlaunch_* include with pto_launch_* declaration. |
| csrc/host/torch_tri_inv_rec_unroll.h | Replaces generated aclrtlaunch_* includes with pto_launch_* declarations. |
| csrc/host/torch_tri_inv_ns.h | Replaces generated aclrtlaunch_* include with pto_launch_* declaration. |
| csrc/host/torch_swiglu.h | Replaces generated aclrtlaunch_* include with pto_launch_* declaration. |
| csrc/host/torch_simple_matmul.h | Replaces generated aclrtlaunch_* includes with pto_launch_* declarations. |
| csrc/host/torch_scan_ul1.h | Replaces generated aclrtlaunch_* includes with pto_launch_* declarations. |
| csrc/host/torch_kda_wy.h | Replaces generated aclrtlaunch_* include with pto_launch_* declaration. |
| csrc/host/torch_kda_kkt.h | Replaces generated aclrtlaunch_* include with pto_launch_* declaration. |
| csrc/host/torch_kda_gate_cumsum.h | Replaces generated aclrtlaunch_* include with pto_launch_* declaration. |
| csrc/host/torch_kda_chunk_o.h | Replaces generated aclrtlaunch_* include with pto_launch_* declaration. |
| csrc/host/torch_kda_chunk_h.h | Replaces generated aclrtlaunch_* include with pto_launch_* declaration. |
| csrc/host/torch_gdn_wy_fast.h | Replaces generated aclrtlaunch_* include with pto_launch_* declaration. |
| csrc/host/torch_gdn_scaled_dot_kkt.h | Replaces generated aclrtlaunch_* include with pto_launch_* declaration. |
| csrc/host/torch_gdn_chunk_o.h | Replaces generated aclrtlaunch_* include with pto_launch_* declaration. |
| csrc/host/torch_gdn_chunk_h.h | Replaces generated aclrtlaunch_* include with pto_launch_* declaration. |
| csrc/host/torch_gdn_chunk_cumsum.h | Replaces generated aclrtlaunch_* include with pto_launch_* declaration. |
| csrc/host/torch_csr_gather.h | Replaces generated aclrtlaunch_* includes with pto_launch_* declarations. |
| csrc/host/torch_batch_matrix_square.h | Replaces generated aclrtlaunch_* includes with pto_launch_* declarations. |
| csrc/host/torch_abs.h | Replaces generated aclrtlaunch_* includes with pto_launch_* declarations. |
| CMakeLists.txt | Adds USE_ASC_LANGUAGE option, SOC→arch mapping, ASC-language toolchain setup, and ASC-specific target options. |
💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.
Comment on lines
+425
to
+426
| swiglu_fp16<<<blockDim, nullptr, stream>>>((GM_ADDR)x, (GM_ADDR)y, batch, | ||
| input_n); |
Comment on lines
+141
to
+144
| extern "C" void pto_launch_vabs_fp16(uint32_t blockDim, void* stream, void* x, | ||
| void* z, uint32_t in_length) { | ||
| vabs_fp16<<<blockDim, nullptr, stream>>>((GM_ADDR)x, (GM_ADDR)z, in_length); | ||
| } |
Comment on lines
+146
to
+149
| extern "C" void pto_launch_vabs_fp32(uint32_t blockDim, void* stream, void* x, | ||
| void* z, uint32_t in_length) { | ||
| vabs_fp32<<<blockDim, nullptr, stream>>>((GM_ADDR)x, (GM_ADDR)z, in_length); | ||
| } |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
This MR updates the Ascend CMake infrastructure to the new
9.0.0version whereASClanguage feature is enabled.In short, this remove the
aclrtXXXobsolete calls withkernel<<<block_dim,...,stream>kernel invocation.Tests pass on 910B4 and 910B2:
make wheel install test