openeuler/tilelang repository overview

⁠Quick reference

• The official TileLang Ascend docker images

• Maintained by: openEuler CloudNative SIG⁠

• Where to get help: openEuler CloudNative SIG⁠, openEuler⁠

⁠TileLang | openEuler

Current TileLang docker images are built on the openEuler⁠. This repository is free to use and exempted from per-user rate limits.

TileLang⁠ is an AI programming language/DSL for writing high-performance GPU/NPU kernels with a tile-based abstraction. The Ascend adaptation lives in the ecosystem repository tilelang-ascend⁠ (Ascend C & PTO route): it builds a wheel from source so that TileLang kernels are JIT-compiled for and executed on the Ascend NPU.

This image ships tilelang-ascend built from source with torch_npu⁠, and keeps the upstream examples under /opt/tilelang-ascend/examples. The software stack is validated on a single Ascend 910B4 card with CANN 9.1.0, torch==2.9.0 and torch_npu==2.9.0.post8 on Python 3.12, following the Ascend quick start guard⁠.

Read more about Ascend at hiascend.com⁠ and explore the TileLang documentation at tilelang.com⁠.

⁠Supported tags and respective Dockerfile links

The tag of each TileLang docker image consists of the TileLang version and the version of the software stack. The details are as follows

⁠Usage

⁠Quick start 1: supported devices

• Atlas A2 Training series (Atlas 800T A2, Atlas 900 A2 PoD, Atlas 200T A2 Box16, Atlas 300T A2)

• Atlas 800I A2 Inference series (Atlas 800I A2)

⁠Quick start 2: setup environment using container

# Update DEVICE according to your device (/dev/davinci[0-7])
export DEVICE=/dev/davinci0
# Update the tilelang image
export IMAGE=openeuler/tilelang:0.1.2.000-torch_npu2.9.0-cann9.1.0-python3.12-oe2403sp4
docker run --rm \
--name tilelang \
--device $DEVICE \
--device /dev/davinci_manager \
--device /dev/devmm_svm \
--device /dev/hisi_hdc \
-v /usr/local/dcmi:/usr/local/dcmi \
-v /usr/local/bin/npu-smi:/usr/local/bin/npu-smi \
-v /usr/local/Ascend/driver/lib64/:/usr/local/Ascend/driver/lib64/ \
-v /usr/local/Ascend/driver/version.info:/usr/local/Ascend/driver/version.info \
-v /etc/ascend_install.info:/etc/ascend_install.info \
-v /root/.cache:/root/.cache \
-it $IMAGE bash

⁠Quick start 3: run a GEMM kernel on Ascend NPU

Load the CANN environment and validate the stack inside the container:

source /usr/local/Ascend/ascend-toolkit/set_env.sh
python3 -c "import torch, torch_npu, tilelang; print('tilelang', tilelang.__version__); print('is_available:', torch.npu.is_available()); print('count:', torch.npu.device_count())"

Run the upstream GEMM example (1024x1024x1024 fp16 by default; JIT compilation, NPU execution and numerical check against PyTorch):

cd /opt/tilelang-ascend
python3 examples/gemm/example_gemm.py

The script prints init successful! after initialization and Kernel Output Match! once the NPU kernel output matches the PyTorch reference. More operator examples (FlashAttention, Softmax, normalization, DeepSeek V4 operators, etc.) are available in the upstream examples/ directory.

⁠Question and answering

If you have any questions or want to use some special features, please submit an issue or a pull request on openeuler-docker-images⁠.