tilelang-api-best-practices

Documents TileLang Ascend API usage for memory, compute, scheduling, and synchronization primitives.

Updated Sep 15, 2026
One-click install
npx skills add https://github.com/WangWindow/CANN-BatchMatMulMaxsum --skill tilelang-api-best-practices-wangwindow
Or copy as Structured Prompt for Agent
Please help me install this Agent Skill.
Skill: tilelang-api-best-practices
Source: https://github.com/WangWindow/CANN-BatchMatMulMaxsum/tree/main/.agents/skills/tilelang-api-best-practices
Command: npx skills add https://github.com/WangWindow/CANN-BatchMatMulMaxsum --skill tilelang-api-best-practices-wangwindow

SYSTEM DOCUMENTATION & REQUIREMENTS

💡 This Skill includes references (resource) components.

What problem does it solve? Writing correct TileLang kernels for Ascend NPU requires knowing dozens of APIs across memory allocation, data movement, GEMM, reduction, scheduling, and synchronization, and misuse causes silent data corruption or compile errors. This Skill provides an indexed reference of correct API usage patterns and constraints. ## Core Features & Use Cases - API Quick Reference: Lookup tables for kernel definition, memory allocation (Developer vs Expert modes), data movement (T.copy), compute (T.gemm_v0, T.reduce_sum/max, T.tile.xxx), scheduling (T.Pipelined, T.Persistent), and synchronization primitives. - Scenario Index: Maps common tasks like GEMM, Softmax, CV-fused operators, multi-core atomic accumulation, sorting, and kernel debugging to the relevant reference documents and key techniques. - Constraint Documentation: Covers fractal size limits for GEMM buffers, memory alignment and capacity limits, T.copy slicing restrictions, T.Parallel SIMD limitations, and dtype-specific hardware path adaptations. - Use Case: When implementing a fused BatchMatmul + ReduceMax + ReduceSum operator on Ascend NPU, consult the GEMM and reduction sections to select correct block sizes, workspace routing between Cube and Vector cores, and pipeline configuration. ## Quick Start Ask how to allocate on-chip memory and move data between GM, L1, UB, and L0 levels when writing a TileLang Ascend kernel.

Frequently Asked Questions about tilelang-api-best-practices

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

FAQPage Schema
How do I allocate on-chip memory in TileLang Ascend kernels?

Use T.alloc_shared and T.alloc_fragment in Developer mode where the compiler infers L1/UB/L0 placement, or T.alloc_ub, T.alloc_L1, and T.alloc_L0A/L0B/L0C in Expert mode for explicit control. Buffers must respect per-level alignment and size limits.

How do I move data between memory levels in TileLang Ascend?

Use T.copy for GM to L1, L1 to L0A/L0B, L0C to GM, and GM to UB transfers. UB and L1 cannot exchange data directly; route through a GM workspace buffer declared with workspace_idx in the jit decorator.

When should I use Developer mode versus Expert mode in TileLang?

Developer mode (T.alloc_shared, T.alloc_fragment) lets the compiler manage storage levels and synchronization, suiting most operators. Expert mode (T.alloc_L1, T.alloc_ub) gives explicit control for performance tuning but requires manual T.barrier_all and flag synchronization.

Why does T.gemm_v0 produce wrong results with small block sizes?

GEMM buffers must satisfy fractal limits: L0A/L0B are fixed 512-byte fractals and L0C is fixed 16x16. For float16, M, N, and K must each be at least 16; a block_N of 8 violates the L0C fractal and silently corrupts results.

Does T.Parallel support if-else branches on Ascend NPU?

No. T.Parallel lowers to SIMD vector instructions that cannot execute per-element conditionals, so if-else causes compile errors. Replace branching with T.tile.compare to build a mask followed by T.tile.select to choose values.

What is the difference between T.Pipelined and T.Persistent?

T.Pipelined(range, num_stages) overlaps data copies with compute across loop iterations to hide memory latency. T.Persistent(domain, wave_size, index) schedules adjacent tiles onto the same AI core for cache-friendly multi-core load balancing.