forked from zhoulin/zl_bench
AharrypotterandClaude Opus 5
c8ea586d0f
Everything about which accelerator a run targets was previously spelled CUDA: device discovery went through nvidia-smi, pinning through CUDA_VISIBLE_DEVICES, the vendor guard only recognised NVIDIA, and the -p plugin list was hardcoded in run_pytest.sh with no way to drop the CUDA-only members. Running the same operator corpus on another vendor's device therefore meant editing the drivers. The selection now hangs off one ZL_BACKEND variable (cuda/musa/rocm/ascend, with the usual aliases): - backends/ owns per-backend device enumeration, visibility env keys and capability flags, so nothing above it names a vendor; - plugin_plan.py derives the -p list from those capabilities instead of a hardcoded array, which is what lets a non-CUDA run omit the graph-capture and cluster-fix plugins rather than watch them fail; - _timing_plugin selects the timer from the same capabilities: CUDA keeps graph capture with its per-shape fallback, other backends stay on upstream do_bench and record timer=direct, because capture-then-fallback on a device without graph support only produces noise; - _device_setup_plugin replaces the NVIDIA-only vendor guard; - providers/flaggems.py resolves the pytest node, so a caller may pass TEST_NODE directly instead of the OP/OP_FILE pair, and SHAPE_MODE=upstream can defer to FlagGems' own core_shapes; - _run_identity.py writes run_identity.json next to each run, recording backend, device, timer, autotune mode and the resolved test node. Two runs whose calibers differ are otherwise indistinguishable after the fact. run_batch gains --backend and --devices (--gpus kept as an alias) and carries backend/device/timer into ops_status.csv. test_backends.py covers the registry, alias resolution, visibility env and plugin planning. Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
48 lines
1.9 KiB
Python
48 lines
1.9 KiB
Python
"""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.
|
|
"""
|
|
|
|
from __future__ import annotations
|
|
|
|
import os
|
|
import sys
|
|
|
|
from backends import get_backend
|
|
from _run_identity import update_identity
|
|
|
|
|
|
def _graph_requested(backend) -> bool:
|
|
raw = os.environ.get("ZL_ENABLE_CUDAGRAPH", "auto").strip().lower()
|
|
if raw in ("0", "false", "no", "off", "never", "direct"):
|
|
return False
|
|
if raw in ("1", "true", "yes", "on", "always"):
|
|
return bool(backend.supports_cuda_graph)
|
|
return bool(backend.supports_cuda_graph)
|
|
|
|
|
|
def pytest_configure(config):
|
|
from _term_style import tag
|
|
backend = get_backend()
|
|
raw = os.environ.get("ZL_ENABLE_CUDAGRAPH", "auto").strip().lower()
|
|
force_on = raw in ("1", "true", "yes", "on", "always")
|
|
if force_on and not backend.supports_cuda_graph:
|
|
print(f"{tag('[timing-plugin]')} ZL_ENABLE_CUDAGRAPH=1 ignored on "
|
|
f"backend={backend.name} (no CUDA Graph); using timer=direct",
|
|
file=sys.stderr, flush=True)
|
|
if _graph_requested(backend):
|
|
import _cudagraph_plugin as _cg
|
|
_cg.pytest_configure(config)
|
|
update_identity(timer="cuda_graph", graph_capture="attempted")
|
|
print(f"{tag('[timing-plugin]')} timer=cuda_graph backend={backend.name}",
|
|
file=sys.stderr, flush=True)
|
|
return
|
|
update_identity(timer="direct", graph_capture="disabled")
|
|
print(f"{tag('[timing-plugin]')} timer=direct backend={backend.name} "
|
|
"(upstream do_bench; graph capture not selected)",
|
|
file=sys.stderr, flush=True)
|