diff --git a/TODO.md b/TODO.md index 2239e85..0c3a691 100644 --- a/TODO.md +++ b/TODO.md @@ -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()`，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）**，其余都是收尾工作。