forked from metax-maca/op_optimization
Delete quickstart_maca.md
This commit is contained in:
yyyymmm
parent
88f1635344
commit
3ce480c976
|
|
@ -1,304 +0,0 @@
|
|||
# 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 或在社区讨论。
|
||||
Loading…
Reference in New Issue