diff --git a/README.md b/README.md
index 0dd1ad2..d152b38 100644
--- a/README.md
+++ b/README.md
@@ -15,8 +15,13 @@ bash run_pytest.sh
 # 换算子：OP=测试函数名去掉 test_ 前缀，OP_FILE=benchmark 文件名去掉 test_ 前缀/.py 后缀
 # （benchmark 文件为 $FLAGGEMS_DIR/benchmark/test_<OP_FILE>.py::test_<OP>）
 OP=softmax OP_FILE=softmax SHAPE_FILE=my_shapes.yaml bash run_pytest.sh
+
+# 测 FlagGems-vllm 仓库的算子：换 FLAGGEMS_DIR 即可（插件自动适配其包名 flaggems_vllm）
+FLAGGEMS_DIR=/workspace/dev/FlagGems-vllm OP=fused_marlin_moe OP_FILE=fused_marlin_moe bash run_pytest.sh
 ```
 
+上游现为两个仓库，benchmark 体系同构、均受支持：主仓 `/workspace/dev/FlagGems`（Python 包名 `flag_gems`，默认；`fused_marlin_moe_mxfp4` 仅此仓有）与 `/workspace/dev/FlagGems-vllm`（包名 `flaggems_vllm`）。两包需各自 `pip install -e <仓库> --no-deps` 安装（`--no-deps` 避免动 FlagTree 的 triton）。
+
 shape yaml 顶层 key 必须是 op 名：
 
 ```yaml
@@ -41,7 +46,7 @@ softmax:
 | `OP` | `fused_marlin_moe_mxfp4` | 测试函数名（`test_` 之后的部分） |
 | `OP_FILE` | `fused_marlin_moe` | benchmark 文件名（`test_` 与 `.py` 之间的部分） |
 | `SHAPE_FILE` | 空（用脚本内置 yaml） | shape yaml 路径 |
-| `FLAGGEMS_DIR` | `/workspace/FlagGems-dev` | FlagGems 仓库路径 |
+| `FLAGGEMS_DIR` | `/workspace/dev/FlagGems` | FlagGems 仓库路径（测 vllm 仓库时指向 `/workspace/dev/FlagGems-vllm`） |
 
 **怎么测**
 
@@ -100,6 +105,8 @@ A/B 两侧的 `USE_FLAGTUNE` 必须取同值：该开关决定调优空间，两
 
 对比 `run.log` 的 latency 表看性能差异；diff 两侧 `ttgir/<shape>/<kernel>/` 下的同名文件看 IR 差异。
 
+带 `pre_hook` 的 config（如 hopper mm 的 TMA configs）：libtuner kernel 正常 record/replay——注入的 config 会由上游按 kwargs 匹配自动接回 pre_hook（与其自身 ConfigCache 的 DB 回读同一条路径）；普通 `@triton.autotune` kernel 无此恢复机制，这类 config 不落 record，replay 时表现为下述 `key_missing` fallback。
+
 replay 的兜底行为：B 侧遇到记录中没有的 key、或记录的 config 在新编译器下编译失败时，自动回退到现场 autotune 并在 `run.log` 打 `AUTOTUNE_REPLAY_FALLBACK reason=...` 标记——出现该标记的测量点不再满足"同 config"前提，解读时注意。标记以 `_no_evict` 结尾时更弱一层：坏 config 存在 libtuner 的 sqlite 缓存里、没有 `__delitem__` 可摘除，重试会再读到同一个 config——这类测量点按未验证处理。另外 replay 模式的 run 目录不产生 `autotune_records/`，后续 run 的 `REPLAY_FROM` 应始终指向最初 record 的那次 A 侧目录，不要链式指向 replay 产物。
 
 ### 为什么需要 replay（以及它管不到什么）
@@ -111,6 +118,8 @@ replay 的兜底行为：B 侧遇到记录中没有的 key、或记录的 config
 | Triton `@triton.autotune` | 仅进程内存 | 每次进程重新 sweep | 计时噪声可能让 A/B 选中**不同 config** → 用 REPLAY_FROM 固定（影响比预期大，见下） |
 | FlagGems `@libtuner`（含 FlagTune 扩展空间） | `~/.flaggems/config_cache/*.db`（sqlite，跨进程持久） | **不失效**（表名只含 kernel 源码与 config 空间的 hash），A/B 自动命中同一 winner | 反向风险：B 侧沿用 A 侧选的旧 winner，测的是"旧 config 下的编译器差异"而非"各自最优" |
 
