9.4 KiB

Raw Blame History

TileLang on MetaX (MACA) 新手指南

本文面向第一次接触 tilelang-metax 的新同学，带你从环境确认 → 源码构建 → 跑通测试 → 运行样例，走完一整条可用路径。

适用硬件：MetaX C500（已在该型号上验证）
适用 MACA 版本：>= 3.3.0（本文验证环境为 3.5.3.20）
适用 Python 版本：>= 3.10（本文验证环境为 3.10）
操作系统：Linux（Ubuntu/Debian 系列）

1. 确认 GPU 环境到位

在动手之前，先确认机器上已经能正确识别 MetaX 显卡：

mx-smi

正常应看到类似下面的输出（示例为双卡 C500）：

=================== MetaX System Management Interface Log ===================
Attached GPUs                                     : 2
+---------------------------------------------------------------------------------+
| MX-SMI 2.2.12                      Kernel Mode Driver Version: 3.6.11           |
| MACA Version: 3.5.3.20             BIOS Version: 1.31.1.0                       |
|------------------+-----------------+---------------------+----------------------|
| Board       Name | GPU   Persist-M | Bus-id              | GPU-Util      sGPU-M |
| Pwr:Usage/Cap    | Temp       Perf | Memory-Usage        | GPU-State            |
|==================+=================+=====================+======================|
| 0     MetaX C500 | 0           Off | 0000:0f:00.0        | 0%          Disabled |
| 56W / 350W       | 37C          P0 | 826/65536 MiB       | Available            |
+------------------+-----------------+---------------------+----------------------+
| 1     MetaX C500 | 1           Off | 0000:10:00.0        | 0%          Disabled |
| 59W / 350W       | 40C          P0 | 826/65536 MiB       | Available            |
+------------------+-----------------+---------------------+----------------------+

如果 mx-smi 不存在，说明驱动或 MACA SDK 尚未安装，请先参考 Installation_maca 完成前置步骤。

2. 环境准备（二选一）

方式 A：Docker 镜像（推荐，最省心）

MetaX 提供了预装好 PyTorch + MACA 的 Docker 镜像，能省去大量手动配环境的功夫。

# 1. 登录并拉取镜像（以 maca-pytorch:3.3.0.4-torch2.6-py310-ubuntu24.04-amd64 为例）
docker login --username=cr_temp_user --password=eyJpbnN0YW5jZUlkIjoiY3JpLXpxYTIzejI2YTU5M3R3M2QiLCJ0aW1lIjoiMTc3MDg5NTI0MzAwMCIsInR5cGUiOiJzdWIiLCJ1c2VySWQiOiIyMDcwOTQwMTA1NjYzNDE3OTIifQ:91ecedb8bd5c4af6858745f0329d069263e1bf82 cr.metax-tech.com
docker pull cr.metax-tech.com/public-library/maca-pytorch:3.3.0.4-torch2.6-py310-ubuntu24.04-amd64

# 2. 启动容器（需要映射 GPU 设备）
docker run -it --net=host --device=/dev/dri --device=/dev/mxcd --group-add video \
  --name tilelang-maca \
  cr.metax-tech.com/public-library/maca-pytorch:3.3.0.4-torch2.6-py310-ubuntu24.04-amd64 \
  /bin/bash

# 3. 进入容器后补装构建工具
apt-get update
apt-get install -y cmake git

注意：宿主机必须已经安装好 MetaX 驱动；容器内只需 SDK + PyTorch。

方式 B：原生 Linux 手动安装

如果你不想用 Docker，需要在宿主机上完成：

系统依赖

apt-get update
apt-get install -y python3 python3-dev python3-setuptools gcc \
  zlib1g-dev build-essential cmake libedit-dev git

MACA SDK（驱动 + maca-sdk + cu-bridge + PyTorch）
请参考 Installation_maca 的 "MACA SDK" 和 "Install pytorch" 章节完成安装。
Python 依赖
```
pip install z3-solver cython psutil cloudpickle tqdm torch-c-dlpack-ext
```
z3-solver 版本需 >= 4.13.0，且安装路径下能找到 include 和 lib 目录。

3. 从 0 到 1 构建 TileLang

以下步骤在** Docker 容器内或原生 Linux** 中都相同。

3.1 克隆仓库（必须加 `--recursive`）

git clone --recursive https://github.com/tile-ai/tilelang-metax.git
cd tilelang-metax

子模块包含 TVM，漏掉 --recursive 会导致后续编译报错。

3.2 配置 git（构建脚本会检查 committer 信息）

git config --global user.email "you@example.com"
git config --global user.name "Your Name"

3.3 编译

# 通过环境变量启用 MACA 后端
USE_MACA=ON cmake -B build

# 并行编译（-j 后跟核心数，或直接写 -j$(nproc)）
make -C build -j$(nproc)

正常结束时，你会看到：

[100%] Built target tilelang

并在 build/lib/ 下生成如下关键产物：

libtilelang.so
libtvm.so
tilelang_cython_wrapper.cpython-310-x86_64-linux-gnu.so（具体文件名随 Python 版本变化）

3.4 安装 TVM FFI

TileLang 依赖 TVM 的 FFI 包，需要从子模块源码安装：

