diff --git a/benchmarks/microbenchmarks/BENCHMARKING_NOTES.md b/benchmarks/microbenchmarks/BENCHMARKING_NOTES.md
new file mode 100644
index 0000000000..812c817a1f
--- /dev/null
+++ b/benchmarks/microbenchmarks/BENCHMARKING_NOTES.md
@@ -0,0 +1,210 @@
+# Microbenchmark Notes
+
+How to read the microbenchmark results (including the TE dashboard, which is fed by
+`--dashboard-run`), and why the harness measures the way it does. For how to run the
+benchmarks and their options, see [README.md](README.md).
+
+- [Reading the results](#reading-the-results) is for anyone looking at the numbers.
+- [Measurement setup](#measurement-setup) and [Known limitations](#known-limitations)
+  are for people changing the harness or chasing a surprising result.
+
+## Reading the results
+
+### What a row measures
+
+Each row is one case: model shape x precision x backend (x direction where it applies), timed
+in one process on one GPU. Times are per call and are medians over many samples.
+
+| Family | What one call is | Throughput |
+|---|---|---|
+| `gemm` | `te.Linear` forward, or its backward, under `te.autocast` for the quantized recipes | TFLOPS from GEMM FLOPs: `2*M*N*K` forward, twice that backward (dgrad + wgrad) |
+| `grouped_gemm` | `te.GroupedLinear` forward or backward with `B` experts of `M` tokens each (balanced) | TFLOPS from `2*B*M*N*K` forward, twice that backward |
+| `casting` | quantize one BF16 tensor, or dequantize one quantized tensor back to BF16 | GB/s from BF16 bytes + quantized bytes (data + block scales) |
+| `normalization` | LayerNorm / RMSNorm forward, writing its output in the target precision | GB/s from BF16 input read + output write |
+
+- GEMM rows time the whole TE layer call, not just the GEMM kernel: casts TE does inside
+  the layer (e.g. quantizing the input under autocast) and its Python overhead are included,
+  but the throughput counts GEMM FLOPs only.
+- The `fp8` recipe is per-tensor delayed scaling (HYBRID format); `mxfp8`, `mxfp4` and
+  `nvfp4` are TE's block-scaling recipes.
+- `B` in grouped GEMM is the number of experts per GPU for a given expert-parallel size
+  (e.g. 256 routed experts at EP 32 gives `B = 8`).
+
+### Wall time and kernel time
+
+- **Wall time** is measured on the host around back-to-back calls, so it is throughput
+  (calls pipeline on the GPU, as in training), not single-call latency. It includes Python
+  dispatch and kernel launch overhead.
+- **Kernel time** (`--kernel-profile`) is GPU device time per call from `torch.profiler`:
+  the summed device time of the call's kernels, without host overhead.
+- **Kernel time well below wall time** means the case is launch- or host-bound (common at
+  small M and for higher-overhead backends). For GPU-bound cases the two nearly match.
+- **Kernel time slightly above wall time** (a few percent) happens for some GPU-bound cases:
+  the profiler itself changes how the GPU runs back-to-back kernels, in either direction.
+  For GPU-bound cases, wall time is the more faithful number.
+- **Grouped GEMM kernel time is not device-only.** It is the event span between consecutive
+  calls on the main stream, because the per-expert path overlaps kernels on several streams
+  (a per-kernel sum would double-count). The span includes GPU idle time while the host
+  submits the next call, so launch-bound grouped cases read at or slightly above wall time.
+  Compare it with dense kernel time only for GPU-bound cases.
+- **An empty kernel cell** means the profiled kernels could not be matched to calls for that
+  case; the run log names it.
+- **The TE dashboard plots wall time** for every case. Kernel time is only in the raw CSVs.
+
+### Forward and backward
+
+The backward row times only the backward: a block of forwards runs untimed, then their
+backwards run back to back and are timed, so the host can run ahead as in the forward
+timing. No forward time is subtracted.
+
+### Inputs
+
+Activations are rotated through a ring of buffers larger than twice the last-level cache, so
+consecutive calls read memory that is not cache-resident (close to a steady-state training
+workload). Weights are not rotated.
+
+### Backends and skipped cases
+
+- A backend is selected with TE's environment variables; if the installed TE build doesn't
+  honor them, the case is skipped rather than silently timing the default path.
+- Each FlyDSL case first runs once with FlyDSL's fallback warning enabled and is skipped if
+  FlyDSL hands the shape back to the default GEMM, so every `flydsl` row is a FlyDSL
+  measurement. CK and HipKittens grouped GEMM are probed the same way for fallbacks to the
+  per-expert path.
+- Grouped GEMM backends:
+  - `hipblaslt`: per-expert hipBLASLt GEMMs on TE's compute streams (TE's default for BF16,
+    FP8 and NVFP4).
+  - `hipkittens_per_expert`: dense HipKittens GEMM per expert on the compute streams (TE's
+    default for MXFP8 on gfx950).
+  - `hipkittens_grouped`: one HipKittens grouped kernel on the main stream (opt-in in TE).
+  - `ck_tile`, `triton`: CK and Triton grouped kernels.
+- Dense MXFP4 GEMM runs only the `aiter` backend (AITER a4w4, TE's default); TE's opt-in
+  hipBLASLt MXFP4 path fails in the `te.Linear` backward.
+- HipKittens needs gfx950 and N and K divisible by 256; MXFP8 and MXFP4 GEMMs need
+  32-aligned dimensions (hipBLASLt MXFP8 on gfx950 needs K divisible by 128). Unaligned
+  shapes are skipped.
+- Known backend bugs on specific shapes are marked xfail (failing configs) or skipped
+  (configs that hang); the reasons are in the benchmark files.
+
+### Noise and outliers
+
+- GPU-bound cases typically repeat within a few percent between runs. Launch-bound cases
+  depend on host speed and are more sensitive, which is why the harness pins CPUs (below).
+- Occasionally a stretch of cases runs slower for reasons outside the harness. Two patterns:
+  - host-side: wall time up (often ~2x), kernel time unchanged;
+  - GPU-side: wall and kernel time both up by the same factor, uniformly for the whole case.
+- To judge a change, first check whether kernel time moved with wall time (GPU-side) or not
+  (host-side).
+
+## Measurement setup
+
+### Wall timing
+
+- Forward and single-op cases use `torch.utils.benchmark.Timer` (`blocked_autorange` for
+  casting, `adaptive_autorange` otherwise) with a 0.2 s minimum run time, and report the
+  median.
+- Each case first makes one untimed call. The Timer chooses its block size (calls per
+  sample) from its first timed calls; a slow one-off first call (lazy kernel load, JIT
+  compilation, algorithm selection, autotuning) can make it pick a block of one call, so
+  every sample includes a device sync and measures latency instead of throughput. Without
+  the one-off, blocks last at least several milliseconds and the per-block sync cost is
+  negligible.
+
+### Backward timing
+
+- Per block, N forwards (N up to 20, fewer if their saved state would exceed 16 GB) run
+  untimed, then the N backwards run back to back and are timed. The first block is warmup;
+  timing continues for at least 15 blocks and 0.2 s. GC is off while timing (N live
+  autograd graphs make collections expensive).
+- Rejected alternatives:
+  - `(forward + backward) - forward`: its median is stable, but individual runs produced
+    occasional large outliers.
+  - Synchronizing around each backward: repeatable, but measures latency plus a fixed sync
+    cost rather than pipelined throughput, and reads high for GPU-bound shapes.
+  - Repeating the backward on one graph (`retain_graph`): TE frees saved state in backward,
+    so this fails for `Linear` and `GroupedLinear`.
+- The short blocks leave a small fixed per-block cost in the result (around 1% for the
+  fastest backward cases). It is constant, so trends are unaffected.
+
+### Kernel time
+
+- `--kernel-profile` runs a separate pass after the wall timing: 100 warmup calls, then one
+  profiler session of 100 calls (for backward, sessions of one backward block each). The
+  kernel list is split per call by matching the kernel sequence of a single profiled call;
+  only fully matched calls count, and the median is reported.
+- One long session rather than many short ones: short sessions let the GPU clock down
+  between them and bias large GEMMs high. A brief busy loop before each session brings the
+  clocks back up.
+- The profiler on ROCm can drop kernel records that are still in flight when a session
+  stops; a short sleep before stopping avoids that. Turning off the profiler's queue
+  interposition also avoids it but adds launch overhead to every call, so it isn't used.
+- Grouped GEMM uses the event-span method described above. `NVTE_MICROBENCH_KERNEL_TIME`
+  (`profiler` or `event`) forces one method for all families, for comparisons.
+
+### Input rotation
+
+The ring is sized to the `--rotating` budget, or by default to twice a 256 MB last-level
+cache (the AMD Infinity Cache size on gfx942/gfx950), and capped at 1000 buffers.
+
+### CPU pinning
+
+- Launch-bound cases depend on CPU clocks. With the default `schedutil` governor, the wall
+  time of launch-bound backward cases (which run on PyTorch's autograd thread) varied up to
+  2x between runs.
+- The harness therefore pins the main thread to 4 physical cores on the GPU's NUMA node
+  (offset by GPU id, so runs on other GPUs use other cores) and sets them to the
+  `performance` governor. Threads the main thread creates later, such as the autograd
+  thread, inherit the cores.
+- HIP/ROCr helper threads that exist after HIP init go to those cores' SMT siblings. One of
+  them polls events continuously, and the multi-stream grouped GEMM path ran about 1.6x
+  slower when it landed on another CCD or socket than the main thread (and slower still
+  when it shared the main thread's cores).
+- The previous governors are restored at the end of the session, or by a small watchdog
+  process if the benchmark is killed or crashes. `--no-cpu-pin` disables all of this.
+
+### HW queues
+
+- `GPU_MAX_HW_QUEUES` defaults to 2 (as in Megatron's launch scripts); an explicit value
+  overrides it.
+- Kernels on the main stream slow down while a cross-queue event wait is pending on a HW
+  queue that shares the main queue's hardware pipe. With HIP's default of 4 queues, TE's
+  high-priority compute streams get an extra queue that lands on the main queue's pipe, and
+  the per-expert grouped GEMM path ran up to about 2x slower and bimodally. With 2 queues it
+  is stable.
+
+### Isolation
+
+- TE's FP8 global state is reset after each test: modules from earlier cases otherwise stay
+  registered and slow down every later autocast exit.
+- At startup the harness warns (or, with `--abort-on-gpu-interference`, aborts) if another
+  process is using the benchmark GPU. The direct-run shim enters pytest before torch
+  touches the GPU, so the benchmark's own process is not reported.
+
+### Dashboard runs
+
+`--dashboard-run` writes all families of a run into
+`results/<arch>_<host>_gpu<id>_<pci>/<date>_<commit>/`:
+
+- `benchmark_<family>.csv`, each row tagged with the run timestamp, commit SHA and commit
+  date;
+- `run_info.txt` with the host, GPU, software versions, CPU pinning and HW-queue setting;
+- `samples/` with per-sample times when `--csv-samples` is given.
+
+## Known limitations
+
+- **Profiler perturbation:** kernel time of GPU-bound cases can differ from the unprofiled
+  run by a few percent in either direction.
+- **Grouped GEMM kernel time** includes GPU idle time for launch-bound cases (see above). A
+  device-busy measurement (the union of kernel intervals per call) would fix this; not
+  implemented.
+- **No per-case GPU clock or throttle record:** GPU-side transients are found only by
+  comparing with other runs.
+- **Configuration-specific results:** the numbers hold for pinned CPUs and
+  `GPU_MAX_HW_QUEUES=2`. Without them, launch-bound and per-expert grouped GEMM cases can
+  be up to about 2x slower.
+- **Single measurement per run:** each case is timed once per run, in one process, so
+  process-level effects (e.g. memory placement) are not averaged out.
+- **Upstream issues:** some results are shaped by issues outside TE that are worked around
+  or marked in the harness: the ROCm runtime's cross-queue wait slowdown (HW-queue setting),
+  profiler record loss at session stop (sleep before stop), and known backend failures or
+  hangs on specific shapes (xfail/skip).
diff --git a/benchmarks/microbenchmarks/README.md b/benchmarks/microbenchmarks/README.md
index 99aba90e77..94c843e802 100644
--- a/benchmarks/microbenchmarks/README.md
+++ b/benchmarks/microbenchmarks/README.md
@@ -2,6 +2,8 @@
 
 This directory contains lightweight Python microbenchmarks for selected
 Transformer Engine kernels and helper scripts for comparing benchmark CSVs.
+For how to read the results and why the harness measures the way it does, see
+[BENCHMARKING_NOTES.md](BENCHMARKING_NOTES.md).
 
 ## Benchmarks
 
@@ -9,11 +11,28 @@ Transformer Engine kernels and helper scripts for comparing benchmark CSVs.
   low-precision recipes (FP8, MXFP8, MXFP4, NVFP4) via `autocast`
 - `benchmark_grouped_gemm.py`: grouped GEMM benchmark for MoE-style shapes
 - `benchmark_casting.py`: quantize / dequantize benchmark across FP8, MXFP8, NVFP4, and MXFP4
-- `benchmark_normalization.py`: LayerNorm / RMSNorm forward benchmark across BF16 and quantized (FP8, MXFP8) output
+  (MXFP4 quantize only)
+- `benchmark_normalization.py`: LayerNorm / RMSNorm forward benchmark with BF16 output and each
+  supported quantized output (FP8, MXFP8, MXFP4, NVFP4)
 
-Run a benchmark directly from this directory. Pass `--csv` to write results.
-When no filename is provided, `run_benchmarks` derives the CSV name from the
-benchmark script file name.
+## Running
+
+Run a benchmark with `python`:
+
+```bash
+python benchmark_gemm.py                 # run, print a results table
+```
+
+Each benchmark is a pytest module under the hood (wired up by `conftest.py`),
+and the file forwards its arguments to pytest, so any pytest flag (`-k`, `-v`,
+...) and the custom options below behave the same whether you launch it with
+`python benchmark_gemm.py` or `pytest benchmark_gemm.py`. To run every family in
+one go you can use `pytest .`.
+
+### CSV output
+
+Pass `--csv` to write results. With no filename the CSV is named after the
+module (`benchmark_gemm.py` -> `benchmark_gemm.csv`):
 
 ```bash
 python benchmark_gemm.py --csv
@@ -28,8 +47,91 @@ python benchmark_gemm.py --csv --csv-samples
 python benchmark_gemm.py --csv --csv-samples gemm_samples.csv
 ```
 
-The samples CSV contains one row per timing sample with columns for all
-benchmark parameters plus `label`, `sample_idx`, and `time_ms`.
+The samples CSV (`<module>_samples.csv` by default) contains one row per timing
+sample with columns for all benchmark parameters plus `label`, `sample_idx`,
+and `time_ms`.
+
+The `benchmark_gemm.py` CSV also has a `Kernels` column with the GEMM kernels the timed
+direction launched (`;`-separated, in launch order): hipBLASLt kernel names (Tensile names
+are cut after the tile and matrix-instruction sizes and get a `~<crc32>` of the full name)
+and, for Triton, the autotuned kernel with the config it picked, e.g.
+`matmul_kernel[128x256x32 GROUP_SIZE_M=4 waves_per_eu=0 num_warps=8 num_stages=3]`. It is
+empty for the other backends.
+
+### Selecting cases
+
+Use pytest's `-k` to run a subset by parametrize id:
+
+```bash
+python benchmark_gemm.py -k "bf16 and QKV"   # select shapes/precisions
+python benchmark_gemm.py -k triton           # select the Triton backend
+python benchmark_gemm.py -k "not flydsl"     # skip FlyDSL (long compile times)
+```
+
+### Other options
+
+- `--kernel-profile`: also measure GPU kernel (device) time (see
+  [Kernel profiling](#kernel-profiling)).
+- `--no-gpu-interference-check` / `--abort-on-gpu-interference`: control the
+  shared-GPU neighbor check, which warns (or aborts) when another process is
+  using the benchmark GPU (requires the `amdsmi` package).
+- `--no-cpu-pin`: disable CPU pinning. By default the main thread is pinned to 4 physical
+  cores on the GPU's NUMA node (offset by GPU id), the HIP/ROCr helper threads to the cores'
+  SMT siblings, and the cores are set to the performance governor for the session (needs
+  root; restored at the end, also if the process is killed or crashes, by a small watchdog
+  process). The header and `run_info.txt` show the cores used; see
+  [CPU pinning](BENCHMARKING_NOTES.md#cpu-pinning) for why.
+- `--dashboard-run`: write all families of the session into one run directory for the TE
+  dashboard instead of the `--csv` outputs:
+  `<dashboard-out>/<arch>_<host>_gpu<id>_<pci>/<date>_<commit>/` with
+  `benchmark_<family>.csv` (each row tagged with the run timestamp and commit),
+  `run_info.txt` (host, GPU, software versions, CPU pinning, HW queues) and, with
+  `--csv-samples`, `samples/`. `--dashboard-out DIR` sets the base directory (default
+  `results`); `--gpu-model NAME` overrides the detected GPU model (the run is refused if
+  no model can be detected and none is given). A full dashboard run:
+
+  ```bash
+  pytest . --dashboard-run --kernel-profile --csv-samples
+  ```
+
+### Environment
+
+- `GPU_MAX_HW_QUEUES` defaults to `2`, as in Megatron's launch scripts; set it explicitly to
+  override (with `4`, HIP's default, the per-expert grouped GEMM path runs up to ~2x slower;
+  see [HW queues](BENCHMARKING_NOTES.md#hw-queues)). The header and `run_info.txt` show the
+  value used.
+- The `hipblaslt` rows of `benchmark_gemm.py` run with TE's hipBLASLt tuning
+  (`TE_HIPBLASLT_TUNING_RUN_COUNT=10`, `TE_HIPBLASLT_TUNING_ALGO_COUNT=10`), which Megatron's
+  training scripts enable for dense GEMMs. Each shape is tuned at its first call (about a
+  minute per MXFP8 shape, whose hipBLASLt kernels are compiled on the fly); the `Kernels`
+  CSV column (see [CSV output](#csv-output)) shows the kernels it picked.
+
+### Kernel profiling
+
+Pass `--kernel-profile` to report GPU kernel (device) time alongside the default
+host wall-clock time. It adds `<label> Kernel Time (ms)` / `<label> Kernel <unit>`
+columns to the console table and the CSV:
+
+```bash
+python benchmark_gemm.py --kernel-profile --csv
+```
+
+- **Wall time** is the end-to-end host-side time per call over back-to-back calls (median):
+  Python dispatch + kernel launch/queue latency + GPU execution. Forward is timed with
+  `torch.utils.benchmark.Timer`; backward runs a block of forwards untimed and then times
+  their backwards back to back, so no forward time is subtracted.
+- **Kernel time** is the GPU device time only (`torch.profiler`, summing the device
+  time of each call's kernels, reported as the median over a separate profiling pass),
+  so it excludes host and launch overhead. If no call's kernels can be matched, the
+  kernel cells are left empty and a warning names the case.
+- **Grouped GEMM kernel time is not device-only.** It is the CUDA-event span between
+  consecutive calls on the main stream, which includes GPU idle time in launch-bound cases.
+
+See [Wall time and kernel time](BENCHMARKING_NOTES.md#wall-time-and-kernel-time) for how to
+interpret the two (launch-bound gaps, profiler effects, grouped GEMM).
+
+Because it runs an extra warmup plus a profiled loop under `torch.profiler`,
+`--kernel-profile` makes each benchmark take longer.
 
 ### Rotating input buffers
 
@@ -67,24 +169,55 @@ expert count `B` in addition to `M`.
 
 ## Adding a benchmark
 
-Use `run_benchmarks(test_cases, bench_fn, param_columns)`.
+Each benchmark is a pytest module named `benchmark_<family>.py` (`conftest.py` collects that
+pattern, and the family and CSV names come from it) with three pieces:
+
+1. A case generator plus `pytest_generate_tests`, which parametrizes a `case`:
 
-- `test_cases` is a list of dictionaries containing benchmark inputs.
-- `param_columns` lists the case fields that should appear in stdout headers
-  and CSV output.
-- `bench_fn(**case)` must return a list of metric records created by
-  `make_metric_record(...)` or `make_forward_backward_metric_records(...)`.
+   ```python
+   def pytest_generate_tests(metafunc):
+       if "case" in metafunc.fixturenames:
+           cases = generate_cases()
+           metafunc.parametrize("case", cases, ids=[_case_id(c) for c in cases])
+   ```
 
-Each metric record represents one benchmark line such as `GEMM Forward`. The
-runner prints that line to stdout and expands it into two CSV columns:
+2. A `bench_*` function that runs the kernel and returns a list of metric
+   records built with `make_metric_record(...)`, or with `direction_records(...)`,
+   which times the forward or the backward of a forward/backward callable pair.
+
+3. A `@pytest.mark.benchmark` test that applies any backend env vars and hands
+   the work to the `microbench` fixture:
+
+   ```python
+   @pytest.mark.benchmark
+   def test_gemm(microbench, case, monkeypatch):
+       apply_backend_env(monkeypatch, BACKENDS[case["Backend"]])
+       microbench.run(case, lambda: bench_gemm(**case))
+   ```
+
+Start the file, right after the module docstring and before importing torch, with
+the direct-run shim so `python benchmark_x.py` works. Entering pytest before torch
+touches the GPU keeps this process out of the GPU-interference snapshot:
+
+```python
+if __name__ == "__main__":
+    import sys
+
+    import pytest
+
+    raise SystemExit(pytest.main([__file__, *sys.argv[1:]]))
+```
 
-- `<label> Time (ms)`
-- `<label> <unit>`
+`microbench.run(case, bench_callable)` prints the case, records its metrics, and
+`conftest.py` collects them per family and writes the CSV. Each metric record
+represents one line such as `Normalization Forward` and expands into two CSV columns:
 
-For example, a `GEMM Forward` metric with unit `TFLOPS` becomes:
+- `<label> Wall Time (ms)`
+- `<label> Wall <unit>`
 
-- `GEMM Forward Time (ms)`
-- `GEMM Forward TFLOPS`
+For example, a `Normalization Forward` metric with unit `GB/s` becomes
+`Normalization Forward Wall Time (ms)` and `Normalization Forward Wall GB/s`. Passing
+`--kernel-profile` adds matching `<label> Kernel Time (ms)` / `<label> Kernel <unit>` columns.
 
 ## Comparing results
 
@@ -96,4 +229,6 @@ python compare_results.py baseline.csv candidate.csv --bench-name GEMM
 ```
 
 The script auto-detects metric columns, computes speedups for overlapping rows,
-and reports rows that exist only in the baseline or only in the candidate.
+and reports rows that exist only in the baseline or only in the candidate. It ignores the
+run-tag columns of `--dashboard-run` CSVs (`run_ts`, `commit_sha`, `commit_date`), so two
+dashboard runs can be compared directly.
diff --git a/benchmarks/microbenchmarks/benchmark_casting.py b/benchmarks/microbenchmarks/benchmark_casting.py
index 2221d1736f..e098d0e8ed 100755
--- a/benchmarks/microbenchmarks/benchmark_casting.py
+++ b/benchmarks/microbenchmarks/benchmark_casting.py
@@ -7,6 +7,9 @@
 """
 Low-precision casting micro-benchmark.
 
+Run with ``python benchmark_casting.py`` (a pytest module under the hood; see
+conftest.py).
+
 Benchmarks quantization (BF16 -> low precision) and dequantization
 (low precision -> BF16) for the formats used by TE training recipes:
 
@@ -25,6 +28,15 @@
 Output: benchmark_casting.csv (written to cwd)
 """
 
+if __name__ == "__main__":
+    # Enter pytest before torch touches the GPU, so conftest's neighbor snapshot excludes us.
+    import sys
+
+    import pytest
+
+    raise SystemExit(pytest.main([__file__, *sys.argv[1:]]))
+
+import pytest
 import torch
 import transformer_engine
 import transformer_engine_torch as tex
@@ -38,8 +50,8 @@
 )
 from utils import (
     MODEL_HIDDEN_SIZES, M_SIZE_LIST,
-    time_func, compute_gbps, make_metric_record, run_benchmarks,
-    make_input, rotating,
+    apply_backend_env, time_func_dual, compute_gbps, make_metric_record,
+    make_input, rotating, te_honors_env,
 )
 
 TE_FP8_E4M3 = tex.DType.kFloat8E4M3
@@ -69,24 +81,24 @@ def build():
 #   MXFP4 : 0.5 data + E8M0 1 byte / 32-elem block  -> 0.5 + 1/32
 # MXFP4 has no packed-FP4 dequantize kernel yet, so it runs the quantize direction only.
 _CAST_FORMATS = (
-    ("FP8-E4M3", _fp8_quantizer(TE_FP8_E4M3), 1.0, check_fp8_support, True),
-    ("FP8-E5M2", _fp8_quantizer(TE_FP8_E5M2), 1.0, check_fp8_support, True),
+    ("fp8-e4m3", _fp8_quantizer(TE_FP8_E4M3), 1.0, check_fp8_support, True),
+    ("fp8-e5m2", _fp8_quantizer(TE_FP8_E5M2), 1.0, check_fp8_support, True),
     (
-        "MXFP8-E4M3",
+        "mxfp8-e4m3",
         lambda: MXFP8Quantizer(TE_FP8_E4M3, rowwise=True, columnwise=False),
         1.0 + 1.0 / 32,
         check_mxfp8_support,
         True,
     ),
     (
-        "MXFP8-E5M2",
+        "mxfp8-e5m2",
         lambda: MXFP8Quantizer(TE_FP8_E5M2, rowwise=True, columnwise=False),
         1.0 + 1.0 / 32,
         check_mxfp8_support,
         True,
     ),
     (
-        "NVFP4",
+        "nvfp4",
         lambda: NVFP4Quantizer(
             fp4_dtype=TE_FP4_E2M1, rowwise=True, columnwise=False, with_rht=False
         ),
@@ -95,7 +107,7 @@ def build():
         True,
     ),
     (
-        "MXFP4",
+        "mxfp4",
         lambda: MXFP4Quantizer(fp4_dtype=TE_FP4_E2M1, rowwise=True, columnwise=False),
         0.5 + 1.0 / 32,
         check_mxfp4_support,
@@ -118,37 +130,73 @@ def _active_formats():
     return formats
 
 
-def _generate_test_cases():
-    test_cases = []
-    active = _active_formats()
+# Backend axis (None unsets, so "hip" is the native C++ path even if the ambient
+# env has a toggle set). "triton" flips the Triton kernel for the op being timed:
+# quantize -> NVTE_USE_CAST_TRANSPOSE_TRITON, dequantize -> NVTE_USE_DEQUANTIZE_TRITON.
+CAST_BACKENDS = {
+    "hip": {"NVTE_USE_CAST_TRANSPOSE_TRITON": None, "NVTE_USE_DEQUANTIZE_TRITON": None},
+    "triton": {"NVTE_USE_CAST_TRANSPOSE_TRITON": "1", "NVTE_USE_DEQUANTIZE_TRITON": "1"},
+}
+
+_FORMATS = None
+
+
+def _triton_applies(fmt, direction):
+    # Cast-transpose Triton covers FP8/MXFP8/MXFP4 quantize (not NVFP4); the
+    # dequantize Triton path exists only for MXFP8 (mxfp8_tensor_storage).
+    if direction == "quantize":
+        return fmt != "nvfp4"
+    return fmt.startswith("mxfp8")
+
+
+def _backends_for(fmt, direction):
+    return ["hip", "triton"] if _triton_applies(fmt, direction) else ["hip"]
+
+
+def _formats():
+    """{format_name: (quantizer_factory, quantized_bytes/elem, dequant_supported)}."""
+    global _FORMATS
+    if _FORMATS is None:
+        _FORMATS = {
+            name: (make_quantizer, q_bytes, dequant_supported)
+            for name, make_quantizer, q_bytes, dequant_supported in _active_formats()
+        }
+    return _FORMATS
+
+
+def generate_cases():
+    """Cross models x cast format x direction x backend x M."""
+    cases = []
     for model_name, hidden in MODEL_HIDDEN_SIZES:
-        for fmt_name, make_quantizer, q_bytes_per_elem, dequant_supported in active:
+        for fmt_name, (_mk, _qb, dequant_supported) in _formats().items():
             for direction in DIRECTIONS:
                 if direction == "dequantize" and not dequant_supported:
                     continue
-                cast_name = (
-                    f"BF16-to-{fmt_name}" if direction == "quantize" else f"{fmt_name}-to-BF16"
-                )
-                for M in M_SIZE_LIST:
-                    test_cases.append({
-                        "Case": f"{model_name}/{cast_name}",
-                        "M": M,
-                        "hidden_size": hidden,
-                        "direction": direction,
-                        "make_quantizer": make_quantizer,
-                        "q_bytes_per_elem": q_bytes_per_elem,
-                        "dtype_str": cast_name,
-                    })
-    return test_cases
-
-
-def bench_cast(Case, M, hidden_size, direction, make_quantizer, q_bytes_per_elem, dtype_str):
+                for backend in _backends_for(fmt_name, direction):
+                    for M in M_SIZE_LIST:
+                        cases.append({
+                            "Case": model_name,
+                            "Format": fmt_name,
+                            "Direction": direction,
+                            "Backend": backend,
+                            "M": M,
+                            "hidden_size": hidden,
+                        })
+    return cases
+
+
+def _case_id(c):
+    return f"{c['Case']}-{c['Format']}-{c['Direction']}-{c['Backend']}-M{c['M']}"
+
+
+def bench_cast(Format, Direction, M, hidden_size):
     device = "cuda"
 
+    make_quantizer, q_bytes_per_elem, _deq = _formats()[Format]
     numel = M * hidden_size
     quantizer = make_quantizer()
 
-    if direction == "quantize":
+    if Direction == "quantize":
         next_x = make_input((M, hidden_size), torch.bfloat16, device=device)
         out = quantizer(next_x())
         cast_func = lambda: quantizer.quantize(next_x(), out=out)
@@ -164,15 +212,29 @@ def bench_cast(Case, M, hidden_size, direction, make_quantizer, q_bytes_per_elem
         cast_func = lambda: next_q().dequantize()
         total_bytes = int(numel * (q_bytes_per_elem + 2))  # quantized read + BF16 write
 
-    ms, measurement = time_func(cast_func, method="blocked")
-    gbps = compute_gbps(total_bytes, ms)
-
-    return [make_metric_record(CAST_LABEL, ms, "GB/s", gbps, measurement=measurement)]
-
-
-if __name__ == "__main__":
-    run_benchmarks(
-        test_cases=_generate_test_cases(),
-        bench_fn=bench_cast,
-        param_columns=["Case", "M", "hidden_size", "dtype_str"],
+    ms, measurement, kernel_ms = time_func_dual(cast_func, method="blocked")
+    return [make_metric_record(
+        CAST_LABEL, ms, "GB/s", compute_gbps(total_bytes, ms), measurement=measurement,
+        kernel_ms=kernel_ms,
+        kernel_throughput=compute_gbps(total_bytes, kernel_ms) if kernel_ms else None,
+    )]
+
+
+def pytest_generate_tests(metafunc):
+    if "case" in metafunc.fixturenames:
+        cases = generate_cases()
+        metafunc.parametrize("case", cases, ids=[_case_id(c) for c in cases])
+
+
+@pytest.mark.benchmark
+def test_cast(microbench, case, monkeypatch):
+    if case["Backend"] == "triton":
+        var = ("NVTE_USE_CAST_TRANSPOSE_TRITON" if case["Direction"] == "quantize"
+               else "NVTE_USE_DEQUANTIZE_TRITON")
+        if not te_honors_env(var):
+            pytest.skip("Triton cast backend not available in this TE build")
+    apply_backend_env(monkeypatch, CAST_BACKENDS[case["Backend"]])
+    microbench.run(
+        case,
+        lambda: bench_cast(case["Format"], case["Direction"], case["M"], case["hidden_size"]),
     )
diff --git a/benchmarks/microbenchmarks/benchmark_gemm.py b/benchmarks/microbenchmarks/benchmark_gemm.py
index 29445179f9..7cd07867c0 100755
--- a/benchmarks/microbenchmarks/benchmark_gemm.py
+++ b/benchmarks/microbenchmarks/benchmark_gemm.py
@@ -4,40 +4,205 @@
 #
 # See LICENSE for license information.
 ###############################################################################
-"""Dense GEMM micro-benchmark using te.Linear across precisions.
+"""Dense GEMM micro-benchmark using te.Linear across precisions and backends.
 
-Sweeps the shared model GEMM shapes over BF16 (the high-precision baseline)
-plus every supported low-precision recipe (FP8, MXFP8, MXFP4, NVFP4) via
-te.autocast. Precisions whose hardware/runtime support is unavailable on the
-current device are skipped automatically.
+Run with ``python benchmark_gemm.py`` (a pytest module under the hood; see
+conftest.py). Sweeps the shared model GEMM shapes over BF16 (the high-precision
+baseline) plus every supported low-precision recipe (FP8, MXFP8, MXFP4, NVFP4)
+via te.autocast, crossed with a selectable kernel Backend. Precisions whose
+hardware/runtime support is unavailable on the current device are skipped
+automatically.
 
-Output: benchmark_gemm.csv (written to cwd)
+Examples::
+
+    python benchmark_gemm.py --csv                   # -> benchmark_gemm.csv
+    python benchmark_gemm.py -k "bf16 and QKV"       # select shapes/precisions
+    python benchmark_gemm.py -k triton               # select the Triton backend
+    python benchmark_gemm.py -k "not flydsl"         # skip FlyDSL (long compile times)
+
+Output: benchmark_gemm.csv (written to cwd when --csv is passed).
 """
 
+if __name__ == "__main__":
+    # Enter pytest before torch touches the GPU, so conftest's neighbor snapshot excludes us.
+    import sys
+
+    import pytest
+
+    raise SystemExit(pytest.main([__file__, *sys.argv[1:]]))
+
+import functools
+import importlib.util
+import os
+import warnings
+
+import pytest
 import torch
 import transformer_engine.pytorch as te
 from utils import (
     build_recipes,
     generate_gemm_test_cases,
-    time_func, compute_tflops, make_forward_backward_metric_records, run_benchmarks,
-    make_input,
+    apply_backend_env, compute_tflops, direction_records,
+    gemm_kernel_names, make_input, te_honors_env,
 )
 
 BENCHMARK_LABEL = "GEMM"
 
 RECIPES = build_recipes()
 
+# Env recipes to force a dense-GEMM kernel backend (None unsets the var). Per the
+# dispatch: bf16 defaults to hipBLASLt; Triton is NVTE_GEMM_BACKEND=TRITON (PR #676,
+# unified selector) or the legacy NVTE_USE_GEMM_TRITON=1 toggle (PR #667), resolved
+# per build by _triton_gemm_env. mxfp8 defaults to HipKittens, forced to hipBLASLt via
+# NVTE_ROCM_USE_HIPBLASLT_MXFP8 (rocm_gemm.cu). fp8 has a single backend. FlyDSL
+# (gfx950-only, BF16/FP8/MXFP8) is NVTE_GEMM_BACKEND=FLYDSL (cpp_extensions/gemm.py).
+# mxfp4 runs AITER a4w4 unless NVTE_ROCM_USE_HIPBLASLT_MXFP4=1 (cpp_extensions/gemm.py);
+# that hipBLASLt path fails in the te.Linear backward, so it isn't benchmarked.
+_GEMM_TRITON = "NVTE_USE_GEMM_TRITON"
+_HIPBLASLT_MXFP8 = "NVTE_ROCM_USE_HIPBLASLT_MXFP8"
+_HIPBLASLT_MXFP4 = "NVTE_ROCM_USE_HIPBLASLT_MXFP4"
+_GEMM_BACKEND = "NVTE_GEMM_BACKEND"
+
+_HIPBLASLT_TUNING = {"TE_HIPBLASLT_TUNING_RUN_COUNT": "10", "TE_HIPBLASLT_TUNING_ALGO_COUNT": "10"}
+
+# Triton's env is resolved at runtime by _triton_gemm_env (mechanism varies by build).
+GEMM_BACKENDS = {
+    "hipblaslt":  {_GEMM_TRITON: None, _HIPBLASLT_MXFP8: "1", _GEMM_BACKEND: None,
+                   **_HIPBLASLT_TUNING},
+    "hipkittens": {_GEMM_TRITON: None, _HIPBLASLT_MXFP8: None, _GEMM_BACKEND: None},
+    "flydsl":     {_GEMM_TRITON: None, _HIPBLASLT_MXFP8: None, _GEMM_BACKEND: "FLYDSL"},
+    "aiter":      {_GEMM_TRITON: None, _HIPBLASLT_MXFP4: None, _GEMM_BACKEND: None},
+}
+
+# Backends with a real choice per precision (the supported-backends table).
+_BACKENDS_BY_PRECISION = {
+    "bf16": ["hipblaslt", "triton", "flydsl"],
+    "fp8": ["hipblaslt", "flydsl"],
+    "mxfp8": ["hipblaslt", "hipkittens", "flydsl"],
+    "mxfp4": ["aiter"],
+}
+
+
+def _backends_for(precision):
+    return _BACKENDS_BY_PRECISION.get(precision, ["hipblaslt"])
+
+
+@functools.lru_cache(maxsize=1)
+def _triton_gemm_env():
+    """Env that forces the Triton GEMM backend on the installed build, or None.
+
+    PR #676 unified selection under NVTE_GEMM_BACKEND=TRITON; older builds (PR
+    #667) use the legacy NVTE_USE_GEMM_TRITON=1 toggle. Prefer the unified selector.
+    """
+    if te_honors_env(_GEMM_BACKEND):
+        return {_GEMM_BACKEND: "TRITON", _GEMM_TRITON: None, _HIPBLASLT_MXFP8: None}
+    if te_honors_env(_GEMM_TRITON):
+        return {_GEMM_TRITON: "1", _GEMM_BACKEND: None, _HIPBLASLT_MXFP8: None}
+    return None
+
+
+def _triton_gemm_supported():
+    return _triton_gemm_env() is not None
+
+
+@functools.lru_cache(maxsize=1)
+def _flydsl_gemm_supported():
+    if not te_honors_env(_GEMM_BACKEND):
+        return False
+    try:
+        from transformer_engine.pytorch.utils import get_device_compute_capability
+        if get_device_compute_capability() != (9, 5):
+            return False
+        from transformer_engine.pytorch.flydsl_kernels.gemm import (  # noqa: F401
+            te_generic_gemm_flydsl,
+        )
+        return True
+    except Exception:
+        return False
+
 
-def generate_precision_gemm_test_cases():
-    """Cross the shared dense GEMM shapes with each supported precision."""
-    test_cases = []
-    for base_case in generate_gemm_test_cases():
+@functools.lru_cache(maxsize=1)
+def _device_cc():
+    """(major, minor) compute capability of the current GPU, or None."""
+    try:
+        from transformer_engine.pytorch.utils import get_device_compute_capability
+        return get_device_compute_capability()
+    except Exception:
+        return None
+
+
+def _mxfp8_gemm_k_multiple():
+    """Multiple hipBLASLt requires for an MXFP8 GEMM's contraction dim: 128 on
+    gfx950 (swizzled scales unsupported), 32 on gfx1250 (see rocm_gemm.cu)."""
+    return 32 if _device_cc() == (12, 5) else 128
+
+
+def _flydsl_fell_back(fn):
+    """Run *fn* once with FlyDSL fallback warnings on; True if FlyDSL fell back to C++."""
+    prev = os.environ.get("NVTE_FLYDSL_GEMM_WARN_FALLBACK")
+    os.environ["NVTE_FLYDSL_GEMM_WARN_FALLBACK"] = "1"
+    try:
+        with warnings.catch_warnings(record=True) as caught:
+            warnings.simplefilter("always")
+            fn()
+        return any("FlyDSL GEMM does not support" in str(w.message) for w in caught)
+    finally:
+        if prev is None:
+            os.environ.pop("NVTE_FLYDSL_GEMM_WARN_FALLBACK", None)
+        else:
+            os.environ["NVTE_FLYDSL_GEMM_WARN_FALLBACK"] = prev
+
+
+def generate_cases():
+    """Cross the shared dense GEMM shapes with each precision, backend, direction."""
+    cases = []
+    for base in generate_gemm_test_cases():
         for precision in RECIPES:
-            test_cases.append({**base_case, "Precision": precision})
-    return test_cases
+            for backend in _backends_for(precision):
+                for direction in ("fwd", "bwd"):
+                    cases.append({
+                        "Case": base["Case"],
+                        "Precision": precision,
+                        "Backend": backend,
+                        "Direction": direction,
+                        "M": base["M"],
+                        "N": base["N"],
+                        "K": base["K"],
+                        "dtype": base["dtype"],
+                    })
+    return cases
+
+
+def _case_id(c):
+    return f"{c['Case']}-{c['Precision']}-{c['Backend']}-{c['Direction']}-M{c['M']}"
 
 
-def bench_gemm(Case, Precision, M, N, K, dtype):
+def _known_hipblaslt_qkv_xfail(c):
+    """Llama3.1-405B/TP8-QKV bf16/nvfp4 hipBLASLt GEMM hits a hipBLASLt algo bug --
+    the top heuristic algo returns HIPBLASLT status 6 (INTERNAL_ERROR) at small M on
+    gfx950. nvfp4 dequantizes to bf16, same GEMM."""
+    return (
+        c["Case"] == "Llama3.1-405B/TP8-QKV"
+        and c["Backend"] == "hipblaslt"
+        and c["Precision"] in ("bf16", "nvfp4")
+        and c["M"] in (1024, 2048)
+        and _device_cc() == (9, 5)
+    )
+
+
+def _marks_for(c):
+    """pytest marks for one case: a temporary xfail for the known hipBLASLt QKV bug
+    (remove once hipBLASLt / hipblaslt_gemm is fixed)."""
+    marks = []
+    if _known_hipblaslt_qkv_xfail(c):
+        marks.append(pytest.mark.xfail(
+            reason="known hipBLASLt algo bug on Llama3.1-405B/TP8-QKV bf16/nvfp4 at "
+                   "small M (gfx950, HIPBLASLT status 6)",
+            raises=RuntimeError, strict=False))
+    return tuple(marks)
+
+
+def bench_gemm(Case, Precision, Direction, M, N, K, dtype, kernel_names=False):
     device = "cuda"
 
     recipe = RECIPES[Precision]
@@ -47,48 +212,80 @@ def bench_gemm(Case, Precision, M, N, K, dtype):
     next_x = make_input((M, K), dtype, device=device, requires_grad=True)
 
     def fwd_func():
+        xb = next_x()
         with te.autocast(enabled=use_fp8, recipe=recipe):
-            return linear(next_x())
+            return linear(xb), xb
 
-    out = fwd_func()
-    grad_out = torch.randn_like(out)
+    grad_out = torch.randn_like(fwd_func()[0])
 
-    def fwd_bwd_func():
-        xb = next_x()
-        with te.autocast(enabled=use_fp8, recipe=recipe):
-            out = linear(xb)
-            out.backward(grad_out)
+    def bwd_func(fwd_out):
+        out, xb = fwd_out
+        out.backward(grad_out)
         xb.grad = None
         linear.weight.grad = None
 
-    fwd_bwd_func()
+    if os.environ.get("NVTE_GEMM_BACKEND") == "FLYDSL" and _flydsl_fell_back(
+        (lambda: bwd_func(fwd_func())) if Direction == "bwd" else fwd_func
+    ):
+        pytest.skip("FlyDSL GEMM fell back to C++ for this shape/direction")
 
     fwd_flops = 2 * M * N * K
-    bwd_flops = 2 * fwd_flops  # dX + dW
-
-    fwd_ms, fwd_measurement = time_func(fwd_func)
-    fwd_bwd_ms, fwd_bwd_measurement = time_func(fwd_bwd_func)
-    bwd_ms = fwd_bwd_ms - fwd_ms
-
-    fwd_tflops = compute_tflops(fwd_flops, fwd_ms)
-    bwd_tflops = compute_tflops(bwd_flops, bwd_ms)
-
-    return make_forward_backward_metric_records(
-        BENCHMARK_LABEL,
-        "TFLOPS",
-        fwd_ms,
-        fwd_tflops,
-        bwd_ms,
-        bwd_tflops,
-        backward_derived=True,
-        fwd_measurement=fwd_measurement,
-        fwd_bwd_measurement=fwd_bwd_measurement,
+    records = direction_records(
+        Direction, BENCHMARK_LABEL, "TFLOPS", compute_tflops,
+        fwd_func, bwd_func, fwd_flops, 2 * fwd_flops,
     )
+    timed, prepare = (fwd_func, None) if Direction == "fwd" else (bwd_func, fwd_func)
+    records[0]["kernel_names"] = gemm_kernel_names(timed, prepare) if kernel_names else ""
+    return records
 
 
-if __name__ == "__main__":
-    run_benchmarks(
-        test_cases=generate_precision_gemm_test_cases(),
-        bench_fn=bench_gemm,
-        param_columns=["Case", "Precision", "M", "N", "K", "dtype"],
+def pytest_generate_tests(metafunc):
+    if "case" in metafunc.fixturenames:
+        cases = generate_cases()
+        params = [
+            pytest.param(c, id=_case_id(c), marks=_marks_for(c))
+            for c in cases
+        ]
+        metafunc.parametrize("case", params)
+
+
+@pytest.mark.benchmark
+def test_gemm(microbench, case, monkeypatch):
+    if case["Precision"] == "mxfp4" and any(
+        dim % 32 for dim in (case["M"], case["N"], case["K"])
+    ):
+        pytest.skip("MXFP4 GEMM needs M/N/K divisible by 32")
+    if case["Precision"] == "mxfp8":
+        # Block-quantize needs every operand dim divisible by 32; the GEMM needs
+        # each contraction dim divisible by _mxfp8_gemm_k_multiple() -- K for fwd,
+        # plus N and M for the bwd dgrad/wgrad.
+        if any(dim % 32 for dim in (case["M"], case["N"], case["K"])):
+            pytest.skip("MXFP8 GEMM needs M/N/K divisible by 32")
+        kmul = _mxfp8_gemm_k_multiple()
+        contraction = ((case["K"],) if case["Direction"] == "fwd"
+                       else (case["K"], case["N"], case["M"]))
+        if any(dim % kmul for dim in contraction):
+            pytest.skip(f"MXFP8 GEMM needs the contraction dim divisible by {kmul}")
+    if case["Backend"] == "triton" and not _triton_gemm_supported():
+        pytest.skip("Triton GEMM backend not available in this TE build")
+    if case["Backend"] == "triton" and _device_cc() == (12, 5):
+        pytest.skip("Triton GEMM hangs or crashes on many shapes on gfx1250")
+    if case["Backend"] == "hipkittens" and not te_honors_env(_HIPBLASLT_MXFP8):
+        pytest.skip("HipKittens GEMM backend not available in this TE build")
+    if case["Backend"] == "hipkittens" and (case["N"] % 256 or case["K"] % 256):
+        pytest.skip("HipKittens GEMM needs 256-aligned N/K")
+    if case["Backend"] == "flydsl" and not _flydsl_gemm_supported():
+        pytest.skip("FlyDSL GEMM backend not available (needs gfx950 + flydsl package)")
+    if case["Backend"] == "aiter" and importlib.util.find_spec("aiter") is None:
+        pytest.skip("aiter is not installed (required for the MXFP4 GEMM backend)")
+    # Triton's selector varies by TE version (NVTE_GEMM_BACKEND=TRITON or legacy toggle).
+    backend_env = _triton_gemm_env() if case["Backend"] == "triton" else GEMM_BACKENDS[case["Backend"]]
+    apply_backend_env(monkeypatch, backend_env)
+    microbench.run(
+        case,
+        lambda: bench_gemm(
+            case["Case"], case["Precision"], case["Direction"],
+            case["M"], case["N"], case["K"], case["dtype"],
+            kernel_names=case["Backend"] in ("hipblaslt", "triton"),
+        ),
     )
diff --git a/benchmarks/microbenchmarks/benchmark_grouped_gemm.py b/benchmarks/microbenchmarks/benchmark_grouped_gemm.py
index d95e184c4c..b476efbebf 100755
--- a/benchmarks/microbenchmarks/benchmark_grouped_gemm.py
+++ b/benchmarks/microbenchmarks/benchmark_grouped_gemm.py
@@ -4,26 +4,184 @@
 #
 # See LICENSE for license information.
 ###############################################################################
+"""Grouped GEMM micro-benchmark using te.GroupedLinear across precisions and backends.
 
+Run with ``python benchmark_grouped_gemm.py`` (a pytest module under the hood;
+see conftest.py). Sweeps MoE grouped-GEMM shapes over BF16, FP8, MXFP8, and
+NVFP4, crossed with the selectable kernel backend for each precision (hipBLASLt /
+CK_Tile / Triton / HipKittens) and forward/backward direction.
+
+    python benchmark_grouped_gemm.py --csv
+    python benchmark_grouped_gemm.py -k "mxfp8 and hipkittens_grouped"
+"""
+
+if __name__ == "__main__":
+    # Enter pytest before torch touches the GPU, so conftest's neighbor snapshot excludes us.
+    import sys
+
+    import pytest
+
+    raise SystemExit(pytest.main([__file__, *sys.argv[1:]]))
+
+import os
+import sys
+import tempfile
+
+import pytest
 import torch
 import transformer_engine.pytorch as te
+from transformer_engine.pytorch.utils import get_device_compute_capability
 from utils import (
     DTYPE_LIST,
+    apply_backend_env,
     build_recipes,
-    time_func,
     compute_tflops,
-    make_forward_backward_metric_records,
-    run_benchmarks,
+    direction_records,
     make_input,
+    te_honors_env,
 )
 
 BENCHMARK_LABEL = "Grouped GEMM"
 
-# Same precision sweep as benchmark_gemm.py, minus MXFP4 (GroupedLinear has no
-# MXFP4 grouped kernel). Each test case carries a recipe label; bf16 maps to
-# None (plain path) and unsupported precisions are skipped by build_recipes().
+# FIXME: Add mxfp4 when https://github.com/ROCm/TransformerEngine/pull/745 is merged
 RECIPES = build_recipes(names=("bf16", "fp8", "mxfp8", "nvfp4"))
 
+# Env recipes to force a grouped-GEMM kernel backend (None unsets the var). Per the
+# C++ dispatch (cublaslt_gemm.cu / rocm_gemm.cu): all-unset -> one GEMM per expert on
+# the compute streams; CUTLASS+CK -> CK; CUTLASS+HK -> HipKittens; NVTE_USE_GROUPED_GEMM_TRITON
+# routes bf16 to the Triton grouped GEMM. Per-expert MXFP8 GEMMs default to dense HipKittens
+# on gfx950 ("hipkittens_per_expert", TE's default MXFP8 path), so "hipblaslt" forces hipBLASLt
+# via NVTE_ROCM_USE_HIPBLASLT_MXFP8.
+_CUTLASS = "NVTE_USE_CUTLASS_GROUPED_GEMM"
+_CK = "NVTE_USE_CK_GROUPED_GEMM"
+_HK = "NVTE_USE_HIPKITTENS_GROUPED_GEMM"
+_TRITON = "NVTE_USE_GROUPED_GEMM_TRITON"
+_HIPBLASLT_MXFP8 = "NVTE_ROCM_USE_HIPBLASLT_MXFP8"
+
+GROUPED_BACKENDS = {
+    "hipblaslt":  {_CUTLASS: None, _CK: None, _HK: None, _TRITON: None, _HIPBLASLT_MXFP8: "1"},
+    "hipkittens_per_expert": {_CUTLASS: None, _CK: None, _HK: None, _TRITON: None, _HIPBLASLT_MXFP8: None},
+    "ck_tile":    {_CUTLASS: "1", _CK: "1", _HK: None, _TRITON: None, _HIPBLASLT_MXFP8: None},
+    "hipkittens_grouped": {_CUTLASS: "1", _CK: None, _HK: "1", _TRITON: None, _HIPBLASLT_MXFP8: None},
+    "triton":     {_CUTLASS: None, _CK: None, _HK: None, _TRITON: "1", _HIPBLASLT_MXFP8: None},
+}
+
+_BACKENDS_BY_PRECISION = {
+    "bf16": ["hipblaslt", "ck_tile", "triton"],
+    "fp8": ["hipblaslt", "ck_tile"],
+    "mxfp8": ["hipblaslt", "hipkittens_per_expert", "hipkittens_grouped", "ck_tile"],
+    "nvfp4": ["hipblaslt"],
+}
+
+
+# (Case, recipe, backend, direction, B, M) grouped-GEMM configs that hang hipBLASLt on
+# gfx950
+_GFX950_HANG_CONFIGS = {
+    ("DSV3-GateUP", "bf16", "hipblaslt", "fwd", 8, 1024),
+    ("DSV3-GateUP", "bf16", "hipblaslt", "bwd", 8, 1024),
+    ("DSV3-GateUP", "bf16", "hipblaslt", "fwd", 16, 1024),
+    ("DSV3-GateUP", "bf16", "hipblaslt", "bwd", 16, 1024),
+    ("DSV3-GateUP", "bf16", "hipblaslt", "fwd", 32, 1024),
+    ("DSV3-GateUP", "bf16", "hipblaslt", "bwd", 32, 1024),
+    ("DSV3-Down", "nvfp4", "hipblaslt", "bwd", 8, 512),
+    ("DSV3-Down", "nvfp4", "hipblaslt", "bwd", 16, 512),
+    ("DSV3-Down", "nvfp4", "hipblaslt", "bwd", 32, 512),
+}
+
+
+# (Case, recipe, backend, direction, B, M) nvfp4 grouped configs that fail on gfx950 with
+# 'NVFP4 GEMM requires ... workspace': FP4 dequant scratch exceeds the fixed grouped-GEMM
+# workspace.
+_GFX950_XFAIL_CONFIGS = {
+    ("DSV2-GateUP", "nvfp4", "hipblaslt", "fwd", 5, 1024),
+    ("DSV2-GateUP", "nvfp4", "hipblaslt", "bwd", 5, 1024),
+    ("DSV2-Down", "nvfp4", "hipblaslt", "bwd", 5, 2048),
+    ("DSV2-GateUP", "nvfp4", "hipblaslt", "fwd", 5, 4096),
+    ("DSV2-GateUP", "nvfp4", "hipblaslt", "bwd", 5, 4096),
+    ("DSV2-GateUP", "nvfp4", "hipblaslt", "fwd", 10, 1024),
+    ("DSV2-GateUP", "nvfp4", "hipblaslt", "bwd", 10, 1024),
+    ("DSV2-Down", "nvfp4", "hipblaslt", "bwd", 10, 2048),
+    ("DSV2-GateUP", "nvfp4", "hipblaslt", "fwd", 10, 4096),
+    ("DSV2-GateUP", "nvfp4", "hipblaslt", "bwd", 10, 4096),
+    ("DSV2-GateUP", "nvfp4", "hipblaslt", "fwd", 20, 1024),
+    ("DSV2-GateUP", "nvfp4", "hipblaslt", "bwd", 20, 1024),
+    ("DSV2-Down", "nvfp4", "hipblaslt", "bwd", 20, 2048),
+    ("DSV2-GateUP", "nvfp4", "hipblaslt", "fwd", 20, 4096),
+    ("DSV2-GateUP", "nvfp4", "hipblaslt", "bwd", 20, 4096),
+    ("DSV3-GateUP", "nvfp4", "hipblaslt", "fwd", 8, 1024),
+    ("DSV3-GateUP", "nvfp4", "hipblaslt", "bwd", 8, 1024),
+    ("DSV3-Down", "nvfp4", "hipblaslt", "bwd", 8, 1024),
+    ("DSV3-GateUP", "nvfp4", "hipblaslt", "fwd", 8, 2048),
+    ("DSV3-GateUP", "nvfp4", "hipblaslt", "bwd", 8, 2048),
+    ("DSV3-GateUP", "nvfp4", "hipblaslt", "fwd", 8, 4096),
+    ("DSV3-GateUP", "nvfp4", "hipblaslt", "bwd", 8, 4096),
+    ("DSV3-Down", "nvfp4", "hipblaslt", "bwd", 8, 4096),
+    ("DSV3-GateUP", "nvfp4", "hipblaslt", "fwd", 16, 1024),
+    ("DSV3-GateUP", "nvfp4", "hipblaslt", "bwd", 16, 1024),
+    ("DSV3-Down", "nvfp4", "hipblaslt", "bwd", 16, 1024),
+    ("DSV3-GateUP", "nvfp4", "hipblaslt", "fwd", 16, 2048),
+    ("DSV3-GateUP", "nvfp4", "hipblaslt", "bwd", 16, 2048),
+    ("DSV3-GateUP", "nvfp4", "hipblaslt", "fwd", 16, 4096),
+    ("DSV3-GateUP", "nvfp4", "hipblaslt", "bwd", 16, 4096),
+    ("DSV3-Down", "nvfp4", "hipblaslt", "bwd", 16, 4096),
+    ("DSV3-GateUP", "nvfp4", "hipblaslt", "fwd", 32, 1024),
+    ("DSV3-GateUP", "nvfp4", "hipblaslt", "bwd", 32, 1024),
+    ("DSV3-Down", "nvfp4", "hipblaslt", "bwd", 32, 1024),
+    ("DSV3-GateUP", "nvfp4", "hipblaslt", "fwd", 32, 2048),
+    ("DSV3-GateUP", "nvfp4", "hipblaslt", "bwd", 32, 2048),
+    ("DSV3-GateUP", "nvfp4", "hipblaslt", "fwd", 32, 4096),
+    ("DSV3-GateUP", "nvfp4", "hipblaslt", "bwd", 32, 4096),
+    ("DSV3-Down", "nvfp4", "hipblaslt", "bwd", 32, 4096),
+    ("Grok-V2-GateUP", "nvfp4", "hipblaslt", "fwd", 1, 512),
+    ("Grok-V2-GateUP", "nvfp4", "hipblaslt", "bwd", 1, 512),
+    ("Grok-V2-Down", "nvfp4", "hipblaslt", "fwd", 1, 512),
+    ("Grok-V2-Down", "nvfp4", "hipblaslt", "bwd", 1, 512),
+    ("Grok-V2-GateUP", "nvfp4", "hipblaslt", "fwd", 1, 1024),
+    ("Grok-V2-GateUP", "nvfp4", "hipblaslt", "bwd", 1, 1024),
+    ("Grok-V2-Down", "nvfp4", "hipblaslt", "fwd", 1, 1024),
+    ("Grok-V2-Down", "nvfp4", "hipblaslt", "bwd", 1, 1024),
+    ("Grok-V2-GateUP", "nvfp4", "hipblaslt", "fwd", 1, 2048),
+    ("Grok-V2-GateUP", "nvfp4", "hipblaslt", "bwd", 1, 2048),
+    ("Grok-V2-Down", "nvfp4", "hipblaslt", "fwd", 1, 2048),
+    ("Grok-V2-Down", "nvfp4", "hipblaslt", "bwd", 1, 2048),
+    ("Grok-V2-GateUP", "nvfp4", "hipblaslt", "fwd", 1, 4096),
+    ("Grok-V2-GateUP", "nvfp4", "hipblaslt", "bwd", 1, 4096),
+    ("Grok-V2-Down", "nvfp4", "hipblaslt", "fwd", 1, 4096),
+    ("Grok-V2-Down", "nvfp4", "hipblaslt", "bwd", 1, 4096),
+}
+
+
+def _backends_for(recipe):
+    return _BACKENDS_BY_PRECISION.get(recipe, ["hipblaslt"])
+
+
+def _grouped_fallback_count(fn):
+    """Run *fn* once with grouped-GEMM fallback warnings on; return how many of its
+    grouped GEMMs fell back to the per-expert path. Counts the dispatcher's per-call
+    NVTE_WARN (HipKittens' own messages are latched once per process in C++); it is C++
+    stderr, so capture at the fd level rather than via ``warnings``."""
+    prev = os.environ.get("NVTE_CUTLASS_GROUPED_GEMM_WARN_FALLBACK")
+    os.environ["NVTE_CUTLASS_GROUPED_GEMM_WARN_FALLBACK"] = "1"
+    saved = os.dup(2)
+    try:
+        with tempfile.TemporaryFile(mode="w+b") as tmp:
+            sys.stderr.flush()
+            os.dup2(tmp.fileno(), 2)
+            try:
+                fn()
+            finally:
+                sys.stderr.flush()
+                os.dup2(saved, 2)
+            tmp.seek(0)
+            out = tmp.read().decode(errors="ignore")
+        return out.count("Fallback to cuBLAS grouped GEMM")
+    finally:
+        os.close(saved)
+        if prev is None:
+            os.environ.pop("NVTE_CUTLASS_GROUPED_GEMM_WARN_FALLBACK", None)
+        else:
+            os.environ["NVTE_CUTLASS_GROUPED_GEMM_WARN_FALLBACK"] = prev
+
 def generate_grouped_gemm_group_lens(b, m, balance: bool):
     if balance:
         return torch.full((b,), m, dtype=torch.int64)
@@ -82,10 +240,8 @@ def _generate_moe_test_cases(
 
 
 def generate_deepseekv3_test_cases():
-    # DSV3-GateUP hangs on some hardware; only benchmark DSV3-Down.
     return _generate_moe_test_cases(
         "DSV3", n_routed_experts=256, moe_intermediate_size=2048, hidden_size=7168,
-        skip_shapes=["GateUP"],
     )
 
 
@@ -107,7 +263,7 @@ def generate_grok_v2_test_cases():
     )
 
 
-def bench_grouped_gemm(Case, B, M, N, K, dtype, recipe):
+def bench_grouped_gemm(Case, B, M, N, K, dtype, recipe, Direction):
     device = "cuda"
 
     fp8_recipe = RECIPES[recipe]
@@ -119,68 +275,109 @@ def bench_grouped_gemm(Case, B, M, N, K, dtype, recipe):
     sum_M = sum(m_splits)
 
     grouped_linear = te.GroupedLinear(
-        B,
-        K,
-        N,
-        bias=False,
-        params_dtype=dtype,
-        device=device,
+        B, K, N, bias=False, params_dtype=dtype, device=device,
     )
-    # Rotate the activation buffer (on by default) so back-to-back grouped GEMMs 
+    # Rotate the activation buffer (on by default) so back-to-back grouped GEMMs
     # read different memory; GroupedLinear splits it internally per m_splits.
     next_x = make_input((sum_M, K), dtype, device=device, requires_grad=True)
 
-    def fwd_func_te():
+    def fwd_func():
+        xb = next_x()
         with te.autocast(enabled=use_fp8, recipe=fp8_recipe):
-            return grouped_linear(next_x(), m_splits, m_splits_tensor=m_splits_tensor)
+            return grouped_linear(xb, m_splits, m_splits_tensor=m_splits_tensor), xb
 
-    out_te = fwd_func_te()
-    grad_out = torch.randn_like(out_te)
+    grad_out = torch.randn_like(fwd_func()[0])
 
-    def fwd_bwd_func_te():
-        xb = next_x()
-        with te.autocast(enabled=use_fp8, recipe=fp8_recipe):
-            out = grouped_linear(xb, m_splits, m_splits_tensor=m_splits_tensor)
-            out.backward(grad_out)
+    def bwd_func(fwd_out):
+        out, xb = fwd_out
+        out.backward(grad_out)
         xb.grad = None
         for param in grouped_linear.parameters():
             param.grad = None
 
-    fwd_bwd_func_te()
+    if os.environ.get(_CUTLASS) == "1":
+        # Only fallbacks beyond the forward pass count for the backward direction.
+        fwd_fallbacks = _grouped_fallback_count(fwd_func)
+        fell_back = (fwd_fallbacks > 0 if Direction == "fwd"
+                     else _grouped_fallback_count(lambda: bwd_func(fwd_func())) > fwd_fallbacks)
+        if fell_back:
+            name = "HipKittens" if os.environ.get(_HK) == "1" else "CK"
+            pytest.skip(f"{name} grouped GEMM fell back to per-expert GEMMs for this config")
 
     fwd_total_flops = 2 * sum_M * N * K
-    bwd_total_flops = 2 * fwd_total_flops
-
-    fwd_te_ms, fwd_measurement = time_func(fwd_func_te)
-    fwd_bwd_te_ms, fwd_bwd_measurement = time_func(fwd_bwd_func_te)
-    bwd_te_ms = fwd_bwd_te_ms - fwd_te_ms
-
-    fwd_te_tflops = compute_tflops(fwd_total_flops, fwd_te_ms)
-    bwd_te_tflops = compute_tflops(bwd_total_flops, bwd_te_ms)
-
-    return make_forward_backward_metric_records(
-        BENCHMARK_LABEL,
-        "TFLOPS",
-        fwd_te_ms,
-        fwd_te_tflops,
-        bwd_te_ms,
-        bwd_te_tflops,
-        backward_derived=True,
-        fwd_measurement=fwd_measurement,
-        fwd_bwd_measurement=fwd_bwd_measurement,
+    # The per-expert path overlaps the experts on several streams, which a per-kernel sum
+    # double-counts, so measure elapsed device time by makespan.
+    return direction_records(
+        Direction, BENCHMARK_LABEL, "TFLOPS", compute_tflops,
+        fwd_func, bwd_func, fwd_total_flops, 2 * fwd_total_flops,
+        kernel_method="event",
     )
 
 
-if __name__ == "__main__":
-    test_cases = (
+def generate_cases():
+    """MoE grouped-GEMM cases crossed with per-precision backend and direction."""
+    base = (
         generate_deepseekv2_lite_test_cases()
         + generate_deepseekv2_test_cases()
         + generate_deepseekv3_test_cases()
         + generate_grok_v2_test_cases()
     )
-
-    run_benchmarks(
-        test_cases=test_cases,
-        bench_fn=bench_grouped_gemm,
-        param_columns=["Case", "B", "M", "N", "K", "dtype", "recipe"],
+    cases = []
+    for b in base:
+        for backend in _backends_for(b["recipe"]):
+            for direction in ("fwd", "bwd"):
+                cases.append({**b, "Backend": backend, "Direction": direction})
+    return cases
+
+
+def _case_id(c):
+    return f"{c['Case']}-{c['recipe']}-{c['Backend']}-{c['Direction']}-B{c['B']}-M{c['M']}"
+
+
+def pytest_generate_tests(metafunc):
+    if "case" in metafunc.fixturenames:
+        cases = generate_cases()
+        metafunc.parametrize("case", cases, ids=[_case_id(c) for c in cases])
+
+
+@pytest.mark.benchmark
+def test_grouped_gemm(request, microbench, case, monkeypatch):
+    backend = case["Backend"]
+    sig = (case["Case"], case["recipe"], backend, case["Direction"], case["B"], case["M"])
+    if get_device_compute_capability() == (9, 5) and sig in _GFX950_XFAIL_CONFIGS:
+        request.node.add_marker(pytest.mark.xfail(
+            reason="nvfp4 grouped FP4 dequant workspace too small on gfx950",
+            raises=RuntimeError, strict=False))
+    if get_device_compute_capability() == (9, 5) and sig in _GFX950_HANG_CONFIGS:
+        pytest.skip("known hipBLASLt grouped-GEMM hang on gfx950")
+    if backend in ("ck_tile", "hipkittens_grouped") and case["B"] <= 1:
+        pytest.skip(f"{backend} grouped GEMM needs num_groups > 1")
+    # Dense HipKittens per expert silently falls back to hipBLASLt for unaligned GEMMs.
+    if backend.startswith("hipkittens") and (case["N"] % 256 or case["K"] % 256):
+        pytest.skip("HipKittens GEMM needs 256-aligned expert dims")
+    # HK-vs-CK grouped MXFP8 is latched per process (rocm_gemm.cu); arch gates keep them apart.
+    if backend.startswith("hipkittens") and get_device_compute_capability() != (9, 5):
+        pytest.skip("HipKittens GEMM is gfx950-only")
+    if backend == "triton" and get_device_compute_capability() == (12, 5):
+        pytest.skip("Triton grouped GEMM hangs or crashes on many configs on gfx1250")
+    if backend == "hipkittens_per_expert" and not te_honors_env(_HIPBLASLT_MXFP8):
+        pytest.skip("HipKittens GEMM backend not available in this TE build")
+    if backend == "ck_tile" and case["recipe"] == "mxfp8" and get_device_compute_capability() != (12, 5):
+        pytest.skip("CK MXFP8 grouped GEMM is gfx1250-only")
+    # Skip a forced backend when the build doesn't honor the toggles it enables,
+    # so old builds show no data instead of silently measuring hipBLASLt. Builds that
+    # don't read NVTE_ROCM_USE_HIPBLASLT_MXFP8 already run hipBLASLt per expert.
+    required = [k for k, v in GROUPED_BACKENDS[backend].items()
+                if v is not None and k != _HIPBLASLT_MXFP8]
+    if required and not all(te_honors_env(k) for k in required):
+        pytest.skip(f"{backend} grouped GEMM backend not available in this TE build")
+    # The dense-GEMM selector also drives autocast's recipe check (TRITON rejects fp8 HYBRID).
+    monkeypatch.delenv("NVTE_GEMM_BACKEND", raising=False)
+    apply_backend_env(monkeypatch, GROUPED_BACKENDS[backend])
+    microbench.run(
+        case,
+        lambda: bench_grouped_gemm(
+            case["Case"], case["B"], case["M"], case["N"], case["K"],
+            case["dtype"], case["recipe"], case["Direction"],
+        ),
     )
diff --git a/benchmarks/microbenchmarks/benchmark_normalization.py b/benchmarks/microbenchmarks/benchmark_normalization.py
index ff9f7b28b4..31d23b1cf0 100755
--- a/benchmarks/microbenchmarks/benchmark_normalization.py
+++ b/benchmarks/microbenchmarks/benchmark_normalization.py
@@ -7,8 +7,11 @@
 """
 Normalization micro-benchmark using the fusible-ops LayerNorm / RMSNorm.
 
-Sweeps BF16 plus the quantized-output precisions (FP8, MXFP8) that TE training
-recipes produce. In FP8/FP4 training the norm is fused with the following
+Run with ``python benchmark_normalization.py`` (a pytest module under the hood;
+see conftest.py).
+
+Sweeps BF16 plus the quantized-output precisions (FP8, MXFP8, MXFP4, NVFP4)
+that TE training recipes produce. In FP8/FP4 training the norm is fused with the following
 Linear (LayerNormLinear / LayerNormMLP) and writes its output already
 quantized. That is reproduced here with ``ops.Sequential(Norm, Quantize)``: the
 op fuser threads the Quantize op's input quantizer into the norm, so the norm
@@ -16,22 +19,36 @@
 identity outside ``autocast``, so the bf16 baseline uses the same harness.
 
 Forward only: the quantize epilogue is a forward-pass phenomenon; the norm
-backward reads/writes high precision and is precision-independent. NVFP4/MXFP4
-norm outputs are not swept (no validated norm->FP4 cast path). Precisions
+backward reads/writes high precision and is precision-independent. On ROCm,
+MXFP8/MXFP4/NVFP4 norms run unfused (bf16 norm + separate cast kernel), so
+those rows measure norm + cast as LayerNormLinear sees it. Precisions
 unsupported on the current device are skipped.
 
 These are memory-bound; we report GB/s (input read + output write).
 Output: benchmark_normalization.csv (written to cwd)
 """
 
+if __name__ == "__main__":
+    # Enter pytest before torch touches the GPU, so conftest's neighbor snapshot excludes us.
+    import sys
+
+    import pytest
+
+    raise SystemExit(pytest.main([__file__, *sys.argv[1:]]))
+
+import functools
+import sys
+from pathlib import Path
+
+import pytest
 import torch
 import transformer_engine.pytorch as te
 from transformer_engine.pytorch import ops
 from utils import (
     MODEL_HIDDEN_SIZES, M_SIZE_LIST,
     build_recipes,
-    time_func, compute_gbps, make_metric_record, run_benchmarks,
-    make_input,
+    apply_backend_env, time_func_dual, compute_gbps, make_metric_record,
+    make_input, te_honors_env,
 )
 
 NORM_TYPES = [
@@ -41,10 +58,8 @@
 
 BENCHMARK_LABEL = "Normalization Forward"
 
-# Quantized-output precisions validated for the norm cast; tests/pytorch/
-# triton_kernels/test_norms.py exercises fp8 and mxfp8 norm quantizers. bf16 is
-# the plain baseline (Quantize is an identity outside autocast).
-RECIPES = build_recipes(names=("bf16", "fp8", "mxfp8"))
+# bf16 is the plain baseline (Quantize is an identity outside autocast).
+RECIPES = build_recipes()
 
 # Forward output bytes/elem by precision. Under autocast the norm quantizes
 # through the recipe-created quantizers, which default to columnwise usage on, so
@@ -54,35 +69,68 @@
 #   bf16  : 2.0                    (single bf16 output; Quantize is identity)
 #   fp8   : 2.0                    (rowwise + columnwise E4M3/E5M2 data; scale ~ 0)
 #   mxfp8 : 2.0 + 2/32 = 2 + 1/16  (rowwise + columnwise data + both E8M0 scales)
+#   mxfp4 : 1.0 + 2/32 = 1 + 1/16  (packed E2M1 data x2 + E8M0 scale per 32)
+#   nvfp4 : 1.0 + 2/16 = 1 + 1/8   (packed E2M1 data x2 + E4M3 scale per 16)
 _FWD_WRITE_BYTES = {
     "bf16": 2.0,
     "fp8": 2.0,
     "mxfp8": 2.0 + 1.0 / 16,
+    "mxfp4": 1.0 + 1.0 / 16,
+    "nvfp4": 1.0 + 1.0 / 8,
+}
+
+
+# Backend axis (None unsets, so "hip" is the native C++ path even if the ambient
+# env has a toggle set). "triton" flips the Triton RMSNorm/LayerNorm kernel (each
+# NormType reads only its own toggle).
+NORM_BACKENDS = {
+    "hip": {"NVTE_USE_RMSNORM_TRITON": None, "NVTE_USE_LAYERNORM_TRITON": None},
+    "triton": {"NVTE_USE_RMSNORM_TRITON": "1", "NVTE_USE_LAYERNORM_TRITON": "1"},
 }
 
+_NORM_CLS = {name: cls for name, cls in NORM_TYPES}
 
-def _generate_test_cases():
-    test_cases = []
+
+@functools.lru_cache(maxsize=1)
+def _triton_norm_handles_mxfp4():
+    """False on TE builds whose Triton norm predates MXFP4 output support."""
+    try:
+        src = (Path(te.__file__).parent / "triton_kernels/norms_common.py").read_text()
+    except OSError:
+        return False
+    return "MXFP4Quantizer" in src
+
+
+def generate_cases():
+    """Cross models x norm type x precision x backend x M (forward only)."""
+    cases = []
     for model_name, hidden in MODEL_HIDDEN_SIZES:
-        for norm_name, norm_op_cls in NORM_TYPES:
+        for norm_name in _NORM_CLS:
             for precision in RECIPES:
-                for M in M_SIZE_LIST:
-                    test_cases.append({
-                        "Case": f"{model_name}/{norm_name}",
-                        "Precision": precision,
-                        "M": M,
-                        "hidden_size": hidden,
-                        "norm_op_cls": norm_op_cls,
-                        "dtype": torch.bfloat16,
-                    })
-    return test_cases
-
-
-def bench_norm(Case, Precision, M, hidden_size, norm_op_cls, dtype):
+                for backend in NORM_BACKENDS:
+                    for M in M_SIZE_LIST:
+                        cases.append({
+                            "Case": model_name,
+                            "NormType": norm_name,
+                            "Precision": precision,
+                            "Backend": backend,
+                            "M": M,
+                            "hidden_size": hidden,
+                        })
+    return cases
+
+
+def _case_id(c):
+    return f"{c['Case']}-{c['NormType']}-{c['Precision']}-{c['Backend']}-M{c['M']}"
+
+
+def bench_norm(NormType, Precision, M, hidden_size):
     device = "cuda"
+    dtype = torch.bfloat16
 
     recipe = RECIPES[Precision]
     use_fp8 = recipe is not None
+    norm_op_cls = _NORM_CLS[NormType]
 
     # Norm followed by Quantize so the norm writes its output directly in the
     # target precision under autocast (identity when use_fp8 is False).
@@ -101,17 +149,39 @@ def fwd_func():
     # BF16 read + quantized write.
     fwd_bytes = int(M * hidden_size * (2 + _FWD_WRITE_BYTES[Precision]))
 
-    fwd_ms, fwd_measurement = time_func(fwd_func)
-    fwd_gbps = compute_gbps(fwd_bytes, fwd_ms)
-
+    fwd_ms, fwd_measurement, fwd_kernel_ms = time_func_dual(fwd_func)
     return [make_metric_record(
-        BENCHMARK_LABEL, fwd_ms, "GB/s", fwd_gbps, measurement=fwd_measurement,
+        BENCHMARK_LABEL, fwd_ms, "GB/s", compute_gbps(fwd_bytes, fwd_ms),
+        measurement=fwd_measurement,
+        kernel_ms=fwd_kernel_ms,
+        kernel_throughput=compute_gbps(fwd_bytes, fwd_kernel_ms) if fwd_kernel_ms else None,
     )]
 
 
-if __name__ == "__main__":
-    run_benchmarks(
-        test_cases=_generate_test_cases(),
-        bench_fn=bench_norm,
-        param_columns=["Case", "Precision", "M", "hidden_size", "dtype"],
+def pytest_generate_tests(metafunc):
+    if "case" in metafunc.fixturenames:
+        cases = generate_cases()
+        metafunc.parametrize("case", cases, ids=[_case_id(c) for c in cases])
+
+
+@pytest.mark.benchmark
+def test_norm(request, microbench, case, monkeypatch):
+    if case["Backend"] == "triton":
+        var = ("NVTE_USE_RMSNORM_TRITON" if case["NormType"] == "RMSNorm"
+               else "NVTE_USE_LAYERNORM_TRITON")
+        if not te_honors_env(var):
+            pytest.skip("Triton norm backend not available in this TE build")
+        if case["Precision"] == "mxfp4" and not _triton_norm_handles_mxfp4():
+            # Imperative xfail (not run): the GPU memory fault would abort the whole session.
+            pytest.xfail("Triton norm lacks MXFP4 support in this TE build (GPU memory fault)")
+    if case["Backend"] == "hip" and case["hidden_size"] == 16384:
+        # Llama3.1-405B: the hip norm at hidden 16384 is a multi-CTA cooperative launch, which
+        # makes HIP's exit-time teardown segfault (see pytest_unconfigure in conftest.py).
+        request.config._skip_hip_teardown_session = request.session
+    # The dense-GEMM selector also drives autocast's recipe check (TRITON rejects fp8 HYBRID).
+    monkeypatch.delenv("NVTE_GEMM_BACKEND", raising=False)
+    apply_backend_env(monkeypatch, NORM_BACKENDS[case["Backend"]])
+    microbench.run(
+        case,
+        lambda: bench_norm(case["NormType"], case["Precision"], case["M"], case["hidden_size"]),
     )
diff --git a/benchmarks/microbenchmarks/compare_results.py b/benchmarks/microbenchmarks/compare_results.py
index 4a7e1dab8c..a63c992a0d 100755
--- a/benchmarks/microbenchmarks/compare_results.py
+++ b/benchmarks/microbenchmarks/compare_results.py
@@ -21,7 +21,7 @@
 import numpy as np
 import pandas as pd
 
-SKIP_COLS = {"TestID", "Label"}
+SKIP_COLS = {"TestID", "Label", "run_ts", "commit_sha", "commit_date", "Kernels"}
 DEFAULT_MIN_BASELINE_METRIC = 0.5
 
 
diff --git a/benchmarks/microbenchmarks/conftest.py b/benchmarks/microbenchmarks/conftest.py
new file mode 100644
index 0000000000..41261b3972
--- /dev/null
+++ b/benchmarks/microbenchmarks/conftest.py
@@ -0,0 +1,254 @@
+# Copyright (c) 2026, Advanced Micro Devices, Inc. All rights reserved.
+# License for AMD contributions = MIT. See LICENSE for more information
+"""pytest glue for the microbenchmarks.
+
+Thin shim: the option/CSV/timing logic lives in utils.py; this file only wires
+those helpers into pytest hooks and exposes the ``microbench`` fixture.  Each
+benchmark is run canonically with ``python benchmark_gemm.py`` (the file
+forwards to pytest), e.g.::
+
+    python benchmark_gemm.py --csv                 # write benchmark_gemm.csv
+    python benchmark_gemm.py -k "bf16 and QKV"     # select by parametrize id
+    python benchmark_gemm.py -k triton             # select the triton backend
+"""
+
+import ctypes
+import os
+import sys
+
+# Snapshot GPU neighbors BEFORE importing torch: on ROCm the first CUDA call (which
+# `from utils import` can trigger via `import torch.utils.benchmark`) registers this
+# process on every visible GPU, which would otherwise look like a neighbor. Taken
+# here, the snapshot is free of our own PID.
+import gpu_neighbors
+
+_GPU_SNAPSHOT = gpu_neighbors.snapshot_gpu_neighbors()
+
+from pathlib import Path
+
+import pytest
+
+from utils import (
+    configure_kernel_profile,
+    configure_rotating,
+    dashboard_run_plan,
+    format_aggregate_table,
+    format_results_table,
+    pin_cpus,
+    print_case,
+    record_bench,
+    write_bench_outputs,
+    write_dashboard_run,
+)
+from transformer_engine.pytorch.quantization import FP8GlobalStateManager
+
+
+def pytest_addoption(parser):
+    group = parser.getgroup("microbench", "TE GPU microbenchmarks")
+    group.addoption(
+        "--csv", nargs="?", const=True, default=None, metavar="FILE",
+        help="Write results to CSV (one per family; default name from the module).",
+    )
+    group.addoption(
+        "--csv-samples", nargs="?", const=True, default=None, metavar="FILE",
+        help="Write per-sample timing data to a CSV.",
+    )
+    group.addoption(
+        "--kernel-profile", action="store_true", default=False,
+        help="Also measure GPU kernel (device) time alongside wall time, adding "
+             "Kernel Time / Kernel <unit> columns to the CSV.",
+    )
+    group.addoption(
+        "--rotating", nargs="?", type=int, const=0, default=None, metavar="MB",
+        help="Rotate inputs through a ring of buffers (optional MB budget). On by default.",
+    )
+    group.addoption(
+        "--no-rotating", action="store_true", default=False,
+        help="Disable input buffer rotation.",
+    )
+    group.addoption(
+        "--abort-on-gpu-interference", action="store_true", default=False,
+        help="Abort the run if another process is using the benchmark GPU(s).",
+    )
+    group.addoption(
+        "--no-gpu-interference-check", action="store_true", default=False,
+        help="Disable the shared-GPU interference check.",
+    )
+    group.addoption(
+        "--no-cpu-pin", action="store_true", default=False,
+        help="Don't pin the run to a few GPU-local cores at the performance governor "
+             "(on by default; the governor change needs root).",
+    )
+    group.addoption(
+        "--dashboard-run", action="store_true", default=False,
+        help="Collect all families into one run dir (born-tagged CSVs + samples/ + "
+             "run_info.txt) under --dashboard-out, for the TE dashboard.",
+    )
+    group.addoption(
+        "--dashboard-out", default="results", metavar="DIR",
+        help="Base output dir for --dashboard-run; a "
+             "<arch>_<host>_gpu<id>_<pci>/<date>_<commit>/ run dir is created under it "
+             "(default: results).",
+    )
+    group.addoption(
+        "--gpu-model", default=None, metavar="NAME",
+        help="Override the detected GPU model/arch name (e.g. MI400X) used in the "
+             "--dashboard-run dir + run_info; required when it can't be detected.",
+    )
+
+
+def pytest_configure(config):
+    config.addinivalue_line("markers", "benchmark: TE GPU microbenchmark")
+    configure_rotating(config.getoption("--rotating"), config.getoption("--no-rotating"))
+    configure_kernel_profile(config.getoption("--kernel-profile"))
+    config._microbench_store = {}
+    if not config.getoption("--no-gpu-interference-check"):
+        status, foreign = gpu_neighbors.detect_gpu_interference(_GPU_SNAPSHOT)
+        if foreign:
+            print(gpu_neighbors.format_gpu_interference(foreign))
+            if config.getoption("--abort-on-gpu-interference"):
+                pytest.exit(
+                    "GPU interference detected; aborting (--abort-on-gpu-interference).",
+                    returncode=3,
+                )
+        elif status == "unavailable":
+            print("WARNING: GPU interference check skipped -- could not list the processes on "
+                  "the benchmark GPU via amdsmi (pip install amdsmi if missing, or pass "
+                  "--no-gpu-interference-check to silence).")
+
+    # Resolve the dashboard run dir up front so pytest_report_header can show it
+    # and write_dashboard_run can reuse the same metadata at session end.
+    plan = None
+    if config.getoption("--dashboard-run"):
+        plan = dashboard_run_plan(
+            config.getoption("--dashboard-out"), model=config.getoption("--gpu-model"))
+        if not plan.meta["model"]:
+            raise pytest.UsageError(
+                "--dashboard-run: could not determine the GPU model from the device name; "
+                "pass --gpu-model NAME (e.g. MI355X).")
+
+    config._cpu_pin = None if config.getoption("--no-cpu-pin") else pin_cpus()
+
+    if plan is not None:
+        config._dashboard_plan = plan
+        if config._cpu_pin is not None:
+            plan.meta["cpus"] = config._cpu_pin.description
+
+
+def pytest_report_header(config):
+    pin = getattr(config, "_cpu_pin", None)
+    lines = [f"cpus: {pin.description if pin else 'unpinned'}",
+             f"GPU_MAX_HW_QUEUES: {os.environ.get('GPU_MAX_HW_QUEUES', 'HIP default')}"]
+    plan = getattr(config, "_dashboard_plan", None)
+    if plan is None:
+        return lines
+    extras = [name for name, opt in (
+        ("kernel-profile", "--kernel-profile"),
+        ("samples", "--csv-samples"),
+    ) if config.getoption(opt)]
+    m = plan.meta
+    return [
+        f"dashboard-run -> {plan.run_dir}",
+        f"  {m['model']} on {m['host']} gpu{m['gpu']} "
+        f"({m['gpu_bdf'] or 'no pci'}) @ {m['short']}; "
+        f"extras: {', '.join(extras) or 'none (wall only)'}",
+        *lines,
+    ]
+
+
+def pytest_collect_file(parent, file_path):
+    # Collect benchmark_*.py like test files so `pytest .` finds them without a
+    # rename; skip init paths so an explicitly-passed file isn't double-collected.
+    if (
+        file_path.suffix == ".py"
+        and file_path.name.startswith("benchmark_")
+        and not parent.session.isinitpath(file_path)
+    ):
+        return pytest.Module.from_parent(parent, path=file_path)
+    return None
+
+
+class _MicroBench:
+    """Handed to each test as the ``microbench`` fixture."""
+
+    def __init__(self, request):
+        self._request = request
+        self._config = request.config
+
+    def run(self, case, bench_callable):
+        """Time *bench_callable* (returns metric records) and record it under *case*."""
+        records = bench_callable()
+        print_case(case, records)
+        family = Path(self._request.module.__file__).stem
+        record_bench(
+            self._config._microbench_store, family, case, records,
+            self._request.node.name,
+        )
+        return records
+
+
+@pytest.fixture
+def microbench(request):
+    return _MicroBench(request)
+
+
+@pytest.fixture(autouse=True)
+def _reset_fp8_global_state():
+    yield
+    # Every autocast exit reduces all DelayedScaling amaxes ever registered, dead modules included.
+    FP8GlobalStateManager.reset()
+
+
+def pytest_terminal_summary(terminalreporter, exitstatus, config):
+    store = getattr(config, "_microbench_store", None)
+    if not store:
+        return
+    table = format_results_table(store)
+    if table:
+        terminalreporter.write_line("")
+        for row in table.splitlines():
+            terminalreporter.write_line(row)
+    aggregate = format_aggregate_table(store)
+    if aggregate:
+        terminalreporter.write_line("")
+        for row in aggregate.splitlines():
+            terminalreporter.write_line(row)
+
+
+def pytest_sessionfinish(session, exitstatus):
+    config = session.config
+    if getattr(config, "_cpu_pin", None) is not None:
+        config._cpu_pin.restore()
+    store = getattr(config, "_microbench_store", None)
+    if not store:
+        return
+    if config.getoption("--dashboard-run"):
+        plan = getattr(config, "_dashboard_plan", None)
+        run_dir = write_dashboard_run(
+            store,
+            out_base=config.getoption("--dashboard-out"),
+            csv_samples=config.getoption("--csv-samples"),
+            meta=plan.meta if plan else None,
+        )
+        print(f"microbench: dashboard run -> {run_dir}")
+        return
+    written = write_bench_outputs(
+        store,
+        csv=config.getoption("--csv"),
+        csv_samples=config.getoption("--csv-samples"),
+    )
+    for path in written:
+        print(f"microbench: wrote {path}")
+
+
+@pytest.hookimpl(trylast=True)
+def pytest_unconfigure(config):
+    # Armed by benchmarks that trigger ROCm 7.14's segfault in HIP's exit-time hsa_shut_down (as in
+    # tests/pytorch/triton_kernels/conftest.py). Not under rocprofv3, which writes its trace there.
+    session = getattr(config, "_skip_hip_teardown_session", None)
+    if session is None or "ROCP_TOOL_LIBRARIES" in os.environ:
+        return
+    sys.stdout.flush()
+    sys.stderr.flush()
+    ctypes.CDLL(None).fflush(None)
+    os._exit(int(session.exitstatus))
diff --git a/benchmarks/microbenchmarks/gpu_neighbors.py b/benchmarks/microbenchmarks/gpu_neighbors.py
new file mode 100644
index 0000000000..8950a96aeb
--- /dev/null
+++ b/benchmarks/microbenchmarks/gpu_neighbors.py
@@ -0,0 +1,198 @@
+# Copyright (c) 2026, Advanced Micro Devices, Inc. All rights reserved.
+# License for AMD contributions = MIT. See LICENSE for more information
+"""GPU interference detection: warn when another process shares the
+benchmark GPU, so timings aren't silently skewed by a neighbor.
+
+This module is deliberately free of a top-level ``torch`` import. On ROCm the
+first CUDA call in a process (e.g. importing ``torch.utils.benchmark`` or
+``torch.cuda.get_device_properties``) opens every visible device and registers
+this process on all of them under its *host* PID -- which, inside a container,
+can't be matched to our container PID. So the neighbor snapshot must be taken
+*before* torch touches the GPU; conftest calls :func:`snapshot_gpu_neighbors`
+at import time, ahead of ``from utils import ...`` (which pulls torch). With no
+CUDA context yet we are absent from the process list, so every entry is a
+genuine neighbor and there is no self to exclude across the PID namespace.
+:func:`detect_gpu_interference` then maps torch's ``cuda:0`` to its amdsmi
+handle by PCI address and reports the pre-existing processes on that GPU.
+"""
+
+import os
+
+
+def _proc_cmdline(pid):
+    """Full argv for *pid* from /proc, or None if unreadable (e.g. another pid namespace)."""
+    try:
+        with open(f"/proc/{pid}/cmdline", "rb") as fh:
+            raw = fh.read()
+    except OSError:
+        return None
+    parts = [p.decode("utf-8", "replace") for p in raw.split(b"\x00") if p]
+    if parts:
+        return " ".join(parts)
+    try:  # kernel thread / hidden argv -> fall back to comm
+        with open(f"/proc/{pid}/comm") as fh:
+            return fh.read().strip()
+    except OSError:
+        return None
+
+
+# Monitoring tools open every GPU but aren't real compute interference. Matched on the
+# executable's basename after stripping non-alphanumerics, so e.g. "/opt/rocm/bin/rocm-smi"
+# -> "rocmsmi".
+_IGNORED_PROC_NAMES = ("nvtop", "amdsmi", "rocmsmi")
+
+
+def _is_ignored_proc(name):
+    exe = os.path.basename(str(name or "").strip())
+    base = "".join(ch for ch in exe.lower() if ch.isalnum())
+    return bool(base) and any(base.startswith(n) for n in _IGNORED_PROC_NAMES)
+
+
+def _amdsmi_attr(obj, key):
+    if isinstance(obj, dict):
+        return obj.get(key)
+    return getattr(obj, key, None)
+
+
+def _amdsmi_gpu_processes(amdsmi, handles):
+    """{handle_index: {pid: name}} of compute processes across *handles*.
+
+    A GPU whose process list can't be queried maps to None (unknown), not to {}.
+    """
+    out = {}
+    for idx, handle in enumerate(handles):
+        procs = {}
+        try:
+            plist = amdsmi.amdsmi_get_gpu_process_list(handle)
+        except Exception:
+            out[idx] = None
+            continue
+        for p in plist:
+            pid = _amdsmi_attr(p, "pid")
+            name = _amdsmi_attr(p, "name") or _amdsmi_attr(p, "process_name")
+            if pid is None:  # older amdsmi returns opaque process handles -> fetch details
+                try:
+                    info = amdsmi.amdsmi_get_gpu_process_info(handle, p)
+                    pid = _amdsmi_attr(info, "pid")
+                    name = name or _amdsmi_attr(info, "name")
+                except Exception:
+                    pass
+            try:
+                pid = int(pid)
+            except (TypeError, ValueError):
+                continue
+            procs[pid] = str(name) if name else ""
+        out[idx] = procs
+    return out
+
+
+def _amdsmi_bdf(amdsmi, handle):
+    """amdsmi BDF for *handle* normalized to ``'domain:bus:device'`` (function dropped), or None."""
+    try:
+        bdf = str(amdsmi.amdsmi_get_gpu_device_bdf(handle)).strip().lower()
+    except Exception:
+        return None
+    return bdf.split(".")[0] or None
+
+
+def snapshot_gpu_neighbors():
+    """Snapshot GPU processes + PCI addresses via amdsmi, before any CUDA init.
+
+    Must be called before this process makes its first CUDA call, so the process
+    list is free of our own (soon-to-be-registered) host PID. Returns a dict with
+    ``status`` (``"ok"`` / ``"unavailable"``), ``neighbors`` (``{idx: {pid:
+    name}}``) and ``bdfs`` (``{idx: "domain:bus:device"}``). amdsmi is shut down
+    before returning, so nothing GPU-related is held for the rest of the session.
+    """
+    try:
+        import amdsmi
+    except Exception:
+        return {"status": "unavailable", "neighbors": {}, "bdfs": {}}
+    try:
+        amdsmi.amdsmi_init()
+    except Exception:
+        return {"status": "unavailable", "neighbors": {}, "bdfs": {}}
+    try:
+        handles = amdsmi.amdsmi_get_processor_handles()
+        neighbors = _amdsmi_gpu_processes(amdsmi, handles)
+        bdfs = {idx: _amdsmi_bdf(amdsmi, h) for idx, h in enumerate(handles)}
+        return {"status": "ok", "neighbors": neighbors, "bdfs": bdfs}
+    except Exception:
+        return {"status": "unavailable", "neighbors": {}, "bdfs": {}}
+    finally:
+        try:
+            amdsmi.amdsmi_shut_down()
+        except Exception:
+            pass
+
+
+def _torch_cuda_bdf(index=0):
+    """PCI address of torch's ``cuda:index`` as ``'domain:bus:device'`` (lowercase hex), or None.
+
+    Reading device properties lazily initializes CUDA, so call this only *after*
+    the neighbor snapshot. It respects HIP_VISIBLE_DEVICES, so cuda:0 is the GPU
+    the benchmark actually uses.
+    """
+    try:
+        import torch
+        p = torch.cuda.get_device_properties(index)
+        return f"{p.pci_domain_id:04x}:{p.pci_bus_id:02x}:{p.pci_device_id:02x}"
+    except Exception:
+        return None
+
+
+def detect_gpu_interference(snapshot):
+    """Find foreign compute processes on *our* GPU.
+
+    Uses the pre-CUDA *snapshot* from :func:`snapshot_gpu_neighbors` (genuine
+    neighbors only, no self), then maps torch's ``cuda:0`` to its amdsmi handle
+    by matching PCI addresses (immune to the amdsmi<->HIP index mismatch), and 
+    reports the pre-existing processes on that GPU, filtering out monitoring tools 
+    (nvtop/amd-smi/rocm-smi). If we can't locate our GPU we report ``"unavailable"``
+    rather than risk a false positive.
+    Returns ``(status, foreign)`` with *status* in ``"ok"`` / ``"unavailable"`` /
+    ``"not_amd"`` and *foreign* a list of ``(gpu, pid, cmdline)``.
+    """
+    import torch
+
+    if not getattr(torch.version, "hip", None):
+        return "not_amd", []
+    if not snapshot or snapshot.get("status") != "ok":
+        return "unavailable", []
+
+    neighbors = snapshot.get("neighbors", {})
+    bdfs = snapshot.get("bdfs", {})
+    target = _torch_cuda_bdf(0)
+    our_idx = None
+    if target is not None:
+        for idx, bdf in bdfs.items():
+            if bdf == target:
+                our_idx = idx
+                break
+
+    if our_idx is None:
+        return "unavailable", []
+    procs = neighbors.get(our_idx, {})
+    if procs is None:  # amdsmi couldn't list our GPU's processes
+        return "unavailable", []
+    foreign = []
+    for pid, name in sorted(procs.items()):
+        if _is_ignored_proc(name):
+            continue
+        foreign.append((our_idx, pid, _proc_cmdline(pid) or name or f"pid {pid}"))
+    return "ok", foreign
+
+
+def format_gpu_interference(foreign):
+    """Render the foreign-process list as a warning banner."""
+    bar = "=" * 74
+    gpu = foreign[0][0] if foreign else "?"
+    lines = ["", bar,
+             f"GPU INTERFERENCE: {len(foreign)} other process(es) on benchmark GPU {gpu}",
+             "-" * 74]
+    for _gpu, pid, cmd in foreign:
+        if len(cmd) > 150:
+            cmd = cmd[:147] + "..."
+        lines.append(f"  PID {pid:<8}  {cmd}")
+    lines.append(bar)
+    return "\n".join(lines)
diff --git a/benchmarks/microbenchmarks/run_benchmarks.sh b/benchmarks/microbenchmarks/run_benchmarks.sh
deleted file mode 100755
index 5a185177b1..0000000000
--- a/benchmarks/microbenchmarks/run_benchmarks.sh
+++ /dev/null
@@ -1,38 +0,0 @@
-#!/bin/bash
-###############################################################################
-# Copyright (c) 2026, Advanced Micro Devices, Inc. All rights reserved.
-#
-# See LICENSE for license information.
-###############################################################################
-# Run every microbenchmark.
-#
-#   ./run_all.sh                  # defaults (rotation on)
-#   ./run_all.sh --rotating 512   # rotate within a 512 MB budget
-#   ./run_all.sh --no-rotating    # disable input rotation
-#   ./run_all.sh --csv            # also write per-benchmark CSVs
-#
-# Set PYTHON to pick a specific interpreter (default: python).
-
-shopt -s nullglob
-
-cd "$(dirname "$0")"
-PYTHON="${PYTHON:-python}"
-
-failed=()
-for bench in benchmark_*.py; do
-    echo
-    echo "############################################################"
-    echo "# ${bench} $*"
-    echo "############################################################"
-    if ! "$PYTHON" "$bench" "$@"; then
-        echo "!!! ${bench} FAILED" >&2
-        failed+=("$bench")
-    fi
-done
-
-echo
-if (( ${#failed[@]} )); then
-    echo "FAILED: ${failed[*]}" >&2
-    exit 1
-fi
-echo "All benchmarks completed."
diff --git a/benchmarks/microbenchmarks/utils.py b/benchmarks/microbenchmarks/utils.py
index 45101b60f7..8136bd07e3 100644
--- a/benchmarks/microbenchmarks/utils.py
+++ b/benchmarks/microbenchmarks/utils.py
@@ -4,12 +4,34 @@
 #
 # See LICENSE for license information.
 ###############################################################################
-"""Shared utilities for microbenchmarks: model configs, timing, throughput, runner."""
+"""Shared utilities for microbenchmarks: model configs, timing, throughput, CSV output."""
 
-import argparse
-import importlib.util
+import csv
+import functools
+import gc
 import itertools
 import math
+import mmap
+import os
+import re
+import statistics
+import subprocess
+import sys
+import threading
+import time
+import warnings
+import zlib
+from pathlib import Path
+from types import SimpleNamespace
+
+# MXFP8 is gated off by default on ROCm (NVTE_ROCM_ENABLE_MXFP8=0). Enable it here, before any
+# build_recipes() call caches check_mxfp8_support(): a directly run `python benchmark_*.py` builds
+# its RECIPES at import time, before pytest loads conftest.py.
+os.environ.setdefault("NVTE_ROCM_ENABLE_MXFP8", "1")
+# Megatron's launch scripts default to 2 HW queues. With HIP's default 4, TE's high-priority compute
+# streams get a 5th queue whose pending event waits slow the main stream (per-expert grouped GEMM ~2x).
+os.environ.setdefault("GPU_MAX_HW_QUEUES", "2")
+
 import torch
 import torch.utils.benchmark as benchmark
 
@@ -47,9 +69,9 @@
 
 # Unique (model_name, hidden_size) pairs for element-wise benchmarks
 MODEL_HIDDEN_SIZES = [
-    ("Llama3-8B",   4096),
-    ("Llama3-70B",  8192),
-    ("Llama3-405B", 16384),
+    ("Llama3.1-8B",   4096),
+    ("Llama3.1-70B",  8192),
+    ("Llama3.1-405B", 16384),
     ("Qwen2.5-7B",  3584),
     ("Qwen2.5-72B", 8192),
 ]
@@ -95,22 +117,6 @@ def generate_gemm_test_cases(configs=None, m_sizes=None, dtypes=None):
 # offline tooling importable).
 
 
-def _check_mxfp4_support_with_aiter():
-    """MXFP4 gate: device support plus the aiter FP4 GEMM backend.
-
-    The MXFP4 GEMM path calls into aiter's a4w4 kernels, so a missing aiter
-    package would crash at benchmark time even on supported hardware.
-    """
-    from transformer_engine.pytorch.quantization import check_mxfp4_support
-
-    supported, reason = check_mxfp4_support()
-    if not supported:
-        return supported, reason
-    if importlib.util.find_spec("aiter") is None:
-        return False, "aiter is not installed (required for the MXFP4 GEMM backend)."
-    return True, ""
-
-
 def _precision_specs():
     """Ordered sweep of (name, recipe factory | None, support check | None).
 
@@ -127,6 +133,7 @@ def _precision_specs():
     )
     from transformer_engine.pytorch.quantization import (
         check_fp8_support,
+        check_mxfp4_support,
         check_mxfp8_support,
         check_nvfp4_support,
     )
@@ -143,7 +150,7 @@ def _precision_specs():
             check_fp8_support,
         ),
         ("mxfp8", MXFP8BlockScaling, check_mxfp8_support),
-        ("mxfp4", MXFP4BlockScaling, _check_mxfp4_support_with_aiter),
+        ("mxfp4", MXFP4BlockScaling, check_mxfp4_support),
         ("nvfp4", NVFP4BlockScaling, check_nvfp4_support),
     )
 
@@ -177,7 +184,7 @@ def build_recipes(names=None):
 # ---------------------------------------------------------------------------
 
 def time_func(fn, method="adaptive", min_run_time=DEFAULT_MIN_RUN_TIME_SECONDS):
-    """Time *fn* and return ``(mean_ms, measurement)``.
+    """Time *fn* and return ``(median_ms, measurement)``.
 
     The ``Measurement`` object carries per-sample times accessible via
     ``measurement.times`` (total wall time per run) and
@@ -186,12 +193,286 @@ def time_func(fn, method="adaptive", min_run_time=DEFAULT_MIN_RUN_TIME_SECONDS):
     method: "adaptive" uses adaptive_autorange (good for compute-bound),
             "blocked"  uses blocked_autorange  (good for memory-bound).
     """
+    # Untimed first (pre-warmup) call: the timers size their blocks from their first timed calls,
+    # and a slow one-off first call (e.g. lazy kernel load, JIT compilation, algorithm selection,
+    # autotuning) can make them pick a too-small block, so each sample includes a sync and
+    # measures latency instead of throughput.
+    fn()
+    torch.cuda.synchronize()
     timer = benchmark.Timer(stmt="fn()", globals={"fn": fn})
     if method == "blocked":
         m = timer.blocked_autorange(min_run_time=min_run_time)
     else:
         m = timer.adaptive_autorange(min_run_time=min_run_time)
-    return m.mean * 1e3, m
+    return m.median * 1e3, m
+
+
+# Dual wall + GPU-kernel timing. Off by default; enabled per run by
+# --kernel-profile (set via configure_kernel_profile from conftest).
+_KERNEL_PROFILE = False
+
+
+def configure_kernel_profile(enabled):
+    """Enable/disable dual wall+kernel timing (set from --kernel-profile)."""
+    global _KERNEL_PROFILE
+    _KERNEL_PROFILE = bool(enabled)
+
+
+def _prepared_calls(fn, prepare, n):
+    """*n* zero-arg calls of *fn*, each bound to its own ``prepare()`` result when given."""
+    if prepare is None:
+        return [fn] * n
+    return [functools.partial(fn, prepare()) for _ in range(n)]
+
+
+def _kernel_time_profiler_ms(fn, warmup=100, iters=100, prepare=None, block=None):
+    """Median GPU kernel (device) time per call, in ms, via torch.profiler.
+
+    Profiles ``iters`` back-to-back calls and splits the in-order kernel list into
+    per-call groups (matched against one call's kernel sequence), summing the device
+    time of each call's kernels; this excludes host launch overhead and host-side
+    timing noise. Accurate for single-stream ops; unreliable for concurrent
+    multi-stream ops (use the "event" method there). Returns ``None`` (with a
+    warning) when no call's kernels fully match. With *prepare*, each call is
+    ``fn(prepare())`` with ``prepare`` run before the profiled region, in sessions
+    of at most *block* calls.
+    """
+    from torch.autograd import DeviceType
+    from torch.profiler import profile, ProfilerActivity
+
+    def profiled_kernels(n):
+        # The previous session's pause lets the GPU clock down; bring it back up first.
+        start = time.perf_counter()
+        while time.perf_counter() - start < 0.02:
+            _prepared_calls(fn, prepare, 1)[0]()
+            torch.cuda.synchronize()
+        calls = _prepared_calls(fn, prepare, n)
+        torch.cuda.synchronize()
+        with profile(activities=[ProfilerActivity.CUDA]) as prof:
+            for call in calls:
+                call()
+            torch.cuda.synchronize()
+            # ROCm's profiler drops kernel records still being processed when the session stops.
+            time.sleep(0.05)
+        return sorted((e for e in prof.events() if e.device_type != DeviceType.CPU),
+                      key=lambda e: e.time_range.start)
+
+    for _ in range(warmup):
+        _prepared_calls(fn, prepare, 1)[0]()
+    torch.cuda.synchronize()
+    pattern = max(([e.name for e in profiled_kernels(1)] for _ in range(3)), key=len)
+    if not pattern:
+        warnings.warn("No kernel records captured; kernel time left empty.")
+        return None
+    per_call_us, n_kernels = [], 0
+    for start in range(0, iters, block or iters):
+        kernels = profiled_kernels(min(block or iters, iters - start))
+        n_kernels += len(kernels)
+        # ROCm's profiler can still drop the odd kernel record, so only fully matched calls count.
+        group = []
+        for e in kernels:
+            if e.name != pattern[len(group)]:
+                group = []
+                if e.name != pattern[0]:
+                    continue
+            group.append(e)
+            if len(group) == len(pattern):
+                per_call_us.append(sum(k.device_time_total for k in group))
+                group = []
+    if not per_call_us:
+        warnings.warn(f"No profiled call matched the {len(pattern)}-kernel pattern "
+                      f"({n_kernels} kernel records from {iters} calls); "
+                      "kernel time left empty.")
+        return None
+    return statistics.median(per_call_us) / 1e3
+
+
+def _kernel_time_event_ms(fn, warmup=100, iters=100, prepare=None, block=None):
+    """Median elapsed GPU device time per call, in ms, via CUDA events.
+
+    Records an event on the current stream between consecutive calls of a
+    saturated loop, so concurrent multi-stream kernels are measured by their
+    overlapped span rather than a per-kernel sum. Handles the multi-stream
+    grouped-GEMM path (where a per-kernel sum double-counts overlapping kernels);
+    converges to the wall time for device-bound ops. *prepare* / *block* are as in
+    :func:`_kernel_time_profiler_ms`.
+    """
+    for _ in range(warmup):
+        _prepared_calls(fn, prepare, 1)[0]()
+    torch.cuda.synchronize()
+    spans = []
+    for start in range(0, iters, block or iters):
+        calls = _prepared_calls(fn, prepare, min(block or iters, iters - start))
+        torch.cuda.synchronize()
+        events = [torch.cuda.Event(enable_timing=True) for _ in range(len(calls) + 1)]
+        events[0].record()
+        for call, event in zip(calls, events[1:]):
+            call()
+            event.record()
+        torch.cuda.synchronize()
+        spans += [a.elapsed_time(b) for a, b in zip(events, events[1:])]
+    return statistics.median(spans)
+
+
+# Kernel-time method: "profiler" (default) or "event". A benchmark opts into the
+# event makespan for concurrent multi-stream ops (grouped GEMM); setting
+# NVTE_MICROBENCH_KERNEL_TIME overrides the method globally for A/B testing.
+_KERNEL_TIME_METHOD_ENV = "NVTE_MICROBENCH_KERNEL_TIME"
+
+
+def _resolve_kernel_method(method):
+    env = os.environ.get(_KERNEL_TIME_METHOD_ENV, "").strip().lower()
+    if env in ("profiler", "event"):
+        return env
+    return method or "profiler"
+
+
+def _kernel_time_ms(fn, warmup=100, iters=100, method=None, prepare=None, block=None):
+    """Median GPU device time per call (ms) using the resolved kernel-time method."""
+    if _resolve_kernel_method(method) == "event":
+        return _kernel_time_event_ms(fn, warmup, iters, prepare, block)
+    return _kernel_time_profiler_ms(fn, warmup, iters, prepare, block)
+
+
+# Tensile kernel names spell out every tuning parameter (~900 chars): keep them up to the
+# tile and matrix-instruction sizes, plus a checksum of the full name.
+_TENSILE_NAME_RE = re.compile(r"^(Cijk_.*?_MT\d+x\d+x\d+_MI\d+x\d+x\d+)_")
+
+
+def _short_kernel_name(name):
+    name = name.removesuffix(".kd")
+    m = _TENSILE_NAME_RE.match(name)
+    return f"{m.group(1)}~{zlib.crc32(name.encode()):08x}" if m else name
+
+
+def _triton_config_name(name, config):
+    """``name[<M>x<N>x<K> key=value ...]`` for a Triton autotuner config."""
+    kw = dict(config.kwargs)
+    tile = [f"BLOCK_SIZE_{d}" for d in "MNK"]
+    parts = ["x".join(str(kw.pop(k)) for k in tile)] if all(k in kw for k in tile) else []
+    kw.update(num_warps=config.num_warps, num_stages=config.num_stages)
+    return f"{name}[{' '.join(parts + [f'{k}={v}' for k, v in kw.items()])}]"
+
+
+def gemm_kernel_names(fn, prepare=None):
+    """``;``-joined GEMM kernels that one call of *fn* launches, in launch order: hipBLASLt
+    kernels (Tensile ``Cijk_*``, rocRoller ``RR_GEMM_*``) and Triton autotuned kernels with the
+    config they run. *prepare* is as in :func:`_kernel_time_profiler_ms`.
+    """
+    from torch.autograd import DeviceType
+    from torch.profiler import profile, ProfilerActivity
+
+    # A Triton kernel has one name for all its configs: record the autotuner's pick per launch.
+    autotuner = sys.modules.get("triton.runtime.autotuner")
+    run = autotuner.Autotuner.run if autotuner else None
+    best = (False, [])
+    for _ in range(3):  # ROCm's profiler can drop the odd kernel record
+        launches = []
+
+        def tracked_run(self, *args, **kwargs):
+            out = run(self, *args, **kwargs)
+            base, config = getattr(self, "base_fn", None), getattr(self, "best_config", None)
+            if base is not None and config is not None:
+                launches.append((base.__name__, _triton_config_name(base.__name__, config)))
+            return out
+
+        call = _prepared_calls(fn, prepare, 1)[0]
+        torch.cuda.synchronize()
+        if autotuner:
+            autotuner.Autotuner.run = tracked_run
+        try:
+            with profile(activities=[ProfilerActivity.CUDA]) as prof:
+                call()
+                torch.cuda.synchronize()
+                time.sleep(0.05)
+        finally:
+            if autotuner:
+                autotuner.Autotuner.run = run
+        triton_configs = {}
+        for name, config in launches:
+            triton_configs.setdefault(name, []).append(config)
+        names = []
+        for e in sorted((e for e in prof.events() if e.device_type != DeviceType.CPU),
+                        key=lambda e: e.time_range.start):
+            name = e.name.removesuffix(".kd")
+            if name.startswith(("Cijk_", "RR_GEMM_")):
+                names.append(_short_kernel_name(name))
+            elif triton_configs.get(name):
+                names.append(triton_configs[name].pop(0))
+        # Prefer a pass where every Triton launch matched its kernel record.
+        best = max(best, (not any(triton_configs.values()), names),
+                   key=lambda r: (r[0], len(r[1])))
+    return ";".join(best[1])
+
+
+def time_func_dual(fn, method="adaptive", min_run_time=DEFAULT_MIN_RUN_TIME_SECONDS,
+                   kernel_method=None):
+    """Time *fn* and return ``(wall_ms, measurement, kernel_ms)``.
+
+    ``wall_ms`` / ``measurement`` are the host wall-clock timing from
+    :func:`time_func`. ``kernel_ms`` is the median GPU kernel (device) time per
+    call from :func:`_kernel_time_ms` when kernel profiling is enabled
+    (``--kernel-profile``); otherwise it is ``None`` and no profiler pass runs.
+    It is also ``None`` when the profiler pass can't attribute kernels to calls.
+    *kernel_method* selects the measurement ("profiler" default, or "event" for
+    concurrent multi-stream ops such as grouped GEMM).
+    """
+    wall_ms, measurement = time_func(fn, method=method, min_run_time=min_run_time)
+    kernel_ms = _kernel_time_ms(fn, method=kernel_method) if _KERNEL_PROFILE else None
+    return wall_ms, measurement, kernel_ms
+
+
+# A backward block keeps this many forward graphs alive (fewer if they'd exceed the byte cap).
+_BWD_BLOCK_CALLS = 20
+_BWD_BLOCK_BYTES = 16 * 1024**3
+_BWD_MIN_BLOCKS = 15
+
+
+def _bwd_block_calls(fwd_func):
+    torch.cuda.synchronize()
+    before = torch.cuda.memory_allocated()
+    fwd_out = fwd_func()
+    graph_bytes = torch.cuda.memory_allocated() - before
+    del fwd_out
+    return max(1, min(_BWD_BLOCK_CALLS, _BWD_BLOCK_BYTES // max(graph_bytes, 1)))
+
+
+def time_backward_dual(fwd_func, bwd_func, min_run_time=DEFAULT_MIN_RUN_TIME_SECONDS,
+                       kernel_method=None):
+    """Time the backward of *fwd_func* and return ``(wall_ms, measurement, kernel_ms)``.
+
+    *bwd_func* takes one *fwd_func* result, runs its backward and clears the gradients
+    it produced. Each block first runs its forwards, then times their backwards back to
+    back, so the host runs ahead of the GPU as in the forward timing and no forward time
+    is subtracted. The first block is warmup; GC is off while timing, as in
+    :func:`time_func`. ``measurement.times`` holds each block's time per call;
+    ``kernel_ms`` is as in :func:`time_func_dual`.
+    """
+    calls = _bwd_block_calls(fwd_func)
+    times = []
+    for block in itertools.count():
+        fwd_outs = [fwd_func() for _ in range(calls)]
+        torch.cuda.synchronize()
+        gc_enabled = gc.isenabled()
+        gc.disable()
+        try:
+            t0 = time.perf_counter()
+            for fwd_out in fwd_outs:
+                bwd_func(fwd_out)
+            torch.cuda.synchronize()
+            elapsed = time.perf_counter() - t0
+        finally:
+            if gc_enabled:
+                gc.enable()
+        del fwd_outs, fwd_out
+        if block == 0:
+            continue
+        times.append(elapsed / calls)
+        if len(times) >= _BWD_MIN_BLOCKS and sum(times) * calls >= min_run_time:
+            break
+    kernel_ms = (_kernel_time_ms(bwd_func, method=kernel_method, prepare=fwd_func, block=calls)
+                 if _KERNEL_PROFILE else None)
+    return statistics.median(times) * 1e3, SimpleNamespace(times=times), kernel_ms
 
 
 # ---------------------------------------------------------------------------
@@ -199,7 +480,7 @@ def time_func(fn, method="adaptive", min_run_time=DEFAULT_MIN_RUN_TIME_SECONDS):
 # ---------------------------------------------------------------------------
 # Benchmark inputs are cycled through a ring of buffers so that back-to-back
 # kernel launches read different input memory and don't benefit from artificial
-# cache residency.  Populated by run_benchmarks() from the parsed CLI args.
+# cache residency.  Configured from the CLI options via conftest.
 _ROTATE_BUFFERS = True
 _ROTATE_MB = 0  # rotation memory budget in MB; 0 => auto-size to exceed the LLC
 # Ceiling on the rotation ring size. hipBLASLt-bench caps its rotating block
@@ -314,87 +595,72 @@ def compute_gbps(nbytes, ms):
     return nbytes / (ms * 1e-3) / 1e9
 
 
-def make_metric_record(label, ms, unit, throughput, derived=False,
-                       ms_precision=3, throughput_precision=2,
-                       measurement=None, samples_only=False):
+def make_metric_record(label, ms, unit, throughput,
+                       ms_precision=4, throughput_precision=2,
+                       measurement=None,
+                       kernel_ms=None, kernel_throughput=None):
     """Create a structured metric record for stdout and CSV generation.
 
     Each record describes one benchmark line item such as "GEMM Forward".
-    ``run_benchmarks`` formats these records for stdout and expands them into
-    ``<label> Time (ms)`` and ``<label> <unit>`` CSV columns.
+    The harness formats these records for stdout and expands them into
+    ``<label> Wall Time (ms)`` / ``<label> Wall <unit>`` columns, plus
+    ``<label> Kernel Time (ms)`` / ``<label> Kernel <unit>`` when kernel timing
+    is enabled (``--kernel-profile``; the cells are empty when ``kernel_ms`` is None).
 
     If *measurement* is provided (a ``torch.utils.benchmark.Measurement``),
     the per-sample times are available for the ``--csv-samples`` output.
-    Records with *samples_only=True* are excluded from stdout and the main
-    CSV but their samples are still written to the samples CSV.
     """
     return {
         "label": label,
         "ms": ms,
         "unit": unit,
         "throughput": throughput,
-        "derived": derived,
         "ms_precision": ms_precision,
         "throughput_precision": throughput_precision,
         "measurement": measurement,
-        "samples_only": samples_only,
+        "kernel_ms": kernel_ms,
+        "kernel_throughput": kernel_throughput,
     }
 
 
-def make_forward_backward_metric_records(label_prefix, unit,
-                                         forward_ms, forward_throughput,
-                                         backward_ms, backward_throughput,
-                                         backward_derived=False,
-                                         ms_precision=3,
-                                         throughput_precision=2,
-                                         fwd_measurement=None,
-                                         bwd_measurement=None,
-                                         fwd_bwd_measurement=None):
-    """Create standard forward/backward metric records for a benchmark.
-
-    When *backward_derived* is True and *fwd_bwd_measurement* is provided,
-    an extra samples-only record for "Forward+Backward" is emitted so that
-    the raw timing samples are preserved in the ``--csv-samples`` output.
+def direction_records(direction, label, unit, throughput,
+                      fwd_func, bwd_func, fwd_work, bwd_work, kernel_method=None):
+    """Metric records for a forward or a backward timing.
+
+    *direction* is ``"fwd"`` or ``"bwd"``. *throughput* is ``compute_tflops`` or
+    ``compute_gbps`` and *fwd_work* / *bwd_work* the matching flops / bytes.
+    Forward times *fwd_func* (:func:`time_func_dual`); backward times *bwd_func* on
+    *fwd_func* results (:func:`time_backward_dual`). *kernel_method* is forwarded
+    ("event" for concurrent multi-stream ops).
     """
-    records = [
-        make_metric_record(
-            f"{label_prefix} Forward",
-            forward_ms,
-            unit,
-            forward_throughput,
-            ms_precision=ms_precision,
-            throughput_precision=throughput_precision,
-            measurement=fwd_measurement,
-        ),
-        make_metric_record(
-            f"{label_prefix} Backward",
-            backward_ms,
-            unit,
-            backward_throughput,
-            derived=backward_derived,
-            ms_precision=ms_precision,
-            throughput_precision=throughput_precision,
-            measurement=bwd_measurement,
-        ),
-    ]
-    if fwd_bwd_measurement is not None:
-        records.append(make_metric_record(
-            f"{label_prefix} Forward+Backward",
-            forward_ms + backward_ms,
-            unit,
-            0,
-            samples_only=True,
-            measurement=fwd_bwd_measurement,
-        ))
-    return records
+    if direction == "fwd":
+        ms, measurement, kernel_ms = time_func_dual(fwd_func, kernel_method=kernel_method)
+        work = fwd_work
+    else:
+        ms, measurement, kernel_ms = time_backward_dual(
+            fwd_func, bwd_func, kernel_method=kernel_method)
+        work = bwd_work
+    return [make_metric_record(
+        label, ms, unit, throughput(work, ms), measurement=measurement,
+        kernel_ms=kernel_ms,
+        kernel_throughput=throughput(work, kernel_ms) if kernel_ms else None,
+    )]
 
 
 def _metric_time_key(metric):
-    return f"{metric['label']} Time (ms)"
+    return f"{metric['label']} Wall Time (ms)"
 
 
 def _metric_throughput_key(metric):
-    return f"{metric['label']} {metric['unit']}"
+    return f"{metric['label']} Wall {metric['unit']}"
+
+
+def _metric_kernel_time_key(metric):
+    return f"{metric['label']} Kernel Time (ms)"
+
+
+def _metric_kernel_throughput_key(metric):
+    return f"{metric['label']} Kernel {metric['unit']}"
 
 
 def _format_metric_number(value, precision):
@@ -404,295 +670,673 @@ def _format_metric_number(value, precision):
 def _metric_row_from_records(metric_records):
     row = {}
     for metric in metric_records:
-        if metric.get("samples_only"):
-            continue
         row[_metric_time_key(metric)] = _format_metric_number(
-            metric["ms"], metric.get("ms_precision", 3)
+            metric["ms"], metric.get("ms_precision", 4)
         )
         row[_metric_throughput_key(metric)] = _format_metric_number(
             metric["throughput"], metric.get("throughput_precision", 2)
         )
+        # Kernel columns appear only under --kernel-profile, empty where unmeasured.
+        if _KERNEL_PROFILE or metric.get("kernel_ms") is not None:
+            km = metric.get("kernel_ms")
+            row[_metric_kernel_time_key(metric)] = (
+                _format_metric_number(km, metric.get("ms_precision", 4))
+                if km is not None else ""
+            )
+            kt = metric.get("kernel_throughput")
+            row[_metric_kernel_throughput_key(metric)] = (
+                _format_metric_number(kt, metric.get("throughput_precision", 2))
+                if kt is not None else ""
+            )
+        if "kernel_names" in metric:
+            row["Kernels"] = metric["kernel_names"]
     return row
 
 
 def _print_metric_records(metric_records):
-    printable = [m for m in metric_records if not m.get("samples_only")]
-    if not printable:
+    if not metric_records:
         return
-    label_width = max(24, *(len(metric["label"]) for metric in printable))
-    for metric in printable:
-        ms_str = _format_metric_number(metric["ms"], metric.get("ms_precision", 3))
+    label_width = max(24, *(len(metric["label"]) for metric in metric_records))
+    for metric in metric_records:
+        ms_str = _format_metric_number(metric["ms"], metric.get("ms_precision", 4))
         throughput_str = _format_metric_number(
             metric["throughput"], metric.get("throughput_precision", 2)
         )
-        derived_suffix = " (derived)" if metric.get("derived", False) else ""
-        print(
+        line = (
             f"  {metric['label']:<{label_width}} {ms_str} ms | "
-            f"{throughput_str} {metric['unit']}{derived_suffix}"
+            f"{throughput_str} {metric['unit']}"
         )
-
-
-def _default_csv_name(bench_fn):
-    import inspect
-    from pathlib import Path
-    return Path(inspect.getfile(bench_fn)).with_suffix(".csv").name
+        kernel_ms = metric.get("kernel_ms")
+        if kernel_ms is not None:
+            k_ms = _format_metric_number(kernel_ms, metric.get("ms_precision", 4))
+            kt = metric.get("kernel_throughput")
+            k_thr = (_format_metric_number(kt, metric.get("throughput_precision", 2))
+                     if kt is not None else "-")
+            line += f"  ||  kernel {k_ms} ms | {k_thr} {metric['unit']}"
+        print(line)
 
 
 # ---------------------------------------------------------------------------
-# Benchmark runner
+# pytest-based execution support
 # ---------------------------------------------------------------------------
+# conftest.py drives the microbenchmarks under pytest via the framework-agnostic
+# helpers below (no pytest import here, so importing utils.py never requires it).
+# Results are collected per family (test module) and written with the CSV / samples
+# schema the dashboard ingest consumes.
+
+# Backward runs on the autograd device thread; spread over cores that schedutil keeps at
+# partial-load clocks, launch-bound backward wall time flips up to 2x between runs.
+_PIN_CPUS = 4
+_CPU_SYSFS = Path("/sys/devices/system/cpu")
+
+# Runs as a separate process: EOF on stdin without a prior message means the benchmark process
+# died without calling restore(), so put the governors ("path=old" args) back.
+# Reads one byte, not up to EOF: forked workers may hold the pipe open after restore().
+_GOVERNOR_GUARD = r"""
+import sys
+if not sys.stdin.buffer.read(1):
+    for arg in sys.argv[1:]:
+        path, _, old = arg.rpartition("=")
+        try:
+            open(path, "w").write(old)
+        except OSError:
+            pass
+"""
+
+
+def _parse_cpulist(text):
+    cpus = set()
+    for part in text.strip().split(","):
+        if part:
+            lo, _, hi = part.partition("-")
+            cpus.update(range(int(lo), int(hi or lo) + 1))
+    return cpus
+
+
+def _gpu_local_cores():
+    """Physical cores (one per SMT pair) on the NUMA node of cuda:0 that we may run on."""
+    allowed = os.sched_getaffinity(0)
+    cpus = allowed
+    try:
+        node = int(Path(f"/sys/bus/pci/devices/{_gpu_pci()}/numa_node").read_text())
+        if node >= 0:
+            node_cpus = Path(f"/sys/devices/system/node/node{node}/cpulist").read_text()
+            cpus = (_parse_cpulist(node_cpus) & allowed) or allowed
+    except (OSError, ValueError):
+        pass
+    cores = []
+    for cpu in sorted(cpus):
+        try:
+            siblings = (_CPU_SYSFS / f"cpu{cpu}/topology/thread_siblings_list").read_text()
+            if cpu != min(_parse_cpulist(siblings) & cpus):
+                continue
+        except (OSError, ValueError):
+            pass
+        cores.append(cpu)
+    return cores
+
+
+def pin_cpus():
+    """Pin the run to a few GPU-local physical cores and set them to the performance governor.
+
+    The main thread gets the cores; threads it creates later (the autograd device thread)
+    inherit them. The threads that already exist after HIP init (HIP/ROCr helpers, incl. a
+    busy event-polling one) go to the cores' SMT siblings: the multi-stream grouped GEMM
+    path runs ~1.6x slower when that helper lands on another CCD or socket than the main
+    thread, and slower still when it shares the main thread's cores (the fallback when the
+    process may not use the siblings).
+    The cores are split into groups of ``_PIN_CPUS`` and the GPU id picks the group, so runs
+    on other GPUs of the node use different ones (wrapping around if there are fewer groups).
+    The governor needs root; without it only the pinning is applied. Returns a namespace
+    with ``description`` and ``restore()``, which puts the previous governors back; if the
+    process dies first (killed, crashed), a watchdog process puts them back instead.
+    """
+    torch.cuda.init()
+    allowed = os.sched_getaffinity(0)
+    cores = _gpu_local_cores()
+    try:
+        gpu = int((os.environ.get("HIP_VISIBLE_DEVICES", "") or "0").split(",")[0])
+    except ValueError:
+        gpu = 0
+    groups = len(cores) // _PIN_CPUS
+    if groups < 1:
+        cpus = cores
+    else:
+        start = _PIN_CPUS * ((gpu + 1) % groups)
+        cpus = cores[start:start + _PIN_CPUS]
+    siblings = set()
+    for cpu in cpus:
+        try:
+            siblings |= _parse_cpulist((_CPU_SYSFS / f"cpu{cpu}/topology/thread_siblings_list").read_text())
+        except (OSError, ValueError):
+            pass
+    helpers = sorted((siblings & allowed) - set(cpus)) or cpus
+    main = threading.get_native_id()
+    for tid in os.listdir("/proc/self/task"):
+        if int(tid) != main:
+            try:
+                os.sched_setaffinity(int(tid), helpers)
+            except OSError:
+                pass
+    os.sched_setaffinity(0, cpus)
+
+    previous, error = {}, None
+    for cpu in cpus:
+        # CPUs that share a cpufreq policy link to the same governor file.
+        path = (_CPU_SYSFS / f"cpu{cpu}/cpufreq/scaling_governor").resolve()
+        if path in previous:
+            continue
+        try:
+            old = path.read_text().strip()
+            path.write_text("performance")
+            previous[path] = old
+        except OSError as e:
+            error = e
+    guard = None
+
+    def restore():
+        for path, old in previous.items():
+            try:
+                path.write_text(old)
+            except OSError:
+                pass
+        if guard is not None:
+            guard.communicate(b"restored")
+
+    if previous:
+        try:
+            # Own session, so a terminal Ctrl-C doesn't kill the guard along with the benchmark.
+            guard = subprocess.Popen(
+                [sys.executable, "-c", _GOVERNOR_GUARD, *(f"{p}={old}" for p, old in previous.items())],
+                stdin=subprocess.PIPE, start_new_session=True)
+        except BaseException:
+            restore()
+            raise
+
+    governor = "performance governor" if error is None else f"governor unchanged: {error}"
+    return SimpleNamespace(
+        description=f"main thread pinned to cpus {','.join(map(str, cpus))}, helper threads to "
+                    f"{','.join(map(str, helpers))} ({governor})",
+        restore=restore)
+
+def configure_rotating(rotating, no_rotating):
+    """Set module-level input-rotation state from parsed options."""
+    global _ROTATE_BUFFERS, _ROTATE_MB
+    if rotating is not None and rotating < 0:
+        raise ValueError("--rotating expects a non-negative size in MB")
+    _ROTATE_BUFFERS = not no_rotating
+    _ROTATE_MB = rotating or 0
 
-def make_parser(**kwargs):
-    """Return an :class:`~argparse.ArgumentParser` with ``--csv``, ``--csv-samples``, and ``--kernel-profile`` flags.
 
-    Any *kwargs* are forwarded to the ``ArgumentParser`` constructor, so
-    callers can set ``description``, ``parents``, etc.
+def apply_backend_env(monkeypatch, env):
+    """Force a kernel backend for one test by setting/unsetting env vars.
+
+    A ``None`` value unsets the var, so forcing one backend cleanly clears the
+    toggles that would select a competing one; pytest restores them afterwards.
     """
-    parser = argparse.ArgumentParser(**kwargs)
-    parser.add_argument(
-        "--csv", nargs="?", const=True, default=None, metavar="FILE",
-        help="Write results to CSV. Optional filename; default derived from script name.",
-    )
-    parser.add_argument(
-        "--csv-samples", nargs="?", const=True, default=None, metavar="FILE",
-        help=(
-            "Write per-sample timing data to a CSV for downstream analysis. "
-            "Optional filename; default derived from script name."
-        ),
-    )
-    parser.add_argument(
-        "--kernel-profile", action="store_true", default=False,
-        help=(
-            "Profile GPU kernels using torch.profiler in addition to normal "
-            "timing. Prints per-kernel CUDA times output. "
-            "Use with --csv to write kernel-level data to CSV. "
-            "--csv-samples is ignored in this mode."
-        ),
-    )
-    rotating_group = parser.add_mutually_exclusive_group()
-    rotating_group.add_argument(
-        "--rotating", nargs="?", type=int, const=0, default=None, metavar="MB",
-        help=(
-            "Rotate benchmark inputs through a ring of buffers so back-to-back "
-            "launches touch different memory (avoids artificial cache "
-            "residency), like hipBLASLt-bench --rotating. Optionally pass the "
-            "rotating memory budget in MB; omit it to auto-size the ring to "
-            "exceed the last-level cache (the 256 MB Infinity Cache on "
-            "gfx942/gfx950, not just L2). On by default; disable with "
-            "--no-rotating."
-        ),
-    )
-    rotating_group.add_argument(
-        "--no-rotating", action="store_true", default=False,
-        help=(
-            "Disable input buffer rotation (see --rotating) and time a single "
-            "cached input buffer."
-        ),
-    )
-    return parser
+    for key, value in env.items():
+        if value is None:
+            monkeypatch.delenv(key, raising=False)
+        else:
+            monkeypatch.setenv(key, value)
 
 
-_KERNEL_NAME_MAX_WIDTH = 80
+@functools.lru_cache(maxsize=1)
+def _te_install_root():
+    import transformer_engine as _te
+    return Path(_te.__file__).resolve().parent
 
 
-def _shorten_kernel_name(name):
-    """Shorten verbose C++/HIP kernel names for readable terminal output.
+@functools.lru_cache(maxsize=1)
+def _te_py_source():
+    chunks = []
+    for p in _te_install_root().rglob("*.py"):
+        try:
+            chunks.append(p.read_text(errors="ignore"))
+        except OSError:
+            pass
+    return "\n".join(chunks)
 
-    Strips ``void `` prefix and template arguments (``<...>``) from
-    fully-qualified kernel names while preserving the function name and
-    any top-level namespace.
+
+@functools.lru_cache(maxsize=1)
+def _te_loaded_libs():
+    """Paths of libtransformer_engine*.so mapped into this process.
+
+    Catches the compiled lib that a source/editable install keeps at the repo
+    root -- outside the package dir _te_install_root scans.
     """
+    libs = []
+    seen = set()
+    try:
+        with open("/proc/self/maps") as fh:
+            for line in fh:
+                parts = line.split(None, 5)
+                path = parts[5].strip() if len(parts) >= 6 else ""
+                if path.endswith(".so") and "libtransformer_engine" in path and path not in seen:
+                    seen.add(path)
+                    libs.append(Path(path))
+    except OSError:
+        pass
+    return libs
+
+
+@functools.lru_cache(maxsize=None)
+def te_honors_env(varname):
+    """True if the installed TE build reads *varname* (python dispatch or compiled .so).
+
+    Lets a benchmark skip a forced backend on builds that predate its dispatch,
+    instead of silently measuring the default (which shows up as a phantom
+    regression in a long-running history sweep).
+    """
+    try:
+        root = _te_install_root()
+    except Exception:
+        return False
+    if varname in _te_py_source():
+        return True
+    needle = varname.encode()
+    # Include the compiled TE lib actually loaded into the process: a source/
+    # editable install keeps libtransformer_engine.so at the repo root, outside
+    # the package dir _te_install_root scans.
+    for so in list(root.rglob("*.so")) + _te_loaded_libs():
+        try:
+            with open(so, "rb") as fh, mmap.mmap(fh.fileno(), 0, access=mmap.ACCESS_READ) as mm:
+                if mm.find(needle) != -1:
+                    return True
+        except (OSError, ValueError):
+            pass
+    return False
+
+
+class _FamilyResults:
+    """Accumulated rows / samples for one benchmark family."""
+
+    def __init__(self):
+        self.param_columns = None
+        self.metric_columns = None
+        self.rows = []
+        self.case_metrics = []
+
+
+def _stringify_params(case_params):
+    return {k: (str(v) if isinstance(v, torch.dtype) else v) for k, v in case_params.items()}
+
+
+def record_bench(store, family, case_params, metric_records, node_name=""):
+    """Record one benchmark case into *store* (a dict keyed by *family*)."""
+    fam = store.setdefault(family, _FamilyResults())
+    metric_row = _metric_row_from_records(metric_records)
+    metric_columns = list(metric_row.keys())
+    if fam.param_columns is None:
+        fam.param_columns = list(case_params.keys())
+        fam.metric_columns = metric_columns
+    elif metric_columns != fam.metric_columns:
+        raise ValueError(
+            f"Inconsistent metric columns for {family}: "
+            f"expected {fam.metric_columns}, got {metric_columns}"
+        )
+    row = _stringify_params(case_params)
+    row.update(metric_row)
+    fam.rows.append(row)
+    fam.case_metrics.append((_stringify_params(case_params), metric_records, node_name))
+
+
+def print_case(case_params, metric_records):
+    """Print a case header and its metric lines (reused stdout format)."""
+    label = "  ".join(f"{k}={v}" for k, v in case_params.items())
+    print(f"\n{'='*60}\nTesting: {label}\n{'='*60}")
+    _print_metric_records(metric_records)
+
+
+def _write_csv(path, rows, columns):
+    """Write list-of-dict *rows* to *path* with header *columns*."""
+    with open(path, "w", newline="") as f:
+        writer = csv.DictWriter(f, fieldnames=columns, extrasaction="ignore", lineterminator="\n")
+        writer.writeheader()
+        writer.writerows(rows)
+
+
+def write_bench_outputs(store, *, csv=None, csv_samples=None):
+    """Write per-family CSV / samples outputs; return paths written."""
+    # When an explicit filename is given but several families run in one session
+    # (e.g. `pytest .`), insert the family name so they don't overwrite each other.
+    multi = sum(1 for fam in store.values() if fam.rows) > 1
+
+    def _dest(explicit, family, default_name):
+        if not isinstance(explicit, str):
+            return default_name
+        if not multi:
+            return explicit
+        p = Path(explicit)
+        return str(p.with_name(f"{p.stem}-{family}{p.suffix}"))
+
+    written = []
+    for family, fam in store.items():
+        if not fam.rows:
+            continue
+        if csv is not None:
+            out = _dest(csv, family, f"{family}.csv")
+            _write_csv(out, fam.rows, fam.param_columns + fam.metric_columns)
+            written.append(out)
+        if csv_samples is not None:
+            sout = _dest(csv_samples, family, f"{family}_samples.csv")
+            sample_rows = _sample_rows(fam)
+            if sample_rows:
+                _write_csv(sout, sample_rows,
+                           fam.param_columns + ["label", "sample_idx", "time_ms"])
+                written.append(sout)
+    return written
+
+
+def _sample_rows(fam):
+    """Flatten a family's per-iteration timing measurements into long rows."""
+    rows = []
+    for case_params, records, _node in fam.case_metrics:
+        for metric in records:
+            m = metric.get("measurement")
+            if m is None:
+                continue
+            for i, t in enumerate(m.times):
+                sr = dict(case_params)
+                sr["label"] = metric["label"]
+                sr["sample_idx"] = i
+                sr["time_ms"] = round(t * 1e3, 6)
+                rows.append(sr)
+    return rows
+
+
+# ---------------------------------------------------------------------------
+# Dashboard run: one pytest session -> one run dir of born-tagged CSVs + samples/
+# + run_info.txt, ready for the TE-dashboard ingest. Replaces the shell runner.
+# ---------------------------------------------------------------------------
+
+def _detect_gpu_model():
+    """Short GPU token (e.g. ``MI355X``) from torch's device name, else None."""
     import re
-    s = name
-    # Strip leading "void "
-    if s.startswith("void "):
-        s = s[5:]
-    # Remove balanced template args (handles one level of nesting)
-    s = re.sub(r"<[^<>]*(?:<[^<>]*>[^<>]*)*>", "", s)
-    # Collapse whitespace
-    s = " ".join(s.split())
-    if len(s) > _KERNEL_NAME_MAX_WIDTH:
-        s = s[: _KERNEL_NAME_MAX_WIDTH - 3] + "..."
-    return s
-
-
-def run_benchmarks(test_cases, bench_fn, param_columns, default_csv=None,
-                   args=None):
-    """Iterate *test_cases*, call *bench_fn*, and optionally write a CSV.
-
-    Parameters
-    ----------
-    test_cases : list[dict]
-        Each dict has at least the keys in *param_columns* plus any extra
-        keys the bench_fn needs (passed as **case).
-    bench_fn : callable
-        Called as ``bench_fn(**case)`` and must return a list of metric
-        records created by ``make_metric_record``. Each record corresponds to
-        one stdout line and expands to a time column plus a throughput column in
-        the CSV output.
-    param_columns : list[str]
-        Column names to pull from each test case into the output row.
-    default_csv : str or None
-        Default CSV filename used when ``--csv`` is passed without a
-        filename. If omitted, the CSV name is derived from the caller's
-        file name. CSV output is only written when the caller passes
-        ``--csv`` on the command line.
-    args : argparse.Namespace or None
-        Pre-parsed arguments.  When a benchmark script needs its own CLI
-        flags it can call ``parser = make_parser()``, add custom
-        arguments, run ``args = parser.parse_args()``, and then pass
-        *args* here.  If *None*, a default parser with only
-        ``--csv`` / ``--csv-samples`` is created and ``parse_args()``
-        is called automatically.
+    try:
+        name = torch.cuda.get_device_name(0)
+    except Exception:
+        return None
+    m = re.search(r"MI\s?\d{3,4}[A-Za-z]*", name)
+    return m.group(0).replace(" ", "").upper() if m else None
+
+
+def _gpu_pci():
+    """PCI BDF of the active GPU (torch device 0), e.g. ``0000:75:00.0``; '' if unknown."""
+    try:
+        p = torch.cuda.get_device_properties(0)
+        return "%04x:%02x:%02x.0" % (p.pci_domain_id, p.pci_bus_id, p.pci_device_id)
+    except Exception:
+        return ""
+
+
+def _dashboard_run_meta(model=None):
+    """Run metadata: run timestamp, git sha/date, GPU model, visible GPU id, host.
+
+    *model* overrides the auto-detected GPU model when given (e.g. --gpu-model).
     """
-    if args is None:
-        args = make_parser().parse_args()
+    import datetime
+    import subprocess
+
+    here = Path(__file__).resolve().parent
+
+    def git(*args):
+        # safe.directory=* bypasses git's dubious-ownership guard when the repo
+        # is bind-mounted into a container under a different uid.
+        try:
+            return subprocess.check_output(
+                ["git", "-c", "safe.directory=*", "-C", str(here), *args],
+                text=True, stderr=subprocess.DEVNULL,
+            ).strip()
+        except (OSError, subprocess.CalledProcessError):
+            return ""
+
+    sha = git("rev-parse", "HEAD")
+    gpu = (os.environ.get("HIP_VISIBLE_DEVICES", "") or "0").split(",")[0].strip() or "0"
+    now = datetime.datetime.now().astimezone()
+    return {
+        "run_ts": now.isoformat(timespec="seconds"),
+        "stamp": now.strftime("%Y-%m-%dT%H-%M-%S"),
+        "sha": sha,
+        "short": sha[:12] or "unknown",
+        "cdate": git("show", "-s", "--format=%cI", "HEAD"),
+        "model": model or _detect_gpu_model(),
+        "gpu": gpu,
+        "gpu_bdf": _gpu_pci(),
+        "host": os.uname().nodename.split(".")[0],
+    }
 
-    global _ROTATE_BUFFERS, _ROTATE_MB
-    _rotating = getattr(args, "rotating", None)
-    if _rotating is not None and _rotating < 0:
-        raise ValueError("--rotating expects a non-negative size in MB")
-    _ROTATE_BUFFERS = not getattr(args, "no_rotating", False)
-    _ROTATE_MB = _rotating or 0
 
-    if args.kernel_profile:
-        from torch.profiler import profile, ProfilerActivity
+def _write_run_info(path, meta):
+    """Write a small version/machine manifest alongside a run's CSVs."""
+    import datetime
+    import importlib
+
+    def ver(mod):
+        try:
+            return getattr(importlib.import_module(mod), "__version__", "?")
+        except Exception:
+            return "?"
+
+    def read(p):
+        try:
+            return Path(p).read_text().strip()
+        except OSError:
+            return ""
+
+    rocm = read("/opt/rocm/.info/version") or getattr(torch.version, "hip", "") or "?"
+    info = {
+        "date": datetime.datetime.now().astimezone().isoformat(timespec="seconds"),
+        "host": meta["host"],
+        "gpu_id": meta["gpu"],
+        "gpu_bdf": meta.get("gpu_bdf") or "?",
+        "arch": meta["model"],
+        "te_commit": meta["sha"],
+        "te": ver("transformer_engine"),
+        "pytorch": ver("torch"),
+        "triton": ver("triton"),
+        "jax": ver("jax"),
+        "rocm": rocm,
+        "amdgpu_drv": read("/sys/module/amdgpu/version") or "?",
+        "kernel": os.uname().release,
+        "cpus": meta.get("cpus") or "unpinned",
+        "hw_queues": os.environ.get("GPU_MAX_HW_QUEUES", "HIP default"),
+    }
+    path.write_text("".join(f"{k + ':':12}{v}\n" for k, v in info.items()))
 
-    rows = []
-    all_case_metrics = []
-    all_kernel_rows = []
-    resolved_metric_columns = None
-
-    for case in test_cases:
-        label = "  ".join(f"{k}={case[k]}" for k in param_columns)
-        print(f"\n{'='*60}")
-        print(f"Testing: {label}")
-        print(f"{'='*60}")
-
-        metric_records = bench_fn(**case)
-        metric_row = _metric_row_from_records(metric_records)
-        _print_metric_records(metric_records)
-        current_metric_columns = list(metric_row.keys())
-
-        if resolved_metric_columns is None:
-            resolved_metric_columns = current_metric_columns
-        elif current_metric_columns != resolved_metric_columns:
-            raise ValueError(
-                f"Inconsistent metric columns for case {case}: "
-                f"expected {resolved_metric_columns}, got {current_metric_columns}"
-            )
 
-        case_params = {k: (str(case[k]) if isinstance(case[k], torch.dtype) else case[k])
-                       for k in param_columns}
-        row = dict(case_params)
-        row.update(metric_row)
-        rows.append(row)
-        all_case_metrics.append((case_params, metric_records))
-
-        if args.kernel_profile:
-            with profile(
-                activities=[ProfilerActivity.CUDA],
-            ) as prof:
-                bench_fn(**case)
-                torch.cuda.synchronize()
+def _dashboard_run_dir(out_base, meta):
+    """Run dir path from run metadata: ``<out_base>/<node>/<datetime>_<commit>/``."""
+    node = f"{meta['model']}_{meta['host']}_gpu{meta['gpu']}"
+    if meta["gpu_bdf"]:
+        node += "_" + meta["gpu_bdf"].replace(":", "-").replace(".", "-")
+    return Path(out_base) / node / f"{meta['stamp']}_{meta['short']}"
 
-            averages = prof.key_averages()
-            gpu_events = [e for e in averages if e.self_device_time_total > 0]
-            gpu_events.sort(key=lambda e: e.self_device_time_total, reverse=True)
-
-            if gpu_events:
-                total_cuda_us = sum(e.self_device_time_total for e in gpu_events)
-                w = _KERNEL_NAME_MAX_WIDTH
-                print(
-                    f"\n  | {'Kernel':<{w}} | "
-                    f"{'Total (us)':>11} | {'Calls':>6} | "
-                    f"{'Avg (us)':>10} | {'%':>6} |"
-                )
-                print(
-                    f"  | {'-'*w} | "
-                    f"{'-'*11} | {'-'*6} | "
-                    f"{'-'*10} | {'-'*6} |"
-                )
-                for e in gpu_events:
-                    avg_us = e.self_device_time_total / e.count if e.count > 0 else 0
-                    pct = (
-                        100.0 * e.self_device_time_total / total_cuda_us
-                        if total_cuda_us > 0
-                        else 0
-                    )
-                    short = _shorten_kernel_name(e.key)
-                    print(
-                        f"  | {short:<{w}} | {e.self_device_time_total:>11.1f} | "
-                        f"{e.count:>6} | {avg_us:>10.2f} | {pct:>5.1f}% |"
-                    )
-                print(
-                    f"  | {'TOTAL':<{w}} | {total_cuda_us:>11.1f} | "
-                    f"{'---':>6} | {'---':>10} | {'---':>6} |"
-                )
-
-            for e in gpu_events:
-                kr = dict(case_params)
-                kr["kernel_name"] = e.key
-                kr["cuda_time_total_us"] = round(e.self_device_time_total, 1)
-                kr["num_calls"] = e.count
-                kr["cuda_time_avg_us"] = (
-                    round(e.self_device_time_total / e.count, 2) if e.count > 0 else 0
-                )
-                all_kernel_rows.append(kr)
-
-    if args.csv is not None:
-        import pandas as pd
-        out_csv = args.csv if isinstance(args.csv, str) else (
-            default_csv or _default_csv_name(bench_fn)
-        )
-        columns = param_columns + (resolved_metric_columns or [])
-        results = pd.DataFrame(rows, columns=columns)
-        results.to_csv(out_csv, index=False)
-        print(f"\nResults saved to {out_csv}")
-
-    if args.kernel_profile and args.csv is not None and all_kernel_rows:
-        import pandas as pd
-        from pathlib import Path
-        base = default_csv or _default_csv_name(bench_fn)
-        out_csv_name = args.csv if isinstance(args.csv, str) else (
-            Path(base).stem + "_kernel_profile.csv"
-        )
-        # Don't overwrite the main CSV if --csv was given a filename
-        if isinstance(args.csv, str):
-            out_csv_name = Path(args.csv).stem + "_kernel_profile.csv"
-        kernel_columns = param_columns + [
-            "kernel_name", "cuda_time_total_us", "num_calls", "cuda_time_avg_us",
+
+def dashboard_run_plan(out_base="results", model=None):
+    """Resolve (meta, run_dir) for a dashboard run without writing anything."""
+    meta = _dashboard_run_meta(model=model)
+    return SimpleNamespace(meta=meta, run_dir=_dashboard_run_dir(out_base, meta))
+
+
+def write_dashboard_run(store, *, out_base="results", csv_samples=None, meta=None, model=None):
+    """Write one dashboard run and return the run dir.
+
+    Layout: ``<out_base>/<arch>_<host>_gpu<id>_<pci>/<datetime>_<commit>/``. Per-family
+    CSVs are born-tagged with run_ts/commit_sha/commit_date (so dashboard_ingest.py
+    consumes them directly); per-iteration samples land untagged under ``samples/``;
+    ``run_info.txt`` records versions + machine. Pass *meta* from
+    :func:`dashboard_run_plan` to reuse the metadata resolved at session start.
+    """
+    if meta is None:
+        meta = _dashboard_run_meta(model=model)
+    run_dir = _dashboard_run_dir(out_base, meta)
+    (run_dir / "samples").mkdir(parents=True, exist_ok=True)
+    _write_run_info(run_dir / "run_info.txt", meta)
+
+    tag = {"run_ts": meta["run_ts"], "commit_sha": meta["sha"], "commit_date": meta["cdate"]}
+    tag_cols = list(tag)
+    for family, fam in store.items():
+        if not fam.rows:
+            continue
+        rows = [{**tag, **row} for row in fam.rows]
+        _write_csv(run_dir / f"{family}.csv", rows,
+                   tag_cols + fam.param_columns + fam.metric_columns)
+        if csv_samples is not None:
+            sample_rows = _sample_rows(fam)
+            if sample_rows:
+                _write_csv(run_dir / "samples" / f"{family}_samples.csv", sample_rows,
+                           fam.param_columns + ["label", "sample_idx", "time_ms"])
+    return run_dir
+
+
+def _result_rows(store):
+    """Flatten *store* into (suite, name, wall_ms, wall_thr, kernel_ms, kernel_thr, unit) rows."""
+    rows = []
+    for family, fam in store.items():
+        suite = family[len("benchmark_"):] if family.startswith("benchmark_") else family
+        for _case_params, records, node_name in fam.case_metrics:
+            base = node_name
+            if base.endswith("]") and "[" in base:
+                base = base[base.index("[") + 1 : -1]  # keep the parametrize id
+            for m in records:
+                name = base if len(records) == 1 else f"{base} {m['label']}"
+                rows.append((
+                    suite, name, m["ms"], m["throughput"],
+                    m.get("kernel_ms"), m.get("kernel_throughput"), m["unit"],
+                ))
+    return rows
+
+
+def format_results_table(store):
+    """Render the results as a Markdown table (times in ms) with a caption line."""
+    rows = _result_rows(store)
+    if not rows:
+        return ""
+    has_kernel = any(r[4] is not None for r in rows)
+
+    def ms_cell(v):
+        return "-" if v is None else f"{v:.4f}"
+
+    def thr_cell(v, unit):
+        return "-" if v is None else f"{v:.2f} {unit}"
+
+    headers = ["Benchmark", "Config", "Wall Median (ms)", "Wall Throughput"]
+    if has_kernel:
+        headers += ["Kernel Median (ms)", "Kernel Throughput"]
+    body = []
+    for suite, name, wall_ms, wall_thr, kernel_ms, kernel_thr, unit in sorted(
+        rows, key=lambda r: (r[0], r[1])
+    ):
+        cells = [suite, name, ms_cell(wall_ms), thr_cell(wall_thr, unit)]
+        if has_kernel:
+            cells += [ms_cell(kernel_ms), thr_cell(kernel_thr, unit)]
+        body.append(cells)
+    widths = [max(len(headers[i]), *(len(r[i]) for r in body)) for i in range(len(headers))]
+
+    def row(cells):
+        # Left-align the text columns (Benchmark, Config); right-align the numerics.
+        padded = [
+            c.ljust(widths[i]) if i <= 1 else c.rjust(widths[i]) for i, c in enumerate(cells)
         ]
-        df = pd.DataFrame(all_kernel_rows, columns=kernel_columns)
-        df.to_csv(out_csv_name, index=False)
-        print(f"Kernel profile saved to {out_csv_name}")
-
-    if args.csv_samples is not None:
-        import pandas as pd
-        from pathlib import Path
-        base = default_csv or _default_csv_name(bench_fn)
-        samples_csv = args.csv_samples if isinstance(args.csv_samples, str) else (
-            Path(base).stem + "_samples.csv"
-        )
-        sample_rows = []
-        for case_params, records in all_case_metrics:
-            for metric in records:
-                measurement = metric.get("measurement")
-                if measurement is None:
+        return "| " + " | ".join(padded) + " |"
+
+    align = [
+        ":" + "-" * (widths[i] - 1) if i <= 1 else "-" * (widths[i] - 1) + ":"
+        for i in range(len(headers))
+    ]
+    caption = f"benchmark: {len(body)} tests"
+    return "\n".join(
+        [caption, "", row(headers), "| " + " | ".join(align) + " |", *(row(r) for r in body)]
+    )
+
+
+_THROUGHPUT_UNITS = frozenset({"TFLOPS", "GB/s"})
+_AGG_COLLAPSE_ALWAYS = frozenset({"Case", "dtype"})
+
+
+def _aggregate_rows(store):
+    """Work-weighted harmonic-mean throughput per (suite, group), collapsing shape/size axes."""
+    agg = {}
+    order = []
+    for family, fam in store.items():
+        suite = family[len("benchmark_"):] if family.startswith("benchmark_") else family
+        params = fam.param_columns or []
+        for case_params, records, _node in fam.case_metrics:
+            group_cols = [
+                c for c in params
+                if c not in _AGG_COLLAPSE_ALWAYS and isinstance(case_params.get(c), str)
+            ]
+            group = tuple((c, case_params[c]) for c in group_cols)
+            for m in records:
+                if m["unit"] not in _THROUGHPUT_UNITS:
                     continue
-                lbl = metric["label"]
-                for i, t in enumerate(measurement.times):
-                    sr = dict(case_params)
-                    sr["label"] = lbl
-                    sr["sample_idx"] = i
-                    # measurement.times is already per-iteration (raw block time
-                    # divided by number_per_run); convert seconds -> ms only.
-                    sr["time_ms"] = t * 1e3
-                    sample_rows.append(sr)
-        if sample_rows:
-            df = pd.DataFrame(
-                sample_rows,
-                columns=param_columns + ["label", "sample_idx", "time_ms"],
-            )
-            df.to_csv(samples_csv, index=False)
-            print(f"Samples saved to {samples_csv}")
+                thr = m["throughput"]
+                wall_ms = m["ms"]
+                if thr is None or not (thr > 0) or not (wall_ms > 0):
+                    continue
+                key = (suite, m["label"], group, m["unit"])
+                a = agg.get(key)
+                if a is None:
+                    a = agg[key] = {"wall_work": 0.0, "wall_time": 0.0,
+                                    "kern_work": 0.0, "kern_time": 0.0, "n": 0}
+                    order.append(key)
+                a["wall_work"] += thr * wall_ms
+                a["wall_time"] += wall_ms
+                a["n"] += 1
+                k_ms, k_thr = m.get("kernel_ms"), m.get("kernel_throughput")
+                if k_ms and k_thr and k_ms > 0:
+                    a["kern_work"] += k_thr * k_ms
+                    a["kern_time"] += k_ms
+    return agg, order
+
+
+def format_aggregate_table(store):
+    """Render per-group aggregate throughput (work-weighted harmonic mean) below the table."""
+    agg, order = _aggregate_rows(store)
+    if not agg:
+        return ""
+    # Include the metric label in the group text only when a suite emits more than one.
+    labels_per_suite = {}
+    for suite, label, _group, _unit in order:
+        labels_per_suite.setdefault(suite, set()).add(label)
+
+    has_kernel = any(a["kern_time"] > 0 for a in agg.values())
+
+    def thr_cell(work, time, unit):
+        return f"{work / time:.2f} {unit}" if time > 0 else "-"
+
+    headers = ["Benchmark", "Group", "n", "Wall Throughput"]
+    if has_kernel:
+        headers += ["Kernel Throughput"]
+    body = []
+    for key in order:
+        suite, label, group, unit = key
+        a = agg[key]
+        parts = [f"{c}={v}" for c, v in group]
+        if len(labels_per_suite[suite]) > 1:
+            parts.insert(0, label)
+        group_text = "  ".join(parts) or "all"
+        cells = [suite, group_text, str(a["n"]), thr_cell(a["wall_work"], a["wall_time"], unit)]
+        if has_kernel:
+            cells.append(thr_cell(a["kern_work"], a["kern_time"], unit))
+        body.append(cells)
+    body.sort(key=lambda r: (r[0], r[1]))
+    widths = [max(len(headers[i]), *(len(r[i]) for r in body)) for i in range(len(headers))]
+
+    def row(cells):
+        # Left-align the text columns (Benchmark, Group); right-align n / throughputs.
+        padded = [
+            c.ljust(widths[i]) if i <= 1 else c.rjust(widths[i]) for i, c in enumerate(cells)
+        ]
+        return "| " + " | ".join(padded) + " |"
+
+    align = [
+        ":" + "-" * (widths[i] - 1) if i <= 1 else "-" * (widths[i] - 1) + ":"
+        for i in range(len(headers))
+    ]
+    caption = (f"aggregate throughput (work-weighted harmonic mean over shapes): "
+               f"{len(body)} groups")
+    return "\n".join(
+        [caption, "", row(headers), "| " + " | ".join(align) + " |", *(row(r) for r in body)]
+    )