+新版 libtuner 在同一 db 里还持久化了 **BenchmarkCache**（sweep 中每个 config 的实测 latency）：db 已热时 record 模式的 sweep 也不重测，直接按历史延迟选 winner——"record 每次现场 sweep"仅在冷 db 下严格成立。下文"删 db 换各自最优口径"的操作会同时清掉 winner 与延迟两层，仍然有效。
+
 **这不是理论风险，量级足以吞掉被测优化本身。** 同一份 `shapes.yaml`、同一口径、相隔十几分钟的两次 record，平均 speedup 可以差出 10% 量级，且逐 shape 单向偏移（不是随机噪声）。diff 两侧 `autotune_records/*.json` 能看到差异往往不是微调而是换挡——`num_stages`、`BLOCK_SIZE_*`、`num_warps` 整档跳变。怀疑遇到这种情况时，先 diff 两侧的 config 再看 latency。
 
 所以 record 模式的数字只用来给 replay 提供 config 基准。**若某次结论只有 record 数据支撑，按未验证处理、重跑补 replay。**
@@ -209,6 +218,6 @@ shape 多时这部分会主导墙上时间，可设 `PARALLEL_WARMUP_GPUS=N` 并
 
 ## 依赖假设
 
-- FlagGems benchmark 体系（`benchmark/base.py` 的 `Benchmark` 类、conftest 的 `--shape_file/--level/--mode` 选项）；
+- FlagGems benchmark 体系（`benchmark/base.py` 的 `Benchmark` 类、conftest 的 `--shape_file/--level/--mode` 选项）；主仓与 FlagGems-vllm 仓库同构，`_autotune_record_plugin` 会按实际被 import 的包（`flag_gems` / `flaggems_vllm`）挂 LibTuner 补丁；
 - Triton 需支持 `knobs.compilation.listener`、`kernel_load_end_hook`、`launch_enter_hook`（当前 FlagTree 的 triton 3.6 满足）；
 - 插件通过 monkeypatch 挂钩上游内部结构，FlagGems/Triton 大版本升级后若行为异常，优先检查各插件 pytest_configure 输出的注册日志是否还正常打印。
diff --git a/_autotune_record_plugin.py b/_autotune_record_plugin.py
index 4b016a9..829c77a 100644
--- a/_autotune_record_plugin.py
+++ b/_autotune_record_plugin.py
@@ -18,6 +18,7 @@ FLAGGEMS_PERF_CURRENT_OP):
 from __future__ import annotations
 
 import atexit
+import importlib
 import json
 import os
 import sys
@@ -30,12 +31,17 @@ _RECORD_DIR_ENV = "FLAGGEMS_PERF_AUTOTUNE_RECORD_DIR"
 _REPLAY_DIR_ENV = "FLAGGEMS_PERF_AUTOTUNE_REPLAY_DIR"
 _OP_ENV = "FLAGGEMS_PERF_CURRENT_OP"
 
+# Upstream split into two repos with different package names; either may host
+# LibTuner-decorated kernels depending on which repo's benchmark is running.
+_LIBTUNER_PKGS = ("flag_gems", "flaggems_vllm")
+
 # Module-global aggregates, accumulated as run() is called and dumped at exit.
 # Lock guards merges in case a kernel uses threads internally.
 _record_map: Dict[str, Dict[str, Dict[str, Any]]] = {}
 _record_lock = threading.Lock()
 _replay_map: Dict[str, Dict[str, Dict[str, Any]]] = {}
 _dump_done = False
+_wrap = None  # set in pytest_configure when record/replay is active
 
 
 def _op_name() -> str:
@@ -65,10 +71,16 @@ def _serialize_key(key: Tuple[Any, ...]) -> str:
     return json.dumps([repr(x) for x in key])
 
 
