mirror of
https://github.com/deepseek-ai/DeepInfer.git
synced 2026-10-31 00:36:20 +08:00
docs(h20): the compat layer works; end to end is blocked by other tenants on the GPUs
Status table updated: the routing layer exists, all 23 APIs execute with TileLang blocked, and the two backends agree on 18 cases. End to end still has not completed, and the doc now says why with evidence rather than a shrug: six of eight GPUs hold 82 GiB from other containers, so the run OOMs in allocate_userbuffer and the rest wait at a barrier. Also records the command to rerun once the GPUs free up, and two environment traps hit on the way -- /tmp filling with docker data (a bind mount elsewhere means deleting the obvious directory frees nothing) and the compiled _C.so being lost by replacing the source tree.
This commit is contained in:
loopyt
parent
39b086d781
commit
13a63ede19
81
TODO.md
81
TODO.md
@ -595,19 +595,45 @@ import 路径上。修完之后 `import deepinfer` 在纯 Triton 环境不再需
|
||||
把热路径逐个换成 Triton。§2.10b 说明了大头在哪（模型前向直接调 `tile_kernels` 的
|
||||
quant/MoE/MHC/engram，25 处模块级 import / 23 个 API）。
|
||||
|
||||
### 6.0 ⚠️ 当前到哪一步了（先说清楚，别误读成「已能跑」）
|
||||
### 6.0 当前到哪一步了（2026-10-22 晚更新）
|
||||
|
||||
| # | 事项 | 状态 |
|
||||
|---|---|---|
|
||||
| 1 | `import tile_kernels.torch`（无 tilelang） | ✅ 通 |
|
||||
| 2 | `import deepinfer`（无 tilelang） | ✅ 通（lazy init 的副作用，不等于能跑） |
|
||||
| 3 | 算子数值等价（5 个） | ✅ 4 个 bitwise + rope 1 ULP |
|
||||
| 4 | **模型前向实际调用 `tile_kernels.quant.*`** | ❌ **仍然抛 `ModuleNotFoundError: tilelang`** |
|
||||
| 5 | **端到端推理** | ❌ **从未跑过** |
|
||||
| 6 | **性能** | ❓ **未测**（跑不起来就无从测，退化是推测） |
|
||||
| 2 | `import deepinfer`（无 tilelang） | ✅ 通 |
|
||||
| 3 | 算子数值等价（全部 23 个 API 的实际调用） | ✅ 见 6.2 |
|
||||
| 4 | **模型前向调用 `tile_kernels.*`** | ✅ **已通**（`tk_compat` 路由层 + 25 处 import 替换） |
|
||||
| 5 | 无 tilelang 下 23 个 API 冒烟 | ✅ `ALL_CALLS_OK`（唯一 skip 是 `force_random`，见下） |
|
||||
| 6 | compat 两模式逐函数对比 | ✅ **`ALL_MATCH` 18 用例，12 项逐位相同** |
|
||||
| 7 | **端到端推理** | ⚠️ **被 GPU 占用阻塞，不是代码问题**（见 6.4） |
|
||||
| 8 | **性能** | ❓ 仍未测 |
|
||||
|
||||
第 4 行是当前的断点：deepinfer 里 25 处 `import tile_kernels` 指向的还是**编译版**，
|
||||
`tk_compat` 路由层没写，所以调用一进去就要求 tilelang。**「导进去了」和「调得动」是两件事。**
|
||||
**「导进去了」和「调得动」现在都是 ✅**，且两模式在真实调用签名下逐函数一致。
|
||||
|
||||
### 6.4 端到端为什么还没跑完（环境问题，不是代码）
|
||||
|
||||
第 4 次端到端（`DEEPINFER_KERNELS=ref` + 屏蔽 tilelang）**通过了权重加载**（31s，80 GiB）、
|
||||
YOCO/EPLB 初始化，然后卡在 `allocate_userbuffer` 18 分钟无输出。日志里是：
|
||||
|
||||
```
|
||||
torch.OutOfMemoryError: CUDA out of memory. Tried to allocate 540.00 MiB.
|
||||
GPU 0 has a total capacity of 139.80 GiB of which 423.62 MiB is free.
|
||||
```
|
||||
|
||||
查 `nvidia-smi`：**GPU 1/2/4/5/6/7 各被别的容器占了 82 GiB**（`zjr-flWithTree` 等 Up 3 周），
|
||||
只有 GPU 0/3 空闲。8 卡 EP 每卡要约 80 GiB，装不下 → 部分 rank OOM 退出，其余 rank 卡在
|
||||
通信 barrier。**与 ref 实现无关。**
|
||||
|
||||
两次环境干扰记录（下次遇到先查这两个）：
|
||||
- `/tmp`（`/dev/vda`，252G）被 `/tmp/docker-data`（217G）占满 → 容器 `docker exec` 全部失败
|
||||
（`write /tmp/runc-process...: no space left`）。清理 `/etc/dsw-logs` 下轮转日志（同设备）
|
||||
释放 ~2.6G 才恢复。**注意 `/data` 是 bind mount，删它不释放 `/tmp`。**
|
||||
- 丢弃 `deepinfer/` 目录重新同步会**丢掉编译产物 `_C*.so`**（它不在源码里）。
|
||||
可从 `build/nvidia-x86_64/lib.linux-x86_64-cpython-312/deepinfer/_C*.so` 复制回来。
|
||||
|
||||
端到端命令（GPU 空闲时可直接执行）：
|
||||
`PYTHONPATH=scripts/simulate_no_tilelang:... DEEPINFER_SIMULATE_NO_TILELANG=1
|
||||
DEEPINFER_KERNELS=ref DEEPINFER_OPS_STRICT=1 bash scripts/h20_e2e_report.sh`
|
||||
|
||||
### 6.1 第一步已通：`tile_kernels.torch` 能在无 TileLang 下导入
|
||||
|
||||
@ -642,22 +668,37 @@ use_tma_aligned_col_major_sf=True`）：
|
||||
| `quant.cast_back` | **rel 0.000e+00** |
|
||||
| `quant.swiglu_forward_and_per_token_cast` | **bitwise 一致** |
|
||||
| `engram.engram_hash` | **bitwise 一致**（溢出敏感，必须是 int64 池算术） |
|
||||
| `engram.fused_weight` | **bitwise 一致**（fp32 乘，不能在 bf16 里乘再转） |
|
||||
| `engram.engram_gate_fwd` | rel ~3e-7；`inplace=True` 的别名语义也对 |
|
||||
| `transform.apply_rotary` | 65536 元素里 1 个差 1.17/1.33 ULP（bf16 末位舍入），可接受 |
|
||||
| `mhc` 8 个 | 全部 ≤2e-7；`expand`/`normw_merge`/`big_fuse.layer_input` 逐位相同 |
|
||||
| `moe.moe_topk_gate` | **bitwise 一致**（5 组，含 4096×384 / 384 专家 / topk6） |
|
||||
|
||||
**这是整条桥的前提**：量化链 bitwise 相同意味着「先跑通」不会引入 token 级偏差。
|
||||
|
||||
### 6.3 还没做的（下一步，按顺序）
|
||||
两个「看起来对但不对」的判据，都踩过：
|
||||
- 相对误差不能除以**逐元素**预期值。`comb`、`mixes` 天然接近 0，除以它会把一个 2.4e-7 的
|
||||
绝对一致报成 2.6e-3。要用「该行/该张量的最大幅值」作分母。
|
||||
- `mhc_reduce_partials` 的 fp32 累加顺序与 kernel 不同，误差 2.3e-7；对照 fp64 真值才能
|
||||
确认这是表示误差而不是逻辑错（实测 kernel 与朴素 fp32 求和的差距只有 2.09×）。
|
||||
|
||||
1. **路由层（唯一的硬阻塞）**：把 deepinfer 里 25 处 `import tile_kernels` 改成走
|
||||
`deepinfer/tk_compat/`，按有无 tilelang 或 `DEEPINFER_KERNELS=ref` 选编译版还是 `_ref`。
|
||||
最小改法：只换 import 那一行（`import tile_kernels` → `from deepinfer import tk_compat as
|
||||
tile_kernels`），`tile_kernels.quant.xxx(...)` 的调用点一行都不用动。
|
||||
**ref 侧覆盖率已核对**：11 个主 API 里 8 个有直接对应（另 3 个是 `config` 开关，本来就在
|
||||
`tile_kernels.config` 里，不需要 ref）；8 个 `mhc_*_ref` 全在。唯一不明的是
|
||||
`mhc_pre_big_fuse_fwd`（只有个 `reduce_fused` 名字像，未确认）。
|
||||
2. **MHC 8 个 + MoE gate 的数值 A/B**：ref 存在 ≠ 数值等价，还没验过，做之前不能假设。
|
||||
3. **端到端**：路由层 + 上面的 A/B 都过之后才能真正跑，然后再谈 token 一致性。
|
||||
4. **性能**：现在说「退化严重」只是推测，没测过。跑通后才有基准。
|
||||
5. **上游**：以上 3 个补丁应该提到 FlagOS，否则每次重装都要重跑脚本。
|
||||
### 6.3 已完成与剩余
|
||||
|
||||
已完成：
|
||||
1. ~~路由层~~ → `deepinfer/tk_compat/`（config/engram/mhc/moe/quant/torch/transform/utils 八个模块），
|
||||
`DEEPINFER_KERNELS=auto|cuda|ref`，25 处 import 由
|
||||
`scripts/tile_kernels_patch/replace_imports.py` 批量替换（幂等，`--revert` 可回滚）。
|
||||
适配器不是简单转发：`num_per_channels→block_size=(1,n)`、`out=` 缓冲回写、
|
||||
`apply_rotary` 拆 cache + 原地写回、`engram_gate` 参数重排、`big_fuse` 用 ref 组件组合。
|
||||
2. ~~MHC 8 个 + MoE gate 的 A/B~~ → 全绿。
|
||||
3. ~~无 tilelang 冒烟~~ → `ALL_CALLS_OK`（`scripts/.../smoke_compat.py`，在 `/data/ftbuild/repo/`）。
|
||||
4. ~~compat 两模式对比~~ → `ALL_MATCH` 18 用例（`scripts/tile_kernels_patch/compat_ab.py`）。
|
||||
|
||||
剩余：
|
||||
1. **端到端**：被 GPU 占用阻塞（见 6.4），空闲时可直接跑。
|
||||
2. **性能**：仍未测。
|
||||
3. **`force_random` 缺口**：ref 侧没有这个参数，`args.force_load_balance` 打开时
|
||||
`tk_compat.moe` 会**显式报错**而不是静默降级（静默会给出一套看起来正常、但路由不同的 token）。
|
||||
4. **上游**：3 个 tile_kernels 补丁应提到 FlagOS，否则每次重装都要重跑脚本。
|
||||
|
||||
（`flag_gems` 不是替代品：903 个 ATen 级算子替换，没有这些定制融合算子。）
|
||||
|
||||
Loading…
x
Reference in New Issue
Block a user