tilelang-ascend-optimization

Optimizes TileLang Ascend operator kernels using intra-core, inter-core, and pipeline techniques.

6|1|Updated Apr 19, 2026
One-click install
npx skills add https://github.com/xchang1121/op-autoresearch --skill tilelang-ascend-optimization-xchang1121
Or copy as Structured Prompt for Agent▼
Please help me install this Agent Skill.
Skill: tilelang-ascend-optimization
Source: https://github.com/xchang1121/op-autoresearch/tree/main/skills/tilelang-ascend/fundamentals/tilelang-ascend-optimization
Command: npx skills add https://github.com/xchang1121/op-autoresearch --skill tilelang-ascend-optimization-xchang1121

SYSTEM DOCUMENTATION & REQUIREMENTS

What problem does it solve? TileLang kernels on Ascend NPUs often underperform due to serial memory transfers, unvectorized scalar loops, redundant data loads, and poor Cube/Vector core coordination. This Skill provides a systematic optimization playbook that maps operator types (Cube, Vector, CV-fused) to concrete optimization techniques with before/after code patterns. ## Core Features & Use Cases - Intra-core optimization: Split-K partitioning, Double Buffer (Ping-Pong), MTE2 prefetch, Full-Load L1 residency, small-block merged loading, instruction vectorization, AXPY instruction fusion, and sparse memory access optimization for Gather workloads. - Inter-core optimization: num_stages tuning for T.Pipelined, cross-core synchronization interval adjustment, and Fixed Core mode launching by physical core count to reduce workspace memory. - Decision guidance: Optimization priority ordering by operator type (detected via IS_ASCEND_AIC / IS_ASCEND_AIV), Developer vs Expert mode selection, and a troubleshooting table mapping symptoms like Cube bubbles or memory overflow to fixes. - Use Case: A GEMM kernel on Atlas A2 shows low MTE2 bandwidth utilization; apply Split-K with Double Buffer and Full-Load to overlap L1-to-L0 transfers with MMA computation. ## Quick Start Ask the agent to analyze my TileLang Ascend kernel source and apply the appropriate intra-core and inter-core optimizations based on its operator type.

Frequently Asked Questions about tilelang-ascend-optimization

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

FAQPage Schema
How do I optimize a TileLang kernel on Ascend NPU?▼

First identify the operator type via IS_ASCEND_AIC or IS_ASCEND_AIV in the kernel source, then apply intra-core optimizations like Double Buffer and Split-K, followed by inter-core tuning for CV-fused kernels. Adjust pass_configs only as a last resort.

What is Double Buffer in TileLang Ascend kernels?▼

Double Buffer allocates two buffer copies and alternates them with side = k % 2 so memory transfers overlap with computation in a pipelined fashion. Vector cores need at least 128 elements per block and Cube cores at least 256 for the overlap to be effective.

When should I use Fixed Core mode in TileLang?▼

Use Fixed Core mode when logical task count greatly exceeds physical core count, workspace memory grows linearly with block_num, or the kernel has heavy initialization. Launch with T.Kernel(core_num, is_npu=True) and manually distribute tasks per core.

Does TileLang Ascend support automatic synchronization?▼

Yes, Developer mode enables TL_ASCEND_AUTO_SYNC so the compiler inserts synchronization automatically. If the generated Ascend C code shows unexpected sync placement, disable it and insert T.set_flag and T.wait_flag manually in Expert mode.

Why does my CV-fused kernel have Cube core bubbles?▼

Cube core bubbles usually mean the Vector core is slower and num_stages is too small in T.Pipelined. Increase num_stages gradually from 2 up to the loop count, but reduce it if memory planning reports overflow.

What are the limitations of pass_configs tuning?▼

Changing pass_configs like disabling auto-sync reduces reliance on Developer mode features and requires manual synchronization, which is error-prone. It should only be used after other optimization techniques fail to meet performance targets.