ADD file via upload

2026-05-15 15:33:16 +08:00 · 2026-05-15 15:33:16 +08:00 · 760df0e5e8
parent a9d2674b1b
commit 760df0e5e8
1 changed files with 304 additions and 0 deletions
--- a/quickstart_maca.md
+++ 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 或在社区讨论。