cd 3rdparty/tvm/3rdparty/tvm-ffi && pip install . && cd -

3.5 设置环境变量

编译完成后，每次新打开终端都需要导出以下环境（建议写进 ~/.bashrc）：

export MACA_PATH=/opt/maca
export LD_LIBRARY_PATH=${MACA_PATH}/lib:${MACA_PATH}/mxgpu_llvm/lib:$LD_LIBRARY_PATH
export PATH=${MACA_PATH}/mxgpu_llvm/bin:${PATH}
export PYTHONPATH=/path/to/mcTileLang:$PYTHONPATH

把 /path/to/mcTileLang 替换为你实际克隆的路径，例如 /data/mcTileLang。

4. 验证安装

执行下面的最小命令，确认 Python 能正常导入 tilelang：

python -c "import tilelang; print(tilelang.__version__)"

预期输出示例：

Loading tilelang libs from dev root: /data/mcTileLang/build
0.1.9+cuda.gitd32a0150

只要没有 ImportError，说明构建和环境变量都已就绪。

5. 跑测试

TileLang 针对 MetaX 有独立的测试目录 testing/maca/，建议优先跑这里面的用例；testing/python/ 下的部分用例面向 CUDA 特性，在 MACA 上可能出现 skip 或 fail。

5.1 跑 MACA 专属测试（推荐）

# 跑全部 MACA 测试（约 700+ 条，耗时较长）
python -m pytest testing/maca/ -x

# 或先跑一个子集做快速验证
python -m pytest testing/maca/maca/ -x -v
python -m pytest testing/maca/language/test_tilelang_language_copy.py -x -v

以 testing/maca/maca/test_maca_f32x2_intrinsics.py 为例，实际验证结果：

============================== 54 passed in 28.55s ==============================

5.2 跑通用测试（按需）

# JIT 相关测试
python -m pytest testing/python/jit/ -x -v

# 语言特性测试（部分用例在 MACA 上会自动 skip）
python -m pytest testing/python/language/test_tilelang_language_copy.py -x -v

常见现象：

SKIPPED：该用例依赖 CUDA 特有特性（如 TMA、WGMMA），在 MACA 上自动跳过，属于正常行为。

FAILED：如果看到 test_gemm_i8i8i32_nt 这类失败，通常是已知兼容性问题，不影响大部分 FP16/BF16 GEMM 功能。

5.3 跑单个测试文件 / 单个用例

# 单个文件
python -m pytest testing/maca/kernel/test_tilelang_kernel_gemm.py -x -v

# 单个用例（通过 -k 过滤）
python -m pytest testing/maca/kernel/test_tilelang_kernel_gemm.py -x -v -k "test_gemm_fp16"

6. 跑样例

测试通过后，说明编译器链路已经打通，可以尝试运行官方样例。

6.1 极简入门：quickstart.py

python examples/quickstart.py

这个脚本会：

定义一个带 ReLU 的 GEMM kernel；
在 MetaX C500 上执行；
与 PyTorch 参考结果做精度对比；
打印 kernel 源码和延迟（示例延迟约 0.12 ms 级别）。

6.2 GEMM 系列样例

# 基础 GEMM
python examples/gemm/example_gemm.py

# 使用 intrinsics 的 GEMM
python examples/gemm/example_gemm_intrinsics.py

# Persistent kernel GEMM
python examples/gemm/example_gemm_persistent.py

# Auto-tune GEMM（搜索最优 tile 配置）
python examples/gemm/example_gemm_autotune.py

6.3 其他高阶样例

样例目录	说明
`examples/flash_attention/`	FlashAttention 实现
`examples/linear_attention/`	Linear Attention（RetNet / Mamba）
`examples/deepseek_mla/`	DeepSeek MLA Decoding
`examples/dequantize_gemm/`	反量化 GEMM

进入对应目录后，通常直接 python xxx.py 即可运行。

7. 常见问题速查

Q1: `mx-smi` 正常，但编译时提示找不到 MACA

确认 USE_MACA=ON 是作为环境变量写在 cmake 前面，而不是 cmake 的 -D 参数：

# 正确
USE_MACA=ON cmake -B build

# 错误（cmake 不会识别）
cmake -DUSE_MACA=ON -B build

Q2: 运行时报 `libmxc-runtime64.so: cannot open shared object file`

环境变量 LD_LIBRARY_PATH 没有包含 /opt/maca/lib。请重新执行第 3.5 节的环境变量导出。

Q3: `git clone` 后编译报错，提示缺少 TVM 文件

大概率是忘了加 --recursive。补救：

git submodule update --init --recursive

Q4: 测试里出现很多 `PytestUnknownMarkWarning`

这是 pytest 对 gpu、cuda 等自定义 mark 的提示，不影响测试结果，可忽略。

8. 下一步

学习语法：阅读 examples/gemm/README.md 和官方编程指南。
调试工具：TileLang 支持 T.print 打印变量/缓冲区，以及 layout 可视化工具，详见 examples/plot_layout/。
性能调优：尝试 examples/gemm/example_gemm_autotune.py 了解自动调优机制。

如果在配置或运行过程中遇到文档未覆盖的问题，欢迎提交 Issue 或在社区讨论。

9.4 KiB Raw Blame History Unescape Escape