docs(h20): end to end passes with the reference backend, and is not token-identical

Records the result honestly in both directions. It runs: exit=0, 8 DP ranks agree, answers are
coherent. It is also not equivalent to the compiled path: 16 of 16 samples differ, and a control
run of the compiled backend twice is bit-identical, so the difference is reproducible rather than
sampling noise. One prompt diverges at token 16, the other at token 1.

Lists what could account for it instead of asserting a cause: batched_transpose_weight_sf is the
only path operator never A/B'd and it rewrites every weight's scale factor; apply_rotary is one
bf16 ULP off on 1 element in 65536 and runs every layer; engram_gate is 3e-7. A first-token
divergence argues for comparing prefill logits next rather than guessing further.

Also rewrites the earlier wrong diagnosis: the eight GPUs were not held by colleagues' containers.
They were held by this container's own deadlocked leftovers from a run I killed -- `pkill` took the
parents, `multiprocessing.spawn` children were orphaned to PID 1 and kept the CUDA contexts for two
days, spinning six cores and writing 1.4 GB of retry logs. The lesson worth keeping is the check
itself: map PID to cgroup before blaming another tenant.

TODO.md

@@ -595,45 +595,66 @@ import 路径上。修完之后 `import deepinfer` 在纯 Triton 环境不再需
 把热路径逐个换成 Triton。§2.10b 说明了大头在哪（模型前向直接调 `tile_kernels` 的
 quant/MoE/MHC/engram，25 处模块级 import / 23 个 API）。
 
-### 6.0 当前到哪一步了（2026-10-22 晚更新）
+### 6.0 当前到哪一步了（2026-10-24 晚更新，**端到端已跑通**）
 
 | # | 事项 | 状态 |
 |---|---|---|
 | 1 | `import tile_kernels.torch`（无 tilelang） | ✅ 通 |
-| 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 | **性能** | ❓ 仍未测 |
+| 2 | 算子数值等价（23 个 API 实际调用 + 18 用例两模式对比） | ✅ 见 6.2 |
+| 3 | 路由层 + 25 处 import 替换 | ✅ 已通 |
+| 4 | **无 tilelang 下 23 个 API 冒烟** | ✅ `ALL_CALLS_OK` |
+| 5 | **端到端（`DEEPINFER_KERNELS=ref`）** | ✅ **exit=0，8 DP 一致，回答正确** |
+| 6 | **与编译版逐 token 一致** | ❌ **16/16 样本不同**（见 6.5，已确认非采样噪声） |
+| 7 | 完全不用 tilelang（含 attention） | ❌ sparse MLA 仍只有 TileLang 实现（见 6.6） |
+| 8 | 性能 | ❓ 仍未测 |
 
-**「导进去了」和「调得动」现在都是 ✅**，且两模式在真实调用签名下逐函数一致。
+### 6.5 端到端结果：功能对，但**不逐 token 一致**
 
-### 6.4 端到端为什么还没跑完（环境问题，不是代码）
+`DEEPINFER_KERNELS=ref`（tilelang 仍可导入，只服务 attention）跑 16 个样本（2 prompt × 8 DP）：
 
-第 4 次端到端（`DEEPINFER_KERNELS=ref` + 屏蔽 tilelang）**通过了权重加载**（31s，80 GiB）、
-YOCO/EPLB 初始化，然后卡在 `allocate_userbuffer` 18 分钟无输出。日志里是：
+- `exit=0`，`hard failures: none`，8 个 DP 逐 token 互相一致，回答语义正确
+  （「问题：中国的首都是哪里？答案：北京。」）
+- 但和 `DEEPINFER_KERNELS=cuda` 比：**16/16 样本不同**。
+  - prompt A：前 15 个 token 相同，第 16 个分叉
+  - prompt B：第 1 个 token 就分叉
 
-```
-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.
-```
+**控制实验（关键）**：`cuda` 连跑两次，16 个样本**逐位相同（0 差异）**。所以 harness 在
+贪心解码下是确定性的，上面的差异**不是采样噪声**，是真实数值差异。
 
-查 `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 实现无关。**
+未归因的候选（按可能性）：
+1. `batched_transpose_weight_sf` —— **唯一没做 A/B 的路径算子**，且作用于**所有权重的 scale factor**。
+   尝试补测时我的输入布局不对，编译版直接 device-side assert；需按 `hugging_face/parameters.py`
+   的真实 3-D `(num_groups, -1, n)` 布局重做。
+2. `apply_rotary` 1 bf16 ULP（1/65536 元素）—— 每层都过，会累积。
+3. `engram_gate` 相对 3e-7。
+4. prompt B 第 1 个 token 就分叉，暗示 prefill 阶段已存在实质差异，需要对比首个 token 的 logits
+   才能定位，而不是继续猜。
 
-两次环境干扰记录（下次遇到先查这两个）：
-- `/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` 复制回来。
+### 6.6 还差什么才能真正脱离 tilelang
 
-端到端命令（GPU 空闲时可直接执行）：
-`PYTHONPATH=scripts/simulate_no_tilelang:... DEEPINFER_SIMULATE_NO_TILELANG=1
-DEEPINFER_KERNELS=ref DEEPINFER_OPS_STRICT=1 bash scripts/h20_e2e_report.sh`
+`h20_ops/sparse_mla_sm90.py` 是 `@tilelang.jit` 实现，且 `h20-sm90-triton` 分支上也是如此
+（该分支只是没屏蔽 tilelang，不是真的无依赖）。**SM90 稀疏 MLA 没有 Triton 版本**，这是
+「国产卡完全不用 tilelang」的最后一个硬缺口，工作量远大于这次的 compat 层。
+`tilelang_ops/cuda/` 那 29 个模块已有 `tilelang_ops/triton/` 对应版本，靠
+`DEEPINFER_OPS_BACKEND=triton` 切换，不是缺口。
+
+### 6.4 端到端曾经为什么跑不起来（环境 + 排查记录）
+
+**先说结论：10-22 我判断「GPU 被同事的容器占用」是错的。** 当时只看 `nvidia-smi` + `docker ps`
+列表就推断，没做 PID→cgroup 映射。10-24 复查发现那 6×82 GB 是**我自己 `h20-ft-build` 容器**
+里两天前 e2e 被 `pkill` 后遗留的死锁进程：
+
+- `state=R`，CPU 时间 ≈ 存活时间（2 天 1 小时）→ 每个进程 100% 空转（NCCL barrier 自旋）
+- 日志实证 `TCPStore server has shut down` + `Broken pipe` 无限重试 → 永不恢复
+- 副作用：`e2e-ref.log` 被刷到 **1.4 GB**，6 个 CPU 核白烧 2 天
+- 根因：`pkill -f h20_ask` 只杀父进程，`multiprocessing.spawn` 的 worker 被 PID 1 收养，
+  **继续持有 CUDA context**。清完父进程后还有 6 个 `state=S` 的子进程持有 nvidia fd，必须一起清
+
+判断规则：**看 GPU 被谁占用必须走 PID→cgroup→容器名**，不能靠「谁 Up 得久」。
+`nvidia-smi --query-compute-apps=pid` 拿 PID → `cat /proc/<pid>/cgroup`。
+
+另一个坑：`/tmp`（`/dev/vda` 252G）会满。`/data` 是 **bind mount，删它不释放 `/tmp`**；
+清 `/etc/dsw-logs` 同设备轮转日志才有效。
 
 ### 6.1 第一步已通：`tile_kernels.torch` 能在无 TileLang 下导入