From 5b080b17803989f023297bd1346bf29eb60c4bc2 Mon Sep 17 00:00:00 2001
From: loopyt <zhsong@baai.ac.cn>
Date: Mon, 26 Oct 2026 15:39:13 +0800
Subject: [PATCH] docs(h20): a runbook for the FlagGems configuration, plus the
 Dockerfile that reproduces it

Everything a second person needs to get from a clean H20 host to the verified result: how to start the
container, how to start inference, which FlagGems operators are excluded and on what evidence, and
seven test cases each with a pass criterion. The image is published at
`harbor.baai.ac.cn/flagrelease-public/deepseek-v4.1-flash-h20-gems5.0.2-...:580.95.05`
(digest sha256:1d249feb...), 46.5 GB, and was re-verified by running the full end-to-end inside it --
same 8/8 ranks, same tokens, `hard failures: none`.

The Dockerfile is not a `docker commit` of the working container, and it could not be: `/data` is a
bind mount on the build host, so the venv and the source tree are invisible to commit. They are tarred
in instead. Three details are load-bearing and commented in place:

* the tree moves `/data/ftbuild` -> `/opt/ftbuild`. Keeping the original paths would collide with the
  `/data` model mount at run time and shadow both venv and source. Every absolute path recorded in the
  venv has to be rewritten, and `bin/python*` is a symlink into /usr/bin that `sed -i` would turn into
  a patched copy and break, so the rewrite is `find -type f`.
* Triton's own CUDA 12 runtime is copied to a fixed path and put on `LD_LIBRARY_PATH`. Triton
  downloads `libcudart.so.12` and does not add it to the search path; anything loading a Triton kernel
  fails without it, `import flag_gems` included.
* `flag_gems` and `deepinfer` cannot be imported during the build -- flag_gems probes for a device and
  `docker build` has no GPU ("No device were detected on your machine"). The build verifies the
  relocation offline instead, and the runbook covers what needs a GPU.

`DEEPINFER_OPS_STRICT=1` is set in the image on purpose. Without it the router imports the compiled
backend, finds 16 of 29 kernels unported, and prints `falling back to 'cuda'` -- which reads as
"needs the compiled path" even though those 16 are never called. Measured both ways in the image: with
STRICT the same run passes, `kernels served by the default backend` is empty, and the tokens are
identical. Strict is both the honest reading and a louder failure if a missing kernel ever is called.

Also corrected in this commit rather than left standing: an earlier note in this branch claimed "0
fallbacks" from a run that had STRICT unset, where the fallback line was merely suppressed from the
summary.
---
 H20-FLAGGEMS-RUNBOOK.md        | 267 +++++++++++++++++++++++++++++++++
 deploy/h20-flaggems/Dockerfile | 107 +++++++++++++
 deploy/h20-flaggems/build.sh   |  48 ++++++
 3 files changed, 422 insertions(+)
 create mode 100644 H20-FLAGGEMS-RUNBOOK.md
 create mode 100644 deploy/h20-flaggems/Dockerfile
 create mode 100755 deploy/h20-flaggems/build.sh

