diff --git a/README.md b/README.md
index 3391c47..9453782 100644
--- a/README.md
+++ b/README.md
@@ -94,7 +94,7 @@ softmax:
 | `ZL_DIRECT_REPEAT` | `1` | direct 计时的重复次数，取中位。只作用于 direct 路径：整体 `timer=direct` 的运行，以及图捕获运行里逐行回退的那些点（标记会带上 `repeat=N`）；捕获成功的行不受影响。`do_bench` 自己已在调用内取中位，压住的是迭代间抖动；跨调用重复补的是它看不见的那层——时钟爬坡、邻居进程进来这类漂移。调用方要 `quantiles` 或 `return_mode="all"` 时返回序列、跨调用中位无定义，自动不重复。在 MThreads 上这不是可选项：cvt 消除量最大的 `sort`/`dropout`/`rand` 恰好都无法捕获，那批结论正是由这条路径产出的。取值写进 `run_identity.json` 的 `direct_repeat` |
 | `DTYPES` | 空（上游默认 dtype 扫描） | 空格分隔的 dtype 白名单（如 `"bfloat16 float16"`），逐个转为上游 `--dtypes`。算子不支持时上游报 `can't be supported by this op` |
 | `PARALLEL_WARMUP_GPUS` | 空（串行） | 设为 `N`（≥2）时把 autotune sweep 分片到 N 卡并行，测量仍单卡串行；不设或 `<2` 则完全串行（见「加速：多卡并行 sweep」） |
-| `ZL_ENABLE_CUDAGRAPH` | `auto` | `auto`/`1` 仅在 CUDA 上走 Graph 计时；非 CUDA 即使设 `1` 也拒绝并保持 `timer=direct` |
+| `ZL_ENABLE_CUDAGRAPH` | `auto` | `auto`/`1` 在**声明了 `supports_cuda_graph` 的后端**上走 Graph 计时（目前 cuda 与 musa），其余后端即使设 `1` 也拒绝并保持 `timer=direct`。注意 `timer=cuda_graph` 只表示**尝试过**捕获，不表示有任何一行捕获成功——S5000 上 `dropout`/`rand`/`masked_select` 是逐 shape 全部回退（每行 `BENCHMARK_DIRECT_NO_CUDAGRAPH`，零毒化、不丢行），整张表实际是 direct 计时。这种情况下设 `0` 走的是同一条路径，只是记录更诚实 |
 | `ZL_ENABLE_AUTOTUNE` | `1` | `0` 关闭 autotune record/replay 插件 |
 | `ZL_ENABLE_IR` | `1` | `0` 关闭 TTGIR / 编译元数据插件 |
 | `ZL_ENABLE_MM_CLUSTER` | `0` | `1` 且后端支持时加载 Hopper mm cluster 规避 |
@@ -380,8 +380,8 @@ python run_batch.py --ops-file my_ops.txt --replay-root runs/batch_xxx
 | `_shape_iter_inject_plugin` | 覆盖在 `get_input_iter` 里硬编码 shape 的类（conv/pool 等） | 不关 |
 | `_bespoke_shape_plugin` | 覆盖特殊输入构造的算子（upsample/flash_mla/cutlass 等，按类名注册） | 测这些算子之外可关 |
 | `_autotune_record_plugin` | record/replay Triton autotune 选择；`autotune_records/meta.json` 记录 backend，跨后端 replay 拒绝 | `ZL_ENABLE_AUTOTUNE=0` |
-| `_timing_plugin` | 按后端选择计时路径：CUDA 可走 Graph（内部调用 `_cudagraph_plugin`），其余后端一开始就是上游 `do_bench`（`timer=direct`） | `ZL_ENABLE_CUDAGRAPH=0` 强制 direct |
-| `_cudagraph_plugin` | 仅由 `_timing_plugin` 在 CUDA 上调用。`do_bench` → `do_bench_cudagraph`，capture 前按 warmup 预算显式预热；无法 capture 时回退并打 `BENCHMARK_DIRECT_NO_CUDAGRAPH` | 不要直接 `-p` 加到非 CUDA 运行 |
+| `_timing_plugin` | 按后端的 `supports_cuda_graph` 选计时路径：声明了的走 Graph（内部调用 `_cudagraph_plugin`，逐 shape 回退），没声明的一开始就是上游 `do_bench`（`timer=direct`） | `ZL_ENABLE_CUDAGRAPH=0` 强制 direct |
+| `_cudagraph_plugin` | 仅由 `_timing_plugin` 在声明了 `supports_cuda_graph` 的后端上调用。`do_bench` → `do_bench_cudagraph`，capture 前按 warmup 预算显式预热；无法 capture 时回退并打 `BENCHMARK_DIRECT_NO_CUDAGRAPH` | 不要直接 `-p` 加到非 CUDA 运行 |
 | `_musa_benchmark_compat_plugin` | 剥掉 FlagGems 上游若干 benchmark 上过时的 `torch.cuda.is_available()` 跳过门——被测代码本就经 `flag_gems.device` 分发，在 MThreads 上那些门把可用算子变成「未测量」。每条覆盖都钉死到一个文件与一条确切的 skip 理由，且仅在 `flag_gems.vendor_name == "mthreads"` 时生效 | 由 `backends/musa.py` 的 `compat_plugins` 声明，其他后端不加载 |
 | `_pretty_report_plugin` | 结果表整理 + 着色：全表相同的输入折叠成表头下一行图例，每行 Size Detail 只留随行变化的部分（shape 保持 `torch.Size([...])` 原样，MoE 类算子单行从 ~400 字符缩到一屏内）；SUCCESS 绿 / FAILED 红。列名与 `SUCCESS` 字样保持上游原文，`run.log` 的 grep/解析不受影响 | 需要与上游原始表格逐字对齐时去掉 |
 | `_ir_meta_plugin` | 编译期记录每个变体的 constexpr/签名/特化；launch 钩子统计每个 (kernel, shape) 的实际使用；退出时按上述结构 dump ttgir。Triton knobs 缺失时 fail-soft | `ZL_ENABLE_IR=0` |
diff --git a/_timing_plugin.py b/_timing_plugin.py
index 7013bd6..7427ae2 100644
--- a/_timing_plugin.py
+++ b/_timing_plugin.py
@@ -1,10 +1,15 @@
 """Select the timing path from backend capabilities, before any shape runs.
 
-CUDA keeps the existing `_cudagraph_plugin` behaviour (graph capture, with
-per-shape fallback to plain do_bench). Other backends do not attempt CUDA
-Graph at all: they leave `triton.testing.do_bench` in place and record
-`timer=direct`. Trying capture-then-fallback on a non-CUDA device would
-produce FATAL_DEVICE noise and mixed-caliber rows.
+Whether capture is attempted comes from the backend's `supports_cuda_graph`,
+not from "is this CUDA". CUDA and MUSA both declare it, so both get graph
+capture with per-shape fallback to plain do_bench; a backend that declares
+nothing leaves `triton.testing.do_bench` in place and records `timer=direct`.
+
+`timer=cuda_graph` therefore means capture was *attempted*, not that any row
+was captured. On S5000 every shape of dropout/rand/masked_select falls back
+(BENCHMARK_DIRECT_NO_CUDAGRAPH on every row, nothing poisoned, no rows lost),
+so the whole table is direct-timed under a cuda_graph label. Setting
+ZL_ENABLE_CUDAGRAPH=0 there selects the same timing path and says so plainly.
 """
 
 from __future__ import annotations