-def _serialize_config(cfg: Any) -> Optional[Dict[str, Any]]:
-    # Drop configs with a pre_hook (un-serializable callable) rather than replay
-    # without it and compute incorrectly.
-    if getattr(cfg, "pre_hook", None) is not None:
+def _serialize_config(cfg: Any, allow_pre_hook: bool = False) -> Optional[Dict[str, Any]]:
+    # pre_hook is an un-serializable callable. For plain triton Autotuner drop
+    # such configs rather than replay without the hook and compute incorrectly.
+    # For LibTuner callers pass allow_pre_hook=True: upstream's own ConfigCache
+    # round-trips configs pre_hook-less and LibTuner.run re-attaches the hook by
+    # matching all_kwargs against self.configs, and replay-injected configs go
+    # through that same path (hopper mm's TMA configs carry a pre_hook even
+    # with USE_FLAGTUNE=0 — dropping them would record nothing for those
+    # kernels and silently defeat replay).
+    if getattr(cfg, "pre_hook", None) is not None and not allow_pre_hook:
         return None
     try:
         return {
@@ -141,7 +153,9 @@ def _record_run(original):
             cfg = self.configs[0]
         if cfg is None:
             return result
-        entry = _serialize_config(cfg)
+        # get_key marks LibTuner (both flag_gems and flaggems_vllm); plain
+        # triton Autotuner has no pre_hook re-attach on cache read, LibTuner does.
+        entry = _serialize_config(cfg, allow_pre_hook=hasattr(self, "get_key"))
         if entry is not None:
             kid = _kernel_id(self)
             with _record_lock:
@@ -240,6 +254,7 @@ def _dump_record(record_dir: str) -> None:
 
 
 def pytest_configure(config):
+    global _wrap
     from _term_style import tag
     record_dir = os.environ.get(_RECORD_DIR_ENV, "").strip()
     replay_dir = os.environ.get(_REPLAY_DIR_ENV, "").strip()
@@ -252,26 +267,44 @@ def pytest_configure(config):
         return
 
     import triton.runtime.autotuner as _autotuner
-    wrap = _record_run if record_dir else _replay_run
-    _autotuner.Autotuner.run = wrap(_autotuner.Autotuner.run)
-    patched = ["Autotuner"]
-    try:
-        from flag_gems.utils.libentry import LibTuner
-        if "run" in LibTuner.__dict__:
-            LibTuner.run = wrap(LibTuner.__dict__["run"])
-            patched.append("LibTuner")
-    except Exception:
-        pass
+    _wrap = _record_run if record_dir else _replay_run
+    _autotuner.Autotuner.run = _wrap(_autotuner.Autotuner.run)
 
     if record_dir:
         # atexit (not sessionfinish): persist whatever was recorded even if an op
         # crash kills the session; a later replay run falls back for missing keys.
         atexit.register(_dump_record, record_dir)
         print(f"{tag('[autotune-record-plugin]')} recording autotune configs to "
-              f"{_record_path(record_dir)} ({'+'.join(patched)})",
+              f"{_record_path(record_dir)} (Autotuner)",
               file=sys.stderr, flush=True)
     else:
         loaded = _load_replay_dir(replay_dir)
         print(f"{tag('[autotune-record-plugin]')} replaying {loaded} recorded entries from "
-              f"{_record_path(replay_dir)} ({'+'.join(patched)})",
+              f"{_record_path(replay_dir)} (Autotuner)",
               file=sys.stderr, flush=True)
+
+
+def pytest_collection_finish(session):
+    # LibTuner is patched here rather than in pytest_configure: only after
+    # collection (which imports the benchmark module) do we know which FlagGems
+    # package is actually in use, and importing the unused one just to patch it
+    # would initialize a second runtime in this process for nothing.
+    if _wrap is None:
+        return
+    from _term_style import tag
+    patched = []
+    for pkg in _LIBTUNER_PKGS:
+        if pkg not in sys.modules:
+            continue
+        try:
+            libentry = importlib.import_module(f"{pkg}.utils.libentry")
+            tuner_cls = libentry.LibTuner
+            if "run" in tuner_cls.__dict__:
+                tuner_cls.run = _wrap(tuner_cls.__dict__["run"])
+                patched.append(f"LibTuner[{pkg}]")
+        except Exception as exc:
+            print(f"{tag('[autotune-record-plugin]')} warning: LibTuner patch "
+                  f"failed for {pkg}: {exc}", file=sys.stderr, flush=True)
+    print(f"{tag('[autotune-record-plugin]')} libtuner coverage: "
+          f"{', '.join(patched) or 'none (no FlagGems package imported?)'}",
+          file=sys.stderr, flush=True)
diff --git a/run_pytest.sh b/run_pytest.sh
index 9d537e8..c5a04e6 100644
--- a/run_pytest.sh
+++ b/run_pytest.sh
@@ -11,7 +11,9 @@ set -euo pipefail
 OP="${OP:-fused_marlin_moe_mxfp4}"
 OP_FILE="${OP_FILE:-fused_marlin_moe}"
 
-FLAGGEMS_DIR="${FLAGGEMS_DIR:-/workspace/FlagGems-dev}"
+# 上游已拆分为两个仓库：主仓 /workspace/dev/FlagGems（包名 flag_gems）与
+# /workspace/dev/FlagGems-vllm（包名 flaggems_vllm）。测 vllm 仓库算子时覆盖本变量即可。
+FLAGGEMS_DIR="${FLAGGEMS_DIR:-/workspace/dev/FlagGems}"
 
 # shape 来源：非空则用该 yaml，否则用 INLINE_YAML。顶层 key 必须是 op 名。
 SHAPE_FILE="${SHAPE_FILE:-}"