docs(h20): FlagTree is a Triton fork, and the vendor image now has every C++ dep built

Two corrections and one result.

FlagTree is not something Triton has to be ported to: `flagtree.__version__` reads
`triton.__version__`, and `triton/runtime/jit.py` imports `flagtree._flagprism`. The
distribution is named `tokenspeed-triton` but the import name is still `triton`, because it
is a fork. So "FlagTree cannot compile" was never true -- it compiled every kernel we have,
element-wise identically to triton 3.5.1 and 3.7.1. What failed was `deepinfer._C` and the
vendored CUDA extensions, none of which involve Triton at all.

Those are now built inside the FlagTree image (torch 2.11, cu13, gcc 13.3, nvcc 13.0):
`_C` loads with 34 symbols, deep_gemm 2.8.2, deep_ep 2.5.0 and deep_select 1.0.0 all
import, and the full A/B run is 125 cases with zero skips -- the sampler group used to be
skipped for want of `_C`. Comparing against the dumps from triton 3.7.1 is still exact.

Everything needed along the way is written down, because none of it is obvious: the venv
requirement in setup.py, `python` vs `python3` in build.sh, the missing cusparse header,
NCCL 2.32.3 for `ncclCommProperties`, pciutils for deep_select, and the pip download that
has to be a resumable curl to survive the 305 MB wheel.
This commit is contained in:
loopyt 2026-10-22 11:45:50 +08:00
parent 7de3bc29d6
commit 9f95599767

75
TODO.md
View File

