From 374488bb3abc856b356b79fbfbcc0e15bbb968c4 Mon Sep 17 00:00:00 2001 From: fxyan <1350821504@qq.com> Date: Sun, 18 Oct 2026 18:08:17 +0800 Subject: [PATCH 01/11] feat(h3): add h3_sigma_schedule with CPU reference and NVIDIA CUDA (Triton) backend Reproduces MiniMaxH3Scheduler.set_timesteps (diffusers f53d552) byte-for-byte: FP32 grid, shift transform, consecutive dedup, timesteps = 1 - sigmas[:-1]. Returns num_model_evals on the deduplicated grid (M - 1). Also wires the CPU-safe H3 suites into CI and states ROCm as unsupported rather than unvalidated: the device division is inline NVIDIA PTX. Refs #420 Signed-off-by: fxyan <1350821504@qq.com> --- .github/workflows/ci.yml | 4 + .../models/benchmark_h3_sigma_schedule.py | 317 ++ docs/.nav.yml | 1 + docs/operators/h3-sigma-schedule.md | 166 + .../experiments/h3-sigma-schedule/README.md | 82 + .../h3-sigma-schedule/benchmark-4090d.json | 2708 +++++++++++++++++ .../h3-sigma-schedule/cuda-4090d.xml | 2 + .../h3-sigma-schedule/source-provenance.json | 61 + .../backends/cuda/model_specific/__init__.py | 4 + .../model_specific/minimax_h3/__init__.py | 8 + .../minimax_h3/sigma_schedule.py | 331 ++ rl_engine/reference/minimax_h3/__init__.py | 6 + .../reference/minimax_h3/sigma_schedule.py | 102 + rl_engine/runtime/registry.py | 60 + .../cuda/test_h3_sigma_schedule_cuda.py | 366 +++ tests/reference/minimax_h3/__init__.py | 2 + .../h3_sigma_schedule_upstream_fixture.json | 201 ++ .../minimax_h3/make_upstream_fixture.py | 154 + .../minimax_h3/test_h3_sigma_schedule.py | 185 ++ .../test_h3_sigma_schedule_dispatch.py | 173 ++ 20 files changed, 4933 insertions(+) create mode 100644 benchmarks/models/benchmark_h3_sigma_schedule.py create mode 100644 docs/operators/h3-sigma-schedule.md create mode 100644 reports/experiments/h3-sigma-schedule/README.md create mode 100644 reports/experiments/h3-sigma-schedule/benchmark-4090d.json create mode 100644 reports/experiments/h3-sigma-schedule/cuda-4090d.xml create mode 100644 reports/experiments/h3-sigma-schedule/source-provenance.json create mode 100644 rl_engine/backends/cuda/model_specific/__init__.py create mode 100644 rl_engine/backends/cuda/model_specific/minimax_h3/__init__.py create mode 100644 rl_engine/backends/cuda/model_specific/minimax_h3/sigma_schedule.py create mode 100644 rl_engine/reference/minimax_h3/__init__.py create mode 100644 rl_engine/reference/minimax_h3/sigma_schedule.py create mode 100644 tests/backends/cuda/test_h3_sigma_schedule_cuda.py create mode 100644 tests/reference/minimax_h3/__init__.py create mode 100644 tests/reference/minimax_h3/h3_sigma_schedule_upstream_fixture.json create mode 100644 tests/reference/minimax_h3/make_upstream_fixture.py create mode 100644 tests/reference/minimax_h3/test_h3_sigma_schedule.py create mode 100644 tests/runtime/test_h3_sigma_schedule_dispatch.py diff --git a/.github/workflows/ci.yml b/.github/workflows/ci.yml index 4634569cd..328dd8530 100644 --- a/.github/workflows/ci.yml +++ b/.github/workflows/ci.yml @@ -71,6 +71,10 @@ jobs: - name: Check refactor packaging and compatibility run: python ci/run.py layout -q + - name: Run H3 Sigma Schedule Tests (CPU-safe) + run: | + python -m pytest tests/reference/minimax_h3 tests/runtime/test_h3_sigma_schedule_dispatch.py -q + - name: Run Mocked Hardware Discovery Tests run: | python -m pytest tests/runtime/test_dispatch.py -v diff --git a/benchmarks/models/benchmark_h3_sigma_schedule.py b/benchmarks/models/benchmark_h3_sigma_schedule.py new file mode 100644 index 000000000..f68237f0d --- /dev/null +++ b/benchmarks/models/benchmark_h3_sigma_schedule.py @@ -0,0 +1,317 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors + +"""Benchmark the h3_sigma_schedule operator across its backends (RFC #420). + +Operator-only microbenchmark: builds sigma grids for N in {25, 50, 1000, +1000000} x shift in {3.0, 12.0} on three paths: + +* path A — CPU reference, grid built on CPU, outputs stay on CPU; +* path B — CPU reference, grid built on CPU, sigmas and timesteps moved to + the target GPU via the operator's ``device`` argument (what a user without + the GPU backend gets: build + host-to-device copy); +* path C — Triton candidate, grid built on the device. + +B vs C is the comparison that matters when the consumer lives on the GPU; A +is an auxiliary CPU baseline and A/C ratios must not be presented as a +speedup under equal output conditions. + +Timing semantics (deliberately conservative): **per-call end-to-end host +wall latency**. Each call is timed individually with ``time.perf_counter`` +around the whole ``op()`` call — including Python dispatch, host +allocations, kernel launches, and the device path's single host +synchronization (the deduplicated-length read) — and, for the GPU paths, +one ``torch.cuda.synchronize()`` before starting each sample (so prior +async work cannot leak into the window) and one after the call (so the +call's device work is complete before the clock stops). Reported per case: +raw per-call samples, mean, median, standard deviation, min, and max. The +median is the headline number; mean and raw samples are recorded so +outliers stay visible. No CUDA events are used, because they cannot see the +host-side part of the call. Triton kernel compilation happens inside the +pre-timing correctness gate (the gate's calls are this process's first +Triton invocations): the gate's total wall time is reported once as +``correctness_gate_wall_ms`` and no later call is labeled "first" or +"compile". Timed samples start only after per-case warmup, so steady-state +numbers exclude compilation by construction. The benchmark calls the op +classes directly (registry bypassed for timing); registry-entry identity +and correctness are pinned by tests/runtime/ and tests/backends/cuda/. + +Before any timing on a GPU host, path C is verified bitwise against path A +for every measured case (sigmas, timesteps, and num_model_evals); a +mismatch aborts with exit code 2 and no performance numbers are produced. + +Exit codes: 0 on success; 2 on the pre-timing correctness gate; nonzero on +invalid arguments, and on any Triton compile, launch, or runtime failure +after the backend was confirmed available. The Triton section is skipped +only for a *confirmed* absent backend: no CUDA/HIP device +(``torch.cuda.is_available()`` false) or a confirmed missing Triton import +on a GPU host. The script runs from an uninstalled checkout (it bootstraps +the repository root onto ``sys.path``). + +The grid is tiny (4 KB at N=1000), so GPU launch overhead may exceed the +arithmetic — the numbers are what they are, and this benchmark does not +assume the device path wins. N=1000000 is a dedup stress case, not a +typical workload. No model, checkpoint, or tokenizer is loaded. + +Example: + + python benchmarks/models/benchmark_h3_sigma_schedule.py --samples 100 + CUDA_VISIBLE_DEVICES=0 python benchmarks/models/benchmark_h3_sigma_schedule.py \ + --samples 100 --json /tmp/h3_sigma_bench.json +""" + +from __future__ import annotations + +import argparse +import json +import statistics +import sys +import time +from functools import partial +from pathlib import Path + +sys.path.insert(0, str(Path(__file__).resolve().parents[2])) # uninstalled checkout + +import torch # noqa: E402 + +_SIZES = [25, 50, 1000, 1000000] +_SHIFTS = [3.0, 12.0] +_SPEEDUP_NOTE = "speedup_b_over_c = latency(B) / latency(C); > 1 means C (Triton) is faster" + + +def _positive_int(value: str) -> int: + try: + parsed = int(value) + except ValueError as exc: + raise argparse.ArgumentTypeError(f"expected an integer, got {value!r}") from exc + if parsed < 1: + raise argparse.ArgumentTypeError(f"expected a positive integer, got {value!r}") + return parsed + + +def _int32_bits(tensor: torch.Tensor) -> list[int]: + return tensor.detach().contiguous().view(torch.int32).tolist() + + +def _time_calls(call, samples: int, *, device_sync: bool) -> dict: + """Per-call end-to-end host wall latency; raw samples plus summary stats.""" + + raw_ms: list[float] = [] + for _ in range(samples): + if device_sync: + torch.cuda.synchronize() # no prior async work leaks into this window + start = time.perf_counter() + call() + if device_sync: + torch.cuda.synchronize() # the call's device work is done before stop + raw_ms.append((time.perf_counter() - start) * 1e3) + return { + "samples_raw_ms": [round(value, 6) for value in raw_ms], + "mean_ms": statistics.fmean(raw_ms), + "median_ms": statistics.median(raw_ms), + "stdev_ms": statistics.stdev(raw_ms) if len(raw_ms) > 1 else 0.0, + "min_ms": min(raw_ms), + "max_ms": max(raw_ms), + } + + +def _reference_op(): + from rl_engine.reference.minimax_h3.sigma_schedule import NativeH3SigmaScheduleOp + + return NativeH3SigmaScheduleOp() + + +def _time_reference(samples: int, warmup: int, *, device) -> dict[str, dict]: + # Path A when device is None (outputs stay on CPU); path B when device is + # "cuda" (same CPU build, then sigmas/timesteps moved by the operator). + op = _reference_op() + results = {} + for n in _SIZES: + for shift in _SHIFTS: + for _ in range(warmup): + op(n, shift=shift, device=device) + if device is not None: + torch.cuda.synchronize() + call = partial(op, n, shift=shift, device=device) + results[f"N={n},shift={shift:g}"] = _time_calls( + call, samples, device_sync=device is not None + ) + return results + + +def _verify_triton_matches_reference() -> float: + """Bitwise gate: path C must equal path A before any timing happens. + + This gate's calls are the FIRST Triton invocations of the process — the + kernels are compiled here, not later. Returns the gate's wall time in ms + (compilation + all case comparisons) so the cold-start cost stays on the + record instead of being mislabeled somewhere else. + """ + + from rl_engine.backends.cuda.model_specific.minimax_h3.sigma_schedule import ( + H3SigmaScheduleCudaOp, + ) + + cpu_op = _reference_op() + gpu_op = H3SigmaScheduleCudaOp() + start = time.perf_counter() + for n in _SIZES: + for shift in _SHIFTS: + expected = cpu_op(n, shift=shift) + actual = gpu_op(n, shift=shift) + if ( + _int32_bits(actual.sigmas.cpu()) != _int32_bits(expected.sigmas) + or _int32_bits(actual.timesteps.cpu()) != _int32_bits(expected.timesteps) + or actual.num_model_evals != expected.num_model_evals + ): + first_sigma = next( + ( + idx + for idx, (a, e) in enumerate( + zip(_int32_bits(actual.sigmas.cpu()), _int32_bits(expected.sigmas)) + ) + if a != e + ), + None, + ) + print( + f"correctness gate FAILED at N={n}, shift={shift}: Triton result " + f"differs from the CPU reference (first mismatching sigma index: " + f"{first_sigma}); refusing to benchmark" + ) + sys.exit(2) # documented gate exit code; SystemExit(str) would exit 1 + return (time.perf_counter() - start) * 1e3 + + +def _time_triton(samples: int, warmup: int) -> dict[str, dict]: + # The caller ran the correctness gate first, so every kernel has already + # been compiled and executed; no "first call" exists here to re-time and + # nothing in this function is labeled as compilation. Any failure below + # (launch, host read) propagates and fails the command. + from rl_engine.backends.cuda.model_specific.minimax_h3.sigma_schedule import ( + H3SigmaScheduleCudaOp, + ) + + op = H3SigmaScheduleCudaOp() + + results = {} + for n in _SIZES: + for shift in _SHIFTS: + for _ in range(warmup): + op(n, shift=shift) + torch.cuda.synchronize() + call = partial(op, n, shift=shift) + results[f"N={n},shift={shift:g}"] = _time_calls(call, samples, device_sync=True) + return results + + +def _summarize(report: dict) -> dict: + """Speedup B/C per case from medians and means; A medians for context.""" + + summary = {"definition": _SPEEDUP_NOTE, "cases": {}} + path_b = report.get("path_b_reference_cpu_build_then_copy_to_gpu") + path_c = report.get("path_c_triton_device_build") + if not path_b or not path_c or "skipped" in path_b or "skipped" in path_c: + return summary + for key, case_b in path_b.items(): + case_c = path_c[key] + summary["cases"][key] = { + "b_median_ms": case_b["median_ms"], + "c_median_ms": case_c["median_ms"], + "speedup_b_over_c_median": case_b["median_ms"] / case_c["median_ms"], + "speedup_b_over_c_mean": case_b["mean_ms"] / case_c["mean_ms"], + "separated_from_dispersion": abs(case_b["median_ms"] - case_c["median_ms"]) + > (case_b["stdev_ms"] + case_c["stdev_ms"]), + } + return summary + + +def main(argv: list[str] | None = None) -> int: + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--samples", type=_positive_int, default=100) + parser.add_argument("--warmup", type=_positive_int, default=5) + parser.add_argument("--json", default=None, help="optional output path") + args = parser.parse_args(argv) + + has_gpu = torch.cuda.is_available() + report: dict = { + "operator": "h3_sigma_schedule", + "timing": ( + "per-call end-to-end host wall latency; GPU paths sync before each " + "sample and after each call; Triton first-call compile excluded" + ), + "entry": "direct op construction; registry bypassed for timing (identity pinned by tests)", + "torch": torch.__version__, + "torch_cuda_runtime": torch.version.cuda, + "hip": torch.version.hip, + "triton": None, + "device": torch.cuda.get_device_name(0) if has_gpu else "cpu-only", + "n_gpu_devices": torch.cuda.device_count() if has_gpu else 0, + "samples": args.samples, + "warmup": args.warmup, + "sizes": _SIZES, + "shifts": _SHIFTS, + "stress_case_note": "N=1000000 is a dedup stress case, not a typical H3 workload", + } + + print( + f"path A (reference, cpu build, cpu outputs): {args.samples} calls x {_SIZES} x {_SHIFTS}" + ) + report["path_a_reference_cpu"] = _time_reference(args.samples, args.warmup, device=None) + + if not has_gpu: + skip = "no CUDA/HIP device (torch.cuda.is_available() is False)" + report["path_b_reference_cpu_build_then_copy_to_gpu"] = {"skipped": skip} + report["path_c_triton_device_build"] = {"skipped": skip} + print("paths B and C: skipped — no CUDA/HIP device in this process") + else: + try: + import triton + + report["triton"] = triton.__version__ + except ImportError as exc: + skip = f"triton import failed on a GPU host: {exc}" + report["path_b_reference_cpu_build_then_copy_to_gpu"] = {"skipped": skip} + report["path_c_triton_device_build"] = {"skipped": skip} + print(f"paths B and C: skipped — triton missing ({exc})") + else: + print("verifying path C against path A bitwise before timing...") + gate_ms = _verify_triton_matches_reference() + report["correctness_gate_wall_ms"] = gate_ms + print( + "correctness gate passed " + f"(includes Triton kernel compilation: {gate_ms:.1f} ms wall, " + "reported once, excluded from all timed samples)" + ) + + print(f"path B (reference, cpu build, outputs moved to gpu): {args.samples} calls") + report["path_b_reference_cpu_build_then_copy_to_gpu"] = _time_reference( + args.samples, args.warmup, device="cuda" + ) + print(f"path C (triton, device build): {args.samples} calls") + report["path_c_triton_device_build"] = _time_triton(args.samples, args.warmup) + + report["summary"] = _summarize(report) + + for section in ( + "path_a_reference_cpu", + "path_b_reference_cpu_build_then_copy_to_gpu", + "path_c_triton_device_build", + ): + for key, value in report[section].items(): + if isinstance(value, dict) and "median_ms" in value: + print(f"{section:>44s} {key:>20s}: {value['median_ms']:.4f} ms/call (median)") + for key, case in report["summary"]["cases"].items(): + separated = case["separated_from_dispersion"] + note = "separated from dispersion" if separated else "within dispersion" + print(f"speedup B/C {key:>20s}: {case['speedup_b_over_c_median']:.2f}x ({note})") + + if args.json: + with open(args.json, "w", encoding="utf-8") as handle: + json.dump(report, handle, indent=2) + print(f"wrote {args.json}") + return 0 + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/docs/.nav.yml b/docs/.nav.yml index 1037edce5..389e7f02a 100644 --- a/docs/.nav.yml +++ b/docs/.nav.yml @@ -30,6 +30,7 @@ nav: - operators/sampling.md - operators/det-gemm.md - operators/embedding.md + - operators/h3-sigma-schedule.md - Developer Guide: - contributing/README.md - Contributor Guide: contributing/contributor-guide.md diff --git a/docs/operators/h3-sigma-schedule.md b/docs/operators/h3-sigma-schedule.md new file mode 100644 index 000000000..49193fe2a --- /dev/null +++ b/docs/operators/h3-sigma-schedule.md @@ -0,0 +1,166 @@ +# H3 Sigma Schedule + +Device-side construction of the MiniMax-H3 rectified-flow sigma grid (RFC #420, +`h3_sigma_schedule`): `linspace(1, 0, N)` in fp32, the exponential shift +`shift·s / (1 + (shift−1)·s)`, `unique_consecutive` deduplication, and +`timesteps = 1 − sigmas[:-1]`. Video uses `shift=12`, audio `shift=3`; the two +grids are independent per-request inputs, never a shared step index. + +## Entry Point + +```python +from rl_engine.runtime.registry import kernel_registry + +op = kernel_registry.get_op("sigma_schedule_h3", device="cuda") +schedule = op(25, shift=12.0, device="cuda") # device= selects backend AND output +schedule.sigmas # fp32 [M], strictly decreasing, terminal exactly 0.0 +schedule.timesteps # fp32 [M-1] = 1 - sigmas[:-1] +schedule.num_model_evals # M - 1 (deduplicated length drives the count) +``` + +The semantic path resolves with provenance (`requested_backend` +`rlkernel.h3_sigma_schedule.triton` / `.reference`): + +```python +from rl_engine.runtime.registry import kernel_registry +from rl_engine.runtime.semantic_registry import OperatorRequirements + +session = kernel_registry.semantic.session() +resolution = session.resolve( + semantic_op="sigma_schedule_h3", + requested_backend="rlkernel.h3_sigma_schedule.triton", + target="rollout", + requirements=OperatorRequirements(device="cuda", dtype="float32"), +) +op = session.instantiate(resolution) +provenance = session.instance_provenance(resolution, op) +``` + +## Backends + +| Backend | Wrapper | Native symbol | Status | +| --- | --- | --- | --- | +| CUDA (Triton) | `rl_engine.backends.cuda.model_specific.minimax_h3.sigma_schedule.H3SigmaScheduleCudaOp` | 3 Triton kernels (grid / dedup scan / timesteps) | Validated on RTX 4090D (torch 2.5.1+cu121, Triton 3.1.0): bitwise-identical to the pinned upstream oracle across the full device suite | +| ROCm | — | — | Not supported: the inline PTX `div.rn.f32` is NVIDIA-only and the backend refuses construction on HIP. The registry serves the reference on ROCm. | +| PyTorch reference | `rl_engine.reference.minimax_h3.sigma_schedule.NativeH3SigmaScheduleOp` | `torch.linspace` + `unique_consecutive` | Active (byte-frozen oracle; fixture generated on Linux CPU under torch 2.14.1, reproduced on 2.4.1 and 2.5.1) | + +## Tensor Contract + +| Argument | Shape | Dtype | Requirements | +| --- | --- | --- | --- | +| `num_inference_steps` | scalar | int | `>= 2`; terminal zero included in the count | +| `shift` | scalar | python float | `> 0` (video 12.0 / audio 3.0) | +| `device` | — | — | CUDA backend requires an NVIDIA CUDA device (NVIDIA-only; refuses construction on HIP/ROCm); CPU requests use the reference | +| → `sigmas` | `[M]`, `M ≤ N` | float32 | strictly decreasing; `sigmas[0] == 1.0` and `sigmas[-1] == 0.0` exactly for released shifts | +| → `timesteps` | `[M-1]` | float32 | `1 − sigmas[:-1]`, strictly increasing from 0.0 | +| → `num_model_evals` | scalar | int | `M − 1` | + +## Dispatch Behavior + +The registry walks the CUDA candidate list +`[Triton H3, PyTorch reference]`; on ROCm/MUSA/NPU/CPU only the reference is +registered. The Triton backend raises `ImportError` at construction +when Triton is not importable, so the first-available-candidate loop skips it +at selection time (see +[Runtime Dispatch](../architecture/runtime-dispatch.md)) — forward-time errors +are never used for selection. When the Triton backend is skipped on CUDA, the +reference builds the grid on CPU and moves the outputs to the requested CUDA +device; the result stays device-resident, but the arithmetic is the CPU path. + +## Accuracy + +The acceptance oracle is the pinned fixture +`tests/reference/minimax_h3/h3_sigma_schedule_upstream_fixture.json`, +generated by executing the original upstream +`MiniMaxH3Scheduler.set_timesteps` (diffusers commit `f53d552`, blob +`ca69ae6a`); six cases N∈{2,4,25} × shift∈{3,12} are compared as fp32 bit +patterns. The device kernels reproduce the pinned CPU construction bit law: +fp32 step, two-sided fp64 accumulation with single rounding, pre-rounded +python-float scalars, correctly-rounded division via libdevice `div_rn` for +the linspace step and inline PTX `div.rn.f32` for the shift transform +(on Triton 3.0.0 / RTX 4090D the +`tl.fdiv(..., ieee_rounding=True)` flag and the `/` operator were measured +bit-identical and NOT correctly rounded, so they are not used), +`enable_fp_fusion=False` on floating-point launches, and IEEE-equality dedup +(NaN rows are kept, matching `torch.unique_consecutive`). NaN and near-zero +parity follows the torch-CPU oracle (validated on RTX 4090D, independent +audit 2026-10-16): a NaN `fl32(shift)` fills the grid with its own quieted +NaN bits — sign included, so `-NaN` yields `0xffc00000` and `+NaN` +`0x7fc00000` (host-side scalar fast path; the device kernel's `0x7fffffff` +canonicalization would otherwise differ); an infinite `fl32(shift)` makes +every element an invalid operation, filled with the x86 default quiet NaN +`0xffc00000`; a subnormal `fl32(shift)` (e.g. `1e-45` → the minimum +subnormal `2^-149`) takes the normal device path with a `+inf` first element +and subnormal quotients (upstream keeps 10 of 25 rows at `1e-45`); a shift +whose `fl32` underflows to `+0.0` (e.g. `1e-46`) makes element 0 the invalid +`+0.0/+0.0`, canonicalized in-kernel to `0xffc00000`. Odd inputs follow +upstream parity: float `N` truncates via `int()`, non-finite or +fp32-overflowing shifts yield NaN grids, extreme finite shifts can collapse +the grid (e.g. `1e30` → `[1.0, 0.0]`). + +## Performance Notes + +```bash +python benchmarks/models/benchmark_h3_sigma_schedule.py --samples 100 +``` + +Measures **per-call end-to-end host wall latency** on three paths — A: +reference built on CPU, outputs on CPU; B: reference built on CPU, sigmas and +timesteps moved to the GPU via `device="cuda"` (what a consumer without the +GPU backend gets); C: the Triton backend building on the device. Each call is +timed individually (GPU paths synchronize before each sample and after each +call); raw samples, mean/median/stdev/min/max are reported per case, and the +median is the headline number. Before any timing on a GPU host the script +verifies path C bitwise against path A for every measured case (exit 2 on +mismatch); Triton compilation happens inside that gate and its wall time is +recorded once as `correctness_gate_wall_ms`, outside all timed samples. The +script runs from an uninstalled checkout. The grid is tiny (≤4 KB for +N=1000): launch and host-sync overhead may well exceed the arithmetic — the +benchmark reports what is measured, it does not assume the device path wins. +N=1000000 is a dedup stress case, not a typical workload. + +Measured on RTX 4090D (torch 2.5.1+cu121, Triton 3.1.0, 100 samples/case): +at N=25/50/1000 the device path C is 2.5–3.2x slower than B (B/C median +0.315–0.403, every case separated from dispersion). C wins only at the N=1e6 +dedup stress case: 51.98x at shift=12 (separated), while shift=3 is 7.9x and +within dispersion, so it is not claimed. The N=1e6 CPU medians are noisy +(B stdev ≈ 44 ms at shift=12), so the 51.98x ratio passes the dispersion +check only by a small margin. See Known Limitations and the raw JSON, CUDA +test results, source hashes, and reproduction commands in the +[4090D evidence report](../../reports/experiments/h3-sigma-schedule/). + +## Tests + +```bash +python -m pytest tests/reference/minimax_h3 tests/runtime/test_h3_sigma_schedule_dispatch.py -v +python -m pytest tests/backends/cuda/test_h3_sigma_schedule_cuda.py -v # NVIDIA CUDA hardware +``` + +The device suite also carries the upstream-parity edge matrix (NaN/±inf and +extreme shifts, float-`N` truncation) against the CPU reference, which is +itself bit-identical to the executed upstream method on those inputs. + +## Known Limitations + +- Validated on an x86_64 host with NVIDIA RTX 4090D only (driver 580, torch 2.5.1+cu121, + Triton 3.1.0); other GPUs/drivers/Triton versions are not validated. On + Triton 3.0.0 the `tl.fdiv(..., ieee_rounding=True)` flag was measured NOT + correctly rounded, which is why the linspace step uses libdevice `div_rn` + and the shift transform uses inline PTX `div.rn.f32`. +- Invalid-operation NaN bits are pinned to that x86 CPU oracle (`0xffc00000`). + A live CPU reference on ARM can produce `0x7fc00000` instead for infinite, + fp32-overflowing, or underflowing shifts. Bitwise parity for those edge cases + is not portable across host architectures; the released shifts (3 and 12) + do not take these paths. NVIDIA GPUs on ARM hosts have not been validated. +- ROCm is not supported: the Triton backend refuses construction on HIP + because its division is emitted as inline NVIDIA PTX. The registry serves the + PyTorch reference on ROCm. +- Measured end-to-end latency on the 4090D: for the typical H3 sizes + (N = 25/50/1000) the device-side Triton build is 2.5–3.2x SLOWER than + building on CPU and copying to the GPU (B/C median 0.315–0.403, all separated + from dispersion). The device path wins only at the N = 1e6 dedup stress + case: 51.98x at shift=12 (separated) and 7.9x at shift=3 (within dispersion, + not claimed). These numbers are single-operator microbenchmark results on + one card and say nothing about whole-model speed. +- `N` beyond 2^31 is not supported (int32 indexing); the shift arithmetic is + internal-fp64, which is slow on consumer GPUs but irrelevant at these sizes. diff --git a/reports/experiments/h3-sigma-schedule/README.md b/reports/experiments/h3-sigma-schedule/README.md new file mode 100644 index 000000000..71327badf --- /dev/null +++ b/reports/experiments/h3-sigma-schedule/README.md @@ -0,0 +1,82 @@ +# H3 sigma schedule: RTX 4090D evidence + +These are historical operator measurements from October 16–17, 2026, +rechecked against candidate commit `73e1aac` on October 21. No new GPU run +was performed during the source-binding check. + +## Correctness and source identity + +The final NVIDIA closing run (`run-20261017-002253-9749`) produced +**34 passed, 1 skipped, 0 failures/errors**. The skipped device-guard test +requires two GPUs; this host had one. [CUDA JUnit](cuda-4090d.xml) preserves +all outcomes; its hostname attribute was removed and the XML reserialized. + +[Source provenance](source-provenance.json) records SHA256 hashes for the +tested archive, selected source files, and included results. The reference, +CUDA backend, fixture, CPU/GPU/dispatch tests, and benchmark in the final +closing snapshot are byte-identical to those files in `73e1aac`. + +The performance run (`run-20261016-235221-8055`) predates the final HIP +rejection guard. Its reference and benchmark are byte-identical to the +candidate; CUDA kernels and the scalar helper have identical ASTs. The +forward method differs only in the error text for a rejected non-CUDA +device. The final NVIDIA correctness run tested the updated CUDA file. + +Hardware/software: one NVIDIA GeForce RTX 4090D, x86_64 Linux, driver +580.76.05, Python 3.12.3, torch 2.5.1+cu121, CUDA runtime 12.1, Triton 3.1.0. + +## Comparable latency + +[Raw benchmark JSON](benchmark-4090d.json) is copied without modification. +It contains 100 samples per case and path (2,400 samples total), after +five warmup calls. Path A builds on CPU and leaves outputs there; path B +builds on CPU and copies both outputs to CUDA; path C constructs on CUDA. +Only B/C compares equal output placement. Timing is per-call host wall +latency with synchronization before and after each GPU sample; correctness +checking and compilation happen before timed samples. + +| Workload | B median (ms) | C median (ms) | B/C | +| --- | ---: | ---: | ---: | +| N=25, shift=3 | 0.04158 | 0.11552 | 0.360 | +| N=25, shift=12 | 0.04188 | 0.13293 | 0.315 | +| N=50, shift=3 | 0.04523 | 0.11281 | 0.401 | +| N=50, shift=12 | 0.04495 | 0.11351 | 0.396 | +| N=1000, shift=3 | 0.04609 | 0.11484 | 0.401 | +| N=1000, shift=12 | 0.04672 | 0.11597 | 0.403 | +| N=1000000, shift=3 | 7.91859 | 1.00316 | 7.894 | +| N=1000000, shift=12 | 48.37110 | 0.93057 | 51.980 | + +C is approximately 2.5–3.2 times slower at the typical sizes. N=1000000 +is a deduplication stress case: the shift=12 gain exceeds the benchmark's +dispersion check; shift=3 does not and is not claimed as a reliable gain. +These are single-operator results, with no whole-model speedup claim. + +## Reproduction + +From a checkout with the declared project dependencies installed, run: + +```bash +python -m pytest tests/reference/minimax_h3 tests/runtime/test_h3_sigma_schedule_dispatch.py -q +python -m pytest tests/backends/cuda/test_h3_sigma_schedule_cuda.py -v +python benchmarks/models/benchmark_h3_sigma_schedule.py --samples 100 --warmup 5 --json /tmp/h3-sigma.json +``` + +The second command and benchmark paths B/C require NVIDIA CUDA hardware +and Triton. The CPU fixture needs neither Diffusers nor network access. + +## Evidence limits + +ROCm is unsupported by this device backend, which uses inline NVIDIA PTX. +Other GPUs, ARM-hosted NVIDIA systems, multi-GPU behavior, and full-model +execution are unvalidated. Invalid-operation NaN payloads follow the pinned +x86 CPU convention; an ARM live CPU oracle may differ for extreme shifts. + +The historical broad dispatch suite also recorded two attention failures +(`test_registered_backend_resolves_with_provenance` and +`test_implementation_kind_policy_filters_without_marking_fallback`). The +recorded baseline/candidate A/B runs contained the same failures. The H3 +CUDA result above is not an all-green historical pipeline claim. + +The existing Actions workflow runs on `main` and `test`; a PR targeting +`test-h3` has no automatic coverage from it. CPU checks for this submission +are reported separately from the historical GPU evidence. diff --git a/reports/experiments/h3-sigma-schedule/benchmark-4090d.json b/reports/experiments/h3-sigma-schedule/benchmark-4090d.json new file mode 100644 index 000000000..4aee5d792 --- /dev/null +++ b/reports/experiments/h3-sigma-schedule/benchmark-4090d.json @@ -0,0 +1,2708 @@ +{ + "operator": "h3_sigma_schedule", + "timing": "per-call end-to-end host wall latency; GPU paths sync before each sample and after each call; Triton first-call compile excluded", + "entry": "direct op construction; registry bypassed for timing (identity pinned by tests)", + "torch": "2.5.1+cu121", + "torch_cuda_runtime": "12.1", + "hip": null, + "triton": "3.1.0", + "device": "NVIDIA GeForce RTX 4090 D", + "n_gpu_devices": 1, + "samples": 100, + "warmup": 5, + "sizes": [ + 25, + 50, + 1000, + 1000000 + ], + "shifts": [ + 3.0, + 12.0 + ], + "stress_case_note": "N=1000000 is a dedup stress case, not a typical H3 workload", + "path_a_reference_cpu": { + "N=25,shift=3": { + "samples_raw_ms": [ + 0.028618, + 0.025094, + 0.026226, + 0.025433, + 0.023674, + 0.025608, + 0.026569, + 0.02386, + 0.025179, + 0.025492, + 0.023644, + 0.025481, + 0.025216, + 0.02351, + 0.025377, + 0.025131, + 0.023674, + 0.025317, + 0.026468, + 0.023615, + 0.025317, + 0.026006, + 0.023726, + 0.02503, + 0.025183, + 0.023499, + 0.025332, + 0.035562, + 0.02487, + 0.024267, + 0.024058, + 0.023831, + 0.023644, + 0.023521, + 0.02354, + 0.023488, + 0.023488, + 0.023372, + 0.023488, + 0.0236, + 0.023458, + 0.02351, + 0.023525, + 0.023305, + 0.023544, + 0.023521, + 0.023507, + 0.023406, + 0.023589, + 0.023495, + 0.023577, + 0.023488, + 0.023976, + 0.025813, + 0.023443, + 0.023402, + 0.023447, + 0.02354, + 0.023533, + 0.036407, + 0.024244, + 0.023976, + 0.031028, + 0.02424, + 0.026822, + 0.024844, + 0.024289, + 0.024002, + 0.023872, + 0.024028, + 0.02367, + 0.024088, + 0.024196, + 0.02411, + 0.024099, + 0.024091, + 0.024144, + 0.024449, + 0.024144, + 0.024281, + 0.023976, + 0.024177, + 0.024013, + 0.024106, + 0.024028, + 0.02417, + 0.024132, + 0.024006, + 0.023987, + 0.025552, + 0.02617, + 0.027087, + 0.024132, + 0.024322, + 0.024095, + 0.024408, + 0.024185, + 0.038289, + 0.024766, + 0.024401 + ], + "mean_ms": 0.024801157414913177, + "median_ms": 0.024102628231048584, + "stdev_ms": 0.002426409127135147, + "min_ms": 0.023305416107177734, + "max_ms": 0.03828853368759155 + }, + "N=25,shift=12": { + "samples_raw_ms": [ + 0.025388, + 0.024661, + 0.024199, + 0.024106, + 0.024032, + 0.024095, + 0.023961, + 0.024367, + 0.02398, + 0.024121, + 0.025757, + 0.024423, + 0.024091, + 0.024192, + 0.023857, + 0.024155, + 0.023976, + 0.024021, + 0.023849, + 0.024058, + 0.024099, + 0.023987, + 0.023998, + 0.037283, + 0.024822, + 0.02462, + 0.024803, + 0.024412, + 0.024524, + 0.024233, + 0.024229, + 0.024218, + 0.024095, + 0.024423, + 0.024155, + 0.024099, + 0.024103, + 0.024009, + 0.023931, + 0.024091, + 0.024285, + 0.023969, + 0.024039, + 0.023901, + 0.024002, + 0.023946, + 0.02392, + 0.023972, + 0.02411, + 0.023957, + 0.02398, + 0.023976, + 0.024054, + 0.023913, + 0.037011, + 0.024766, + 0.024602, + 0.024263, + 0.024226, + 0.024088, + 0.024229, + 0.023842, + 0.024121, + 0.023972, + 0.023972, + 0.023838, + 0.024121, + 0.023969, + 0.023928, + 0.023842, + 0.02382, + 0.023726, + 0.024162, + 0.023812, + 0.02392, + 0.023913, + 0.023883, + 0.024166, + 0.024043, + 0.024058, + 0.023898, + 0.023834, + 0.023946, + 0.023872, + 0.03672, + 0.024613, + 0.024386, + 0.024546, + 0.024762, + 0.023898, + 0.023901, + 0.02389, + 0.024419, + 0.024404, + 0.02386, + 0.023846, + 0.024121, + 0.023954, + 0.023972, + 0.023995 + ], + "mean_ms": 0.024525783956050873, + "median_ms": 0.02405792474746704, + "stdev_ms": 0.0022290060888030773, + "min_ms": 0.02372637391090393, + "max_ms": 0.037282705307006836 + }, + "N=50,shift=3": { + "samples_raw_ms": [ + 0.025947, + 0.025101, + 0.024661, + 0.024606, + 0.02455, + 0.024606, + 0.024673, + 0.036839, + 0.025421, + 0.025164, + 0.024822, + 0.025023, + 0.024784, + 0.024747, + 0.024967, + 0.024844, + 0.024505, + 0.024699, + 0.024404, + 0.024617, + 0.024956, + 0.024602, + 0.024684, + 0.02487, + 0.024471, + 0.024453, + 0.02487, + 0.024457, + 0.024442, + 0.024512, + 0.024468, + 0.024408, + 0.036255, + 0.025496, + 0.025243, + 0.024822, + 0.025112, + 0.024613, + 0.024874, + 0.024792, + 0.024606, + 0.024371, + 0.024535, + 0.024799, + 0.024509, + 0.024378, + 0.024728, + 0.024702, + 0.024658, + 0.024483, + 0.024658, + 0.024471, + 0.024445, + 0.024695, + 0.02443, + 0.024527, + 0.024464, + 0.033382, + 0.025444, + 0.025127, + 0.025041, + 0.025168, + 0.024941, + 0.024669, + 0.024837, + 0.024758, + 0.024613, + 0.024803, + 0.024851, + 0.024743, + 0.024695, + 0.024971, + 0.024784, + 0.024423, + 0.024602, + 0.024281, + 0.024375, + 0.024654, + 0.024442, + 0.024363, + 0.024334, + 0.024475, + 0.033513, + 0.02531, + 0.025067, + 0.024855, + 0.024986, + 0.024788, + 0.025094, + 0.024743, + 0.024412, + 0.024762, + 0.024635, + 0.024609, + 0.024617, + 0.024393, + 0.024624, + 0.024479, + 0.024516, + 0.024669 + ], + "mean_ms": 0.02513684332370758, + "median_ms": 0.02467818558216095, + "stdev_ms": 0.0020672659325772047, + "min_ms": 0.024281442165374756, + "max_ms": 0.03683939576148987 + }, + "N=50,shift=12": { + "samples_raw_ms": [ + 0.025507, + 0.025053, + 0.034623, + 0.024945, + 0.025239, + 0.0249, + 0.024877, + 0.024587, + 0.02471, + 0.024609, + 0.024449, + 0.024572, + 0.024423, + 0.024483, + 0.024386, + 0.024352, + 0.024814, + 0.024494, + 0.024386, + 0.024557, + 0.024375, + 0.02424, + 0.02465, + 0.024427, + 0.024345, + 0.024412, + 0.033643, + 0.025369, + 0.024982, + 0.025064, + 0.02468, + 0.024609, + 0.024684, + 0.024606, + 0.024769, + 0.024535, + 0.024226, + 0.024479, + 0.024423, + 0.024501, + 0.024319, + 0.024557, + 0.024334, + 0.024281, + 0.02436, + 0.024278, + 0.029784, + 0.024673, + 0.02433, + 0.024505, + 0.024188, + 0.033673, + 0.025056, + 0.024986, + 0.024863, + 0.024833, + 0.024632, + 0.024889, + 0.024669, + 0.024728, + 0.024464, + 0.024475, + 0.024538, + 0.024542, + 0.02443, + 0.024498, + 0.024278, + 0.024423, + 0.024747, + 0.024505, + 0.024382, + 0.024412, + 0.024658, + 0.024538, + 0.024412, + 0.024434, + 0.03399, + 0.025112, + 0.025056, + 0.024687, + 0.024687, + 0.024639, + 0.024803, + 0.02481, + 0.024755, + 0.024468, + 0.024639, + 0.024535, + 0.024568, + 0.024445, + 0.024382, + 0.024736, + 0.024382, + 0.024695, + 0.024617, + 0.024494, + 0.024471, + 0.02468, + 0.024494, + 0.024483 + ], + "mean_ms": 0.025032833218574524, + "median_ms": 0.02456270158290863, + "stdev_ms": 0.0019250050694447909, + "min_ms": 0.024188309907913208, + "max_ms": 0.03462284803390503 + }, + "N=1000,shift=3": { + "samples_raw_ms": [ + 0.028893, + 0.028472, + 0.028215, + 0.028256, + 0.027999, + 0.02791, + 0.028025, + 0.027794, + 0.028312, + 0.027798, + 0.027884, + 0.027508, + 0.027683, + 0.02807, + 0.027679, + 0.027474, + 0.027861, + 0.027318, + 0.027314, + 0.027675, + 0.027582, + 0.027664, + 0.027575, + 0.027567, + 0.027411, + 0.027366, + 0.027332, + 0.027604, + 0.027325, + 0.027463, + 0.028174, + 0.027921, + 0.031084, + 0.03067, + 0.030499, + 0.030592, + 0.03022, + 0.03029, + 0.030555, + 0.030432, + 0.030145, + 0.030112, + 0.030436, + 0.030242, + 0.030268, + 0.030234, + 0.030369, + 0.030302, + 0.030532, + 0.030573, + 0.030246, + 0.030492, + 0.030085, + 0.030193, + 0.030119, + 0.030152, + 0.0301, + 0.029918, + 0.029784, + 0.029817, + 0.029765, + 0.029717, + 0.029605, + 0.029858, + 0.030577, + 0.030473, + 0.030484, + 0.030324, + 0.030421, + 0.030421, + 0.03038, + 0.030398, + 0.030287, + 0.030376, + 0.030465, + 0.030439, + 0.030506, + 0.033755, + 0.031125, + 0.038955, + 0.030927, + 0.03054, + 0.030622, + 0.030473, + 0.030816, + 0.03032, + 0.030536, + 0.030469, + 0.03048, + 0.030454, + 0.030421, + 0.030398, + 0.030439, + 0.03067, + 0.030447, + 0.030454, + 0.030573, + 0.030413, + 0.030473, + 0.030365 + ], + "mean_ms": 0.029662102460861206, + "median_ms": 0.030256807804107666, + "stdev_ms": 0.001616210717547429, + "min_ms": 0.027313828468322754, + "max_ms": 0.038955360651016235 + }, + "N=1000,shift=12": { + "samples_raw_ms": [ + 0.030871, + 0.03054, + 0.030603, + 0.030383, + 0.030085, + 0.030488, + 0.030041, + 0.030193, + 0.030138, + 0.030339, + 0.02997, + 0.03019, + 0.030085, + 0.030223, + 0.030126, + 0.030238, + 0.030152, + 0.030112, + 0.030044, + 0.030063, + 0.030175, + 0.030238, + 0.030152, + 0.030126, + 0.030126, + 0.030193, + 0.030216, + 0.030138, + 0.030067, + 0.03019, + 0.030145, + 0.030216, + 0.030037, + 0.03022, + 0.030152, + 0.030119, + 0.030048, + 0.030171, + 0.030167, + 0.030138, + 0.030056, + 0.030268, + 0.030704, + 0.030771, + 0.030614, + 0.030767, + 0.030756, + 0.030756, + 0.030793, + 0.030808, + 0.030715, + 0.030681, + 0.030614, + 0.030737, + 0.030696, + 0.030704, + 0.030622, + 0.030551, + 0.0307, + 0.030596, + 0.030603, + 0.030689, + 0.030596, + 0.030767, + 0.030596, + 0.030551, + 0.030838, + 0.0307, + 0.030845, + 0.030689, + 0.030581, + 0.030708, + 0.030559, + 0.030749, + 0.030551, + 0.030234, + 0.030112, + 0.030208, + 0.030082, + 0.030041, + 0.030026, + 0.030197, + 0.030238, + 0.030268, + 0.030182, + 0.030134, + 0.030093, + 0.030085, + 0.030063, + 0.030302, + 0.030182, + 0.030231, + 0.030261, + 0.030421, + 0.030279, + 0.030383, + 0.03013, + 0.030138, + 0.030171, + 0.030227 + ], + "mean_ms": 0.030355416238307953, + "median_ms": 0.030232593417167664, + "stdev_ms": 0.00026597119961990735, + "min_ms": 0.0299699604511261, + "max_ms": 0.030871480703353882 + }, + "N=1000000,shift=3": { + "samples_raw_ms": [ + 98.036628, + 8.730412, + 93.520977, + 9.203587, + 97.394422, + 93.301896, + 9.268787, + 96.477598, + 94.346248, + 6.778911, + 92.860915, + 6.882805, + 96.026655, + 6.864298, + 95.050234, + 6.832108, + 93.638495, + 94.85922, + 6.765481, + 95.271535, + 6.557032, + 94.918385, + 94.463106, + 7.024731, + 95.403858, + 6.742172, + 92.721418, + 6.972056, + 94.761349, + 6.573778, + 95.462278, + 94.235215, + 7.376682, + 95.555358, + 6.778695, + 92.313088, + 6.678641, + 95.238689, + 6.066728, + 91.907859, + 6.779984, + 93.20775, + 7.191882, + 93.075141, + 6.610051, + 93.257032, + 6.699257, + 93.414858, + 6.857354, + 94.965789, + 6.682359, + 99.277865, + 95.631044, + 94.418861, + 7.148165, + 95.317729, + 97.422332, + 7.461969, + 94.627522, + 6.66241, + 95.223997, + 6.799322, + 95.014967, + 94.363861, + 7.39412, + 95.063124, + 7.123426, + 92.349082, + 6.768007, + 93.703873, + 6.953746, + 95.613036, + 8.824546, + 95.512893, + 95.289852, + 6.379906, + 95.462427, + 6.266601, + 94.293647, + 94.966922, + 6.477915, + 93.44428, + 6.584033, + 95.159143, + 7.033534, + 95.8784, + 6.378505, + 94.424304, + 95.119867, + 7.265992, + 96.060917, + 6.968558, + 95.030703, + 97.042311, + 8.166842, + 93.609802, + 99.965513, + 95.552105, + 7.990137, + 96.968003 + ], + "mean_ms": 58.05059902369976, + "median_ms": 93.27946417033672, + "stdev_ms": 43.607488298814054, + "min_ms": 6.066728383302689, + "max_ms": 99.96551275253296 + }, + "N=1000000,shift=12": { + "samples_raw_ms": [ + 6.64027, + 100.297894, + 93.073688, + 6.873686, + 93.18278, + 6.77117, + 94.959993, + 6.888814, + 95.084712, + 94.147202, + 7.517833, + 94.432641, + 6.975438, + 96.075386, + 6.638095, + 188.340988, + 7.480063, + 94.538264, + 6.921433, + 96.149787, + 95.841412, + 9.611689, + 95.688105, + 95.287181, + 6.816771, + 93.938541, + 6.749317, + 97.705033, + 94.681785, + 6.856363, + 94.323538, + 17.187905, + 87.005608, + 94.767582, + 6.828763, + 96.779134, + 6.679688, + 95.006149, + 94.833337, + 9.139318, + 95.59337, + 94.859775, + 7.134549, + 96.737124, + 6.490875, + 98.469645, + 91.793936, + 6.359782, + 95.519081, + 6.945428, + 94.569005, + 6.170601, + 95.248934, + 94.516166, + 6.834071, + 95.688961, + 6.512713, + 91.937918, + 6.760281, + 95.568947, + 7.04392, + 95.335443, + 95.046874, + 6.781477, + 95.643755, + 6.677359, + 95.740765, + 94.966419, + 7.003978, + 95.238145, + 7.038787, + 95.555391, + 95.020026, + 7.163744, + 95.524229, + 6.524671, + 95.992979, + 99.751059, + 96.447758, + 7.191516, + 95.296003, + 94.574295, + 7.302854, + 95.558546, + 6.635245, + 95.679991, + 6.654218, + 90.029359, + 7.00159, + 95.267821, + 6.61581, + 95.024548, + 95.202584, + 6.915841, + 95.877212, + 6.268453, + 93.390584, + 7.578775, + 95.402561, + 95.191807 + ], + "mean_ms": 59.99584909528494, + "median_ms": 94.23536993563175, + "stdev_ms": 45.21561488384047, + "min_ms": 6.170600652694702, + "max_ms": 188.34098801016808 + } + }, + "correctness_gate_wall_ms": 1041.1223843693733, + "path_b_reference_cpu_build_then_copy_to_gpu": { + "N=25,shift=3": { + "samples_raw_ms": [ + 0.051226, + 0.046548, + 0.043381, + 0.042167, + 0.041589, + 0.04198, + 0.047166, + 0.041559, + 0.042368, + 0.041533, + 0.041269, + 0.041626, + 0.042211, + 0.041243, + 0.041258, + 0.0415, + 0.041343, + 0.041123, + 0.042085, + 0.041217, + 0.041407, + 0.041112, + 0.040948, + 0.041362, + 0.041109, + 0.041496, + 0.04109, + 0.041015, + 0.041183, + 0.04144, + 0.041001, + 0.041164, + 0.041183, + 0.041686, + 0.041168, + 0.040319, + 0.040885, + 0.040639, + 0.04087, + 0.040367, + 0.04134, + 0.040863, + 0.041027, + 0.040784, + 0.041071, + 0.04087, + 0.04087, + 0.040822, + 0.040714, + 0.040371, + 0.040576, + 0.040401, + 0.040423, + 0.046521, + 0.041556, + 0.041787, + 0.041597, + 0.041544, + 0.041671, + 0.0416, + 0.041712, + 0.041492, + 0.041936, + 0.041272, + 0.046469, + 0.042014, + 0.041645, + 0.041246, + 0.042159, + 0.041198, + 0.041578, + 0.041712, + 0.042267, + 0.042021, + 0.042539, + 0.041801, + 0.041932, + 0.042342, + 0.042159, + 0.042044, + 0.041779, + 0.041436, + 0.041962, + 0.04166, + 0.041712, + 0.041705, + 0.041805, + 0.041779, + 0.041857, + 0.042263, + 0.041887, + 0.041477, + 0.041805, + 0.042472, + 0.042148, + 0.042368, + 0.041869, + 0.041697, + 0.042092, + 0.042733 + ], + "mean_ms": 0.041822902858257294, + "median_ms": 0.041583552956581116, + "stdev_ms": 0.0014984721829570426, + "min_ms": 0.040318816900253296, + "max_ms": 0.05122646689414978 + }, + "N=25,shift=12": { + "samples_raw_ms": [ + 0.044722, + 0.043213, + 0.042688, + 0.042889, + 0.042114, + 0.041831, + 0.042126, + 0.042085, + 0.042044, + 0.04195, + 0.041399, + 0.041775, + 0.041209, + 0.041146, + 0.041313, + 0.042174, + 0.041869, + 0.042014, + 0.042252, + 0.042044, + 0.041645, + 0.041835, + 0.041664, + 0.042602, + 0.041567, + 0.041809, + 0.041351, + 0.041861, + 0.041474, + 0.041414, + 0.046808, + 0.042133, + 0.041649, + 0.041604, + 0.04185, + 0.041679, + 0.041354, + 0.041425, + 0.04153, + 0.041377, + 0.041731, + 0.042435, + 0.041392, + 0.041347, + 0.041191, + 0.041094, + 0.041448, + 0.04185, + 0.041965, + 0.041924, + 0.041693, + 0.041723, + 0.041399, + 0.041798, + 0.041611, + 0.045709, + 0.042472, + 0.041727, + 0.041459, + 0.041686, + 0.041556, + 0.04195, + 0.041496, + 0.041708, + 0.041641, + 0.041619, + 0.041887, + 0.041928, + 0.041716, + 0.041872, + 0.041604, + 0.042412, + 0.0416, + 0.041991, + 0.042025, + 0.041556, + 0.041503, + 0.043716, + 0.044368, + 0.044845, + 0.044726, + 0.044703, + 0.044484, + 0.045061, + 0.044562, + 0.044867, + 0.044718, + 0.044897, + 0.044692, + 0.044737, + 0.044417, + 0.044566, + 0.044327, + 0.045042, + 0.044469, + 0.04486, + 0.044517, + 0.044897, + 0.044391, + 0.045542 + ], + "mean_ms": 0.04256613552570343, + "median_ms": 0.04187971353530884, + "stdev_ms": 0.0013992164613653418, + "min_ms": 0.041093677282333374, + "max_ms": 0.04680827260017395 + }, + "N=50,shift=3": { + "samples_raw_ms": [ + 0.04784, + 0.046276, + 0.045851, + 0.046272, + 0.046026, + 0.045475, + 0.046663, + 0.04565, + 0.04565, + 0.046074, + 0.045445, + 0.046071, + 0.045497, + 0.044905, + 0.045463, + 0.045221, + 0.04518, + 0.044998, + 0.045128, + 0.045661, + 0.04565, + 0.04461, + 0.044979, + 0.044968, + 0.045661, + 0.044752, + 0.044897, + 0.045013, + 0.045449, + 0.045169, + 0.045374, + 0.045359, + 0.045229, + 0.045445, + 0.044733, + 0.048399, + 0.045829, + 0.045091, + 0.045396, + 0.045393, + 0.044707, + 0.045083, + 0.044703, + 0.045709, + 0.045709, + 0.046048, + 0.04568, + 0.045136, + 0.044927, + 0.045255, + 0.044964, + 0.045061, + 0.045296, + 0.044886, + 0.044797, + 0.045188, + 0.044703, + 0.044804, + 0.044797, + 0.045232, + 0.045452, + 0.045329, + 0.045121, + 0.045251, + 0.044979, + 0.045236, + 0.045013, + 0.045087, + 0.045829, + 0.045039, + 0.044908, + 0.045009, + 0.045124, + 0.044979, + 0.044648, + 0.045057, + 0.045355, + 0.045583, + 0.044908, + 0.044953, + 0.045076, + 0.044897, + 0.044804, + 0.045042, + 0.045691, + 0.04581, + 0.084665, + 0.047699, + 0.046015, + 0.045773, + 0.045352, + 0.045523, + 0.045996, + 0.045385, + 0.045098, + 0.045016, + 0.045571, + 0.044923, + 0.045028, + 0.044886 + ], + "mean_ms": 0.04576537758111954, + "median_ms": 0.04523061215877533, + "stdev_ms": 0.003977969846116417, + "min_ms": 0.04461035132408142, + "max_ms": 0.08466467261314392 + }, + "N=50,shift=12": { + "samples_raw_ms": [ + 0.046805, + 0.046898, + 0.045687, + 0.046402, + 0.0454, + 0.045277, + 0.045173, + 0.045381, + 0.045292, + 0.04578, + 0.045374, + 0.045929, + 0.045616, + 0.045039, + 0.045158, + 0.04518, + 0.044603, + 0.051565, + 0.045385, + 0.04638, + 0.045519, + 0.04518, + 0.045255, + 0.045285, + 0.045121, + 0.045013, + 0.044964, + 0.04524, + 0.044454, + 0.045311, + 0.045016, + 0.045151, + 0.044633, + 0.044994, + 0.04505, + 0.045657, + 0.045221, + 0.044879, + 0.044808, + 0.044733, + 0.045199, + 0.045128, + 0.044946, + 0.04499, + 0.044748, + 0.044849, + 0.045296, + 0.045005, + 0.044722, + 0.044908, + 0.044946, + 0.045322, + 0.045583, + 0.044659, + 0.044819, + 0.045046, + 0.045039, + 0.044685, + 0.044972, + 0.051372, + 0.045281, + 0.044998, + 0.045106, + 0.044528, + 0.044692, + 0.044748, + 0.044338, + 0.041921, + 0.042036, + 0.042185, + 0.042316, + 0.042628, + 0.042632, + 0.043016, + 0.042509, + 0.042923, + 0.042848, + 0.074446, + 0.043917, + 0.042643, + 0.042345, + 0.042684, + 0.042535, + 0.04252, + 0.042733, + 0.041902, + 0.042301, + 0.042286, + 0.042282, + 0.042338, + 0.041891, + 0.042271, + 0.04261, + 0.042185, + 0.042215, + 0.041705, + 0.041615, + 0.042263, + 0.041857, + 0.042338 + ], + "mean_ms": 0.0447063148021698, + "median_ms": 0.04494562745094299, + "stdev_ms": 0.003458302957096138, + "min_ms": 0.04161521792411804, + "max_ms": 0.07444620132446289 + }, + "N=1000,shift=3": { + "samples_raw_ms": [ + 0.050128, + 0.048701, + 0.048608, + 0.047378, + 0.047337, + 0.047266, + 0.047173, + 0.047013, + 0.0473, + 0.047129, + 0.04904, + 0.046127, + 0.046212, + 0.046019, + 0.046235, + 0.045858, + 0.04632, + 0.046272, + 0.046324, + 0.046074, + 0.046015, + 0.045795, + 0.046138, + 0.046127, + 0.045497, + 0.046, + 0.046611, + 0.045843, + 0.046223, + 0.045892, + 0.046041, + 0.04581, + 0.051867, + 0.046246, + 0.046302, + 0.046369, + 0.046056, + 0.045966, + 0.045788, + 0.045877, + 0.045571, + 0.046194, + 0.04575, + 0.046097, + 0.046156, + 0.04537, + 0.045642, + 0.046033, + 0.046179, + 0.046209, + 0.046041, + 0.045888, + 0.046074, + 0.046004, + 0.045836, + 0.045441, + 0.045925, + 0.045806, + 0.045661, + 0.04594, + 0.045691, + 0.045698, + 0.046156, + 0.045929, + 0.046112, + 0.046276, + 0.046134, + 0.04603, + 0.046201, + 0.046141, + 0.046451, + 0.045817, + 0.045821, + 0.045903, + 0.046339, + 0.04578, + 0.045728, + 0.046041, + 0.045527, + 0.045925, + 0.045672, + 0.04565, + 0.046089, + 0.045892, + 0.045672, + 0.04578, + 0.045687, + 0.045691, + 0.046313, + 0.046819, + 0.046976, + 0.046916, + 0.04673, + 0.046857, + 0.046507, + 0.046927, + 0.046473, + 0.046577, + 0.046786, + 0.046719 + ], + "mean_ms": 0.04633225500583649, + "median_ms": 0.04609301686286926, + "stdev_ms": 0.0009310644898782102, + "min_ms": 0.04537031054496765, + "max_ms": 0.05186721682548523 + }, + "N=1000,shift=12": { + "samples_raw_ms": [ + 0.048593, + 0.047561, + 0.046931, + 0.047106, + 0.047754, + 0.054777, + 0.047673, + 0.046801, + 0.047021, + 0.047088, + 0.047039, + 0.046827, + 0.047166, + 0.046715, + 0.046529, + 0.046879, + 0.046868, + 0.046689, + 0.047192, + 0.046834, + 0.046488, + 0.046488, + 0.046592, + 0.046443, + 0.046514, + 0.046637, + 0.04657, + 0.046831, + 0.046659, + 0.046521, + 0.046339, + 0.046916, + 0.046726, + 0.046682, + 0.046644, + 0.046778, + 0.046343, + 0.046041, + 0.046536, + 0.046689, + 0.046406, + 0.046264, + 0.046767, + 0.046436, + 0.046682, + 0.046752, + 0.046331, + 0.046533, + 0.046633, + 0.046376, + 0.046853, + 0.04622, + 0.046417, + 0.046711, + 0.047121, + 0.046827, + 0.046324, + 0.046555, + 0.046708, + 0.04692, + 0.046596, + 0.046186, + 0.046629, + 0.046968, + 0.048578, + 0.046618, + 0.046752, + 0.04676, + 0.046659, + 0.046499, + 0.047114, + 0.046827, + 0.04689, + 0.046611, + 0.046775, + 0.046823, + 0.047028, + 0.04667, + 0.046492, + 0.046775, + 0.047103, + 0.046782, + 0.046574, + 0.046831, + 0.050016, + 0.047192, + 0.047095, + 0.046931, + 0.046793, + 0.04717, + 0.046656, + 0.04698, + 0.046633, + 0.0466, + 0.046637, + 0.046488, + 0.046581, + 0.046846, + 0.046466, + 0.046737 + ], + "mean_ms": 0.04688676446676254, + "median_ms": 0.046720728278160095, + "stdev_ms": 0.0009437531693891809, + "min_ms": 0.046040862798690796, + "max_ms": 0.054776668548583984 + }, + "N=1000000,shift=3": { + "samples_raw_ms": [ + 5.909961, + 5.378716, + 92.432741, + 5.453605, + 93.805902, + 4.140414, + 93.927026, + 5.54869, + 4.009236, + 93.959201, + 4.926607, + 92.965908, + 5.785037, + 93.763903, + 3.998835, + 93.091145, + 5.636867, + 3.856339, + 94.48389, + 4.664458, + 3.497154, + 94.027806, + 5.124994, + 92.996407, + 6.021373, + 93.4713, + 4.106585, + 94.282214, + 7.681195, + 91.760397, + 6.96284, + 94.414458, + 5.055442, + 6.595064, + 90.845264, + 4.717436, + 95.677175, + 93.781177, + 100.194272, + 8.155979, + 94.301883, + 4.566252, + 94.976161, + 6.205395, + 4.482169, + 5.961504, + 86.18037, + 4.866268, + 95.735755, + 5.587362, + 93.37955, + 6.820064, + 94.416, + 4.625402, + 95.494561, + 93.911625, + 5.221579, + 94.621304, + 6.408781, + 4.700273, + 93.274202, + 5.950678, + 94.100304, + 95.60075, + 6.723024, + 4.732061, + 95.200874, + 94.833065, + 5.302772, + 95.854856, + 6.415736, + 92.477214, + 6.514691, + 94.009709, + 4.591063, + 95.086969, + 6.260347, + 93.097046, + 7.315353, + 94.315913, + 98.918114, + 7.38354, + 94.045345, + 4.801679, + 95.142819, + 6.185293, + 92.060506, + 7.359698, + 94.141532, + 4.977733, + 95.471524, + 12.729816, + 88.717606, + 94.859857, + 7.076494, + 93.199823, + 7.155597, + 5.510379, + 91.828998, + 6.551679 + ], + "mean_ms": 48.093439266085625, + "median_ms": 7.918586954474449, + "stdev_ms": 44.30672562663738, + "min_ms": 3.4971535205841064, + "max_ms": 100.19427165389061 + }, + "N=1000000,shift=12": { + "samples_raw_ms": [ + 95.096972, + 5.005348, + 94.998203, + 6.540809, + 93.503602, + 6.706927, + 94.361749, + 4.820429, + 94.38327, + 6.078709, + 92.825301, + 7.319123, + 94.303988, + 4.480623, + 94.582248, + 6.044071, + 4.382923, + 93.941677, + 93.684979, + 6.067503, + 95.580429, + 5.902249, + 4.594665, + 94.222538, + 93.457267, + 4.940972, + 5.95488, + 94.276771, + 4.79243, + 95.473621, + 5.280193, + 93.536034, + 6.25319, + 93.93701, + 4.542962, + 94.970919, + 5.409606, + 93.62283, + 7.063594, + 96.615449, + 4.467163, + 95.063824, + 94.277587, + 4.935823, + 94.028231, + 5.816832, + 4.548877, + 97.71299, + 5.387995, + 87.853402, + 6.311495, + 93.871828, + 7.169869, + 96.491978, + 5.440231, + 94.174277, + 6.260362, + 93.757551, + 4.406735, + 97.565111, + 5.394496, + 92.624117, + 5.987786, + 93.953013, + 4.81841, + 95.019691, + 5.749509, + 91.080904, + 6.572045, + 93.732946, + 5.133312, + 94.456311, + 5.632665, + 4.72495, + 94.40209, + 5.289201, + 93.607001, + 94.887368, + 6.529916, + 94.201077, + 6.654497, + 5.414564, + 92.984322, + 94.0901, + 8.888789, + 96.398532, + 6.19087, + 94.228171, + 4.379727, + 95.124967, + 5.628973, + 93.718838, + 7.226855, + 93.808323, + 4.500203, + 93.121968, + 5.658548, + 92.840414, + 6.946221, + 93.162011 + ], + "mean_ms": 49.93861924856901, + "median_ms": 48.3710952103138, + "stdev_ms": 44.49452530771314, + "min_ms": 4.379726946353912, + "max_ms": 97.71298989653587 + } + }, + "path_c_triton_device_build": { + "N=25,shift=3": { + "samples_raw_ms": [ + 0.155285, + 0.145249, + 0.145938, + 0.14814, + 0.144746, + 0.154782, + 0.157494, + 0.146277, + 0.141907, + 0.144202, + 0.141993, + 0.142422, + 0.149243, + 0.155918, + 0.14548, + 0.142518, + 0.147298, + 0.127181, + 0.125401, + 0.124663, + 0.121403, + 0.119124, + 0.122614, + 0.139251, + 0.122216, + 0.117227, + 0.116564, + 0.125062, + 0.117298, + 0.116095, + 0.115246, + 0.116419, + 0.117179, + 0.115782, + 0.11497, + 0.11481, + 0.114217, + 0.113681, + 0.117671, + 0.116438, + 0.116643, + 0.116639, + 0.114638, + 0.115581, + 0.122376, + 0.11915, + 0.117455, + 0.116892, + 0.114888, + 0.112247, + 0.115719, + 0.115462, + 0.112291, + 0.113927, + 0.11282, + 0.118829, + 0.115339, + 0.113405, + 0.113182, + 0.113372, + 0.113517, + 0.113256, + 0.111878, + 0.113152, + 0.113931, + 0.111245, + 0.117503, + 0.113338, + 0.111233, + 0.114672, + 0.114784, + 0.118479, + 0.114903, + 0.114799, + 0.114955, + 0.112664, + 0.111289, + 0.113085, + 0.114579, + 0.112664, + 0.114165, + 0.11174, + 0.113674, + 0.113394, + 0.112686, + 0.113007, + 0.114143, + 0.113208, + 0.114847, + 0.116624, + 0.116061, + 0.115935, + 0.117764, + 0.115313, + 0.114907, + 0.114597, + 0.115171, + 0.113029, + 0.113584, + 0.112869 + ], + "mean_ms": 0.12132901698350906, + "median_ms": 0.11552125215530396, + "stdev_ms": 0.01271071310115929, + "min_ms": 0.11123344302177429, + "max_ms": 0.15749409794807434 + }, + "N=25,shift=12": { + "samples_raw_ms": [ + 0.151448, + 0.146642, + 0.145812, + 0.141289, + 0.138737, + 0.148252, + 0.142615, + 0.138938, + 0.140391, + 0.140619, + 0.138797, + 0.149731, + 0.137251, + 0.137094, + 0.137046, + 0.139013, + 0.135299, + 0.136029, + 0.133689, + 0.134237, + 0.132434, + 0.134874, + 0.133228, + 0.134178, + 0.134792, + 0.133984, + 0.133853, + 0.138108, + 0.134494, + 0.137489, + 0.13461, + 0.131417, + 0.131514, + 0.133123, + 0.135351, + 0.134997, + 0.135619, + 0.134159, + 0.132114, + 0.143316, + 0.137232, + 0.132475, + 0.132732, + 0.136334, + 0.13575, + 0.135846, + 0.134196, + 0.134803, + 0.134494, + 0.138767, + 0.13886, + 0.135593, + 0.134863, + 0.137564, + 0.135053, + 0.143025, + 0.116829, + 0.113569, + 0.113465, + 0.113185, + 0.112783, + 0.112329, + 0.113562, + 0.112358, + 0.114009, + 0.114281, + 0.112321, + 0.111464, + 0.120807, + 0.112921, + 0.110328, + 0.112142, + 0.113502, + 0.113368, + 0.115659, + 0.113454, + 0.113964, + 0.114772, + 0.113662, + 0.115097, + 0.113808, + 0.117321, + 0.113957, + 0.11304, + 0.11418, + 0.112157, + 0.112899, + 0.113096, + 0.113368, + 0.113633, + 0.111606, + 0.112277, + 0.11291, + 0.113346, + 0.114202, + 0.115864, + 0.111572, + 0.115123, + 0.113554, + 0.11304 + ], + "mean_ms": 0.12684956192970276, + "median_ms": 0.13292767107486725, + "stdev_ms": 0.012286099921595904, + "min_ms": 0.11032819747924805, + "max_ms": 0.15144795179367065 + }, + "N=50,shift=3": { + "samples_raw_ms": [ + 0.127301, + 0.120342, + 0.120047, + 0.117846, + 0.117783, + 0.119328, + 0.119749, + 0.115916, + 0.114933, + 0.117399, + 0.116017, + 0.114854, + 0.114333, + 0.113685, + 0.114154, + 0.112347, + 0.114679, + 0.115175, + 0.114497, + 0.113137, + 0.114221, + 0.114605, + 0.112802, + 0.113424, + 0.119157, + 0.115998, + 0.112757, + 0.122033, + 0.114638, + 0.115413, + 0.112839, + 0.112675, + 0.114694, + 0.113662, + 0.112984, + 0.110947, + 0.11554, + 0.112884, + 0.11155, + 0.112437, + 0.112817, + 0.113506, + 0.113823, + 0.111211, + 0.110533, + 0.111662, + 0.111971, + 0.111416, + 0.112291, + 0.11453, + 0.114013, + 0.113983, + 0.11174, + 0.110816, + 0.112724, + 0.110108, + 0.110112, + 0.109378, + 0.109635, + 0.110898, + 0.11367, + 0.111502, + 0.111699, + 0.111602, + 0.111699, + 0.115257, + 0.113007, + 0.112075, + 0.111487, + 0.112005, + 0.112422, + 0.110731, + 0.11288, + 0.113171, + 0.112958, + 0.110734, + 0.111807, + 0.111111, + 0.114508, + 0.111118, + 0.111438, + 0.110812, + 0.112273, + 0.113849, + 0.112217, + 0.111535, + 0.111923, + 0.109475, + 0.110365, + 0.109959, + 0.111338, + 0.111867, + 0.110742, + 0.111494, + 0.132043, + 0.1157, + 0.112273, + 0.113484, + 0.112485, + 0.111494 + ], + "mean_ms": 0.1135815680027008, + "median_ms": 0.11280924081802368, + "stdev_ms": 0.003421334969553872, + "min_ms": 0.10937824845314026, + "max_ms": 0.13204291462898254 + }, + "N=50,shift=12": { + "samples_raw_ms": [ + 0.123005, + 0.117045, + 0.118066, + 0.117175, + 0.116773, + 0.1195, + 0.117112, + 0.121471, + 0.117552, + 0.121746, + 0.119243, + 0.116345, + 0.11459, + 0.113774, + 0.11408, + 0.115313, + 0.113286, + 0.115585, + 0.114433, + 0.126291, + 0.116166, + 0.113435, + 0.114262, + 0.113744, + 0.114039, + 0.113808, + 0.114132, + 0.113413, + 0.114538, + 0.111565, + 0.113081, + 0.111975, + 0.11155, + 0.111762, + 0.111386, + 0.111114, + 0.111707, + 0.111748, + 0.112865, + 0.112813, + 0.113118, + 0.11351, + 0.114508, + 0.112981, + 0.113558, + 0.112329, + 0.112697, + 0.11399, + 0.114117, + 0.112817, + 0.11351, + 0.113666, + 0.124104, + 0.115287, + 0.114273, + 0.110622, + 0.116132, + 0.111785, + 0.113364, + 0.115026, + 0.115164, + 0.11263, + 0.11399, + 0.11326, + 0.11443, + 0.114124, + 0.113111, + 0.113349, + 0.11415, + 0.112407, + 0.114482, + 0.112459, + 0.113912, + 0.113122, + 0.112876, + 0.111587, + 0.11383, + 0.111889, + 0.113498, + 0.111569, + 0.112128, + 0.111386, + 0.11244, + 0.111286, + 0.111613, + 0.111569, + 0.116903, + 0.112165, + 0.11364, + 0.113562, + 0.115614, + 0.112735, + 0.111137, + 0.110015, + 0.112057, + 0.111979, + 0.113875, + 0.112563, + 0.111703, + 0.111122 + ], + "mean_ms": 0.11408209800720215, + "median_ms": 0.11350959539413452, + "stdev_ms": 0.002837490376946739, + "min_ms": 0.11001527309417725, + "max_ms": 0.12629106640815735 + }, + "N=1000,shift=3": { + "samples_raw_ms": [ + 0.124503, + 0.119358, + 0.11678, + 0.114378, + 0.116825, + 0.121158, + 0.123974, + 0.122245, + 0.119444, + 0.11801, + 0.115279, + 0.127267, + 0.120673, + 0.116654, + 0.116844, + 0.116274, + 0.115469, + 0.116307, + 0.120327, + 0.116773, + 0.116404, + 0.116616, + 0.117715, + 0.116672, + 0.117328, + 0.116311, + 0.116572, + 0.114828, + 0.11538, + 0.119612, + 0.116579, + 0.117064, + 0.117086, + 0.117764, + 0.114758, + 0.115562, + 0.117801, + 0.115629, + 0.114851, + 0.11437, + 0.113413, + 0.11513, + 0.114311, + 0.12454, + 0.118393, + 0.115171, + 0.116363, + 0.118326, + 0.114385, + 0.112984, + 0.114094, + 0.113245, + 0.11332, + 0.112522, + 0.111848, + 0.112008, + 0.112087, + 0.112031, + 0.11551, + 0.112504, + 0.111353, + 0.112817, + 0.11237, + 0.114582, + 0.115391, + 0.11516, + 0.119209, + 0.112794, + 0.11377, + 0.122625, + 0.116087, + 0.113651, + 0.113565, + 0.112861, + 0.110269, + 0.112019, + 0.118174, + 0.113782, + 0.11152, + 0.111528, + 0.112522, + 0.112314, + 0.112101, + 0.113349, + 0.112463, + 0.113677, + 0.114202, + 0.11462, + 0.114705, + 0.114132, + 0.113431, + 0.112914, + 0.112675, + 0.112817, + 0.113573, + 0.114415, + 0.114609, + 0.11497, + 0.114422, + 0.113081 + ], + "mean_ms": 0.11554151773452759, + "median_ms": 0.11483952403068542, + "stdev_ms": 0.003185604418510141, + "min_ms": 0.11026859283447266, + "max_ms": 0.12726709246635437 + }, + "N=1000,shift=12": { + "samples_raw_ms": [ + 0.124302, + 0.128299, + 0.124693, + 0.127688, + 0.199676, + 0.120614, + 0.119023, + 0.116929, + 0.118993, + 0.117302, + 0.115868, + 0.116505, + 0.114031, + 0.11272, + 0.114989, + 0.114296, + 0.114512, + 0.114225, + 0.114746, + 0.114206, + 0.115432, + 0.119224, + 0.116285, + 0.116367, + 0.115406, + 0.113763, + 0.116974, + 0.153415, + 0.12337, + 0.116806, + 0.11785, + 0.116702, + 0.115078, + 0.137202, + 0.118595, + 0.116348, + 0.115193, + 0.114128, + 0.114843, + 0.115193, + 0.115007, + 0.113666, + 0.115313, + 0.113938, + 0.113741, + 0.11396, + 0.114553, + 0.113048, + 0.114072, + 0.11358, + 0.114698, + 0.115022, + 0.116371, + 0.116032, + 0.155821, + 0.118699, + 0.116408, + 0.117786, + 0.118155, + 0.11719, + 0.117883, + 0.115935, + 0.118557, + 0.11763, + 0.11792, + 0.124704, + 0.117902, + 0.117891, + 0.118371, + 0.118021, + 0.116583, + 0.114623, + 0.11484, + 0.11519, + 0.115316, + 0.115961, + 0.115495, + 0.115216, + 0.116587, + 0.117257, + 0.116695, + 0.114694, + 0.115152, + 0.114415, + 0.117175, + 0.115566, + 0.115976, + 0.115272, + 0.116032, + 0.115797, + 0.114392, + 0.114568, + 0.114474, + 0.114966, + 0.114676, + 0.11602, + 0.113849, + 0.119939, + 0.114884, + 0.221096 + ], + "mean_ms": 0.1193637028336525, + "median_ms": 0.11596828699111938, + "stdev_ms": 0.014619546948211822, + "min_ms": 0.1127198338508606, + "max_ms": 0.2210959792137146 + }, + "N=1000000,shift=3": { + "samples_raw_ms": [ + 1.011424, + 1.021761, + 1.027875, + 1.011312, + 1.003955, + 1.006901, + 1.011718, + 1.002684, + 1.018606, + 1.006205, + 1.013059, + 1.004808, + 1.007494, + 1.008116, + 1.004089, + 1.002021, + 1.001995, + 1.005597, + 1.003798, + 1.002349, + 1.001589, + 1.003049, + 1.004931, + 1.005035, + 1.003128, + 1.003951, + 1.000587, + 1.001064, + 1.003038, + 1.003418, + 1.003422, + 1.000844, + 0.99951, + 1.003146, + 1.003016, + 1.007352, + 1.010727, + 1.007643, + 1.004528, + 1.003016, + 1.003105, + 1.00363, + 1.003299, + 1.001835, + 1.003936, + 1.002222, + 1.002409, + 1.001131, + 1.002453, + 1.003176, + 1.002159, + 1.001637, + 1.001984, + 1.001995, + 1.003928, + 1.0009, + 1.001604, + 1.002364, + 1.001753, + 1.003206, + 1.003519, + 1.007866, + 1.007207, + 1.00185, + 1.002032, + 1.001116, + 1.003575, + 1.004055, + 1.004644, + 1.003701, + 1.001924, + 1.000058, + 1.000959, + 1.002003, + 1.002315, + 1.000766, + 1.002207, + 1.00388, + 1.00264, + 1.003563, + 1.000747, + 1.001306, + 1.001503, + 1.001757, + 1.001492, + 1.001563, + 1.002159, + 1.001731, + 1.010127, + 1.033016, + 1.029123, + 1.023781, + 1.023207, + 1.032684, + 1.025092, + 1.024444, + 1.025222, + 1.003843, + 1.001365, + 1.001809 + ], + "mean_ms": 1.0060333460569382, + "median_ms": 1.003161072731018, + "stdev_ms": 0.007624140766711798, + "min_ms": 0.9995102882385254, + "max_ms": 1.0330155491828918 + }, + "N=1000000,shift=12": { + "samples_raw_ms": [ + 1.010083, + 1.006849, + 1.006499, + 1.005478, + 1.005758, + 1.006301, + 1.015104, + 1.030091, + 1.029763, + 1.027219, + 1.029611, + 1.027722, + 1.026507, + 1.02558, + 1.020867, + 1.008432, + 1.007777, + 1.003552, + 1.003243, + 0.955857, + 0.940233, + 0.930272, + 0.932276, + 0.933845, + 0.93215, + 0.928815, + 0.930589, + 0.931319, + 0.931215, + 0.928909, + 0.930414, + 0.931349, + 0.930544, + 0.927761, + 0.929087, + 0.927754, + 0.926659, + 0.931311, + 0.931047, + 0.929467, + 0.93301, + 0.929397, + 0.927184, + 0.929143, + 0.929829, + 0.929281, + 0.930861, + 0.926059, + 0.935603, + 0.929657, + 0.929531, + 0.926126, + 0.92667, + 0.927754, + 0.927828, + 0.929594, + 0.926983, + 0.931788, + 0.929631, + 0.926953, + 0.926614, + 0.9313, + 0.929434, + 0.927225, + 0.92712, + 0.92705, + 0.931915, + 0.930324, + 0.928361, + 0.927716, + 0.931736, + 0.928953, + 0.927351, + 0.926185, + 0.931945, + 0.927564, + 0.928517, + 0.928134, + 0.929844, + 0.930544, + 0.929352, + 0.932917, + 1.011267, + 1.009226, + 0.983991, + 0.937138, + 0.937056, + 0.935983, + 0.934776, + 0.924367, + 0.930279, + 0.931066, + 0.933498, + 0.931684, + 0.929028, + 0.925358, + 0.929277, + 0.927821, + 0.928801, + 0.936691 + ], + "mean_ms": 0.9486759826540947, + "median_ms": 0.9305663406848907, + "stdev_ms": 0.035319901142147496, + "min_ms": 0.9243674576282501, + "max_ms": 1.0300911962985992 + } + }, + "summary": { + "definition": "speedup_b_over_c = latency(B) / latency(C); > 1 means C (Triton) is faster", + "cases": { + "N=25,shift=3": { + "b_median_ms": 0.041583552956581116, + "c_median_ms": 0.11552125215530396, + "speedup_b_over_c_median": 0.35996452757175107, + "speedup_b_over_c_mean": 0.3447065170233913, + "separated_from_dispersion": true + }, + "N=25,shift=12": { + "b_median_ms": 0.04187971353530884, + "c_median_ms": 0.13292767107486725, + "speedup_b_over_c_median": 0.3150564001961746, + "speedup_b_over_c_mean": 0.3355639142789681, + "separated_from_dispersion": true + }, + "N=50,shift=3": { + "b_median_ms": 0.04523061215877533, + "c_median_ms": 0.11280924081802368, + "speedup_b_over_c_median": 0.40094775774387426, + "speedup_b_over_c_mean": 0.40292961600980276, + "separated_from_dispersion": true + }, + "N=50,shift=12": { + "b_median_ms": 0.04494562745094299, + "c_median_ms": 0.11350959539413452, + "speedup_b_over_c_median": 0.3959632425336396, + "speedup_b_over_c_mean": 0.39187844178100084, + "separated_from_dispersion": true + }, + "N=1000,shift=3": { + "b_median_ms": 0.04609301686286926, + "c_median_ms": 0.11483952403068542, + "speedup_b_over_c_median": 0.4013689298342362, + "speedup_b_over_c_mean": 0.401000920831689, + "separated_from_dispersion": true + }, + "N=1000,shift=12": { + "b_median_ms": 0.046720728278160095, + "c_median_ms": 0.11596828699111938, + "speedup_b_over_c_median": 0.4028750401541921, + "speedup_b_over_c_mean": 0.3928058811320961, + "separated_from_dispersion": true + }, + "N=1000000,shift=3": { + "b_median_ms": 7.918586954474449, + "c_median_ms": 1.003161072731018, + "speedup_b_over_c_median": 7.893634601387383, + "speedup_b_over_c_mean": 47.80501506692969, + "separated_from_dispersion": false + }, + "N=1000000,shift=12": { + "b_median_ms": 48.3710952103138, + "c_median_ms": 0.9305663406848907, + "speedup_b_over_c_median": 51.98027598409909, + "speedup_b_over_c_mean": 52.640332591594216, + "separated_from_dispersion": true + } + } + } +} \ No newline at end of file diff --git a/reports/experiments/h3-sigma-schedule/cuda-4090d.xml b/reports/experiments/h3-sigma-schedule/cuda-4090d.xml new file mode 100644 index 000000000..a1de68e48 --- /dev/null +++ b/reports/experiments/h3-sigma-schedule/cuda-4090d.xml @@ -0,0 +1,2 @@ + +/root/autodl-tmp/h3-nvidia-closing/repo/tests/backends/cuda/test_h3_sigma_schedule_cuda.py:352: multi-GPU device-guard test needs 2 devices diff --git a/reports/experiments/h3-sigma-schedule/source-provenance.json b/reports/experiments/h3-sigma-schedule/source-provenance.json new file mode 100644 index 000000000..e86f3b4d9 --- /dev/null +++ b/reports/experiments/h3-sigma-schedule/source-provenance.json @@ -0,0 +1,61 @@ +{ + "evidence_kind": "Historical hardware results; source binding rechecked 2026-10-21; no new GPU run", + "candidate_commit": "73e1aac95a6c3e0745792b83733ba79432c12b3b", + "historical_snapshot_base_commit": "0e8a18aff441895082aee603959cd25cb66d7e4b", + "environment": { + "python": "3.12.3", + "python_impl": "CPython", + "platform": "Linux-5.15.0-94-generic-x86_64-with-glibc2.35", + "torch": "2.5.1+cu121", + "torch_cuda_runtime": "12.1", + "torch_hip": null, + "cuda_available": true, + "cuda_device_count": 1, + "triton": "3.1.0", + "cuda_device_name": "NVIDIA GeForce RTX 4090 D", + "cuda_capability": "8.9" + }, + "nvidia_driver": "580.76.05", + "cuda_validation": { + "run": "run-20261017-002253-9749", + "passed": 34, + "skipped": 1, + "failures": 0, + "errors": 0, + "skip_reason": "multi-GPU device guard requires two GPUs; host had one", + "snapshot_sha256": "5d7ef73a86d8b1ab2086c6b09e45ce0799484fcac6a0d5964ed6372b02592fe5", + "source_files_sha256": { + "rl_engine/reference/minimax_h3/sigma_schedule.py": "18b36cdbac9e95d4b52782aa17a7228fdf35bc7e3397508a651c1fb07cb36509", + "rl_engine/backends/cuda/model_specific/minimax_h3/sigma_schedule.py": "784fb85f2e3dbce450fd5564c99584eb75a83c8b01374d938f036c7036494262", + "tests/reference/minimax_h3/test_h3_sigma_schedule.py": "1304b4738196250bd09c5532f39dd15b14e1ff46910ef4a807067c5a8b1faeee", + "tests/reference/minimax_h3/h3_sigma_schedule_upstream_fixture.json": "97a929a39eb781ce00bd351d8c30c85d9bfd054e7ca5405dc3a3b8ed794956c1", + "tests/runtime/test_h3_sigma_schedule_dispatch.py": "1750573eb6f609490e637d4088ca7f7ab6dfe01c6077d68cf1428b2a239c8e0e", + "tests/backends/cuda/test_h3_sigma_schedule_cuda.py": "c9ce44388a4fb31e0415026836c74c2818514ec995229d94315462c2894ae72a", + "benchmarks/models/benchmark_h3_sigma_schedule.py": "04975954e863050d602479d737b5218b88e03f4cc87aca7ca90cb47ebf0f3486" + }, + "junit_original_sha256": "9e2d476cdaf24043e66f239e9f5d5acbb58a2949bb474b4662ca09b8a96e7240", + "junit_redaction": "Removed testsuite hostname attribute; XML reserialized. All test outcomes preserved." + }, + "benchmark": { + "run": "run-20261016-235221-8055", + "snapshot_sha256": "0b7c25c6a9bf4e109fe0e053a90fed1a977831d859f2bf99575cfcfe08da9b34", + "binding": "Benchmark and reference byte-identical. CUDA kernels and scalar helper ASTs identical. Forward AST identical after normalizing only the non-CUDA device error text; final constructor adds HIP rejection.", + "cuda_function_asts_identical": [ + "_shift_scalars", + "_shifted_grid_kernel", + "_unique_consecutive_kernel", + "_timesteps_kernel", + "__call__" + ], + "forward_ast_equal_after_normalizing_non_cuda_error": true + }, + "included_files_sha256": { + "benchmark-4090d.json": "84ba715d03e6ada96c122f7cff3992d81e569acb4c56b810bb191fc84f64e6e9", + "cuda-4090d.xml": "c5377af475a7020c02f0c1784bd87df6412ca4f8cfc42f4687dd8f8fac1900e2" + }, + "limits": [ + "One x86_64 NVIDIA RTX 4090D; no ROCm device backend.", + "No fresh GPU measurement, multi-GPU test, or full-model run.", + "Historical broad dispatch suite had two attention failures also present in the recorded baseline A/B comparison; not an all-green pipeline claim." + ] +} diff --git a/rl_engine/backends/cuda/model_specific/__init__.py b/rl_engine/backends/cuda/model_specific/__init__.py new file mode 100644 index 000000000..56bc59826 --- /dev/null +++ b/rl_engine/backends/cuda/model_specific/__init__.py @@ -0,0 +1,4 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors + +"""Triton CUDA/HIP backend for model-specific operators (RFC #420 MiniMax-H3).""" diff --git a/rl_engine/backends/cuda/model_specific/minimax_h3/__init__.py b/rl_engine/backends/cuda/model_specific/minimax_h3/__init__.py new file mode 100644 index 000000000..bd451c764 --- /dev/null +++ b/rl_engine/backends/cuda/model_specific/minimax_h3/__init__.py @@ -0,0 +1,8 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors + +"""MiniMax-H3 (RFC #420) model-specific CUDA/HIP backends.""" + +from rl_engine.backends.cuda.model_specific.minimax_h3.sigma_schedule import H3SigmaScheduleCudaOp + +__all__ = ["H3SigmaScheduleCudaOp"] diff --git a/rl_engine/backends/cuda/model_specific/minimax_h3/sigma_schedule.py b/rl_engine/backends/cuda/model_specific/minimax_h3/sigma_schedule.py new file mode 100644 index 000000000..8f7997267 --- /dev/null +++ b/rl_engine/backends/cuda/model_specific/minimax_h3/sigma_schedule.py @@ -0,0 +1,331 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors + +"""Triton CUDA backend for the MiniMax-H3 sigma grid (RFC #420). + +NVIDIA only: the grid-transform division is emitted as inline PTX +``div.rn.f32`` (see below), which makes the kernels NVIDIA-specific. HIP is +not supported by this backend; the registry does not offer it on rocm. + +The grid is built **on the device**: no CPU construction followed by a copy, +and no torch factory functions whose GPU kernels differ bitwise from the CPU +reference. The bit-level target is the pinned CPU fixture in +``tests/reference/minimax_h3/`` (upstream diffusers ``set_timesteps`` at commit +``f53d552``), reproduced here by: + +* ``linspace`` law: ``step = fl32(-1 / (n-1))``; for ``idx < n//2`` the value + is ``fl32(1.0 + fp64(step) * fp64(idx))``, otherwise + ``fl32(0.0 - fp64(step) * fp64(n-1-idx))`` (two-sided, single rounding); +* shift arithmetic with the scalar roundings torch applies to python-float + operands: ``fl32(shift)`` and ``fl32(shift - 1.0)`` computed first, then + mul / mul+add / div — each fp32 step evaluated through an exact fp64 + intermediate rounded once to fp32 (bit-identical to the correctly rounded, + non-FTZ fp32 op that torch's CPU kernels perform; fp32 x fp32 products and + sums are exact in fp64). Plain Triton fp32 mul/add may compile + flush-to-zero, which destroys the subnormal numerators of subnormal shifts + and broke parity (independent audit 2026-10-16); +* correctly-rounded, non-FTZ fp32 division emitted as inline PTX + ``div.rn.f32``. Two measured findings forced this: (1) on Triton 3.0.0 / + RTX 4090D ``tl.fdiv(..., ieee_rounding=True)`` and the ``/`` operator were + bit-identical to each other and NOT correctly rounded (64707 mismatched + elements vs the CPU division across the fingerprint grids — probe + artifacts/.../diag/div_probe.py); (2) on Triton 3.1.0 the libdevice + ``div_rn`` binding flushed subnormal DIVIDENDS — ``div_rn(2^-149, +0.0)`` + returned NaN instead of +inf and subnormal/normal returned +0, collapsing + the ``1e-45`` grid (probe artifacts/.../diag/subnormal_probe.py). The bare + PTX instruction carries no ``.ftz`` modifier, so IEEE subnormal semantics + hold; ``enable_fp_fusion=False`` on every floating-point launch, so no + multiply-add is contracted into an FMA (fused and separate roundings + differ on real grids, e.g. N=50, shift=12); +* ``unique_consecutive`` semantics: keep an element when it is not IEEE-equal + to its predecessor (NaN never equals, so NaN rows are kept, matching torch); +* ``timesteps = fl32(1.0 - sigma)`` for the first ``M - 1`` sigmas, with NaN + sigmas passed through unchanged (torch CPU returns the quiet NaN operand + itself from ``1.0 - NaN``, sign and payload included). + +NaN and near-zero parity (independent GPU audit 2026-10-16): + +* A NaN ``fl32(shift)`` propagates its own quieted bits through the whole + grid on the CPU oracle — sign included (``+NaN`` -> ``0x7fc00000``, + ``-NaN`` -> ``0xffc00000``; quieting = setting bit 22). ``forward`` + detects NaN/inf scalars on the host and fills the grid directly with the + oracle's pattern; the device kernel canonicalizes NaNs to ``0x7fffffff``, + which is why the fill exists. +* An infinite ``fl32(shift)`` (``inf``, or a finite shift above the fp32 + range such as ``1e39`` / ``3.5e38``) makes every element an invalid + operation (``inf/inf``, or ``inf*0`` at the terminal zero), which x86/SSE + evaluates to the default quiet NaN ``0xffc00000`` regardless of sign. +* A subnormal ``fl32(shift)`` (e.g. ``1e-45`` rounds up to the minimum + subnormal ``2^-149``) takes the normal kernel path: element 0 is + ``2^-149 / +0.0 = +inf`` and the remaining elements are subnormal + quotients — upstream keeps e.g. 10 of 25 rows for ``1e-45``. +* A shift whose ``fl32`` underflows to ``+0.0`` (e.g. ``1e-46``) makes + element 0 the invalid ``+0.0 / +0.0``: the kernel canonicalizes NaN + division results to the CPU's ``0xffc00000`` (finite shifts CAN produce + exactly this one NaN — the denominator ``1 + fl32(shift-1)*s`` is exactly + ``+0.0`` at ``s = 1`` when ``fl32(shift)`` is tiny). + +Dedup runs as one single-program kernel (sequential block scan), so the +compaction order is deterministic for any ``N``. The output length ``M`` is +read back with one host synchronization of a device counter — the only sync; +sigma values never leave the device. + +Availability: constructing the op raises ``ImportError`` when Triton is not +importable, so the registry's first-available-candidate selection skips this +backend at construction time instead of failing later in ``forward``. +""" + +from __future__ import annotations + +import math + +import torch + +try: + import triton + import triton.language as tl + + # Correctly-rounded fp32 division (PTX div.rn.f32). The import path moved + # across Triton versions; both known layouts are accepted. + try: + from triton.language.extra import libdevice as _libdevice + except ImportError: # pragma: no cover - alternative Triton layout + from triton.language.extra.cuda import libdevice as _libdevice + + _TRITON_AVAILABLE = True +except ImportError: # pragma: no cover - exercised only without a Triton wheel + _TRITON_AVAILABLE = False + +from rl_engine.reference.minimax_h3.sigma_schedule import H3SigmaSchedule + +_BLOCK = 1024 + + +def _shift_scalars(shift: float) -> tuple[float, float]: + """Round the two python-float scalars exactly as torch would. + + ``fl32(shift)`` and ``fl32(shift - 1.0)`` with the subtraction in fp64 + first — pinned on the CPU device so an ambient default device (meta, + CUDA) cannot move the factory or add a synchronization. These are the + only host-side roundings; the grid itself is built by the kernels. + """ + + shift_f32 = torch.tensor(shift, dtype=torch.float32, device="cpu").item() + shift_m1_f32 = torch.tensor(shift - 1.0, dtype=torch.float32, device="cpu").item() + return shift_f32, shift_m1_f32 + + +if _TRITON_AVAILABLE: + + @triton.jit + def _shifted_grid_kernel( + out_ptr, + n, + shift_f32, + shift_m1_f32, + BLOCK: tl.constexpr, + ): + offs = tl.program_id(0) * BLOCK + tl.arange(0, BLOCK) + mask = offs < n + + # Base grid, pinned CPU linspace law: fp32 step, two-sided fp64 + # accumulate with a single rounding to fp32. The step division uses + # div_rn like the main transform (tl.fdiv's ieee_rounding flag is not + # correctly rounded on the validated Triton 3.0.0, see module docs). + step32 = _libdevice.div_rn(-1.0, (n - 1).to(tl.float32)) + step64 = step32.to(tl.float64) + half = n // 2 + start_side = 1.0 + step64 * offs.to(tl.float64) + end_side = 0.0 - step64 * (n - 1 - offs).to(tl.float64) + base = tl.where(offs < half, start_side, end_side).to(tl.float32) + + # Shift transform with pre-rounded scalars. Every fp32 step is + # computed through an exact fp64 intermediate rounded once to fp32: + # an fp32 x fp32 product or sum is exactly representable in fp64, so + # this is bit-identical to the correctly rounded (round-to-nearest, + # NON-FTZ) fp32 op that torch's CPU (SSE) kernels perform. Plain + # Triton fp32 mul/add may compile flush-to-zero, which zeros the + # subnormal numerators of subnormal shifts (fl32(1e-45) = 2^-149) and + # broke upstream parity (independent audit 2026-10-16). + base64 = base.to(tl.float64) + numerator = (shift_f32.to(tl.float64) * base64).to(tl.float32) + inner = (shift_m1_f32.to(tl.float64) * base64).to(tl.float32) + denominator = (1.0 + inner.to(tl.float64)).to(tl.float32) + + # Division via inline PTX div.rn.f32: correctly rounded AND non-FTZ + # (no .ftz modifier — IEEE subnormal handling, nonzero/0 -> inf). + # The libdevice div_rn binding was measured on Triton 3.1.0 / RTX + # 4090D to flush subnormal DIVIDENDS: div_rn(2^-149, +0.0) returned + # NaN instead of +inf and subnormal/normal returned +0, collapsing + # the 1e-45 grid (probe: remote/diag/subnormal_probe.py). This makes + # the kernel NVIDIA-PTX-specific — see the module docstring. + value = tl.inline_asm_elementwise( + "div.rn.f32 $0, $1, $2;", + "=r,r,r", + [numerator, denominator], + dtype=tl.float32, + is_pure=True, + pack=1, + ) + + # NaN canonicalization for finite fl32(shift): the only NaN source is + # the invalid 0/+0 at element 0 — a shift whose fl32 underflows to + # +0.0 (e.g. 1e-46) makes (shift-1)*1 round to -1.0 and the + # denominator exactly +0.0 there while the numerator is +0.0. torch + # CPU evaluates 0/+0 to the x86 default quiet NaN 0xffc00000 (sign + # bit SET); the device division yields 0x7fffffff. Rewrite NaN + # results to the CPU bit pattern. (A subnormal fl32(shift) such as + # 1e-45 takes numerator 2^-149 / +0.0 = +inf — a valid IEEE result — + # and flows through untouched.) + value_int = value.to(tl.int32, bitcast=True) + cpu_qnan = tl.full([BLOCK], -4194304, tl.int32) # 0xffc00000 signed + out_int = tl.where(value != value, cpu_qnan, value_int) + tl.store(out_ptr + offs, out_int.to(tl.float32, bitcast=True), mask=mask) + + @triton.jit + def _unique_consecutive_kernel( + src_ptr, + dst_ptr, + count_ptr, + n, + BLOCK: tl.constexpr, + ): + # Single program: sequential over blocks, so the scan order — and + # therefore the compaction result — is deterministic for any n. + running = tl.zeros((1,), dtype=tl.int32) + for start in range(0, n, BLOCK): + offs = start + tl.arange(0, BLOCK) + mask = offs < n + cur = tl.load(src_ptr + offs, mask=mask, other=0.0) + # offs == 0 sees NaN as "previous" and is always kept; IEEE + # equality keeps NaN rows (NaN != NaN), matching torch. + prev = tl.load(src_ptr + offs - 1, mask=mask & (offs > 0), other=float("nan")) + keep = mask & (cur != prev) + keep_i32 = keep.to(tl.int32) + dest = running + tl.cumsum(keep_i32, axis=0) - keep_i32 + tl.store(dst_ptr + dest, cur, mask=keep) + running = running + tl.sum(keep_i32, axis=0) + tl.store(count_ptr + tl.arange(0, 1), running) + + @triton.jit + def _timesteps_kernel(sig_ptr, out_ptr, m, BLOCK: tl.constexpr): + offs = tl.program_id(0) * BLOCK + tl.arange(0, BLOCK) + mask = offs < m - 1 + sigma = tl.load(sig_ptr + offs, mask=mask, other=0.0) + # 1 - sigma, except NaN sigmas pass through unchanged: torch CPU + # (SSE) returns the quiet NaN operand itself from 1.0 - NaN — payload + # and sign included — while the device subtract would canonicalize to + # 0x7fffffff. The grid's NaNs are already in CPU bit form, so copying + # sigma reproduces the CPU result for both NaN signs. + timesteps = tl.where(sigma != sigma, sigma, 1.0 - sigma) + tl.store(out_ptr + offs, timesteps, mask=mask) + + +class H3SigmaScheduleCudaOp: + """Device-side H3 sigma grid via Triton (NVIDIA CUDA only).""" + + def __init__(self) -> None: + # Platform gate first: the kernels emit inline PTX (div.rn.f32) and + # cannot run on AMD. Raising ImportError — like the Triton check + # below — keeps the registry's candidate-skip semantics + # (docs/architecture/runtime-dispatch.md) intact if this class ever + # appears in a HIP device table; construction fails here, long + # before any Triton compilation is attempted. + if torch.version.hip is not None: + raise ImportError( + "H3SigmaScheduleCudaOp is NVIDIA-only: its kernels emit inline " + "PTX div.rn.f32 for exact IEEE division and cannot run on " + "HIP/ROCm. Use the PyTorch reference " + "(rl_engine.reference.minimax_h3.sigma_schedule) on AMD." + ) + if not _TRITON_AVAILABLE: + # Construction-time failure lets the registry's candidate loop + # skip this backend (docs/architecture/runtime-dispatch.md); + # forward-time errors are never used for selection. + raise ImportError("triton is required for H3SigmaScheduleCudaOp") + + def __call__( + self, + num_inference_steps: int, + *, + shift: float, + device: str | torch.device | None = None, + ) -> H3SigmaSchedule: + return self.forward(num_inference_steps, shift=shift, device=device) + + def forward( + self, + num_inference_steps: int, + *, + shift: float, + device: str | torch.device | None = None, + ) -> H3SigmaSchedule: + """Build the grid on the target device; see the module docstring.""" + # Same checks and messages as the CPU reference (upstream's own + # failures; no extra policy). + if num_inference_steps < 2: + raise ValueError( + "num_inference_steps must be >= 2 (terminal zero included), got " + f"{num_inference_steps}." + ) + if shift <= 0: + raise ValueError(f"`shift` must be positive, got {shift}.") + + target = torch.device(device) if device is not None else torch.device("cuda") + if target.type != "cuda": + raise ValueError( + "H3SigmaScheduleCudaOp builds the grid on an NVIDIA CUDA device " + "(the kernels are NVIDIA-only); pass device='cuda' or use the " + "PyTorch reference for CPU or AMD platforms." + ) + + n = int(num_inference_steps) + shift_f32, shift_m1_f32 = _shift_scalars(shift) + + # NaN/inf scalar fast path (module docstring "NaN payloads"). A NaN + # fl32(shift) propagates its own quieted bits through the whole grid + # on the CPU oracle — SIGN INCLUDED, so -NaN yields 0xffc00000 while + # +NaN yields 0x7fc00000 (quieting = setting bit 22). An infinite + # fl32(shift) makes every element an invalid operation (inf/inf, or + # inf*0 at the terminal zero), which x86 evaluates to the default + # quiet NaN 0xffc00000 regardless of the shift's sign. The device + # kernel canonicalizes NaNs to 0x7fffffff instead, so these grids are + # filled directly with the oracle's bit patterns. All NaN rows + # survive unique_consecutive (IEEE !=), so M = n. + if math.isnan(shift_f32) or math.isinf(shift_f32): + if math.isinf(shift_f32): + payload = 0xFFC00000 + else: + raw = int( + torch.tensor(shift, dtype=torch.float32, device="cpu").view(torch.int32).item() + ) + payload = raw & 0xFFFFFFFF | 0x00400000 # quiet the NaN bits + signed = payload if payload < 0x80000000 else payload - 0x1_0000_0000 + with torch.cuda.device(target): + bits = torch.full((n,), signed, dtype=torch.int32, device=target) + sigmas = bits.view(torch.float32) + timesteps = sigmas[: n - 1] + return H3SigmaSchedule(sigmas=sigmas, timesteps=timesteps, num_model_evals=n - 1) + + # Device guard: allocations, Triton launches (which use the current + # device/stream) and the count read must all target `target`, also + # when it is not the current device (cuda:k on multi-GPU hosts). + with torch.cuda.device(target): + grid = torch.empty(n, dtype=torch.float32, device=target) + compacted = torch.empty(n, dtype=torch.float32, device=target) + count = torch.empty(1, dtype=torch.int32, device=target) + + blocks = triton.cdiv(n, _BLOCK) + _shifted_grid_kernel[(blocks,)]( + grid, n, shift_f32, shift_m1_f32, BLOCK=_BLOCK, enable_fp_fusion=False + ) + _unique_consecutive_kernel[(1,)](grid, compacted, count, n, BLOCK=_BLOCK) + # The single allowed host synchronization: the deduplicated length. + m = int(count.item()) + sigmas = compacted[:m] + timesteps = torch.empty(m - 1, dtype=torch.float32, device=target) + if m > 1: + _timesteps_kernel[(triton.cdiv(m - 1, _BLOCK),)]( + compacted, timesteps, m, BLOCK=_BLOCK, enable_fp_fusion=False + ) + return H3SigmaSchedule(sigmas=sigmas, timesteps=timesteps, num_model_evals=m - 1) diff --git a/rl_engine/reference/minimax_h3/__init__.py b/rl_engine/reference/minimax_h3/__init__.py new file mode 100644 index 000000000..75e1157c9 --- /dev/null +++ b/rl_engine/reference/minimax_h3/__init__.py @@ -0,0 +1,6 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors + +from rl_engine.reference.minimax_h3.sigma_schedule import H3SigmaSchedule, NativeH3SigmaScheduleOp + +__all__ = ["H3SigmaSchedule", "NativeH3SigmaScheduleOp"] diff --git a/rl_engine/reference/minimax_h3/sigma_schedule.py b/rl_engine/reference/minimax_h3/sigma_schedule.py new file mode 100644 index 000000000..f2a22ed90 --- /dev/null +++ b/rl_engine/reference/minimax_h3/sigma_schedule.py @@ -0,0 +1,102 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors + +"""FP32 reference for the MiniMax-H3 sigma grid (RFC #420 ``h3_sigma_schedule``). + +Upstream: ``MiniMaxH3Scheduler.set_timesteps`` in diffusers +``schedulers/scheduling_minimax_h3.py``, pinned to git blob sha +``ca69ae6a2f6aae5bb68498216316a71eb07c1f7b``. The file was added once, in +commit f53d552036a0d1bd5570782a39cd40cfabf112bc ("Add MiniMax-H3 (#14355)", +2026-08-05), and has not changed since: +https://github.com/huggingface/diffusers/blob/f53d552/src/diffusers/schedulers/scheduling_minimax_h3.py +The construction order is byte-frozen against that source: + + linspace(1, 0, N, float32) # CPU + -> shift * s / (1 + (shift - 1) * s) # python-float scalar operands + -> unique_consecutive # collapse fp32 collisions near s=1 + -> timesteps = 1 - sigmas[:-1] # t = 1 - sigma, t = 1 is clean + +Released checkpoints run one grid per modality — shift=12.0 for video, +shift=3.0 for audio — so the grid is a per-call argument, not scheduler state. +Every arithmetic step runs on CPU in float32; tensors move to ``device`` only +after the grid is fully built. Consumers take sigma as a per-row input +(PR #471 ``h3_ode_step``), so this operator carries no step index. +""" + +from __future__ import annotations + +from dataclasses import dataclass + +import torch +from torch import Tensor + + +@dataclass(frozen=True) +class H3SigmaSchedule: + """One MiniMax-H3 sigma grid. + + Attributes: + sigmas: float32, shape ``[M]`` with ``M <= num_inference_steps``. + Strictly decreasing; ``sigmas[0] == 1.0`` exactly for the released + shifts (3.0, 12.0) and ``sigmas[-1] == 0.0`` exactly (the shift + maps 0 to 0). + timesteps: float32, shape ``[M - 1]``; ``timesteps = 1 - sigmas[:-1]``, + strictly increasing from ``0.0`` towards ``1.0``. + num_model_evals: ``M - 1``, the number of transformer evaluations the + grid drives (dedup included). + """ + + sigmas: Tensor + timesteps: Tensor + num_model_evals: int + + +class NativeH3SigmaScheduleOp: + """Pure PyTorch reference for the H3 sigma grid — forward only.""" + + def __call__( + self, + num_inference_steps: int, + *, + shift: float, + device: str | torch.device | None = None, + ) -> H3SigmaSchedule: + return self.forward(num_inference_steps, shift=shift, device=device) + + def forward( + self, + num_inference_steps: int, + *, + shift: float, + device: str | torch.device | None = None, + ) -> H3SigmaSchedule: + """Build the grid in the frozen order; see the module docstring.""" + # Both checks mirror upstream's own failures and add no extra policy: + # num_inference_steps < 2 is set_timesteps' check (pinned blob, + # lines 149-153) and shift <= 0 the constructor's (lines 75-76). + if num_inference_steps < 2: + raise ValueError( + "num_inference_steps must be >= 2 (terminal zero included), got " + f"{num_inference_steps}." + ) + if shift <= 0: + raise ValueError(f"`shift` must be positive, got {shift}.") + + # Scalar operands must stay python floats in exactly this expression: + # pre-casting shift to a tensor or reassociating the division changes + # fp32 rounding and breaks byte equality with upstream (lines 156-159). + # device="cpu" is explicit because factory functions without a device + # follow the ambient default device (meta/CUDA context) instead. + base = torch.linspace(1.0, 0.0, int(num_inference_steps), dtype=torch.float32, device="cpu") + shift = float(shift) + sigmas = shift * base / (1 + (shift - 1) * base) + # The shift compresses spacing near sigma=1; drop the fp32 collisions + # it creates before deriving timesteps and the evaluation count. + sigmas = torch.unique_consecutive(sigmas) + timesteps = 1.0 - sigmas[:-1] + + return H3SigmaSchedule( + sigmas=sigmas.to(device=device), + timesteps=timesteps.to(device=device), + num_model_evals=int(sigmas.numel()) - 1, + ) diff --git a/rl_engine/runtime/registry.py b/rl_engine/runtime/registry.py index 528c77b98..b1af59756 100644 --- a/rl_engine/runtime/registry.py +++ b/rl_engine/runtime/registry.py @@ -204,6 +204,14 @@ class OpBackend(Enum, metaclass=_KernelEnumMeta): CUDA_SM90_LM_HEAD = "rl_engine.kernels.ops.cuda.linear.lm_head.SM90LMHeadOp" CUDA_SM90_EMBEDDING = "rl_engine.kernels.ops.cuda.linear.embedding.SM90EmbeddingOp" + # MiniMax-H3 (RFC #420) scheduler grid + PYTORCH_H3_SIGMA_SCHEDULE = ( + "rl_engine.reference.minimax_h3.sigma_schedule.NativeH3SigmaScheduleOp" + ) + CUDA_H3_SIGMA_SCHEDULE = ( + "rl_engine.backends.cuda.model_specific.minimax_h3." "sigma_schedule.H3SigmaScheduleCudaOp" + ) + def _default_semantic_descriptors() -> tuple[OperatorBackendDescriptor, ...]: return ( @@ -412,6 +420,51 @@ def _default_semantic_descriptors() -> tuple[OperatorBackendDescriptor, ...]: fallback_policy=OperatorFallbackPolicy.RUNTIME_MANAGED, version_or_build_fingerprint="runtime-native-ffn-unresolved-v1", ), + # MiniMax-H3 sigma grid (RFC #420). Byte-frozen against the pinned + # upstream construction; see rl_engine/reference/minimax_h3/. + OperatorBackendDescriptor( + semantic_op="sigma_schedule_h3", + backend_id="rlkernel.h3_sigma_schedule.reference", + supported_targets=frozenset({"rollout", "training"}), + supported_devices=frozenset({"cpu", "cuda", "rocm"}), + supported_dtypes=frozenset({"float32"}), + supported_topologies={"*": "*"}, + determinism_or_alignment_properties={ + "algorithm": "upstream MiniMaxH3Scheduler.set_timesteps, fp32 pinned", + "batch_invariant": True, + "deterministic": True, + "strict_observable": True, + }, + lifecycle=OperatorLifecycle.ENGINE_CONSTRUCTION, + implementation_class_or_factory=( + "rl_engine.reference.minimax_h3.sigma_schedule.NativeH3SigmaScheduleOp" + ), + fallback_policy=OperatorFallbackPolicy.ERROR, + version_or_build_fingerprint="NativeH3SigmaScheduleOp-v1", + ), + OperatorBackendDescriptor( + semantic_op="sigma_schedule_h3", + backend_id="rlkernel.h3_sigma_schedule.triton", + supported_targets=frozenset({"rollout", "training"}), + # NVIDIA-only backend (inline PTX div.rn.f32); HIP/ROCm stays + # on the reference — the kernels cannot run on AMD. + supported_devices=frozenset({"cuda"}), + supported_dtypes=frozenset({"float32"}), + supported_topologies={"*": "*"}, + determinism_or_alignment_properties={ + "algorithm": "device-side pinned law; ieee division, no FMA fusion", + "batch_invariant": True, + "deterministic": True, + "strict_observable": True, + }, + lifecycle=OperatorLifecycle.ENGINE_CONSTRUCTION, + implementation_class_or_factory=( + "rl_engine.backends.cuda.model_specific.minimax_h3." + "sigma_schedule.H3SigmaScheduleCudaOp" + ), + fallback_policy=OperatorFallbackPolicy.ERROR, + version_or_build_fingerprint="H3SigmaScheduleTritonOp-v1", + ), ) @@ -673,6 +726,10 @@ def __init__(self): OpBackend.TRITON_ROPE, OpBackend.PYTORCH_NATIVE_ROPE, ], + "sigma_schedule_h3": [ + OpBackend.CUDA_H3_SIGMA_SCHEDULE, + OpBackend.PYTORCH_H3_SIGMA_SCHEDULE, + ], }, "rocm": { "logp": [ @@ -696,6 +753,7 @@ def __init__(self): "kv_cache_attention": [OpBackend.PYTORCH_NATIVE_KV_CACHE_ATTN], "grpo_loss": [OpBackend.TRITON_GRPO_LOSS, OpBackend.PYTORCH_GRPO_LOSS], "rope": [OpBackend.TRITON_ROPE, OpBackend.PYTORCH_NATIVE_ROPE], + "sigma_schedule_h3": [OpBackend.PYTORCH_H3_SIGMA_SCHEDULE], "linear_logp": [ OpBackend.TRITON_LINEAR_LOGP, OpBackend.PYTORCH_LINEAR_LOGP, @@ -739,6 +797,7 @@ def __init__(self): OpBackend.PYTORCH_GRPO_LOSS, ], "rope": [OpBackend.TRITON_ROPE, OpBackend.PYTORCH_NATIVE_ROPE], + "sigma_schedule_h3": [OpBackend.PYTORCH_H3_SIGMA_SCHEDULE], "linear_logp": [ OpBackend.TRITON_LINEAR_LOGP, OpBackend.PYTORCH_LINEAR_LOGP, @@ -778,6 +837,7 @@ def __init__(self): "kv_cache_attention": [OpBackend.PYTORCH_NATIVE_KV_CACHE_ATTN], "grpo_loss": [OpBackend.PYTORCH_GRPO_LOSS], "rope": [OpBackend.PYTORCH_NATIVE_ROPE], + "sigma_schedule_h3": [OpBackend.PYTORCH_H3_SIGMA_SCHEDULE], "linear_logp": [OpBackend.PYTORCH_LINEAR_LOGP], "ratio_kl": [OpBackend.PYTORCH_RATIO_KL], "pack": [OpBackend.PYTORCH_PACK], diff --git a/tests/backends/cuda/test_h3_sigma_schedule_cuda.py b/tests/backends/cuda/test_h3_sigma_schedule_cuda.py new file mode 100644 index 000000000..0f66f91ae --- /dev/null +++ b/tests/backends/cuda/test_h3_sigma_schedule_cuda.py @@ -0,0 +1,366 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors + +"""Device tests for the Triton NVIDIA CUDA h3_sigma_schedule backend (RFC #420). + +The acceptance oracle is the pinned CPU fixture generated from the upstream +diffusers scheduler (see tests/reference/minimax_h3/). GPU results are copied +to the host only for comparison; no value is recomputed on the CPU. + +The backend is NVIDIA-only (its kernels emit inline PTX ``div.rn.f32``); +the suite runs on NVIDIA CUDA hardware and skips on HIP/ROCm builds, where +``H3SigmaScheduleCudaOp`` refuses construction and the registry serves the +PyTorch reference instead. +""" + +from __future__ import annotations + +import json +import pathlib + +import pytest +import torch + +pytestmark = pytest.mark.skipif( + not torch.cuda.is_available() or torch.version.hip is not None, + reason="NVIDIA CUDA device required: the Triton path is NVIDIA-only " + "(skipped without a GPU and on HIP/ROCm builds)", +) + +_FIXTURE_PATH = ( + pathlib.Path(__file__).resolve().parents[2] + / "reference" + / "minimax_h3" + / "h3_sigma_schedule_upstream_fixture.json" +) +_EXPECTED_CASES = {(2, 3.0), (2, 12.0), (4, 3.0), (4, 12.0), (25, 3.0), (25, 12.0)} + + +def _cuda_op(): + from rl_engine.backends.cuda.model_specific.minimax_h3.sigma_schedule import ( + H3SigmaScheduleCudaOp, + ) + + return H3SigmaScheduleCudaOp() + + +def _int32_bits(tensor: torch.Tensor) -> list[int]: + return tensor.detach().contiguous().view(torch.int32).tolist() + + +def _fixture_bits(bits: list[str]) -> list[int]: + return [int(value, 16) for value in bits] + + +def test_matches_pinned_cpu_fixture_bitwise(): + fixture = json.loads(_FIXTURE_PATH.read_text(encoding="utf-8")) + assert fixture["upstream"]["blob_sha1"] == "ca69ae6a2f6aae5bb68498216316a71eb07c1f7b" + + actual_cases = [(case["num_inference_steps"], case["shift"]) for case in fixture["cases"]] + assert set(actual_cases) == _EXPECTED_CASES + + op = _cuda_op() + for case in fixture["cases"]: + schedule = op(case["num_inference_steps"], shift=case["shift"]) + + assert schedule.sigmas.is_cuda + assert schedule.sigmas.dtype is torch.float32 + assert tuple(schedule.sigmas.shape) == (case["sigmas_count"],) + assert _int32_bits(schedule.sigmas.cpu()) == _fixture_bits(case["sigmas_bits"]) + + assert schedule.timesteps.is_cuda + assert schedule.timesteps.dtype is torch.float32 + assert tuple(schedule.timesteps.shape) == (case["timesteps_count"],) + expected = _fixture_bits(case["timesteps_bits"]) + assert _int32_bits(schedule.timesteps.cpu()) == expected + + assert schedule.num_model_evals == case["num_model_evals"] + + +@pytest.mark.parametrize( + ("num_inference_steps", "shift"), + [ + # N=50, shift=12 covers the fusion-sensitive elements (e.g. index 16) + # where a contracted FMA and separate mul/add roundings differ in the + # CPU simulation; the launch disables fp fusion and this comparison is + # the hardware verdict. + (50, 12.0), + (50, 3.0), + (1000, 12.0), + ], +) +def test_matches_live_cpu_reference_bitwise(num_inference_steps, shift): + from rl_engine.reference.minimax_h3.sigma_schedule import NativeH3SigmaScheduleOp + + expected = NativeH3SigmaScheduleOp()(num_inference_steps, shift=shift) + actual = _cuda_op()(num_inference_steps, shift=shift) + + assert _int32_bits(actual.sigmas.cpu()) == _int32_bits(expected.sigmas) + assert _int32_bits(actual.timesteps.cpu()) == _int32_bits(expected.timesteps) + assert actual.num_model_evals == expected.num_model_evals + + +def test_dense_grid_dedup_matches_cpu_reference(): + # shift=12 creates fp32 collisions near sigma=1 at this density: the grid + # collapses by ~4%, so the compaction path is genuinely exercised. + n = 1_000_000 + from rl_engine.reference.minimax_h3.sigma_schedule import NativeH3SigmaScheduleOp + + expected = NativeH3SigmaScheduleOp()(n, shift=12.0) + actual = _cuda_op()(n, shift=12.0) + + assert actual.sigmas.numel() < n + assert actual.sigmas.numel() == expected.sigmas.numel() + assert _int32_bits(actual.sigmas.cpu()) == _int32_bits(expected.sigmas) + assert _int32_bits(actual.timesteps.cpu()) == _int32_bits(expected.timesteps) + assert actual.num_model_evals == expected.num_model_evals + + +def test_reference_with_cuda_device_preserves_fixture_bytes(): + # Path-B oracle (the benchmark's no-GPU-backend baseline): same CPU build, + # then the operator moves sigmas/timesteps to the GPU. The move must + # preserve the pinned bytes — this also guards against a later + # "optimization" building the grid directly on the device inside the + # reference, which would silently change its arithmetic. + fixture = json.loads(_FIXTURE_PATH.read_text(encoding="utf-8")) + from rl_engine.reference.minimax_h3.sigma_schedule import NativeH3SigmaScheduleOp + + op = NativeH3SigmaScheduleOp() + for case in fixture["cases"]: + schedule = op(case["num_inference_steps"], shift=case["shift"], device="cuda") + + assert schedule.sigmas.is_cuda and schedule.timesteps.is_cuda + assert _int32_bits(schedule.sigmas.cpu()) == _fixture_bits(case["sigmas_bits"]) + assert _int32_bits(schedule.timesteps.cpu()) == _fixture_bits(case["timesteps_bits"]) + assert schedule.num_model_evals == case["num_model_evals"] + + +@pytest.mark.parametrize("shift", [3.0, 12.0]) +@pytest.mark.parametrize("num_inference_steps", [2, 4, 25]) +def test_grid_invariants_on_device(num_inference_steps, shift): + schedule = _cuda_op()(num_inference_steps, shift=shift) + sigmas = schedule.sigmas + timesteps = schedule.timesteps + + assert sigmas.is_cuda and timesteps.is_cuda + assert sigmas.dtype is torch.float32 and timesteps.dtype is torch.float32 + assert sigmas[0].item() == 1.0 + assert sigmas[-1].item() == 0.0 + assert bool((sigmas[1:] < sigmas[:-1]).all()) + assert bool((timesteps[1:] > timesteps[:-1]).all()) + assert timesteps[0].item() == 0.0 + assert schedule.num_model_evals == sigmas.numel() - 1 == timesteps.numel() + + +@pytest.mark.parametrize("num_inference_steps", [1, 0, -4]) +def test_rejects_num_inference_steps_below_two(num_inference_steps): + with pytest.raises(ValueError, match="num_inference_steps"): + _cuda_op()(num_inference_steps, shift=12.0) + + +# --- Edge-consistency regression (independent GPU audit 2026-10-16) ------- +# +# These four cases come from the independent audit's counterexamples, each +# confirmed against the AST-executed pinned upstream set_timesteps +# (upstream-counterexamples.json): the CPU reference matches upstream on all +# of them, so the device path must match the reference bitwise. Landmark +# bits/lengths from the upstream execution are asserted directly so the +# oracle does not depend on the local reference staying correct. + + +def _assert_bitwise_match(reference, actual): + assert tuple(actual.sigmas.shape) == tuple(reference.sigmas.shape) + assert tuple(actual.timesteps.shape) == tuple(reference.timesteps.shape) + assert actual.num_model_evals == reference.num_model_evals + assert _int32_bits(actual.sigmas.cpu()) == _int32_bits(reference.sigmas) + assert _int32_bits(actual.timesteps.cpu()) == _int32_bits(reference.timesteps) + + +def test_subnormal_shift_1e_45_matches_reference_upstream_parity(): + # fl32(1e-45) rounds UP to the minimum subnormal 2^-149 (bit pattern 1): + # upstream grid = [+inf, subnormal quotients..., 0.0] with N=25 collapsing + # to 10 rows. A flush-to-zero device multiply turns the numerator into 0 + # and produced [NaN, 0.0] (length 2) before the fix. + from rl_engine.reference.minimax_h3.sigma_schedule import NativeH3SigmaScheduleOp + + expected = NativeH3SigmaScheduleOp()(25, shift=1e-45) + actual = _cuda_op()(25, shift=1e-45) + + _assert_bitwise_match(expected, actual) + # Landmarks from the pinned upstream execution: + assert expected.sigmas.numel() == 10 + assert _int32_bits(actual.sigmas.cpu())[0] == 0x7F800000 # +inf + + +def test_underflow_shift_1e_46_matches_reference_upstream_parity(): + # fl32(1e-46) underflows to +0.0: element 0 is 0/+0, an invalid operation + # whose torch-CPU (x86) result is the DEFAULT QUIET NaN WITH SIGN BIT SET + # (0xffc00000); the rest of the grid is +0.0, so upstream keeps 2 rows. + from rl_engine.reference.minimax_h3.sigma_schedule import NativeH3SigmaScheduleOp + + expected = NativeH3SigmaScheduleOp()(25, shift=1e-46) + actual = _cuda_op()(25, shift=1e-46) + + _assert_bitwise_match(expected, actual) + assert expected.sigmas.numel() == 2 + assert _int32_bits(actual.sigmas.cpu())[0] == -4194304 # 0xffc00000, signed + + +def test_positive_nan_shift_propagates_positive_quiet_nan(): + # shift=+NaN: torch CPU propagates the (quieted) input NaN bits through + # the whole grid -> 0x7fc00000 everywhere, sign included. + from rl_engine.reference.minimax_h3.sigma_schedule import NativeH3SigmaScheduleOp + + expected = NativeH3SigmaScheduleOp()(25, shift=float("nan")) + actual = _cuda_op()(25, shift=float("nan")) + + _assert_bitwise_match(expected, actual) + assert set(_int32_bits(actual.sigmas.cpu())) == {0x7FC00000} + + +def test_negative_nan_shift_propagates_negative_quiet_nan(): + # shift=-NaN: same propagation, but the SIGN BIT travels with the payload + # -> 0xffc00000 everywhere (the fixed positive fill was wrong here). + from rl_engine.reference.minimax_h3.sigma_schedule import NativeH3SigmaScheduleOp + + expected = NativeH3SigmaScheduleOp()(25, shift=-float("nan")) + actual = _cuda_op()(25, shift=-float("nan")) + + _assert_bitwise_match(expected, actual) + assert set(_int32_bits(actual.sigmas.cpu())) == {-4194304} # 0xffc00000, signed + + +_EDGE_SHIFTS = [ + float("nan"), # all-NaN grid; IEEE != keeps every row (matches upstream) + float("inf"), # all-NaN grid via fp32 scalar overflow semantics + 1e-10, # +inf first sigma: (1e-10 - 1.0) rounds to exactly -1.0 in fp32 + 1e39, # fp32-overflowing shift -> NaN grid + 3.5e38, # just above float32 max -> NaN grid + 1e30, # extreme finite shift collapses the grid to [1.0, 0.0] +] + + +@pytest.mark.parametrize("shift", _EDGE_SHIFTS) +def test_edge_shifts_match_reference_upstream_parity(shift): + # contract-audit.md §2 proved the CPU reference bit-identical to the + # actually-executed upstream set_timesteps on exactly these inputs, so + # the device path must reproduce the reference bytes, NaN grids included. + from rl_engine.reference.minimax_h3.sigma_schedule import NativeH3SigmaScheduleOp + + expected = NativeH3SigmaScheduleOp()(8, shift=shift) + actual = _cuda_op()(8, shift=shift) + + assert actual.num_model_evals == expected.num_model_evals + assert tuple(actual.sigmas.shape) == tuple(expected.sigmas.shape) + assert tuple(actual.timesteps.shape) == tuple(expected.timesteps.shape) + assert _int32_bits(actual.sigmas.cpu()) == _int32_bits(expected.sigmas) + assert _int32_bits(actual.timesteps.cpu()) == _int32_bits(expected.timesteps) + + +@pytest.mark.parametrize(("num_inference_steps", "truncated_to"), [(2.9, 2), (3.1, 3)]) +def test_float_num_steps_truncate_like_reference(num_inference_steps, truncated_to): + # Upstream parity: int() truncation, device side must agree with both the + # reference call and the truncated integer grid. + from rl_engine.reference.minimax_h3.sigma_schedule import NativeH3SigmaScheduleOp + + expected = NativeH3SigmaScheduleOp()(num_inference_steps, shift=12.0) + integer_grid = NativeH3SigmaScheduleOp()(truncated_to, shift=12.0) + actual = _cuda_op()(num_inference_steps, shift=12.0) + + assert _int32_bits(expected.sigmas) == _int32_bits(integer_grid.sigmas) + assert _int32_bits(actual.sigmas.cpu()) == _int32_bits(expected.sigmas) + assert _int32_bits(actual.timesteps.cpu()) == _int32_bits(expected.timesteps) + assert actual.num_model_evals == expected.num_model_evals + + +@pytest.mark.parametrize("shift", [0.0, -1.5, float("-inf")]) +def test_rejects_non_positive_shift(shift): + with pytest.raises(ValueError, match="positive"): + _cuda_op()(8, shift=shift) + + +def test_rejects_cpu_device_instead_of_silent_fallback(): + with pytest.raises(ValueError, match="NVIDIA CUDA device"): + _cuda_op()(25, shift=12.0, device="cpu") + + +def test_repeated_and_interleaved_calls_are_stable(): + from rl_engine.reference.minimax_h3.sigma_schedule import NativeH3SigmaScheduleOp + + cuda_op = _cuda_op() + cpu_op = NativeH3SigmaScheduleOp() + + first = cuda_op(25, shift=12.0) + cpu_op(25, shift=3.0) # interleaved CPU call must not pollute device state + second = cuda_op(25, shift=3.0) + third = cuda_op(25, shift=12.0) + + assert _int32_bits(first.sigmas.cpu()) == _int32_bits(third.sigmas.cpu()) + assert _int32_bits(first.timesteps.cpu()) == _int32_bits(third.timesteps.cpu()) + assert _int32_bits(second.timesteps.cpu()) != _int32_bits(first.timesteps.cpu()) + + +def test_backend_identity_is_device_path_not_cpu_fallback(): + from rl_engine.backends.cuda.model_specific.minimax_h3.sigma_schedule import ( + H3SigmaScheduleCudaOp, + ) + from rl_engine.runtime.registry import OpBackend + + # The registry lists the NVIDIA Triton implementation ahead of the PyTorch + # reference for device "cuda"; outputs above prove the device path ran. + assert OpBackend.CUDA_H3_SIGMA_SCHEDULE.value.endswith( + "model_specific.minimax_h3.sigma_schedule.H3SigmaScheduleCudaOp" + ) + assert OpBackend.PYTORCH_H3_SIGMA_SCHEDULE.value.endswith( + "reference.minimax_h3.sigma_schedule.NativeH3SigmaScheduleOp" + ) + + schedule = H3SigmaScheduleCudaOp()(8, shift=3.0) + assert schedule.sigmas.device.type == "cuda" + assert torch.version.hip is None # NVIDIA-only backend + + +def test_public_entry_selects_device_backend_and_reports_provenance(): + # The registry must hand out the Triton class itself (device identity is + # the instance type, not just `.is_cuda`, which a CPU-build-then-copy + # would also have), and the semantic path must report the same backend. + from rl_engine.backends.cuda.model_specific.minimax_h3.sigma_schedule import ( + H3SigmaScheduleCudaOp, + ) + from rl_engine.runtime.registry import kernel_registry + from rl_engine.runtime.semantic_registry import OperatorRequirements + + op = kernel_registry.get_op("sigma_schedule_h3", device="cuda") + assert isinstance(op, H3SigmaScheduleCudaOp) + assert op(25, shift=12.0).sigmas.is_cuda + + session = kernel_registry.semantic.session() + resolution = session.resolve( + semantic_op="sigma_schedule_h3", + requested_backend="rlkernel.h3_sigma_schedule.triton", + target="rollout", + requirements=OperatorRequirements(device="cuda", dtype="float32"), + ) + instance = session.instantiate(resolution) + assert isinstance(instance, H3SigmaScheduleCudaOp) + provenance = session.instance_provenance(resolution, instance) + assert provenance.backend_id == "rlkernel.h3_sigma_schedule.triton" + assert provenance.concrete_implementation.endswith("H3SigmaScheduleCudaOp") + assert torch.version.hip is None # NVIDIA-only backend + + +@pytest.mark.skipif( + torch.cuda.device_count() < 2, reason="multi-GPU device-guard test needs 2 devices" +) +def test_target_device_guard_with_mismatched_current_device(): + from rl_engine.reference.minimax_h3.sigma_schedule import NativeH3SigmaScheduleOp + + other = "cuda:1" if torch.cuda.current_device() == 0 else "cuda:0" + expected = NativeH3SigmaScheduleOp()(25, shift=12.0) + with torch.cuda.device(0 if other == "cuda:1" else 1): + # Current device deliberately differs from the requested target. + schedule = _cuda_op()(25, shift=12.0, device=other) + + assert schedule.sigmas.device == torch.device(other) + assert _int32_bits(schedule.sigmas.cpu()) == _int32_bits(expected.sigmas) + assert _int32_bits(schedule.timesteps.cpu()) == _int32_bits(expected.timesteps) diff --git a/tests/reference/minimax_h3/__init__.py b/tests/reference/minimax_h3/__init__.py new file mode 100644 index 000000000..86cf4c9da --- /dev/null +++ b/tests/reference/minimax_h3/__init__.py @@ -0,0 +1,2 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors diff --git a/tests/reference/minimax_h3/h3_sigma_schedule_upstream_fixture.json b/tests/reference/minimax_h3/h3_sigma_schedule_upstream_fixture.json new file mode 100644 index 000000000..e4a5ba76a --- /dev/null +++ b/tests/reference/minimax_h3/h3_sigma_schedule_upstream_fixture.json @@ -0,0 +1,201 @@ +{ + "upstream": { + "url": "https://github.com/huggingface/diffusers/blob/f53d552/src/diffusers/schedulers/scheduling_minimax_h3.py", + "commit": "f53d552036a0d1bd5570782a39cd40cfabf112bc", + "blob_sha1": "ca69ae6a2f6aae5bb68498216316a71eb07c1f7b", + "fetch_note": "fetched 2026-10-12 via the api.github.com contents endpoint; the blob matched at commit f53d552 and at main at fetch time", + "class": "MiniMaxH3Scheduler", + "method": "set_timesteps", + "execution": "the method body is AST-extracted and executed as-is against a namespace shim; no formula is transcribed", + "torch": "2.14.1", + "device": "cpu", + "dtype": "float32" + }, + "generator": "make_upstream_fixture.py --upstream-file --out ", + "cases": [ + { + "num_inference_steps": 2, + "shift": 3.0, + "sigmas_count": 2, + "sigmas_bits": [ + "3f800000", + "00000000" + ], + "timesteps_count": 1, + "timesteps_bits": [ + "00000000" + ], + "num_model_evals": 1 + }, + { + "num_inference_steps": 2, + "shift": 12.0, + "sigmas_count": 2, + "sigmas_bits": [ + "3f800000", + "00000000" + ], + "timesteps_count": 1, + "timesteps_bits": [ + "00000000" + ], + "num_model_evals": 1 + }, + { + "num_inference_steps": 4, + "shift": 3.0, + "sigmas_count": 4, + "sigmas_bits": [ + "3f800000", + "3f5b6db6", + "3f199999", + "00000000" + ], + "timesteps_count": 3, + "timesteps_bits": [ + "00000000", + "3e124928", + "3eccccce" + ], + "num_model_evals": 3 + }, + { + "num_inference_steps": 4, + "shift": 12.0, + "sigmas_count": 4, + "sigmas_bits": [ + "3f800000", + "3f75c28f", + "3f5b6db6", + "00000000" + ], + "timesteps_count": 3, + "timesteps_bits": [ + "00000000", + "3d23d710", + "3e124928" + ], + "num_model_evals": 3 + }, + { + "num_inference_steps": 25, + "shift": 3.0, + "sigmas_count": 25, + "sigmas_bits": [ + "3f800000", + "3f7c57c6", + "3f787878", + "3f745d17", + "3f700001", + "3f6b5ad6", + "3f666666", + "3f611a7d", + "3f5b6db6", + "3f555555", + "3f4ec4ed", + "3f47ae14", + "3f400000", + "3f37a6f4", + "3f2e8ba3", + "3f249249", + "3f199999", + "3f0d7944", + "3f000000", + "3ee1e1e1", + "3ec00000", + "3e99999a", + "3e5b6db7", + "3dec4ec4", + "00000000" + ], + "timesteps_count": 24, + "timesteps_bits": [ + "00000000", + "3c6a0e80", + "3cf0f100", + "3d3a2e90", + "3d7ffff0", + "3da52950", + "3dccccd0", + "3df72c18", + "3e124928", + "3e2aaaac", + "3e44ec4c", + "3e6147b0", + "3e800000", + "3e90b218", + "3ea2e8ba", + "3eb6db6e", + "3eccccce", + "3ee50d78", + "3f000000", + "3f0f0f10", + "3f200000", + "3f333333", + "3f492492", + "3f627628" + ], + "num_model_evals": 24 + }, + { + "num_inference_steps": 25, + "shift": 12.0, + "sigmas_count": 25, + "sigmas_bits": [ + "3f800000", + "3f7f1369", + "3f7e133f", + "3f7cfcfd", + "3f7bcda5", + "3f7a81a5", + "3f7914c2", + "3f7781d3", + "3f75c28f", + "3f73cf3d", + "3f719e34", + "3f6f2341", + "3f6c4ec5", + "3f690c5b", + "3f6540f4", + "3f60c7ce", + "3f5b6db6", + "3f54e931", + "3f4ccccd", + "3f426e2d", + "3f34b4b4", + "3f21af28", + "3f0590b2", + "3eaf8af8", + "00000000" + ], + "timesteps_count": 24, + "timesteps_bits": [ + "00000000", + "3b6c9700", + "3bf66080", + "3c40c0c0", + "3c864b60", + "3cafcb60", + "3cdd67c0", + "3d07e2d0", + "3d23d710", + "3d430c30", + "3d661cc0", + "3d86e5f8", + "3d9d89d8", + "3db79d28", + "3dd5f860", + "3df9c190", + "3e124928", + "3e2c5b3c", + "3e4ccccc", + "3e76474c", + "3e969698", + "3ebca1b0", + "3ef4de9c", + "3f283a84" + ], + "num_model_evals": 24 + } + ] +} diff --git a/tests/reference/minimax_h3/make_upstream_fixture.py b/tests/reference/minimax_h3/make_upstream_fixture.py new file mode 100644 index 000000000..0dcb5d2a7 --- /dev/null +++ b/tests/reference/minimax_h3/make_upstream_fixture.py @@ -0,0 +1,154 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors + +"""Regenerate the pinned-upstream fixture for ``test_h3_sigma_schedule``. + +Executes the ORIGINAL ``MiniMaxH3Scheduler.set_timesteps`` extracted from the +upstream diffusers source file — no formula is transcribed by hand. The +method's AST node is compiled and run against a namespace shim that supplies +the scheduler attributes it touches, so the construction arithmetic (linspace, +exponential shift, ``unique_consecutive``, timestep conversion) is upstream's +own code. + +The upstream file's git blob sha must equal ``UPSTREAM_BLOB_SHA`` below, so +this script cannot silently regenerate the fixture from a different revision. + +Usage:: + + python make_upstream_fixture.py \ + --upstream-file \ + --out h3_sigma_schedule_upstream_fixture.json +""" + +from __future__ import annotations + +import argparse +import ast +import copy +import hashlib +import json +import types + +import torch + +UPSTREAM_URL = ( + "https://github.com/huggingface/diffusers/blob/f53d552/src/diffusers/schedulers/" + "scheduling_minimax_h3.py" +) +UPSTREAM_COMMIT = "f53d552036a0d1bd5570782a39cd40cfabf112bc" +UPSTREAM_BLOB_SHA = "ca69ae6a2f6aae5bb68498216316a71eb07c1f7b" +UPSTREAM_FETCH_NOTE = ( + "fetched 2026-10-12 via the api.github.com contents endpoint; the blob " + "matched at commit f53d552 and at main at fetch time" +) + +# RFC #420 fingerprint matrix: video shift 12.0, audio shift 3.0. +CASES = [ + (2, 3.0), + (2, 12.0), + (4, 3.0), + (4, 12.0), + (25, 3.0), + (25, 12.0), +] + + +def _git_blob_sha(data: bytes) -> str: + header = b"blob %d\x00" % len(data) + return hashlib.sha1(header + data).hexdigest() + + +def _extract_set_timesteps(source: str) -> ast.FunctionDef: + """Return the ``set_timesteps`` FunctionDef of ``MiniMaxH3Scheduler``.""" + for node in ast.parse(source).body: + if isinstance(node, ast.ClassDef) and node.name == "MiniMaxH3Scheduler": + for item in node.body: + if isinstance(item, ast.FunctionDef) and item.name == "set_timesteps": + return item + raise SystemExit("MiniMaxH3Scheduler.set_timesteps not found in upstream source") + + +def _run_upstream_set_timesteps( + func: ast.FunctionDef, num_inference_steps: int, shift: float +) -> tuple[torch.Tensor, torch.Tensor]: + """Execute the extracted upstream method once against a shim instance.""" + module = ast.Module(body=[copy.deepcopy(func)], type_ignores=[]) + code = compile(module, "", "exec") + namespace: dict = {"torch": torch} + exec(code, namespace) # the point is to run upstream's own code + + scheduler = types.SimpleNamespace( + num_inference_steps=None, + sigmas=None, + timesteps=None, + _shift=float(shift), + _step_index=None, + _begin_index=None, + ) + namespace["set_timesteps"](scheduler, num_inference_steps=num_inference_steps, device=None) + return scheduler.sigmas, scheduler.timesteps + + +def _bits(tensor: torch.Tensor) -> list[str]: + values = tensor.detach().cpu().contiguous().view(torch.int32).tolist() + return [f"{value & 0xFFFFFFFF:08x}" for value in values] + + +def main() -> None: + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--upstream-file", required=True) + parser.add_argument("--out", required=True) + args = parser.parse_args() + + with open(args.upstream_file, "rb") as handle: + data = handle.read() + blob_sha = _git_blob_sha(data) + if blob_sha != UPSTREAM_BLOB_SHA: + raise SystemExit( + f"upstream file blob sha {blob_sha} != pinned {UPSTREAM_BLOB_SHA}; " + "refusing to regenerate the fixture from an unknown revision" + ) + + func = _extract_set_timesteps(data.decode("utf-8")) + cases = [] + for num_inference_steps, shift in CASES: + sigmas, timesteps = _run_upstream_set_timesteps(func, num_inference_steps, shift) + cases.append( + { + "num_inference_steps": num_inference_steps, + "shift": shift, + "sigmas_count": int(sigmas.numel()), + "sigmas_bits": _bits(sigmas), + "timesteps_count": int(timesteps.numel()), + "timesteps_bits": _bits(timesteps), + "num_model_evals": int(timesteps.numel()), + } + ) + + fixture = { + "upstream": { + "url": UPSTREAM_URL, + "commit": UPSTREAM_COMMIT, + "blob_sha1": UPSTREAM_BLOB_SHA, + "fetch_note": UPSTREAM_FETCH_NOTE, + "class": "MiniMaxH3Scheduler", + "method": "set_timesteps", + "execution": ( + "the method body is AST-extracted and executed as-is against a " + "namespace shim; no formula is transcribed" + ), + "torch": torch.__version__, + "device": "cpu", + "dtype": "float32", + }, + "generator": "make_upstream_fixture.py --upstream-file --out ", + "cases": cases, + } + with open(args.out, "w", encoding="utf-8") as handle: + json.dump(fixture, handle, indent=2) + handle.write("\n") + print(f"wrote {args.out} with {len(cases)} cases") + + +if __name__ == "__main__": + main() diff --git a/tests/reference/minimax_h3/test_h3_sigma_schedule.py b/tests/reference/minimax_h3/test_h3_sigma_schedule.py new file mode 100644 index 000000000..4ac2af4a2 --- /dev/null +++ b/tests/reference/minimax_h3/test_h3_sigma_schedule.py @@ -0,0 +1,185 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors + +"""Focused tests for the MiniMax-H3 sigma grid reference (RFC #420). + +Covers only the four-step construction pinned by +``MiniMaxH3Scheduler.set_timesteps`` (diffusers ``scheduling_minimax_h3.py``, +lines 148-168): ``linspace(1, 0, N)`` in fp32 -> exponential shift -> +``unique_consecutive`` -> ``timesteps = 1 - sigmas[:-1]``. + +The primary oracle is ``h3_sigma_schedule_upstream_fixture.json``: fp32 bit +patterns produced by actually executing the upstream method (AST-extracted +and run as-is, see ``make_upstream_fixture.py``) from the source file with +git blob sha ``ca69ae6a2f6aae5bb68498216316a71eb07c1f7b`` (diffusers commit +``f53d552``, the only revision that touched the file). That check needs +no network, no installed diffusers, and never skips. +""" + +from __future__ import annotations + +import json +import pathlib + +import pytest +import torch + +from rl_engine.reference.minimax_h3.sigma_schedule import NativeH3SigmaScheduleOp + +_OP = NativeH3SigmaScheduleOp() + +# RFC #420 fingerprint shifts: video 12.0, audio 3.0 (H12 — one grid per +# modality; grids must never be shared implicitly). +_CASES = [ + (2, 12.0), + (2, 3.0), + (8, 12.0), + (8, 3.0), + (25, 12.0), + (25, 3.0), + (50, 12.0), + (50, 3.0), + (1000, 12.0), + (1000, 3.0), +] + + +_FIXTURE_PATH = pathlib.Path(__file__).parent / "h3_sigma_schedule_upstream_fixture.json" +_UPSTREAM_BLOB_SHA = "ca69ae6a2f6aae5bb68498216316a71eb07c1f7b" +_EXPECTED_CASES = {(2, 3.0), (2, 12.0), (4, 3.0), (4, 12.0), (25, 3.0), (25, 12.0)} + + +def _int32_bits(tensor: torch.Tensor) -> list[int]: + return tensor.detach().contiguous().view(torch.int32).tolist() + + +def _bits_to_int32(bits: list[str]) -> torch.Tensor: + return torch.tensor([int(value, 16) for value in bits], dtype=torch.int32) + + +def test_matches_pinned_upstream_set_timesteps_bitwise(): + # Non-skippable: the expected fp32 bit patterns were generated by actually + # executing the upstream set_timesteps method from the pinned source (see + # make_upstream_fixture.py) — no network, no diffusers install, and no + # hand-copied formula are involved. + fixture = json.loads(_FIXTURE_PATH.read_text(encoding="utf-8")) + assert fixture["upstream"]["blob_sha1"] == _UPSTREAM_BLOB_SHA + + actual_cases = [(case["num_inference_steps"], case["shift"]) for case in fixture["cases"]] + assert len(actual_cases) == len(set(actual_cases)), "duplicate fixture cases" + assert set(actual_cases) == _EXPECTED_CASES + + for case in fixture["cases"]: + schedule = _OP(case["num_inference_steps"], shift=case["shift"]) + + assert schedule.sigmas.dtype is torch.float32 + assert tuple(schedule.sigmas.shape) == (case["sigmas_count"],) + assert _int32_bits(schedule.sigmas) == _bits_to_int32(case["sigmas_bits"]).tolist() + + assert schedule.timesteps.dtype is torch.float32 + assert tuple(schedule.timesteps.shape) == (case["timesteps_count"],) + expected_timesteps = _bits_to_int32(case["timesteps_bits"]).tolist() + assert _int32_bits(schedule.timesteps) == expected_timesteps + + assert schedule.num_model_evals == case["num_model_evals"] + + +@pytest.mark.parametrize(("num_inference_steps", "shift"), _CASES) +def test_grid_invariants(num_inference_steps, shift): + schedule = _OP(num_inference_steps, shift=shift) + sigmas, timesteps = schedule.sigmas, schedule.timesteps + + assert sigmas.dtype is torch.float32 + assert timesteps.dtype is torch.float32 + assert sigmas.numel() <= num_inference_steps + # First sigma is exactly 1.0 for the released shifts (3.0, 12.0); the + # shift maps the terminal 0.0 to exactly 0.0. + assert sigmas[0].item() == 1.0 + assert sigmas[-1].item() == 0.0 + assert bool((sigmas[1:] < sigmas[:-1]).all()) + assert bool((timesteps[1:] > timesteps[:-1]).all()) + assert timesteps[0].item() == 0.0 + # Call count follows the deduplicated grid, not the requested N. + assert schedule.num_model_evals == sigmas.numel() - 1 == timesteps.numel() + + +def test_unique_consecutive_collapses_shift_created_collisions(): + # shift=12 compresses spacing near sigma=1 enough that a dense grid holds + # fp32 collisions: N=1e6 collapses ~4% of rows in a few ms. The base + # linspace itself stays duplicate-free at this density, so the collapse is + # attributable to the shift step, as the upstream docstring claims. + n = 1_000_000 + base = torch.linspace(1.0, 0.0, n, dtype=torch.float32) + assert torch.unique_consecutive(base).numel() == n # precondition + + schedule = _OP(n, shift=12.0) + + assert schedule.sigmas.numel() < n + assert schedule.num_model_evals == schedule.sigmas.numel() - 1 + assert schedule.timesteps.numel() == schedule.num_model_evals + assert bool((schedule.sigmas[1:] < schedule.sigmas[:-1]).all()) + assert schedule.sigmas[-1].item() == 0.0 + + +def test_video_and_audio_shifts_give_distinct_grids(): + video = _OP(25, shift=12.0) + audio = _OP(25, shift=3.0) + + assert not torch.equal(video.sigmas, audio.sigmas) + assert not torch.equal(video.timesteps, audio.timesteps) + + +@pytest.mark.parametrize("num_inference_steps", [1, 0, -4]) +def test_rejects_num_inference_steps_below_two(num_inference_steps): + with pytest.raises(ValueError, match="num_inference_steps"): + _OP(num_inference_steps, shift=12.0) + + +@pytest.mark.parametrize("shift", [0.0, -1.5]) +def test_rejects_non_positive_shift(shift): + with pytest.raises(ValueError, match="positive"): + _OP(8, shift=shift) + + +def test_device_argument_preserves_bytes_and_dtype(): + local = _OP(25, shift=12.0) + moved = _OP(25, shift=12.0, device="cpu") + + assert _int32_bits(moved.sigmas) == _int32_bits(local.sigmas) + assert _int32_bits(moved.timesteps) == _int32_bits(local.timesteps) + assert moved.sigmas.dtype is torch.float32 + assert moved.sigmas.device.type == "cpu" + + +def test_builds_on_cpu_under_non_cpu_default_device(): + # torch.linspace without an explicit device follows the ambient default + # device, so the reference must pin "cpu" itself. Under a meta default the + # unfixed implementation raised NotImplementedError inside + # unique_consecutive (no meta kernel); on real accelerators it would + # silently build the grid off-CPU. + with torch.device("meta"): + schedule = _OP(25, shift=12.0) + + assert schedule.sigmas.device.type == "cpu" + assert schedule.timesteps.device.type == "cpu" + reference = _OP(25, shift=12.0) + assert _int32_bits(schedule.sigmas) == _int32_bits(reference.sigmas) + assert _int32_bits(schedule.timesteps) == _int32_bits(reference.timesteps) + + +def test_matches_installed_diffusers_scheduler_when_available(): + diffusers = pytest.importorskip("diffusers") + scheduler_cls = getattr(diffusers, "MiniMaxH3Scheduler", None) + if scheduler_cls is None: + pytest.skip("installed diffusers has no MiniMaxH3Scheduler") + + for num_inference_steps, shift in ((25, 12.0), (50, 3.0)): + scheduler = scheduler_cls(shift=shift) + scheduler.set_timesteps(num_inference_steps, device="cpu") + schedule = _OP(num_inference_steps, shift=shift) + + assert torch.equal(scheduler.sigmas, schedule.sigmas) + assert torch.equal(scheduler.timesteps, schedule.timesteps) + # Upstream stores len(timesteps) as num_inference_steps, i.e. the + # model-evaluation count. + assert scheduler.num_inference_steps == schedule.num_model_evals diff --git a/tests/runtime/test_h3_sigma_schedule_dispatch.py b/tests/runtime/test_h3_sigma_schedule_dispatch.py new file mode 100644 index 000000000..47e4e1f84 --- /dev/null +++ b/tests/runtime/test_h3_sigma_schedule_dispatch.py @@ -0,0 +1,173 @@ +# SPDX-License-Identifier: Apache-2.0 +# Copyright (c) 2026 RL-Kernel Contributors + +"""Registry and semantic-entry tests for ``h3_sigma_schedule`` (RFC #420). + +Covers the public selection paths on any machine: legacy ``get_op`` candidate +fallback when the Triton backend cannot construct, semantic descriptor +registration, session resolve / instantiate / provenance for the reference +backend, and the CPU-pinned scalar-rounding helper. The device-path identity +(legacy ``get_op`` returning the Triton class, plus result bytes) is asserted +on hardware in tests/backends/cuda/test_h3_sigma_schedule_cuda.py. +""" + +from __future__ import annotations + +import importlib + +import pytest +import torch + +from rl_engine.runtime.registry import KernelRegistry, kernel_registry +from rl_engine.runtime.semantic_registry import ( + OperatorInstantiationError, + OperatorRequirements, + OperatorResolutionError, +) + +_BACKEND_MODULE = importlib.import_module( + "rl_engine.backends.cuda.model_specific.minimax_h3.sigma_schedule" +) +_REFERENCE_PATH = "rl_engine.reference.minimax_h3.sigma_schedule.NativeH3SigmaScheduleOp" + + +def test_get_op_cpu_selects_reference_and_builds_pinned_grid(): + op = kernel_registry.get_op("sigma_schedule_h3", device="cpu") + assert f"{type(op).__module__}.{type(op).__name__}" == _REFERENCE_PATH + + schedule = op(4, shift=3.0, device="cpu") + assert [f"{v:08x}" for v in schedule.sigmas.view(torch.int32).tolist()] == [ + "3f800000", + "3f5b6db6", + "3f199999", + "00000000", + ] + assert schedule.num_model_evals == 3 + + +def test_get_op_cuda_without_triton_skips_unavailable_backend(monkeypatch): + # Old behavior (review 2026-10-13): the unavailable GPU class was selected + # and failed only inside forward. Construction now raises ImportError, so + # the registry's first-available-candidate loop skips it at selection time + # (docs/architecture/runtime-dispatch.md) instead of the registry having + # to catch arbitrary forward errors. + monkeypatch.setattr(_BACKEND_MODULE, "_TRITON_AVAILABLE", False) + registry = KernelRegistry() + + op = registry.get_op("sigma_schedule_h3", device="cuda") + + assert f"{type(op).__module__}.{type(op).__name__}" == _REFERENCE_PATH + + +def test_cuda_backend_construction_fails_loudly_without_triton(monkeypatch): + monkeypatch.setattr(_BACKEND_MODULE, "_TRITON_AVAILABLE", False) + + with pytest.raises(ImportError, match="triton is required"): + _BACKEND_MODULE.H3SigmaScheduleCudaOp() + + +def test_shift_scalars_pinned_to_cpu_under_meta_default(): + # Old behavior: the factory calls followed the ambient default device and + # failed under a meta default (RuntimeError on .item()). + with torch.device("meta"): + shift_f32, shift_m1_f32 = _BACKEND_MODULE._shift_scalars(3.0) + assert shift_f32 == 3.0 + assert shift_m1_f32 == 2.0 + + +def test_shift_scalars_keep_fp64_subtract_then_round_semantics(): + # Pinned artifact of the upstream scalar handling: (1e-10 - 1.0) is + # subtracted in fp64 first and only then rounded to fp32, landing on + # exactly -1.0 (this feeds the division-by-zero +inf first sigma). + _, shift_m1_f32 = _BACKEND_MODULE._shift_scalars(1e-10) + assert shift_m1_f32 == -1.0 + + +def test_cuda_backend_construction_rejected_on_hip_build(monkeypatch): + # The kernels emit inline PTX div.rn.f32 and cannot run on AMD: a HIP + # build must be rejected AT CONSTRUCTION (long before any Triton + # compilation), with the same ImportError contract the registry's + # candidate loop already handles for missing Triton. Simulated without + # any AMD hardware by monkeypatching torch.version.hip. + monkeypatch.setattr(torch.version, "hip", "6.2") + + with pytest.raises(ImportError, match="NVIDIA-only.*inline.*PTX"): + _BACKEND_MODULE.H3SigmaScheduleCudaOp() + + +def test_get_op_falls_back_to_reference_on_hip_build(monkeypatch): + # End-to-end selection semantics under the same simulation: the registry + # walks the cuda candidate list, the Triton class refuses construction + # with ImportError, and the PyTorch reference is served instead — no + # error reaches the caller and nothing is silently misattributed. + monkeypatch.setattr(torch.version, "hip", "6.2") + registry = KernelRegistry() + + op = registry.get_op("sigma_schedule_h3", device="cuda") + + assert f"{type(op).__module__}.{type(op).__name__}" == _REFERENCE_PATH + + +def test_semantic_descriptors_declared_truthfully(): + reference = kernel_registry.semantic.backend_descriptor( + "sigma_schedule_h3", "rlkernel.h3_sigma_schedule.reference" + ) + triton = kernel_registry.semantic.backend_descriptor( + "sigma_schedule_h3", "rlkernel.h3_sigma_schedule.triton" + ) + assert reference is not None and triton is not None + assert reference.supported_devices == frozenset({"cpu", "cuda", "rocm"}) + # NVIDIA-only backend (inline PTX); HIP/ROCm stays on the reference. + assert triton.supported_devices == frozenset({"cuda"}) + assert reference.supported_dtypes == triton.supported_dtypes == frozenset({"float32"}) + assert reference.implementation_reference == _REFERENCE_PATH + + +def test_semantic_reference_path_resolves_instantiates_and_reports_provenance(): + session = kernel_registry.semantic.session() + resolution = session.resolve( + semantic_op="sigma_schedule_h3", + requested_backend="rlkernel.h3_sigma_schedule.reference", + target="rollout", + requirements=OperatorRequirements(device="cpu", dtype="float32"), + ) + assert resolution.trace.status == "resolved" + assert resolution.trace.concrete_backend == "rlkernel.h3_sigma_schedule.reference" + assert resolution.trace.implementation_reference == _REFERENCE_PATH + + instance = session.instantiate(resolution) + assert type(instance).__name__ == "NativeH3SigmaScheduleOp" + assert instance(4, shift=3.0, device="cpu").num_model_evals == 3 + + provenance = session.instance_provenance(resolution, instance) + assert provenance.semantic_op == "sigma_schedule_h3" + assert provenance.backend_id == "rlkernel.h3_sigma_schedule.reference" + assert provenance.concrete_implementation == _REFERENCE_PATH + assert provenance.factory_reference == _REFERENCE_PATH + + +def test_semantic_triton_backend_fails_explicitly_when_unavailable(monkeypatch): + # Resolution matches the declared capability (device "cuda"); the missing + # runtime dependency then surfaces as an explicit instantiation error — + # never as a silent swap to another backend's arithmetic. + monkeypatch.setattr(_BACKEND_MODULE, "_TRITON_AVAILABLE", False) + session = kernel_registry.semantic.session() + resolution = session.resolve( + semantic_op="sigma_schedule_h3", + requested_backend="rlkernel.h3_sigma_schedule.triton", + target="rollout", + requirements=OperatorRequirements(device="cuda", dtype="float32"), + ) + with pytest.raises(OperatorInstantiationError, match="triton is required"): + session.instantiate(resolution) + + +def test_semantic_unknown_backend_is_rejected(): + session = kernel_registry.semantic.session() + with pytest.raises(OperatorResolutionError): + session.resolve( + semantic_op="sigma_schedule_h3", + requested_backend="rlkernel.h3_sigma_schedule.does-not-exist", + target="rollout", + requirements=OperatorRequirements(device="cpu", dtype="float32"), + ) From 7961a28a42f9072e0774780f15ee8b5438fcb100 Mon Sep 17 00:00:00 2001 From: fxyan <1350821504@qq.com> Date: Wed, 21 Oct 2026 15:47:45 +0800 Subject: [PATCH 02/11] docs(h3): rebind evidence report and align ROCm scope Point the evidence report at 374488b instead of the amended-away 73e1aac, and describe the device backend as CUDA-only with the ROCm column open, matching the H3 stack. Refs #420 Signed-off-by: fxyan <1350821504@qq.com> --- docs/operators/h3-sigma-schedule.md | 8 ++++---- reports/experiments/h3-sigma-schedule/README.md | 4 ++-- .../experiments/h3-sigma-schedule/source-provenance.json | 2 +- rl_engine/backends/cuda/model_specific/__init__.py | 2 +- .../backends/cuda/model_specific/minimax_h3/__init__.py | 2 +- 5 files changed, 9 insertions(+), 9 deletions(-) diff --git a/docs/operators/h3-sigma-schedule.md b/docs/operators/h3-sigma-schedule.md index 49193fe2a..654367395 100644 --- a/docs/operators/h3-sigma-schedule.md +++ b/docs/operators/h3-sigma-schedule.md @@ -41,7 +41,7 @@ provenance = session.instance_provenance(resolution, op) | Backend | Wrapper | Native symbol | Status | | --- | --- | --- | --- | | CUDA (Triton) | `rl_engine.backends.cuda.model_specific.minimax_h3.sigma_schedule.H3SigmaScheduleCudaOp` | 3 Triton kernels (grid / dedup scan / timesteps) | Validated on RTX 4090D (torch 2.5.1+cu121, Triton 3.1.0): bitwise-identical to the pinned upstream oracle across the full device suite | -| ROCm | — | — | Not supported: the inline PTX `div.rn.f32` is NVIDIA-only and the backend refuses construction on HIP. The registry serves the reference on ROCm. | +| ROCm | — | — | **CUDA only.** The ROCm column stays open. The inline PTX `div.rn.f32` is NVIDIA-only, so the backend refuses construction on HIP; ROCm dispatches the PyTorch reference. | | PyTorch reference | `rl_engine.reference.minimax_h3.sigma_schedule.NativeH3SigmaScheduleOp` | `torch.linspace` + `unique_consecutive` | Active (byte-frozen oracle; fixture generated on Linux CPU under torch 2.14.1, reproduced on 2.4.1 and 2.5.1) | ## Tensor Contract @@ -152,9 +152,9 @@ itself bit-identical to the executed upstream method on those inputs. fp32-overflowing, or underflowing shifts. Bitwise parity for those edge cases is not portable across host architectures; the released shifts (3 and 12) do not take these paths. NVIDIA GPUs on ARM hosts have not been validated. -- ROCm is not supported: the Triton backend refuses construction on HIP - because its division is emitted as inline NVIDIA PTX. The registry serves the - PyTorch reference on ROCm. +- **CUDA only.** The ROCm column stays open: the Triton backend refuses + construction on HIP because its division is emitted as inline NVIDIA PTX; + ROCm dispatches the PyTorch reference. - Measured end-to-end latency on the 4090D: for the typical H3 sizes (N = 25/50/1000) the device-side Triton build is 2.5–3.2x SLOWER than building on CPU and copying to the GPU (B/C median 0.315–0.403, all separated diff --git a/reports/experiments/h3-sigma-schedule/README.md b/reports/experiments/h3-sigma-schedule/README.md index 71327badf..95907ff50 100644 --- a/reports/experiments/h3-sigma-schedule/README.md +++ b/reports/experiments/h3-sigma-schedule/README.md @@ -1,7 +1,7 @@ # H3 sigma schedule: RTX 4090D evidence These are historical operator measurements from October 16–17, 2026, -rechecked against candidate commit `73e1aac` on October 21. No new GPU run +rechecked against candidate commit `374488b` on October 21. No new GPU run was performed during the source-binding check. ## Correctness and source identity @@ -14,7 +14,7 @@ all outcomes; its hostname attribute was removed and the XML reserialized. [Source provenance](source-provenance.json) records SHA256 hashes for the tested archive, selected source files, and included results. The reference, CUDA backend, fixture, CPU/GPU/dispatch tests, and benchmark in the final -closing snapshot are byte-identical to those files in `73e1aac`. +closing snapshot are byte-identical to those files in `374488b`. The performance run (`run-20261016-235221-8055`) predates the final HIP rejection guard. Its reference and benchmark are byte-identical to the diff --git a/reports/experiments/h3-sigma-schedule/source-provenance.json b/reports/experiments/h3-sigma-schedule/source-provenance.json index e86f3b4d9..391daedfb 100644 --- a/reports/experiments/h3-sigma-schedule/source-provenance.json +++ b/reports/experiments/h3-sigma-schedule/source-provenance.json @@ -1,6 +1,6 @@ { "evidence_kind": "Historical hardware results; source binding rechecked 2026-10-21; no new GPU run", - "candidate_commit": "73e1aac95a6c3e0745792b83733ba79432c12b3b", + "candidate_commit": "374488bb3abc856b356b79fbfbcc0e15bbb968c4", "historical_snapshot_base_commit": "0e8a18aff441895082aee603959cd25cb66d7e4b", "environment": { "python": "3.12.3", diff --git a/rl_engine/backends/cuda/model_specific/__init__.py b/rl_engine/backends/cuda/model_specific/__init__.py index 56bc59826..9ed43ae29 100644 --- a/rl_engine/backends/cuda/model_specific/__init__.py +++ b/rl_engine/backends/cuda/model_specific/__init__.py @@ -1,4 +1,4 @@ # SPDX-License-Identifier: Apache-2.0 # Copyright (c) 2026 RL-Kernel Contributors -"""Triton CUDA/HIP backend for model-specific operators (RFC #420 MiniMax-H3).""" +"""Triton CUDA backend for model-specific operators (RFC #420 MiniMax-H3).""" diff --git a/rl_engine/backends/cuda/model_specific/minimax_h3/__init__.py b/rl_engine/backends/cuda/model_specific/minimax_h3/__init__.py index bd451c764..572c4ac13 100644 --- a/rl_engine/backends/cuda/model_specific/minimax_h3/__init__.py +++ b/rl_engine/backends/cuda/model_specific/minimax_h3/__init__.py @@ -1,7 +1,7 @@ # SPDX-License-Identifier: Apache-2.0 # Copyright (c) 2026 RL-Kernel Contributors -"""MiniMax-H3 (RFC #420) model-specific CUDA/HIP backends.""" +"""MiniMax-H3 (RFC #420) model-specific CUDA backends.""" from rl_engine.backends.cuda.model_specific.minimax_h3.sigma_schedule import H3SigmaScheduleCudaOp From ed316d863f9ceb243a21f531f62e9dad20434e60 Mon Sep 17 00:00:00 2001 From: fxyan <1350821504@qq.com> Date: Wed, 21 Oct 2026 16:14:44 +0800 Subject: [PATCH 03/11] docs(h3): fix sigma evidence report link Signed-off-by: fxyan <1350821504@qq.com> --- docs/operators/h3-sigma-schedule.md | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/docs/operators/h3-sigma-schedule.md b/docs/operators/h3-sigma-schedule.md index 654367395..5296ccb6d 100644 --- a/docs/operators/h3-sigma-schedule.md +++ b/docs/operators/h3-sigma-schedule.md @@ -127,7 +127,7 @@ within dispersion, so it is not claimed. The N=1e6 CPU medians are noisy (B stdev ≈ 44 ms at shift=12), so the 51.98x ratio passes the dispersion check only by a small margin. See Known Limitations and the raw JSON, CUDA test results, source hashes, and reproduction commands in the -[4090D evidence report](../../reports/experiments/h3-sigma-schedule/). +[4090D evidence report](https://github.com/RL-Align/RL-Kernel/tree/test-h3/reports/experiments/h3-sigma-schedule). ## Tests From a66c999e0ccfc1362943c692a06846c7075c94ac Mon Sep 17 00:00:00 2001 From: fxyan <1350821504@qq.com> Date: Sun, 25 Oct 2026 23:41:34 +0800 Subject: [PATCH 04/11] refactor(h3): default sigma_schedule_h3 to the PyTorch reference on CUDA Default backend change per review (PR #566, registry.py inline comment): the RTX 4090D evidence shows the device-side Triton build is 2.5-3.2x slower than CPU construction plus host-to-device copy at typical H3 sizes (N=25/50/1000), and nothing in RFC #420 or any consumer requires on-device construction. The legacy "cuda" candidate list now carries only PYTORCH_H3_SIGMA_SCHEDULE, matching every other platform; the Triton backend stays available as an explicit opt-in through the semantic requested_backend="rlkernel.h3_sigma_schedule.triton" or direct construction of H3SigmaScheduleCudaOp. Unchanged: the CPU reference, the Triton kernels and forward(), the upstream fixture, and every existing bitwise assertion. The CUDA suite gains an on-hardware assertion that the default reference path and the Triton path produce identical sigma/timestep bits; the dispatch tests pin the reference-only list on all platforms. Docs, benchmark docstring and the evidence report describe the new default; source-provenance.json is left as recorded (three test/benchmark hashes now stale) pending the SM90 evidence run. CPU verification on macOS (no GPU): 32 passed / 1 skipped for tests/reference/minimax_h3 + dispatch; CUDA suite 36 skipped. black, isort, flake8 (pre-commit args), mypy (3.10 target) and mkdocs --strict clean. CUDA hardware tests were not run on this host. Signed-off-by: fxyan <1350821504@qq.com> --- .../models/benchmark_h3_sigma_schedule.py | 7 +- docs/architecture/runtime-dispatch.md | 9 ++ docs/operators/h3-sigma-schedule.md | 88 +++++++++++++------ .../experiments/h3-sigma-schedule/README.md | 23 ++++- rl_engine/runtime/registry.py | 17 ++-- .../cuda/test_h3_sigma_schedule_cuda.py | 48 +++++++--- .../test_h3_sigma_schedule_dispatch.py | 63 ++++++++----- 7 files changed, 185 insertions(+), 70 deletions(-) diff --git a/benchmarks/models/benchmark_h3_sigma_schedule.py b/benchmarks/models/benchmark_h3_sigma_schedule.py index f68237f0d..04da2af01 100644 --- a/benchmarks/models/benchmark_h3_sigma_schedule.py +++ b/benchmarks/models/benchmark_h3_sigma_schedule.py @@ -8,9 +8,10 @@ * path A — CPU reference, grid built on CPU, outputs stay on CPU; * path B — CPU reference, grid built on CPU, sigmas and timesteps moved to - the target GPU via the operator's ``device`` argument (what a user without - the GPU backend gets: build + host-to-device copy); -* path C — Triton candidate, grid built on the device. + the target GPU via the operator's ``device`` argument (the registry default + path on every platform: build + host-to-device copy); +* path C — Triton opt-in backend (semantic ``requested_backend`` or direct + construction), grid built on the device. B vs C is the comparison that matters when the consumer lives on the GPU; A is an auxiliary CPU baseline and A/C ratios must not be presented as a diff --git a/docs/architecture/runtime-dispatch.md b/docs/architecture/runtime-dispatch.md index c1321a496..b3ee8b9c0 100644 --- a/docs/architecture/runtime-dispatch.md +++ b/docs/architecture/runtime-dispatch.md @@ -37,6 +37,15 @@ the legacy SM90 LogP backend when `RL_KERNEL_ENABLE_EXPERIMENTAL_SM90_LOGP=1` is set. The fused linear logp SM90 backend is gated separately and remains the default linear logp backend when the extension is built on Hopper. +## H3 Sigma Schedule + +`sigma_schedule_h3` lists only the PyTorch reference on every platform, CUDA +included; the NVIDIA Triton backend is opt-in through the semantic +`requested_backend="rlkernel.h3_sigma_schedule.triton"` or direct construction, +because its device-side build measured 2.5–3.2x slower than CPU construction +plus transfer at typical H3 sizes (PR #566 evidence). See +[H3 Sigma Schedule](../operators/h3-sigma-schedule.md). + ## Relevant Files - `rl_engine/runtime/registry.py` diff --git a/docs/operators/h3-sigma-schedule.md b/docs/operators/h3-sigma-schedule.md index 5296ccb6d..9a6404f66 100644 --- a/docs/operators/h3-sigma-schedule.md +++ b/docs/operators/h3-sigma-schedule.md @@ -1,25 +1,35 @@ # H3 Sigma Schedule -Device-side construction of the MiniMax-H3 rectified-flow sigma grid (RFC #420, -`h3_sigma_schedule`): `linspace(1, 0, N)` in fp32, the exponential shift +The MiniMax-H3 rectified-flow sigma grid (RFC #420, `h3_sigma_schedule`): +`linspace(1, 0, N)` in fp32, the exponential shift `shift·s / (1 + (shift−1)·s)`, `unique_consecutive` deduplication, and `timesteps = 1 − sigmas[:-1]`. Video uses `shift=12`, audio `shift=3`; the two -grids are independent per-request inputs, never a shared step index. +grids are independent per-request inputs, never a shared step index. The +default backend on every platform is the PyTorch reference (CPU construction, +outputs copied to the requested device); a device-side NVIDIA Triton backend +is available as an explicit opt-in. ## Entry Point ```python from rl_engine.runtime.registry import kernel_registry -op = kernel_registry.get_op("sigma_schedule_h3", device="cuda") -schedule = op(25, shift=12.0, device="cuda") # device= selects backend AND output +op = kernel_registry.get_op("sigma_schedule_h3", device="cuda") # PyTorch reference +schedule = op(25, shift=12.0, device="cuda") # device= selects output placement only schedule.sigmas # fp32 [M], strictly decreasing, terminal exactly 0.0 schedule.timesteps # fp32 [M-1] = 1 - sigmas[:-1] schedule.num_model_evals # M - 1 (deduplicated length drives the count) ``` -The semantic path resolves with provenance (`requested_backend` -`rlkernel.h3_sigma_schedule.triton` / `.reference`): +`get_op` returns the PyTorch reference on every platform, CUDA included: the +grid is built on CPU and both outputs are moved to `device`. + +## Opt-in: device-side Triton backend + +Two explicit routes select `H3SigmaScheduleCudaOp`; neither is reachable +through a legacy `get_op` candidate list. The semantic path resolves with +provenance (`requested_backend` `rlkernel.h3_sigma_schedule.triton`; the +reference is `.reference`): ```python from rl_engine.runtime.registry import kernel_registry @@ -36,13 +46,27 @@ op = session.instantiate(resolution) provenance = session.instance_provenance(resolution, op) ``` +Or construct the class directly: + +```python +from rl_engine.backends.cuda.model_specific.minimax_h3.sigma_schedule import ( + H3SigmaScheduleCudaOp, +) + +op = H3SigmaScheduleCudaOp() # raises ImportError without Triton or on HIP/ROCm +schedule = op(25, shift=12.0, device="cuda") +``` + +Both routes produce the same bytes as the default reference path (the CUDA +suite asserts this); opting in only moves the arithmetic onto the device. + ## Backends | Backend | Wrapper | Native symbol | Status | | --- | --- | --- | --- | -| CUDA (Triton) | `rl_engine.backends.cuda.model_specific.minimax_h3.sigma_schedule.H3SigmaScheduleCudaOp` | 3 Triton kernels (grid / dedup scan / timesteps) | Validated on RTX 4090D (torch 2.5.1+cu121, Triton 3.1.0): bitwise-identical to the pinned upstream oracle across the full device suite | +| CUDA (Triton) | `rl_engine.backends.cuda.model_specific.minimax_h3.sigma_schedule.H3SigmaScheduleCudaOp` | 3 Triton kernels (grid / dedup scan / timesteps) | **Opt-in** (semantic `requested_backend` or direct construction; not a legacy `get_op` candidate). Validated on RTX 4090D (torch 2.5.1+cu121, Triton 3.1.0): bitwise-identical to the pinned upstream oracle across the full device suite | | ROCm | — | — | **CUDA only.** The ROCm column stays open. The inline PTX `div.rn.f32` is NVIDIA-only, so the backend refuses construction on HIP; ROCm dispatches the PyTorch reference. | -| PyTorch reference | `rl_engine.reference.minimax_h3.sigma_schedule.NativeH3SigmaScheduleOp` | `torch.linspace` + `unique_consecutive` | Active (byte-frozen oracle; fixture generated on Linux CPU under torch 2.14.1, reproduced on 2.4.1 and 2.5.1) | +| PyTorch reference | `rl_engine.reference.minimax_h3.sigma_schedule.NativeH3SigmaScheduleOp` | `torch.linspace` + `unique_consecutive` | **Default on every platform** (byte-frozen oracle; fixture generated on Linux CPU under torch 2.14.1, reproduced on 2.4.1 and 2.5.1). Builds on CPU and copies the outputs to the requested device | ## Tensor Contract @@ -50,22 +74,27 @@ provenance = session.instance_provenance(resolution, op) | --- | --- | --- | --- | | `num_inference_steps` | scalar | int | `>= 2`; terminal zero included in the count | | `shift` | scalar | python float | `> 0` (video 12.0 / audio 3.0) | -| `device` | — | — | CUDA backend requires an NVIDIA CUDA device (NVIDIA-only; refuses construction on HIP/ROCm); CPU requests use the reference | +| `device` | — | — | Output placement. The default reference accepts any device (builds on CPU, copies to `device`); the opt-in Triton backend requires an NVIDIA CUDA device (refuses construction on HIP/ROCm) | | → `sigmas` | `[M]`, `M ≤ N` | float32 | strictly decreasing; `sigmas[0] == 1.0` and `sigmas[-1] == 0.0` exactly for released shifts | | → `timesteps` | `[M-1]` | float32 | `1 − sigmas[:-1]`, strictly increasing from 0.0 | | → `num_model_evals` | scalar | int | `M − 1` | ## Dispatch Behavior -The registry walks the CUDA candidate list -`[Triton H3, PyTorch reference]`; on ROCm/MUSA/NPU/CPU only the reference is -registered. The Triton backend raises `ImportError` at construction -when Triton is not importable, so the first-available-candidate loop skips it -at selection time (see -[Runtime Dispatch](../architecture/runtime-dispatch.md)) — forward-time errors -are never used for selection. When the Triton backend is skipped on CUDA, the -reference builds the grid on CPU and moves the outputs to the requested CUDA -device; the result stays device-resident, but the arithmetic is the CPU path. +The legacy candidate list for `sigma_schedule_h3` is `[PyTorch reference]` on +every platform — CUDA, ROCm, MUSA, NPU and CPU alike — so `get_op` always +builds the grid on CPU and moves the outputs to the requested device. This is a +deliberate default (review 2026-10-25, PR #566): the 4090D measurements below +show the device-side build is 2.5–3.2x slower at typical H3 sizes, and no +consumer requires on-device construction. The Triton backend is not a legacy +candidate anywhere; it is selected only through the semantic +`requested_backend="rlkernel.h3_sigma_schedule.triton"` or by direct +construction (see Opt-in above). Its construction-time guards are unchanged — +`H3SigmaScheduleCudaOp()` raises `ImportError` when Triton is not importable +or on a HIP/ROCm build, so an opt-in that cannot be honored fails explicitly +(semantic `instantiate` raises `OperatorInstantiationError`) instead of +silently swapping arithmetic (see +[Runtime Dispatch](../architecture/runtime-dispatch.md)). ## Accuracy @@ -106,8 +135,8 @@ python benchmarks/models/benchmark_h3_sigma_schedule.py --samples 100 Measures **per-call end-to-end host wall latency** on three paths — A: reference built on CPU, outputs on CPU; B: reference built on CPU, sigmas and -timesteps moved to the GPU via `device="cuda"` (what a consumer without the -GPU backend gets); C: the Triton backend building on the device. Each call is +timesteps moved to the GPU via `device="cuda"` (the registry default path on +every platform); C: the opt-in Triton backend building on the device. Each call is timed individually (GPU paths synchronize before each sample and after each call); raw samples, mean/median/stdev/min/max are reported per case, and the median is the headline number. Before any timing on a GPU host the script @@ -129,6 +158,13 @@ check only by a small margin. See Known Limitations and the raw JSON, CUDA test results, source hashes, and reproduction commands in the [4090D evidence report](https://github.com/RL-Align/RL-Kernel/tree/test-h3/reports/experiments/h3-sigma-schedule). +These numbers are the basis for the default: the reference path (B) is the +legacy `get_op` backend on CUDA because it is 2.5–3.2x faster at every typical +H3 size and produces identical bytes. The trade reverses only at the N=1e6 +dedup stress case — at shift=12 the reference median is 48.37 ms against +0.93 ms for Triton (51.98x), so a consumer building grids of that size should +opt into the Triton backend explicitly; it is not a released-model workload. + ## Tests ```bash @@ -158,9 +194,11 @@ itself bit-identical to the executed upstream method on those inputs. - Measured end-to-end latency on the 4090D: for the typical H3 sizes (N = 25/50/1000) the device-side Triton build is 2.5–3.2x SLOWER than building on CPU and copying to the GPU (B/C median 0.315–0.403, all separated - from dispersion). The device path wins only at the N = 1e6 dedup stress - case: 51.98x at shift=12 (separated) and 7.9x at shift=3 (within dispersion, - not claimed). These numbers are single-operator microbenchmark results on - one card and say nothing about whole-model speed. + from dispersion) — which is why it is opt-in rather than the default. The + device path wins only at the N = 1e6 dedup stress case: 51.98x at shift=12 + (separated) and 7.9x at shift=3 (within dispersion, not claimed); the default + reference path is correspondingly slower there (48.37 ms vs 0.93 ms at + shift=12). These numbers are single-operator microbenchmark results on one + card and say nothing about whole-model speed. - `N` beyond 2^31 is not supported (int32 indexing); the shift arithmetic is internal-fp64, which is slow on consumer GPUs but irrelevant at these sizes. diff --git a/reports/experiments/h3-sigma-schedule/README.md b/reports/experiments/h3-sigma-schedule/README.md index 95907ff50..2ed71b14b 100644 --- a/reports/experiments/h3-sigma-schedule/README.md +++ b/reports/experiments/h3-sigma-schedule/README.md @@ -13,8 +13,24 @@ all outcomes; its hostname attribute was removed and the XML reserialized. [Source provenance](source-provenance.json) records SHA256 hashes for the tested archive, selected source files, and included results. The reference, -CUDA backend, fixture, CPU/GPU/dispatch tests, and benchmark in the final -closing snapshot are byte-identical to those files in `374488b`. +CUDA backend, fixture and CPU reference test in the final closing snapshot are +byte-identical to those files in `374488b` and to the current branch. + +The dispatch test, the CUDA test and the benchmark docstring were revised +after the 4090D runs for the review-driven default backend change (the legacy +`sigma_schedule_h3` candidate list is now reference-only on CUDA; the Triton +backend is opt-in through the semantic `requested_backend` or direct +construction). Those edits touch candidate order, test selection logic and +prose only: every kernel bit assertion is unchanged, and the CUDA suite gained +an assertion that the default reference path matches the Triton path bit for +bit. No GPU run backs the revised tests yet. `source-provenance.json` still +carries the `374488b` hashes; its `source_files_sha256` entries for +`tests/runtime/test_h3_sigma_schedule_dispatch.py`, +`tests/backends/cuda/test_h3_sigma_schedule_cuda.py` and +`benchmarks/models/benchmark_h3_sigma_schedule.py`, and the benchmark +`binding` note ("Benchmark and reference byte-identical"), are therefore stale. +They are deliberately left as recorded and will be refreshed together with the +pending SM90 (Hopper) evidence run rather than piecemeal. The performance run (`run-20261016-235221-8055`) predates the final HIP rejection guard. Its reference and benchmark are byte-identical to the @@ -30,7 +46,8 @@ Hardware/software: one NVIDIA GeForce RTX 4090D, x86_64 Linux, driver [Raw benchmark JSON](benchmark-4090d.json) is copied without modification. It contains 100 samples per case and path (2,400 samples total), after five warmup calls. Path A builds on CPU and leaves outputs there; path B -builds on CPU and copies both outputs to CUDA; path C constructs on CUDA. +builds on CPU and copies both outputs to CUDA (the registry default path); +path C constructs on CUDA (the Triton opt-in backend). Only B/C compares equal output placement. Timing is per-call host wall latency with synchronization before and after each GPU sample; correctness checking and compilation happen before timed samples. diff --git a/rl_engine/runtime/registry.py b/rl_engine/runtime/registry.py index b1af59756..218979b25 100644 --- a/rl_engine/runtime/registry.py +++ b/rl_engine/runtime/registry.py @@ -204,7 +204,12 @@ class OpBackend(Enum, metaclass=_KernelEnumMeta): CUDA_SM90_LM_HEAD = "rl_engine.kernels.ops.cuda.linear.lm_head.SM90LMHeadOp" CUDA_SM90_EMBEDDING = "rl_engine.kernels.ops.cuda.linear.embedding.SM90EmbeddingOp" - # MiniMax-H3 (RFC #420) scheduler grid + # MiniMax-H3 (RFC #420) scheduler grid. The PyTorch reference is the + # legacy ``get_op`` default on every platform; the Triton CUDA backend is + # opt-in only — select it through the semantic path with + # requested_backend="rlkernel.h3_sigma_schedule.triton", or construct + # H3SigmaScheduleCudaOp directly (default backend change per review, + # PR #566). PYTORCH_H3_SIGMA_SCHEDULE = ( "rl_engine.reference.minimax_h3.sigma_schedule.NativeH3SigmaScheduleOp" ) @@ -726,10 +731,12 @@ def __init__(self): OpBackend.TRITON_ROPE, OpBackend.PYTORCH_NATIVE_ROPE, ], - "sigma_schedule_h3": [ - OpBackend.CUDA_H3_SIGMA_SCHEDULE, - OpBackend.PYTORCH_H3_SIGMA_SCHEDULE, - ], + # Reference only, like every other platform: the device-side + # Triton build is slower than CPU build + copy at typical H3 + # sizes and nothing requires on-device construction. Triton is + # opt-in via the semantic path or direct construction (see + # docs/operators/h3-sigma-schedule.md). + "sigma_schedule_h3": [OpBackend.PYTORCH_H3_SIGMA_SCHEDULE], }, "rocm": { "logp": [ diff --git a/tests/backends/cuda/test_h3_sigma_schedule_cuda.py b/tests/backends/cuda/test_h3_sigma_schedule_cuda.py index 0f66f91ae..72690ae88 100644 --- a/tests/backends/cuda/test_h3_sigma_schedule_cuda.py +++ b/tests/backends/cuda/test_h3_sigma_schedule_cuda.py @@ -9,8 +9,12 @@ The backend is NVIDIA-only (its kernels emit inline PTX ``div.rn.f32``); the suite runs on NVIDIA CUDA hardware and skips on HIP/ROCm builds, where -``H3SigmaScheduleCudaOp`` refuses construction and the registry serves the -PyTorch reference instead. +``H3SigmaScheduleCudaOp`` refuses construction. It is opt-in everywhere: the +legacy registry serves the PyTorch reference on CUDA too (default backend +change per review, 2026-10-25), and the Triton class is reached through the +semantic ``requested_backend`` or direct construction. The suite therefore +constructs the backend directly for kernel checks and asserts that the default +reference path matches it bit for bit. """ from __future__ import annotations @@ -306,8 +310,9 @@ def test_backend_identity_is_device_path_not_cpu_fallback(): ) from rl_engine.runtime.registry import OpBackend - # The registry lists the NVIDIA Triton implementation ahead of the PyTorch - # reference for device "cuda"; outputs above prove the device path ran. + # Both backends stay registered as enum members; the legacy "cuda" list + # carries only the reference (Triton is opt-in). The outputs above came + # from the directly constructed Triton class, proving the device path ran. assert OpBackend.CUDA_H3_SIGMA_SCHEDULE.value.endswith( "model_specific.minimax_h3.sigma_schedule.H3SigmaScheduleCudaOp" ) @@ -320,19 +325,40 @@ def test_backend_identity_is_device_path_not_cpu_fallback(): assert torch.version.hip is None # NVIDIA-only backend -def test_public_entry_selects_device_backend_and_reports_provenance(): - # The registry must hand out the Triton class itself (device identity is - # the instance type, not just `.is_cuda`, which a CPU-build-then-copy - # would also have), and the semantic path must report the same backend. +def test_public_entry_defaults_to_reference_and_matches_triton_bits(): + # Default backend change per review (2026-10-25): legacy ``get_op`` hands + # out the PyTorch reference on CUDA devices (CPU build + host-to-device + # copy) and the Triton class is opt-in. The switch is only safe because the + # two paths are bit-identical — pin that equality here, on hardware, for + # both outputs and the evaluation count. from rl_engine.backends.cuda.model_specific.minimax_h3.sigma_schedule import ( H3SigmaScheduleCudaOp, ) + from rl_engine.reference.minimax_h3.sigma_schedule import NativeH3SigmaScheduleOp from rl_engine.runtime.registry import kernel_registry - from rl_engine.runtime.semantic_registry import OperatorRequirements op = kernel_registry.get_op("sigma_schedule_h3", device="cuda") - assert isinstance(op, H3SigmaScheduleCudaOp) - assert op(25, shift=12.0).sigmas.is_cuda + assert isinstance(op, NativeH3SigmaScheduleOp) + + default = op(25, shift=12.0, device="cuda") + device_side = H3SigmaScheduleCudaOp()(25, shift=12.0) + assert default.sigmas.is_cuda and default.timesteps.is_cuda + assert _int32_bits(default.sigmas.cpu()) == _int32_bits(device_side.sigmas.cpu()) + assert _int32_bits(default.timesteps.cpu()) == _int32_bits(device_side.timesteps.cpu()) + assert default.num_model_evals == device_side.num_model_evals + assert torch.version.hip is None # NVIDIA-only backend + + +def test_public_entry_opt_in_selects_triton_and_reports_provenance(): + # The explicit opt-in is unchanged: the semantic path must hand out the + # Triton class itself (device identity is the instance type, not just + # `.is_cuda`, which the CPU-build-then-copy default also has) and report + # that backend in the provenance. + from rl_engine.backends.cuda.model_specific.minimax_h3.sigma_schedule import ( + H3SigmaScheduleCudaOp, + ) + from rl_engine.runtime.registry import kernel_registry + from rl_engine.runtime.semantic_registry import OperatorRequirements session = kernel_registry.semantic.session() resolution = session.resolve( diff --git a/tests/runtime/test_h3_sigma_schedule_dispatch.py b/tests/runtime/test_h3_sigma_schedule_dispatch.py index 47e4e1f84..fd05062d1 100644 --- a/tests/runtime/test_h3_sigma_schedule_dispatch.py +++ b/tests/runtime/test_h3_sigma_schedule_dispatch.py @@ -3,12 +3,13 @@ """Registry and semantic-entry tests for ``h3_sigma_schedule`` (RFC #420). -Covers the public selection paths on any machine: legacy ``get_op`` candidate -fallback when the Triton backend cannot construct, semantic descriptor -registration, session resolve / instantiate / provenance for the reference -backend, and the CPU-pinned scalar-rounding helper. The device-path identity -(legacy ``get_op`` returning the Triton class, plus result bytes) is asserted -on hardware in tests/backends/cuda/test_h3_sigma_schedule_cuda.py. +Covers the public selection paths on any machine: legacy ``get_op`` serving +the PyTorch reference on every platform (the Triton backend is opt-in only and +never a legacy candidate), semantic descriptor registration, session resolve / +instantiate / provenance for both backends, and the CPU-pinned scalar-rounding +helper. The device-path parity (the default reference path producing the same +bytes as the directly constructed Triton class) is asserted on hardware in +tests/backends/cuda/test_h3_sigma_schedule_cuda.py. """ from __future__ import annotations @@ -18,7 +19,7 @@ import pytest import torch -from rl_engine.runtime.registry import KernelRegistry, kernel_registry +from rl_engine.runtime.registry import KernelRegistry, OpBackend, kernel_registry from rl_engine.runtime.semantic_registry import ( OperatorInstantiationError, OperatorRequirements, @@ -45,18 +46,22 @@ def test_get_op_cpu_selects_reference_and_builds_pinned_grid(): assert schedule.num_model_evals == 3 -def test_get_op_cuda_without_triton_skips_unavailable_backend(monkeypatch): - # Old behavior (review 2026-10-13): the unavailable GPU class was selected - # and failed only inside forward. Construction now raises ImportError, so - # the registry's first-available-candidate loop skips it at selection time - # (docs/architecture/runtime-dispatch.md) instead of the registry having - # to catch arbitrary forward errors. - monkeypatch.setattr(_BACKEND_MODULE, "_TRITON_AVAILABLE", False) - registry = KernelRegistry() +def test_get_op_cuda_defaults_to_reference_regardless_of_triton(monkeypatch): + # Default backend change per review (2026-10-25): the CUDA candidate list + # carries only the PyTorch reference (CPU build + host-to-device copy) — + # the device-side Triton build measured 2.5-3.2x slower at typical H3 + # sizes and nothing requires on-device construction. Triton is opt-in via + # the semantic path or direct construction, so whether it is importable + # must not influence the legacy selection in either direction. + for triton_importable in (False, True): + monkeypatch.setattr(_BACKEND_MODULE, "_TRITON_AVAILABLE", triton_importable) + registry = KernelRegistry() - op = registry.get_op("sigma_schedule_h3", device="cuda") - - assert f"{type(op).__module__}.{type(op).__name__}" == _REFERENCE_PATH + assert registry._priority_map["cuda"]["sigma_schedule_h3"] == [ + OpBackend.PYTORCH_H3_SIGMA_SCHEDULE + ] + op = registry.get_op("sigma_schedule_h3", device="cuda") + assert f"{type(op).__module__}.{type(op).__name__}" == _REFERENCE_PATH def test_cuda_backend_construction_fails_loudly_without_triton(monkeypatch): @@ -95,11 +100,12 @@ def test_cuda_backend_construction_rejected_on_hip_build(monkeypatch): _BACKEND_MODULE.H3SigmaScheduleCudaOp() -def test_get_op_falls_back_to_reference_on_hip_build(monkeypatch): - # End-to-end selection semantics under the same simulation: the registry - # walks the cuda candidate list, the Triton class refuses construction - # with ImportError, and the PyTorch reference is served instead — no - # error reaches the caller and nothing is silently misattributed. +def test_get_op_on_hip_build_serves_reference(monkeypatch): + # Same simulation, end to end: ``device="cuda"`` resolves to the rocm + # platform, whose list has always been reference-only, and the NVIDIA-only + # Triton class is not a legacy candidate anywhere — the caller gets the + # reference with no error and nothing silently misattributed, exactly as + # on an NVIDIA build. monkeypatch.setattr(torch.version, "hip", "6.2") registry = KernelRegistry() @@ -108,6 +114,17 @@ def test_get_op_falls_back_to_reference_on_hip_build(monkeypatch): assert f"{type(op).__module__}.{type(op).__name__}" == _REFERENCE_PATH +def test_triton_backend_is_not_a_legacy_candidate_on_any_platform(): + # Opt-in only: the Triton class is reachable through the semantic path + # (requested_backend="rlkernel.h3_sigma_schedule.triton") or direct + # construction, never through a legacy ``get_op`` candidate list. + registry = KernelRegistry() + + for platform in ("cuda", "rocm", "musa", "cpu", "npu"): + candidates = registry._priority_map[platform]["sigma_schedule_h3"] + assert candidates == [OpBackend.PYTORCH_H3_SIGMA_SCHEDULE], platform + + def test_semantic_descriptors_declared_truthfully(): reference = kernel_registry.semantic.backend_descriptor( "sigma_schedule_h3", "rlkernel.h3_sigma_schedule.reference" From ace9c20cbe6ca5a85768d228fe1fd3084e728aa6 Mon Sep 17 00:00:00 2001 From: fxyan <1350821504@qq.com> Date: Mon, 26 Oct 2026 02:15:34 +0800 Subject: [PATCH 05/11] fix(h3): address CodeRabbit review on sigma schedule benchmark and docs MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Benchmark: on a GPU host without Triton, path B (CPU build + copy, the registry default path) no longer shares path C's skip — it is measured and only path C is skipped. On a HIP/ROCm host where Triton imports, the NVIDIA-only backend is now recorded as unsupported before the correctness gate instead of the gate's construction ImportError crashing the script; path B is measured there too. The NVIDIA+Triton path (gate, then B, then C), timing semantics, sample counts and JSON keys are unchanged — the H800 evidence run exercised exactly that path. _summarize treats a path marked skipped or unsupported as absent instead of indexing into it. Docs: Dispatch Behavior now spells out the two selection layers — legacy candidate fallback vs. the semantic path's strict no-fallback resolve/ instantiate — and that the semantic reference descriptor declares cpu/cuda/rocm only (legacy lists also serve MUSA/NPU). Entry Point notes the per-backend device defaults when device= is omitted (reference -> CPU tensors, Triton -> cuda); both defaults are intentional and unchanged. The new GPU-without-Triton and HIP branches were verified by CPU-only execution (--samples 3), _summarize simulations and static review only; they have not run on GPU or HIP hardware. Signed-off-by: fxyan <1350821504@qq.com> --- .../models/benchmark_h3_sigma_schedule.py | 82 +++++++++++++------ docs/operators/h3-sigma-schedule.md | 24 +++++- 2 files changed, 81 insertions(+), 25 deletions(-) diff --git a/benchmarks/models/benchmark_h3_sigma_schedule.py b/benchmarks/models/benchmark_h3_sigma_schedule.py index 04da2af01..45184db0e 100644 --- a/benchmarks/models/benchmark_h3_sigma_schedule.py +++ b/benchmarks/models/benchmark_h3_sigma_schedule.py @@ -37,17 +37,21 @@ classes directly (registry bypassed for timing); registry-entry identity and correctness are pinned by tests/runtime/ and tests/backends/cuda/. -Before any timing on a GPU host, path C is verified bitwise against path A -for every measured case (sigmas, timesteps, and num_model_evals); a -mismatch aborts with exit code 2 and no performance numbers are produced. +Before any timing on a GPU host where the Triton backend is available, path C +is verified bitwise against path A for every measured case (sigmas, +timesteps, and num_model_evals); a mismatch aborts with exit code 2 and no +performance numbers are produced. Exit codes: 0 on success; 2 on the pre-timing correctness gate; nonzero on invalid arguments, and on any Triton compile, launch, or runtime failure -after the backend was confirmed available. The Triton section is skipped -only for a *confirmed* absent backend: no CUDA/HIP device -(``torch.cuda.is_available()`` false) or a confirmed missing Triton import -on a GPU host. The script runs from an uninstalled checkout (it bootstraps -the repository root onto ``sys.path``). +after the backend was confirmed available. Path B needs only torch and a +device, so it is measured on every GPU host. Path C (and its gate) is +skipped only for a *confirmed* absent backend — a missing Triton import on a +GPU host — or recorded as unsupported on a HIP/ROCm build, where the +NVIDIA-only kernels (inline PTX) refuse construction. Without any CUDA/HIP +device (``torch.cuda.is_available()`` false) both B and C are skipped. The +script runs from an uninstalled checkout (it bootstraps the repository root +onto ``sys.path``). The grid is tiny (4 KB at N=1000), so GPU launch overhead may exceed the arithmetic — the numbers are what they are, and this benchmark does not @@ -212,7 +216,16 @@ def _summarize(report: dict) -> dict: summary = {"definition": _SPEEDUP_NOTE, "cases": {}} path_b = report.get("path_b_reference_cpu_build_then_copy_to_gpu") path_c = report.get("path_c_triton_device_build") - if not path_b or not path_c or "skipped" in path_b or "skipped" in path_c: + # A path that was skipped or recorded as unsupported carries no per-case + # timings, so there is no ratio to compute (B alone is still reported). + if ( + not path_b + or not path_c + or "skipped" in path_b + or "unsupported" in path_b + or "skipped" in path_c + or "unsupported" in path_c + ): return summary for key, case_b in path_b.items(): case_c = path_c[key] @@ -271,26 +284,47 @@ def main(argv: list[str] | None = None) -> int: report["triton"] = triton.__version__ except ImportError as exc: + # Path B is the registry default path and needs only torch plus a + # device, so a missing Triton skips path C alone. skip = f"triton import failed on a GPU host: {exc}" - report["path_b_reference_cpu_build_then_copy_to_gpu"] = {"skipped": skip} - report["path_c_triton_device_build"] = {"skipped": skip} - print(f"paths B and C: skipped — triton missing ({exc})") - else: - print("verifying path C against path A bitwise before timing...") - gate_ms = _verify_triton_matches_reference() - report["correctness_gate_wall_ms"] = gate_ms - print( - "correctness gate passed " - f"(includes Triton kernel compilation: {gate_ms:.1f} ms wall, " - "reported once, excluded from all timed samples)" - ) - print(f"path B (reference, cpu build, outputs moved to gpu): {args.samples} calls") report["path_b_reference_cpu_build_then_copy_to_gpu"] = _time_reference( args.samples, args.warmup, device="cuda" ) - print(f"path C (triton, device build): {args.samples} calls") - report["path_c_triton_device_build"] = _time_triton(args.samples, args.warmup) + report["path_c_triton_device_build"] = {"skipped": skip} + print(f"path C: skipped — triton missing ({exc})") + else: + if torch.version.hip is not None: + # Triton may import on a HIP host, but the backend is NVIDIA-only + # (inline PTX) and refuses construction: record that instead of + # letting the ImportError escape from inside the correctness gate. + unsupported = ( + "HIP/ROCm build: the Triton backend is NVIDIA-only (inline PTX " + "div.rn.f32) and refuses construction; the bitwise correctness " + "gate was not run" + ) + print(f"path B (reference, cpu build, outputs moved to gpu): {args.samples} calls") + report["path_b_reference_cpu_build_then_copy_to_gpu"] = _time_reference( + args.samples, args.warmup, device="cuda" + ) + report["path_c_triton_device_build"] = {"unsupported": unsupported} + print("path C: unsupported — NVIDIA-only backend on a HIP/ROCm build") + else: + print("verifying path C against path A bitwise before timing...") + gate_ms = _verify_triton_matches_reference() + report["correctness_gate_wall_ms"] = gate_ms + print( + "correctness gate passed " + f"(includes Triton kernel compilation: {gate_ms:.1f} ms wall, " + "reported once, excluded from all timed samples)" + ) + + print(f"path B (reference, cpu build, outputs moved to gpu): {args.samples} calls") + report["path_b_reference_cpu_build_then_copy_to_gpu"] = _time_reference( + args.samples, args.warmup, device="cuda" + ) + print(f"path C (triton, device build): {args.samples} calls") + report["path_c_triton_device_build"] = _time_triton(args.samples, args.warmup) report["summary"] = _summarize(report) diff --git a/docs/operators/h3-sigma-schedule.md b/docs/operators/h3-sigma-schedule.md index 9a6404f66..2c064254a 100644 --- a/docs/operators/h3-sigma-schedule.md +++ b/docs/operators/h3-sigma-schedule.md @@ -22,7 +22,11 @@ schedule.num_model_evals # M - 1 (deduplicated length drives the count) ``` `get_op` returns the PyTorch reference on every platform, CUDA included: the -grid is built on CPU and both outputs are moved to `device`. +grid is built on CPU and both outputs are moved to `device`. When `device` is +omitted, the reference returns CPU tensors while the opt-in Triton backend +builds on `cuda` by default — callers should always pass `device=` explicitly. +Both per-backend defaults are intentional and pinned by the existing tests; +the default-backend switch did not change them. ## Opt-in: device-side Triton backend @@ -96,6 +100,24 @@ or on a HIP/ROCm build, so an opt-in that cannot be honored fails explicitly silently swapping arithmetic (see [Runtime Dispatch](../architecture/runtime-dispatch.md)). +The two selection layers differ in fallback semantics and device scope: + +- Candidate-list fallback belongs to the legacy `kernel_registry.get_op` path + only: it moves to the next candidate when a backend cannot be constructed. + For `sigma_schedule_h3` every platform list holds the reference alone, so + there is nothing to fall back from or to. +- The semantic path never falls back. + `session.resolve(requested_backend="rlkernel.h3_sigma_schedule.triton")` + requires that exact backend to be registered and to match the request + (declared `supported_devices={"cuda"}`, `fallback_policy=ERROR`). When + Triton is not importable, resolution still succeeds but + `session.instantiate` raises `OperatorInstantiationError` + (`"triton is required"`); the reference is not substituted. +- The semantic reference descriptor declares + `supported_devices={"cpu", "cuda", "rocm"}` and `fallback_policy=ERROR` — + it does not cover MUSA or NPU, whereas the legacy `get_op` lists serve the + reference on those platforms too. + ## Accuracy The acceptance oracle is the pinned fixture From bcddcd8eab2dda9a454040734464139d63999276 Mon Sep 17 00:00:00 2001 From: fxyan <1350821504@qq.com> Date: Mon, 26 Oct 2026 02:22:35 +0800 Subject: [PATCH 06/11] docs(h3): add SM90 (H800) evidence for sigma schedule MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Adds the 2026-10-26 SM90 hardware evidence under reports/experiments/ h3-sigma-schedule/sm90-h800/: environment snapshot, the three pytest runs (reference+dispatch+CUDA 67 passed / 2 skipped; CUDA+dispatch subset 47 / 1; the two registry tests introduced by the default-backend change, 2 / 0 — their first hardware execution) with logs and JUnit XML, the benchmark log and raw JSON (bitwise gate C==A 8/8; B/C 0.60–0.68x at N=25/50/1000, 4.12–4.34x at N=1e6), and the full collection report. Tested commit: 40070d37495d193e2d8764837321d3a68dad9483 (40070d3). GPU: NVIDIA H800 PCIe, compute capability 9.0, driver 580.82.07, torch 2.5.1+cu124, Triton 3.1.0. Versus the 4090D record: one more CUDA test passed because a66c999 split the registry-identity test; same single skip (multi-GPU guard) for the same reason; Triton version identical. Redaction before committing: GPU UUID, container hostname and /root/... paths removed (JUnit hostname attribute dropped like the 4090D XML; host paths additionally tokenized as REPO_ROOT/EVIDENCE_DIR/MINICONDA_PYTHON); a final newline appended to the three XMLs and the JSON for the end-of-file hook. No credentials, IPs or account names are included. source-provenance.json: cuda_validation keeps the historical 4090D hashes; the new sm90_validation block records the tested-commit source hashes (7/7), the committed (redacted) artifact hashes (10/10 re-verified), the pre-redaction original hashes, the redaction rules, and the binding note that commit ace9c20 changed only the benchmark's backend-unavailable/HIP branches and docs after the run — the executed NVIDIA+Triton path is unchanged, and those branches have not run on GPU or HIP hardware. Latency numbers are supplementary (reduced H800 FP64 throughput, different host CPU) and not comparable with the 4090D table. Signed-off-by: fxyan <1350821504@qq.com> --- .../experiments/h3-sigma-schedule/README.md | 90 +- .../sm90-h800/benchmark-h800.json | 2708 +++++++++++++++++ .../h3-sigma-schedule/sm90-h800/env.txt | 6 + .../sm90-h800/pr566_sm90_evidence.md | 449 +++ .../h3-sigma-schedule/sm90-h800/run_a.xml | 1 + .../h3-sigma-schedule/sm90-h800/run_b.xml | 1 + .../h3-sigma-schedule/sm90-h800/run_c.xml | 1 + .../h3-sigma-schedule/source-provenance.json | 86 +- 8 files changed, 3329 insertions(+), 13 deletions(-) create mode 100644 reports/experiments/h3-sigma-schedule/sm90-h800/benchmark-h800.json create mode 100644 reports/experiments/h3-sigma-schedule/sm90-h800/env.txt create mode 100644 reports/experiments/h3-sigma-schedule/sm90-h800/pr566_sm90_evidence.md create mode 100644 reports/experiments/h3-sigma-schedule/sm90-h800/run_a.xml create mode 100644 reports/experiments/h3-sigma-schedule/sm90-h800/run_b.xml create mode 100644 reports/experiments/h3-sigma-schedule/sm90-h800/run_c.xml diff --git a/reports/experiments/h3-sigma-schedule/README.md b/reports/experiments/h3-sigma-schedule/README.md index 2ed71b14b..0a553150a 100644 --- a/reports/experiments/h3-sigma-schedule/README.md +++ b/reports/experiments/h3-sigma-schedule/README.md @@ -23,14 +23,12 @@ backend is opt-in through the semantic `requested_backend` or direct construction). Those edits touch candidate order, test selection logic and prose only: every kernel bit assertion is unchanged, and the CUDA suite gained an assertion that the default reference path matches the Triton path bit for -bit. No GPU run backs the revised tests yet. `source-provenance.json` still -carries the `374488b` hashes; its `source_files_sha256` entries for -`tests/runtime/test_h3_sigma_schedule_dispatch.py`, -`tests/backends/cuda/test_h3_sigma_schedule_cuda.py` and -`benchmarks/models/benchmark_h3_sigma_schedule.py`, and the benchmark -`binding` note ("Benchmark and reference byte-identical"), are therefore stale. -They are deliberately left as recorded and will be refreshed together with the -pending SM90 (Hopper) evidence run rather than piecemeal. +bit. No GPU run backed the revised tests at the time; the SM90 (H800) run +below now does. `source-provenance.json` has been refreshed accordingly: its +`cuda_validation` block keeps the historical `374488b` hashes of the +4090D-tested snapshot, and the new `sm90_validation` block records the hashes +at the SM90-tested commit (`40070d3`), the post-run changes, and the SHA256 +of every included SM90 artifact. The performance run (`run-20261016-235221-8055`) predates the final HIP rejection guard. Its reference and benchmark are byte-identical to the @@ -68,6 +66,82 @@ is a deduplication stress case: the shift=12 gain exceeds the benchmark's dispersion check; shift=3 does not and is not claimed as a reliable gain. These are single-operator results, with no whole-model speedup claim. +## SM90 (H800) evidence + +Collected on 2026-10-26 on a rented SM90 host (one **NVIDIA H800 PCIe**, per +`nvidia-smi`), with the same three pytest invocations and the same benchmark +command as above, at tested commit +`40070d37495d193e2d8764837321d3a68dad9483` (`40070d3`). Raw logs, JUnit XML, +the benchmark JSON and the full collection report are in +[sm90-h800/](sm90-h800/). GPU UUID, container hostname and `/root/…` paths +were redacted: the JUnit `hostname` attribute was removed (as for the 4090D +XML, which keeps the historical `/root/…` skip-message paths — the SM90 files +additionally replace them with `REPO_ROOT` / `EVIDENCE_DIR` / +`MINICONDA_PYTHON` tokens). `benchmark-h800.json` and `env.txt` required no +redaction; a final newline was appended to the three XMLs and the JSON to +satisfy the repository's end-of-file hook. `source-provenance.json` records +both the committed (redacted) and the original server-side SHA256 of every +file. + +Hardware/software: NVIDIA H800 PCIe (compute capability 9.0), driver +580.82.07, x86_64 Linux 5.15.0-78-generic, 2× Intel Xeon Platinum 8458P, +Python 3.12.3, torch 2.5.1+cu124, CUDA runtime 12.4, Triton 3.1.0 — the same +Triton as the 4090D run; torch differs only in the CUDA build (cu124 vs +cu121), the driver is 580.82.07 vs 580.76.05. + +| Run | Scope (`-v -rs`) | Collected | Passed | Skipped | Failed | +| --- | --- | ---: | ---: | ---: | ---: | +| A | reference + dispatch + CUDA suites | 69 | 67 | 2 | 0 | +| B | CUDA + dispatch subset | 48 | 47 | 1 | 0 | +| C | the two registry tests introduced after the 4090D runs | 2 | 2 | 0 | 0 | + +Skips, with the reasons printed by `-rs`: the optional Diffusers parity test +(`diffusers` not installed on the host; the pinned fixture test passed) and +the multi-GPU device-guard test (the host has one GPU) — the same test, for +the same reason, as on the 4090D. Run C was the first hardware execution of +`test_public_entry_defaults_to_reference_and_matches_triton_bits` and +`test_public_entry_opt_in_selects_triton_and_reports_provenance`; both +passed — on SM90 the legacy `get_op(device="cuda")` reference path and the +directly constructed Triton backend produce identical sigma/timestep bits. + +Versus the 4090D record (34 passed / 1 skipped): one more CUDA test passed +because the default-backend change split the registry-identity test into +two; no test that passed on the 4090D failed here; the single skip is the +same test for the same reason. + +Benchmark (100 samples, 5 warmups, [raw JSON](sm90-h800/benchmark-h800.json)): +the pre-timing bitwise gate (path C vs path A) passed for all 8 cases; Triton +compilation 1143.6 ms, reported once and excluded from timed samples. + +| Workload | B median (ms) | C median (ms) | B/C | +| --- | ---: | ---: | ---: | +| N=25, shift=3 | 0.0551 | 0.0926 | 0.60x | +| N=25, shift=12 | 0.0554 | 0.0913 | 0.61x | +| N=50, shift=3 | 0.0556 | 0.0918 | 0.61x | +| N=50, shift=12 | 0.0554 | 0.0887 | 0.62x | +| N=1000, shift=3 | 0.0598 | 0.0907 | 0.66x | +| N=1000, shift=12 | 0.0609 | 0.0892 | 0.68x | +| N=1000000, shift=3 | 5.5054 | 1.3377 | 4.12x | +| N=1000000, shift=12 | 5.8109 | 1.3374 | 4.34x | + +All eight ratios pass the benchmark's dispersion check. These latency numbers +are supplementary only and not comparable with the 4090D table: the H800 has +reduced FP64 throughput (the Triton kernel evaluates its intermediates in +fp64), and paths A/B run on a different host CPU (at N=1e6 the CPU path is +about 5.5 ms here versus about 48 ms on the 4090D host). On this host the +device build is 1.5–1.65x slower than CPU build plus copy at N=25/50/1000 +and 4.1–4.3x faster at the N=1e6 dedup stress case — directionally the same +default-relevant conclusion as the 4090D, at a smaller margin. + +Post-run changes: after the tested commit `40070d3`, commit `ace9c20` +("fix(h3): address CodeRabbit review on sigma schedule benchmark and docs") +modified only the benchmark's GPU-without-Triton and HIP/ROCm branches, its +docstring, and documentation. The H800 run executed the NVIDIA + Triton +branch, whose logic, timing order, sample counts and JSON keys that commit +leaves unchanged; the new branches were verified by CPU-only execution, +`_summarize` simulations and static review, and have not run on GPU or HIP +hardware. + ## Reproduction From a checkout with the declared project dependencies installed, run: diff --git a/reports/experiments/h3-sigma-schedule/sm90-h800/benchmark-h800.json b/reports/experiments/h3-sigma-schedule/sm90-h800/benchmark-h800.json new file mode 100644 index 000000000..ed2a0dfe1 --- /dev/null +++ b/reports/experiments/h3-sigma-schedule/sm90-h800/benchmark-h800.json @@ -0,0 +1,2708 @@ +{ + "operator": "h3_sigma_schedule", + "timing": "per-call end-to-end host wall latency; GPU paths sync before each sample and after each call; Triton first-call compile excluded", + "entry": "direct op construction; registry bypassed for timing (identity pinned by tests)", + "torch": "2.5.1+cu124", + "torch_cuda_runtime": "12.4", + "hip": null, + "triton": "3.1.0", + "device": "NVIDIA H800 PCIe", + "n_gpu_devices": 1, + "samples": 100, + "warmup": 5, + "sizes": [ + 25, + 50, + 1000, + 1000000 + ], + "shifts": [ + 3.0, + 12.0 + ], + "stress_case_note": "N=1000000 is a dedup stress case, not a typical H3 workload", + "path_a_reference_cpu": { + "N=25,shift=3": { + "samples_raw_ms": [ + 0.037838, + 0.034582, + 0.03184, + 0.032533, + 0.032764, + 0.031274, + 0.033908, + 0.032909, + 0.030935, + 0.032555, + 0.032168, + 0.030652, + 0.034053, + 0.032242, + 0.030402, + 0.032399, + 0.032403, + 0.030871, + 0.032302, + 0.032403, + 0.031132, + 0.033367, + 0.032492, + 0.030559, + 0.032444, + 0.030715, + 0.032466, + 0.032306, + 0.030492, + 0.032187, + 0.032309, + 0.030369, + 0.051644, + 0.043776, + 0.033252, + 0.031613, + 0.030953, + 0.030562, + 0.030536, + 0.030629, + 0.030488, + 0.030402, + 0.030518, + 0.030365, + 0.030447, + 0.030328, + 0.030443, + 0.030585, + 0.030633, + 0.030279, + 0.030536, + 0.030395, + 0.031631, + 0.030708, + 0.030566, + 0.030473, + 0.030592, + 0.030588, + 0.030573, + 0.030462, + 0.030588, + 0.030532, + 0.030551, + 0.030417, + 0.030544, + 0.03051, + 0.030596, + 0.050612, + 0.032309, + 0.031516, + 0.031527, + 0.03146, + 0.030793, + 0.030778, + 0.030726, + 0.030477, + 0.030618, + 0.031006, + 0.030767, + 0.030912, + 0.030547, + 0.030503, + 0.030492, + 0.030521, + 0.030447, + 0.030361, + 0.030406, + 0.030521, + 0.030503, + 0.03048, + 0.030458, + 0.030443, + 0.030488, + 0.030793, + 0.030693, + 0.030763, + 0.030492, + 0.030499, + 0.030406, + 0.030603 + ], + "mean_ms": 0.03173504024744034, + "median_ms": 0.03063119947910309, + "stdev_ms": 0.0032713277358520274, + "min_ms": 0.030279159545898438, + "max_ms": 0.051643699407577515 + }, + "N=25,shift=12": { + "samples_raw_ms": [ + 0.032358, + 0.032112, + 0.0313, + 0.030957, + 0.030804, + 0.030871, + 0.030339, + 0.030708, + 0.030313, + 0.030376, + 0.033069, + 0.030559, + 0.030465, + 0.030059, + 0.03041, + 0.030063, + 0.030331, + 0.031155, + 0.030305, + 0.030857, + 0.030488, + 0.030644, + 0.03022, + 0.029992, + 0.030119, + 0.03029, + 0.030108, + 0.030138, + 0.030119, + 0.030078, + 0.030447, + 0.030134, + 0.030208, + 0.030272, + 0.051066, + 0.030968, + 0.030816, + 0.030667, + 0.030439, + 0.030313, + 0.030249, + 0.030458, + 0.030771, + 0.030443, + 2.686162, + 0.272006, + 0.03279, + 0.031814, + 0.030681, + 0.03079, + 0.0307, + 0.030696, + 0.030503, + 0.032492, + 0.031654, + 0.031717, + 0.030711, + 0.031848, + 0.031769, + 0.030823, + 0.031393, + 0.031695, + 0.031039, + 0.031628, + 0.030253, + 0.030711, + 0.030827, + 0.031084, + 0.03095, + 0.056971, + 0.031833, + 0.031535, + 0.031088, + 0.030939, + 0.030205, + 0.030499, + 0.03019, + 0.030834, + 0.039294, + 0.037517, + 0.030413, + 0.0301, + 0.030525, + 0.030048, + 0.02997, + 0.029892, + 0.030082, + 0.029996, + 0.030022, + 0.030082, + 0.030007, + 0.03003, + 0.030138, + 0.030965, + 0.030022, + 0.030495, + 0.030361, + 0.030238, + 0.031918, + 0.030376 + ], + "mean_ms": 0.0603017583489418, + "median_ms": 0.030541792511940002, + "stdev_ms": 0.26635107595218305, + "min_ms": 0.0298917293548584, + "max_ms": 2.686161547899246 + }, + "N=50,shift=3": { + "samples_raw_ms": [ + 0.032939, + 0.031751, + 0.031073, + 0.03152, + 0.030801, + 0.033367, + 0.031594, + 0.030868, + 0.030711, + 0.030648, + 0.030681, + 0.031445, + 0.030696, + 0.030756, + 0.030749, + 0.030696, + 0.031199, + 0.030939, + 0.030749, + 0.03089, + 0.030708, + 0.030417, + 0.030581, + 0.030637, + 0.030506, + 0.030722, + 0.031095, + 0.030611, + 0.046067, + 0.031214, + 0.03127, + 0.031061, + 0.033855, + 0.031311, + 0.031166, + 0.031158, + 0.030778, + 0.030972, + 0.030614, + 0.030428, + 0.030626, + 0.030905, + 0.030886, + 0.030614, + 0.030503, + 0.030525, + 0.031069, + 0.030674, + 0.030603, + 0.030663, + 0.030767, + 0.030607, + 0.030488, + 0.030614, + 0.030689, + 0.030622, + 0.03105, + 0.044879, + 0.031509, + 0.031158, + 0.031225, + 0.031445, + 0.031359, + 0.03111, + 0.031073, + 0.031218, + 0.030831, + 0.030577, + 0.030722, + 0.030473, + 0.030886, + 0.030678, + 0.0306, + 0.0306, + 0.030633, + 0.030909, + 0.0307, + 0.030831, + 0.030424, + 0.030596, + 0.030555, + 0.03041, + 0.030622, + 0.030462, + 0.030421, + 0.040594, + 0.045575, + 0.032853, + 0.032116, + 0.031419, + 0.031181, + 0.031255, + 0.03095, + 0.031099, + 0.031367, + 0.030801, + 0.030953, + 0.031319, + 0.030987, + 0.030782 + ], + "mean_ms": 0.03151305019855499, + "median_ms": 0.0308305025100708, + "stdev_ms": 0.0027166237925082674, + "min_ms": 0.030409544706344604, + "max_ms": 0.04606693983078003 + }, + "N=50,shift=12": { + "samples_raw_ms": [ + 0.032045, + 0.031587, + 0.030905, + 0.030711, + 0.030987, + 0.033181, + 0.031073, + 0.031095, + 0.031244, + 0.031192, + 0.04537, + 0.031762, + 0.031557, + 0.031292, + 0.031274, + 0.031576, + 0.031266, + 0.031613, + 0.03149, + 0.030849, + 0.031348, + 0.030819, + 0.031006, + 0.031441, + 0.030696, + 0.030879, + 0.030685, + 0.0306, + 0.030685, + 0.031255, + 0.030916, + 0.030678, + 0.030585, + 0.030603, + 0.030737, + 0.03086, + 0.03076, + 0.030719, + 0.030953, + 0.044774, + 0.031244, + 0.031393, + 0.031304, + 0.031162, + 0.031125, + 0.031289, + 0.031102, + 0.031207, + 0.030741, + 0.03117, + 0.030886, + 0.030991, + 0.030708, + 0.030749, + 0.030886, + 0.030745, + 0.030871, + 0.030678, + 0.030871, + 0.031061, + 0.03067, + 0.030719, + 0.030544, + 0.030659, + 0.030804, + 0.030771, + 0.030916, + 0.031032, + 0.044439, + 0.031795, + 0.031687, + 0.031464, + 0.031274, + 0.031218, + 0.03105, + 0.031114, + 0.03092, + 0.031173, + 0.031177, + 0.030864, + 0.031229, + 0.030778, + 0.030924, + 0.031147, + 0.030782, + 0.031121, + 0.030987, + 0.030968, + 0.031177, + 0.030834, + 0.030793, + 0.030924, + 0.030842, + 0.0313, + 0.030909, + 0.030927, + 0.031084, + 0.044756, + 0.038069, + 0.032388 + ], + "mean_ms": 0.03169480711221695, + "median_ms": 0.031018629670143127, + "stdev_ms": 0.002812785952845524, + "min_ms": 0.030543655157089233, + "max_ms": 0.04537031054496765 + }, + "N=1000,shift=3": { + "samples_raw_ms": [ + 0.037182, + 0.036236, + 0.035934, + 0.035975, + 0.035644, + 0.036888, + 0.035163, + 0.034999, + 0.035509, + 0.035387, + 0.034757, + 0.035331, + 0.034951, + 0.035122, + 0.034671, + 0.034954, + 0.034891, + 0.035074, + 0.0346, + 0.034753, + 0.034429, + 0.034787, + 0.034735, + 0.034835, + 0.035245, + 0.034831, + 0.035275, + 0.034831, + 0.03517, + 0.035141, + 0.034921, + 0.034615, + 0.035211, + 0.034578, + 0.034623, + 0.035118, + 0.03501, + 0.034831, + 0.03475, + 0.035003, + 0.034723, + 0.034299, + 0.034656, + 0.034865, + 0.03453, + 0.034705, + 0.034578, + 0.035554, + 0.035349, + 0.034653, + 0.034552, + 0.034839, + 0.034519, + 0.035152, + 0.034392, + 0.034723, + 0.035036, + 0.0346, + 0.034515, + 0.034727, + 0.034846, + 0.034697, + 0.034567, + 0.03456, + 0.03485, + 0.034656, + 0.034545, + 0.034556, + 0.03453, + 0.034578, + 0.034317, + 0.034351, + 0.034407, + 0.034928, + 0.034917, + 0.034351, + 0.034451, + 0.038326, + 0.035465, + 0.034664, + 0.034798, + 0.034444, + 0.034921, + 0.034608, + 0.034381, + 0.034504, + 0.034433, + 0.034295, + 0.034615, + 0.034411, + 0.034455, + 0.034556, + 0.034481, + 0.034671, + 0.034783, + 0.034489, + 0.034597, + 0.034463, + 0.034735, + 0.034459 + ], + "mean_ms": 0.034885890781879425, + "median_ms": 0.034730881452560425, + "stdev_ms": 0.0005986840381473487, + "min_ms": 0.03429502248764038, + "max_ms": 0.03832578659057617 + }, + "N=1000,shift=12": { + "samples_raw_ms": [ + 0.035159, + 0.035308, + 0.035092, + 0.034455, + 0.03485, + 0.03631, + 0.034999, + 0.034284, + 0.034627, + 0.03488, + 0.034537, + 0.03431, + 0.034388, + 0.034936, + 0.034869, + 0.034586, + 0.034425, + 0.034705, + 0.03491, + 0.035301, + 0.034288, + 0.034925, + 0.035167, + 0.034943, + 0.034627, + 0.034854, + 0.03479, + 0.035092, + 0.035081, + 0.034705, + 0.034664, + 0.035044, + 0.035003, + 0.034831, + 0.03504, + 0.034574, + 0.03447, + 0.034656, + 0.034887, + 0.035007, + 0.03488, + 0.034861, + 0.034448, + 0.03479, + 0.034768, + 0.035189, + 0.034481, + 0.034608, + 0.03428, + 0.034347, + 0.034146, + 0.034112, + 0.03418, + 0.034288, + 0.034269, + 0.034083, + 0.034176, + 0.034858, + 0.034384, + 0.034191, + 0.034418, + 0.034533, + 0.034079, + 0.034295, + 0.034388, + 0.034582, + 0.034731, + 0.0345, + 0.034254, + 0.034258, + 0.034567, + 0.034705, + 0.034478, + 0.034247, + 0.034526, + 0.034232, + 0.03428, + 0.044417, + 0.035044, + 0.03485, + 0.034276, + 0.03428, + 0.034515, + 0.034157, + 0.034202, + 0.034384, + 0.034302, + 0.034414, + 0.034358, + 0.034317, + 0.034336, + 0.034351, + 0.034057, + 0.034902, + 0.034843, + 0.034727, + 0.034668, + 0.034556, + 0.034507, + 0.034362 + ], + "mean_ms": 0.03470815718173981, + "median_ms": 0.03456138074398041, + "stdev_ms": 0.0010438846934519405, + "min_ms": 0.03405660390853882, + "max_ms": 0.0444166362285614 + }, + "N=1000000,shift=3": { + "samples_raw_ms": [ + 5.361874, + 5.65834, + 5.872671, + 5.757738, + 5.676169, + 5.505249, + 5.514953, + 5.293507, + 5.593728, + 5.615912, + 5.507495, + 5.496029, + 8.408848, + 5.396958, + 5.217146, + 5.568784, + 5.486324, + 5.46227, + 5.47206, + 5.362313, + 5.541779, + 5.188312, + 5.583975, + 5.5346, + 5.447749, + 5.46075, + 5.460072, + 5.366024, + 5.345449, + 5.22073, + 5.449198, + 5.493548, + 5.421292, + 5.410228, + 5.455185, + 5.614895, + 6.810188, + 5.617775, + 5.48723, + 5.477309, + 5.434804, + 5.453546, + 5.378246, + 5.171869, + 5.559742, + 5.550258, + 5.618561, + 5.566526, + 6.193709, + 5.438566, + 5.382642, + 5.216405, + 5.467933, + 5.489282, + 5.454388, + 5.474135, + 5.568266, + 5.366448, + 5.166207, + 5.503662, + 5.432479, + 5.381357, + 5.280394, + 5.429726, + 5.335905, + 5.09141, + 8.170862, + 5.520195, + 5.426619, + 5.506072, + 5.41712, + 5.375735, + 5.329158, + 5.776018, + 5.886864, + 5.708899, + 5.563445, + 5.504258, + 5.477723, + 5.475532, + 5.242143, + 5.685389, + 5.651239, + 5.584341, + 8.296784, + 5.548671, + 5.399782, + 5.201422, + 5.424861, + 5.478892, + 5.338028, + 5.386479, + 5.370464, + 5.383082, + 5.351197, + 5.602982, + 5.339969, + 5.388834, + 5.380187, + 5.431764 + ], + "mean_ms": 5.566461347043514, + "median_ms": 5.465101450681686, + "stdev_ms": 0.52512210414059, + "min_ms": 5.091410130262375, + "max_ms": 8.408848196268082 + }, + "N=1000000,shift=12": { + "samples_raw_ms": [ + 5.727381, + 5.48099, + 5.72215, + 5.575776, + 5.591355, + 5.791593, + 5.785979, + 5.561594, + 5.49572, + 5.610313, + 5.440947, + 5.400874, + 5.664624, + 5.553484, + 5.759057, + 8.177653, + 5.502317, + 5.602334, + 5.268849, + 5.678028, + 5.730242, + 5.587954, + 5.67773, + 5.560193, + 5.574163, + 5.520288, + 5.769704, + 5.698543, + 5.571235, + 5.769622, + 5.472884, + 5.544968, + 6.228525, + 5.955953, + 5.897775, + 5.695306, + 5.763184, + 5.66797, + 5.725149, + 5.389154, + 5.637657, + 5.674887, + 5.596112, + 5.684637, + 5.598892, + 5.511619, + 5.389471, + 5.824927, + 5.621884, + 8.258492, + 5.752951, + 5.454056, + 5.399726, + 5.26803, + 5.511761, + 5.416594, + 5.517431, + 5.437512, + 5.404018, + 5.334832, + 5.174622, + 5.43889, + 5.469266, + 5.632438, + 5.481295, + 5.500466, + 5.627308, + 5.258638, + 7.602617, + 5.54863, + 5.552225, + 5.589299, + 5.539946, + 5.848981, + 5.547348, + 5.573068, + 5.536903, + 5.511921, + 5.591687, + 5.578797, + 5.445872, + 5.334768, + 5.764473, + 5.819, + 5.644642, + 6.066658, + 5.552158, + 5.409602, + 5.289368, + 5.667087, + 5.695198, + 5.800877, + 6.829537, + 8.392874, + 5.510505, + 5.460434, + 5.553387, + 5.550217, + 5.552404, + 5.725563 + ], + "mean_ms": 5.701560154557228, + "median_ms": 5.577286705374718, + "stdev_ms": 0.5389929450092558, + "min_ms": 5.174621939659119, + "max_ms": 8.392874151468277 + } + }, + "correctness_gate_wall_ms": 1143.569815903902, + "path_b_reference_cpu_build_then_copy_to_gpu": { + "N=25,shift=3": { + "samples_raw_ms": [ + 0.073742, + 0.069104, + 0.061151, + 0.057418, + 0.056166, + 0.056311, + 0.056222, + 0.060938, + 0.057295, + 0.055633, + 0.056069, + 0.055559, + 0.056248, + 0.055764, + 0.054941, + 0.055809, + 0.055347, + 0.055254, + 0.05506, + 0.055276, + 0.055008, + 0.055056, + 0.05535, + 0.056021, + 0.054829, + 0.05506, + 0.05519, + 0.055134, + 0.05471, + 0.05535, + 0.054941, + 0.055, + 0.05449, + 0.054881, + 0.054438, + 0.055112, + 0.054084, + 0.055626, + 0.054527, + 0.054847, + 0.054393, + 0.055503, + 0.055034, + 0.054803, + 0.054527, + 0.054613, + 0.054453, + 0.055667, + 0.055142, + 0.068955, + 0.056643, + 0.0556, + 0.054516, + 0.055306, + 0.054799, + 0.05538, + 0.054646, + 0.055209, + 0.054561, + 0.055268, + 0.05459, + 0.054676, + 0.054218, + 0.055533, + 0.054553, + 0.054531, + 0.054214, + 0.054609, + 0.054285, + 0.054814, + 0.054799, + 0.054635, + 0.053909, + 0.054743, + 0.054311, + 0.055391, + 0.05427, + 0.094466, + 0.056751, + 0.055436, + 0.055157, + 0.054762, + 0.054799, + 0.055578, + 0.054721, + 0.054516, + 0.054751, + 0.055131, + 0.061873, + 0.057012, + 0.056118, + 0.055466, + 0.05535, + 0.055201, + 0.055369, + 0.055369, + 0.054963, + 0.055123, + 0.054903, + 0.061948 + ], + "mean_ms": 0.056268200278282166, + "median_ms": 0.0551324337720871, + "stdev_ms": 0.0048781186614028456, + "min_ms": 0.05390867590904236, + "max_ms": 0.09446591138839722 + }, + "N=25,shift=12": { + "samples_raw_ms": [ + 0.059947, + 0.06073, + 0.056688, + 0.056695, + 0.063073, + 0.058059, + 0.056956, + 0.055294, + 0.055853, + 0.055477, + 0.055544, + 0.055969, + 0.055347, + 0.055674, + 0.055719, + 0.062454, + 0.057034, + 0.055417, + 0.056483, + 0.055585, + 0.055287, + 0.055224, + 0.055805, + 0.05519, + 0.055, + 0.055827, + 0.056006, + 0.054784, + 0.055496, + 0.063196, + 0.057079, + 0.055742, + 0.055742, + 0.055209, + 0.055499, + 0.05544, + 0.055406, + 0.055123, + 0.055887, + 0.055283, + 0.055384, + 0.05535, + 0.055034, + 0.063378, + 0.056215, + 0.055261, + 0.055838, + 0.056263, + 0.055343, + 0.05554, + 0.054672, + 0.055708, + 0.055131, + 0.055075, + 0.05471, + 0.059985, + 0.05509, + 0.054747, + 0.054192, + 0.054475, + 0.054136, + 0.054557, + 0.054646, + 0.054605, + 0.054639, + 0.055492, + 0.05668, + 0.05582, + 0.055078, + 0.056058, + 0.05509, + 0.054978, + 0.054456, + 0.055004, + 0.054725, + 0.054777, + 0.054792, + 0.074245, + 0.057973, + 0.055943, + 0.05598, + 0.055183, + 0.055317, + 0.055578, + 0.055224, + 0.055473, + 0.055436, + 0.054963, + 0.055566, + 0.055172, + 0.054926, + 0.055805, + 0.055194, + 0.054948, + 0.054836, + 0.0543, + 0.054784, + 0.054087, + 0.05503, + 0.054624 + ], + "mean_ms": 0.05605731159448624, + "median_ms": 0.05539506673812866, + "stdev_ms": 0.0025914793282457386, + "min_ms": 0.05408748984336853, + "max_ms": 0.07424503564834595 + }, + "N=50,shift=3": { + "samples_raw_ms": [ + 0.058305, + 0.057299, + 0.055715, + 0.055555, + 0.055641, + 0.076547, + 0.05858, + 0.055876, + 0.056185, + 0.056192, + 0.055715, + 0.055164, + 0.055335, + 0.0556, + 0.055902, + 0.056002, + 0.0556, + 0.055719, + 0.055563, + 0.055172, + 0.055403, + 0.055272, + 0.06232, + 0.057839, + 0.055723, + 0.055857, + 0.05592, + 0.055835, + 0.055321, + 0.055525, + 0.056054, + 0.061613, + 0.055149, + 0.055362, + 0.055574, + 0.055518, + 0.055317, + 0.055317, + 0.055794, + 0.055309, + 0.054989, + 0.05576, + 0.055365, + 0.055797, + 0.055067, + 0.056487, + 0.056263, + 0.055853, + 0.055544, + 0.055965, + 0.054967, + 0.055447, + 0.05538, + 0.055738, + 0.055559, + 0.055365, + 0.054739, + 0.055317, + 0.055276, + 0.055425, + 0.05563, + 0.055417, + 0.05576, + 0.055797, + 0.054833, + 0.055283, + 0.055008, + 0.055436, + 0.055302, + 0.055585, + 0.05579, + 0.055723, + 0.055239, + 0.055496, + 0.055056, + 0.055499, + 0.055172, + 0.055782, + 0.090294, + 0.057507, + 0.055984, + 0.055961, + 0.055782, + 0.055786, + 0.055607, + 0.05557, + 0.056401, + 0.066824, + 0.057545, + 0.055801, + 0.055708, + 0.055697, + 0.05535, + 0.056457, + 0.055544, + 0.055563, + 0.055481, + 0.055432, + 0.055335, + 0.055484 + ], + "mean_ms": 0.05650915205478668, + "median_ms": 0.05559995770454407, + "stdev_ms": 0.004274637427571502, + "min_ms": 0.054739415645599365, + "max_ms": 0.09029358625411987 + }, + "N=50,shift=12": { + "samples_raw_ms": [ + 0.057478, + 0.056051, + 0.055581, + 0.055239, + 0.05525, + 0.057001, + 0.055306, + 0.054765, + 0.054955, + 0.054847, + 0.055358, + 0.055045, + 0.055153, + 0.054512, + 0.055235, + 0.054602, + 0.054646, + 0.054825, + 0.054713, + 0.055026, + 0.054862, + 0.054356, + 0.055533, + 0.05506, + 0.054799, + 0.054821, + 0.054941, + 0.054445, + 0.054777, + 0.055298, + 0.055399, + 0.055313, + 0.054836, + 0.060797, + 0.05528, + 0.055075, + 0.055429, + 0.055224, + 0.05566, + 0.055376, + 0.055302, + 0.054944, + 0.055041, + 0.055328, + 0.055421, + 0.062045, + 0.057876, + 0.05646, + 0.056285, + 0.056047, + 0.055891, + 0.055678, + 0.055894, + 0.056006, + 0.056539, + 0.055585, + 0.056107, + 0.055257, + 0.055589, + 0.055578, + 0.055451, + 0.055607, + 0.056148, + 0.055477, + 0.055932, + 0.055734, + 0.055481, + 0.055682, + 0.055574, + 0.055633, + 0.055708, + 0.055391, + 0.05636, + 0.055291, + 0.055481, + 0.055309, + 0.056416, + 0.103641, + 0.057999, + 0.056032, + 0.056434, + 0.055935, + 0.056356, + 0.05598, + 0.055339, + 0.055358, + 0.056207, + 0.055205, + 0.055999, + 0.055321, + 0.055205, + 0.055369, + 0.056036, + 0.05535, + 0.055529, + 0.055384, + 0.055268, + 0.054713, + 0.055276, + 0.055511 + ], + "mean_ms": 0.05611862987279892, + "median_ms": 0.05539506673812866, + "stdev_ms": 0.004914717693255326, + "min_ms": 0.05435571074485779, + "max_ms": 0.10364130139350891 + }, + "N=1000,shift=3": { + "samples_raw_ms": [ + 0.064794, + 0.080377, + 0.066295, + 0.062879, + 0.06295, + 0.063401, + 0.062127, + 0.061866, + 0.061288, + 0.06146, + 0.060804, + 0.061095, + 0.060733, + 0.060484, + 0.060435, + 0.061229, + 0.060517, + 0.060644, + 0.063952, + 0.059702, + 0.059679, + 0.059642, + 0.060178, + 0.059795, + 0.059452, + 0.05969, + 0.059724, + 0.05956, + 0.059601, + 0.059985, + 0.059564, + 0.059493, + 0.059687, + 0.059359, + 0.059512, + 0.059575, + 0.059418, + 0.059951, + 0.059456, + 0.059433, + 0.059485, + 0.059627, + 0.059355, + 0.059508, + 0.059295, + 0.05924, + 0.059403, + 0.059541, + 0.059348, + 0.059389, + 0.059649, + 0.059206, + 0.059187, + 0.059612, + 0.059661, + 0.059478, + 0.059213, + 0.059318, + 0.059474, + 0.059355, + 0.059176, + 0.05899, + 0.069093, + 0.06013, + 0.060134, + 0.059996, + 0.060096, + 0.059314, + 0.059567, + 0.059795, + 0.059936, + 0.060033, + 0.059817, + 0.060085, + 0.059653, + 0.060029, + 0.059839, + 0.060413, + 0.061087, + 0.060648, + 0.060417, + 0.060022, + 0.060219, + 0.059403, + 0.059515, + 0.059813, + 0.059932, + 0.059783, + 0.059925, + 0.059664, + 0.059716, + 0.05959, + 0.059541, + 0.060145, + 0.059668, + 0.059418, + 0.060171, + 0.060383, + 0.061072, + 0.060719 + ], + "mean_ms": 0.060450807213783264, + "median_ms": 0.0597536563873291, + "stdev_ms": 0.0025028076165445503, + "min_ms": 0.05898997187614441, + "max_ms": 0.08037686347961426 + }, + "N=1000,shift=12": { + "samples_raw_ms": [ + 0.065081, + 0.063077, + 0.061702, + 0.061788, + 0.06159, + 0.061926, + 0.061221, + 0.061147, + 0.061255, + 0.061322, + 0.059925, + 0.060678, + 0.061266, + 0.061076, + 0.060935, + 0.072021, + 0.061978, + 0.061266, + 0.061255, + 0.060707, + 0.060923, + 0.061005, + 0.060428, + 0.060752, + 0.060637, + 0.061203, + 0.060983, + 0.060577, + 0.061054, + 0.061009, + 0.060562, + 0.061199, + 0.061151, + 0.060976, + 0.06146, + 0.060599, + 0.060879, + 0.060599, + 0.060529, + 0.061356, + 0.061292, + 0.060912, + 0.060834, + 0.061065, + 0.061307, + 0.061035, + 0.060804, + 0.061158, + 0.061233, + 0.060759, + 0.060707, + 0.06067, + 0.060372, + 0.061009, + 0.060655, + 0.060745, + 0.060696, + 0.061024, + 0.061106, + 0.063062, + 0.060428, + 0.060897, + 0.060804, + 0.060666, + 0.061084, + 0.060894, + 0.060782, + 0.060838, + 0.060655, + 0.060946, + 0.060413, + 0.060875, + 0.060517, + 0.060648, + 0.060916, + 0.0665, + 0.061866, + 0.061233, + 0.06089, + 0.060897, + 0.060517, + 0.060886, + 0.060949, + 0.060674, + 0.06083, + 0.060715, + 0.060823, + 0.060983, + 0.061061, + 0.060834, + 0.061192, + 0.060342, + 0.060827, + 0.060961, + 0.0608, + 0.060976, + 0.060547, + 0.060607, + 0.060868, + 0.060383 + ], + "mean_ms": 0.06118062883615494, + "median_ms": 0.06091408431529999, + "stdev_ms": 0.001370676881341059, + "min_ms": 0.05992501974105835, + "max_ms": 0.07202103734016418 + }, + "N=1000000,shift=3": { + "samples_raw_ms": [ + 5.420301, + 5.573802, + 7.34169, + 6.922781, + 5.630899, + 5.271818, + 5.573958, + 5.579349, + 5.228311, + 5.624823, + 5.456261, + 5.547434, + 5.642328, + 5.290475, + 5.615506, + 5.367339, + 5.729217, + 5.421154, + 5.361885, + 5.70387, + 5.244575, + 5.617555, + 5.282879, + 5.545091, + 5.599212, + 5.375799, + 5.702991, + 5.39916, + 5.93612, + 5.573813, + 5.326714, + 5.865097, + 5.263448, + 5.71344, + 5.534325, + 5.382109, + 5.472254, + 5.239766, + 5.599029, + 5.361926, + 5.351588, + 5.505066, + 5.474094, + 5.751725, + 5.332585, + 5.716708, + 5.760171, + 5.388729, + 5.633354, + 5.343515, + 5.623125, + 5.505346, + 5.437244, + 5.408697, + 5.314242, + 5.53897, + 5.53057, + 5.336247, + 5.406167, + 5.358256, + 5.525365, + 5.406685, + 5.451351, + 8.264519, + 5.398445, + 5.717769, + 5.36618, + 5.523179, + 5.447816, + 5.392302, + 5.760994, + 6.553531, + 5.695596, + 5.69338, + 5.324505, + 5.451925, + 5.55541, + 5.444765, + 5.536642, + 5.418479, + 5.745929, + 5.225446, + 5.52218, + 5.494945, + 5.313542, + 5.850047, + 5.202089, + 5.426206, + 5.405411, + 5.535334, + 5.630612, + 5.307063, + 5.680021, + 5.187072, + 5.54166, + 5.575329, + 5.505398, + 5.791351, + 5.235676, + 7.981692 + ], + "mean_ms": 5.591407492756844, + "median_ms": 5.50537183880806, + "stdev_ms": 0.4726045374094776, + "min_ms": 5.187071859836578, + "max_ms": 8.264519274234772 + }, + "N=1000000,shift=12": { + "samples_raw_ms": [ + 5.8144, + 5.566522, + 5.728371, + 5.45799, + 5.871754, + 5.484313, + 5.945504, + 5.828053, + 5.445063, + 5.824745, + 5.454518, + 5.839102, + 6.108165, + 5.487736, + 5.986929, + 5.710568, + 6.1781, + 5.624902, + 5.586345, + 5.754508, + 5.526345, + 5.759485, + 5.699798, + 5.677782, + 5.771004, + 5.553622, + 6.049965, + 5.541608, + 5.961142, + 8.272201, + 5.578995, + 5.891763, + 5.506337, + 5.931854, + 5.696416, + 5.585905, + 5.625974, + 5.693983, + 5.777933, + 5.891308, + 6.078996, + 6.179873, + 6.293595, + 6.164681, + 5.785927, + 8.556936, + 6.764196, + 5.982038, + 6.021906, + 5.57629, + 5.719669, + 5.656384, + 5.579263, + 5.901765, + 5.698279, + 5.967963, + 5.824465, + 5.784102, + 5.811989, + 5.63607, + 5.837403, + 5.790506, + 5.440705, + 5.977049, + 5.725887, + 5.75798, + 5.93657, + 9.06219, + 5.736075, + 5.809858, + 6.007906, + 5.756225, + 5.859595, + 5.892143, + 5.480286, + 5.88192, + 5.751912, + 6.020565, + 5.975105, + 7.755134, + 5.964935, + 6.091718, + 5.667616, + 6.122351, + 5.820289, + 5.859446, + 5.65305, + 5.789768, + 5.70032, + 5.747419, + 5.798645, + 5.770408, + 5.854588, + 5.633131, + 6.108124, + 5.952206, + 5.944513, + 6.242868, + 6.384339, + 5.952809 + ], + "mean_ms": 5.921889469027519, + "median_ms": 5.810923874378204, + "stdev_ms": 0.5644538440448857, + "min_ms": 5.440704524517059, + "max_ms": 9.062189608812332 + } + }, + "path_c_triton_device_build": { + "N=25,shift=3": { + "samples_raw_ms": [ + 0.103951, + 0.111766, + 0.10309, + 0.098165, + 0.095345, + 0.103746, + 0.104621, + 0.103086, + 0.096597, + 0.098553, + 0.097677, + 0.096571, + 0.095259, + 0.097714, + 0.093605, + 0.100791, + 0.092756, + 0.103936, + 0.097848, + 0.089973, + 0.089932, + 0.095066, + 0.091497, + 0.093587, + 0.093643, + 0.100166, + 0.09859, + 0.095081, + 0.094574, + 0.093251, + 0.094414, + 0.093527, + 0.09146, + 0.09191, + 0.092145, + 0.090931, + 0.091583, + 0.092193, + 0.098307, + 0.092398, + 0.094507, + 0.098903, + 0.092566, + 0.092674, + 0.096779, + 0.093237, + 0.091247, + 0.091303, + 0.092566, + 0.092238, + 0.093929, + 0.092812, + 0.098154, + 0.090521, + 0.093959, + 0.091925, + 0.092518, + 0.090931, + 0.091672, + 0.091903, + 0.093237, + 0.094127, + 0.091758, + 0.090003, + 0.091545, + 0.092093, + 0.091113, + 0.091329, + 0.091556, + 0.093363, + 0.093646, + 0.091545, + 0.09257, + 0.092369, + 0.092331, + 0.092775, + 0.09105, + 0.092003, + 0.097167, + 0.093352, + 0.092313, + 0.091318, + 0.090856, + 0.092816, + 0.093054, + 0.092097, + 0.091065, + 0.090789, + 0.090852, + 0.091314, + 0.091083, + 0.090361, + 0.090566, + 0.089448, + 0.088938, + 0.089183, + 0.088762, + 0.090048, + 0.08899, + 0.08953 + ], + "mean_ms": 0.09385962039232254, + "median_ms": 0.09256601333618164, + "stdev_ms": 0.004073192358592415, + "min_ms": 0.08876249194145203, + "max_ms": 0.11176615953445435 + }, + "N=25,shift=12": { + "samples_raw_ms": [ + 0.105191, + 0.095554, + 0.097599, + 0.094187, + 0.094991, + 0.098132, + 0.098217, + 0.096351, + 0.094704, + 0.096109, + 0.094272, + 0.095949, + 0.096161, + 0.095271, + 0.093851, + 0.095166, + 0.09745, + 0.095543, + 0.095367, + 0.096168, + 0.093751, + 0.0934, + 0.091754, + 0.092141, + 0.090647, + 0.091922, + 0.090957, + 0.090167, + 0.089161, + 0.089977, + 0.090145, + 0.09042, + 0.089716, + 0.090588, + 0.091426, + 0.090789, + 0.092391, + 0.092197, + 0.091914, + 0.09229, + 0.09314, + 0.092704, + 0.0907, + 0.089705, + 0.088841, + 0.088818, + 0.088889, + 0.090413, + 0.089772, + 0.093449, + 0.092477, + 0.091095, + 0.091735, + 0.092182, + 0.090916, + 0.09172, + 0.089988, + 0.090539, + 0.090964, + 0.092115, + 0.090085, + 0.090484, + 0.089683, + 0.089429, + 0.089288, + 0.088889, + 0.089061, + 0.087962, + 0.088438, + 0.089824, + 0.089481, + 0.089534, + 0.088744, + 0.089496, + 0.088342, + 0.089064, + 0.089452, + 0.08934, + 0.08909, + 0.089455, + 0.09004, + 0.091311, + 0.089694, + 0.090603, + 0.090465, + 0.089716, + 0.088461, + 0.100542, + 0.120722, + 0.097133, + 0.093628, + 0.093393, + 0.091262, + 0.093538, + 0.105828, + 0.093546, + 0.092402, + 0.093728, + 0.091046, + 0.09124 + ], + "mean_ms": 0.09245585650205612, + "median_ms": 0.09125098586082458, + "stdev_ms": 0.004356443106898552, + "min_ms": 0.08796155452728271, + "max_ms": 0.12072175741195679 + }, + "N=50,shift=3": { + "samples_raw_ms": [ + 0.106525, + 0.097033, + 0.095535, + 0.094634, + 0.09504, + 0.095066, + 0.096597, + 0.095472, + 0.094999, + 0.093792, + 0.092577, + 0.092089, + 0.094403, + 0.094097, + 0.09438, + 0.092272, + 0.092562, + 0.09096, + 0.091966, + 0.091854, + 0.093408, + 0.092272, + 0.092555, + 0.091773, + 0.096358, + 0.09165, + 0.091601, + 0.102084, + 0.094928, + 0.092417, + 0.092573, + 0.091564, + 0.091884, + 0.093143, + 0.091717, + 0.091195, + 0.09143, + 0.091724, + 0.091985, + 0.091761, + 0.090174, + 0.090003, + 0.090759, + 0.090707, + 0.091325, + 0.090186, + 0.090051, + 0.089768, + 0.091299, + 0.090837, + 0.090729, + 0.090715, + 0.090856, + 0.091489, + 0.090946, + 0.092722, + 0.091843, + 0.091847, + 0.091568, + 0.09165, + 0.092532, + 0.092443, + 0.09146, + 0.090703, + 0.091873, + 0.092179, + 0.091124, + 0.091448, + 0.102494, + 0.093229, + 0.091042, + 0.092421, + 0.091482, + 0.091672, + 0.091102, + 0.090536, + 0.09153, + 0.091556, + 0.09121, + 0.090621, + 0.091206, + 0.09089, + 0.092573, + 0.089545, + 0.092268, + 0.091739, + 0.09048, + 0.092365, + 0.091545, + 0.091139, + 0.090674, + 0.091158, + 0.090178, + 0.090484, + 0.09089, + 0.091985, + 0.09276, + 0.093378, + 0.092745, + 0.093505 + ], + "mean_ms": 0.09243514388799667, + "median_ms": 0.09175017476081848, + "stdev_ms": 0.0025312190311252967, + "min_ms": 0.08954480290412903, + "max_ms": 0.10652467608451843 + }, + "N=50,shift=12": { + "samples_raw_ms": [ + 0.100508, + 0.108246, + 0.096697, + 0.093598, + 0.093032, + 0.094451, + 0.09485, + 0.10528, + 0.093866, + 0.097193, + 0.095006, + 0.094298, + 0.092424, + 0.089772, + 0.09061, + 0.089962, + 0.088762, + 0.088856, + 0.089537, + 0.088587, + 0.089373, + 0.089306, + 0.089779, + 0.089362, + 0.087984, + 0.089746, + 0.088811, + 0.088871, + 0.09007, + 0.089336, + 0.08826, + 0.088952, + 0.089221, + 0.090141, + 0.088673, + 0.088479, + 0.087593, + 0.088457, + 0.088867, + 0.087403, + 0.087153, + 0.087999, + 0.087012, + 0.096496, + 0.090342, + 0.088532, + 0.088356, + 0.088599, + 0.088774, + 0.088766, + 0.08817, + 0.087295, + 0.086319, + 0.08728, + 0.088375, + 0.08703, + 0.086956, + 0.086833, + 0.087183, + 0.087306, + 0.089619, + 0.089247, + 0.088334, + 0.08953, + 0.088152, + 0.087988, + 0.088017, + 0.087067, + 0.086699, + 0.088621, + 0.08823, + 0.089779, + 0.089463, + 0.089344, + 0.089139, + 0.08944, + 0.089128, + 0.089146, + 0.087559, + 0.087917, + 0.087172, + 0.087366, + 0.086132, + 0.085946, + 0.086136, + 0.0875, + 0.096213, + 0.089101, + 0.089038, + 0.087943, + 0.088662, + 0.088185, + 0.088193, + 0.088811, + 0.087984, + 0.086695, + 0.086986, + 0.086717, + 0.086591, + 0.086181 + ], + "mean_ms": 0.08958965539932251, + "median_ms": 0.08871778845787048, + "stdev_ms": 0.0036410057723217574, + "min_ms": 0.08594617247581482, + "max_ms": 0.10824576020240784 + }, + "N=1000,shift=3": { + "samples_raw_ms": [ + 0.100907, + 0.095759, + 0.093251, + 0.095285, + 0.095952, + 0.09536, + 0.096578, + 0.095107, + 0.095624, + 0.094082, + 0.093598, + 0.09447, + 0.09362, + 0.09476, + 0.094548, + 0.093557, + 0.093285, + 0.095617, + 0.095878, + 0.106394, + 0.094552, + 0.100516, + 0.093702, + 0.094552, + 0.09764, + 0.095058, + 0.09577, + 0.093881, + 0.093065, + 0.092398, + 0.092916, + 0.095256, + 0.09295, + 0.092428, + 0.092234, + 0.092961, + 0.093162, + 0.096999, + 0.092503, + 0.091515, + 0.091013, + 0.08966, + 0.091393, + 0.091352, + 0.0907, + 0.091203, + 0.091191, + 0.090383, + 0.090424, + 0.090096, + 0.090323, + 0.089664, + 0.090856, + 0.090227, + 0.089712, + 0.090178, + 0.089683, + 0.089135, + 0.089802, + 0.091266, + 0.101447, + 0.089027, + 0.088811, + 0.089195, + 0.087988, + 0.087444, + 0.088334, + 0.087444, + 0.086397, + 0.094682, + 0.090681, + 0.089746, + 0.089154, + 0.089146, + 0.09004, + 0.089094, + 0.088654, + 0.08814, + 0.087809, + 0.088129, + 0.087768, + 0.087153, + 0.088148, + 0.087082, + 0.087492, + 0.088166, + 0.087608, + 0.087377, + 0.089224, + 0.088066, + 0.088807, + 0.089917, + 0.089712, + 0.088863, + 0.088152, + 0.089582, + 0.088535, + 0.088681, + 0.08874, + 0.089489 + ], + "mean_ms": 0.09165875613689423, + "median_ms": 0.09069032967090607, + "stdev_ms": 0.0035835712922101495, + "min_ms": 0.08639693260192871, + "max_ms": 0.10639429092407227 + }, + "N=1000,shift=12": { + "samples_raw_ms": [ + 0.10056, + 0.09796, + 0.096533, + 0.093345, + 0.093963, + 0.094704, + 0.092544, + 0.091843, + 0.094969, + 0.095494, + 0.093628, + 0.093844, + 0.094566, + 0.093557, + 0.093792, + 0.091303, + 0.091195, + 0.090074, + 0.090335, + 0.089291, + 0.089295, + 0.089157, + 0.088632, + 0.088055, + 0.088532, + 0.088818, + 0.091296, + 0.088859, + 0.087544, + 0.088513, + 0.08782, + 0.087544, + 0.087991, + 0.087876, + 0.087071, + 0.087887, + 0.098597, + 0.093199, + 0.090774, + 0.090249, + 0.090163, + 0.089966, + 0.089552, + 0.089012, + 0.088107, + 0.08896, + 0.089824, + 0.089005, + 0.089217, + 0.089001, + 0.088017, + 0.089861, + 0.089265, + 0.090931, + 0.089824, + 0.088617, + 0.089422, + 0.089269, + 0.087224, + 0.088476, + 0.087798, + 0.088163, + 0.088226, + 0.088029, + 0.087727, + 0.088986, + 0.088859, + 0.089172, + 0.089053, + 0.090096, + 0.089325, + 0.088807, + 0.088237, + 0.088394, + 0.087421, + 0.08617, + 0.086464, + 0.088885, + 0.087317, + 0.10702, + 0.088554, + 0.088684, + 0.089664, + 0.088759, + 0.08852, + 0.089802, + 0.09013, + 0.089575, + 0.090152, + 0.089671, + 0.089891, + 0.088923, + 0.0888, + 0.08991, + 0.091184, + 0.08944, + 0.088844, + 0.089306, + 0.089198, + 0.088882 + ], + "mean_ms": 0.09016964584589005, + "median_ms": 0.08924119174480438, + "stdev_ms": 0.003103408004480225, + "min_ms": 0.08616968989372253, + "max_ms": 0.10702013969421387 + }, + "N=1000000,shift=3": { + "samples_raw_ms": [ + 1.346868, + 1.343202, + 1.344576, + 1.340825, + 1.339346, + 1.344271, + 1.340635, + 1.340043, + 1.344305, + 1.343489, + 1.339734, + 1.341939, + 1.340128, + 1.340102, + 1.341369, + 1.338579, + 1.339316, + 1.338154, + 1.338232, + 1.336016, + 1.337826, + 1.33704, + 1.336992, + 1.336593, + 1.337945, + 1.336839, + 1.339357, + 1.335349, + 1.335498, + 1.343962, + 1.349226, + 1.356393, + 1.337171, + 1.334772, + 1.335219, + 1.33815, + 1.343101, + 1.337707, + 1.337178, + 1.337744, + 1.338683, + 1.340717, + 1.339, + 1.338344, + 1.336578, + 1.33612, + 1.342952, + 1.336519, + 1.336772, + 1.3377, + 1.336586, + 1.340091, + 1.335457, + 1.334585, + 1.334488, + 1.336001, + 1.334991, + 1.3358, + 1.336623, + 1.338292, + 1.336835, + 1.336332, + 1.336429, + 1.336154, + 1.336616, + 1.337755, + 1.336738, + 1.337308, + 1.336675, + 1.335952, + 1.336798, + 1.336128, + 1.336168, + 1.33564, + 1.33767, + 1.336019, + 1.336593, + 1.336034, + 1.334552, + 1.336396, + 1.336023, + 1.335826, + 1.357414, + 1.357533, + 1.340996, + 1.340307, + 1.33799, + 1.335617, + 1.345411, + 1.345601, + 1.337957, + 1.337353, + 1.337498, + 1.336988, + 1.337774, + 1.338392, + 1.337975, + 1.338411, + 1.337651, + 1.336746 + ], + "mean_ms": 1.3388974964618683, + "median_ms": 1.337684690952301, + "stdev_ms": 0.004346946042564277, + "min_ms": 1.3344883918762207, + "max_ms": 1.357533037662506 + }, + "N=1000000,shift=12": { + "samples_raw_ms": [ + 1.344826, + 1.342811, + 1.341563, + 1.338191, + 1.339726, + 1.339354, + 1.337297, + 1.339965, + 1.338299, + 1.339275, + 1.341406, + 1.338795, + 1.33954, + 1.339663, + 1.337919, + 1.337104, + 1.336828, + 1.337096, + 1.3384, + 1.338288, + 1.337435, + 1.337677, + 1.338221, + 1.336727, + 1.337383, + 1.345214, + 1.337145, + 1.336083, + 1.349531, + 1.337845, + 1.343396, + 1.338195, + 1.338288, + 1.339983, + 1.337457, + 1.335952, + 1.339074, + 1.337267, + 1.335349, + 1.337521, + 1.338232, + 1.336712, + 1.339134, + 1.336753, + 1.335461, + 1.337487, + 1.335699, + 1.336042, + 1.337364, + 1.338124, + 1.335025, + 1.338486, + 1.335803, + 1.336493, + 1.336809, + 1.337592, + 1.336183, + 1.335129, + 1.334485, + 1.334824, + 1.33504, + 1.333695, + 1.333684, + 1.334105, + 1.334365, + 1.335446, + 1.335602, + 1.335338, + 1.334786, + 1.335274, + 1.33587, + 1.336608, + 1.336042, + 1.337789, + 1.336578, + 1.336321, + 1.337025, + 1.337081, + 1.337364, + 1.337558, + 1.33691, + 1.337189, + 1.351085, + 1.345713, + 1.33995, + 1.338493, + 1.336526, + 1.340214, + 1.349073, + 1.340814, + 1.342051, + 1.339089, + 1.339469, + 1.339812, + 1.33799, + 1.337036, + 1.33878, + 1.336373, + 1.335192, + 1.336206 + ], + "mean_ms": 1.338074654340744, + "median_ms": 1.3374090194702148, + "stdev_ms": 0.003123706782920795, + "min_ms": 1.333683729171753, + "max_ms": 1.3510845601558685 + } + }, + "summary": { + "definition": "speedup_b_over_c = latency(B) / latency(C); > 1 means C (Triton) is faster", + "cases": { + "N=25,shift=3": { + "b_median_ms": 0.0551324337720871, + "c_median_ms": 0.09256601333618164, + "speedup_b_over_c_median": 0.5956012556342563, + "speedup_b_over_c_mean": 0.5994931584326411, + "separated_from_dispersion": true + }, + "N=25,shift=12": { + "b_median_ms": 0.05539506673812866, + "c_median_ms": 0.09125098586082458, + "speedup_b_over_c_median": 0.6070626658501735, + "speedup_b_over_c_mean": 0.6063143397870051, + "separated_from_dispersion": true + }, + "N=50,shift=3": { + "b_median_ms": 0.05559995770454407, + "c_median_ms": 0.09175017476081848, + "speedup_b_over_c_median": 0.6059929351577409, + "speedup_b_over_c_mean": 0.6113383901177091, + "separated_from_dispersion": true + }, + "N=50,shift=12": { + "b_median_ms": 0.05539506673812866, + "c_median_ms": 0.08871778845787048, + "speedup_b_over_c_median": 0.6243963888305689, + "speedup_b_over_c_mean": 0.6263963135326815, + "separated_from_dispersion": true + }, + "N=1000,shift=3": { + "b_median_ms": 0.0597536563873291, + "c_median_ms": 0.09069032967090607, + "speedup_b_over_c_median": 0.6588757214155148, + "speedup_b_over_c_mean": 0.6595202658379822, + "separated_from_dispersion": true + }, + "N=1000,shift=12": { + "b_median_ms": 0.06091408431529999, + "c_median_ms": 0.08924119174480438, + "speedup_b_over_c_median": 0.6825781135856066, + "speedup_b_over_c_mean": 0.6785058127068553, + "separated_from_dispersion": true + }, + "N=1000000,shift=3": { + "b_median_ms": 5.50537183880806, + "c_median_ms": 1.337684690952301, + "speedup_b_over_c_median": 4.11559755153419, + "speedup_b_over_c_mean": 4.176128125963738, + "separated_from_dispersion": true + }, + "N=1000000,shift=12": { + "b_median_ms": 5.810923874378204, + "c_median_ms": 1.3374090194702148, + "speedup_b_over_c_median": 4.3449115340048134, + "speedup_b_over_c_mean": 4.425679426642436, + "separated_from_dispersion": true + } + } + } +} diff --git a/reports/experiments/h3-sigma-schedule/sm90-h800/env.txt b/reports/experiments/h3-sigma-schedule/sm90-h800/env.txt new file mode 100644 index 000000000..ee5c97f32 --- /dev/null +++ b/reports/experiments/h3-sigma-schedule/sm90-h800/env.txt @@ -0,0 +1,6 @@ +Sun Oct 25 17:32:03 UTC 2026 +NVIDIA H800 PCIe, 9.0, 580.82.07 +Python 3.12.3 +torch 2.5.1+cu124 cuda 12.4 triton 3.1.0 cc (9, 0) dev NVIDIA H800 PCIe +pytest 9.1.1 +HEAD 40070d37495d193e2d8764837321d3a68dad9483 diff --git a/reports/experiments/h3-sigma-schedule/sm90-h800/pr566_sm90_evidence.md b/reports/experiments/h3-sigma-schedule/sm90-h800/pr566_sm90_evidence.md new file mode 100644 index 000000000..f27ff2f39 --- /dev/null +++ b/reports/experiments/h3-sigma-schedule/sm90-h800/pr566_sm90_evidence.md @@ -0,0 +1,449 @@ +# PR #566 `h3_sigma_schedule` — SM90 CUDA bitwise evidence (NVIDIA H800 PCIe) + +Collected 2026-10-26 01:27–01:34 CST (2026-10-25 17:27–17:34 UTC) on a rented +AutoDL container (`CONTAINER-NAME-REDACTED`). Evidence-only run: +no source, test or documentation file was modified; nothing was pushed or +fetched from GitHub. The GPU name below is copied from `nvidia-smi` as printed. + +## Environment + +| Item | Value | +| --- | --- | +| GPU | `NVIDIA H800 PCIe` (nvidia-smi; 1 device, 81559 MiB; UUID GPU-UUID-REDACTED) | +| Compute capability | `9.0` (`nvidia-smi --query-gpu=compute_cap`); `torch.cuda.get_device_capability()` = `(9, 0)` | +| NVIDIA driver | `580.82.07` (nvidia-smi header reports driver-side CUDA Version 13.0) | +| CUDA runtime seen by torch | `12.4` | +| torch | `2.5.1+cu124` (pre-installed in the image; not touched) | +| Triton | `3.1.0` (pre-installed; not touched; identical to the RTX 4090D run) | +| Python | `3.12.3` CPython, `MINICONDA_PYTHON` | +| pytest | `9.1.1` (the only package installed for this run) | +| Host | x86_64, Linux 5.15.0-78-generic; 2× Intel Xeon Platinum 8458P (176 logical CPUs); ~1 TB RAM | +| Repository HEAD | `40070d37495d193e2d8764837321d3a68dad9483` — "Merge remote-tracking branch 'fork/h3-sigma-schedule' into h3-sigma-schedule" | +| Working tree | clean (`git status --short` printed nothing) | +| Repo path on host | `REPO_ROOT` | + +## Code transfer and version gates (step 2) + +The code was transferred from the developer machine, not fetched from GitHub: +a `git clone --branch h3-sigma-schedule --single-branch` of the local +repository was tarred (`rlk-pr566-40070d3.tgz`, 66 MB, sha256 +`3fef727cb277d412a1d2f31fbf56708c423ca5b2c0c4b632de14bbd20b8fd4c0`), copied +with scp, and the sha256 recomputed on the host matched. Extraction printed +harmless `LIBARCHIVE.xattr.com.apple.provenance` header warnings (macOS tar). +Because the tarball preserved the developer's uid, git refused with +"detected dubious ownership"; `chown -R root:root` on the extracted tree +cleared it (ownership only — no file content changed). + +``` +$ git log --oneline -8 +40070d3 Merge remote-tracking branch 'fork/h3-sigma-schedule' into h3-sigma-schedule +eae0118 Merge branch 'test-h3' into h3-sigma-schedule +a66c999 refactor(h3): default sigma_schedule_h3 to the PyTorch reference on CUDA +d22d4b4 Merge pull request #576 from fusheng-ji/feat/h3-flash-attention-v1 +ae8dff6 Merge branch 'test-h3' into feat/h3-flash-attention-v1 +8542cf7 docs(h3): document the missing-Triton RuntimeError for h3-flash-attention-v1 +4c2348f fix(h3): address CodeRabbit review on h3-flash-attention-v1 +6f59021 Merge branch 'test-h3' into feat/h3-flash-attention-v1 + +$ git status --short +(empty) + +$ git rev-parse HEAD +40070d37495d193e2d8764837321d3a68dad9483 + +$ git log --oneline | grep -c 'refactor(h3): default sigma_schedule_h3 to the PyTorch reference on CUDA' +1 + +$ grep -n sigma_schedule_h3 -A2 rl_engine/runtime/registry.py +431: semantic_op="sigma_schedule_h3", # semantic descriptor: reference +451: semantic_op="sigma_schedule_h3", # semantic descriptor: triton +739: "sigma_schedule_h3": [OpBackend.PYTORCH_H3_SIGMA_SCHEDULE], # cuda +763: "sigma_schedule_h3": [OpBackend.PYTORCH_H3_SIGMA_SCHEDULE], # rocm +807: "sigma_schedule_h3": [OpBackend.PYTORCH_H3_SIGMA_SCHEDULE], # musa +847: "sigma_schedule_h3": [OpBackend.PYTORCH_H3_SIGMA_SCHEDULE], # cpu +``` + +All gates satisfied: HEAD is `40070d3…`, the default-backend commit `a66c999` +is in the history, and every legacy platform list for `sigma_schedule_h3` +contains only `PYTORCH_H3_SIGMA_SCHEDULE`. + +## Test files (step 3) + +`git diff --stat origin/test-h3...HEAD -- tests` fails on this copy +(`fatal: bad revision` — the single-branch clone carries no `origin/test-h3` +ref). Fallbacks used: + +``` +$ git diff --stat 97a7ab9 HEAD -- tests # 97a7ab9 = PR base before the test-h3 merge + tests/backends/cuda/test_h3_flash_attention.py | 129 +++++++ <- PR #576 via test-h3, NOT this PR + tests/backends/cuda/test_h3_sigma_schedule_cuda.py | 392 +++++++++++++++++++++ + tests/reference/minimax_h3/__init__.py | 2 + + .../h3_sigma_schedule_upstream_fixture.json | 201 +++++++++++ + .../reference/minimax_h3/make_upstream_fixture.py | 154 ++++++++ + .../reference/minimax_h3/test_h3_sigma_schedule.py | 185 ++++++++++ + tests/runtime/test_h3_sigma_schedule_dispatch.py | 190 ++++++++++ + +$ git ls-files | grep -i h3_sigma +benchmarks/models/benchmark_h3_sigma_schedule.py +tests/backends/cuda/test_h3_sigma_schedule_cuda.py +tests/reference/minimax_h3/h3_sigma_schedule_upstream_fixture.json +tests/reference/minimax_h3/test_h3_sigma_schedule.py +tests/runtime/test_h3_sigma_schedule_dispatch.py +``` + +Sigma-schedule test files run: `tests/reference/minimax_h3/test_h3_sigma_schedule.py`, +`tests/runtime/test_h3_sigma_schedule_dispatch.py`, +`tests/backends/cuda/test_h3_sigma_schedule_cuda.py`. +`test_h3_flash_attention.py` belongs to PR #576 and was not run. + +All runs: `python -m pytest … -v -rs` from the repository root, same +interpreter, no `-x`. Logs and JUnit XML are in `EVIDENCE_DIR/`. + +## Results + +| Run | Command | Collected | Passed | Skipped | Failed | Errors | Wall | +| --- | --- | ---: | ---: | ---: | ---: | ---: | ---: | +| A — all H3 sigma tests | `python -m pytest tests/reference/minimax_h3 tests/runtime/test_h3_sigma_schedule_dispatch.py tests/backends/cuda/test_h3_sigma_schedule_cuda.py -v -rs` | 69 | **67** | 2 | **0** | 0 | 4.86 s | +| B — CUDA/Triton/bitwise + dispatch subset | `python -m pytest tests/backends/cuda/test_h3_sigma_schedule_cuda.py tests/runtime/test_h3_sigma_schedule_dispatch.py -v -rs` | 48 | **47** | 1 | **0** | 0 | 2.89 s | +| C — the two new/renamed registry tests | `python -m pytest "tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_public_entry_defaults_to_reference_and_matches_triton_bits" "tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_public_entry_opt_in_selects_triton_and_reports_provenance" -v -rs` | 2 | **2** | 0 | **0** | 0 | 2.51 s | + +Run A by file: `test_h3_sigma_schedule.py` 20 passed / 1 skipped; +`test_h3_sigma_schedule_dispatch.py` 12 passed; `test_h3_sigma_schedule_cuda.py` +35 passed / 1 skipped. + +### Skips (both runs, same two) + +| Test | Reason printed by `-rs` | Explanation | +| --- | --- | --- | +| `tests/reference/minimax_h3/test_h3_sigma_schedule.py:171` `test_matches_installed_diffusers_scheduler_when_available` | `could not import 'diffusers': No module named 'diffusers'` | Optional dependency not installed on this host (it was not installed to avoid touching the pinned torch). The pinned fixture test covers the upstream parity and passed. | +| `tests/backends/cuda/test_h3_sigma_schedule_cuda.py:378` `test_target_device_guard_with_mismatched_current_device` | `multi-GPU device-guard test needs 2 devices` | This host has exactly one GPU (`nvidia-smi` lists GPU 0 only; `--query-gpu=count` = 1). This is the same test skipped for the same reason on the RTX 4090D host. | + +### Run C — singled out (first execution on any GPU) + +Both tests were introduced/renamed in `a66c999` after the RTX 4090D runs, so +this is their first execution on hardware. + +- `test_public_entry_defaults_to_reference_and_matches_triton_bits` — **PASSED**. + `kernel_registry.get_op("sigma_schedule_h3", device="cuda")` returned + `NativeH3SigmaScheduleOp`; its outputs built on CUDA (`N=25, shift=12.0`) + compared bit-for-bit equal (`_int32_bits`) to `H3SigmaScheduleCudaOp()`'s + `sigmas` and `timesteps`, and `num_model_evals` matched. +- `test_public_entry_opt_in_selects_triton_and_reports_provenance` — **PASSED**. + Semantic `requested_backend="rlkernel.h3_sigma_schedule.triton"` + instantiated `H3SigmaScheduleCudaOp` and provenance reported that backend. + +No bitwise test failed, so step 5 (mismatch forensics) was not triggered. + +## Benchmark (step 3d, optional) + +``` +python benchmarks/models/benchmark_h3_sigma_schedule.py --samples 100 --warmup 5 --json ~/pr566_evidence/benchmark-h800.json +``` + +Pre-timing correctness gate: path C (Triton, device build) bitwise equal to +path A (CPU reference) for all 8 cases — **passed**; Triton compilation +1143.6 ms wall, reported once and excluded from timed samples. Raw JSON: +`EVIDENCE_DIR/benchmark-h800.json` (52 KB, 100 samples per case and +path after 5 warmups). + +| Workload | A median (ms) | B median (ms) | C median (ms) | B/C | +| --- | ---: | ---: | ---: | ---: | +| N=25, shift=3 | 0.0306 | 0.0551 | 0.0926 | 0.60x (separated) | +| N=25, shift=12 | 0.0305 | 0.0554 | 0.0913 | 0.61x (separated) | +| N=50, shift=3 | 0.0308 | 0.0556 | 0.0918 | 0.61x (separated) | +| N=50, shift=12 | 0.0310 | 0.0554 | 0.0887 | 0.62x (separated) | +| N=1000, shift=3 | 0.0347 | 0.0598 | 0.0907 | 0.66x (separated) | +| N=1000, shift=12 | 0.0346 | 0.0609 | 0.0892 | 0.68x (separated) | +| N=1000000, shift=3 | 5.4651 | 5.5054 | 1.3377 | 4.12x (separated) | +| N=1000000, shift=12 | 5.5773 | 5.8109 | 1.3374 | 4.34x (separated) | + +A = reference, CPU build, CPU outputs; B = reference, CPU build, outputs +copied to CUDA (the registry default path); C = Triton opt-in, device build. +"separated" = the benchmark's own dispersion check passed. + +Caveats (why these numbers are supplementary only): the H800 PCIe has reduced +FP64 throughput and the Triton kernel evaluates its intermediates in fp64; +paths A/B run on this host's Xeon 8458P, which is a different CPU from the +4090D host (at N=1e6 the CPU path here is ~5.5 ms versus ~48 ms there); +single-operator microbenchmark, one card. On this host the device build is +about 1.5–1.65x slower than CPU build plus copy at N=25/50/1000 and about +4.1–4.3x faster at the N=1e6 dedup stress case. The reviewer's question is +answered by the bitwise results above, not by this table. + +## Observations (not failures) + +- `WARNING … Core binary extension (_C) unavailable … Falling back to native code` + was printed once during the benchmark. The repository was not `pip install -e .`-built + (its declared `transformers==5.13.1`/`accelerate` dependencies were not installed to + keep torch untouched); `h3_sigma_schedule` does not use `_C`. No test printed this + warning. +- Only `pytest` was installed; torch 2.5.1+cu124 and Triton 3.1.0 are the image's. +- `git diff --stat origin/test-h3...HEAD` is unavailable on this copy (no such ref); + `97a7ab9` and `git ls-files` fallbacks were used as instructed. + +## Factual conclusion + +On NVIDIA H800 PCIe (SM90, compute capability 9.0, driver 580.82.07, torch +2.5.1+cu124, CUDA 12.4, Triton 3.1.0) at HEAD `40070d37495d193e2d8764837321d3a68dad9483`: + +- All 35 runnable tests in `tests/backends/cuda/test_h3_sigma_schedule_cuda.py` + passed — pinned-fixture bitwise match, live-CPU-reference bitwise match + (N=50/1000, shift 3/12), dense-grid dedup, grid invariants, subnormal + (1e-45) / underflow (1e-46) / ±NaN / inf / 1e39 / 3.5e38 / 1e30 / 1e-10 + edge parity, float-N truncation, input rejections, repeated/interleaved + stability, backend identity, default-reference-vs-Triton bit equality, and + semantic Triton opt-in. 1 test skipped (needs 2 GPUs; host has 1). +- All 12 dispatch tests passed. 20 reference tests passed, 1 skipped + (`diffusers` not installed). +- 0 failures and 0 errors across runs A, B and C. +- Compared with the historical RTX 4090D CUDA run (34 passed / 1 skipped, + torch 2.5.1+cu121, Triton 3.1.0, driver 580.76.05): one more test passed + here because `a66c999` split the registry-identity test into two; the one + skip is the same test for the same reason; no test that passed on the + 4090D failed here. Triton version is identical; torch differs only in the + CUDA build (cu124 vs cu121); driver 580.82.07 vs 580.76.05. +- Benchmark correctness gate passed on all 8 cases; latency ratios differ + from the 4090D (see caveats) and are reported as supplementary data only. + +Evidence files: `EVIDENCE_DIR/{env.txt, run_a.log, run_a.xml, run_b.log, +run_b.xml, run_c.log, run_c.xml, bench.log, benchmark-h800.json}` on the host; +a copy was pulled to the developer machine. + +## Raw outputs + +
env.txt + +```text +Sun Oct 25 17:32:03 UTC 2026 +NVIDIA H800 PCIe, 9.0, 580.82.07 +Python 3.12.3 +torch 2.5.1+cu124 cuda 12.4 triton 3.1.0 cc (9, 0) dev NVIDIA H800 PCIe +pytest 9.1.1 +HEAD 40070d37495d193e2d8764837321d3a68dad9483 +``` + +
+ +
run_a.log + +```text +============================= test session starts ============================== +platform linux -- Python 3.12.3, pytest-9.1.1, pluggy-1.7.0 -- MINICONDA_PYTHON +cachedir: .pytest_cache +rootdir: REPO_ROOT +configfile: pyproject.toml +plugins: anyio-4.6.2.post1 +collecting ... collected 69 items + +tests/reference/minimax_h3/test_h3_sigma_schedule.py::test_matches_pinned_upstream_set_timesteps_bitwise PASSED [ 1%] +tests/reference/minimax_h3/test_h3_sigma_schedule.py::test_grid_invariants[2-12.0] PASSED [ 2%] +tests/reference/minimax_h3/test_h3_sigma_schedule.py::test_grid_invariants[2-3.0] PASSED [ 4%] +tests/reference/minimax_h3/test_h3_sigma_schedule.py::test_grid_invariants[8-12.0] PASSED [ 5%] +tests/reference/minimax_h3/test_h3_sigma_schedule.py::test_grid_invariants[8-3.0] PASSED [ 7%] +tests/reference/minimax_h3/test_h3_sigma_schedule.py::test_grid_invariants[25-12.0] PASSED [ 8%] +tests/reference/minimax_h3/test_h3_sigma_schedule.py::test_grid_invariants[25-3.0] PASSED [ 10%] +tests/reference/minimax_h3/test_h3_sigma_schedule.py::test_grid_invariants[50-12.0] PASSED [ 11%] +tests/reference/minimax_h3/test_h3_sigma_schedule.py::test_grid_invariants[50-3.0] PASSED [ 13%] +tests/reference/minimax_h3/test_h3_sigma_schedule.py::test_grid_invariants[1000-12.0] PASSED [ 14%] +tests/reference/minimax_h3/test_h3_sigma_schedule.py::test_grid_invariants[1000-3.0] PASSED [ 15%] +tests/reference/minimax_h3/test_h3_sigma_schedule.py::test_unique_consecutive_collapses_shift_created_collisions PASSED [ 17%] +tests/reference/minimax_h3/test_h3_sigma_schedule.py::test_video_and_audio_shifts_give_distinct_grids PASSED [ 18%] +tests/reference/minimax_h3/test_h3_sigma_schedule.py::test_rejects_num_inference_steps_below_two[1] PASSED [ 20%] +tests/reference/minimax_h3/test_h3_sigma_schedule.py::test_rejects_num_inference_steps_below_two[0] PASSED [ 21%] +tests/reference/minimax_h3/test_h3_sigma_schedule.py::test_rejects_num_inference_steps_below_two[-4] PASSED [ 23%] +tests/reference/minimax_h3/test_h3_sigma_schedule.py::test_rejects_non_positive_shift[0.0] PASSED [ 24%] +tests/reference/minimax_h3/test_h3_sigma_schedule.py::test_rejects_non_positive_shift[-1.5] PASSED [ 26%] +tests/reference/minimax_h3/test_h3_sigma_schedule.py::test_device_argument_preserves_bytes_and_dtype PASSED [ 27%] +tests/reference/minimax_h3/test_h3_sigma_schedule.py::test_builds_on_cpu_under_non_cpu_default_device PASSED [ 28%] +tests/reference/minimax_h3/test_h3_sigma_schedule.py::test_matches_installed_diffusers_scheduler_when_available SKIPPED [ 30%] +tests/runtime/test_h3_sigma_schedule_dispatch.py::test_get_op_cpu_selects_reference_and_builds_pinned_grid PASSED [ 31%] +tests/runtime/test_h3_sigma_schedule_dispatch.py::test_get_op_cuda_defaults_to_reference_regardless_of_triton PASSED [ 33%] +tests/runtime/test_h3_sigma_schedule_dispatch.py::test_cuda_backend_construction_fails_loudly_without_triton PASSED [ 34%] +tests/runtime/test_h3_sigma_schedule_dispatch.py::test_shift_scalars_pinned_to_cpu_under_meta_default PASSED [ 36%] +tests/runtime/test_h3_sigma_schedule_dispatch.py::test_shift_scalars_keep_fp64_subtract_then_round_semantics PASSED [ 37%] +tests/runtime/test_h3_sigma_schedule_dispatch.py::test_cuda_backend_construction_rejected_on_hip_build PASSED [ 39%] +tests/runtime/test_h3_sigma_schedule_dispatch.py::test_get_op_on_hip_build_serves_reference PASSED [ 40%] +tests/runtime/test_h3_sigma_schedule_dispatch.py::test_triton_backend_is_not_a_legacy_candidate_on_any_platform PASSED [ 42%] +tests/runtime/test_h3_sigma_schedule_dispatch.py::test_semantic_descriptors_declared_truthfully PASSED [ 43%] +tests/runtime/test_h3_sigma_schedule_dispatch.py::test_semantic_reference_path_resolves_instantiates_and_reports_provenance PASSED [ 44%] +tests/runtime/test_h3_sigma_schedule_dispatch.py::test_semantic_triton_backend_fails_explicitly_when_unavailable PASSED [ 46%] +tests/runtime/test_h3_sigma_schedule_dispatch.py::test_semantic_unknown_backend_is_rejected PASSED [ 47%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_matches_pinned_cpu_fixture_bitwise PASSED [ 49%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_matches_live_cpu_reference_bitwise[50-12.0] PASSED [ 50%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_matches_live_cpu_reference_bitwise[50-3.0] PASSED [ 52%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_matches_live_cpu_reference_bitwise[1000-12.0] PASSED [ 53%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_dense_grid_dedup_matches_cpu_reference PASSED [ 55%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_reference_with_cuda_device_preserves_fixture_bytes PASSED [ 56%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_grid_invariants_on_device[2-3.0] PASSED [ 57%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_grid_invariants_on_device[2-12.0] PASSED [ 59%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_grid_invariants_on_device[4-3.0] PASSED [ 60%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_grid_invariants_on_device[4-12.0] PASSED [ 62%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_grid_invariants_on_device[25-3.0] PASSED [ 63%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_grid_invariants_on_device[25-12.0] PASSED [ 65%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_rejects_num_inference_steps_below_two[1] PASSED [ 66%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_rejects_num_inference_steps_below_two[0] PASSED [ 68%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_rejects_num_inference_steps_below_two[-4] PASSED [ 69%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_subnormal_shift_1e_45_matches_reference_upstream_parity PASSED [ 71%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_underflow_shift_1e_46_matches_reference_upstream_parity PASSED [ 72%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_positive_nan_shift_propagates_positive_quiet_nan PASSED [ 73%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_negative_nan_shift_propagates_negative_quiet_nan PASSED [ 75%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_edge_shifts_match_reference_upstream_parity[nan] PASSED [ 76%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_edge_shifts_match_reference_upstream_parity[inf] PASSED [ 78%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_edge_shifts_match_reference_upstream_parity[1e-10] PASSED [ 79%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_edge_shifts_match_reference_upstream_parity[1e+39] PASSED [ 81%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_edge_shifts_match_reference_upstream_parity[3.5e+38] PASSED [ 82%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_edge_shifts_match_reference_upstream_parity[1e+30] PASSED [ 84%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_float_num_steps_truncate_like_reference[2.9-2] PASSED [ 85%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_float_num_steps_truncate_like_reference[3.1-3] PASSED [ 86%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_rejects_non_positive_shift[0.0] PASSED [ 88%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_rejects_non_positive_shift[-1.5] PASSED [ 89%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_rejects_non_positive_shift[-inf] PASSED [ 91%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_rejects_cpu_device_instead_of_silent_fallback PASSED [ 92%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_repeated_and_interleaved_calls_are_stable PASSED [ 94%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_backend_identity_is_device_path_not_cpu_fallback PASSED [ 95%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_public_entry_defaults_to_reference_and_matches_triton_bits PASSED [ 97%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_public_entry_opt_in_selects_triton_and_reports_provenance PASSED [ 98%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_target_device_guard_with_mismatched_current_device SKIPPED [100%] + +-------------- generated xml file: EVIDENCE_DIR/run_a.xml -------------- +=========================== short test summary info ============================ +SKIPPED [1] tests/reference/minimax_h3/test_h3_sigma_schedule.py:171: could not import 'diffusers': No module named 'diffusers' +SKIPPED [1] tests/backends/cuda/test_h3_sigma_schedule_cuda.py:378: multi-GPU device-guard test needs 2 devices +======================== 67 passed, 2 skipped in 4.86s ========================= +``` + +
+ +
run_b.log + +```text +============================= test session starts ============================== +platform linux -- Python 3.12.3, pytest-9.1.1, pluggy-1.7.0 -- MINICONDA_PYTHON +cachedir: .pytest_cache +rootdir: REPO_ROOT +configfile: pyproject.toml +plugins: anyio-4.6.2.post1 +collecting ... collected 48 items + +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_matches_pinned_cpu_fixture_bitwise PASSED [ 2%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_matches_live_cpu_reference_bitwise[50-12.0] PASSED [ 4%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_matches_live_cpu_reference_bitwise[50-3.0] PASSED [ 6%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_matches_live_cpu_reference_bitwise[1000-12.0] PASSED [ 8%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_dense_grid_dedup_matches_cpu_reference PASSED [ 10%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_reference_with_cuda_device_preserves_fixture_bytes PASSED [ 12%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_grid_invariants_on_device[2-3.0] PASSED [ 14%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_grid_invariants_on_device[2-12.0] PASSED [ 16%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_grid_invariants_on_device[4-3.0] PASSED [ 18%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_grid_invariants_on_device[4-12.0] PASSED [ 20%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_grid_invariants_on_device[25-3.0] PASSED [ 22%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_grid_invariants_on_device[25-12.0] PASSED [ 25%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_rejects_num_inference_steps_below_two[1] PASSED [ 27%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_rejects_num_inference_steps_below_two[0] PASSED [ 29%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_rejects_num_inference_steps_below_two[-4] PASSED [ 31%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_subnormal_shift_1e_45_matches_reference_upstream_parity PASSED [ 33%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_underflow_shift_1e_46_matches_reference_upstream_parity PASSED [ 35%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_positive_nan_shift_propagates_positive_quiet_nan PASSED [ 37%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_negative_nan_shift_propagates_negative_quiet_nan PASSED [ 39%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_edge_shifts_match_reference_upstream_parity[nan] PASSED [ 41%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_edge_shifts_match_reference_upstream_parity[inf] PASSED [ 43%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_edge_shifts_match_reference_upstream_parity[1e-10] PASSED [ 45%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_edge_shifts_match_reference_upstream_parity[1e+39] PASSED [ 47%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_edge_shifts_match_reference_upstream_parity[3.5e+38] PASSED [ 50%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_edge_shifts_match_reference_upstream_parity[1e+30] PASSED [ 52%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_float_num_steps_truncate_like_reference[2.9-2] PASSED [ 54%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_float_num_steps_truncate_like_reference[3.1-3] PASSED [ 56%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_rejects_non_positive_shift[0.0] PASSED [ 58%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_rejects_non_positive_shift[-1.5] PASSED [ 60%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_rejects_non_positive_shift[-inf] PASSED [ 62%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_rejects_cpu_device_instead_of_silent_fallback PASSED [ 64%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_repeated_and_interleaved_calls_are_stable PASSED [ 66%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_backend_identity_is_device_path_not_cpu_fallback PASSED [ 68%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_public_entry_defaults_to_reference_and_matches_triton_bits PASSED [ 70%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_public_entry_opt_in_selects_triton_and_reports_provenance PASSED [ 72%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_target_device_guard_with_mismatched_current_device SKIPPED [ 75%] +tests/runtime/test_h3_sigma_schedule_dispatch.py::test_get_op_cpu_selects_reference_and_builds_pinned_grid PASSED [ 77%] +tests/runtime/test_h3_sigma_schedule_dispatch.py::test_get_op_cuda_defaults_to_reference_regardless_of_triton PASSED [ 79%] +tests/runtime/test_h3_sigma_schedule_dispatch.py::test_cuda_backend_construction_fails_loudly_without_triton PASSED [ 81%] +tests/runtime/test_h3_sigma_schedule_dispatch.py::test_shift_scalars_pinned_to_cpu_under_meta_default PASSED [ 83%] +tests/runtime/test_h3_sigma_schedule_dispatch.py::test_shift_scalars_keep_fp64_subtract_then_round_semantics PASSED [ 85%] +tests/runtime/test_h3_sigma_schedule_dispatch.py::test_cuda_backend_construction_rejected_on_hip_build PASSED [ 87%] +tests/runtime/test_h3_sigma_schedule_dispatch.py::test_get_op_on_hip_build_serves_reference PASSED [ 89%] +tests/runtime/test_h3_sigma_schedule_dispatch.py::test_triton_backend_is_not_a_legacy_candidate_on_any_platform PASSED [ 91%] +tests/runtime/test_h3_sigma_schedule_dispatch.py::test_semantic_descriptors_declared_truthfully PASSED [ 93%] +tests/runtime/test_h3_sigma_schedule_dispatch.py::test_semantic_reference_path_resolves_instantiates_and_reports_provenance PASSED [ 95%] +tests/runtime/test_h3_sigma_schedule_dispatch.py::test_semantic_triton_backend_fails_explicitly_when_unavailable PASSED [ 97%] +tests/runtime/test_h3_sigma_schedule_dispatch.py::test_semantic_unknown_backend_is_rejected PASSED [100%] + +-------------- generated xml file: EVIDENCE_DIR/run_b.xml -------------- +=========================== short test summary info ============================ +SKIPPED [1] tests/backends/cuda/test_h3_sigma_schedule_cuda.py:378: multi-GPU device-guard test needs 2 devices +======================== 47 passed, 1 skipped in 2.89s ========================= +``` + +
+ +
run_c.log + +```text +============================= test session starts ============================== +platform linux -- Python 3.12.3, pytest-9.1.1, pluggy-1.7.0 -- MINICONDA_PYTHON +cachedir: .pytest_cache +rootdir: REPO_ROOT +configfile: pyproject.toml +plugins: anyio-4.6.2.post1 +collecting ... collected 2 items + +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_public_entry_defaults_to_reference_and_matches_triton_bits PASSED [ 50%] +tests/backends/cuda/test_h3_sigma_schedule_cuda.py::test_public_entry_opt_in_selects_triton_and_reports_provenance PASSED [100%] + +-------------- generated xml file: EVIDENCE_DIR/run_c.xml -------------- +============================== 2 passed in 2.51s =============================== +``` + +
+ +
bench.log + +```text +path A (reference, cpu build, cpu outputs): 100 calls x [25, 50, 1000, 1000000] x [3.0, 12.0] +verifying path C against path A bitwise before timing... +WARNING 10-26 01:33:03 [RL-Kernel]: Core binary extension (_C) unavailable: cannot import name '_C' from 'rl_engine' (REPO_ROOT/rl_engine/__init__.py). Falling back to native code. +correctness gate passed (includes Triton kernel compilation: 1143.6 ms wall, reported once, excluded from all timed samples) +path B (reference, cpu build, outputs moved to gpu): 100 calls +path C (triton, device build): 100 calls + path_a_reference_cpu N=25,shift=3: 0.0306 ms/call (median) + path_a_reference_cpu N=25,shift=12: 0.0305 ms/call (median) + path_a_reference_cpu N=50,shift=3: 0.0308 ms/call (median) + path_a_reference_cpu N=50,shift=12: 0.0310 ms/call (median) + path_a_reference_cpu N=1000,shift=3: 0.0347 ms/call (median) + path_a_reference_cpu N=1000,shift=12: 0.0346 ms/call (median) + path_a_reference_cpu N=1000000,shift=3: 5.4651 ms/call (median) + path_a_reference_cpu N=1000000,shift=12: 5.5773 ms/call (median) + path_b_reference_cpu_build_then_copy_to_gpu N=25,shift=3: 0.0551 ms/call (median) + path_b_reference_cpu_build_then_copy_to_gpu N=25,shift=12: 0.0554 ms/call (median) + path_b_reference_cpu_build_then_copy_to_gpu N=50,shift=3: 0.0556 ms/call (median) + path_b_reference_cpu_build_then_copy_to_gpu N=50,shift=12: 0.0554 ms/call (median) + path_b_reference_cpu_build_then_copy_to_gpu N=1000,shift=3: 0.0598 ms/call (median) + path_b_reference_cpu_build_then_copy_to_gpu N=1000,shift=12: 0.0609 ms/call (median) + path_b_reference_cpu_build_then_copy_to_gpu N=1000000,shift=3: 5.5054 ms/call (median) + path_b_reference_cpu_build_then_copy_to_gpu N=1000000,shift=12: 5.8109 ms/call (median) + path_c_triton_device_build N=25,shift=3: 0.0926 ms/call (median) + path_c_triton_device_build N=25,shift=12: 0.0913 ms/call (median) + path_c_triton_device_build N=50,shift=3: 0.0918 ms/call (median) + path_c_triton_device_build N=50,shift=12: 0.0887 ms/call (median) + path_c_triton_device_build N=1000,shift=3: 0.0907 ms/call (median) + path_c_triton_device_build N=1000,shift=12: 0.0892 ms/call (median) + path_c_triton_device_build N=1000000,shift=3: 1.3377 ms/call (median) + path_c_triton_device_build N=1000000,shift=12: 1.3374 ms/call (median) +speedup B/C N=25,shift=3: 0.60x (separated from dispersion) +speedup B/C N=25,shift=12: 0.61x (separated from dispersion) +speedup B/C N=50,shift=3: 0.61x (separated from dispersion) +speedup B/C N=50,shift=12: 0.62x (separated from dispersion) +speedup B/C N=1000,shift=3: 0.66x (separated from dispersion) +speedup B/C N=1000,shift=12: 0.68x (separated from dispersion) +speedup B/C N=1000000,shift=3: 4.12x (separated from dispersion) +speedup B/C N=1000000,shift=12: 4.34x (separated from dispersion) +wrote EVIDENCE_DIR/benchmark-h800.json +``` + +
diff --git a/reports/experiments/h3-sigma-schedule/sm90-h800/run_a.xml b/reports/experiments/h3-sigma-schedule/sm90-h800/run_a.xml new file mode 100644 index 000000000..45e368791 --- /dev/null +++ b/reports/experiments/h3-sigma-schedule/sm90-h800/run_a.xml @@ -0,0 +1 @@ +REPO_ROOT/tests/reference/minimax_h3/test_h3_sigma_schedule.py:171: could not import 'diffusers': No module named 'diffusers'REPO_ROOT/tests/backends/cuda/test_h3_sigma_schedule_cuda.py:378: multi-GPU device-guard test needs 2 devices diff --git a/reports/experiments/h3-sigma-schedule/sm90-h800/run_b.xml b/reports/experiments/h3-sigma-schedule/sm90-h800/run_b.xml new file mode 100644 index 000000000..2d76b25c0 --- /dev/null +++ b/reports/experiments/h3-sigma-schedule/sm90-h800/run_b.xml @@ -0,0 +1 @@ +REPO_ROOT/tests/backends/cuda/test_h3_sigma_schedule_cuda.py:378: multi-GPU device-guard test needs 2 devices diff --git a/reports/experiments/h3-sigma-schedule/sm90-h800/run_c.xml b/reports/experiments/h3-sigma-schedule/sm90-h800/run_c.xml new file mode 100644 index 000000000..069554b98 --- /dev/null +++ b/reports/experiments/h3-sigma-schedule/sm90-h800/run_c.xml @@ -0,0 +1 @@ + diff --git a/reports/experiments/h3-sigma-schedule/source-provenance.json b/reports/experiments/h3-sigma-schedule/source-provenance.json index 391daedfb..5db435dee 100644 --- a/reports/experiments/h3-sigma-schedule/source-provenance.json +++ b/reports/experiments/h3-sigma-schedule/source-provenance.json @@ -1,5 +1,5 @@ { - "evidence_kind": "Historical hardware results; source binding rechecked 2026-10-21; no new GPU run", + "evidence_kind": "Historical 4090D operator results (source binding rechecked 2026-10-21) plus an SM90 (H800) hardware run on 2026-10-26.", "candidate_commit": "374488bb3abc856b356b79fbfbcc0e15bbb968c4", "historical_snapshot_base_commit": "0e8a18aff441895082aee603959cd25cb66d7e4b", "environment": { @@ -39,7 +39,7 @@ "benchmark": { "run": "run-20261016-235221-8055", "snapshot_sha256": "0b7c25c6a9bf4e109fe0e053a90fed1a977831d859f2bf99575cfcfe08da9b34", - "binding": "Benchmark and reference byte-identical. CUDA kernels and scalar helper ASTs identical. Forward AST identical after normalizing only the non-CUDA device error text; final constructor adds HIP rejection.", + "binding": "Benchmark and reference byte-identical. CUDA kernels and scalar helper ASTs identical. Forward AST identical after normalizing only the non-CUDA device error text; final constructor adds HIP rejection. Post-run note (2026-10-26): the benchmark file later changed twice — the docstring / path-B wording (a66c999) and the GPU-without-Triton + HIP/ROCm branch fix (ace9c20). Both 4090D measurements above ran on the pre-change file; the NVIDIA + Triton execution path is unchanged by both commits.", "cuda_function_asts_identical": [ "_shift_scalars", "_shifted_grid_kernel", @@ -54,8 +54,84 @@ "cuda-4090d.xml": "c5377af475a7020c02f0c1784bd87df6412ca4f8cfc42f4687dd8f8fac1900e2" }, "limits": [ - "One x86_64 NVIDIA RTX 4090D; no ROCm device backend.", - "No fresh GPU measurement, multi-GPU test, or full-model run.", + "One x86_64 NVIDIA RTX 4090D (October 2026) and one x86_64 NVIDIA H800 PCIe SM90 host (2026-10-26); no ROCm device backend.", + "No multi-GPU test or full-model run on either host.", + "The benchmark's GPU-without-Triton and HIP/ROCm branches (ace9c20) have not run on GPU or HIP hardware; the H800 latency table is supplementary (reduced FP64 throughput, different host CPU) and not comparable with the 4090D numbers.", "Historical broad dispatch suite had two attention failures also present in the recorded baseline A/B comparison; not an all-green pipeline claim." - ] + ], + "sm90_validation": { + "run_date": "2026-10-26", + "tested_commit": "40070d37495d193e2d8764837321d3a68dad9483", + "environment": { + "gpu": "NVIDIA H800 PCIe", + "compute_capability": "9.0", + "driver": "580.82.07", + "platform": "x86_64 Linux 5.15.0-78-generic; 2x Intel Xeon Platinum 8458P", + "python": "3.12.3", + "torch": "2.5.1+cu124", + "torch_cuda_runtime": "12.4", + "torch_hip": null, + "triton": "3.1.0", + "pytest": "9.1.1 (the only package installed on the host)" + }, + "run_a": { + "collected": 69, + "passed": 67, + "skipped": 2, + "failures": 0, + "errors": 0, + "skip_reasons": [ + "could not import 'diffusers' (optional parity test; fixture test passed)", + "multi-GPU device-guard test needs 2 devices; host had 1" + ] + }, + "run_b": { + "collected": 48, + "passed": 47, + "skipped": 1, + "failures": 0, + "errors": 0 + }, + "run_c": { + "collected": 2, + "passed": 2, + "note": "first hardware execution of the two registry tests introduced by a66c999; the default reference path and the Triton backend produced identical bits" + }, + "source_files_sha256_at_tested_commit": { + "rl_engine/reference/minimax_h3/sigma_schedule.py": "18b36cdbac9e95d4b52782aa17a7228fdf35bc7e3397508a651c1fb07cb36509", + "rl_engine/backends/cuda/model_specific/minimax_h3/sigma_schedule.py": "784fb85f2e3dbce450fd5564c99584eb75a83c8b01374d938f036c7036494262", + "tests/reference/minimax_h3/test_h3_sigma_schedule.py": "1304b4738196250bd09c5532f39dd15b14e1ff46910ef4a807067c5a8b1faeee", + "tests/reference/minimax_h3/h3_sigma_schedule_upstream_fixture.json": "97a929a39eb781ce00bd351d8c30c85d9bfd054e7ca5405dc3a3b8ed794956c1", + "tests/runtime/test_h3_sigma_schedule_dispatch.py": "c14a503dcdc82ccfc2c3c3a8da695038bfdaaa65db1d5f4ae87fbd4bb19e686f", + "tests/backends/cuda/test_h3_sigma_schedule_cuda.py": "82a584cadadebdf139ed9871ba2165d492c94727d024cffbf48bc079013c176b", + "benchmarks/models/benchmark_h3_sigma_schedule.py": "5c1b135c5b723cfcdfc6a9dc9c542601bbc3067a62f28ae75bd8fa2772059ec2" + }, + "included_files_sha256": { + "sm90-h800/env.txt": "13d42ff62485307c55d4af5efe7fb6e7633a8f37a869d975d70249b1acaaaebf", + "sm90-h800/run_a.log": "e94695df1a4bab87c590e51b2231b179d91dc935674ad39a9388508f5854b593", + "sm90-h800/run_a.xml": "4797238671384fbeb1ad9ddd54fa18ba48d760d18d69207f354ca291c3882604", + "sm90-h800/run_b.log": "62bf01571a0e70e8e095ed7e59fe5dc0e68b760e51f9abeca658cbd33b938aee", + "sm90-h800/run_b.xml": "86cfec27c58060990317720777ce63b51dc208d09962445f1bca3a8121e5a658", + "sm90-h800/run_c.log": "7b46d39f310343f35776b76dc26405bb67bc9c047abd8bc1c987724de04cc12b", + "sm90-h800/run_c.xml": "e2ea3433d80e7634976819856123659d6dafa8b3c747b964afc192e90ca802d2", + "sm90-h800/bench.log": "30acf2fb3ae50aed57efad735565b4225021ed13b9c81e1165995c8cf09b5224", + "sm90-h800/benchmark-h800.json": "769b89b16679bdc5db365a2cc10f1dc7cc1966d72fbcf5ce0aae5a9d61b73d4c", + "sm90-h800/pr566_sm90_evidence.md": "5ecf4b0672fd41bf74f9ce9f0484ed07fce252e3c6aadc490a5dd8b1080a9e81" + }, + "original_server_sha256": { + "sm90-h800/env.txt": "13d42ff62485307c55d4af5efe7fb6e7633a8f37a869d975d70249b1acaaaebf", + "sm90-h800/run_a.log": "565ae81515f8d1111f7043b1e195c7ad8196cd031a90d38bb73c407e21be2240", + "sm90-h800/run_a.xml": "66889f04f9f162c2c36d8e21b148c5d9d33596550cd71c13ed91b5d339ceb9d0", + "sm90-h800/run_b.log": "ad200760d36aad6802506b43b5ec7bc043eb55ef25907bd285c50962a25fee2a", + "sm90-h800/run_b.xml": "1d5767d44b77c4bb489eb7a3c602791bc824c97082f9ac363bc2482f0211807e", + "sm90-h800/run_c.log": "6f3a91c8d97cf81e8e00a38b559fe0f3dc52bb1aa8313a1723de754a993ea7e1", + "sm90-h800/run_c.xml": "bda3e575dc09ec32463dfbc8fb2cdc9a6259f68ae14cdc53a529cdc5be2d73c6", + "sm90-h800/bench.log": "da43426479f55f16cb4d1202691b1234a91434265f18d9419af91aa26ecae845", + "sm90-h800/benchmark-h800.json": "50868bae541aee093465fb8e5998af0d6d80f967538763fca6b1ce8a491b97b4", + "sm90-h800/pr566_sm90_evidence.md": "f4edb693161ce5a212c4da47f9e2737f3b35ec1f78ddd2908700f0e76e81e0c6" + }, + "benchmark_gate": "path C (Triton) bitwise equal to path A (CPU reference) for all 8 cases; gate wall incl. Triton compilation 1143.6 ms, excluded from timed samples; benchmark-h800.json is byte-identical to the server original except an appended final newline", + "redaction": "GPU UUID, container hostname and /root/... paths removed from the committed logs, JUnit XML and report (JUnit hostname attribute dropped, as for the 4090D XML; host paths additionally replaced by REPO_ROOT / EVIDENCE_DIR / MINICONDA_PYTHON tokens, one step beyond the 4090D precedent). A final newline was appended to the three XMLs and to benchmark-h800.json for the end-of-file-fixer hook; env.txt needed no change. original_server_sha256 records the pre-redaction originals (verified 10/10 against the host before it was released); included_files_sha256 refers to the committed, redacted files.", + "binding": "Tested tree 40070d3 = a66c999 (default-backend change) plus the test-h3 merge (eae0118). Reference, CUDA backend, fixture and reference test are byte-identical to the 4090D snapshot; the dispatch and CUDA tests and the benchmark differ (a66c999 split one registry test into two and rewrote dispatch assertions; every kernel bit assertion is unchanged). After the run, commit ace9c20 changed only the benchmark's GPU-without-Triton and HIP/ROCm branches (not executed on the H800 host; the run took the NVIDIA + Triton branch, whose logic, order, samples and JSON keys are unchanged) plus documentation: current sha256 benchmarks/models/benchmark_h3_sigma_schedule.py 04e2403ec4093245583d48f0e8a1b8645d6810743cbb9534a762335251ed4d16, docs/operators/h3-sigma-schedule.md 8f31e87df382178a4001167a29ab53c7062d3eef1b577bf683c7817b6de11956. Those new branches were verified by CPU-only execution, _summarize simulations and static review; they have not run on GPU or HIP hardware." + } } From 1cd229c0e7734ec498640e072f69d3f8b86bfecf Mon Sep 17 00:00:00 2001 From: fxyan <1350821504@qq.com> Date: Mon, 26 Oct 2026 02:23:54 +0800 Subject: [PATCH 07/11] fix(h3): align SM90 evidence listing with the repository *.log ignore rule MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The four standalone console logs (run_a/b/c.log, bench.log) were listed in sm90_validation.included_files_sha256 but are excluded by the repository-wide *.log ignore rule (.gitignore:60) and were silently skipped by git add — the 4090D evidence directory contains no .log files for the same reason. Drop them from included_files_sha256 (their full text is embedded verbatim in sm90-h800/pr566_sm90_evidence.md and their hashes stay recorded in original_server_sha256) and state the file set accurately in the README. No source, test or benchmark change. Signed-off-by: fxyan <1350821504@qq.com> --- reports/experiments/h3-sigma-schedule/README.md | 10 +++++++--- .../h3-sigma-schedule/source-provenance.json | 6 +----- 2 files changed, 8 insertions(+), 8 deletions(-) diff --git a/reports/experiments/h3-sigma-schedule/README.md b/reports/experiments/h3-sigma-schedule/README.md index 0a553150a..fbfce93a8 100644 --- a/reports/experiments/h3-sigma-schedule/README.md +++ b/reports/experiments/h3-sigma-schedule/README.md @@ -71,9 +71,13 @@ These are single-operator results, with no whole-model speedup claim. Collected on 2026-10-26 on a rented SM90 host (one **NVIDIA H800 PCIe**, per `nvidia-smi`), with the same three pytest invocations and the same benchmark command as above, at tested commit -`40070d37495d193e2d8764837321d3a68dad9483` (`40070d3`). Raw logs, JUnit XML, -the benchmark JSON and the full collection report are in -[sm90-h800/](sm90-h800/). GPU UUID, container hostname and `/root/…` paths +`40070d37495d193e2d8764837321d3a68dad9483` (`40070d3`). The JUnit XML for each +run, the benchmark JSON, the environment snapshot and the full collection +report (which embeds the raw console logs verbatim) are in +[sm90-h800/](sm90-h800/). The standalone console `.log` files are excluded by +the repository-wide `*.log` ignore rule — as in the 4090D directory; their +hashes are recorded in `source-provenance.json` (`original_server_sha256`). +GPU UUID, container hostname and `/root/…` paths were redacted: the JUnit `hostname` attribute was removed (as for the 4090D XML, which keeps the historical `/root/…` skip-message paths — the SM90 files additionally replace them with `REPO_ROOT` / `EVIDENCE_DIR` / diff --git a/reports/experiments/h3-sigma-schedule/source-provenance.json b/reports/experiments/h3-sigma-schedule/source-provenance.json index 5db435dee..31dc8b905 100644 --- a/reports/experiments/h3-sigma-schedule/source-provenance.json +++ b/reports/experiments/h3-sigma-schedule/source-provenance.json @@ -108,13 +108,9 @@ }, "included_files_sha256": { "sm90-h800/env.txt": "13d42ff62485307c55d4af5efe7fb6e7633a8f37a869d975d70249b1acaaaebf", - "sm90-h800/run_a.log": "e94695df1a4bab87c590e51b2231b179d91dc935674ad39a9388508f5854b593", "sm90-h800/run_a.xml": "4797238671384fbeb1ad9ddd54fa18ba48d760d18d69207f354ca291c3882604", - "sm90-h800/run_b.log": "62bf01571a0e70e8e095ed7e59fe5dc0e68b760e51f9abeca658cbd33b938aee", "sm90-h800/run_b.xml": "86cfec27c58060990317720777ce63b51dc208d09962445f1bca3a8121e5a658", - "sm90-h800/run_c.log": "7b46d39f310343f35776b76dc26405bb67bc9c047abd8bc1c987724de04cc12b", "sm90-h800/run_c.xml": "e2ea3433d80e7634976819856123659d6dafa8b3c747b964afc192e90ca802d2", - "sm90-h800/bench.log": "30acf2fb3ae50aed57efad735565b4225021ed13b9c81e1165995c8cf09b5224", "sm90-h800/benchmark-h800.json": "769b89b16679bdc5db365a2cc10f1dc7cc1966d72fbcf5ce0aae5a9d61b73d4c", "sm90-h800/pr566_sm90_evidence.md": "5ecf4b0672fd41bf74f9ce9f0484ed07fce252e3c6aadc490a5dd8b1080a9e81" }, @@ -131,7 +127,7 @@ "sm90-h800/pr566_sm90_evidence.md": "f4edb693161ce5a212c4da47f9e2737f3b35ec1f78ddd2908700f0e76e81e0c6" }, "benchmark_gate": "path C (Triton) bitwise equal to path A (CPU reference) for all 8 cases; gate wall incl. Triton compilation 1143.6 ms, excluded from timed samples; benchmark-h800.json is byte-identical to the server original except an appended final newline", - "redaction": "GPU UUID, container hostname and /root/... paths removed from the committed logs, JUnit XML and report (JUnit hostname attribute dropped, as for the 4090D XML; host paths additionally replaced by REPO_ROOT / EVIDENCE_DIR / MINICONDA_PYTHON tokens, one step beyond the 4090D precedent). A final newline was appended to the three XMLs and to benchmark-h800.json for the end-of-file-fixer hook; env.txt needed no change. original_server_sha256 records the pre-redaction originals (verified 10/10 against the host before it was released); included_files_sha256 refers to the committed, redacted files.", + "redaction": "GPU UUID, container hostname and /root/... paths removed from the committed logs, JUnit XML and report (JUnit hostname attribute dropped, as for the 4090D XML; host paths additionally replaced by REPO_ROOT / EVIDENCE_DIR / MINICONDA_PYTHON tokens, one step beyond the 4090D precedent). A final newline was appended to the three XMLs and to benchmark-h800.json for the end-of-file-fixer hook; env.txt needed no change. The four standalone console logs (run_a/b/c.log, bench.log) are excluded by the repository-wide *.log ignore rule — as in the 4090D evidence directory — so they are not committed; their full text is embedded verbatim in sm90-h800/pr566_sm90_evidence.md and their hashes are recorded in original_server_sha256. included_files_sha256 covers the committed, redacted files.", "binding": "Tested tree 40070d3 = a66c999 (default-backend change) plus the test-h3 merge (eae0118). Reference, CUDA backend, fixture and reference test are byte-identical to the 4090D snapshot; the dispatch and CUDA tests and the benchmark differ (a66c999 split one registry test into two and rewrote dispatch assertions; every kernel bit assertion is unchanged). After the run, commit ace9c20 changed only the benchmark's GPU-without-Triton and HIP/ROCm branches (not executed on the H800 host; the run took the NVIDIA + Triton branch, whose logic, order, samples and JSON keys are unchanged) plus documentation: current sha256 benchmarks/models/benchmark_h3_sigma_schedule.py 04e2403ec4093245583d48f0e8a1b8645d6810743cbb9534a762335251ed4d16, docs/operators/h3-sigma-schedule.md 8f31e87df382178a4001167a29ab53c7062d3eef1b577bf683c7817b6de11956. Those new branches were verified by CPU-only execution, _summarize simulations and static review; they have not run on GPU or HIP hardware." } } From 6d80c7b9738c79ce04f77cda305cff58186cdcbd Mon Sep 17 00:00:00 2001 From: fxyan <1350821504@qq.com> Date: Tue, 27 Oct 2026 18:49:19 +0800 Subject: [PATCH 08/11] fix(h3): copy sigmas before slicing timesteps in the NaN/inf fast path Signed-off-by: fxyan <1350821504@qq.com> --- .../model_specific/minimax_h3/sigma_schedule.py | 2 +- tests/backends/cuda/test_h3_sigma_schedule_cuda.py | 14 ++++++++++++++ 2 files changed, 15 insertions(+), 1 deletion(-) diff --git a/rl_engine/backends/cuda/model_specific/minimax_h3/sigma_schedule.py b/rl_engine/backends/cuda/model_specific/minimax_h3/sigma_schedule.py index 8f7997267..9bf31dc02 100644 --- a/rl_engine/backends/cuda/model_specific/minimax_h3/sigma_schedule.py +++ b/rl_engine/backends/cuda/model_specific/minimax_h3/sigma_schedule.py @@ -304,7 +304,7 @@ def forward( with torch.cuda.device(target): bits = torch.full((n,), signed, dtype=torch.int32, device=target) sigmas = bits.view(torch.float32) - timesteps = sigmas[: n - 1] + timesteps = sigmas[: n - 1].clone() return H3SigmaSchedule(sigmas=sigmas, timesteps=timesteps, num_model_evals=n - 1) # Device guard: allocations, Triton launches (which use the current diff --git a/tests/backends/cuda/test_h3_sigma_schedule_cuda.py b/tests/backends/cuda/test_h3_sigma_schedule_cuda.py index 72690ae88..4223cee6f 100644 --- a/tests/backends/cuda/test_h3_sigma_schedule_cuda.py +++ b/tests/backends/cuda/test_h3_sigma_schedule_cuda.py @@ -234,6 +234,20 @@ def test_negative_nan_shift_propagates_negative_quiet_nan(): assert set(_int32_bits(actual.sigmas.cpu())) == {-4194304} # 0xffc00000, signed +def test_nan_path_timesteps_do_not_alias_sigmas_storage(): + # The NaN/inf fast path slices sigmas for timesteps; the slice must own its + # storage like every other path here (the kernel path writes into a fresh + # torch.empty) and like the reference, whose ``1.0 - sigmas[:-1]`` is a new + # tensor. An in-place write to timesteps must not bleed into sigmas — the + # bitwise tests above cannot catch aliasing because the two NaN grids carry + # identical bit patterns. + schedule = _cuda_op()(25, shift=float("nan")) + + bits_before = _int32_bits(schedule.sigmas.cpu()) + schedule.timesteps.fill_(123.0) + assert _int32_bits(schedule.sigmas.cpu()) == bits_before + + _EDGE_SHIFTS = [ float("nan"), # all-NaN grid; IEEE != keeps every row (matches upstream) float("inf"), # all-NaN grid via fp32 scalar overflow semantics From 595c76c93063b591161589aec16e16ecc8854d93 Mon Sep 17 00:00:00 2001 From: fxyan <1350821504@qq.com> Date: Tue, 27 Oct 2026 18:52:40 +0800 Subject: [PATCH 09/11] docs(h3): tighten SM90 evidence report wording per review Signed-off-by: fxyan <1350821504@qq.com> --- docs/operators/h3-sigma-schedule.md | 7 ++++++- reports/experiments/h3-sigma-schedule/README.md | 5 +++-- .../sm90-h800/pr566_sm90_evidence.md | 16 +++++++++++----- 3 files changed, 20 insertions(+), 8 deletions(-) diff --git a/docs/operators/h3-sigma-schedule.md b/docs/operators/h3-sigma-schedule.md index 2c064254a..e1ac3e9f5 100644 --- a/docs/operators/h3-sigma-schedule.md +++ b/docs/operators/h3-sigma-schedule.md @@ -62,7 +62,12 @@ schedule = op(25, shift=12.0, device="cuda") ``` Both routes produce the same bytes as the default reference path (the CUDA -suite asserts this); opting in only moves the arithmetic onto the device. +suite asserts this; the byte identity was verified on x86_64 Linux hosts). +Whether it extends to ARM hosts for the extreme shifts (infinite, +fp32-overflowing, underflowing) is unverified — the device fast path fills +those grids with the x86 CPU oracle's NaN bits (see +[Known Limitations](#known-limitations)). Opting in only moves the +arithmetic onto the device. ## Backends diff --git a/reports/experiments/h3-sigma-schedule/README.md b/reports/experiments/h3-sigma-schedule/README.md index fbfce93a8..80b076557 100644 --- a/reports/experiments/h3-sigma-schedule/README.md +++ b/reports/experiments/h3-sigma-schedule/README.md @@ -114,8 +114,9 @@ two; no test that passed on the 4090D failed here; the single skip is the same test for the same reason. Benchmark (100 samples, 5 warmups, [raw JSON](sm90-h800/benchmark-h800.json)): -the pre-timing bitwise gate (path C vs path A) passed for all 8 cases; Triton -compilation 1143.6 ms, reported once and excluded from timed samples. +the pre-timing bitwise gate (path C vs path A) passed for all 8 cases; the +gate took 1143.6 ms wall in total (Triton kernel compilation included), +reported once and excluded from timed samples. | Workload | B median (ms) | C median (ms) | B/C | | --- | ---: | ---: | ---: | diff --git a/reports/experiments/h3-sigma-schedule/sm90-h800/pr566_sm90_evidence.md b/reports/experiments/h3-sigma-schedule/sm90-h800/pr566_sm90_evidence.md index f27ff2f39..e8e67b96e 100644 --- a/reports/experiments/h3-sigma-schedule/sm90-h800/pr566_sm90_evidence.md +++ b/reports/experiments/h3-sigma-schedule/sm90-h800/pr566_sm90_evidence.md @@ -141,8 +141,9 @@ python benchmarks/models/benchmark_h3_sigma_schedule.py --samples 100 --warmup 5 ``` Pre-timing correctness gate: path C (Triton, device build) bitwise equal to -path A (CPU reference) for all 8 cases — **passed**; Triton compilation -1143.6 ms wall, reported once and excluded from timed samples. Raw JSON: +path A (CPU reference) for all 8 cases — **passed**; the gate's total wall +time was 1143.6 ms (Triton kernel compilation included), reported once and +excluded from timed samples. Raw JSON: `EVIDENCE_DIR/benchmark-h800.json` (52 KB, 100 samples per case and path after 5 warmups). @@ -187,12 +188,17 @@ On NVIDIA H800 PCIe (SM90, compute capability 9.0, driver 580.82.07, torch 2.5.1+cu124, CUDA 12.4, Triton 3.1.0) at HEAD `40070d37495d193e2d8764837321d3a68dad9483`: - All 35 runnable tests in `tests/backends/cuda/test_h3_sigma_schedule_cuda.py` - passed — pinned-fixture bitwise match, live-CPU-reference bitwise match - (N=50/1000, shift 3/12), dense-grid dedup, grid invariants, subnormal + passed — pinned-fixture bitwise match, live-CPU-reference bitwise match on + the three combos that test parametrizes (N=50/shift=3, N=50/shift=12 and + N=1000/shift=12; its parametrization has no N=1000/shift=3 case), + dense-grid dedup, grid invariants, subnormal (1e-45) / underflow (1e-46) / ±NaN / inf / 1e39 / 3.5e38 / 1e30 / 1e-10 edge parity, float-N truncation, input rejections, repeated/interleaved stability, backend identity, default-reference-vs-Triton bit equality, and - semantic Triton opt-in. 1 test skipped (needs 2 GPUs; host has 1). + semantic Triton opt-in. 1 test skipped (needs 2 GPUs; host has 1). The + three-combo list is that test's full parametrization — all-8-case coverage + (N=25/50/1000/1e6 × shift 3/12) is provided separately by the benchmark's + pre-timing correctness gate, not by this suite. - All 12 dispatch tests passed. 20 reference tests passed, 1 skipped (`diffusers` not installed). - 0 failures and 0 errors across runs A, B and C. From 4c9c2722ae35896690bacc2dadd82f99fd44d536 Mon Sep 17 00:00:00 2001 From: fxyan <1350821504@qq.com> Date: Thu, 29 Oct 2026 00:32:35 +0800 Subject: [PATCH 10/11] fix(h3): keep the H3 constants in the reference package wildcard exports Signed-off-by: fxyan <1350821504@qq.com> --- rl_engine/reference/minimax_h3/__init__.py | 15 +++++++++++++-- 1 file changed, 13 insertions(+), 2 deletions(-) diff --git a/rl_engine/reference/minimax_h3/__init__.py b/rl_engine/reference/minimax_h3/__init__.py index 06e019ba3..7570ffbd7 100644 --- a/rl_engine/reference/minimax_h3/__init__.py +++ b/rl_engine/reference/minimax_h3/__init__.py @@ -13,7 +13,18 @@ from rl_engine.reference.minimax_h3.sigma_schedule import H3SigmaSchedule, NativeH3SigmaScheduleOp -__all__ = ["H3SigmaSchedule", "NativeH3SigmaScheduleOp"] +__all__ = [ + "H3SigmaSchedule", + "H3_ADALN_CHUNK_NAMES", + "H3_ADALN_CHUNKS", + "H3_FREQ_DIM", + "H3_HIDDEN_SIZE", + "H3_MAX_PERIOD", + "H3_MODALITY_NUM", + "H3_TIME_EMBED_DIM", + "H3_TIME_EMBED_HIDDEN_DIM", + "NativeH3SigmaScheduleOp", +] H3_FREQ_DIM = 256 H3_MAX_PERIOD = 10000 @@ -31,4 +42,4 @@ "scale_mlp", "gate_mlp", ) -H3_ADALN_CHUNKS = len(H3_ADALN_CHUNK_NAMES) \ No newline at end of file +H3_ADALN_CHUNKS = len(H3_ADALN_CHUNK_NAMES) From 1853a28a06a554e5807ae6c58d1bb3033c8df633 Mon Sep 17 00:00:00 2001 From: fxyan <1350821504@qq.com> Date: Thu, 29 Oct 2026 00:55:41 +0800 Subject: [PATCH 11/11] test(h3): skip the sigma CUDA suite when the Triton wheel is missing Signed-off-by: fxyan <1350821504@qq.com> --- .../backends/cuda/test_h3_sigma_schedule_cuda.py | 16 +++++++++++----- 1 file changed, 11 insertions(+), 5 deletions(-) diff --git a/tests/backends/cuda/test_h3_sigma_schedule_cuda.py b/tests/backends/cuda/test_h3_sigma_schedule_cuda.py index 4223cee6f..c6cecdb17 100644 --- a/tests/backends/cuda/test_h3_sigma_schedule_cuda.py +++ b/tests/backends/cuda/test_h3_sigma_schedule_cuda.py @@ -8,8 +8,11 @@ to the host only for comparison; no value is recomputed on the CPU. The backend is NVIDIA-only (its kernels emit inline PTX ``div.rn.f32``); -the suite runs on NVIDIA CUDA hardware and skips on HIP/ROCm builds, where -``H3SigmaScheduleCudaOp`` refuses construction. It is opt-in everywhere: the +the suite runs on NVIDIA CUDA hardware with a Triton wheel and skips +without a GPU, on HIP/ROCm builds, and when the backend module cannot +import Triton (same guard shape as the sibling Triton suite +``tests/backends/cuda/test_h3_flash_attention.py``). It is opt-in +everywhere: the legacy registry serves the PyTorch reference on CUDA too (default backend change per review, 2026-10-25), and the Triton class is reached through the semantic ``requested_backend`` or direct construction. The suite therefore @@ -25,10 +28,13 @@ import pytest import torch +from rl_engine.backends.cuda.model_specific.minimax_h3.sigma_schedule import _TRITON_AVAILABLE + pytestmark = pytest.mark.skipif( - not torch.cuda.is_available() or torch.version.hip is not None, - reason="NVIDIA CUDA device required: the Triton path is NVIDIA-only " - "(skipped without a GPU and on HIP/ROCm builds)", + not torch.cuda.is_available() or torch.version.hip is not None or not _TRITON_AVAILABLE, + reason="NVIDIA CUDA device and Triton wheel required: the Triton path is " + "NVIDIA-only (skipped without a GPU, on HIP/ROCm builds, and when Triton " + "is not importable)", ) _FIXTURE_PATH = (