From 760df0e5e85adef364c1fe72b7d4dda39309a746 Mon Sep 17 00:00:00 2001 From: yyyymmm Date: Fri, 15 May 2026 15:33:16 +0800 Subject: [PATCH] ADD file via upload --- quickstart_maca.md | 304 +++++++++++++++++++++++++++++++++++++++++++++ 1 file changed, 304 insertions(+) create mode 100644 quickstart_maca.md diff --git a/quickstart_maca.md b/quickstart_maca.md new file mode 100644 index 0000000..ade922d --- /dev/null +++ b/quickstart_maca.md @@ -0,0 +1,304 @@ +# 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 显卡： + +```bash +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](./Installation_maca.md) 完成前置步骤。 + +--- + +## 2. 环境准备（二选一） + +### 方式 A：Docker 镜像（推荐，最省心） + +MetaX 提供了预装好 PyTorch + MACA 的 Docker 镜像，能省去大量手动配环境的功夫。 + +```bash +# 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，需要在宿主机上完成： + +1. **系统依赖** + ```bash + apt-get update + apt-get install -y python3 python3-dev python3-setuptools gcc \ + zlib1g-dev build-essential cmake libedit-dev git + ``` + +2. **MACA SDK**（驱动 + maca-sdk + cu-bridge + PyTorch） + 请参考 [Installation_maca](./Installation_maca.md) 的 "MACA SDK" 和 "Install pytorch" 章节完成安装。 + +3. **Python 依赖** + ```bash + 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`） + +```bash +git clone --recursive https://github.com/tile-ai/tilelang-metax.git +cd tilelang-metax +``` + +> 子模块包含 TVM，漏掉 `--recursive` 会导致后续编译报错。 + +### 3.2 配置 git（构建脚本会检查 committer 信息） + +```bash +git config --global user.email "you@example.com" +git config --global user.name "Your Name" +``` + +### 3.3 编译 + +```bash +# 通过环境变量启用 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 包，需要从子模块源码安装： + +```bash +cd 3rdparty/tvm/3rdparty/tvm-ffi && pip install . && cd - +``` + +### 3.5 设置环境变量 + +编译完成后，每次新打开终端都需要导出以下环境（建议写进 `~/.bashrc`）： + +```bash +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： + +```bash +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 专属测试（推荐） + +```bash +# 跑全部 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 跑通用测试（按需） + +```bash +# 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 跑单个测试文件 / 单个用例 + +```bash +# 单个文件 +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 + +```bash +python examples/quickstart.py +``` + +这个脚本会： +1. 定义一个带 ReLU 的 GEMM kernel； +2. 在 MetaX C500 上执行； +3. 与 PyTorch 参考结果做精度对比； +4. 打印 kernel 源码和延迟（示例延迟约 **0.12 ms** 级别）。 + +### 6.2 GEMM 系列样例 + +```bash +# 基础 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` 参数： + +```bash +# 正确 +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`。补救： + +```bash +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 或在社区讨论。