diff --git a/H20-FLAGGEMS-RUNBOOK.md b/H20-FLAGGEMS-RUNBOOK.md
new file mode 100644
index 0000000..8479004
--- /dev/null
+++ b/H20-FLAGGEMS-RUNBOOK.md
@@ -0,0 +1,267 @@
+# H20 / SM90 + FlagGems：运行与复现手册
+
+DeepSeek-V4.1-Flash 跑在 8×H20-3e（SM90）上，稀疏 MLA 与引擎 kernel 走 Triton，
+可选启用 FlagGems。本文是唯一需要照着做的东西：起容器、起服务、看结果、判断是否正常，
+以及每个"为什么这么设"的原因。
+
+- 代码分支：`h20-sm90-flaggems`
+- 镜像（46.5 GB）：
+  ```
+  harbor.baai.ac.cn/flagrelease-public/deepseek-v4.1-flash-h20-gems5.0.2-tree3.6.0-cxnone-plugin0.0.0-vllm0.24.0-cp312-pt211-cu130-x64:580.95.05
+  ```
+  digest `sha256:1d249feb594d94bab0b94f90a4ebdc7fe6571b47003eef47583483332ba44093`
+- 权重**不在镜像里**（476 GB），运行时挂载
+
+---
+
+## 1. 已验证的事实
+
+这一版不是"能起来"就算完，下面每条都有对应测法（§6）。
+
+| 项 | 值 |
+|---|---|
+| DP 一致 | **8/8 rank 产生同一 token 序列** |
+| kernel 回退 | **无**（`kernels served by the default backend` 为空，且开了 `DEEPINFER_OPS_STRICT=1`） |
+| 硬错误 | **none** |
+| 回答正确性 | 中国首都→北京、日本首都→东京，语义连贯 |
+| FlagGems 实际派发 | **19 个算子、152 处调用点** |
+| FlagGems 注册 | 37 个请求名 → 96 个变体 → 43 个实际注册 |
+| 确定性 | 同配置连跑两次**逐 token 完全相同** |
+| 镜像内复跑 | 从发布镜像起容器重跑，`exit=0`，token 与本地开发容器**完全一致** |
+
+关于 `DEEPINFER_OPS_STRICT=1`：不设它时，路由会去 import 编译版后端，发现 29 个 kernel 里有 16 个
+没移植，于是打印 `falling back to 'cuda'`——**看起来像依赖编译版，其实那 16 个从未被调用**。
+实测：开 STRICT 后同一次端到端照样通过、`kernels served by the default backend` 为空、token 完全相同。
+所以镜像默认开 STRICT，既如实反映"没走回退"，也能让将来真有缺失时立刻报错而不是被悄悄替换。
+
+镜像内软件版本：
+
+| 组件 | 版本 |
+|---|---|
+| Python | 3.12.3 |
+| torch | 2.11.0+cu130 |
+| Triton | 3.6.0 |
+| CUDA（runtime） | 13.0 |
+| flag_gems | 5.0.2 |
+| tilelang | 0.1.15（**本路径不需要**，仅保留以便对照） |
+| 驱动 / GPU | 580.95.05 / NVIDIA H20-3e，143771 MiB ×8 |
+
+---
+
+## 2. 起容器
+
+```bash
+export PATH=/mnt/workspace/docker-binary/docker:$PATH   # 这台机器的 docker 不在默认 PATH
+
+docker run -d --name h20-fg \
+  --gpus all \
+  --ipc=host --network=host \
+  -v /data/models:/data/models \
+  --entrypoint sleep \
+  harbor.baai.ac.cn/flagrelease-public/deepseek-v4.1-flash-h20-gems5.0.2-tree3.6.0-cxnone-plugin0.0.0-vllm0.24.0-cp312-pt211-cu130-x64:580.95.05 \
+  infinity
+```
+
+为什么这几项：
+
+- `--ipc=host` + `--network=host`：8 个 DP rank 之间走 host IPC 与 NCCL，容器内默认的
+  64 MB `/dev/shm` 不够。
+- `-v /data/models`：权重 476 GB，只能挂。镜像里已把 venv 与源码放在 `/opt/ftbuild`，
+  正是为了避免和这个 `/data` 挂载冲突（见 §7）。
+- `--entrypoint sleep ... infinity`：镜像默认 `CMD bash`，用 `sleep infinity` 让它常驻，
+  后续所有命令用 `docker exec` 进。
+
+进入：
+
+```bash
+docker exec -it h20-fg bash
+# 镜像已设好 PATH/PYTHONPATH/LD_LIBRARY_PATH 与全部 DEEPINFER_* 默认值
+cd /opt/ftbuild/repo && source /opt/ftbuild/venv/bin/activate
+```
+
+---
+
+## 3. 起服务（推理）
+
+最短的一条：单条 prompt、贪心、8 个 token。启动约 3.5 分钟（加载 476 GB 权重）。
+
+```bash
+docker exec h20-fg bash -lc '
+  cd /opt/ftbuild/repo && source /opt/ftbuild/venv/bin/activate
+  export DEEPINFER_HF_MODEL_DIR=/data/models/DeepSeek-V4.1-Flash-1017
+  python scripts/h20_ask.py "问题：中国的首都是哪里？\n答案：" --max_tokens 8
+'
+```
+
+完整带判据的一次跑（推荐，输出被整理成"该看什么"）：
+
+```bash
+docker exec h20-fg bash -lc '
+  cd /opt/ftbuild/repo && source /opt/ftbuild/venv/bin/activate
+  export DEEPINFER_HF_MODEL_DIR=/data/models/DeepSeek-V4.1-Flash-1017
+  export LOG=/tmp/e2e.log
+  bash scripts/h20_e2e_report.sh
+'
+```
+
+`h20_e2e_report.sh` 只打印四件事：token id 去重结果（几个 rank 一致）、
+每条 prompt 的续写、默认后端实际服务的 kernel（应为空）、硬错误（应为 none）。
+
+### 环境变量（镜像已内置，改一处即可对照）
+
+| 变量 | 本镜像默认 | 作用 |
+|---|---|---|
+| `DEEPINFER_SPARSE_MLA_BACKEND` | `triton` | 稀疏 MLA 用 FlagGems 的 Triton kernel；设 `tilelang` 走原实现 |
+| `DEEPINFER_OPS_BACKEND` | `triton` | 引擎 29 个 kernel 的 Triton 版本 |
+| `DEEPINFER_KERNELS` | `ref` | `tile_kernels` 走 torch 参考（不依赖 TileLang） |
+| `DEEPINFER_OPS_STRICT` | `1` | 禁止 kernel 回退；见 §1 末尾说明 |
+| `DEEPINFER_ENABLE_FLAGGEMS` | `1` | 启用 FlagGems（白名单模式） |
+| `DEEPINFER_FLAGGEMS_MODE` | `allow` | `allow` = 只开验证过的；`deny` = 全开但排除坏算子 |
+| `DEEPINFER_USE_NEXTN_DECODE` | `0` | 关投机解码：其 rejection-sampling 是唯一仍要 TileLang 的 kernel |
+| `EP_DISABLE_GIN` | `1` | 容器内无 `/dev/infiniband`，单机走 scale-up |
+| `DEEPINFER_ENFORCE_EAGER` | `1` | Engram 局部 gather 在 CUDA graph capture 里做 D2H |
+
+**想证明"完全不需要 TileLang"**（镜像里仍装了它，只为对照）：
+
+```bash
+docker exec h20-fg bash -lc '
+  cd /opt/ftbuild/repo && source /opt/ftbuild/venv/bin/activate
+  export PYTHONPATH=/opt/ftbuild/repo/scripts/simulate_no_tilelang:/opt/ftbuild/repo
+  export DEEPINFER_SIMULATE_NO_TILELANG=1 DEEPINFER_SPARSE_MLA_BACKEND=triton
+  export DEEPINFER_HF_MODEL_DIR=/data/models/DeepSeek-V4.1-Flash-1017 LOG=/tmp/e2e-notl.log
+  bash scripts/h20_e2e_report.sh
+'
+```
+
+预期：日志里 8 条 `sparse MLA: TileLang unavailable (...); using the Triton backend`、
+8 条 `sparse MLA backend: triton`、0 个 `ModuleNotFoundError` 导致失败。
+
+---
+
+## 4. FlagGems：开了哪些、关了哪些、为什么
+
+启用方式是 FlagGems 官方文档的 `only_enable(include=[...])`
+（`docs/content/en/usage/selective.md`），在 `deepinfer/flaggems_bootstrap.py` 里实现，
+由 `examples/_config.py:create_llm()` 在 `setup_cuda_visible_devices()` **之后**调用。
+
+### 4.1 为什么不"全开"
+
+先按 `enable()` 全开做过 5 轮，每轮在完整引擎里撞出一个新问题，每轮约 5 分钟：
+
+`cat` → `mul` → `topk` → `rms_norm` → `mm`
+
+这条路的性质是"必须先被坑到才能排除"，且**中间任何状态都是不安全的**（可能静默算错）。
+白名单相反：没验证的算子留在 PyTorch，**正确性天然保证**，列表不完整只影响性能。
+
+### 4.2 排除清单（`_BROKEN_OPS`）
+
+| 算子 | 症状 | 后果 |
+|---|---|---|
+| **`mul`** | **广播时返回第一个操作数的形状**：`(1,64,8)*(4,1,8)` → `(1,64,8)`（应为 `(4,64,8)`）；`(64,1)*(1,128)` → `(64,1)`。静默、不报错、少算元素。同形状与 scalar 精确。`add/sub/div/maximum/minimum` 广播都正确，是 `mul` 自身缺陷 | 引擎里 `engram_hash_ref` 结果偏离 PyTorch |
+| `cumsum` | 512 元素 bf16 行偏差最大 **1.0**（容差 ~0.17），bf16 累加 | 数值漂移 |
+| `rms_norm` | PyTorch 允许省略 `weight`，FlagGems 强制必填（`TypeError`）；带 weight 仍差 ~3e-2 | 直接崩或算错 |
+| `topk` | 编译时间随 k 爆炸：k=8→4.8 s，k=64→19 s，k=512→>180 s。采样用 **k=4096** | 启动时永不返回（8 rank 100% CPU、GPU 全空） |
+| `cat` | UE8M0 scale 张量上 `KeyError: 'float8_e8m0fnu'`（Triton 无 e8m0 类型） | 权重加载即挂 |
+
+`mm` 家族不属于本清单，但也不在白名单：隔离探测（M/N/K 扫描、NT 布局、batch、addmm）全部正确，
+引擎里仍撞 `accumulator shape is incompatible`，触发条件未定位。GEMM 留 cuBLAS——本来也更快。
+
+### 4.3 白名单（`_VERIFIED_OPS`，37 个名字）
+
+```
+abs      add       amax        bincount    bitwise_and  clamp      cos       div
+eq       exp       gather      ge          gelu         index_select  lt      masked_fill
+maximum  mean      minimum     neg         nonzero      pow        reciprocal  relu
+remainder repeat_interleave    rsqrt       sigmoid     silu         sin       _softmax
+sort     sqrt      stack       sub         sum          where
+```
+
+两点容易踩：
+
+- **`_softmax`，不是 `softmax`**。FlagGems 用 aten 内部名注册。同理 `eq.Scalar`、`lt.Tensor`、
+  `where.self`、`repeat_interleave.self_int`、`div.true_divide`、`pow.pow_scalar` 都是带后缀的。
+  只写裸名会**静默什么都不注册**。bootstrap 里 `_expand()` 对着 `flag_gems._FULL_CONFIG` 做前缀展开，
+  匹配不到会打印警告。
+- **`bitwise_xor` 和 `square` 在 flag_gems 5.0.2 里根本没实现**。筛查脚本会把这类算子标成
+  `skip（PyTorch answered）`，而不是 `ok`——否则会得出"FlagGems 这个算子是对的"的错误结论。
+
+### 4.4 实际在跑的 19 个
+
+```
+abs  amax  bincount  clamp(clamp_min/clamp_/clamp)  cos  exp  gather  index_select
+neg  nonzero  reciprocal  relu  rsqrt  sigmoid  sin  sqrt  stack
+```
+
+都是点状/reduction 算子——正是 eager PyTorch 一次一个 launch 吃亏、Triton 能合并的地方。
+
+---
+
+## 5. 想改白名单
+
+```bash
+# 加算子（会被前缀展开，不必写全变体）
+export DEEPINFER_FLAGGEMS_INCLUDE=layer_norm
+# 减算子
+export DEEPINFER_FLAGGEMS_EXCLUDE=gelu
+# 退回"全开+排除"模式（不推荐，见 §4.1）
+export DEEPINFER_FLAGGEMS_MODE=deny
+# 记录实际派发了哪些算子（每 rank 一个文件，避免互相覆盖）
+export DEEPINFER_FLAGGEMS_RECORD_DIR=/tmp/fgops
+```
+
+**加算子前必须先过 §6 的 TC-4。** 只把名字加进去、不做 A/B，等于把"可能静默算错"引进来。
+
+---
+
+## 6. 测试用例
+
+每条都给出命令与判据，不需要额外脚本。
+
+| # | 用例 | 命令 | 通过判据 |
+|---|---|---|---|
+| TC-1 | 语义冒烟 | `python scripts/h20_ask.py "问题：中国的首都是哪里？\n答案：" --max_tokens 8` | 续写为"北京。"，不是乱码 |
+| TC-2 | 完整端到端 | `bash scripts/h20_e2e_report.sh` | `exit=0`；DP 一致 `8 [..]`；`kernels served` 为空；`hard failures: none` |
+| TC-3 | 无 TileLang 复现 | §3 最后的 `DEEPINFER_SIMULATE_NO_TILELANG=1` 命令 | 同 TC-2，且 8 条 `backend: triton` |
+| TC-4 | 算子等价性筛查 | `python scripts/flaggems_port/screen_ops.py` | 输出里没有 `BAD`；`not served` 只出现在 `bitwise_xor_`/`sqr` |
+| TC-5 | FlagGems 是否真生效 | TC-2 加 `DEEPINFER_FLAGGEMS_RECORD_DIR=/tmp/fgops`，然后 `cat /tmp/fgops/*.log \| wc -l` | > 0，且 §4.4 的 19 个算子出现在其中 |
+| TC-6 | per-token 量化等价 | `python scripts/flaggems_port/ab_quant.py` | `ALL_MATCH`，量化值 bitwise 相同 |
+| TC-7 | 确定性 | TC-2 跑两次，比对两次的 token id | 逐 token 相同 |
+
+耗时参考：TC-2/TC-3/TC-7 各约 4–5 分钟（含加载权重）；TC-4/TC-6 数秒。
+
+---
+
+## 7. 已知限制（都是实测，不是猜测）
+
+1. **与不开 FlagGems 的版本不逐 token 一致**：第 1 个 token 就分叉。但同配置连跑两次完全相同，
+   所以是 Triton 与 eager PyTorch 的累加顺序差异被贪心解码放大，**不是不确定性**。
+   回答仍然正确连贯。SM90 本身已有 2.7% logits 漂移的旧账（见 `H20-DRIFT-ANALYSIS.md`）。
+2. **性能没测**。本文所有结论都是"能跑、数值对"，没有一条是"更快"。白名单里只有 19 个算子
+   真被用到，收益未知。
+3. **投机解码关着**（`DEEPINFER_USE_NEXTN_DECODE=0`）。开着会走 `fused_sample` 的
+   rejection-sampling 变体，那是唯一仍依赖 TileLang 的 kernel（要复刻它的 Philox 随机流）。
+4. **`mm` 家族留在 cuBLAS**，原因见 §4.2 末尾。
+
+---
+
+## 8. 重建镜像
+
+`/data` 是网络盘、bind mount，`docker commit` 带不走 venv 和源码，所以镜像是用 Dockerfile
+**从 tar 装进去**的，不是 commit 出来的：
+
+```bash
+bash deploy/h20-flaggems/build.sh            # 只构建
+bash deploy/h20-flaggems/build.sh --push     # 构建并推送
+```
+
+Dockerfile 里有两处不是随手写的，改动前请先读注释：
+
+- **把 `/data/ftbuild` 迁到 `/opt/ftbuild`** 并重写 venv 里所有绝对路径（`bin/*` 的 shebang、
+  `pyvenv.cfg`、`.pth`）。不迁会和运行时挂载的 `/data` 冲突，venv 与源码会被遮蔽。
+  `bin/python*` 是指向 `/usr/bin` 的软链，`sed -i` 会把它变成打了补丁的副本从而弄坏解释器，
+  所以重写用 `find -type f`。
+- **把 Triton 自带的 CUDA 12 runtime 拷到 `/opt/ftbuild/cudart` 并加进 `LD_LIBRARY_PATH`**。
+  Triton 下载了自己的 `libcudart.so.12` 却不加入搜索路径，任何加载 Triton kernel 的动作
+  （包括 `import flag_gems`，它会拉 TLE）都会报 `libcudart.so.12: cannot open shared object file`。
+
+镜像里不包含权重，运行时挂 `-v /data/models:/data/models`。
diff --git a/deploy/h20-flaggems/Dockerfile b/deploy/h20-flaggems/Dockerfile
new file mode 100644
index 0000000..291da51
--- /dev/null
+++ b/deploy/h20-flaggems/Dockerfile
@@ -0,0 +1,107 @@
+# DeepSeek-V4.1-Flash on 8x H20 (SM90), sparse MLA and engine kernels on Triton, FlagGems enabled.
+#
+# Two things are in here that the base image does not have, and both used to live on a network
+# volume that `docker commit` cannot capture (`/data` is a bind mount on the build host):
+#
+#   /opt/ftbuild/venv   the Python environment: torch 2.11.0+cu130, triton 3.6.0, flag_gems 5.0.2,
+#                       tilelang 0.1.15, and the compiled deep_gemm / deep_ep / deep_select
+#   /opt/ftbuild/repo   the DeepInfer source tree, including the prebuilt deepinfer/_C .so
+#
+# They are moved from /data/ftbuild to /opt/ftbuild on purpose. Keeping the original paths would
+# collide with the bind mount of the model directory at run time (`/data`), which would shadow the
+# venv and the source. The move means every absolute path baked into the venv has to be rewritten --
+# `bin/*` console scripts carry `#!/data/ftbuild/venv/bin/python` shebangs, and a `.pth` file points
+# at the source tree. `bin/python*` is a symlink into /usr/bin and must not be touched by sed, which
+# is why the rewrite is restricted to `find -type f`.
+#
+# Model weights are NOT in the image (476 GB). Mount them at run time; see H20-FLAGGEMS-RUNBOOK.md.
+#
+# Build:  bash deploy/h20-flaggems/build.sh
+# Run:    see H20-FLAGGEMS-RUNBOOK.md
+
+ARG BASE_IMAGE=harbor.baai.ac.cn/flagtree/flagtree-py312-torch2.11.0_cu130-cuda13.0-ubuntu24.04:202609-3.6-vllm0.24.0
+
+FROM ${BASE_IMAGE}
+
+LABEL org.opencontainers.image.title="DeepSeek-V4.1-Flash H20 SM90 FlagGems" \
+      org.opencontainers.image.description="Sparse MLA + engine kernels on Triton, FlagGems allow-list enabled, no TileLang required" \
+      org.opencontainers.image.source="https://github.com/deepseek-ai/DeepInfer" \
+      org.opencontainers.image.licenses="MIT"
+
+COPY venv.tar.gz /tmp/venv.tar.gz
+COPY repo.tar.gz /tmp/repo.tar.gz
+
+RUN set -eux; \
+    mkdir -p /opt/ftbuild; \
+    tar xzf /tmp/venv.tar.gz -C /opt/ftbuild; \
+    tar xzf /tmp/repo.tar.gz -C /opt/ftbuild; \
+    rm -f /tmp/venv.tar.gz /tmp/repo.tar.gz
+
+# Rewrite the creating prefix everywhere the venv recorded it. `-type f` skips bin/python*, which is
+# a symlink to /usr/bin/python3.12; sed -i would replace the symlink with a patched copy and break
+# the interpreter.
+RUN set -eux; \
+    find /opt/ftbuild/venv/bin -maxdepth 1 -type f -exec sed -i 's|/data/ftbuild|/opt/ftbuild|g' {} +; \
+    find /opt/ftbuild/venv -name '*.pth' -type f -exec sed -i 's|/data/ftbuild|/opt/ftbuild|g' {} +; \
+    sed -i 's|/data/ftbuild|/opt/ftbuild|g' /opt/ftbuild/venv/pyvenv.cfg
+
+# Triton downloads its own CUDA 12 runtime and does not put it on the search path. Anything that
+# loads a Triton kernel -- including `import flag_gems`, which pulls in TLE -- fails without it:
+#   OSError: libcudart.so.12: cannot open shared object file
+# Copied to a fixed location rather than referencing the versioned download directory, so the ENV
+# below does not have to be updated when the base image's Triton moves.
+RUN set -eux; \
+    mkdir -p /opt/ftbuild/cudart; \
+    for d in /root/.triton/nvidia/cudart/*/lib; do \
+        cp -a "$d"/libcudart.so* /opt/ftbuild/cudart/ 2>/dev/null || true; \
+    done; \
+    test -f /opt/ftbuild/cudart/libcudart.so.12
+
+# Engine kernels and sparse MLA on Triton rather than the compiled/TileLang paths; FlagGems enabled
+# in allow-list mode. Single-node H20 needs both EP_DISABLE_GIN (no /dev/infiniband in the container)
+# and DEEPINFER_ENFORCE_EAGER (the Engram local gather does a D2H inside CUDA graph capture).
+# DEEPINFER_USE_NEXTN_DECODE=0 keeps speculative decoding off: its rejection-sampling kernel is the
+# one piece that still requires TileLang.
+#
+# DEEPINFER_OPS_STRICT=1 is set deliberately. Without it the router imports the compiled backend,
+# finds 16 of 29 kernels unported, and prints `falling back to 'cuda'` -- which reads as "depends on
+# the compiled path" even though those 16 are never called. Measured: with STRICT=1 the same
+# end-to-end run passes, `kernels served by the default backend` is empty, and the tokens are
+# identical. Strict is the honest setting here; it also turns any future missing-kernel call into a
+# loud failure instead of a silent substitution.
+ENV VIRTUAL_ENV=/opt/ftbuild/venv \
+    PATH=/opt/ftbuild/venv/bin:${PATH} \
+    PYTHONPATH=/opt/ftbuild/repo:/opt/ftbuild/repo/examples:/opt/ftbuild/repo/scripts \
+    LD_LIBRARY_PATH=/opt/ftbuild/cudart:${LD_LIBRARY_PATH} \
+    PYTHONUNBUFFERED=1 \
+    DEEPINFER_OPS_BACKEND=triton \
+    DEEPINFER_SPARSE_MLA_BACKEND=triton \
+    DEEPINFER_KERNELS=ref \
+    DEEPINFER_OPS_STRICT=1 \
+    DEEPINFER_ENABLE_FLAGGEMS=1 \
+    DEEPINFER_FLAGGEMS_MODE=allow \
+    DEEPINFER_USE_NEXTN_DECODE=0 \
+    EP_DISABLE_GIN=1 \
+    DEEPINFER_ENFORCE_EAGER=1
+
+WORKDIR /opt/ftbuild/repo
+
+# Fail the build rather than the first run if the relocation broke something. `flag_gems` and
+# `deepinfer` cannot be imported here: flag_gems probes for a device at import time and `docker build`
+# has no GPU ("No device were detected on your machine"), so those two are covered by the runbook's
+# TC-1/TC-5 instead. What is checkable offline is that the relocation is complete and torch/triton
+# resolve from the venv.
+RUN set -eux; \
+    python -c "import sys, torch, triton; print(sys.version.split()[0], torch.__version__, triton.__version__)"; \
+    if grep -rlq '/data/ftbuild' /opt/ftbuild/venv/bin /opt/ftbuild/venv/pyvenv.cfg; then \
+        echo "stale /data/ftbuild prefix left in the venv" >&2; exit 1; \
+    fi; \
+    if find /opt/ftbuild/venv -name '*.pth' -print0 | xargs -0 grep -l '/data/ftbuild' 2>/dev/null | grep -q .; then \
+        echo "stale /data/ftbuild prefix left in a .pth file" >&2; exit 1; \
+    fi; \
+    ls /opt/ftbuild/repo/deepinfer/_C*.so >/dev/null; \
+    head -1 /opt/ftbuild/venv/bin/pip | grep -q '/opt/ftbuild/venv/bin/python'; \
+    test -f /opt/ftbuild/repo/examples/_config.py; \
+    python -c "import _config; print('examples importable:', _config.__name__)"
+
+CMD ["bash"]
diff --git a/deploy/h20-flaggems/build.sh b/deploy/h20-flaggems/build.sh
new file mode 100755
index 0000000..58b5901
--- /dev/null
+++ b/deploy/h20-flaggems/build.sh
@@ -0,0 +1,48 @@
+#!/usr/bin/env bash
+# Build and publish the H20 / SM90 / FlagGems image.
+#
+# The build context needs the venv and the source tree, which live on a network volume on the build
+# host (`/data/ftbuild`), so they are tarred into a scratch directory under /data -- never into this
+# git tree, which would leave 2 GB of tars in a working copy.
+#
+#   bash deploy/h20-flaggems/build.sh              # build only
+#   bash deploy/h20-flaggems/build.sh --push       # build and push
+#
+# Env in: SRC_VENV (default /data/ftbuild/venv), SRC_REPO (default /data/ftbuild/repo),
+#         BUILD_DIR (default /data/ftbuild/image-build), IMAGE, TAG (default: the driver version,
+#         which is what the other images in flagrelease-public use as their tag).
+
+set -euo pipefail
+
+SRC_VENV="${SRC_VENV:-/data/ftbuild/venv}"
+SRC_REPO="${SRC_REPO:-/data/ftbuild/repo}"
+BUILD_DIR="${BUILD_DIR:-/data/ftbuild/image-build}"
+IMAGE="${IMAGE:-harbor.baai.ac.cn/flagrelease-public/deepseek-v4.1-flash-h20-gems5.0.2-tree3.6.0-cxnone-plugin0.0.0-vllm0.24.0-cp312-pt211-cu130-x64}"
+TAG="${TAG:-580.95.05}"
+
+HERE="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)"
+
+[ -d "$SRC_VENV" ] || { echo "missing $SRC_VENV" >&2; exit 2; }
+[ -d "$SRC_REPO" ] || { echo "missing $SRC_REPO" >&2; exit 2; }
+
+mkdir -p "$BUILD_DIR"
+cp "$HERE/Dockerfile" "$BUILD_DIR/Dockerfile"
+cd "$BUILD_DIR"
+
+echo "==> packing build context into $BUILD_DIR (reads ~3.6 GB from the network volume)"
+rm -f venv.tar.gz repo.tar.gz
+tar czf venv.tar.gz -C "$(dirname "$SRC_VENV")" --exclude='**/__pycache__' "$(basename "$SRC_VENV")"
+tar czf repo.tar.gz -C "$(dirname "$SRC_REPO")" --exclude='build' --exclude='**/__pycache__' "$(basename "$SRC_REPO")"
+ls -lh venv.tar.gz repo.tar.gz
+
+echo "==> building ${IMAGE}:${TAG}"
+docker build -t "${IMAGE}:${TAG}" .
+
+echo "==> built"
+docker images "${IMAGE}:${TAG}" --format '{{.Repository}}:{{.Tag}}  {{.Size}}  {{.ID}}'
+
+if [ "${1:-}" = "--push" ]; then
+    echo "==> pushing"
+    docker push "${IMAGE}:${TAG}"
+    docker inspect "${IMAGE}:${TAG}" --format 'digest: {{index .RepoDigests 0}}'
+fi