tile-ai/tilelang-ascend
Ascend TileLang adapter
What it solves
TileLang-Ascend addresses the difficulty of creating high-performance AI compute kernels for Huawei Ascend NPUs. It allows developers to write optimized kernels (such as GEMM, attention mechanisms, and vector operations) using a Pythonic syntax, removing the need to manually handle complex low-level NPU optimizations while maintaining state-of-the-art performance.
How it works
It is a domain-specific language (DSL) built on the TVM compiler infrastructure. It maps high-level Pythonic code to the Ascend NPU architecture by managing memory hierarchies (Global, L1, and L0 buffers) and providing specialized compiler backends (Ascend C & PTO and AscendNPU IR). The system supports both a "Developer mode" for automatic synchronization and scope separation, and an "Expert mode" for manual control over execution scopes and synchronization flags.
Who it’s for
AI engineers and kernel developers who need to optimize compute workloads specifically for Huawei Ascend A2 and A3 processors.
Highlights
- Broad Operator Support: Includes implementations for GEMM, Flash Attention, Sparse Flash Attention, Convolution, and Softmax.
- Memory Optimizations: Features automatic buffer reuse to reduce on-chip memory footprint and automatic workspace allocation.
- Performance Tooling: Supports software pipelining (
T.Pipelined), automatic vectorization (T.Parallel), and L2 cache swizzling. - PyTorch Integration: Provides
torch_tl_ascendfor seamless workflows between PyTorch and the NPU.
Related
- Project
- Project
- Project
- Project
- Project