skills.rest
WangWindow
WangWindowCOMMUNITY
Explore all skills by this author →
WebsiteGitHub

ascendc-registry-invoke-to-direct-invoke

Migrates AscendC custom operator kernels from registry invocation to direct <<<>>> kernel launch.

Updated Sep 15, 2026
Search Tags:#tiling#npu#ascendc#kernel launch#dependency decoupling#operator migration#direct invoke
One-click install
npx skills add https://github.com/WangWindow/CANN-BatchMatMulMaxsum --skill ascendc-registry-invoke-to-direct-invoke-wangwindow
Or copy as Structured Prompt for Agent▼
Please help me install this Agent Skill.
Skill: ascendc-registry-invoke-to-direct-invoke
Source: https://github.com/WangWindow/CANN-BatchMatMulMaxsum/tree/main/.agents/skills/ascendc-registry-invoke-to-direct-invoke
Command: npx skills add https://github.com/WangWindow/CANN-BatchMatMulMaxsum --skill ascendc-registry-invoke-to-direct-invoke-wangwindow

SYSTEM DOCUMENTATION & REQUIREMENTS

💡 This Skill includes references (resource) components.

What problem does it solve? Converting an AscendC custom operator project from the registration framework to <<<>>> direct kernel launch is error-prone: developers often accidentally modify kernel computation logic or tiling math while only the framework glue should change, silently breaking correctness or performance. This Skill enforces a fidelity-preserving migration workflow with strict zero-modification principles. ## Core Features & Use Cases - Zero-modification migration rules: Kernel implementation code and tiling math stay untouched; only three framework glue points change (entry functions, TilingData macro-to-struct conversion, host tiling interface replacement). - Dependency decoupling workflow: Traces relative ../ includes, builds a minimal external dependency closure in local_deps.h, and moves the operator's own source files wholesale. - Edge-case guidance: Covers -xasc host/device dual-pass guards for MicroAPI/Reg kernels, bf16 template mangling workarounds, and GM_ADDR ABI consistency across host, launch, and kernel sides. - Use Case: You have an existing AscendC operator project using GET_TILING_DATA and TILING_KEY_IS dispatch and need a standalone directly-launchable kernel; the Skill walks you through entry splitting, TilingData struct conversion, and static verification checklists. ## Quick Start Convert my custom operator project in this directory from registry invocation to <<<>>> direct kernel launch, keeping the kernel and tiling logic unchanged.

Frequently Asked Questions about ascendc-registry-invoke-to-direct-invoke

High-intent search queries and answers about installing and using this skill.

FAQPage Schema
How do I convert an AscendC custom operator to <<<>>> direct kernel launch?▼

Replace only the registration framework glue: split the entry into per-tilingKey __global__ functions taking TilingData by value, convert BEGIN_TILING_DATA_DEF macros to a plain struct, and swap gert::TilingContext for PlatformAscendCManager. Kernel implementation and tiling math stay unchanged.

What code changes are allowed when migrating AscendC kernels to direct launch?▼

Only three change types are allowed: replacing relative includes with local ones, optionally wrapping code in a namespace, and sourcing TilingData from a local POD struct. Modifying kernel member functions, tiling formulas, or merging alignment branches is forbidden.

Why does my templated __global__ kernel fail to link with bfloat16_t?▼

The -xasc compiler mangles GM_ADDR parameters incorrectly for template __global__ functions instantiated with bfloat16_t, causing undefined symbol errors. Use explicitly named non-template entry functions such as op_fp32, op_fp16, and op_bf16 instead.

How do I handle MicroAPI or Reg types with -xasc dual-pass compilation?▼

Wrap kernel implementation headers using MicroAPI or Reg types in #if !defined(__NPU_HOST__) guards, since those types exist only in the device pass. Keep pure type-definition headers like tilingdata.h unguarded, and guard only the body of __global__ entry functions.

Can I use void pointers and cast to GM_ADDR at kernel launch time?▼

No, the AscendC compiler rejects reinterpret_cast<GM_ADDR> at the call site. Declare device buffers as GM_ADDR from aclrtMalloc onward and pass them directly to the kernel launch, keeping GM_ADDR consistent across host, launch, and kernel sides.

When should I not use this registry-to-direct-invoke migration approach?▼

Do not use it when developing a new operator from scratch, since the Skill only covers migrating existing operator projects. It also does not apply to host, graph, or registration integration work beyond the direct-launch conversion boundary.

Related Skills

Explore More Software Engineering Skills→
xchang1121xchang1121
6

ascendc-registry-to-direct-invoke

Migrates registry-based AscendC operators to direct-invoke kernel.py and ascendc_op projects.

Community
Advanced
sgl-projectsgl-project
31.6k

add-sgl-kernel

Guide adding AOT CUDA/C++ kernels to the sgl-kernel library.

Official
Advanced
mindspore-aimindspore-ai
258

triton-ascend-examples-mindspore

Register Triton Ascend custom operators with MindSpore primitives and tensor patterns.

Official
Intermediate
thinkeloquentthinkeloquent

create-atom

Scaffold React and TypeScript atom boilerplate with CSS Modules under src/atoms.

Official
Intermediate