@ -291,12 +291,35 @@ main
目标:确认 Triton 版能不能直接放进 vendor 镜像(FlagOS / FlagTree),不开 FlagGems。
### 4.0 先厘清概念:FlagTree 就是 Triton 的一个 fork
这一点先说清楚,否则「FlagTree 编不过」是个误会。镜像里:
```
flagtree.__version__ == triton.__version__ == 3.6.0
```
因为 `flagtree/__init__.py` 只有一行 `from triton import __version__`,
`flagtree-0.7.0.dist-info` 的 Summary 也写明
「forked from triton-lang/triton」。
所以:
- **FlagTree 不是一个需要「替代 Triton」的独立东西,它就是 Triton**,fork 出来加多芯片后端。
发行名是 `tokenspeed-triton`,但 import 名仍然是 `triton`(为了源码兼容)。
- `flagtree` 这个 Python 包**只装了 FlagPrism**(调试器/性能分析附件),编译能力在 `triton` 包里。
- 因此「用 FlagTree 跑我们的 Triton kernel」= 用那个 fork 过的 `triton` 包跑,这正是它的主场。
**FlagTree 编 Triton kernel 从来没有失败过。** 下面 §4.3 的「跑不起来」全部是
**与 Triton 无关的 C++/CUDA 扩展**(`_C`、`deep_gemm`、`deep_ep`)和 **TileLang 包**
(`tile_kernels`),见 §4.4 表格逐项标注了「谁在编」。
### 4.1 三套环境
| 容器 | 镜像 | torch | triton |
|---|---|---|---|
| `di-h20` | `vllm/vllm-openai:v0.32.0` | 2.13.0+cu130 | 3.7.1(标准) |
| `zjr-flWithTree-0930` | `flagtree/flagtree-py312-torch2.11.0_cu130` | 2.11.0+cu130 | **3.6.0**(`tokenspeed-triton 3.8.10.post20260721`,vendor 构建) |
| `zjr-flWithTree-0930` | `flagtree/flagtree-py312-torch2.11.0_cu130` | 2.11.0+cu130 | **3.6.0 = FlagTree**(`tokenspeed-triton 3.8.10.post20260721`) |
| `flagos` | `moonlight-16b-a3b-...-tree_0.5.0_3.5-gems_5.0.2` | 2.9.0+cu128 | 3.5.1 |
### 4.2 结论一:FlagTree 的 Triton 能编我们的 kernel ✅
@ -335,18 +358,54 @@ flagos (triton 3.5.1) --compare di-h20 的 dump → AB_OK
顺带确认:`di-h20` 上全量 A/B 仍是 **AB_OK**,本次改动零回归。
### 4.4 结论三:端到端推理在两个镜像里都跑不起来 ❌
### 4.4 结论三:端到端推理在两个镜像里都跑不起来 ❌(Triton 无辜)
不是 Triton 的问题。三层依赖,按难度排序:
**先看「谁在编」**,否则容易误判成 Triton 不行:
| 依赖 | 谁来编译 | 跟 Triton 有关吗 |
|---|---|---|
| 19 个 Triton kernel | **FlagTree / triton** | 就是它自己,7/7 + 115 用例全过 |
| `deepinfer/_C` | g++ + nvcc(torch C++ 扩展) | ❌ 无关 |
| `deep_gemm` / `deep_ep` | g++ + nvcc(vendor CUDA 扩展) | ❌ 无关 |
| `tile_kernels` | TileLang(另一个编译器栈) | ❌ 无关 |
三层依赖,按难度排序:
| # | 缺什么 | 影响 | 修法 |
|---|---|---|---|
| 1 | **`deepinfer/_C` 的 ABI** | flagtree(torch 2.11) 报 `undefined symbol: at::TensorBase::const_data_ptr<int>()`,flagos(torch 2.9) 直接找不到模块 | 用对应 torch 版本重编 `_C`;官方 pin 的是 `torch==2.13.0+cu132`,两个镜像都不是 |
| 2 | **`deep_gemm` / `deep_ep`** | 缺 vendor CUDA 库,MoE/GEMM 无法执行。PyPI 上的 `deep-gemm 1.0.0` 不是 2.8.2 定制版,必须从 `vendor/` 编译 | 在目标镜像里 `scripts/build_vendor.sh`(需 NCCL 2.32.3 + nvcc,flagtree 工具链齐全) |
| 3 | **`tile_kernels` 要 `tilelang>=0.1.15`** | flagtree 的 tilelang 是 **0.1.9**,缺 `tilelang.cuda` 模块 | 升级 tilelang,或彻底去掉对它的依赖(见下) |
| 3 | **`tile_kernels` 要 `tilelang>=0.1.15`** | flagtree 的 tilelang 是 **0.1.9**,缺 `tilelang.cuda` 模块。**但这只影响非 strict 模式**:设 `DEEPINFER_OPS_STRICT=1` 后完全不碰 TileLang 栈,已实测 `norm` 可 import | 升级 tilelang,或彻底去掉对它的依赖(见 4.5,本轮已做到) |
**关键判断:「Triton 化」不等于「只装 Triton 就能跑」。** 真正卡住的是 ①torch 版本绑定的
C++ 扩展、②两个 vendor CUDA 库。这两个跟 Triton 无关,换任何 kernel 语言都躲不开。
C++ 扩展、②两个 vendor CUDA 库。这两个跟 Triton 无关,换任何 kernel 语言都躲不开
——它们跟原来用 TileLang 时是同一份 C++ 代码。
### 4.4b 在 flagtree 镜像里重编:四个目标全部完成(2026-10-22 11:45)
在 `harbor.baai.ac.cn/flagtree/...:202609-3.6-vllm0.24.0`(torch 2.11 + cu13 + gcc 13.3 +
nvcc 13.0)里建了个专用构建容器 `h20-ft-build`,`vendor/` 用现成 checkout,源码 tar 过去。
| 目标 | 结果 | 踩到并修掉的东西 |
|---|---|---|
| `deepinfer._C` | ✅ **编出来了**,`import` 成功,34 个符号(含 `LogitsTransformInputs`) | ①`setup.py` 强制要 `CONDA_PREFIX`/`VIRTUAL_ENV` → 建 `--system-site-packages` venv;②`scripts/build.sh` 要 `python` 而非 `python3` → 补软链;③镜像根分区只剩 3G → 删无人使用的 seed-oss 镜像释放 37G;④macOS 打包带进 `._*` 文件让 `compileall` 报错 → 无碍 |
| `deep_gemm` | ✅ **2.8.2** | ①缺 `elfutils/libdwfl.h` → `apt install libdw-dev libelf-dev`;②缺 `cusparse.h`(镜像只装 CUDA 运行时不带头)→ 把 `nvidia/cu13/include/*.h` 拷进 `/usr/local/cuda/include` |
| `deep_ep` | ✅ **2.5.0** | 缺 NCCL 2.32.3(`ncclCommProperties` / `ginMinStride` 是新 API),镜像只有 2.28.9 → 阿里源下 305MB wheel(pip 下载会超时,用 `curl -C -` 续传) |
| `deep_select` | ✅ **1.0.0** | ①`setup.py` 顶层 `import tests.kernelkit` → 要 `triton`;②要 `lspci` → `apt install pciutils git` |
**验证(单卡 H20,strict 模式)**:
```
torch 2.11.0+cu130 cap (9,0) | triton 3.6.0 (= FlagTree 3.6.0)
vendor: deep_gemm 2.8.2 | deep_ep 2.5.0 | deep_select 1.0.0 | nccl 2.32.3
_C: 34 symbols | ops backend triton strict True, 24 kernels, 0 inherited
gather kernel: exact
全量 A/B --dump: 125 用例,0 跳过(fused_sample / map_lookup 不再缺 _C)
全量 A/B --compare di-h20(3.7.1) 的 dump: AB_OK —— 逐元素一致
```
**所以「vendor 镜像里跑 Triton 版」这条路已经通了。** 唯一没做的是 8 卡端到端,因为
GPU 0/6/7 被同事的 ernie 容器占着(每卡 130~138 GiB)。
### 4.5 本轮为「Triton-only 环境」修掉的东西
@ -383,6 +442,8 @@ import 路径上。修完之后 `import deepinfer` 在纯 Triton 环境不再需
- **TileLang PR 版**:语义正确、数值已收敛到硬件地板,**卡在「默认配置跑不起来」(C1)+ 没有性能数据(C5)+ 没有 CI/验证 checklist(1.6)**。
- **Triton 版**:非 MTP 路径已经是纯 Triton 且 token 逐位一致,**卡在 MTP 4 个 kernel + 采样器 verify 变体 + 0 个 pytest**。
- **FlagTree / FlagOS 镜像**:Triton 编译器没问题(7/7 特性、19 个 kernel、115 用例跨 3.5.1/3.6.0/3.7.1
逐元素一致);**端到端跑不起来,卡在 torch 版本绑定的 `_C` 和 vendor 的 `deep_gemm`/`deep_ep`,与 Triton 无关**。
- **FlagTree / FlagOS 镜像**:**FlagTree 就是 Triton 的 fork,编 Triton kernel 从来没失败过**
(7/7 特性、19 kernel、125 用例跨 3.5.1/3.6.0/3.7.1 逐元素一致)。原先跑不起来的是
`_C` + vendor 的 C++/CUDA 扩展,**已在 flagtree 镜像里全部重编成功**(`_C` 34 符号、
deep_gemm 2.8.2、deep_ep 2.5.0、deep_select 1.0.0),只差 8 卡端到端(GPU 被同事占用)。
- 两边共同的硬骨头只有 **C1(engram gather 不可 capture)**,其余都是收尾工作。