tile-ai/tilelang-ascend

Ascend TileLang adapter

What it solves

TileLang-Ascend simplifies the creation of high-performance AI compute kernels for Huawei Ascend NPUs. It removes the need for developers to manually handle complex low-level optimizations while still providing the tools necessary to achieve state-of-the-art performance for operations like GEMM, attention mechanisms, and vector operations.

How it works

It is a domain-specific language (DSL) with a Pythonic syntax built on the TVM compiler infrastructure. It maps GPU-like memory hierarchies (global, shared, and register memory) to the Ascend NPU architecture (global memory, L1/unified buffers, and L0 buffers). The compiler supports two technical routes for code generation: Ascend C & PTO and AscendNPU IR. It offers different programming modes, including a "Developer mode" for automatic scope separation and synchronization, and an "Expert mode" for manual control over execution scopes and synchronization flags.

Who it’s for

AI kernel developers and engineers who need to optimize AI workloads specifically for Huawei Ascend NPU hardware (such as A2 and A3 devices).

Highlights

  • Broad Operator Support: Includes implementations for GEMM, Flash Attention, Sparse Flash Attention, Convolution, and Softmax.
  • Memory Optimization: Features automatic buffer reuse to reduce on-chip memory footprint and automatic workspace allocation.
  • Performance Tools: Supports software pipelining (T.Pipelined), automatic vectorization (T.Parallel), and L2 cache swizzling.
  • Developer Experience: Provides a JIT compiler, debug tools (T.printf, T.dump_tensor), and seamless PyTorch integration via torch_tl_ascend.

Related

  • Project
    tile-ai/tilelang

    A concise domain-specific language and compiler built on TVM for developing high-performance GPU/CPU kernels like GEMM and FlashAttention across multiple hardware backends.

  • Project
    pytorch/helion

    A Python-embedded DSL for authoring machine learning kernels that compiles to Triton, offering higher abstractions and automated performance tuning to simplify GPU programming.

  • Project
    NVIDIA/tilus

    Tilus is a tile-level GPU kernel programming language that enables high-performance, low-precision computation with explicit control over shared memory and registers.

  • Project
    xLLM-AI/xllm

    xLLM is an efficient LLM inference framework optimized for Chinese AI accelerators, designed for high-throughput, low-latency enterprise-grade deployment.

  • Project
    datawhalechina/diy-llm

    A Chinese‑language, open‑source curriculum that teaches you to build, train, and evaluate large language models from scratch, complete with lecture notes, six hands‑on assignments, and runnable PyTorch code.