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..45184db0e
--- /dev/null
+++ b/benchmarks/models/benchmark_h3_sigma_schedule.py
@@ -0,0 +1,352 @@
+# 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 (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
+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 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. 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
+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")
+ # 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]
+ 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:
+ # 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}"
+ 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"] = {"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)
+
+ 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 3a383da15..6e1c7357c 100644
--- a/docs/.nav.yml
+++ b/docs/.nav.yml
@@ -32,6 +32,7 @@ nav:
- operators/det-gemm.md
- operators/embedding.md
- MiniMax-H3:
+ - operators/h3-sigma-schedule.md
- operators/h3-timestep-sinusoid.md
- operators/h3-timestep-mlp.md
- operators/h3-adaln-projection.md
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
new file mode 100644
index 000000000..e1ac3e9f5
--- /dev/null
+++ b/docs/operators/h3-sigma-schedule.md
@@ -0,0 +1,231 @@
+# H3 Sigma Schedule
+
+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. 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") # 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)
+```
+
+`get_op` returns the PyTorch reference on every platform, CUDA included: the
+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
+
+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
+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)
+```
+
+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; 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
+
+| 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) | **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` | **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
+
+| 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` | — | — | 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 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)).
+
+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
+`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"` (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
+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](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
+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.
+- **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
+ 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
new file mode 100644
index 000000000..80b076557
--- /dev/null
+++ b/reports/experiments/h3-sigma-schedule/README.md
@@ -0,0 +1,178 @@
+# H3 sigma schedule: RTX 4090D evidence
+
+These are historical operator measurements from October 16–17, 2026,
+rechecked against candidate commit `374488b` 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 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 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
+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 (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.
+
+| 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.
+
+## 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`). 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` /
+`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; 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 |
+| --- | ---: | ---: | ---: |
+| 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:
+
+```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/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..e8e67b96e
--- /dev/null
+++ b/reports/experiments/h3-sigma-schedule/sm90-h800/pr566_sm90_evidence.md
@@ -0,0 +1,455 @@
+# 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**; 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).
+
+| 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 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). 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.
+- 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
new file mode 100644
index 000000000..31dc8b905
--- /dev/null
+++ b/reports/experiments/h3-sigma-schedule/source-provenance.json
@@ -0,0 +1,133 @@
+{
+ "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": {
+ "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. 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",
+ "_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 (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.xml": "4797238671384fbeb1ad9ddd54fa18ba48d760d18d69207f354ca291c3882604",
+ "sm90-h800/run_b.xml": "86cfec27c58060990317720777ce63b51dc208d09962445f1bca3a8121e5a658",
+ "sm90-h800/run_c.xml": "e2ea3433d80e7634976819856123659d6dafa8b3c747b964afc192e90ca802d2",
+ "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. 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."
+ }
+}
diff --git a/rl_engine/backends/cuda/model_specific/__init__.py b/rl_engine/backends/cuda/model_specific/__init__.py
index 3b09d3c11..9ed43ae29 100644
--- a/rl_engine/backends/cuda/model_specific/__init__.py
+++ b/rl_engine/backends/cuda/model_specific/__init__.py
@@ -1,2 +1,4 @@
# SPDX-License-Identifier: Apache-2.0
-"""CUDA implementations with model-specific arithmetic contracts."""
+# Copyright (c) 2026 RL-Kernel Contributors
+
+"""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 86cf4c9da..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,2 +1,8 @@
# SPDX-License-Identifier: Apache-2.0
# Copyright (c) 2026 RL-Kernel Contributors
+
+"""MiniMax-H3 (RFC #420) model-specific CUDA 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..9bf31dc02
--- /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].clone()
+ 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
index e54c60b77..7570ffbd7 100644
--- a/rl_engine/reference/minimax_h3/__init__.py
+++ b/rl_engine/reference/minimax_h3/__init__.py
@@ -11,6 +11,21 @@
from __future__ import annotations
+from rl_engine.reference.minimax_h3.sigma_schedule import 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
H3_TIME_EMBED_HIDDEN_DIM = 5376
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 1f1f3d0df..3b589a13d 100644
--- a/rl_engine/runtime/registry.py
+++ b/rl_engine/runtime/registry.py
@@ -208,6 +208,18 @@ 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. 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"
+ )
+ CUDA_H3_SIGMA_SCHEDULE = (
+ "rl_engine.backends.cuda.model_specific.minimax_h3." "sigma_schedule.H3SigmaScheduleCudaOp"
+ )
# MiniMax-H3 (RFC #420) conditioning path
PYTORCH_H3_TIMESTEP_SINUSOID = (
"rl_engine.reference.minimax_h3.timestep_sinusoid.NativeH3TimestepSinusoidOp"
@@ -488,6 +500,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",
+ ),
)
@@ -751,6 +808,12 @@ def __init__(self):
OpBackend.TRITON_ROPE,
OpBackend.PYTORCH_NATIVE_ROPE,
],
+ # 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],
"timestep_sinusoid_h3": [
OpBackend.CUDA_H3_TIMESTEP_SINUSOID,
OpBackend.PYTORCH_H3_TIMESTEP_SINUSOID,
@@ -791,6 +854,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,
@@ -839,6 +903,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,
@@ -883,6 +948,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..c6cecdb17
--- /dev/null
+++ b/tests/backends/cuda/test_h3_sigma_schedule_cuda.py
@@ -0,0 +1,412 @@
+# 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 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
+constructs the backend directly for kernel checks and asserts that the default
+reference path matches it bit for bit.
+"""
+
+from __future__ import annotations
+
+import json
+import pathlib
+
+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 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 = (
+ 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
+
+
+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
+ 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
+
+ # 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"
+ )
+ 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_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
+
+ op = kernel_registry.get_op("sigma_schedule_h3", device="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(
+ 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..fd05062d1
--- /dev/null
+++ b/tests/runtime/test_h3_sigma_schedule_dispatch.py
@@ -0,0 +1,190 @@
+# 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`` 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
+
+import importlib
+
+import pytest
+import torch
+
+from rl_engine.runtime.registry import KernelRegistry, OpBackend, 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_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()
+
+ 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):
+ 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_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()
+
+ op = registry.get_op("sigma_schedule_h3", device="cuda")
+
+ 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"
+ )
+ 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"),
+ )