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"), + )