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

tilelang-programming-model-guide

Guides selection and configuration of TileLang Ascend Developer and Expert programming modes.

Updated Sep 15, 2026
Search Tags:#ascend#npu#tilelang#gemm#kernel programming#pass_configs#cv fusion
One-click install
npx skills add https://github.com/WangWindow/CANN-BatchMatMulMaxsum --skill tilelang-programming-model-guide-wangwindow
Or copy as Structured Prompt for Agent▼
Please help me install this Agent Skill.
Skill: tilelang-programming-model-guide
Source: https://github.com/WangWindow/CANN-BatchMatMulMaxsum/tree/main/.agents/skills/tilelang-programming-model-guide
Command: npx skills add https://github.com/WangWindow/CANN-BatchMatMulMaxsum --skill tilelang-programming-model-guide-wangwindow

SYSTEM DOCUMENTATION & REQUIREMENTS

💡 This Skill includes references (resource) components.

What problem does it solve? Choosing between TileLang Ascend Developer and Expert programming modes and configuring the four Ascend-specific pass_configs switches is error-prone, and mixing automatic CV separation with manual T.Scope or missing manual synchronization causes correctness bugs. ## Core Features & Use Cases - Mode Comparison: Contrasts Developer and Expert modes across memory allocation, scope control, synchronization, CV interaction, and pass_configs, including a hybrid mode for most practical operators. - pass_configs Configuration: Explains the four Ascend switches (TL_ASCEND_AUTO_SYNC, TL_ASCEND_MEMORY_PLANNING, TL_ASCEND_AUTO_CV_COMBINE, TL_ASCEND_AUTO_CV_SYNC) with per-scenario configuration tables for pure Vector kernels, GEMM, Flash Attention, and CV fusion. - Expert-to-Developer Conversion: Provides a step-by-step migration checklist and mapping table, including the threads=2 to vid-elimination to workspace-elimination chain for on-chip CV data flow. - Use Case: When porting an Expert-mode Flash Attention kernel to Developer mode, follow the conversion rules to remove T.Scope, manual barriers, and workspace parameters while enabling all four pass_configs. ## Quick Start Ask how to configure pass_configs for a Developer-mode GEMM kernel on Ascend NPU and whether manual T.Scope annotations are needed.

Frequently Asked Questions about tilelang-programming-model-guide

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

FAQPage Schema
What is the difference between TileLang Developer mode and Expert mode?▼

Developer mode uses T.alloc_shared/T.alloc_fragment with compiler-automatic Cube/Vector separation and synchronization, while Expert mode uses T.alloc_L1/T.alloc_ub/T.alloc_L0C with manual T.Scope and T.barrier_all. Developer enables all four pass_configs; Expert disables them for low-level control.

How do I configure pass_configs for TileLang Ascend kernels?▼

Pure Vector kernels need only TL_ASCEND_AUTO_SYNC and TL_ASCEND_MEMORY_PLANNING enabled. Developer GEMM, Flash Attention, and CV fusion kernels should enable all four switches including AUTO_CV_COMBINE and AUTO_CV_SYNC. Expert mode disables all four.

Can I use T.Scope together with TL_ASCEND_AUTO_CV_COMBINE?▼

No, mixing them is not allowed. When AUTO_CV_COMBINE is enabled, the compiler separates Cube and Vector code automatically based on buffer types and primitives, and handwritten T.Scope annotations can cause incorrect compilation. Choose either automatic or manual scope control.

How do I convert an Expert mode kernel to Developer mode in TileLang?▼

Enable all four pass_configs, replace alloc_L1/alloc_ub with alloc_shared and alloc_L0C with alloc_fragment, remove T.Scope blocks and manual synchronization calls, and eliminate workspace and vid by adding threads=2 to T.Kernel with on-chip direct copies.

What happens if I disable AUTO_SYNC without adding manual synchronization?▼

It is a correctness risk, not just a performance issue. With AUTO_SYNC off you must manually insert T.barrier_all or set_flag/wait_flag between data movement and compute, and with AUTO_CV_SYNC off you must manage cross-core flags yourself.

When should I keep workspace and vid instead of eliminating them in CV fusion?▼

Keep explicit workspace and vid only for complex scenarios such as multi-version inter-core pipelining or fine-grained semaphore control that the compiler cannot handle automatically. Standard Developer CV fusion should use threads=2 with on-chip direct copies.

Related Skills

Explore More Software Engineering Skills→
xchang1121xchang1121
6

tilelang-ascend-api

Provides a complete TileLang Ascend API reference for writing NPU kernels.

Community
Intermediate
xchang1121xchang1121
6

tilelang-ascend-attention

Guides implementation of Attention operators using TileLang Cube and Vector fusion on Ascend hardware.

Community
Advanced
dongg622dongg622

ascend-profiling-anomaly

Analyze Huawei Ascend NPU profiling traces to detect performance anomalies and bottlenecks.

Community
Advanced
xchang1121xchang1121
6

tilelang-ascend-debugging

Diagnose common coding pattern errors in TileLang-Ascend operator kernels.

Community
Intermediate