docs: stop claiming non-CUDA backends never attempt graph capture

Both the _timing_plugin docstring and the README env table still described the
behaviour from before MUSA declared supports_cuda_graph: "other backends do
not attempt CUDA Graph at all", "非 CUDA 即使设 1 也拒绝". Capture selection has
been driven by the backend capability, not by "is this CUDA", since the MUSA
fallback policy landed. Reading the old text is most likely why the S5000
campaigns were launched with ZL_ENABLE_CUDAGRAPH=0 without anyone checking
what auto would do.

Also states what timer=cuda_graph does and does not mean, since the two are
easy to conflate: it says capture was attempted. Measured on S5000 with auto
and configs pinned, dropout/rand/masked_select select the graph path and then
fall back on every single shape — BENCHMARK_DIRECT_NO_CUDAGRAPH on all 15
rows, zero poisoned, zero rows lost — so the whole table is direct-timed under
a cuda_graph label. ZL_ENABLE_CUDAGRAPH=0 selects the same path and records it
plainly, which is why the existing campaign data stands as measured.
This commit is contained in:
Aharrypotter committed 2026-09-10 16:00:44 +08:00
1 parent ee5f70fa89
commit 4329b0c835
2 files changed
+13 -8

No files matched your search

+10 -5
View File
@@ -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