From ac17ab87e8d93dabe93dfdc134e9853d47de7af1 Mon Sep 17 00:00:00 2001 From: dandan20031031 <2425926169@qq.com> Date: Wed, 28 Oct 2026 17:12:03 +0800 Subject: [PATCH 1/5] feat(rope): add Qwen-Image multi-axis rotation with WS1 validation Provide stride-aware native CUDA and FlashAttention-derived Triton paths with fixed FP32 arithmetic, dispatcher/fake/autograd registration, standard operator validation, package build checks, and mature-baseline benchmark evidence. Refs #386. Signed-off-by: dandan20031031 <2425926169@qq.com> --- MANIFEST.in | 1 + .../rope/benchmark_multi_axis_rope.py | 290 ++ build_tools/extensions.py | 29 + csrc/bindings/rope_module.cpp | 8 + csrc/cuda/rope/multi_axis_rope.cu | 82 + docs/.nav.yml | 1 + docs/operators/README.md | 1 + docs/operators/multi-axis-rope.md | 199 ++ pyproject.toml | 2 +- .../README.md | 152 + .../baseline-layout.txt | 138 + .../benchmark-bfloat16.json | 2781 +++++++++++++++++ .../installed-wheel.json | 37 + .../installed-wheel.txt | 38 + .../layout.txt | 138 + .../manifest.json | 65 + .../memcheck.txt | 8 + .../operator-cuda.json | 70 + .../operator-triton.json | 70 + .../package-build.txt | 40 + .../package.json | 31 + .../precommit.txt | 7 + .../quality.txt | 58 + .../regression.txt | 10 + .../tests.txt | 5 + .../upstream.json | 335 ++ .../ws1-cuda.json | 1962 ++++++++++++ .../ws1-cuda.txt | 3 + .../ws1-triton.json | 1962 ++++++++++++ .../ws1-triton.txt | 3 + .../backends/cuda/rope/multi_axis_rope.py | 49 + .../triton/rope/FLASH_ATTENTION_LICENSE | 29 + .../shared/triton/rope/_flash_attn_rotary.py | 141 + .../shared/triton/rope/multi_axis_rope.py | 68 + rl_engine/ops/rope/_registration.py | 41 + rl_engine/ops/rope/multi_axis_rope.py | 132 + rl_engine/reference/rope/multi_axis_rope.py | 25 + rl_engine/runtime/registry.py | 14 + .../validation/operators/gradient_adapters.py | 29 + .../validation/operators/operator_inputs.py | 19 + .../validation/operators/operator_specs.py | 14 + .../validation/operators/qwen_image_rope.py | 80 + tests/build/test_multi_axis_rope_build.py | 46 + tests/ops/rope/test_multi_axis_rope.py | 443 +++ .../operators/test_operator_inputs.py | 18 + .../check_multi_axis_rope_install.py | 67 + .../check_qwenimage_rope_upstream.py | 99 + .../operators/validate_multi_axis_rope.py | 85 + 48 files changed, 9924 insertions(+), 1 deletion(-) create mode 100644 benchmarks/operators/rope/benchmark_multi_axis_rope.py create mode 100644 csrc/bindings/rope_module.cpp create mode 100644 csrc/cuda/rope/multi_axis_rope.cu create mode 100644 docs/operators/multi-axis-rope.md create mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/README.md create mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/baseline-layout.txt create mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/benchmark-bfloat16.json create mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/installed-wheel.json create mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/installed-wheel.txt create mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/layout.txt create mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/manifest.json create mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/memcheck.txt create mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/operator-cuda.json create mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/operator-triton.json create mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/package-build.txt create mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/package.json create mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/precommit.txt create mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/quality.txt create mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/regression.txt create mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/tests.txt create mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/upstream.json create mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-cuda.json create mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-cuda.txt create mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-triton.json create mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-triton.txt create mode 100644 rl_engine/backends/cuda/rope/multi_axis_rope.py create mode 100644 rl_engine/backends/shared/triton/rope/FLASH_ATTENTION_LICENSE create mode 100644 rl_engine/backends/shared/triton/rope/_flash_attn_rotary.py create mode 100644 rl_engine/backends/shared/triton/rope/multi_axis_rope.py create mode 100644 rl_engine/ops/rope/_registration.py create mode 100644 rl_engine/ops/rope/multi_axis_rope.py create mode 100644 rl_engine/reference/rope/multi_axis_rope.py create mode 100644 rl_engine/validation/operators/qwen_image_rope.py create mode 100644 tests/build/test_multi_axis_rope_build.py create mode 100644 tests/ops/rope/test_multi_axis_rope.py create mode 100644 tools/validation/operators/check_multi_axis_rope_install.py create mode 100644 tools/validation/operators/check_qwenimage_rope_upstream.py create mode 100644 tools/validation/operators/validate_multi_axis_rope.py diff --git a/MANIFEST.in b/MANIFEST.in index 5ce1ee23f..67165e195 100644 --- a/MANIFEST.in +++ b/MANIFEST.in @@ -2,3 +2,4 @@ include pyproject.toml setup.py README.md LICENSE recursive-include build_tools *.py recursive-include csrc *.cpp *.cu *.cuh *.h *.hip *.asc recursive-include rl_engine *.json *.jsonl *.sh *.patch *.pyi +recursive-include rl_engine FLASH_ATTENTION_LICENSE diff --git a/benchmarks/operators/rope/benchmark_multi_axis_rope.py b/benchmarks/operators/rope/benchmark_multi_axis_rope.py new file mode 100644 index 000000000..316a75ba7 --- /dev/null +++ b/benchmarks/operators/rope/benchmark_multi_axis_rope.py @@ -0,0 +1,290 @@ +# SPDX-License-Identifier: Apache-2.0 +"""Standalone RoPE: contiguous/packed QKV, eager/graph, optional pinned FlashAttention.""" + +from __future__ import annotations + +import argparse +import ast +import gc +import hashlib +import importlib.util +import json +import platform +import statistics +import subprocess +import sys +import time +from pathlib import Path +from typing import Optional, Union + +import torch + +ROOT = Path(__file__).resolve().parents[3] +if str(ROOT) not in sys.path: + sys.path.insert(0, str(ROOT)) + +from rl_engine.ops.rope.multi_axis_rope import ( # noqa: E402 + build_multi_axis_cos_sin, + get_multi_axis_rope_op, + qwen_image_positions, +) + +FLASH_REVISION = "98c1cee0ed59307833841cdbf567fe336ceec13a" +FLASH_SHA256 = "d2037e623854c702fb075784fe4292d4a89d6b9ac1ae33cab3d6788f8f94f240" +FLASH_LAYER_SHA256 = "a79cb42dfc33e126227bdba27ae7a024d7d5cc5c8fd7f7a5f26c68d501024e0a" + + +def upstream_complex(x, cis): + """Diffusers use_real=False arithmetic; table assembly is outside timing.""" + z = torch.view_as_complex(x.float().reshape(*x.shape[:-1], 64, 2)) + return torch.view_as_real(z * cis.unsqueeze(-2)).flatten(-2).to(x.dtype) + + +def load_flash_baselines(source_path, layer_path): + """Load checksum-pinned upstream kernel and autograd class without installing flash-attn.""" + source = source_path.read_text(encoding="utf-8") + layer = layer_path.read_text(encoding="utf-8") + for text, expected in ((source, FLASH_SHA256), (layer, FLASH_LAYER_SHA256)): + if hashlib.sha256(text.encode()).hexdigest() != expected: + raise ValueError("FlashAttention source checksum mismatch (normalized UTF-8 text)") + node = next(n for n in ast.parse(layer).body if getattr(n, "name", None) == "ApplyRotaryEmb") + directory = ROOT / "artifacts/flash_rope_baseline" + directory.mkdir(parents=True, exist_ok=True) + result = {} + assert source.count("BLOCK_H=2,") == 1 + for variant in ("default", "no_fma"): + code = source + if variant == "no_fma": + code = code.replace("BLOCK_H=2,", "BLOCK_H=2,\n enable_fp_fusion=False,") + path = directory / f"flash_rotary_{variant}.py" + path.write_text(code, encoding="utf-8") + name = f"rlk_flash_baseline_{variant}" + spec = importlib.util.spec_from_file_location(name, path) + module = importlib.util.module_from_spec(spec) + sys.modules[name] = module + spec.loader.exec_module(module) + namespace = { + "__name__": name + "_autograd", + "torch": torch, + "Tensor": torch.Tensor, + "Optional": Optional, + "Union": Union, + "apply_rotary": module.apply_rotary, + } + exec(compile(ast.Module(body=[node], type_ignores=[]), str(layer_path), "exec"), namespace) + cls = namespace["ApplyRotaryEmb"] + result[f"flash_{variant}"] = lambda x, c, s, cls=cls: cls.apply(x, c, s, True) + return result + + +def measure(fn, args): + for _ in range(args.warmup): + result = fn() + torch.cuda.synchronize() + del result + wall = [] + for _ in range(args.trials): + start = time.perf_counter() + for _ in range(args.iterations): + result = fn() + torch.cuda.synchronize() + wall.append((time.perf_counter() - start) * 1000 / args.iterations) + del result + + # Collect graph owners before capture, so GC cannot destroy another graph + # during stream capture. Preserve the caller's GC state. + gc.collect() + enabled = gc.isenabled() + gc.disable() + graph = torch.cuda.CUDAGraph() + try: + with torch.cuda.graph(graph, stream=torch.cuda.Stream()): + for _ in range(args.graph_calls): + result = fn() + for _ in range(args.warmup): + graph.replay() + events = [ + (torch.cuda.Event(enable_timing=True), torch.cuda.Event(enable_timing=True)) + for _ in range(args.iterations) + ] + for start, end in events: + start.record() + graph.replay() + end.record() + torch.cuda.synchronize() + times = [start.elapsed_time(end) / args.graph_calls for start, end in events] + del graph, result + finally: + if enabled: + gc.enable() + + before = torch.cuda.memory_allocated() + torch.cuda.reset_peak_memory_stats() + result = fn() + torch.cuda.synchronize() + extra = torch.cuda.max_memory_allocated() - before + del result + return { + "graph_median_ms": statistics.median(times), + "graph_min_ms": min(times), + "graph_max_ms": max(times), + "eager_wall_median_ms": statistics.median(wall), + "eager_wall_trials_ms": wall, + "extra_peak_bytes": extra, + } + + +def byte_differences(a, b): + return int((a.contiguous().view(torch.uint8) != b.contiguous().view(torch.uint8)).sum()) + + +def main(): + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--output", type=Path, required=True) + parser.add_argument("--iterations", type=int, default=30) + parser.add_argument("--warmup", type=int, default=5) + parser.add_argument("--trials", type=int, default=3) + parser.add_argument("--graph-calls", type=int, default=20) + parser.add_argument("--dtype", choices=["float32", "float16", "bfloat16"], default="bfloat16") + parser.add_argument( + "--layouts", + nargs="+", + choices=["contiguous", "packed_qkv"], + default=["contiguous", "packed_qkv"], + ) + parser.add_argument("--flash-source", type=Path) + parser.add_argument("--flash-layer-source", type=Path) + args = parser.parse_args() + if min(args.iterations, args.warmup, args.trials, args.graph_calls) < 1: + parser.error("timing counts must be positive") + if bool(args.flash_source) != bool(args.flash_layer_source): + parser.error("both FlashAttention source files are required") + dtype = getattr(torch, args.dtype) + torch.manual_seed(386) + ours = {name: get_multi_axis_rope_op(name, device="cuda") for name in ("cuda", "triton")} + flash = ( + load_flash_baselines(args.flash_source, args.flash_layer_source) + if args.flash_source + else {} + ) + reference = get_multi_axis_rope_op("reference") + rows = [] + correctness = [] + for height, width in ((2, 3), (64, 64), (83, 83), (104, 58)): + positions, _ = qwen_image_positions(height, width, 0) + cos, sin = build_multi_axis_cos_sin(positions, device="cuda") + cis = torch.complex(cos, sin) + + def upstream_op(x, _cos, _sin, cis=cis): + return upstream_complex(x, cis) + + ops = {"upstream_complex": upstream_op, **ours, **flash} + for layout in args.layouts: + if layout == "packed_qkv": + packed = torch.randn(1, height * width, 3, 24, 128, device="cuda", dtype=dtype) + x = packed[:, :, 0].detach().requires_grad_() + else: + x = torch.randn( + 1, height * width, 24, 128, device="cuda", dtype=dtype, requires_grad=True + ) + grad = torch.randn_like(x) + gold = reference(x, cos, sin) + gold_dx = torch.autograd.grad(gold, x, grad)[0] + for name, op in ops.items(): + result = op(x, cos, sin) + dx = torch.autograd.grad(result, x, grad)[0] + diff = [byte_differences(result, gold), byte_differences(dx, gold_dx)] + if name in ("cuda", "triton", "flash_no_fma"): + assert diff == [0, 0], (name, layout, x.shape, diff) + else: + tol = {torch.float32: 2e-6, torch.float16: 0.004, torch.bfloat16: 0.04}[dtype] + torch.testing.assert_close(result, gold, atol=tol, rtol=tol) + torch.testing.assert_close(dx, gold_dx, atol=tol, rtol=tol) + correctness.append( + { + "backend": name, + "layout": layout, + "shape": list(x.shape), + "different_bytes_y_dx": diff, + } + ) + del result, dx, gold, gold_dx + for name, op in ops.items(): + for phase in ("forward", "forward_backward"): + + def call(op=op, phase=phase, x=x, cos=cos, sin=sin, grad=grad): + if phase == "forward": + with torch.no_grad(): + return op(x, cos, sin) + return torch.autograd.grad(op(x, cos, sin), x, grad)[0] + + stats = measure(call, args) + rows.append( + { + "backend": name, + "phase": phase, + "layout": layout, + "shape": list(x.shape), + "input_stride": list(x.stride()), + **stats, + } + ) + print( + f"{height}x{width} {layout:11s} {name:16s} {phase:16s} " + f"graph={stats['graph_median_ms']:.5f}ms " + f"eager={stats['eager_wall_median_ms']:.5f}ms", + flush=True, + ) + sources = [ + "csrc/cuda/rope/multi_axis_rope.cu", + "rl_engine/backends/shared/triton/rope/_flash_attn_rotary.py", + "rl_engine/backends/shared/triton/rope/multi_axis_rope.py", + "rl_engine/backends/cuda/rope/multi_axis_rope.py", + "rl_engine/ops/rope/_registration.py", + "rl_engine/ops/rope/multi_axis_rope.py", + "rl_engine/reference/rope/multi_axis_rope.py", + "benchmarks/operators/rope/benchmark_multi_axis_rope.py", + ] + import triton + + evidence = { + "gpu": torch.cuda.get_device_name(), + "capability": torch.cuda.get_device_capability(), + "os": platform.platform(), + "python": sys.version, + "torch": torch.__version__, + "cuda_runtime": torch.version.cuda, + "triton": triton.__version__, + "dtype": args.dtype, + "base_commit": subprocess.check_output( + ["git", "rev-parse", "HEAD"], cwd=ROOT, text=True + ).strip(), + "source_sha256": {p: hashlib.sha256((ROOT / p).read_bytes()).hexdigest() for p in sources}, + "timing": "CUDA graph ms/call and synchronized eager wall ms/call; warm allocation; " + "excludes tables, compilation, and initial allocation of inputs", + "memory": "extra peak allocated bytes excluding inputs/tables/grad; includes returned " + "output or gradient; packed QKV backing storage allocated before timing", + "iterations": args.iterations, + "warmup": args.warmup, + "trials": args.trials, + "graph_calls": args.graph_calls, + "flash_baseline": { + "included": bool(flash), + "revision": FLASH_REVISION, + "kernel_sha256_normalized": FLASH_SHA256, + "layer_sha256_normalized": FLASH_LAYER_SHA256, + "mode": "upstream rotary_kernel, apply_rotary and ApplyRotaryEmb, interleaved=True", + "no_fma_change": "only add enable_fp_fusion=False at the original kernel launch", + }, + "trace": {name: op.trace() for name, op in ours.items()}, + "qualification": "Exact bytes vs unfused FP32 gold for CUDA, Triton and flash_no_fma; " + "declared tolerance for upstream complex and default FlashAttention", + "correctness": correctness, + "results": rows, + } + args.output.parent.mkdir(parents=True, exist_ok=True) + args.output.write_text(json.dumps(evidence, indent=2) + "\n", encoding="utf-8") + + +if __name__ == "__main__": + main() diff --git a/build_tools/extensions.py b/build_tools/extensions.py index db079e828..2f04f5cdd 100644 --- a/build_tools/extensions.py +++ b/build_tools/extensions.py @@ -105,6 +105,7 @@ def _native_extension_required() -> bool: """Whether the caller explicitly requested a native extension build.""" return ( envs.env_flag(envs.RL_KERNEL_REQUIRE_EXT) + or envs.env_flag("RL_KERNEL_ROPE_ONLY") or bool(os.environ.get("PYTORCH_ROCM_ARCH", "").strip()) or bool(os.environ.get("TORCH_CUDA_ARCH_LIST", "").strip()) or envs.env_flag("FORCE_CUDA") @@ -182,6 +183,33 @@ def get_extensions(): torch_rpath.append(f"-Wl,-rpath,{torch_lib_dir}") is_rocm = getattr(torch.version, "hip", None) is not None + # A normal wheel build for the independently usable RoPE library. This + # explicit profile avoids compiling unrelated Hopper/IPC operators and is + # also useful on Windows. Default builds still include RoPE in the full _C. + if envs.env_flag("RL_KERNEL_ROPE_ONLY"): + if is_rocm or getattr(torch.version, "cuda", None) is None: + raise RuntimeError("RL_KERNEL_ROPE_ONLY requires CUDA-enabled PyTorch") + from torch.utils.cpp_extension import CUDA_HOME + + library_dirs = [] + if os.name == "nt" and CUDA_HOME and (Path(CUDA_HOME) / "lib/cudart.lib").is_file(): + library_dirs.append(str(Path(CUDA_HOME) / "lib")) + return [ + GpuExtension( + name="rl_engine._rope_cuda", + sources=[ + "csrc/bindings/rope_module.cpp", + "csrc/cuda/rope/multi_axis_rope.cu", + ], + library_dirs=library_dirs, + extra_compile_args={ + "cxx": ["/O2", "/std:c++17"] if os.name == "nt" else ["-O3", "-std=c++17"], + "nvcc": ["-O3", "-std=c++17", "--fmad=false"], + }, + extra_link_args=[] if os.name == "nt" else list(torch_rpath), + ) + ] + if _musa_build_available(torch): extensions.append( GpuExtension( @@ -233,6 +261,7 @@ def get_extensions(): ] ) else: + cuda_sources.append("csrc/cuda/rope/multi_axis_rope.cu") # CUDA IPC and the fixed-tree collective implementation are not # part of the ROCm extension. cuda_sources.append("csrc/cuda/collectives/deterministic_collective.cu") diff --git a/csrc/bindings/rope_module.cpp b/csrc/bindings/rope_module.cpp new file mode 100644 index 000000000..c12e5b90d --- /dev/null +++ b/csrc/bindings/rope_module.cpp @@ -0,0 +1,8 @@ +// SPDX-License-Identifier: Apache-2.0 +// Python module entry point for the optional RoPE-only wheel. Operators are +// registered by multi_axis_rope.cu through the PyTorch dispatcher. +#include + +PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) { + m.doc() = "Qwen-Image RoPE dispatcher library"; +} diff --git a/csrc/cuda/rope/multi_axis_rope.cu b/csrc/cuda/rope/multi_axis_rope.cu new file mode 100644 index 000000000..3de49caba --- /dev/null +++ b/csrc/cuda/rope/multi_axis_rope.cu @@ -0,0 +1,82 @@ +// SPDX-License-Identifier: Apache-2.0 +// Qwen-Image adjacent-pair rotation. Reads strided BSHD inputs/tables directly. +#include +#include +#include +#include +#include +#include +#include + +namespace { +struct Strides { + int64_t xb, xs, xh, xd; + int64_t cb, cs, cd, sb, ss, sd; +}; + +template +__global__ void rotate_pairs(const scalar_t* x, const float* cos, const float* sin, + scalar_t* y, int64_t n, int64_t seq, int64_t heads, + Strides strides, bool transpose) { + const int64_t step = static_cast(blockDim.x) * gridDim.x; + for (int64_t i = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; + i < n; i += step) { + const int64_t pair = i % 64; + const int64_t row = i / 64; + const int64_t head = row % heads; + const int64_t token = row / heads; + const int64_t batch = token / seq; + const int64_t position = token % seq; + const int64_t xi = batch * strides.xb + position * strides.xs + + head * strides.xh + 2 * pair * strides.xd; + const float c = cos[batch * strides.cb + position * strides.cs + pair * strides.cd]; + const float sn = sin[batch * strides.sb + position * strides.ss + pair * strides.sd]; + const float s = transpose ? -sn : sn; + const float a = static_cast(x[xi]); + const float b = static_cast(x[xi + strides.xd]); + y[2 * i] = static_cast(__fsub_rn(__fmul_rn(a, c), __fmul_rn(b, s))); + y[2 * i + 1] = static_cast(__fadd_rn(__fmul_rn(b, c), __fmul_rn(a, s))); + } +} + +at::Tensor multi_axis_rope_apply(const at::Tensor& x, const at::Tensor& cos, + const at::Tensor& sin, bool transpose) { + TORCH_CHECK(x.is_cuda() && cos.is_cuda() && sin.is_cuda(), "multi_axis_rope: CUDA tensors required"); + TORCH_CHECK(x.device() == cos.device() && x.device() == sin.device(), "multi_axis_rope: device mismatch"); + TORCH_CHECK(x.dim() == 4 && x.size(3) == 128 && x.size(2) > 0, + "multi_axis_rope: x must be [B, S, H, 128], H > 0"); + TORCH_CHECK(x.scalar_type() == at::kFloat || x.scalar_type() == at::kHalf || + x.scalar_type() == at::kBFloat16, "multi_axis_rope: unsupported input dtype"); + TORCH_CHECK(cos.scalar_type() == at::kFloat && sin.scalar_type() == at::kFloat, + "multi_axis_rope: tables must be fp32"); + TORCH_CHECK(cos.sizes() == sin.sizes(), "multi_axis_rope: table shape mismatch"); + TORCH_CHECK((cos.dim() == 2 && cos.size(0) == x.size(1) && cos.size(1) == 64) || + (cos.dim() == 3 && cos.size(0) == x.size(0) && cos.size(1) == x.size(1) && cos.size(2) == 64), + "multi_axis_rope: tables must be [S, 64] or [B, S, 64]"); + TORCH_CHECK(!cos.requires_grad() && !sin.requires_grad(), "multi_axis_rope: table gradients unsupported"); + const c10::cuda::CUDAGuard guard(x.device()); + auto y = at::empty(x.sizes(), x.options()); + const int64_t n = x.numel() / 2; + if (n == 0) return y; + const Strides strides{x.stride(0), x.stride(1), x.stride(2), x.stride(3), + cos.dim() == 3 ? cos.stride(0) : 0, cos.stride(-2), cos.stride(-1), + sin.dim() == 3 ? sin.stride(0) : 0, sin.stride(-2), sin.stride(-1)}; + const int blocks = static_cast(std::min((n + 255) / 256, 65535)); + auto stream = c10::cuda::getCurrentCUDAStream(x.get_device()); + AT_DISPATCH_FLOATING_TYPES_AND2(at::ScalarType::Half, at::ScalarType::BFloat16, + x.scalar_type(), "multi_axis_rope", [&] { + rotate_pairs<<>>(x.data_ptr(), cos.data_ptr(), + sin.data_ptr(), y.data_ptr(), n, x.size(1), x.size(2), strides, transpose); + }); + C10_CUDA_KERNEL_LAUNCH_CHECK(); + return y; +} +} // namespace + +// The focused validator and normal package build load this same registration. +TORCH_LIBRARY_FRAGMENT(rl_kernel_native, m) { + m.def("multi_axis_rope(Tensor x, Tensor cos, Tensor sin, bool transpose) -> Tensor"); +} +TORCH_LIBRARY_IMPL(rl_kernel_native, CUDA, m) { + m.impl("multi_axis_rope", &multi_axis_rope_apply); +} diff --git a/docs/.nav.yml b/docs/.nav.yml index 29b5ead37..d03ade331 100644 --- a/docs/.nav.yml +++ b/docs/.nav.yml @@ -29,6 +29,7 @@ nav: - operators/sampling.md - operators/det-gemm.md - operators/embedding.md + - operators/multi-axis-rope.md - Developer Guide: - contributing/README.md - Contributor Guide: contributing/contributor-guide.md diff --git a/docs/operators/README.md b/docs/operators/README.md index 00f4cbb45..2d3b27d5d 100644 --- a/docs/operators/README.md +++ b/docs/operators/README.md @@ -26,6 +26,7 @@ Every operator page should include: - [Fused Linear LogP TP Test Runbook](linear-logp-tp-test.md) - [GRPO Loss](grpo-loss.md) - [RoPE](rope.md) +- [Qwen-Image Multi-Axis RoPE](multi-axis-rope.md) - [LM Head](lm_head.md) - [Policy Ratio + KL Penalty](ratio-kl.md) - [Pack and Pad](pack-and-pad.md) diff --git a/docs/operators/multi-axis-rope.md b/docs/operators/multi-axis-rope.md new file mode 100644 index 000000000..fa908a03e --- /dev/null +++ b/docs/operators/multi-axis-rope.md @@ -0,0 +1,199 @@ +# Qwen-Image Multi-Axis RoPE + +This operator implements the `multi_axis_rope` CUDA row of +[RFC #386](https://github.com/RL-Align/RL-Kernel/issues/386). It rotates **adjacent +channel pairs** with three independently indexed axes of widths `[16, 56, 56]`. +The existing GPT-NeoX rotate-half operator has different semantics. + +## Reference and scope + +Semantics are pinned to Diffusers +[`40f63690da440e8c382ab53e67508835deb1ed2f`](https://github.com/huggingface/diffusers/blob/40f63690da440e8c382ab53e67508835deb1ed2f/src/diffusers/models/transformers/transformer_qwenimage.py): +`QwenEmbedRope(scale_rope=True)` and `apply_rotary_emb_qwen(use_real=False)`. +The position helper covers single-frame text-to-image at packed latent resolution. +For a 1024×1024 image the packed grid is 64×64. Multi-image editing, video, +QwenEmbedLayer3DRope, distributed execution and full-model integration are outside +this operator's qualification. + +Image coordinates are `(0, h, w)` in row-major order. The centered coordinate +range is `[-ceil(size/2), floor(size/2))`, so an odd grid has one extra negative +coordinate. Text starts at `max(height//2, width//2)` on all three axes. The helper +returns image and text positions separately: the caller determines stream order. + +## API + +```python +import torch +from rl_engine.ops.rope.multi_axis_rope import ( + qwen_image_positions, build_multi_axis_cos_sin, get_multi_axis_rope_op, +) + +image_pos, text_pos = qwen_image_positions(64, 64, 128) +cos, sin = build_multi_axis_cos_sin(image_pos, device="cuda") +x = torch.randn(1, 4096, 24, 128, device="cuda", dtype=torch.bfloat16, + requires_grad=True) +op = get_multi_axis_rope_op("cuda", device=x.device) +y = op(x, cos, sin) +y.float().square().mean().backward() +print(op.trace()) +``` + +| Argument | Contract | +| --- | --- | +| `x` | `[B, S, H, 128]`, H > 0; FP16, BF16 or FP32 | +| `cos`, `sin` | Matching FP32 `[S, 64]` shared tables or `[B, S, 64]` per-sample tables | +| Devices | Input and tables on the same device | +| Layout | BSHD, adjacent pairs; input and independent table strides read directly | +| Output | Input shape/dtype/device; contiguous | +| Empty input | B=0 and/or S=0 supported, including backward | +| Gradients | Input gradients and higher-order input VJPs; table gradients rejected | +| Position builder | int32/int64 `[S, 3]` or `[B, S, 3]`, within `[-4096, 4095]` | + +Table construction follows the upstream CPU FP32 `pow` and `polar` operations, +then copies the results to the target device. Build and retain tables before CUDA +Graph capture. Reusing these tables also prevents per-batch transcendental drift. +The CPU factories select their device explicitly, including inside a CUDA or +meta default-device context. Packed QKV views require no materializing copy. + +## Numerical contract + +Contract ID: `qwenimage-interleaved-rope-fp32-v1`. + +For each adjacent pair `(a, b)`, compute `(a*c - b*s, b*c + a*s)` in FP32 with +separate rounded multiplications and additions, followed by one output cast. +There is no cross-element reduction, split-K, atomic accumulation or TF32. +CUDA uses round-to-nearest arithmetic intrinsics; Triton disables FP contraction. +Backward uses the transpose rotation (negate sine), including recursive VJPs. + +CUDA and Triton must agree in raw bytes within a pinned environment. Compare +batch placement, token chunking, padding and graph replay the same way. +Independent FP64 complex arithmetic provides an accuracy oracle. Upstream's GPU +complex multiplication may contract arithmetic, so its comparison uses declared +tolerances, not an unverified bitwise claim. Representative test tolerances are +`8e-7` for FP32, `0.002` for FP16 and `0.02` for BF16 (absolute and relative). +Non-finite propagation follows the explicit arithmetic; NaN payload equality is +not a cross-backend guarantee. + +## Backends and dispatch + +- `cuda`: production function in `csrc/cuda/rope/multi_axis_rope.cu`, registered + in `rl_engine._C` by ordinary CUDA builds. Does not require Hopper-specific PTX. +- `triton`: adapts FlashAttention's existing interleaved rotary kernel, with + explicit shared/per-sample table strides and `enable_fp_fusion=False`. + The upstream tile geometry is retained; this profile is limited to NVIDIA CUDA. +- `reference`: PyTorch FP32 arithmetic, including CPU. +- `auto`: CUDA → Triton → reference on NVIDIA, reference on other platforms. + +Explicit `cuda`/`triton` requests fail when unavailable; they never silently switch +implementations. Strict callers select an explicit backend and retain `op.trace()`. +`auto` is for diagnostic use and is not a strict-route certification. The legacy +registry also exposes `kernel_registry.get_op("multi_axis_rope", device=...)`. +This PR does not add an Ascend backend or certify ROCm. Existing Ascend work in +PR #411 can share semantics once its reference contract is aligned. + +Both GPU paths expose a PyTorch custom operator with fake-tensor and autograd +registration. `torch.compile(op, fullgraph=True)` supports forward and backward. +The native implementation is registered through `TORCH_LIBRARY`; Python does +not attach symbols to an unrelated extension at validation time. + +## Related work and reuse + +The Triton kernel comes from FlashAttention +[`98c1cee`](https://github.com/Dao-AILab/flash-attention/blob/98c1cee0ed59307833841cdbf567fe336ceec13a/flash_attn/ops/triton/rotary.py), +with its BSD-3-Clause copyright and full license retained beside the adapted +source and included in wheels/source distributions. Its default fused arithmetic +can differ in low bits; disabling contraction gives the declared profile. +The benchmark loads the original kernel and autograd wrapper from checksum-pinned +source files to compare the default and contraction-disabled variants. + +[Ascend PR #411](https://github.com/RL-Align/RL-Kernel/pull/411) covers the same +roadmap row on another platform. Its current half-split reference and uncentered +positions differ from the pinned Diffusers contract here. Combining the two +references requires semantic alignment. Existing Qwen-Image implementations in +[SGLang](https://github.com/sgl-project/sglang/tree/a1966e1b5d8b750b494d7e53cca1d2464f960038/python/sglang/multimodal_gen) +and [vLLM-Omni](https://github.com/vllm-project/vllm-omni/tree/c9c288e3f5828f2a2a9b07acb16b3b709906059a) +also motivate stride support and future QK-norm/RoPE fusion. This contribution +qualifies the unfused operator; it makes no new-algorithm or fused-model claim. + +## Validation and benchmarks + +```bash +python -m pytest tests/ops/rope/test_multi_axis_rope.py -q -rs +python tools/validation/operators/check_operator.py --op multi_axis_rope \ + --candidate cuda --device cuda --dtype bf16 --batch 2 --seq 7 --n-heads 3 --check-grad +python -m benchmarks.operators.rope.benchmark_multi_axis_rope --output rope.json +python tools/validation/operators/validate_multi_axis_rope.py --ws1 cuda --output ws1-cuda.json +python tools/validation/operators/validate_multi_axis_rope.py --ws1 triton --output ws1-triton.json +mkdocs build --strict -f mkdocs.yaml +``` + +For operator-only native compilation, without building every unrelated extension: + +```bash +python tools/validation/operators/validate_multi_axis_rope.py -q -rs +python tools/validation/operators/validate_multi_axis_rope.py --benchmark --output rope.json +``` + +To audit against the actual pinned upstream symbols without installing Diffusers, +download the source file linked above and run: + +```bash +python tools/validation/operators/check_qwenimage_rope_upstream.py \ + --source transformer_qwenimage.py --output upstream.json +``` + +The audit checks the source checksum before loading the two relevant definitions, +then compares image/text tables and forward/backward across three dtypes. + +The focused runner loads the **same production CUDA source and dispatcher +registration**, without a generated binding or extension monkeypatch. It reuses +an installed native library when available. Set a matching `CUDA_HOME`, CUDA +architecture and host compiler before invoking it. + +Ordinary CUDA builds include the source in `rl_engine._C`. An explicit operator +wheel profile builds `rl_engine._rope_cuda` with the same dispatcher registration: + +```bash +RL_KERNEL_ROPE_ONLY=1 python -m pip wheel --no-build-isolation --no-deps . -w dist +python -m pip install --force-reinstall --no-deps dist/rl_kernel-*.whl +# Change to a directory outside the source checkout before this command: +python /path/to/RL-Kernel/tools/validation/operators/check_multi_axis_rope_install.py \ + --package-root /path/to/site-packages --output installed-wheel.json +``` + +On Windows set `$env:RL_KERNEL_ROPE_ONLY='1'` in PowerShell and enter the matching +MSVC build environment. Select a CUDA toolkit matching `torch.version.cuda`. +The operator profile does not certify the complete collection of unrelated +native operators. CPU-only/ROCm builds reject this explicit CUDA wheel profile. + +`multi_axis_rope` has an `OP_SPECS` entry, standard input builder and gradient +adapter. `--ws1` runs the existing forward and gradient accuracy/invariance APIs +with C1's elementwise tolerance profile and C2's logical configuration transforms. +This gate qualifies the Qwen-Image operator and does not certify the Dense Qwen3 +chain, distributed execution, or model assembly. + +Tests cover independent forward/VJP accuracy, adjacent-pair layout, shared and +batched positions, odd and rectangular grids, batch/padding/chunk invariance, +non-contiguous tensors, packed QKV views without copies, default-device contexts, +empty inputs, boundary rejection, higher derivatives, CUDA/Triton byte equality, +graph replay, the current CUDA stream, `torch.library.opcheck` and full-graph +compilation. A two-device guard test skips when only one GPU is present. + +The benchmark covers a short synthetic case and 64×64, 83×83 and 104×58 grids, +with 24 heads and both contiguous/packed-QKV layouts. It checks forward and input +gradients before timing and records CUDA Graph and synchronized eager wall time, +extra peak allocation, source hashes and actual-backend traces. Table generation +and compilation are excluded. These measurements establish no model speedup. + +To include the mature FlashAttention baselines, download `flash_attn/ops/triton/rotary.py` +and `flash_attn/layers/rotary.py` at the linked revision, then pass both paths: + +```bash +python tools/validation/operators/validate_multi_axis_rope.py --benchmark \ + --flash-source /path/to/ops/rotary.py --flash-layer-source /path/to/layers/rotary.py \ + --output rope-with-flash.json +``` + +The normalized source checksums are enforced before loading the audited symbols. +The default variant retains upstream arithmetic; `flash_no_fma` changes only the +kernel launch contraction setting and must match the FP32 gold's output bytes. diff --git a/pyproject.toml b/pyproject.toml index 207606282..f4f8b8ae9 100644 --- a/pyproject.toml +++ b/pyproject.toml @@ -41,7 +41,7 @@ where = ["."] include = ["rl_engine*"] [tool.setuptools.package-data] -rl_engine = ["**/*.json", "**/*.jsonl", "**/*.sh", "**/*.patch", "*.pyi"] +rl_engine = ["**/*.json", "**/*.jsonl", "**/*.sh", "**/*.patch", "*.pyi", "**/FLASH_ATTENTION_LICENSE"] [tool.ruff] line-length = 100 diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/README.md b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/README.md new file mode 100644 index 000000000..64232c731 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/README.md @@ -0,0 +1,152 @@ +# Qwen-Image multi-axis RoPE validation — 2026-10-28 + +Evidence for the CUDA `multi_axis_rope` row of +[RFC #386](https://github.com/RL-Align/RL-Kernel/issues/386), based on +`test-qwenimage@16b81175874ecf11d5ff05c00de3b69e9457591b`. +The source fingerprints in `manifest.json` identify the final implementation. +This is single-GPU operator evidence; it does not qualify a model or trajectory. + +## Implementation + +- Adjacent pairs, axes `[16, 56, 56]`, canonical centered image positions and + diagonal text offsets, pinned to Diffusers `40f63690da440e8c382ab53e67508835deb1ed2f`. +- Explicit CPU table factories remain CPU factories inside CUDA/meta default-device contexts. +- Native CUDA reads all input/table strides directly, including packed QKV views. + The shared Triton implementation adapts FlashAttention's existing rotary kernel + at `98c1cee0ed59307833841cdbf567fe336ceec13a`, retaining attribution and BSD-3 terms. +- Both GPU paths have dispatcher, fake and input-autograd registration. The + arithmetic is separate FP32 multiply/add followed by one output cast. +- OP_SPECS, the ordinary input builder and a gradient adapter expose the operator + to existing validation APIs. The operator gate runs all four forward/gradient + accuracy/invariance judgments with explicit backend/kernel provenance. +- Default CUDA builds include the native source in `_C`. The explicit + `RL_KERNEL_ROPE_ONLY=1` build profile also produces an installable companion + library, without compiling unrelated native operators. + +## Environment + +| Item | Value | +| --- | --- | +| GPU | NVIDIA GeForce RTX 5080, SM120, 16 GB | +| Driver | 617.42 | +| OS / Python | Windows / 3.13.9 | +| PyTorch / CUDA toolkit | 2.9.1+cu128 / 12.8 | +| Triton | triton-windows 3.5.1.post24; runtime 3.5.1 | +| Native compiler | MSVC 2022 x64, C++17; CUDA target 12.0 | +| Sanitizer | NVIDIA Compute Sanitizer 13.3, memcheck | +| Benchmark dtype | BF16, FP32 position tables | + +The installed wheel is CPython 3.13/Windows-specific. Toolkit/driver versions +have different roles; the native code was compiled with CUDA 12.8, not 13.3. + +## Validation + +| Check | Result | Evidence | +| --- | --- | --- | +| Public operator suite using the built companion library | 81 passed, 1 two-GPU skip | `tests.txt` | +| Operator/build/registry/validation regressions | 195 passed, 7 Ascend skips | `regression.txt` | +| CUDA WS1 four judgments | Passed | `ws1-cuda.json` | +| Triton WS1 four judgments | Passed | `ws1-triton.json` | +| Standard CLI, BF16 forward and input gradient, CUDA/Triton | Both passed | `operator-{cuda,triton}.json` | +| Pinned actual Diffusers symbols | 30 cases passed; image/text tables exact | `upstream.json` | +| Memory safety on selected accuracy/strided/empty/stream cases | 34 passed, zero memcheck errors | `memcheck.txt` | +| Source distribution + operator wheel build | Passed; native sources and BSD license present | `package-build.txt`, `package.json` | +| Fresh interpreter outside checkout after wheel installation | CUDA/Triton eager and fullgraph forward/input-gradient bytes exact | `installed-wheel.json` | +| Changed Python formatting/static checks; strict docs build | Passed | `quality.txt` | +| Repository layout check | 5 failures, 31 passed, 1 skipped | `layout.txt` | +| Same layout check on clean base | Same 5 failures, 27 passed, 1 skipped | `baseline-layout.txt` | + +Operator tests cover FP32/FP16/BF16, shared/per-sample tables, odd and rectangular +grids, batch/chunk/padding invariance, higher input VJPs, strided gradients, +independent table strides, empty tensors, invalid-input rejection, non-default +streams, graph replay with changed static buffers, `torch.library.opcheck`, and +`torch.compile(..., fullgraph=True)` through backward. A CPU profiler assertion +rejects `contiguous`, `clone` or `copy_` inside either GPU call for a packed view. + +The layout failures reproduce on the clean base: the Windows driver-link flag +assertion, symlink creation without Windows privilege, and three checks affected +by symlink-disabled JSON/source placeholders. The added build tests pass. This +report does not describe the entire repository test suite as green. + +## Mature baseline and performance + +`benchmark-bfloat16.json` contains 40 pre-timing correctness comparisons and 80 +measurements: four geometries, two layouts, five implementations, two phases. +The CUDA, adapted Triton and contraction-disabled upstream FlashAttention paths +all have zero differing output/gradient bytes against the unfused FP32 gold. +Original FlashAttention and Diffusers complex arithmetic use declared tolerances. + +The optional FlashAttention comparison loads the checksum-pinned original +`rotary_kernel`, `apply_rotary` and `ApplyRotaryEmb` class without installing +flash-attn. `flash_no_fma` adds only `enable_fp_fusion=False` to the original +launch; the default baseline retains its fused arithmetic. + +CUDA Graph medians in ms/call; each cell is **forward / forward+backward**. +Batch 1, 24 heads, D=128: + +| Grid / layout | Native CUDA | Adapted Triton | FlashAttention, no FMA | Diffusers complex | +| --- | --- | --- | --- | --- | +| 64×64 contiguous | 0.05771 / 0.12953 | 0.05191 / 0.12463 | 0.05202 / 0.12443 | 0.24277 / 0.50118 | +| 64×64 packed QKV | 0.05866 / 0.13089 | 0.04835 / 0.12314 | 0.04896 / 0.12374 | 0.24599 / 0.50083 | +| 83×83 contiguous | 0.11203 / 0.22478 | 0.10685 / 0.21339 | 0.10693 / 0.21351 | 0.51009 / 1.05054 | +| 83×83 packed QKV | 0.11280 / 0.22514 | 0.10758 / 0.21415 | 0.10746 / 0.21431 | 0.51117 / 1.05461 | +| 104×58 contiguous | 0.09809 / 0.19646 | 0.09356 / 0.18708 | 0.09362 / 0.18709 | 0.43411 / 0.89356 | +| 104×58 packed QKV | 0.09596 / 0.19664 | 0.09341 / 0.18776 | 0.09344 / 0.18786 | 0.43668 / 0.89561 | + +The adapted Triton path is close to its mature upstream baseline. The native +CUDA path avoids requiring Triton but is slower in this graph comparison. +No speed advantage over FlashAttention is established. At 4096 tokens, either +GPU path's forward extra peak is 24 MiB, exactly the returned BF16 output size, +for both layouts. There is no extra packed-input materialization. + +Five warm calls/replays, 20 calls per captured graph, median of 30 replay +measurements. Synchronized eager wall time is also recorded as three trials of +30 calls. It includes Python/dispatcher/autograd overhead and can differ +substantially from graph time. Tables, compilation, input allocation and packed +backing storage are outside timing. Extra peak allocation includes the returned +output or gradient. These are one-machine, one-process microbenchmarks, without +clock pinning; close differences do not establish a stable winner. + +## Reproduction + +Set a toolkit matching PyTorch, the target architecture and a matching host +compiler. On the recorded Windows host, run from an x64 MSVC environment with +`CUDA_HOME`/`CUDA_PATH` selecting CUDA 12.8 and `TORCH_CUDA_ARCH_LIST=12.0`. +Use `$env:RL_KERNEL_ROPE_ONLY='1'` for the operator package profile. + +```bash +python setup.py sdist bdist_wheel +python setup.py build_ext --inplace +python -m pytest tests/ops/rope/test_multi_axis_rope.py -q -rs +python tools/validation/operators/check_operator.py --op multi_axis_rope \ + --candidate cuda --device cuda --dtype bf16 --batch 2 --seq 7 --n-heads 3 --check-grad --json +python tools/validation/operators/validate_multi_axis_rope.py --ws1 cuda --output ws1-cuda.json +python tools/validation/operators/validate_multi_axis_rope.py --ws1 triton --output ws1-triton.json +python tools/validation/operators/validate_multi_axis_rope.py --benchmark \ + --flash-source /path/to/flash_attn/ops/triton/rotary.py \ + --flash-layer-source /path/to/flash_attn/layers/rotary.py --output benchmark.json +python tools/validation/operators/check_qwenimage_rope_upstream.py \ + --source /path/to/transformer_qwenimage.py --output upstream.json +compute-sanitizer --tool memcheck --error-exitcode 1 python \ + tools/validation/operators/validate_multi_axis_rope.py -q \ + -k "forward_backward_accuracy or noncontiguous_input_tables_and_gradient or empty_forward_backward or packed_qkv_and_independent_table_strides or native_entrypoint_checks" +mkdocs build --strict -f mkdocs.yaml +``` + +Install the built wheel, change to a directory outside the checkout, then run +`check_multi_axis_rope_install.py --package-root /path/to/site-packages --output installed-wheel.json`. +This probe neither changes sys.path nor invokes the focused JIT loader. +See `quality.txt` for the regression command and static-check scope. + +## Related work and remaining qualification + +The [operator page](../../../docs/operators/multi-axis-rope.md) documents +FlashAttention reuse and the related Ascend #411, SGLang and vLLM-Omni paths. +The live repository search on 2026-10-28 found #411 for the exact row name; +its head remains `285c4dccab7e9048956584caf431946d82405c09`. Its half-split and +uncentered reference differs from the pinned adjacent-pair/centered contract. + +Linux, A100/H100, multi-GPU guards, the complete monolithic `_C` build, QK-norm +fusion, model/block assembly and trajectories remain unqualified. There is no +Dense Qwen3 or distributed certification implied by the operator WS1 adapter. +Cross-version/device byte equality and fast-math builds are not claimed. diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/baseline-layout.txt b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/baseline-layout.txt new file mode 100644 index 000000000..16a5c97e8 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/baseline-layout.txt @@ -0,0 +1,138 @@ +.FF.............s........F.FF.... [100%] +================================== FAILURES =================================== +________________ test_cuda_build_keeps_existing_ipc_collective ________________ + +monkeypatch = <_pytest.monkeypatch.MonkeyPatch object at 0x0000020FF72168B0> + + def test_cuda_build_keeps_existing_ipc_collective(monkeypatch) -> None: + extension = _load_extension_config(monkeypatch, hip=None) + + assert "csrc/cuda/collectives/deterministic_collective.cu" in extension["sources"] + assert "-DKERNEL_ALIGN_WITH_CUDA" in extension["extra_compile_args"]["cxx"] + assert "-DKERNEL_ALIGN_WITH_ROCM" not in extension["extra_compile_args"]["cxx"] +> assert "-lcuda" in extension["extra_link_args"] +E AssertionError: assert '-lcuda' in ['-Wl,-rpath,$ORIGIN/../torch/lib'] + +tests\build\test_build_platform_collectives.py:61: AssertionError +____________ test_package_data_json_links_resolve_to_real_content _____________ + +tmp_path = WindowsPath('C:/Users/123/AppData/Local/Temp/pytest-of-123/pytest-4/test_package_data_json_links_r0') + + def test_package_data_json_links_resolve_to_real_content(tmp_path) -> None: + target = tmp_path / "configs" / "profile.json" + target.parent.mkdir() + target.write_text('{"name": "profile"}\n', encoding="utf-8") + source = tmp_path / "rl_engine" / "profile.json" + source.parent.mkdir() +> source.symlink_to("../configs/profile.json") + +tests\build\test_build_platform_collectives.py:70: +_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ + +self = WindowsPath('C:/Users/123/AppData/Local/Temp/pytest-of-123/pytest-4/test_package_data_json_links_r0/rl_engine/profile.json') +target = '../configs/profile.json', target_is_directory = False + + def symlink_to(self, target, target_is_directory=False): + """ + Make this path a symlink pointing to the target path. + Note the order of arguments (link, target) is the reverse of os.symlink. + """ +> os.symlink(target, self, target_is_directory) +E OSError: [WinError 1314] �ͻ���û���������Ȩ��: '../configs/profile.json' -> 'C:\\Users\\123\\AppData\\Local\\Temp\\pytest-of-123\\pytest-4\\test_package_data_json_links_r0\\rl_engine\\profile.json' + +D:\anaconda\Lib\pathlib\_local.py:789: OSError +_______ test_recorded_algorithm_sources_resolve_to_real_implementations _______ + + def test_recorded_algorithm_sources_resolve_to_real_implementations(): + from rl_engine.config.workload import load_manifest + + moves = json.loads((ROOT / "tools/migration/layout.json").read_text())["moves"] + for case in load_manifest().representative_cases: + source, symbol = case["provenance_evidence"]["algorithm_source"].rsplit(":", 1) + implementation = ROOT / moves.get(source, source) +> text = implementation.read_text() + ^^^^^^^^^^^^^^^^^^^^^^^^^^ + +tests\runtime\test_layout_compatibility.py:87: +_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ +D:\anaconda\Lib\pathlib\_local.py:546: in read_text + return PathBase.read_text(self, encoding, errors, newline) + ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ + +self = WindowsPath('D:/dandan123/dandan/issue/investigation-rl-kernel/baseline-validation/csrc/cuda/attention/deterministic_attention.cu') +encoding = 'locale', errors = None, newline = None + + def read_text(self, encoding=None, errors=None, newline=None): + """ + Open the file in text mode, read it, and close the file. + """ + with self.open(mode='r', encoding=encoding, errors=errors, newline=newline) as f: +> return f.read() + ^^^^^^^^ +E UnicodeDecodeError: 'gbk' codec can't decode byte 0x92 in position 173: illegal multibyte sequence + +D:\anaconda\Lib\pathlib\_abc.py:633: UnicodeDecodeError +_ test_packaged_contract_bytes_match_main[rl_engine.config-workloads/qwen3_8b.json-80395644b5e84ed1a76729ffc6f78eb11f41b9ab42fb36e2d00f19635d3df0f6] _ + +package = 'rl_engine.config', resource = 'workloads/qwen3_8b.json' +digest = '80395644b5e84ed1a76729ffc6f78eb11f41b9ab42fb36e2d00f19635d3df0f6' + + @pytest.mark.parametrize( + ("package", "resource", "digest"), + [ + ( + "rl_engine.config", + "workloads/qwen3_8b.json", + "80395644b5e84ed1a76729ffc6f78eb11f41b9ab42fb36e2d00f19635d3df0f6", + ), + ( + "rl_engine.contracts", + "profiles/precision/ws1.json", + "044167ec21e372017a2758dc09f805b368283c06b83f64061877628cbd101532", + ), + ], + ) + def test_packaged_contract_bytes_match_main(package, resource, digest): +> assert hashlib.sha256(files(package).joinpath(resource).read_bytes()).hexdigest() == digest +E AssertionError: assert 'c97e966fb075...75e111e052d7d' == '80395644b5e8...f19635d3df0f6' +E +E - 80395644b5e84ed1a76729ffc6f78eb11f41b9ab42fb36e2d00f19635d3df0f6 +E + c97e966fb075dad8ce7dcd7a23a58ffb856e7474dfb3c50510675e111e052d7d + +tests\runtime\test_layout_compatibility.py:118: AssertionError +_ test_packaged_contract_bytes_match_main[rl_engine.contracts-profiles/precision/ws1.json-044167ec21e372017a2758dc09f805b368283c06b83f64061877628cbd101532] _ + +package = 'rl_engine.contracts', resource = 'profiles/precision/ws1.json' +digest = '044167ec21e372017a2758dc09f805b368283c06b83f64061877628cbd101532' + + @pytest.mark.parametrize( + ("package", "resource", "digest"), + [ + ( + "rl_engine.config", + "workloads/qwen3_8b.json", + "80395644b5e84ed1a76729ffc6f78eb11f41b9ab42fb36e2d00f19635d3df0f6", + ), + ( + "rl_engine.contracts", + "profiles/precision/ws1.json", + "044167ec21e372017a2758dc09f805b368283c06b83f64061877628cbd101532", + ), + ], + ) + def test_packaged_contract_bytes_match_main(package, resource, digest): +> assert hashlib.sha256(files(package).joinpath(resource).read_bytes()).hexdigest() == digest +E AssertionError: assert '2020c02e7af4...313d6752cd34e' == '044167ec21e3...7628cbd101532' +E +E - 044167ec21e372017a2758dc09f805b368283c06b83f64061877628cbd101532 +E + 2020c02e7af4e2ca4163e1ecfb9c491baef07b73deba39b0f44313d6752cd34e + +tests\runtime\test_layout_compatibility.py:118: AssertionError +=========================== short test summary info =========================== +FAILED tests/build/test_build_platform_collectives.py::test_cuda_build_keeps_existing_ipc_collective +FAILED tests/build/test_build_platform_collectives.py::test_package_data_json_links_resolve_to_real_content +FAILED tests/runtime/test_layout_compatibility.py::test_recorded_algorithm_sources_resolve_to_real_implementations +FAILED tests/runtime/test_layout_compatibility.py::test_packaged_contract_bytes_match_main[rl_engine.config-workloads/qwen3_8b.json-80395644b5e84ed1a76729ffc6f78eb11f41b9ab42fb36e2d00f19635d3df0f6] +FAILED tests/runtime/test_layout_compatibility.py::test_packaged_contract_bytes_match_main[rl_engine.contracts-profiles/precision/ws1.json-044167ec21e372017a2758dc09f805b368283c06b83f64061877628cbd101532] +5 failed, 27 passed, 1 skipped in 5.70s diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/benchmark-bfloat16.json b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/benchmark-bfloat16.json new file mode 100644 index 000000000..72fb42bbe --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/benchmark-bfloat16.json @@ -0,0 +1,2781 @@ +{ + "gpu": "NVIDIA GeForce RTX 5080", + "capability": [ + 12, + 0 + ], + "os": "Windows-11-10.0.26200-SP0", + "python": "3.13.9 | packaged by Anaconda, Inc. | (main, Oct 21 2025, 19:09:58) [MSC v.1929 64 bit (AMD64)]", + "torch": "2.9.1+cu128", + "cuda_runtime": "12.8", + "triton": "3.5.1", + "dtype": "bfloat16", + "base_commit": "16b81175874ecf11d5ff05c00de3b69e9457591b", + "source_sha256": { + "csrc/cuda/rope/multi_axis_rope.cu": "f5db5cf5d4962053194455d9cde8cf91cf4e824bad7abf72e8410a404bcb0e44", + "rl_engine/backends/shared/triton/rope/_flash_attn_rotary.py": "5c6984da351e5228f7549629d4ff53e5772b9a6ef2a01ba2a38493e769b15ab5", + "rl_engine/backends/shared/triton/rope/multi_axis_rope.py": "5b989bd63efd77913267b48654d10204bc563768c6c5d4e1cde9fb55e1354547", + "rl_engine/backends/cuda/rope/multi_axis_rope.py": "d779d9b60ce5d1d6651925f723049930d898024fdf0548ce29a2019b84753d65", + "rl_engine/ops/rope/_registration.py": "add0375d1639444908091d01e8d29781f7a975beffec64b77bfaa81a5b9a0ae3", + "rl_engine/ops/rope/multi_axis_rope.py": "07bcff1160f22c1ff6730be731206b2f4a153e2fada26afc93fc884e780568cb", + "rl_engine/reference/rope/multi_axis_rope.py": "c9a94730047f51659f7aa0a443e88d6ce7d03d00d6813a9aed548cafbf6347bf", + "benchmarks/operators/rope/benchmark_multi_axis_rope.py": "6961a55803fbd4e566bd060558fb59e658e7f1f704e61ca97f912e5d16b55884" + }, + "timing": "CUDA graph ms/call and synchronized eager wall ms/call; warm allocation; excludes tables, compilation, and initial allocation of inputs", + "memory": "extra peak allocated bytes excluding inputs/tables/grad; includes returned output or gradient; packed QKV backing storage allocated before timing", + "iterations": 30, + "warmup": 5, + "trials": 3, + "graph_calls": 20, + "flash_baseline": { + "included": true, + "revision": "98c1cee0ed59307833841cdbf567fe336ceec13a", + "kernel_sha256_normalized": "d2037e623854c702fb075784fe4292d4a89d6b9ac1ae33cab3d6788f8f94f240", + "layer_sha256_normalized": "a79cb42dfc33e126227bdba27ae7a024d7d5cc5c8fd7f7a5f26c68d501024e0a", + "mode": "upstream rotary_kernel, apply_rotary and ApplyRotaryEmb, interleaved=True", + "no_fma_change": "only add enable_fp_fusion=False at the original kernel launch" + }, + "trace": { + "cuda": { + "contract": "qwenimage-interleaved-rope-fp32-v1", + "actual_backend": "rlkernel.cuda.qwenimage_multi_axis_rope.v1", + "kernel_id": "torch.ops.rl_kernel_native.multi_axis_rope", + "layout": "BSHD, adjacent channel pairs", + "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", + "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" + }, + "triton": { + "contract": "qwenimage-interleaved-rope-fp32-v1", + "actual_backend": "rlkernel.triton.qwenimage_multi_axis_rope.v1", + "kernel_id": "rl_engine.backends.shared.triton.rope._flash_attn_rotary.rotary_kernel", + "layout": "BSHD, adjacent channel pairs", + "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", + "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" + } + }, + "qualification": "Exact bytes vs unfused FP32 gold for CUDA, Triton and flash_no_fma; declared tolerance for upstream complex and default FlashAttention", + "correctness": [ + { + "backend": "upstream_complex", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "cuda", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "triton", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "flash_default", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "flash_no_fma", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "upstream_complex", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "cuda", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "triton", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "flash_default", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "flash_no_fma", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "upstream_complex", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "different_bytes_y_dx": [ + 130, + 139 + ] + }, + { + "backend": "cuda", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "triton", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "flash_default", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "different_bytes_y_dx": [ + 130, + 139 + ] + }, + { + "backend": "flash_no_fma", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "upstream_complex", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "different_bytes_y_dx": [ + 139, + 140 + ] + }, + { + "backend": "cuda", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "triton", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "flash_default", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "different_bytes_y_dx": [ + 139, + 140 + ] + }, + { + "backend": "flash_no_fma", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "upstream_complex", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "different_bytes_y_dx": [ + 234, + 225 + ] + }, + { + "backend": "cuda", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "triton", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "flash_default", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "different_bytes_y_dx": [ + 234, + 225 + ] + }, + { + "backend": "flash_no_fma", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "upstream_complex", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "different_bytes_y_dx": [ + 220, + 226 + ] + }, + { + "backend": "cuda", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "triton", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "flash_default", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "different_bytes_y_dx": [ + 220, + 226 + ] + }, + { + "backend": "flash_no_fma", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "upstream_complex", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "different_bytes_y_dx": [ + 196, + 203 + ] + }, + { + "backend": "cuda", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "triton", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "flash_default", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "different_bytes_y_dx": [ + 196, + 203 + ] + }, + { + "backend": "flash_no_fma", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "upstream_complex", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "different_bytes_y_dx": [ + 230, + 218 + ] + }, + { + "backend": "cuda", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "triton", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "flash_default", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "different_bytes_y_dx": [ + 230, + 218 + ] + }, + { + "backend": "flash_no_fma", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + } + ], + "results": [ + { + "backend": "upstream_complex", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 18432, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.003558399900794029, + "graph_min_ms": 0.0035264000296592714, + "graph_max_ms": 0.015438400208950043, + "eager_wall_median_ms": 0.0339666667059646, + "eager_wall_trials_ms": [ + 0.0356399999873247, + 0.0339666667059646, + 0.03332666662269427 + ], + "extra_peak_bytes": 184320 + }, + { + "backend": "upstream_complex", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 18432, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.007688799872994423, + "graph_min_ms": 0.007651200145483017, + "graph_max_ms": 0.007790400087833405, + "eager_wall_median_ms": 0.21265999997316007, + "eager_wall_trials_ms": [ + 0.21187666667780528, + 0.21265999997316007, + 0.22587333332921844 + ], + "extra_peak_bytes": 187392 + }, + { + "backend": "cuda", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 18432, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.0012016000226140022, + "graph_min_ms": 0.0011567999608814716, + "graph_max_ms": 0.0012384000234305858, + "eager_wall_median_ms": 0.02289333333465038, + "eager_wall_trials_ms": [ + 0.02289333333465038, + 0.02245666667780218, + 0.02458333328831941 + ], + "extra_peak_bytes": 36864 + }, + { + "backend": "cuda", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 18432, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.0022280000150203703, + "graph_min_ms": 0.0021888000890612604, + "graph_max_ms": 0.0022655999287962913, + "eager_wall_median_ms": 0.1764300000407578, + "eager_wall_trials_ms": [ + 0.20031666669334905, + 0.1764300000407578, + 0.15246333338533682 + ], + "extra_peak_bytes": 73728 + }, + { + "backend": "triton", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 18432, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.001142400037497282, + "graph_min_ms": 0.0010591999627649785, + "graph_max_ms": 0.005691200122237205, + "eager_wall_median_ms": 0.057766666698929235, + "eager_wall_trials_ms": [ + 0.07126333330234047, + 0.057766666698929235, + 0.05640333335274287 + ], + "extra_peak_bytes": 36864 + }, + { + "backend": "triton", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 18432, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.002125599980354309, + "graph_min_ms": 0.002091200090944767, + "graph_max_ms": 0.002160000056028366, + "eager_wall_median_ms": 0.21872666663208898, + "eager_wall_trials_ms": [ + 0.22702333335473668, + 0.21872666663208898, + 0.21611666670651175 + ], + "extra_peak_bytes": 73728 + }, + { + "backend": "flash_default", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 18432, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.0010935999918729067, + "graph_min_ms": 0.001033599954098463, + "graph_max_ms": 0.0011536000296473504, + "eager_wall_median_ms": 0.10024000002886169, + "eager_wall_trials_ms": [ + 0.10126666669142044, + 0.09858333329854456, + 0.10024000002886169 + ], + "extra_peak_bytes": 36864 + }, + { + "backend": "flash_default", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 18432, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.0020527999848127367, + "graph_min_ms": 0.002019199915230274, + "graph_max_ms": 0.0021023999899625776, + "eager_wall_median_ms": 0.3111766666431019, + "eager_wall_trials_ms": [ + 0.30222666664485587, + 0.3111766666431019, + 0.333893333360417 + ], + "extra_peak_bytes": 73728 + }, + { + "backend": "flash_no_fma", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 18432, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.0010991999879479408, + "graph_min_ms": 0.001033599954098463, + "graph_max_ms": 0.0011392000131309032, + "eager_wall_median_ms": 0.10421333336125826, + "eager_wall_trials_ms": [ + 0.10421333336125826, + 0.10106999998242827, + 0.12048333337588701 + ], + "extra_peak_bytes": 36864 + }, + { + "backend": "flash_no_fma", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 18432, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.0020503999665379526, + "graph_min_ms": 0.002001599967479706, + "graph_max_ms": 0.00886240005493164, + "eager_wall_median_ms": 0.31822000000829576, + "eager_wall_trials_ms": [ + 0.28550666665978497, + 0.31822000000829576, + 0.33777000001767493 + ], + "extra_peak_bytes": 73728 + }, + { + "backend": "upstream_complex", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 55296, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.00379519984126091, + "graph_min_ms": 0.003745599836111069, + "graph_max_ms": 0.003859199956059456, + "eager_wall_median_ms": 0.10653666668076767, + "eager_wall_trials_ms": [ + 0.10653666668076767, + 0.10879000001295935, + 0.09587666669782872 + ], + "extra_peak_bytes": 184320 + }, + { + "backend": "upstream_complex", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 55296, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.008064000308513642, + "graph_min_ms": 0.008011200278997422, + "graph_max_ms": 0.008102399855852127, + "eager_wall_median_ms": 0.2260500000071867, + "eager_wall_trials_ms": [ + 0.23646333332483968, + 0.2187999999781217, + 0.2260500000071867 + ], + "extra_peak_bytes": 187392 + }, + { + "backend": "cuda", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 55296, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.0012024000287055969, + "graph_min_ms": 0.0011392000131309032, + "graph_max_ms": 0.0025040000677108766, + "eager_wall_median_ms": 0.023706666615908034, + "eager_wall_trials_ms": [ + 0.023706666615908034, + 0.024339999981748406, + 0.023019999995691858 + ], + "extra_peak_bytes": 36864 + }, + { + "backend": "cuda", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 55296, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.002226400002837181, + "graph_min_ms": 0.002174399979412556, + "graph_max_ms": 0.0022671999409794807, + "eager_wall_median_ms": 0.157756666643157, + "eager_wall_trials_ms": [ + 0.165013333329019, + 0.15389999995628992, + 0.157756666643157 + ], + "extra_peak_bytes": 73728 + }, + { + "backend": "triton", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 55296, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.0011008000001311302, + "graph_min_ms": 0.0010688000358641148, + "graph_max_ms": 0.0011359999887645244, + "eager_wall_median_ms": 0.059236666675133165, + "eager_wall_trials_ms": [ + 0.053669999942940194, + 0.06250666665437166, + 0.059236666675133165 + ], + "extra_peak_bytes": 36864 + }, + { + "backend": "triton", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 55296, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.0021215999498963357, + "graph_min_ms": 0.002065600082278252, + "graph_max_ms": 0.0021679999306797983, + "eager_wall_median_ms": 0.25444666662224336, + "eager_wall_trials_ms": [ + 0.25444666662224336, + 0.27043333332888625, + 0.2240899999984928 + ], + "extra_peak_bytes": 73728 + }, + { + "backend": "flash_default", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 55296, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.001101600006222725, + "graph_min_ms": 0.001065600011497736, + "graph_max_ms": 0.0011376000009477138, + "eager_wall_median_ms": 0.12566333334689261, + "eager_wall_trials_ms": [ + 0.12566333334689261, + 0.12918333328949907, + 0.10983666664590903 + ], + "extra_peak_bytes": 36864 + }, + { + "backend": "flash_default", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 55296, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.0020527999848127367, + "graph_min_ms": 0.0020160000771284104, + "graph_max_ms": 0.0021023999899625776, + "eager_wall_median_ms": 0.28767333333235, + "eager_wall_trials_ms": [ + 0.2609133333559536, + 0.28767333333235, + 0.3289566666353494 + ], + "extra_peak_bytes": 73728 + }, + { + "backend": "flash_no_fma", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 55296, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.001101600006222725, + "graph_min_ms": 0.0010223999619483949, + "graph_max_ms": 0.0022975999861955644, + "eager_wall_median_ms": 0.0995366666757036, + "eager_wall_trials_ms": [ + 0.10132333336514421, + 0.0995366666757036, + 0.0990000000456348 + ], + "extra_peak_bytes": 36864 + }, + { + "backend": "flash_no_fma", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 55296, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.0020527999848127367, + "graph_min_ms": 0.0020128000527620316, + "graph_max_ms": 0.0021007999777793883, + "eager_wall_median_ms": 0.3198099999887442, + "eager_wall_trials_ms": [ + 0.33375666668386356, + 0.30899333335886087, + 0.3198099999887442 + ], + "extra_peak_bytes": 73728 + }, + { + "backend": "upstream_complex", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 12582912, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.2427680015563965, + "graph_min_ms": 0.24123198986053468, + "graph_max_ms": 0.3049216032028198, + "eager_wall_median_ms": 0.25121666670505266, + "eager_wall_trials_ms": [ + 0.24960999999166233, + 0.2669566666857766, + 0.25121666670505266 + ], + "extra_peak_bytes": 125829120 + }, + { + "backend": "upstream_complex", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 12582912, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.5011807918548584, + "graph_min_ms": 0.4943295955657959, + "graph_max_ms": 0.5467391967773437, + "eager_wall_median_ms": 0.5078799999561548, + "eager_wall_trials_ms": [ + 0.5254200000005463, + 0.5078799999561548, + 0.5070699999729792 + ], + "extra_peak_bytes": 127926272 + }, + { + "backend": "cuda", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 12582912, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.05770559906959534, + "graph_min_ms": 0.05762240290641785, + "graph_max_ms": 0.07225120067596436, + "eager_wall_median_ms": 0.061780000032740645, + "eager_wall_trials_ms": [ + 0.06511333334249987, + 0.06060333331940152, + 0.061780000032740645 + ], + "extra_peak_bytes": 25165824 + }, + { + "backend": "cuda", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 12582912, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.12953280210494994, + "graph_min_ms": 0.12918239831924438, + "graph_max_ms": 0.14451359510421752, + "eager_wall_median_ms": 0.15664999997776854, + "eager_wall_trials_ms": [ + 0.2333100000517637, + 0.15664999997776854, + 0.14973666669296412 + ], + "extra_peak_bytes": 50331648 + }, + { + "backend": "triton", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 12582912, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.05190559923648834, + "graph_min_ms": 0.05078719854354859, + "graph_max_ms": 0.06408159732818604, + "eager_wall_median_ms": 0.0845666666161075, + "eager_wall_trials_ms": [ + 0.08885000000494377, + 0.08201333330362104, + 0.0845666666161075 + ], + "extra_peak_bytes": 25165824 + }, + { + "backend": "triton", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 12582912, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.12463040351867677, + "graph_min_ms": 0.12331520318984986, + "graph_max_ms": 0.1529088020324707, + "eager_wall_median_ms": 0.2539966667124342, + "eager_wall_trials_ms": [ + 0.2757299999454214, + 0.25390333333537757, + 0.2539966667124342 + ], + "extra_peak_bytes": 50331648 + }, + { + "backend": "flash_default", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 12582912, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.05208719968795776, + "graph_min_ms": 0.05123199820518494, + "graph_max_ms": 0.06861280202865601, + "eager_wall_median_ms": 0.13019999999717888, + "eager_wall_trials_ms": [ + 0.1370399999965836, + 0.12465666665472479, + 0.13019999999717888 + ], + "extra_peak_bytes": 25165824 + }, + { + "backend": "flash_default", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 12582912, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.12441120147705079, + "graph_min_ms": 0.12314239740371705, + "graph_max_ms": 0.139137601852417, + "eager_wall_median_ms": 0.340176666698729, + "eager_wall_trials_ms": [ + 0.35652999995363643, + 0.340176666698729, + 0.33921999996285496 + ], + "extra_peak_bytes": 50331648 + }, + { + "backend": "flash_no_fma", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 12582912, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.052018401026725766, + "graph_min_ms": 0.05074719786643982, + "graph_max_ms": 0.07086880207061767, + "eager_wall_median_ms": 0.10117999997116083, + "eager_wall_trials_ms": [ + 0.10250666667464732, + 0.10117999997116083, + 0.10012000002461718 + ], + "extra_peak_bytes": 25165824 + }, + { + "backend": "flash_no_fma", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 12582912, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.12442959547042846, + "graph_min_ms": 0.12299360036849975, + "graph_max_ms": 0.14664160013198851, + "eager_wall_median_ms": 0.34945666666317265, + "eager_wall_trials_ms": [ + 0.36080999998375773, + 0.3163533333160255, + 0.34945666666317265 + ], + "extra_peak_bytes": 50331648 + }, + { + "backend": "upstream_complex", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 37748736, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.2459879994392395, + "graph_min_ms": 0.24467840194702148, + "graph_max_ms": 0.272545599937439, + "eager_wall_median_ms": 0.2536533333720096, + "eager_wall_trials_ms": [ + 0.27815666665749933, + 0.2529566666756485, + 0.2536533333720096 + ], + "extra_peak_bytes": 125829120 + }, + { + "backend": "upstream_complex", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 37748736, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.5008344173431396, + "graph_min_ms": 0.4996367931365967, + "graph_max_ms": 0.5373487949371338, + "eager_wall_median_ms": 0.5262199999682101, + "eager_wall_trials_ms": [ + 0.5269766666970099, + 0.5070566666593853, + 0.5262199999682101 + ], + "extra_peak_bytes": 127926272 + }, + { + "backend": "cuda", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 37748736, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.05865840017795562, + "graph_min_ms": 0.058550399541854856, + "graph_max_ms": 0.08192480206489564, + "eager_wall_median_ms": 0.06282000000889336, + "eager_wall_trials_ms": [ + 0.06282000000889336, + 0.06640333334265354, + 0.06277333335068154 + ], + "extra_peak_bytes": 25165824 + }, + { + "backend": "cuda", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 37748736, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.1308944046497345, + "graph_min_ms": 0.13030719757080078, + "graph_max_ms": 0.145305597782135, + "eager_wall_median_ms": 0.1688300000144712, + "eager_wall_trials_ms": [ + 0.1688300000144712, + 0.15922666667999388, + 0.17570666665657578 + ], + "extra_peak_bytes": 50331648 + }, + { + "backend": "triton", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 37748736, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.04834639877080918, + "graph_min_ms": 0.047383999824523924, + "graph_max_ms": 0.06508319973945617, + "eager_wall_median_ms": 0.08053666667062014, + "eager_wall_trials_ms": [ + 0.09511666667094687, + 0.07707999999790142, + 0.08053666667062014 + ], + "extra_peak_bytes": 25165824 + }, + { + "backend": "triton", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 37748736, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.1231384038925171, + "graph_min_ms": 0.12216000556945801, + "graph_max_ms": 0.1478927969932556, + "eager_wall_median_ms": 0.2741499999804849, + "eager_wall_trials_ms": [ + 0.28788666665301815, + 0.2741499999804849, + 0.24507333334137607 + ], + "extra_peak_bytes": 50331648 + }, + { + "backend": "flash_default", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 37748736, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.048105600476264956, + "graph_min_ms": 0.0474047988653183, + "graph_max_ms": 0.06945440173149109, + "eager_wall_median_ms": 0.12940666668631215, + "eager_wall_trials_ms": [ + 0.14179000002817096, + 0.12940666668631215, + 0.12683000004471978 + ], + "extra_peak_bytes": 25165824 + }, + { + "backend": "flash_default", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 37748736, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.12347760200500488, + "graph_min_ms": 0.12205439805984497, + "graph_max_ms": 0.14951200485229493, + "eager_wall_median_ms": 0.36926333335334977, + "eager_wall_trials_ms": [ + 0.36926333335334977, + 0.3322200000184239, + 0.36990666667406913 + ], + "extra_peak_bytes": 50331648 + }, + { + "backend": "flash_no_fma", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 37748736, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.04896320104598999, + "graph_min_ms": 0.047470399737358095, + "graph_max_ms": 0.06636160016059875, + "eager_wall_median_ms": 0.14505666664869446, + "eager_wall_trials_ms": [ + 0.14700000004571243, + 0.14505666664869446, + 0.13565333332129134 + ], + "extra_peak_bytes": 25165824 + }, + { + "backend": "flash_no_fma", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 37748736, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.12373999953269958, + "graph_min_ms": 0.1219599962234497, + "graph_max_ms": 0.1542415976524353, + "eager_wall_median_ms": 0.33643999995547347, + "eager_wall_trials_ms": [ + 0.33643999995547347, + 0.3219899999749032, + 0.3583933333478247 + ], + "extra_peak_bytes": 50331648 + }, + { + "backend": "upstream_complex", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 21163008, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.51009202003479, + "graph_min_ms": 0.5082543849945068, + "graph_max_ms": 0.5431327819824219, + "eager_wall_median_ms": 0.526480000007723, + "eager_wall_trials_ms": [ + 0.5092500000197712, + 0.5466633333223095, + 0.526480000007723 + ], + "extra_peak_bytes": 211630080 + }, + { + "backend": "upstream_complex", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 21163008, + 3072, + 128, + 1 + ], + "graph_median_ms": 1.0505383968353272, + "graph_min_ms": 1.0236895561218262, + "graph_max_ms": 1.129139232635498, + "eager_wall_median_ms": 1.044759999967937, + "eager_wall_trials_ms": [ + 1.044759999967937, + 1.0434166666527744, + 1.0474333333453008 + ], + "extra_peak_bytes": 215157248 + }, + { + "backend": "cuda", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 21163008, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.11203439831733704, + "graph_min_ms": 0.1114351987838745, + "graph_max_ms": 0.14540959596633912, + "eager_wall_median_ms": 0.11692666666931473, + "eager_wall_trials_ms": [ + 0.11580999998841435, + 0.11692666666931473, + 0.12043333335896023 + ], + "extra_peak_bytes": 42326016 + }, + { + "backend": "cuda", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 21163008, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.22477760314941406, + "graph_min_ms": 0.22440481185913086, + "graph_max_ms": 0.24935200214385986, + "eager_wall_median_ms": 0.23845333331943644, + "eager_wall_trials_ms": [ + 0.23845333331943644, + 0.2526333333056148, + 0.230336666633472 + ], + "extra_peak_bytes": 84652032 + }, + { + "backend": "triton", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 21163008, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.10685279965400696, + "graph_min_ms": 0.10539519786834717, + "graph_max_ms": 0.12008960247039795, + "eager_wall_median_ms": 0.11027333330275724, + "eager_wall_trials_ms": [ + 0.11027333330275724, + 0.10863333330538201, + 0.11599999997997656 + ], + "extra_peak_bytes": 42326016 + }, + { + "backend": "triton", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 21163008, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.21339280605316163, + "graph_min_ms": 0.2126607894897461, + "graph_max_ms": 0.23967199325561522, + "eager_wall_median_ms": 0.27515000001585577, + "eager_wall_trials_ms": [ + 0.2870199999961187, + 0.2669133333256468, + 0.27515000001585577 + ], + "extra_peak_bytes": 84652032 + }, + { + "backend": "flash_default", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 21163008, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.10692960023880005, + "graph_min_ms": 0.10526080131530761, + "graph_max_ms": 0.12183040380477905, + "eager_wall_median_ms": 0.10955999999472017, + "eager_wall_trials_ms": [ + 0.1204166666866513, + 0.10955999999472017, + 0.10893666670502473 + ], + "extra_peak_bytes": 42326016 + }, + { + "backend": "flash_default", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 21163008, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.21346559524536132, + "graph_min_ms": 0.21225121021270751, + "graph_max_ms": 0.2430032014846802, + "eager_wall_median_ms": 0.338726666692916, + "eager_wall_trials_ms": [ + 0.33888666663794237, + 0.3213533333716138, + 0.338726666692916 + ], + "extra_peak_bytes": 84652032 + }, + { + "backend": "flash_no_fma", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 21163008, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.10693199634552003, + "graph_min_ms": 0.10515999794006348, + "graph_max_ms": 0.12142080068588257, + "eager_wall_median_ms": 0.1100600000427221, + "eager_wall_trials_ms": [ + 0.11902333335456206, + 0.1100600000427221, + 0.10893666664439176 + ], + "extra_peak_bytes": 42326016 + }, + { + "backend": "flash_no_fma", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 21163008, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.21350640058517456, + "graph_min_ms": 0.2123823881149292, + "graph_max_ms": 0.23954401016235352, + "eager_wall_median_ms": 0.3223733333167426, + "eager_wall_trials_ms": [ + 0.29901333333934116, + 0.3223733333167426, + 0.326143333343983 + ], + "extra_peak_bytes": 84652032 + }, + { + "backend": "upstream_complex", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 63489024, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.5111671924591064, + "graph_min_ms": 0.5095888137817383, + "graph_max_ms": 0.5389599800109863, + "eager_wall_median_ms": 0.5120633333717706, + "eager_wall_trials_ms": [ + 0.5099033333560025, + 0.532273333374178, + 0.5120633333717706 + ], + "extra_peak_bytes": 211630080 + }, + { + "backend": "upstream_complex", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 63489024, + 9216, + 128, + 1 + ], + "graph_median_ms": 1.0546072006225586, + "graph_min_ms": 1.0272704124450684, + "graph_max_ms": 1.1091728210449219, + "eager_wall_median_ms": 1.0453999999905743, + "eager_wall_trials_ms": [ + 1.0565066666458733, + 1.0453999999905743, + 1.042836666662576 + ], + "extra_peak_bytes": 215157248 + }, + { + "backend": "cuda", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 63489024, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.11280480027198792, + "graph_min_ms": 0.11215200424194335, + "graph_max_ms": 0.13529759645462036, + "eager_wall_median_ms": 0.12308333334658528, + "eager_wall_trials_ms": [ + 0.11563333331044608, + 0.12308333334658528, + 0.12365000002318993 + ], + "extra_peak_bytes": 42326016 + }, + { + "backend": "cuda", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 63489024, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.2251416087150574, + "graph_min_ms": 0.22450559139251708, + "graph_max_ms": 0.2644239902496338, + "eager_wall_median_ms": 0.2328933333653064, + "eager_wall_trials_ms": [ + 0.23086333330866182, + 0.2328933333653064, + 0.25025666667109664 + ], + "extra_peak_bytes": 84652032 + }, + { + "backend": "triton", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 63489024, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.1075808048248291, + "graph_min_ms": 0.1061360001564026, + "graph_max_ms": 0.12249759435653687, + "eager_wall_median_ms": 0.10984000000462402, + "eager_wall_trials_ms": [ + 0.1298533332980393, + 0.10860666661756113, + 0.10984000000462402 + ], + "extra_peak_bytes": 42326016 + }, + { + "backend": "triton", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 63489024, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.21414560079574585, + "graph_min_ms": 0.21337599754333497, + "graph_max_ms": 0.23953759670257568, + "eager_wall_median_ms": 0.26501666670810664, + "eager_wall_trials_ms": [ + 0.2692166666747653, + 0.23041333333821967, + 0.26501666670810664 + ], + "extra_peak_bytes": 84652032 + }, + { + "backend": "flash_default", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 63489024, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.10766559839248657, + "graph_min_ms": 0.10606240034103394, + "graph_max_ms": 0.13348640203475953, + "eager_wall_median_ms": 0.11105333330003002, + "eager_wall_trials_ms": [ + 0.11928999999023897, + 0.11105333330003002, + 0.11004666666849516 + ], + "extra_peak_bytes": 42326016 + }, + { + "backend": "flash_default", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 63489024, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.21448720693588258, + "graph_min_ms": 0.21337599754333497, + "graph_max_ms": 0.25474560260772705, + "eager_wall_median_ms": 0.30118666666870314, + "eager_wall_trials_ms": [ + 0.32142000000021653, + 0.30118666666870314, + 0.2931233333583805 + ], + "extra_peak_bytes": 84652032 + }, + { + "backend": "flash_no_fma", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 63489024, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.1074567973613739, + "graph_min_ms": 0.10611200332641602, + "graph_max_ms": 0.1223423957824707, + "eager_wall_median_ms": 0.11078333336627111, + "eager_wall_trials_ms": [ + 0.12015666664713838, + 0.10989999997642978, + 0.11078333336627111 + ], + "extra_peak_bytes": 42326016 + }, + { + "backend": "flash_no_fma", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 63489024, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.21430879831314087, + "graph_min_ms": 0.2134079933166504, + "graph_max_ms": 0.23155200481414795, + "eager_wall_median_ms": 0.3170500000123866, + "eager_wall_trials_ms": [ + 0.3170500000123866, + 0.32710333331730607, + 0.31006333332091646 + ], + "extra_peak_bytes": 84652032 + }, + { + "backend": "upstream_complex", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 18530304, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.43411200046539306, + "graph_min_ms": 0.4321375846862793, + "graph_max_ms": 0.46126718521118165, + "eager_wall_median_ms": 0.4498766666680846, + "eager_wall_trials_ms": [ + 0.43220333336648764, + 0.4498766666680846, + 0.46697666663627996 + ], + "extra_peak_bytes": 185303040 + }, + { + "backend": "upstream_complex", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 18530304, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.8935583591461181, + "graph_min_ms": 0.870751953125, + "graph_max_ms": 0.9448016166687012, + "eager_wall_median_ms": 0.8892600000156866, + "eager_wall_trials_ms": [ + 0.8870666666553007, + 0.8892600000156866, + 0.8902699999453034 + ], + "extra_peak_bytes": 188391424 + }, + { + "backend": "cuda", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 18530304, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.09808879792690278, + "graph_min_ms": 0.09759039878845215, + "graph_max_ms": 0.11292320489883423, + "eager_wall_median_ms": 0.11404999998679462, + "eager_wall_trials_ms": [ + 0.11404999998679462, + 0.11419666667886001, + 0.102829999984048 + ], + "extra_peak_bytes": 37060608 + }, + { + "backend": "cuda", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 18530304, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.19646480083465576, + "graph_min_ms": 0.19599839448928832, + "graph_max_ms": 0.21091840267181397, + "eager_wall_median_ms": 0.20405333331533862, + "eager_wall_trials_ms": [ + 0.20405333331533862, + 0.20323666661473302, + 0.2081700000417186 + ], + "extra_peak_bytes": 74809344 + }, + { + "backend": "triton", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 18530304, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.09356399774551391, + "graph_min_ms": 0.09257439970970154, + "graph_max_ms": 0.10800800323486329, + "eager_wall_median_ms": 0.10138999999374694, + "eager_wall_trials_ms": [ + 0.10716000003109609, + 0.10138999999374694, + 0.10003000000627556 + ], + "extra_peak_bytes": 37060608 + }, + { + "backend": "triton", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 18530304, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.18708080053329468, + "graph_min_ms": 0.1862496018409729, + "graph_max_ms": 0.21909279823303224, + "eager_wall_median_ms": 0.26604333331003244, + "eager_wall_trials_ms": [ + 0.30756000002535683, + 0.2481433333135404, + 0.26604333331003244 + ], + "extra_peak_bytes": 74809344 + }, + { + "backend": "flash_default", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 18530304, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.09362079799175263, + "graph_min_ms": 0.09246399998664856, + "graph_max_ms": 0.10841599702835084, + "eager_wall_median_ms": 0.10133000002194119, + "eager_wall_trials_ms": [ + 0.14010333331195093, + 0.10133000002194119, + 0.10039666667580605 + ], + "extra_peak_bytes": 37060608 + }, + { + "backend": "flash_default", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 18530304, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.18697920441627502, + "graph_min_ms": 0.18585760593414308, + "graph_max_ms": 0.21307199001312255, + "eager_wall_median_ms": 0.3559766667118917, + "eager_wall_trials_ms": [ + 0.3559766667118917, + 0.3666866666511244, + 0.35564333332634607 + ], + "extra_peak_bytes": 74809344 + }, + { + "backend": "flash_no_fma", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 18530304, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.09362399876117705, + "graph_min_ms": 0.09233919978141784, + "graph_max_ms": 0.10795680284500123, + "eager_wall_median_ms": 0.1356800000091122, + "eager_wall_trials_ms": [ + 0.14346666666824603, + 0.1356800000091122, + 0.1269966666465431 + ], + "extra_peak_bytes": 37060608 + }, + { + "backend": "flash_no_fma", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 18530304, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.18709120154380798, + "graph_min_ms": 0.18642239570617675, + "graph_max_ms": 0.22162079811096191, + "eager_wall_median_ms": 0.3122899999652873, + "eager_wall_trials_ms": [ + 0.29407333331619157, + 0.3122899999652873, + 0.36742666664698237 + ], + "extra_peak_bytes": 74809344 + }, + { + "backend": "upstream_complex", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 55590912, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.43668241500854493, + "graph_min_ms": 0.43459200859069824, + "graph_max_ms": 0.4815072059631348, + "eager_wall_median_ms": 0.4541133333380761, + "eager_wall_trials_ms": [ + 0.4636933333434475, + 0.43558999999125564, + 0.4541133333380761 + ], + "extra_peak_bytes": 185303040 + }, + { + "backend": "upstream_complex", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 55590912, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.8956063747406007, + "graph_min_ms": 0.8733648300170899, + "graph_max_ms": 0.9540863990783691, + "eager_wall_median_ms": 0.8916500000244317, + "eager_wall_trials_ms": [ + 0.8923599999737538, + 0.8916500000244317, + 0.8882866666681366 + ], + "extra_peak_bytes": 188391424 + }, + { + "backend": "cuda", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 55590912, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.09595520198345184, + "graph_min_ms": 0.09553279876708984, + "graph_max_ms": 0.11051679849624634, + "eager_wall_median_ms": 0.1001133333678202, + "eager_wall_trials_ms": [ + 0.1001133333678202, + 0.09971333335367187, + 0.10664000001270324 + ], + "extra_peak_bytes": 37060608 + }, + { + "backend": "cuda", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 55590912, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.19663520455360411, + "graph_min_ms": 0.19614880084991454, + "graph_max_ms": 0.22351360321044922, + "eager_wall_median_ms": 0.20444999997077198, + "eager_wall_trials_ms": [ + 0.22424333334735516, + 0.20414666669239523, + 0.20444999997077198 + ], + "extra_peak_bytes": 74809344 + }, + { + "backend": "triton", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 55590912, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.0934136003255844, + "graph_min_ms": 0.0916208028793335, + "graph_max_ms": 0.10537919998168946, + "eager_wall_median_ms": 0.09456333336856915, + "eager_wall_trials_ms": [ + 0.09251333334153362, + 0.09570666667665743, + 0.09456333336856915 + ], + "extra_peak_bytes": 37060608 + }, + { + "backend": "triton", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 55590912, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.1877624034881592, + "graph_min_ms": 0.18668639659881592, + "graph_max_ms": 0.21486239433288573, + "eager_wall_median_ms": 0.24588000002646973, + "eager_wall_trials_ms": [ + 0.26271666671770316, + 0.24588000002646973, + 0.23198333334827717 + ], + "extra_peak_bytes": 74809344 + }, + { + "backend": "flash_default", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 55590912, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.09341759979724884, + "graph_min_ms": 0.0921504020690918, + "graph_max_ms": 0.11279360055923462, + "eager_wall_median_ms": 0.14909333337224476, + "eager_wall_trials_ms": [ + 0.15717333335487638, + 0.12570666664638944, + 0.14909333337224476 + ], + "extra_peak_bytes": 37060608 + }, + { + "backend": "flash_default", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 55590912, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.1880079984664917, + "graph_min_ms": 0.18690240383148193, + "graph_max_ms": 0.2143280029296875, + "eager_wall_median_ms": 0.3311766666835562, + "eager_wall_trials_ms": [ + 0.3571433333490859, + 0.3311766666835562, + 0.32030999997611315 + ], + "extra_peak_bytes": 74809344 + }, + { + "backend": "flash_no_fma", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 55590912, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.0934448003768921, + "graph_min_ms": 0.09198240041732789, + "graph_max_ms": 0.10836479663848878, + "eager_wall_median_ms": 0.132580000051045, + "eager_wall_trials_ms": [ + 0.132580000051045, + 0.14571666670235572, + 0.12959000002107737 + ], + "extra_peak_bytes": 37060608 + }, + { + "backend": "flash_no_fma", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 55590912, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.18786399960517886, + "graph_min_ms": 0.18675040006637572, + "graph_max_ms": 0.2141263961791992, + "eager_wall_median_ms": 0.34321999998307245, + "eager_wall_trials_ms": [ + 0.36083333331286366, + 0.3284899999925983, + 0.34321999998307245 + ], + "extra_peak_bytes": 74809344 + } + ] +} diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/installed-wheel.json b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/installed-wheel.json new file mode 100644 index 000000000..468ce07af --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/installed-wheel.json @@ -0,0 +1,37 @@ +{ + "package": "D:\\dandan123\\dandan\\issue\\investigation-rl-kernel\\RL-Kernel\\.venv\\Lib\\site-packages\\rl_engine", + "companion_library": "D:\\dandan123\\dandan\\issue\\investigation-rl-kernel\\RL-Kernel\\.venv\\Lib\\site-packages\\rl_engine\\_rope_cuda.cp313-win_amd64.pyd", + "native_dispatcher_registered": true, + "torch": "2.9.1+cu128", + "gpu": "NVIDIA GeForce RTX 5080", + "license_present": true, + "packed_qkv_stride": [ + 17280, + 1152, + 128, + 1 + ], + "checks": { + "cuda_eager": "passed", + "cuda_compiled": "passed", + "cuda_trace": { + "contract": "qwenimage-interleaved-rope-fp32-v1", + "actual_backend": "rlkernel.cuda.qwenimage_multi_axis_rope.v1", + "kernel_id": "torch.ops.rl_kernel_native.multi_axis_rope", + "layout": "BSHD, adjacent channel pairs", + "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", + "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" + }, + "triton_eager": "passed", + "triton_compiled": "passed", + "triton_trace": { + "contract": "qwenimage-interleaved-rope-fp32-v1", + "actual_backend": "rlkernel.triton.qwenimage_multi_axis_rope.v1", + "kernel_id": "rl_engine.backends.shared.triton.rope._flash_attn_rotary.rotary_kernel", + "layout": "BSHD, adjacent channel pairs", + "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", + "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" + } + }, + "passed": true +} diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/installed-wheel.txt b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/installed-wheel.txt new file mode 100644 index 000000000..01c1dd377 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/installed-wheel.txt @@ -0,0 +1,38 @@ +WARNING 10-28 17:08:06 [RL-Kernel]: Core binary extension (_C) unavailable: cannot import name '_C' from 'rl_engine' (D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\.venv\Lib\site-packages\rl_engine\__init__.py). Falling back to native code. +{ + "package": "D:\\dandan123\\dandan\\issue\\investigation-rl-kernel\\RL-Kernel\\.venv\\Lib\\site-packages\\rl_engine", + "companion_library": "D:\\dandan123\\dandan\\issue\\investigation-rl-kernel\\RL-Kernel\\.venv\\Lib\\site-packages\\rl_engine\\_rope_cuda.cp313-win_amd64.pyd", + "native_dispatcher_registered": true, + "torch": "2.9.1+cu128", + "gpu": "NVIDIA GeForce RTX 5080", + "license_present": true, + "packed_qkv_stride": [ + 17280, + 1152, + 128, + 1 + ], + "checks": { + "cuda_eager": "passed", + "cuda_compiled": "passed", + "cuda_trace": { + "contract": "qwenimage-interleaved-rope-fp32-v1", + "actual_backend": "rlkernel.cuda.qwenimage_multi_axis_rope.v1", + "kernel_id": "torch.ops.rl_kernel_native.multi_axis_rope", + "layout": "BSHD, adjacent channel pairs", + "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", + "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" + }, + "triton_eager": "passed", + "triton_compiled": "passed", + "triton_trace": { + "contract": "qwenimage-interleaved-rope-fp32-v1", + "actual_backend": "rlkernel.triton.qwenimage_multi_axis_rope.v1", + "kernel_id": "rl_engine.backends.shared.triton.rope._flash_attn_rotary.rotary_kernel", + "layout": "BSHD, adjacent channel pairs", + "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", + "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" + } + }, + "passed": true +} diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/layout.txt b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/layout.txt new file mode 100644 index 000000000..0086d019b --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/layout.txt @@ -0,0 +1,138 @@ +.FF.............s............F.FF.... [100%] +================================== FAILURES =================================== +________________ test_cuda_build_keeps_existing_ipc_collective ________________ + +monkeypatch = <_pytest.monkeypatch.MonkeyPatch object at 0x000001D4B0E668B0> + + def test_cuda_build_keeps_existing_ipc_collective(monkeypatch) -> None: + extension = _load_extension_config(monkeypatch, hip=None) + + assert "csrc/cuda/collectives/deterministic_collective.cu" in extension["sources"] + assert "-DKERNEL_ALIGN_WITH_CUDA" in extension["extra_compile_args"]["cxx"] + assert "-DKERNEL_ALIGN_WITH_ROCM" not in extension["extra_compile_args"]["cxx"] +> assert "-lcuda" in extension["extra_link_args"] +E AssertionError: assert '-lcuda' in ['-Wl,-rpath,$ORIGIN/../torch/lib'] + +tests\build\test_build_platform_collectives.py:61: AssertionError +____________ test_package_data_json_links_resolve_to_real_content _____________ + +tmp_path = WindowsPath('C:/Users/123/AppData/Local/Temp/pytest-of-123/pytest-3/test_package_data_json_links_r0') + + def test_package_data_json_links_resolve_to_real_content(tmp_path) -> None: + target = tmp_path / "configs" / "profile.json" + target.parent.mkdir() + target.write_text('{"name": "profile"}\n', encoding="utf-8") + source = tmp_path / "rl_engine" / "profile.json" + source.parent.mkdir() +> source.symlink_to("../configs/profile.json") + +tests\build\test_build_platform_collectives.py:70: +_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ + +self = WindowsPath('C:/Users/123/AppData/Local/Temp/pytest-of-123/pytest-3/test_package_data_json_links_r0/rl_engine/profile.json') +target = '../configs/profile.json', target_is_directory = False + + def symlink_to(self, target, target_is_directory=False): + """ + Make this path a symlink pointing to the target path. + Note the order of arguments (link, target) is the reverse of os.symlink. + """ +> os.symlink(target, self, target_is_directory) +E OSError: [WinError 1314] �ͻ���û���������Ȩ��: '../configs/profile.json' -> 'C:\\Users\\123\\AppData\\Local\\Temp\\pytest-of-123\\pytest-3\\test_package_data_json_links_r0\\rl_engine\\profile.json' + +D:\anaconda\Lib\pathlib\_local.py:789: OSError +_______ test_recorded_algorithm_sources_resolve_to_real_implementations _______ + + def test_recorded_algorithm_sources_resolve_to_real_implementations(): + from rl_engine.config.workload import load_manifest + + moves = json.loads((ROOT / "tools/migration/layout.json").read_text())["moves"] + for case in load_manifest().representative_cases: + source, symbol = case["provenance_evidence"]["algorithm_source"].rsplit(":", 1) + implementation = ROOT / moves.get(source, source) +> text = implementation.read_text() + ^^^^^^^^^^^^^^^^^^^^^^^^^^ + +tests\runtime\test_layout_compatibility.py:87: +_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ +D:\anaconda\Lib\pathlib\_local.py:546: in read_text + return PathBase.read_text(self, encoding, errors, newline) + ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ +_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ + +self = WindowsPath('D:/dandan123/dandan/issue/investigation-rl-kernel/RL-Kernel/csrc/cuda/attention/deterministic_attention.cu') +encoding = 'locale', errors = None, newline = None + + def read_text(self, encoding=None, errors=None, newline=None): + """ + Open the file in text mode, read it, and close the file. + """ + with self.open(mode='r', encoding=encoding, errors=errors, newline=newline) as f: +> return f.read() + ^^^^^^^^ +E UnicodeDecodeError: 'gbk' codec can't decode byte 0x92 in position 173: illegal multibyte sequence + +D:\anaconda\Lib\pathlib\_abc.py:633: UnicodeDecodeError +_ test_packaged_contract_bytes_match_main[rl_engine.config-workloads/qwen3_8b.json-80395644b5e84ed1a76729ffc6f78eb11f41b9ab42fb36e2d00f19635d3df0f6] _ + +package = 'rl_engine.config', resource = 'workloads/qwen3_8b.json' +digest = '80395644b5e84ed1a76729ffc6f78eb11f41b9ab42fb36e2d00f19635d3df0f6' + + @pytest.mark.parametrize( + ("package", "resource", "digest"), + [ + ( + "rl_engine.config", + "workloads/qwen3_8b.json", + "80395644b5e84ed1a76729ffc6f78eb11f41b9ab42fb36e2d00f19635d3df0f6", + ), + ( + "rl_engine.contracts", + "profiles/precision/ws1.json", + "044167ec21e372017a2758dc09f805b368283c06b83f64061877628cbd101532", + ), + ], + ) + def test_packaged_contract_bytes_match_main(package, resource, digest): +> assert hashlib.sha256(files(package).joinpath(resource).read_bytes()).hexdigest() == digest +E AssertionError: assert 'c97e966fb075...75e111e052d7d' == '80395644b5e8...f19635d3df0f6' +E +E - 80395644b5e84ed1a76729ffc6f78eb11f41b9ab42fb36e2d00f19635d3df0f6 +E + c97e966fb075dad8ce7dcd7a23a58ffb856e7474dfb3c50510675e111e052d7d + +tests\runtime\test_layout_compatibility.py:118: AssertionError +_ test_packaged_contract_bytes_match_main[rl_engine.contracts-profiles/precision/ws1.json-044167ec21e372017a2758dc09f805b368283c06b83f64061877628cbd101532] _ + +package = 'rl_engine.contracts', resource = 'profiles/precision/ws1.json' +digest = '044167ec21e372017a2758dc09f805b368283c06b83f64061877628cbd101532' + + @pytest.mark.parametrize( + ("package", "resource", "digest"), + [ + ( + "rl_engine.config", + "workloads/qwen3_8b.json", + "80395644b5e84ed1a76729ffc6f78eb11f41b9ab42fb36e2d00f19635d3df0f6", + ), + ( + "rl_engine.contracts", + "profiles/precision/ws1.json", + "044167ec21e372017a2758dc09f805b368283c06b83f64061877628cbd101532", + ), + ], + ) + def test_packaged_contract_bytes_match_main(package, resource, digest): +> assert hashlib.sha256(files(package).joinpath(resource).read_bytes()).hexdigest() == digest +E AssertionError: assert '2020c02e7af4...313d6752cd34e' == '044167ec21e3...7628cbd101532' +E +E - 044167ec21e372017a2758dc09f805b368283c06b83f64061877628cbd101532 +E + 2020c02e7af4e2ca4163e1ecfb9c491baef07b73deba39b0f44313d6752cd34e + +tests\runtime\test_layout_compatibility.py:118: AssertionError +=========================== short test summary info =========================== +FAILED tests/build/test_build_platform_collectives.py::test_cuda_build_keeps_existing_ipc_collective +FAILED tests/build/test_build_platform_collectives.py::test_package_data_json_links_resolve_to_real_content +FAILED tests/runtime/test_layout_compatibility.py::test_recorded_algorithm_sources_resolve_to_real_implementations +FAILED tests/runtime/test_layout_compatibility.py::test_packaged_contract_bytes_match_main[rl_engine.config-workloads/qwen3_8b.json-80395644b5e84ed1a76729ffc6f78eb11f41b9ab42fb36e2d00f19635d3df0f6] +FAILED tests/runtime/test_layout_compatibility.py::test_packaged_contract_bytes_match_main[rl_engine.contracts-profiles/precision/ws1.json-044167ec21e372017a2758dc09f805b368283c06b83f64061877628cbd101532] +5 failed, 31 passed, 1 skipped in 5.67s diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/manifest.json b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/manifest.json new file mode 100644 index 000000000..190cf7834 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/manifest.json @@ -0,0 +1,65 @@ +{ + "recorded_at_utc": "2026-10-28T09:12:00.361759+00:00", + "base_commit": "16b81175874ecf11d5ff05c00de3b69e9457591b", + "source_sha256": { + "MANIFEST.in": "706ec113547f7012cfa6d15d062817d5d2d9e0c0e18e0fb706eca5ba2335bcc5", + "benchmarks/operators/rope/benchmark_multi_axis_rope.py": "6961a55803fbd4e566bd060558fb59e658e7f1f704e61ca97f912e5d16b55884", + "build_tools/extensions.py": "cf07cd49cba0ba7b4da459c26546380b7b288232fcc032e7641863e09ae3014a", + "csrc/bindings/rope_module.cpp": "c69eb3f200de306ca5885fd7c4c7ccf49c822ed1f2f1f639223ec36f0c51b524", + "csrc/cuda/rope/multi_axis_rope.cu": "f5db5cf5d4962053194455d9cde8cf91cf4e824bad7abf72e8410a404bcb0e44", + "docs/.nav.yml": "814a6d60839465e8f9405a9a26306dbc203b0fcb6aa36db264fd7651624d7ef3", + "docs/operators/README.md": "ab926bf5a3fb719e4a57cb65f0784a52019f6ed4bd49c0b5f1b15a1765fd6d7a", + "docs/operators/multi-axis-rope.md": "79f94810ee15ead5f1ce24715456f65f12b383728f2a3571104614b1a33b47d8", + "pyproject.toml": "a132159ae2a2c077d9fe001194fa2b9ab382acbcd58a9eccdea9387286bfb258", + "rl_engine/backends/cuda/rope/multi_axis_rope.py": "d779d9b60ce5d1d6651925f723049930d898024fdf0548ce29a2019b84753d65", + "rl_engine/backends/shared/triton/rope/FLASH_ATTENTION_LICENSE": "8c9ccb96c065e706135b6cbad279b721da6156e51f3a5f27c6b3329af9416d73", + "rl_engine/backends/shared/triton/rope/_flash_attn_rotary.py": "5c6984da351e5228f7549629d4ff53e5772b9a6ef2a01ba2a38493e769b15ab5", + "rl_engine/backends/shared/triton/rope/multi_axis_rope.py": "5b989bd63efd77913267b48654d10204bc563768c6c5d4e1cde9fb55e1354547", + "rl_engine/ops/rope/_registration.py": "add0375d1639444908091d01e8d29781f7a975beffec64b77bfaa81a5b9a0ae3", + "rl_engine/ops/rope/multi_axis_rope.py": "07bcff1160f22c1ff6730be731206b2f4a153e2fada26afc93fc884e780568cb", + "rl_engine/reference/rope/multi_axis_rope.py": "c9a94730047f51659f7aa0a443e88d6ce7d03d00d6813a9aed548cafbf6347bf", + "rl_engine/runtime/registry.py": "9f9c69bb6c2e941b057a66ac2d388be021cbb4b55cb3c602976844aabc6106e1", + "rl_engine/validation/operators/gradient_adapters.py": "5ee2e9ad619a4cad6e81d8f4a6ccf81ecf45e63a7764d0fcb4c204206b37c16b", + "rl_engine/validation/operators/operator_inputs.py": "55c8d9e7cbde2070662dd1d0ebe9888aef106b45fb05902057aec032d29a9836", + "rl_engine/validation/operators/operator_specs.py": "cef2e5bed8075509d5d5a538b378477b09f1adfe5fd01829156357cefa99f983", + "rl_engine/validation/operators/qwen_image_rope.py": "c3ddfc732f32de3c633de16064810d67202254ea78fd8107d558b471bb15a37f", + "tests/build/test_multi_axis_rope_build.py": "90c3740b2d4946acc6206d952622c19afaa0089cca904919088563427555b956", + "tests/ops/rope/test_multi_axis_rope.py": "1a0950b09eec407e8ffd490b40dda4118f8d9b876f4db4af8aac648af1a76687", + "tests/validation/operators/test_operator_inputs.py": "36365f662194f158276b3c07f6447471a0cf17253ec420ab63c8b9149c511b49", + "tools/validation/operators/check_multi_axis_rope_install.py": "ff8198ff55779ea2ede88d4ecfe07caa627631b4b08bce38139ccfa147abf538", + "tools/validation/operators/check_qwenimage_rope_upstream.py": "cc368d5bebf0f0acb21b7455348592d30442b6a8d349de0f49d5a09c5f20ce4e", + "tools/validation/operators/validate_multi_axis_rope.py": "def4747da9ea9194c761bef24479017ac74ac10de46a3fbe1a9528ec6ee07731" + }, + "evidence_sha256": { + "baseline-layout.txt": "b050d798a5e3a6c2d43c7c22e4e1140529ebf05ebab982df17f501f5b69cff12", + "benchmark-bfloat16.json": "0679a422e8129a8a36ada938b024794334f784cdea28013574f99c70673e13a0", + "installed-wheel.json": "60f6a1a48d35fa25a639a1ff168ce0c37659e0258737422cafda3741d8ef0200", + "installed-wheel.txt": "2be6479a8e30938410095dbcc7724ad114cb61cb35ef86c7f86b8d507a3193f6", + "layout.txt": "83169d7995bdadc14744a0ed0fff69b334f2ed6ff0b92cda005a9cf0f733116d", + "memcheck.txt": "f6f05d4e5b772909f29baf2a1d6be04b4456660ca024e5067c3f64c7efa4b81f", + "operator-cuda.json": "1a231d26f2411a38cffe70fdb59b388227fc0e7da9625370b165d7d590e1e011", + "operator-triton.json": "af2be1f4df1bf9d7afb4802ca02c81826e64acacae891436f83ad6b3342109f5", + "package-build.txt": "34b11d626dc4988f3d862e4122599a0b1a949b1d799436981dd74f3ee5c1b07e", + "package.json": "e267d82707cdeb0c53d120928cd32a544a44190e8e41c43f3b6c1d4668275a7e", + "precommit.txt": "04f9db16780e0ad0b41e6450d362e6f0a7bdb07128596c8ea13f8ea8bb2cc674", + "quality.txt": "d4728355fe488102ca644f510795c474b54f8d42bb92ddfda34557ba521c6afd", + "README.md": "aa2cabc1e8423f4d737d0db90786142c4a919041dbf97335b17991475a4cbaca", + "regression.txt": "5d058ac93516d9ce1f486fbd4f35b84a254a32d8ddf26eedf41cf2388153e004", + "tests.txt": "f6af2230f6a0b5d993cb6bca2724a82f772c6b33df456edce6e728996f2a76d5", + "upstream.json": "8c087d60013f0266303c6d60237d87e5f1ff67b03e8e8c1900de4b20349e2e93", + "ws1-cuda.json": "531124223ae5fe43eca890fb1ccf80e09f33affe34bab132c7c17742081d85a3", + "ws1-cuda.txt": "b4cf09ef5165e535dba9ed19eeec08f15eccaf86b549bb1a253fd08e87be35ab", + "ws1-triton.json": "29eedee8e50b8750d0ce9be95deb39fc6aa9ab47b0867fe20ce2e7ad5cb47548", + "ws1-triton.txt": "a1daefba35f7c1a06d222750ad1ce9c4bdc97cc03410969e7b4f4ce011a6770a" + }, + "layout_failures_equal_clean_base": true, + "layout_failure_ids": [ + "tests/build/test_build_platform_collectives.py::test_cuda_build_keeps_existing_ipc_collective", + "tests/build/test_build_platform_collectives.py::test_package_data_json_links_resolve_to_real_content", + "tests/runtime/test_layout_compatibility.py::test_recorded_algorithm_sources_resolve_to_real_implementations", + "tests/runtime/test_layout_compatibility.py::test_packaged_contract_bytes_match_main[rl_engine.config-workloads/qwen3_8b.json-80395644b5e84ed1a76729ffc6f78eb11f41b9ab42fb36e2d00f19635d3df0f6]", + "tests/runtime/test_layout_compatibility.py::test_packaged_contract_bytes_match_main[rl_engine.contracts-profiles/precision/ws1.json-044167ec21e372017a2758dc09f805b368283c06b83f64061877628cbd101532]" + ], + "package_source_and_license_checks": "passed", + "hash_mode": "SHA256 of recorded working-tree bytes; upstream baselines separately normalize newlines" +} diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/memcheck.txt b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/memcheck.txt new file mode 100644 index 000000000..dd59e40b0 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/memcheck.txt @@ -0,0 +1,8 @@ +========= COMPUTE-SANITIZER +WARNING 10-28 16:46:09 [RL-Kernel]: Core binary extension (_C) unavailable: cannot import name '_C' from 'rl_engine' (D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\rl_engine\__init__.py). Falling back to native code. +W1028 16:46:10.075000 16608 site-packages\torch\utils\cpp_extension.py:480] Error checking compiler version for cl: 'cp1' codec can't decode bytes in position 0--1: No mapping for the Unicode character exists in the target code page. +ninja: no work to do. +Focused CUDA source build; torch=2.9.1+cu128, GPU=NVIDIA GeForce RTX 5080 +.................................. [100%] +34 passed, 48 deselected in 3.96s +========= ERROR SUMMARY: 0 errors diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/operator-cuda.json b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/operator-cuda.json new file mode 100644 index 000000000..03d59bb4b --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/operator-cuda.json @@ -0,0 +1,70 @@ +{ + "suite_name": "multi_axis_rope", + "total_candidates": 1, + "passed_candidates": 1, + "pass_rate": 1.0, + "passed": true, + "candidates": [ + { + "candidate_name": "cuda-multi_axis_rope", + "backend": "cuda", + "total_outputs": 2, + "passed_outputs": 2, + "pass_rate": 1.0, + "passed": true, + "cases": [ + { + "case_name": "multi_axis_rope-torch.bfloat16-2x7x3x128", + "dtype": "torch.bfloat16", + "op_class": "elementwise", + "passed": true, + "outputs": [ + { + "output_index": 0, + "shape": [ + 2, + 7, + 3, + 128 + ], + "candidate_dtype": "torch.bfloat16", + "gold_dtype": "torch.float32", + "atol": 0.02, + "rtol": 0.016, + "max_abs_error": 0.007704973220825195, + "mean_abs_error": 0.0006018931162543595, + "max_rel_error": 0.003828622167930007, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference", + "message": "" + }, + { + "output_index": 1, + "shape": [ + 2, + 7, + 3, + 128 + ], + "candidate_dtype": "torch.bfloat16", + "gold_dtype": "torch.bfloat16", + "atol": 0.02, + "rtol": 0.016, + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference", + "message": "gradient:x" + } + ] + } + ], + "backend_provenance": null + } + ] +} diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/operator-triton.json b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/operator-triton.json new file mode 100644 index 000000000..6f086edac --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/operator-triton.json @@ -0,0 +1,70 @@ +{ + "suite_name": "multi_axis_rope", + "total_candidates": 1, + "passed_candidates": 1, + "pass_rate": 1.0, + "passed": true, + "candidates": [ + { + "candidate_name": "triton-multi_axis_rope", + "backend": "triton", + "total_outputs": 2, + "passed_outputs": 2, + "pass_rate": 1.0, + "passed": true, + "cases": [ + { + "case_name": "multi_axis_rope-torch.bfloat16-2x7x3x128", + "dtype": "torch.bfloat16", + "op_class": "elementwise", + "passed": true, + "outputs": [ + { + "output_index": 0, + "shape": [ + 2, + 7, + 3, + 128 + ], + "candidate_dtype": "torch.bfloat16", + "gold_dtype": "torch.float32", + "atol": 0.02, + "rtol": 0.016, + "max_abs_error": 0.007704973220825195, + "mean_abs_error": 0.0006018931162543595, + "max_rel_error": 0.003828622167930007, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference", + "message": "" + }, + { + "output_index": 1, + "shape": [ + 2, + 7, + 3, + 128 + ], + "candidate_dtype": "torch.bfloat16", + "gold_dtype": "torch.bfloat16", + "atol": 0.02, + "rtol": 0.016, + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference", + "message": "gradient:x" + } + ] + } + ], + "backend_provenance": null + } + ] +} diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/package-build.txt b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/package-build.txt new file mode 100644 index 000000000..d63d79d52 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/package-build.txt @@ -0,0 +1,40 @@ +Selected lines from the completed sdist/wheel build log. +Full local log SHA256: bfbcd342efaf8caa36adcd208acee385db1ba6fbc4326fe8e0325f9523f199b0 +running sdist +copying rl_engine\backends\shared\triton\rope\FLASH_ATTENTION_LICENSE -> rl_kernel-0.1.0\rl_engine\backends\shared\triton\rope +running bdist_wheel +copying rl_engine\backends\shared\triton\rope\FLASH_ATTENTION_LICENSE -> build\lib.win-amd64-cpython-313\rl_engine\backends\shared\triton\rope +copying rl_engine\backends\shared\triton\rope\FLASH_ATTENTION_LICENSE -> build\lib.win-amd64-cpython-313\rl_engine\backends\shared\triton\rope +running build_ext +ninja: no work to do. +"C:\Program Files (x86)\Microsoft Visual Studio\2022\BuildTools\VC\Tools\MSVC\14.44.35207\bin\HostX64\x64\link.exe" /nologo /INCREMENTAL:NO /LTCG /DLL /MANIFEST:EMBED,ID=2 /MANIFESTUAC:NO /LIBPATH:D:\dandan123\dandan\myvllm\.cuda-12.8\Library\lib /LIBPATH:D:\anaconda\Lib\site-packages\torch\lib /LIBPATH:D:\dandan123\dandan\myvllm\.cuda-12.8\Library\lib\x64 /LIBPATH:D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\.venv\libs /LIBPATH:D:\anaconda\libs /LIBPATH:D:\anaconda /LIBPATH:D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\.venv\PCbuild\amd64 "/LIBPATH:C:\Program Files (x86)\Microsoft Visual Studio\2022\BuildTools\VC\Tools\MSVC\14.44.35207\lib\x64" "/LIBPATH:C:\Program Files (x86)\Windows Kits\10\lib\10.0.26100.0\ucrt\x64" "/LIBPATH:C:\Program Files (x86)\Windows Kits\10\\lib\10.0.26100.0\\um\x64" c10.lib torch.lib torch_cpu.lib torch_python.lib cudart.lib c10_cuda.lib torch_cuda.lib /EXPORT:PyInit__rope_cuda D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\build\temp.win-amd64-cpython-313\Release\csrc\bindings\rope_module.obj D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\build\temp.win-amd64-cpython-313\Release\csrc\cuda\rope\multi_axis_rope.obj /OUT:build\lib.win-amd64-cpython-313\rl_engine\_rope_cuda.cp313-win_amd64.pyd /IMPLIB:D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\build\temp.win-amd64-cpython-313\Release\csrc\bindings\_rope_cuda.cp313-win_amd64.lib +copying build\lib.win-amd64-cpython-313\rl_engine\backends\shared\triton\rope\FLASH_ATTENTION_LICENSE -> build\bdist.win-amd64\wheel\.\rl_engine\backends\shared\triton\rope +copying build\lib.win-amd64-cpython-313\rl_engine\_rope_cuda.cp313-win_amd64.pyd -> build\bdist.win-amd64\wheel\.\rl_engine +adding 'rl_engine/_rope_cuda.cp313-win_amd64.pyd' +adding 'rl_engine/backends/shared/triton/rope/FLASH_ATTENTION_LICENSE' + +In-place build log: +D:\anaconda\Lib\site-packages\setuptools\config\_apply_pyprojecttoml.py:82: SetuptoolsDeprecationWarning: `project.license` as a TOML table is deprecated +!! + + ******************************************************************************** + Please use a simple string containing a SPDX expression for `project.license`. You can also use `project.license-files`. (Both options available on setuptools>=77.0.0). + + By 2027-Feb-18, you need to update your project and remove deprecated calls + or your builds will no longer be supported. + + See https://packaging.python.org/en/latest/guides/writing-pyproject-toml/#license for details. + ******************************************************************************** + +!! + corresp(dist, value, root_dir) +running build_ext +W1028 16:52:54.107000 9384 site-packages\torch\utils\cpp_extension.py:480] Error checking compiler version for cl: 'cp1' codec can't decode bytes in position 0--1: No mapping for the Unicode character exists in the target code page. +building 'rl_engine._rope_cuda' extension +W1028 16:52:54.150000 9384 site-packages\torch\utils\cpp_extension.py:480] Error checking compiler version for cl: 'cp1' codec can't decode bytes in position 0--1: No mapping for the Unicode character exists in the target code page. +ninja: no work to do. +"C:\Program Files (x86)\Microsoft Visual Studio\2022\BuildTools\VC\Tools\MSVC\14.44.35207\bin\HostX64\x64\link.exe" /nologo /INCREMENTAL:NO /LTCG /DLL /MANIFEST:EMBED,ID=2 /MANIFESTUAC:NO /LIBPATH:D:\dandan123\dandan\myvllm\.cuda-12.8\Library\lib /LIBPATH:D:\anaconda\Lib\site-packages\torch\lib /LIBPATH:D:\dandan123\dandan\myvllm\.cuda-12.8\Library\lib\x64 /LIBPATH:D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\.venv\libs /LIBPATH:D:\anaconda\libs /LIBPATH:D:\anaconda /LIBPATH:D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\.venv\PCbuild\amd64 "/LIBPATH:C:\Program Files (x86)\Microsoft Visual Studio\2022\BuildTools\VC\Tools\MSVC\14.44.35207\lib\x64" "/LIBPATH:C:\Program Files (x86)\Windows Kits\10\lib\10.0.26100.0\ucrt\x64" "/LIBPATH:C:\Program Files (x86)\Windows Kits\10\\lib\10.0.26100.0\\um\x64" c10.lib torch.lib torch_cpu.lib torch_python.lib cudart.lib c10_cuda.lib torch_cuda.lib /EXPORT:PyInit__rope_cuda D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\build\temp.win-amd64-cpython-313\Release\csrc\bindings\rope_module.obj D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\build\temp.win-amd64-cpython-313\Release\csrc\cuda\rope\multi_axis_rope.obj /OUT:build\lib.win-amd64-cpython-313\rl_engine\_rope_cuda.cp313-win_amd64.pyd /IMPLIB:D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\build\temp.win-amd64-cpython-313\Release\csrc\bindings\_rope_cuda.cp313-win_amd64.lib + 正在创建库 D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\build\temp.win-amd64-cpython-313\Release\csrc\bindings\_rope_cuda.cp313-win_amd64.lib 和对象 D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\build\temp.win-amd64-cpython-313\Release\csrc\bindings\_rope_cuda.cp313-win_amd64.exp +正在生成代码 +已完成代码的生成 +copying build\lib.win-amd64-cpython-313\rl_engine\_rope_cuda.cp313-win_amd64.pyd -> rl_engine diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/package.json b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/package.json new file mode 100644 index 000000000..51e102428 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/package.json @@ -0,0 +1,31 @@ +{ + "wheel": "rl_kernel-0.1.0-cp313-cp313-win_amd64.whl", + "wheel_sha256": "bc5418bbc4be779ec6d5b4fa99d54b3244f2570d60a29adc3e6a1c38be25a864", + "sdist": "rl_kernel-0.1.0.tar.gz", + "sdist_sha256": "56026574fa3307e85e6e278298a09deee13b43d4ec5c61899698ac5e55163a12", + "native_library_member": "rl_engine/_rope_cuda.cp313-win_amd64.pyd", + "native_library_sha256": "c0878eeac05b6b400ed7c3935561de525ce3af8b2a050b5d8bfe5cdd66650940", + "inplace_library_sha256": "aab5bf0777da43609e067e6a41bc37c46fdf6b1dd0c60d1071c15bbc4f707242", + "binary_hash_scope": "Wheel and in-place libraries were linked separately; record both hashes", + "license_in_both_archives": true, + "license_sha256": "8c9ccb96c065e706135b6cbad279b721da6156e51f3a5f27c6b3329af9416d73", + "wheel_sources_match_checkout": [ + "rl_engine/backends/cuda/rope/multi_axis_rope.py", + "rl_engine/backends/shared/triton/rope/_flash_attn_rotary.py", + "rl_engine/backends/shared/triton/rope/multi_axis_rope.py", + "rl_engine/ops/rope/_registration.py", + "rl_engine/ops/rope/multi_axis_rope.py", + "rl_engine/reference/rope/multi_axis_rope.py", + "rl_engine/runtime/registry.py", + "rl_engine/validation/operators/gradient_adapters.py", + "rl_engine/validation/operators/operator_inputs.py", + "rl_engine/validation/operators/operator_specs.py", + "rl_engine/validation/operators/qwen_image_rope.py" + ], + "sdist_required_sources_match_checkout": [ + "csrc/bindings/rope_module.cpp", + "csrc/cuda/rope/multi_axis_rope.cu", + "rl_engine/backends/shared/triton/rope/FLASH_ATTENTION_LICENSE" + ], + "profile": "RL_KERNEL_ROPE_ONLY=1; full monolithic _C build not qualified" +} diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/precommit.txt b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/precommit.txt new file mode 100644 index 000000000..475a8c0d8 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/precommit.txt @@ -0,0 +1,7 @@ +trim trailing whitespace.................................................Passed +fix end of files.........................................................Passed +check yaml...............................................................Passed +check for added large files..............................................Passed +black....................................................................Passed +isort....................................................................Passed +flake8...................................................................Passed diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/quality.txt b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/quality.txt new file mode 100644 index 000000000..2756f4004 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/quality.txt @@ -0,0 +1,58 @@ +Changed-file checks; repository-pinned Black 24.4.2, isort 5.13.2 and Flake8 7.0.0. +Test commands are separate from the static checks below. +Regression command: python -m pytest tests/build/test_multi_axis_rope_build.py tests/validation/operators/test_operator_inputs.py tests/validation/operators/test_gradient_invariance.py tests/validation/operators/test_forward_invariance.py tests/validation/operators/test_elementwise_inventory.py tests/validation/operators/test_four_judgment_matrix.py tests/validation/operators/test_ws1_workload.py tests/runtime/test_kernel_registry.py tests/runtime/test_dispatch.py tests/ops/rope/test_rope.py -q -rs +Layout output copies have only trailing whitespace normalized by pre-commit. + +Command: python -m isort --check-only --profile black --line-length 100 benchmarks/operators/rope/benchmark_multi_axis_rope.py build_tools/extensions.py rl_engine/backends/cuda/rope/multi_axis_rope.py rl_engine/backends/shared/triton/rope/_flash_attn_rotary.py rl_engine/backends/shared/triton/rope/multi_axis_rope.py rl_engine/ops/rope/_registration.py rl_engine/ops/rope/multi_axis_rope.py rl_engine/reference/rope/multi_axis_rope.py rl_engine/runtime/registry.py rl_engine/validation/operators/gradient_adapters.py rl_engine/validation/operators/operator_inputs.py rl_engine/validation/operators/operator_specs.py rl_engine/validation/operators/qwen_image_rope.py tests/build/test_multi_axis_rope_build.py tests/ops/rope/test_multi_axis_rope.py tests/validation/operators/test_operator_inputs.py tools/validation/operators/check_multi_axis_rope_install.py tools/validation/operators/check_qwenimage_rope_upstream.py tools/validation/operators/validate_multi_axis_rope.py +Exit code: 0 + + +Command: python -m black --check --line-length 100 benchmarks/operators/rope/benchmark_multi_axis_rope.py build_tools/extensions.py rl_engine/backends/cuda/rope/multi_axis_rope.py rl_engine/backends/shared/triton/rope/_flash_attn_rotary.py rl_engine/backends/shared/triton/rope/multi_axis_rope.py rl_engine/ops/rope/_registration.py rl_engine/ops/rope/multi_axis_rope.py rl_engine/reference/rope/multi_axis_rope.py rl_engine/runtime/registry.py rl_engine/validation/operators/gradient_adapters.py rl_engine/validation/operators/operator_inputs.py rl_engine/validation/operators/operator_specs.py rl_engine/validation/operators/qwen_image_rope.py tests/build/test_multi_axis_rope_build.py tests/ops/rope/test_multi_axis_rope.py tests/validation/operators/test_operator_inputs.py tools/validation/operators/check_multi_axis_rope_install.py tools/validation/operators/check_qwenimage_rope_upstream.py tools/validation/operators/validate_multi_axis_rope.py +Exit code: 0 +All done! ✨ 🍰 ✨ + +19 files would be left unchanged. + +Command: python -m flake8 --max-line-length=100 --extend-ignore=E203,E704 benchmarks/operators/rope/benchmark_multi_axis_rope.py build_tools/extensions.py rl_engine/backends/cuda/rope/multi_axis_rope.py rl_engine/backends/shared/triton/rope/_flash_attn_rotary.py rl_engine/backends/shared/triton/rope/multi_axis_rope.py rl_engine/ops/rope/_registration.py rl_engine/ops/rope/multi_axis_rope.py rl_engine/reference/rope/multi_axis_rope.py rl_engine/runtime/registry.py rl_engine/validation/operators/gradient_adapters.py rl_engine/validation/operators/operator_inputs.py rl_engine/validation/operators/operator_specs.py rl_engine/validation/operators/qwen_image_rope.py tests/build/test_multi_axis_rope_build.py tests/ops/rope/test_multi_axis_rope.py tests/validation/operators/test_operator_inputs.py tools/validation/operators/check_multi_axis_rope_install.py tools/validation/operators/check_qwenimage_rope_upstream.py tools/validation/operators/validate_multi_axis_rope.py +Exit code: 0 + + +Command: python -m ruff check benchmarks/operators/rope/benchmark_multi_axis_rope.py build_tools/extensions.py rl_engine/backends/cuda/rope/multi_axis_rope.py rl_engine/backends/shared/triton/rope/_flash_attn_rotary.py rl_engine/backends/shared/triton/rope/multi_axis_rope.py rl_engine/ops/rope/_registration.py rl_engine/ops/rope/multi_axis_rope.py rl_engine/reference/rope/multi_axis_rope.py rl_engine/runtime/registry.py rl_engine/validation/operators/gradient_adapters.py rl_engine/validation/operators/operator_inputs.py rl_engine/validation/operators/operator_specs.py rl_engine/validation/operators/qwen_image_rope.py tests/build/test_multi_axis_rope_build.py tests/ops/rope/test_multi_axis_rope.py tests/validation/operators/test_operator_inputs.py tools/validation/operators/check_multi_axis_rope_install.py tools/validation/operators/check_qwenimage_rope_upstream.py tools/validation/operators/validate_multi_axis_rope.py +Exit code: 0 +All checks passed! + +Command: python -m mkdocs build --strict -f mkdocs.yaml +Exit code: 0 +INFO - Cleaning site directory + +INFO - Building documentation to directory: D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\site + +mkdocs_git_revision_date_localized_plugin: 'D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\docs\operators\multi-axis-rope.md' has no git logs, using current timestamp + +INFO - The following pages exist in the docs directory, but are not included in the "nav" configuration: + + - blog\2026-07-08-announcing-rl-kernel-linear-logp-for-vime-zh.md + + - blog\2026-08-09-dsv4-flash-moe-consistency-roadmap.md + + - operators\matmul.md + + - operators\rope.md + + - usage\cross-config-drift-report.md + + - usage\cuda-rocm-consistency-audit.md + + - usage\h100-configurable.md + + - usage\h100-cp-gradient-validation.md + + - usage\h100-matrix-validation.md + + - usage\h100-pcp-validation.md + + - usage\rocm-sparse-performance.md + +INFO - Documentation built in 4.01 seconds + +Repository pre-commit: passed; see precommit.txt. diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/regression.txt b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/regression.txt new file mode 100644 index 000000000..066e09cec --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/regression.txt @@ -0,0 +1,10 @@ +........................................................................ [ 35%] +........................................................................ [ 71%] +...................................................sssssss [100%] +=========================== short test summary info =========================== +SKIPPED [3] tests\ops\rope\test_rope.py:390: rope_apply_ascend is not compiled (requires an Ascend NPU build). +SKIPPED [1] tests\ops\rope\test_rope.py:406: rope_apply_ascend is not compiled (requires an Ascend NPU build). +SKIPPED [1] tests\ops\rope\test_rope.py:413: rope_apply_ascend is not compiled (requires an Ascend NPU build). +SKIPPED [1] tests\ops\rope\test_rope.py:420: rope_apply_ascend is not compiled (requires an Ascend NPU build). +SKIPPED [1] tests\ops\rope\test_rope.py:433: rope_apply_ascend is not compiled (requires an Ascend NPU build). +195 passed, 7 skipped in 4.31s diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/tests.txt b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/tests.txt new file mode 100644 index 000000000..38e8e21e6 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/tests.txt @@ -0,0 +1,5 @@ +........................................................................ [ 87%] +s......... [100%] +=========================== short test summary info =========================== +SKIPPED [1] tests\ops\rope\test_multi_axis_rope.py:336: requires two NVIDIA GPUs +81 passed, 1 skipped in 5.55s diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/upstream.json b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/upstream.json new file mode 100644 index 000000000..1f89880c1 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/upstream.json @@ -0,0 +1,335 @@ +{ + "source_sha256": "0bad11901deae69f05aec01fd464e51e38582d455068f2851ee665036179321f", + "cases": [ + { + "grid": [ + 2, + 3 + ], + "stream": "image", + "dtype": "torch.float32", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 2, + 3 + ], + "stream": "image", + "dtype": "torch.float16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 2, + 3 + ], + "stream": "image", + "dtype": "torch.bfloat16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 2, + 3 + ], + "stream": "text", + "dtype": "torch.float32", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 2, + 3 + ], + "stream": "text", + "dtype": "torch.float16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 2, + 3 + ], + "stream": "text", + "dtype": "torch.bfloat16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 3, + 5 + ], + "stream": "image", + "dtype": "torch.float32", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 3, + 5 + ], + "stream": "image", + "dtype": "torch.float16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 3, + 5 + ], + "stream": "image", + "dtype": "torch.bfloat16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 3, + 5 + ], + "stream": "text", + "dtype": "torch.float32", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 3, + 5 + ], + "stream": "text", + "dtype": "torch.float16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 3, + 5 + ], + "stream": "text", + "dtype": "torch.bfloat16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 64, + 64 + ], + "stream": "image", + "dtype": "torch.float32", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 64, + 64 + ], + "stream": "image", + "dtype": "torch.float16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 64, + 64 + ], + "stream": "image", + "dtype": "torch.bfloat16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 64, + 64 + ], + "stream": "text", + "dtype": "torch.float32", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 64, + 64 + ], + "stream": "text", + "dtype": "torch.float16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 64, + 64 + ], + "stream": "text", + "dtype": "torch.bfloat16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 83, + 83 + ], + "stream": "image", + "dtype": "torch.float32", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 83, + 83 + ], + "stream": "image", + "dtype": "torch.float16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 83, + 83 + ], + "stream": "image", + "dtype": "torch.bfloat16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 83, + 83 + ], + "stream": "text", + "dtype": "torch.float32", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 83, + 83 + ], + "stream": "text", + "dtype": "torch.float16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 83, + 83 + ], + "stream": "text", + "dtype": "torch.bfloat16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 104, + 58 + ], + "stream": "image", + "dtype": "torch.float32", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 104, + 58 + ], + "stream": "image", + "dtype": "torch.float16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 104, + 58 + ], + "stream": "image", + "dtype": "torch.bfloat16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 104, + 58 + ], + "stream": "text", + "dtype": "torch.float32", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 104, + 58 + ], + "stream": "text", + "dtype": "torch.float16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 104, + 58 + ], + "stream": "text", + "dtype": "torch.bfloat16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + } + ] +} diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-cuda.json b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-cuda.json new file mode 100644 index 000000000..20ad4dd78 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-cuda.json @@ -0,0 +1,1962 @@ +{ + "scope": "Qwen-Image multi_axis_rope operator; C1 elementwise and C2 config transforms", + "trace": { + "contract": "qwenimage-interleaved-rope-fp32-v1", + "actual_backend": "rlkernel.cuda.qwenimage_multi_axis_rope.v1", + "kernel_id": "torch.ops.rl_kernel_native.multi_axis_rope", + "layout": "BSHD, adjacent channel pairs", + "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", + "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" + }, + "passed": true, + "forward": { + "op_name": "multi_axis_rope", + "backend_profile": "cuda_bf16", + "accuracy_reports": [ + { + "config_id": "BN/full", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.054684440139681e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/chunked", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/chunked", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.054684440139681e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s0", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/full/s0", + "fp32_reference" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004673302173614502, + "mean_abs_error": 6.917401333339512e-05, + "max_rel_error": 0.003753266530111432, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s1", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/full/s1", + "fp32_reference" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.00048516690731048584, + "mean_abs_error": 7.802647451171651e-05, + "max_rel_error": 0.0037793784867972136, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s2", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/full/s2", + "fp32_reference" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004864335060119629, + "mean_abs_error": 7.697804539930075e-05, + "max_rel_error": 0.003821910358965397, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s3", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/full/s3", + "fp32_reference" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.710371912457049e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s0", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/chunked/s0", + "fp32_reference" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004673302173614502, + "mean_abs_error": 6.917401333339512e-05, + "max_rel_error": 0.003753266530111432, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s1", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/chunked/s1", + "fp32_reference" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.00048516690731048584, + "mean_abs_error": 7.802647451171651e-05, + "max_rel_error": 0.0037793784867972136, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s2", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/chunked/s2", + "fp32_reference" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004864335060119629, + "mean_abs_error": 7.697804539930075e-05, + "max_rel_error": 0.003821910358965397, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s3", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/chunked/s3", + "fp32_reference" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.710371912457049e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/permuted", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/permuted", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.054684440139681e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/padded_right", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/padded_right", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.054684440139681e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/padded_left", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/padded_left", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.054684440139681e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + } + ], + "invariance_reports": [ + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/chunked", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "BN/chunked" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s0", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s0" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s1", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s1" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s2", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s2" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s3", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s3" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s0", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s0" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s1", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s1" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s2", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s2" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s3", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s3" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/permuted", + "transform_kind": "permutation", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "BN/permuted" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/padded_right", + "transform_kind": "padding", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "BN/padded_right" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/padded_left", + "transform_kind": "padding", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "BN/padded_left" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + } + ], + "logprob_smoke": null, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + }, + "candidate_id": "rlkernel.cuda.qwenimage_multi_axis_rope.v1", + "device": "cuda", + "compute_capability": "sm120", + "seed": 20260812, + "fallback_reason": null, + "passed": true, + "provenance_valid": true, + "metadata_valid": true, + "observed_kernel_id": "torch.ops.rl_kernel_native.multi_axis_rope" + }, + "gradient": { + "op_name": "multi_axis_rope", + "backend_profile": "cuda_bf16", + "accuracy_reports": [ + { + "config_id": "BN/full", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.699562850873917e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/chunked", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/chunked", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.699562850873917e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s0", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/full/s0", + "fp32_reference" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 2.7887988835573196e-05, + "mean_abs_error": 3.6821968478761846e-06, + "max_rel_error": 1.8164952993392944, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s1", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/full/s1", + "fp32_reference" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.7113780965446495e-06, + "max_rel_error": 0.4316565990447998, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s2", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/full/s2", + "fp32_reference" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 2.995925024151802e-05, + "mean_abs_error": 4.7168209675874095e-06, + "max_rel_error": 0.7439075112342834, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s3", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/full/s3", + "fp32_reference" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.0211638659238815e-05, + "mean_abs_error": 4.9605891945248e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s0", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/chunked/s0", + "fp32_reference" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 2.7887988835573196e-05, + "mean_abs_error": 3.6821968478761846e-06, + "max_rel_error": 1.8164952993392944, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s1", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/chunked/s1", + "fp32_reference" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.7113780965446495e-06, + "max_rel_error": 0.4316565990447998, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s2", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/chunked/s2", + "fp32_reference" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 2.995925024151802e-05, + "mean_abs_error": 4.7168209675874095e-06, + "max_rel_error": 0.7439075112342834, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s3", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/chunked/s3", + "fp32_reference" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.0211638659238815e-05, + "mean_abs_error": 4.9605891945248e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/permuted", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/permuted", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.699562850873917e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/padded_right", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/padded_right", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.699562850873917e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/padded_left", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/padded_left", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.699562850873917e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + } + ], + "invariance_reports": [ + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/chunked", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "BN/chunked" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s0", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s0" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s1", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s1" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s2", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s2" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s3", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s3" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s0", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s0" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s1", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s1" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s2", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s2" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s3", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s3" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/permuted", + "transform_kind": "permutation", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "BN/permuted" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/padded_right", + "transform_kind": "padding", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "BN/padded_right" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/padded_left", + "transform_kind": "padding", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "BN/padded_left" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + } + ], + "singleton_aggregate_reports": [], + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + }, + "candidate_id": "rlkernel.cuda.qwenimage_multi_axis_rope.v1", + "device": "cuda", + "compute_capability": "sm120", + "seed": 20260812, + "fallback_reason": null, + "passed": true, + "provenance_valid": true, + "metadata_valid": true, + "loss_reduction": "sum_over_active_tokens_then_optional_mean_by_active_count", + "active_token_denominator": 27, + "grad_tensor_names": [ + "dx" + ], + "first_failing_op": null, + "first_failing_tensor": null, + "first_failing_config_pair": null, + "observed_kernel_id": "torch.ops.rl_kernel_native.multi_axis_rope" + } +} diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-cuda.txt b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-cuda.txt new file mode 100644 index 000000000..e69d60986 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-cuda.txt @@ -0,0 +1,3 @@ +WARNING 10-28 16:54:40 [RL-Kernel]: Core binary extension (_C) unavailable: cannot import name '_C' from 'rl_engine' (D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\rl_engine\__init__.py). Falling back to native code. +Using the installed production RoPE dispatcher library +WS1 four judgments backend=cuda: passed=True diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-triton.json b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-triton.json new file mode 100644 index 000000000..d81adfea1 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-triton.json @@ -0,0 +1,1962 @@ +{ + "scope": "Qwen-Image multi_axis_rope operator; C1 elementwise and C2 config transforms", + "trace": { + "contract": "qwenimage-interleaved-rope-fp32-v1", + "actual_backend": "rlkernel.triton.qwenimage_multi_axis_rope.v1", + "kernel_id": "rl_engine.backends.shared.triton.rope._flash_attn_rotary.rotary_kernel", + "layout": "BSHD, adjacent channel pairs", + "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", + "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" + }, + "passed": true, + "forward": { + "op_name": "multi_axis_rope", + "backend_profile": "triton_cuda_bf16", + "accuracy_reports": [ + { + "config_id": "BN/full", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.054684440139681e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/chunked", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/chunked", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.054684440139681e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s0", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/full/s0", + "fp32_reference" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004673302173614502, + "mean_abs_error": 6.917401333339512e-05, + "max_rel_error": 0.003753266530111432, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s1", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/full/s1", + "fp32_reference" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.00048516690731048584, + "mean_abs_error": 7.802647451171651e-05, + "max_rel_error": 0.0037793784867972136, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s2", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/full/s2", + "fp32_reference" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004864335060119629, + "mean_abs_error": 7.697804539930075e-05, + "max_rel_error": 0.003821910358965397, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s3", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/full/s3", + "fp32_reference" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.710371912457049e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s0", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/chunked/s0", + "fp32_reference" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004673302173614502, + "mean_abs_error": 6.917401333339512e-05, + "max_rel_error": 0.003753266530111432, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s1", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/chunked/s1", + "fp32_reference" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.00048516690731048584, + "mean_abs_error": 7.802647451171651e-05, + "max_rel_error": 0.0037793784867972136, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s2", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/chunked/s2", + "fp32_reference" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004864335060119629, + "mean_abs_error": 7.697804539930075e-05, + "max_rel_error": 0.003821910358965397, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s3", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/chunked/s3", + "fp32_reference" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.710371912457049e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/permuted", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/permuted", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.054684440139681e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/padded_right", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/padded_right", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.054684440139681e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/padded_left", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/padded_left", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.054684440139681e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + } + ], + "invariance_reports": [ + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/chunked", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "BN/chunked" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s0", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s0" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s1", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s1" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s2", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s2" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s3", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s3" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s0", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s0" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s1", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s1" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s2", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s2" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s3", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s3" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/permuted", + "transform_kind": "permutation", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "BN/permuted" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/padded_right", + "transform_kind": "padding", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "BN/padded_right" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/padded_left", + "transform_kind": "padding", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "BN/padded_left" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + } + ], + "logprob_smoke": null, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + }, + "candidate_id": "rlkernel.triton.qwenimage_multi_axis_rope.v1", + "device": "cuda", + "compute_capability": "sm120", + "seed": 20260812, + "fallback_reason": null, + "passed": true, + "provenance_valid": true, + "metadata_valid": true, + "observed_kernel_id": "rl_engine.backends.shared.triton.rope._flash_attn_rotary.rotary_kernel" + }, + "gradient": { + "op_name": "multi_axis_rope", + "backend_profile": "triton_cuda_bf16", + "accuracy_reports": [ + { + "config_id": "BN/full", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.699562850873917e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/chunked", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/chunked", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.699562850873917e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s0", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/full/s0", + "fp32_reference" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 2.7887988835573196e-05, + "mean_abs_error": 3.6821968478761846e-06, + "max_rel_error": 1.8164952993392944, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s1", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/full/s1", + "fp32_reference" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.7113780965446495e-06, + "max_rel_error": 0.4316565990447998, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s2", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/full/s2", + "fp32_reference" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 2.995925024151802e-05, + "mean_abs_error": 4.7168209675874095e-06, + "max_rel_error": 0.7439075112342834, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s3", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/full/s3", + "fp32_reference" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.0211638659238815e-05, + "mean_abs_error": 4.9605891945248e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s0", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/chunked/s0", + "fp32_reference" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 2.7887988835573196e-05, + "mean_abs_error": 3.6821968478761846e-06, + "max_rel_error": 1.8164952993392944, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s1", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/chunked/s1", + "fp32_reference" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.7113780965446495e-06, + "max_rel_error": 0.4316565990447998, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s2", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/chunked/s2", + "fp32_reference" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 2.995925024151802e-05, + "mean_abs_error": 4.7168209675874095e-06, + "max_rel_error": 0.7439075112342834, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s3", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/chunked/s3", + "fp32_reference" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.0211638659238815e-05, + "mean_abs_error": 4.9605891945248e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/permuted", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/permuted", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.699562850873917e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/padded_right", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/padded_right", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.699562850873917e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/padded_left", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/padded_left", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.699562850873917e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + } + ], + "invariance_reports": [ + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/chunked", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "BN/chunked" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s0", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s0" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s1", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s1" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s2", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s2" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s3", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s3" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s0", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s0" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s1", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s1" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s2", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s2" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s3", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s3" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/permuted", + "transform_kind": "permutation", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "BN/permuted" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/padded_right", + "transform_kind": "padding", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "BN/padded_right" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/padded_left", + "transform_kind": "padding", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "BN/padded_left" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + } + ], + "singleton_aggregate_reports": [], + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + }, + "candidate_id": "rlkernel.triton.qwenimage_multi_axis_rope.v1", + "device": "cuda", + "compute_capability": "sm120", + "seed": 20260812, + "fallback_reason": null, + "passed": true, + "provenance_valid": true, + "metadata_valid": true, + "loss_reduction": "sum_over_active_tokens_then_optional_mean_by_active_count", + "active_token_denominator": 27, + "grad_tensor_names": [ + "dx" + ], + "first_failing_op": null, + "first_failing_tensor": null, + "first_failing_config_pair": null, + "observed_kernel_id": "rl_engine.backends.shared.triton.rope._flash_attn_rotary.rotary_kernel" + } +} diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-triton.txt b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-triton.txt new file mode 100644 index 000000000..283d1933a --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-triton.txt @@ -0,0 +1,3 @@ +WARNING 10-28 16:54:47 [RL-Kernel]: Core binary extension (_C) unavailable: cannot import name '_C' from 'rl_engine' (D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\rl_engine\__init__.py). Falling back to native code. +Using the installed production RoPE dispatcher library +WS1 four judgments backend=triton: passed=True diff --git a/rl_engine/backends/cuda/rope/multi_axis_rope.py b/rl_engine/backends/cuda/rope/multi_axis_rope.py new file mode 100644 index 000000000..83ef741ca --- /dev/null +++ b/rl_engine/backends/cuda/rope/multi_axis_rope.py @@ -0,0 +1,49 @@ +# SPDX-License-Identifier: Apache-2.0 +"""Architecture-independent CUDA implementation; no Hopper-only instructions.""" + +import importlib.util + +import torch + +from rl_engine.backends import extension # noqa: F401 - loads the normal native library +from rl_engine.ops.rope._registration import register_rotation +from rl_engine.ops.rope.multi_axis_rope import MultiAxisRoPEMetadata, validate_inputs + + +def _apply(x, cos, sin, transpose): + return torch.ops.rl_kernel_native.multi_axis_rope(x, cos, sin, transpose) + + +def native_available(): + if torch.version.hip is not None: + return False + if not hasattr(torch.ops.rl_kernel_native, "multi_axis_rope"): + library = importlib.util.find_spec("rl_engine._rope_cuda") + if library is not None: + torch.ops.load_library(library.origin) + return hasattr(torch.ops.rl_kernel_native, "multi_axis_rope") + + +_rotation = register_rotation( + "qwen_image_rope_cuda", + _apply, + kernel_id="torch.ops.rl_kernel_native.multi_axis_rope", + family="cuda", +) + + +class MultiAxisRoPECudaOp(MultiAxisRoPEMetadata): + backend_id = "rlkernel.cuda.qwenimage_multi_axis_rope.v1" + kernel_id = "torch.ops.rl_kernel_native.multi_axis_rope" + + def __init__(self): + if not native_available(): + raise RuntimeError( + "CUDA multi_axis_rope extension is unavailable; rebuild rl_engine._C" + ) + + def __call__(self, x, cos, sin): + validate_inputs(x, cos, sin) + if x.device.type != "cuda": + raise ValueError("CUDA multi_axis_rope requires CUDA tensors") + return _rotation(x, cos, sin, False) diff --git a/rl_engine/backends/shared/triton/rope/FLASH_ATTENTION_LICENSE b/rl_engine/backends/shared/triton/rope/FLASH_ATTENTION_LICENSE new file mode 100644 index 000000000..5860e4b33 --- /dev/null +++ b/rl_engine/backends/shared/triton/rope/FLASH_ATTENTION_LICENSE @@ -0,0 +1,29 @@ +BSD 3-Clause License + +Copyright (c) 2022, the respective contributors, as shown by the AUTHORS file. +All rights reserved. + +Redistribution and use in source and binary forms, with or without +modification, are permitted provided that the following conditions are met: + +* Redistributions of source code must retain the above copyright notice, this + list of conditions and the following disclaimer. + +* Redistributions in binary form must reproduce the above copyright notice, + this list of conditions and the following disclaimer in the documentation + and/or other materials provided with the distribution. + +* Neither the name of the copyright holder nor the names of its + contributors may be used to endorse or promote products derived from + this software without specific prior written permission. + +THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" +AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE +IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE +DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE LIABLE +FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL +DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR +SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER +CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, +OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE +OF THIS SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. diff --git a/rl_engine/backends/shared/triton/rope/_flash_attn_rotary.py b/rl_engine/backends/shared/triton/rope/_flash_attn_rotary.py new file mode 100644 index 000000000..bd617c8c1 --- /dev/null +++ b/rl_engine/backends/shared/triton/rope/_flash_attn_rotary.py @@ -0,0 +1,141 @@ +# SPDX-License-Identifier: BSD-3-Clause +# Copyright (c) 2025, Tri Dao. +# Adapted from Dao-AILab/flash-attention, revision +# 98c1cee0ed59307833841cdbf567fe336ceec13a, flash_attn/ops/triton/rotary.py. +# Changes: explicit shared/per-sample table strides. Launch policy is in the +# sibling wrapper. See FLASH_ATTENTION_LICENSE for the retained BSD-3 terms. + +import triton +import triton.language as tl + + +@triton.jit +def rotary_kernel( + OUT, # Pointers to matrices + X, + COS, + SIN, + CU_SEQLENS, + SEQLEN_OFFSETS, # this could be int or a pointer + # Matrix dimensions + seqlen, + nheads, + seqlen_ro, + # strides + stride_out_batch, + stride_out_seqlen, + stride_out_nheads, + stride_out_headdim, + stride_x_batch, + stride_x_seqlen, + stride_x_nheads, + stride_x_headdim, + stride_cos_batch, + stride_cos_seqlen, + stride_cos_dim, + stride_sin_batch, + stride_sin_seqlen, + stride_sin_dim, + # Meta-parameters + # We want ROTARY_DIM to be constexpr, otherwise the triton compiler doesn't know that + # the mask is constant every 8 elements, and it will generate LDG.16 instead of LDG.128 + ROTARY_DIM: tl.constexpr, + IS_SEQLEN_OFFSETS_TENSOR: tl.constexpr, + IS_VARLEN: tl.constexpr, + INTERLEAVED: tl.constexpr, + CONJUGATE: tl.constexpr, + BLOCK_H: tl.constexpr, + BLOCK_M: tl.constexpr, +): + BLOCK_K: tl.constexpr = triton.next_power_of_2(ROTARY_DIM) + ROTARY_DIM_HALF = ROTARY_DIM // 2 + pid_head = tl.program_id(axis=0) + pid_m = tl.program_id(axis=1) + pid_batch = tl.program_id(axis=2) + + if not IS_VARLEN: + X = X + pid_batch * stride_x_batch + OUT = OUT + pid_batch * stride_out_batch + else: + start_idx = tl.load(CU_SEQLENS + pid_batch) + seqlen = tl.load(CU_SEQLENS + pid_batch + 1) - start_idx + X = X + start_idx * stride_x_seqlen + OUT = OUT + start_idx * stride_out_seqlen + + if pid_m * BLOCK_M >= seqlen: + return + + rh = pid_head * BLOCK_H + tl.arange(0, BLOCK_H) + rm = pid_m * BLOCK_M + tl.arange(0, BLOCK_M) + if not IS_SEQLEN_OFFSETS_TENSOR: + rm_cs = rm + SEQLEN_OFFSETS + else: + rm_cs = rm + tl.load(SEQLEN_OFFSETS + pid_batch) + + rk_half = tl.arange(0, BLOCK_K // 2) + COS = ( + COS + + pid_batch * stride_cos_batch + + (rm_cs[:, None] * stride_cos_seqlen + rk_half[None, :] * stride_cos_dim) + ) + SIN = ( + SIN + + pid_batch * stride_sin_batch + + (rm_cs[:, None] * stride_sin_seqlen + rk_half[None, :] * stride_sin_dim) + ) + mask_cs = (rm_cs[:, None] < seqlen_ro) & (rk_half[None, :] < ROTARY_DIM_HALF) + cos = tl.load(COS, mask=mask_cs, other=1.0).to(tl.float32) + sin = tl.load(SIN, mask=mask_cs, other=0.0).to(tl.float32) + if CONJUGATE: + sin = -sin + + if not INTERLEAVED: + # Load the 1st and 2nd halves of X, do calculation, then store to 1st and 2nd halves of OUT + X = X + ( + rh[:, None, None] * stride_x_nheads + + rm[None, :, None] * stride_x_seqlen + + rk_half[None, None, :] * stride_x_headdim + ) + OUT = OUT + ( + rh[:, None, None] * stride_out_nheads + + rm[None, :, None] * stride_out_seqlen + + rk_half[None, None, :] * stride_out_headdim + ) + mask = ( + (rh[:, None, None] < nheads) + & (rm[None, :, None] < seqlen) + & (rk_half[None, None, :] < ROTARY_DIM_HALF) + ) + x0 = tl.load(X, mask=mask, other=0.0).to(tl.float32) + x1 = tl.load( + X + ROTARY_DIM_HALF * stride_x_headdim, + mask=mask, + other=0.0, + ).to(tl.float32) + o0 = x0 * cos - x1 * sin + o1 = x0 * sin + x1 * cos + tl.store(OUT, o0, mask=mask) + tl.store(OUT + ROTARY_DIM_HALF * stride_out_headdim, o1, mask=mask) + else: + rk = tl.arange(0, BLOCK_K) + X = X + ( + rh[:, None, None] * stride_x_nheads + + rm[None, :, None] * stride_x_seqlen + + rk[None, None, :] * stride_x_headdim + ) + OUT = OUT + ( + rh[:, None, None] * stride_out_nheads + + rm[None, :, None] * stride_out_seqlen + + rk[None, None, :] * stride_out_headdim + ) + mask = ( + (rh[:, None, None] < nheads) + & (rm[None, :, None] < seqlen) + & (rk[None, None, :] < ROTARY_DIM) + ) + x = tl.load(X, mask=mask, other=0.0).to(tl.float32) + x0, x1 = tl.split(tl.reshape(x, [BLOCK_H, BLOCK_M, BLOCK_K // 2, 2])) + o0 = x0 * cos - x1 * sin + o1 = x0 * sin + x1 * cos + o = tl.reshape(tl.join(o0, o1), [BLOCK_H, BLOCK_M, BLOCK_K]) + tl.store(OUT, o, mask=mask) diff --git a/rl_engine/backends/shared/triton/rope/multi_axis_rope.py b/rl_engine/backends/shared/triton/rope/multi_axis_rope.py new file mode 100644 index 000000000..1a37836fe --- /dev/null +++ b/rl_engine/backends/shared/triton/rope/multi_axis_rope.py @@ -0,0 +1,68 @@ +# SPDX-License-Identifier: Apache-2.0 +"""FlashAttention-derived, stride-aware rotation under a frozen FP32 profile.""" + +import torch +import triton + +from rl_engine.ops.rope._registration import register_rotation +from rl_engine.ops.rope.multi_axis_rope import MultiAxisRoPEMetadata, validate_inputs + +from ._flash_attn_rotary import rotary_kernel + + +def _apply(x, cos, sin, transpose): + output = torch.empty(x.shape, device=x.device, dtype=x.dtype) + batch, seq, heads, _ = x.shape + if batch and seq: + with torch.cuda.device(x.device): + rotary_kernel[(triton.cdiv(heads, 2), triton.cdiv(seq, 8), batch)]( + output, + x, + cos, + sin, + None, + 0, + seq, + heads, + seq, + *output.stride(), + *x.stride(), + cos.stride(0) if cos.ndim == 3 else 0, + cos.stride(-2), + cos.stride(-1), + sin.stride(0) if sin.ndim == 3 else 0, + sin.stride(-2), + sin.stride(-1), + 128, + False, + False, + True, + transpose, + BLOCK_H=2, + BLOCK_M=8, + enable_fp_fusion=False, + ) + return output + + +_rotation = register_rotation( + "qwen_image_rope_triton", + _apply, + kernel_id="rl_engine.backends.shared.triton.rope._flash_attn_rotary.rotary_kernel", + family="triton", +) + + +class MultiAxisRoPETritonOp(MultiAxisRoPEMetadata): + backend_id = "rlkernel.triton.qwenimage_multi_axis_rope.v1" + kernel_id = "rl_engine.backends.shared.triton.rope._flash_attn_rotary.rotary_kernel" + + def __init__(self): + if torch.version.hip is not None or not torch.cuda.is_available(): + raise RuntimeError("this multi_axis_rope Triton profile requires NVIDIA CUDA") + + def __call__(self, x, cos, sin): + validate_inputs(x, cos, sin) + if x.device.type != "cuda": + raise ValueError("Triton multi_axis_rope requires CUDA tensors") + return _rotation(x, cos, sin, False) diff --git a/rl_engine/ops/rope/_registration.py b/rl_engine/ops/rope/_registration.py new file mode 100644 index 000000000..343d41a9c --- /dev/null +++ b/rl_engine/ops/rope/_registration.py @@ -0,0 +1,41 @@ +# SPDX-License-Identifier: Apache-2.0 +"""Dispatcher, fake and recursive input-VJP registration for the RoPE profile.""" + +import torch + +from rl_engine.ops.autograd.backward_runtime import record_backward +from rl_engine.ops.rope.multi_axis_rope import validate_inputs + + +def register_rotation(name, implementation, *, kernel_id, family): + @torch.library.custom_op(f"rl_kernel::{name}", mutates_args=(), device_types="cuda") + def rotation( + x: torch.Tensor, cos: torch.Tensor, sin: torch.Tensor, transpose: bool + ) -> torch.Tensor: + validate_inputs(x, cos, sin) + result = implementation(x, cos, sin, transpose) + if transpose: + record_backward( + "multi_axis_rope", + kernel_id=kernel_id, + impl="interleaved_fp32_transpose", + family=family, + ) + return result + + @rotation.register_fake + def fake(x, cos, sin, transpose): + validate_inputs(x, cos, sin) + return torch.empty(x.shape, dtype=x.dtype, device=x.device) + + def setup_context(ctx, inputs, output): + _, cos, sin, transpose = inputs + ctx.save_for_backward(cos, sin) + ctx.transpose = transpose + + def backward(ctx, grad): + cos, sin = ctx.saved_tensors + return rotation(grad, cos, sin, not ctx.transpose), None, None, None + + rotation.register_autograd(backward, setup_context=setup_context) + return rotation diff --git a/rl_engine/ops/rope/multi_axis_rope.py b/rl_engine/ops/rope/multi_axis_rope.py new file mode 100644 index 000000000..ff95adff2 --- /dev/null +++ b/rl_engine/ops/rope/multi_axis_rope.py @@ -0,0 +1,132 @@ +# SPDX-License-Identifier: Apache-2.0 +"""Qwen-Image interleaved three-axis RoPE (RFC #386). + +Tables are immutable inputs, built before capture/replay. The rotation accepts +BSHD tensors, unlike the legacy BHSD, rotate-half RoPE operator. +""" + +from __future__ import annotations + +import torch + +CONTRACT = "qwenimage-interleaved-rope-fp32-v1" +AXES_DIMS = (16, 56, 56) +UPSTREAM_REVISION = "40f63690da440e8c382ab53e67508835deb1ed2f" + + +def qwen_image_positions( + height: int, width: int, text_length: int +) -> tuple[torch.Tensor, torch.Tensor]: + """CPU int64 positions for one image's *packed latent* H/W, then text. + + Matches QwenEmbedRope(scale_rope=True), single-frame text-to-image. Odd + grids have the extra negative coordinate: H=3 gives [-2, -1, 0]. + Return separate tables so callers explicitly choose their stream order. + """ + for name, value in ( + ("height", height), + ("width", width), + ("text_length", text_length), + ): + if isinstance(value, bool) or not isinstance(value, int): + raise TypeError(f"{name} must be an integer") + if not 1 <= height <= 8192 or not 1 <= width <= 8192 or text_length < 0: + raise ValueError("height/width must be in [1, 8192], text_length must be nonnegative") + start = max(height // 2, width // 2) + if start + text_length > 4096: + raise ValueError("text positions exceed the pinned upstream 4096-entry table") + h = torch.arange(-(height - height // 2), height // 2, device="cpu") + w = torch.arange(-(width - width // 2), width // 2, device="cpu") + hh, ww = torch.meshgrid(h, w, indexing="ij") + image = torch.stack((torch.zeros_like(hh), hh, ww), dim=-1).reshape(-1, 3) + text = ( + torch.arange(start, start + text_length, device="cpu")[:, None].expand(-1, 3).contiguous() + ) + return image, text + + +def build_multi_axis_cos_sin( + positions: torch.Tensor, *, device=None +) -> tuple[torch.Tensor, torch.Tensor]: + """Build canonical FP32 tables on CPU, then copy to the requested device. + + positions: int32/int64 [S, 3] or [B, S, 3], range [-4096, 4095]. + The CPU torch.polar construction follows the pinned Diffusers reference; + no GPU transcendental or per-batch table recomputation enters the hot path. + """ + if positions.dtype not in (torch.int32, torch.int64): + raise TypeError("positions must be int32 or int64") + if positions.ndim not in (2, 3) or positions.shape[-1] != 3: + raise ValueError("positions must be [S, 3] or [B, S, 3]") + cpu = positions.detach().to(device="cpu", dtype=torch.float32) + if torch.any((cpu < -4096) | (cpu >= 4096)): + raise ValueError("positions exceed the pinned upstream [-4096, 4095] range") + phases = [] + for axis, dim in enumerate(AXES_DIMS): + inv_freq = 1.0 / torch.pow(10000, torch.arange(0, dim, 2, device="cpu").float().div(dim)) + phases.append(cpu[..., axis, None] * inv_freq) + phase = torch.cat(phases, dim=-1) + cis = torch.polar(torch.ones_like(phase), phase) + target = positions.device if device is None else device + return cis.real.contiguous().to(target), cis.imag.contiguous().to(target) + + +def validate_inputs(x: torch.Tensor, cos: torch.Tensor, sin: torch.Tensor) -> None: + if x.ndim != 4 or x.shape[-1] != 128 or x.shape[2] <= 0: + raise ValueError("x must be [B, S, H, 128] with H > 0") + if x.dtype not in (torch.float16, torch.bfloat16, torch.float32): + raise TypeError("x must be float16, bfloat16 or float32") + shared = (x.shape[1], 64) + batched = (x.shape[0], x.shape[1], 64) + if cos.shape not in (shared, batched) or sin.shape != cos.shape: + raise ValueError("cos/sin must have matching [S, 64] or [B, S, 64] shapes") + if cos.dtype != torch.float32 or sin.dtype != torch.float32: + raise TypeError("cos/sin must be float32") + if x.device != cos.device or x.device != sin.device: + raise ValueError("x, cos and sin must be on the same device") + if cos.requires_grad or sin.requires_grad: + raise ValueError("cos/sin are fixed tables; table gradients are unsupported") + + +def get_multi_axis_rope_op(backend: str = "auto", *, device="cpu"): + """Resolve an op exposing trace(); explicit backends never silently fall back. + + Strict callers must request cuda or triton and retain the returned trace. + auto is a diagnostic/convenience mode with a PyTorch reference fallback. + """ + if backend == "auto": + from rl_engine.runtime.registry import kernel_registry + + return kernel_registry.get_op("multi_axis_rope", device=device) + if backend == "cuda": + from rl_engine.backends.cuda.rope.multi_axis_rope import MultiAxisRoPECudaOp + + return MultiAxisRoPECudaOp() + if backend == "triton": + from rl_engine.backends.shared.triton.rope.multi_axis_rope import MultiAxisRoPETritonOp + + return MultiAxisRoPETritonOp() + if backend == "reference": + from rl_engine.reference.rope.multi_axis_rope import NativeMultiAxisRoPEOp + + return NativeMultiAxisRoPEOp() + raise ValueError(f"unknown multi_axis_rope backend: {backend!r}") + + +def multi_axis_rope(x, cos, sin, *, backend="auto"): + return get_multi_axis_rope_op(backend, device=x.device)(x, cos, sin) + + +class MultiAxisRoPEMetadata: + backend_id: str + kernel_id: str + + def trace(self) -> dict: + return { + "contract": CONTRACT, + "actual_backend": self.backend_id, + "kernel_id": self.kernel_id, + "layout": "BSHD, adjacent channel pairs", + "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", + "table_revision": UPSTREAM_REVISION, + } diff --git a/rl_engine/reference/rope/multi_axis_rope.py b/rl_engine/reference/rope/multi_axis_rope.py new file mode 100644 index 000000000..353eb6a00 --- /dev/null +++ b/rl_engine/reference/rope/multi_axis_rope.py @@ -0,0 +1,25 @@ +# SPDX-License-Identifier: Apache-2.0 +"""Explicit FP32 reference for the Qwen-Image interleaved rotation.""" + +import torch + +from rl_engine.ops.rope.multi_axis_rope import MultiAxisRoPEMetadata, validate_inputs + + +class NativeMultiAxisRoPEOp(MultiAxisRoPEMetadata): + op_class = "elementwise" + backend_id = "rlkernel.reference.qwenimage_multi_axis_rope.v1" + kernel_id = "rl_engine.reference.rope.multi_axis_rope.NativeMultiAxisRoPEOp" + + def __call__(self, x, cos, sin): + return self.forward(x, cos, sin) + + def forward(self, x, cos, sin): + return self.forward_fp32(x, cos, sin).to(x.dtype) + + def forward_fp32(self, x, cos, sin): + validate_inputs(x, cos, sin) + c, s = cos.unsqueeze(-2), sin.unsqueeze(-2) + even, odd = x[..., 0::2].float(), x[..., 1::2].float() + out = torch.stack((even * c - odd * s, odd * c + even * s), dim=-1) + return out.flatten(-2) diff --git a/rl_engine/runtime/registry.py b/rl_engine/runtime/registry.py index 7751d5901..326081f44 100644 --- a/rl_engine/runtime/registry.py +++ b/rl_engine/runtime/registry.py @@ -167,6 +167,11 @@ class OpBackend(Enum, metaclass=_KernelEnumMeta): PYTORCH_NATIVE_ROPE = "rl_engine.kernels.ops.pytorch.rotary_embedding.rope.NativeRoPEOp" TRITON_ROPE = "rl_engine.kernels.ops.triton.rotary_embedding.rope.TritonRoPEOp" CUDA_ROPE_SM90 = "rl_engine.kernels.ops.cuda.rotary_embedding.rope.RoPESM90Op" + CUDA_MULTI_AXIS_ROPE = "rl_engine.backends.cuda.rope.multi_axis_rope.MultiAxisRoPECudaOp" + TRITON_MULTI_AXIS_ROPE = ( + "rl_engine.backends.shared.triton.rope.multi_axis_rope.MultiAxisRoPETritonOp" + ) + PYTORCH_MULTI_AXIS_ROPE = "rl_engine.reference.rope.multi_axis_rope.NativeMultiAxisRoPEOp" ASCEND_ROPE = "rl_engine.kernels.ops.ascend.rotary_embedding.rope.RoPEAscendOp" PYTORCH_NATIVE_SILU = "rl_engine.kernels.ops.pytorch.activation.swiglu.NativeSiLUOp" PYTORCH_NATIVE_SWIGLU = "rl_engine.kernels.ops.pytorch.activation.swiglu.NativeSwiGLUOp" @@ -737,6 +742,15 @@ def __init__(self): ], }, } + # CUDA qualification for the new interleaved Qwen-Image contract. + # Other platforms retain an explicit diagnostic reference, not logp. + for candidates in self._priority_map.values(): + candidates["multi_axis_rope"] = [OpBackend.PYTORCH_MULTI_AXIS_ROPE] + self._priority_map["cuda"]["multi_axis_rope"] = [ + OpBackend.CUDA_MULTI_AXIS_ROPE, + OpBackend.TRITON_MULTI_AXIS_ROPE, + OpBackend.PYTORCH_MULTI_AXIS_ROPE, + ] # Preserve the former CPU fallback behavior for every operator on NPU, # then override only the operators with an Ascend-specific backend. self._priority_map["npu"] = { diff --git a/rl_engine/validation/operators/gradient_adapters.py b/rl_engine/validation/operators/gradient_adapters.py index f22d80285..6bc6e8104 100644 --- a/rl_engine/validation/operators/gradient_adapters.py +++ b/rl_engine/validation/operators/gradient_adapters.py @@ -220,6 +220,21 @@ def to_dict(self) -> dict[str, Any]: "csrc/cuda/logprob/fused_linear_logp_sm90.cu", ), ), + # Optional Qwen-Image extension: do not certify it as a Dense Qwen3 chain node. + "multi_axis_rope": GradientAdapterSpec( + op_name="multi_axis_rope", + chain_node="multi_axis_rope", + op_class="elementwise", + spec_name="multi_axis_rope", + tensors=(_DX,), + requirement="optional_fused", + source_files=( + "rl_engine/backends/cuda/rope/multi_axis_rope.py", + "rl_engine/backends/shared/triton/rope/multi_axis_rope.py", + "rl_engine/backends/shared/triton/rope/_flash_attn_rotary.py", + "csrc/cuda/rope/multi_axis_rope.cu", + ), + ), "rope": GradientAdapterSpec( op_name="rope", chain_node="rope", @@ -712,6 +727,16 @@ def _row_inputs( "weight": params["weight"], "bias": None, } + if op_name == "multi_axis_rope": + from rl_engine.ops.rope.multi_axis_rope import build_multi_axis_cos_sin + + if head_dim != 128: + raise ValueError("multi_axis_rope gradient adapter requires head_dim=128") + rows = _stack_rows(keys, leading, (n_heads, head_dim), device=device, dtype=dtype) + positions = _row_positions(keys, device=torch.device("cpu")) + coords = torch.stack((positions % 7, positions % 127 - 63, positions % 61 - 30), dim=-1) + cos, sin = build_multi_axis_cos_sin(coords, device=device) + return {"x": rows.unsqueeze(0), "cos": cos, "sin": sin} if op_name == "rope": rows = _stack_rows(keys, leading, (n_heads, head_dim), device=device, dtype=dtype) return { @@ -724,6 +749,10 @@ def _row_inputs( def _to_rows(op_name: str, value: torch.Tensor, n_rows: int) -> torch.Tensor: """Normalize an operator output / input-grad back to (n_rows, *tail).""" + if op_name == "multi_axis_rope": + if value.shape[0] != 1 or value.shape[1] != n_rows: + raise ValueError("multi_axis_rope row adapter expects [1, tokens, heads, 128]") + return value.squeeze(0) if op_name == "rope": # RoPE runs as (1, heads, tokens, head_dim); tokens is the row axis. permuted = value.permute(0, 2, 1, 3) diff --git a/rl_engine/validation/operators/operator_inputs.py b/rl_engine/validation/operators/operator_inputs.py index ca3b7c120..137d8d69b 100644 --- a/rl_engine/validation/operators/operator_inputs.py +++ b/rl_engine/validation/operators/operator_inputs.py @@ -37,6 +37,7 @@ def make_operator_inputs( "linear_logp": _make_linear_logp_inputs, "batch_invariant_logp": _make_batch_invariant_logp_inputs, "rope": _make_rope_inputs, + "multi_axis_rope": _make_multi_axis_rope_inputs, "silu": _make_silu_inputs, "swiglu": _make_swiglu_inputs, "embedding": _make_embedding_inputs, @@ -67,6 +68,7 @@ def operator_shape_name(op_name: str, args: argparse.Namespace) -> str: "linear_logp": f"{batch}x{seq}x{_normalized_dim(args)}x{vocab}", "batch_invariant_logp": f"{batch}x{seq}x{vocab}", "rope": f"{batch}x{DEFAULT_N_HEADS}x{seq}x{DEFAULT_HEAD_DIM}", + "multi_axis_rope": f"{batch}x{seq}x{_arg_int(args, 'n_heads', 24)}x128", "silu": f"{batch}x{seq}x{DEFAULT_INTERMEDIATE}", "swiglu": f"{batch}x{seq}x{DEFAULT_INTERMEDIATE}", "embedding": f"{batch}x{seq}x{vocab}x{_normalized_dim(args)}", @@ -250,6 +252,23 @@ def _make_batch_invariant_logp_inputs( } +def _make_multi_axis_rope_inputs(args, dtype, device): + from rl_engine.ops.rope.multi_axis_rope import build_multi_axis_cos_sin + + batch, seq = _batch_seq(args) + if _arg_int(args, "head_dim", 128) != 128: + raise ValueError("multi_axis_rope requires head_dim=128") + heads = _arg_int(args, "n_heads", 24) + row = torch.arange(seq, device="cpu", dtype=torch.int64) + positions = torch.stack((row % 7, (row // 127) % 4096, row % 127 - 63), dim=-1) + cos, sin = build_multi_axis_cos_sin(positions, device=device) + return { + "x": _floating_tensor((batch, seq, heads, 128), args, dtype, device, offset=0), + "cos": cos, + "sin": sin, + } + + def _make_rope_inputs( args: argparse.Namespace, dtype: torch.dtype, device: torch.device ) -> dict[str, Any]: diff --git a/rl_engine/validation/operators/operator_specs.py b/rl_engine/validation/operators/operator_specs.py index 2323f4787..bc23ab28a 100644 --- a/rl_engine/validation/operators/operator_specs.py +++ b/rl_engine/validation/operators/operator_specs.py @@ -197,6 +197,20 @@ def _load_object(path: str) -> Any: }, grad_input_names=("a", "b"), ), + "multi_axis_rope": OperatorSpec( + name="multi_axis_rope", + op_class="elementwise", + gold_path="rl_engine.reference.rope.multi_axis_rope.NativeMultiAxisRoPEOp", + gold_method="forward_fp32", + candidate_paths={ + "pytorch": "rl_engine.reference.rope.multi_axis_rope.NativeMultiAxisRoPEOp", + "cuda": "rl_engine.backends.cuda.rope.multi_axis_rope.MultiAxisRoPECudaOp", + "triton": ( + "rl_engine.backends.shared.triton.rope.multi_axis_rope.MultiAxisRoPETritonOp" + ), + }, + grad_input_names=("x",), + ), "rope": OperatorSpec( name="rope", op_class="elementwise", diff --git a/rl_engine/validation/operators/qwen_image_rope.py b/rl_engine/validation/operators/qwen_image_rope.py new file mode 100644 index 000000000..bc605bd32 --- /dev/null +++ b/rl_engine/validation/operators/qwen_image_rope.py @@ -0,0 +1,80 @@ +# SPDX-License-Identifier: Apache-2.0 +"""Qwen-Image operator-only four-judgment gate using the standard C3/C4 APIs. + +Reuses C2's logical sample/config transforms and C1's elementwise tolerances. +It does not certify the Dense Qwen3 chain, distributed execution or Qwen-Image +model assembly. The operator has its own OP_SPECS entry and gradient adapter. +""" + +import torch + +from rl_engine.contracts.numerical import BackendProvenance +from rl_engine.ops.rope.multi_axis_rope import get_multi_axis_rope_op +from rl_engine.reference.rope.multi_axis_rope import NativeMultiAxisRoPEOp +from rl_engine.validation.operators.forward_invariance import assert_forward_batch_invariant +from rl_engine.validation.operators.gradient_adapters import ( + get_adapter, + make_forward_runner, + make_gradient_runner, +) +from rl_engine.validation.operators.gradient_invariance import assert_gradient_batch_invariant + + +def run_multi_axis_rope_ws1(backend): + if backend not in ("cuda", "triton"): + raise ValueError("WS1 GPU evidence requires an explicit cuda or triton backend") + device = torch.device("cuda") + profile = "cuda_bf16" if backend == "cuda" else "triton_cuda_bf16" + op = get_multi_axis_rope_op(backend, device=device) + gold = NativeMultiAxisRoPEOp().forward_fp32 + trace = op.trace() + dtype = torch.bfloat16 + provenance = BackendProvenance( + backend_profile=profile, + requested_backend=backend, + actual_backend=backend, + execution_dtype="bfloat16", + accumulation_dtype="float32", + output_dtype="bfloat16", + reference_dtype="float32", + candidate_tf32_enabled=False, + reference_tf32_enabled=False, + ) + metadata = { + "backend_profile": profile, + "provenance": provenance, + "op_class": "elementwise", + "dtype": dtype, + "op_name": "multi_axis_rope", + "candidate_id": trace["actual_backend"], + "device": str(device), + "compute_capability": "sm%d%d" % torch.cuda.get_device_capability(), + "observed_actual_backend": backend, + "observed_kernel_id": trace["kernel_id"], + "observed_output_dtype": "bfloat16", + } + runner_options = {"device": device, "dtype": dtype, "n_heads": 3, "head_dim": 128} + candidate_options = {"backend_family": backend, "kernel_id": trace["kernel_id"]} + forward = assert_forward_batch_invariant( + make_forward_runner( + "multi_axis_rope", op, reference=False, **runner_options, **candidate_options + ), + gold_fn=make_forward_runner("multi_axis_rope", gold, reference=True, **runner_options), + include_logprob_smoke=False, + **metadata, + ) + gradient = assert_gradient_batch_invariant( + make_gradient_runner( + "multi_axis_rope", op, reference=False, **runner_options, **candidate_options + ), + gold_fn=make_gradient_runner("multi_axis_rope", gold, reference=True, **runner_options), + grad_tensors=get_adapter("multi_axis_rope").tensors, + **metadata, + ) + return { + "scope": "Qwen-Image multi_axis_rope operator; C1 elementwise and C2 config transforms", + "trace": trace, + "passed": forward.passed and gradient.passed, + "forward": forward.to_dict(), + "gradient": gradient.to_dict(), + } diff --git a/tests/build/test_multi_axis_rope_build.py b/tests/build/test_multi_axis_rope_build.py new file mode 100644 index 000000000..c9d06791e --- /dev/null +++ b/tests/build/test_multi_axis_rope_build.py @@ -0,0 +1,46 @@ +# SPDX-License-Identifier: Apache-2.0 +"""The explicit operator wheel profile must build native code or fail closed.""" + +from types import SimpleNamespace + +import pytest +import torch + +from build_tools import extensions + + +def test_rope_only_wheel_selects_production_sources(monkeypatch): + monkeypatch.setenv("RL_KERNEL_ROPE_ONLY", "1") + fake_torch = SimpleNamespace( + __file__=torch.__file__, version=SimpleNamespace(cuda="12.8", hip=None) + ) + monkeypatch.setattr( + extensions, + "_load_torch_extension_tools", + lambda: (fake_torch, None, lambda **kwargs: kwargs), + ) + (extension,) = extensions.get_extensions() + assert extension["name"] == "rl_engine._rope_cuda" + assert set(extension["sources"]) == { + "csrc/bindings/rope_module.cpp", + "csrc/cuda/rope/multi_axis_rope.cu", + } + assert "--fmad=false" in extension["extra_compile_args"]["nvcc"] + + +@pytest.mark.parametrize("hip", [None, "6.2"]) +def test_rope_only_wheel_rejects_non_cuda_builds(monkeypatch, hip): + monkeypatch.setenv("RL_KERNEL_ROPE_ONLY", "1") + fake_torch = SimpleNamespace( + __file__=torch.__file__, version=SimpleNamespace(cuda=None, hip=hip) + ) + monkeypatch.setattr(extensions, "_load_torch_extension_tools", lambda: (fake_torch, None, None)) + with pytest.raises(RuntimeError, match="requires CUDA-enabled PyTorch"): + extensions.get_extensions() + + +def test_rope_only_wheel_rejects_missing_torch(monkeypatch): + monkeypatch.setenv("RL_KERNEL_ROPE_ONLY", "1") + monkeypatch.setattr(extensions, "_load_torch_extension_tools", lambda: (None, None, None)) + with pytest.raises(RuntimeError, match="PyTorch is unavailable"): + extensions.get_extensions() diff --git a/tests/ops/rope/test_multi_axis_rope.py b/tests/ops/rope/test_multi_axis_rope.py new file mode 100644 index 000000000..6cbcb2b85 --- /dev/null +++ b/tests/ops/rope/test_multi_axis_rope.py @@ -0,0 +1,443 @@ +# SPDX-License-Identifier: Apache-2.0 +"""Accuracy is independent FP64 complex algebra; invariance compares raw bytes.""" + +import gc +import math + +import pytest +import torch + +from rl_engine.ops.rope.multi_axis_rope import ( + build_multi_axis_cos_sin, + get_multi_axis_rope_op, + multi_axis_rope, + qwen_image_positions, +) + + +def assert_bits(a, b): + assert a.dtype == b.dtype and a.shape == b.shape + assert torch.equal(a.contiguous().view(torch.uint8), b.contiguous().view(torch.uint8)) + + +def complex_reference(x, cos, sin, *, transpose=False): + pairs = x.double().reshape(*x.shape[:-1], 64, 2) + z = torch.view_as_complex(pairs.contiguous()) + phase = torch.complex(cos.double(), sin.double()).unsqueeze(-2) + if transpose: + phase = phase.conj() + return torch.view_as_real(z * phase).flatten(-2) + + +@pytest.fixture(params=["reference", "cuda", "triton"]) +def implementation(request): + backend = request.param + device = "cpu" if backend == "reference" else "cuda" + if device == "cuda" and (not torch.cuda.is_available() or torch.version.hip is not None): + pytest.skip("NVIDIA CUDA required") + if backend == "cuda": + from rl_engine.backends.cuda.rope.multi_axis_rope import native_available + + if not native_available(): + pytest.skip("build native multi_axis_rope or use the focused CUDA validation runner") + if backend == "triton": + pytest.importorskip("triton") + return get_multi_axis_rope_op(backend, device=device), device + + +def inputs(device, dtype=torch.float32, *, batch=2, seq=7, heads=3, batched=False): + generator = torch.Generator().manual_seed(386) + x = torch.randn(batch, seq, heads, 128, generator=generator).to(device, dtype) + shape = (batch, seq, 3) if batched else (seq, 3) + pos = torch.randint(-42, 53, shape, generator=generator) + cos, sin = build_multi_axis_cos_sin(pos, device=device) + return x, cos, sin + + +@pytest.mark.parametrize("height,width", [(2, 4), (3, 5), (64, 64), (83, 83), (104, 58)]) +def test_positions_and_axis_frequencies(height, width): + image, text = qwen_image_positions(height, width, 3) + # Explicit flattened endpoints and a row boundary catch swapped H/W and odd-grid shifts. + assert image[0].tolist() == [0, -((height + 1) // 2), -((width + 1) // 2)] + assert image[-1].tolist() == [0, height // 2 - 1, width // 2 - 1] + assert image[width].tolist() == [0, -((height + 1) // 2) + 1, -((width + 1) // 2)] + assert text[:, 0].tolist() == list( + range(max(height // 2, width // 2), max(height // 2, width // 2) + 3) + ) + cos, sin = build_multi_axis_cos_sin(image[[0, -1]]) + # Independent scalar FP64 trig oracle, per-axis exponents reset at 8 and 36 pairs. + for row, position in enumerate(image[[0, -1]].tolist()): + offset = 0 + for axis, width_axis in enumerate((16, 56, 56)): + for pair in range(width_axis // 2): + angle = position[axis] / (10000.0 ** (2 * pair / width_axis)) + assert abs(cos[row, offset + pair].item() - math.cos(angle)) < 4e-6 + assert abs(sin[row, offset + pair].item() - math.sin(angle)) < 4e-6 + offset += width_axis // 2 + + +@pytest.mark.parametrize("dtype", [torch.float32, torch.float16, torch.bfloat16]) +@pytest.mark.parametrize("batched", [False, True]) +def test_forward_backward_accuracy(implementation, dtype, batched): + op, device = implementation + x, cos, sin = inputs(device, dtype, batched=batched) + x.requires_grad_() + grad = torch.randn_like(x) + result = op(x, cos, sin) + dx = torch.autograd.grad(result, x, grad)[0] + tolerance = {torch.float32: 8e-7, torch.float16: 0.002, torch.bfloat16: 0.02}[dtype] + torch.testing.assert_close( + result.double(), complex_reference(x, cos, sin), atol=tolerance, rtol=tolerance + ) + torch.testing.assert_close( + dx.double(), + complex_reference(grad, cos, sin, transpose=True), + atol=tolerance, + rtol=tolerance, + ) + + +def test_adjacent_pair_layout_not_rotate_half(implementation): + op, device = implementation + x = torch.arange(128, dtype=torch.float32, device=device).reshape(1, 1, 1, 128) + cos = torch.zeros(1, 64, device=device) + sin = torch.ones_like(cos) + y = op(x, cos, sin) + assert_bits(y[..., 0::2], -x[..., 1::2]) + assert_bits(y[..., 1::2], x[..., 0::2]) + + +def test_batch_position_padding_and_chunk_invariance(implementation): + op, device = implementation + x, cos, sin = inputs(device, batched=True, batch=3, seq=19) + x.requires_grad_() + grad = torch.randn_like(x) + y = op(x, cos, sin) + dx = torch.autograd.grad(y, x, grad)[0] + # Reordering samples with different positions must not reuse sample zero's table. + permutation = torch.tensor([2, 0, 1], device=device) + assert_bits(op(x[permutation], cos[permutation], sin[permutation]), y[permutation]) + for b in range(3): + for lo, hi in [(0, 1), (1, 8), (8, 19)]: + xx = x[b : b + 1, lo:hi].detach().requires_grad_() + yy = op(xx, cos[b, lo:hi], sin[b, lo:hi]) + dd = torch.autograd.grad(yy, xx, grad[b : b + 1, lo:hi])[0] + assert_bits(yy, y[b : b + 1, lo:hi]) + assert_bits(dd, dx[b : b + 1, lo:hi]) + padded = torch.cat((x.detach(), torch.randn(3, 5, 3, 128, device=device)), dim=1) + pc = torch.cat((cos, torch.ones(3, 5, 64, device=device)), dim=1) + ps = torch.cat((sin, torch.zeros(3, 5, 64, device=device)), dim=1) + assert_bits(op(padded, pc, ps)[:, :19], y) + + +def test_noncontiguous_input_tables_and_gradient(implementation): + op, device = implementation + base, cos, sin = inputs(device, batched=True, seq=14, heads=4) + x = base[:, ::2, ::2].detach().requires_grad_() + c, s = cos[:, ::2], sin[:, ::2] + grad = torch.randn_like(base)[:, ::2, ::2] + result = op(x, c, s) + dx = torch.autograd.grad(result, x, grad)[0] + xx = x.detach().contiguous().requires_grad_() + yy = op(xx, c.contiguous(), s.contiguous()) + dd = torch.autograd.grad(yy, xx, grad.contiguous())[0] + assert_bits(result, yy) + assert_bits(dx, dd) + + +@pytest.mark.parametrize("batch,seq", [(0, 7), (2, 0), (0, 0)]) +def test_empty_forward_backward(implementation, batch, seq): + op, device = implementation + x, cos, sin = inputs(device, batch=batch, seq=seq) + x.requires_grad_() + result = op(x, cos, sin) + result.sum().backward() + assert result.shape == x.shape and x.grad.shape == x.shape + + +def test_second_derivative(implementation): + op, device = implementation + x, cos, sin = inputs(device, batch=1, seq=2) + x.requires_grad_() + y = op(x, cos, sin) + first = torch.autograd.grad(y.square().sum(), x, create_graph=True)[0] + second = torch.autograd.grad(first.sum(), x)[0] + # R^T R = (cos^2 + sin^2) I, including the small table rounding error. + expected = ( + (2 * (cos.double().square() + sin.double().square())) + .repeat_interleave(2, -1)[None, :, None, :] + .expand_as(x) + ) + torch.testing.assert_close(second.double(), expected, atol=8e-7, rtol=8e-7) + + +@pytest.mark.parametrize( + "problem", ["shape", "dtype", "table_shape", "table_dtype", "table_grad", "device"] +) +def test_reject_invalid_inputs(implementation, problem): + op, device = implementation + x, cos, sin = inputs(device) + if problem == "shape": + x = x[..., :64] + elif problem == "dtype": + x = x.double() + elif problem == "table_shape": + sin = sin[:-1] + elif problem == "table_dtype": + cos = cos.half() + elif problem == "table_grad": + cos.requires_grad_() + elif problem == "device": + sin = torch.empty_like(sin, device="meta") + with pytest.raises((ValueError, TypeError)): + op(x, cos, sin) + + +def test_explicit_backend_and_reference_dispatch(monkeypatch): + from rl_engine.runtime.registry import KernelRegistry + + op = KernelRegistry().get_op("multi_axis_rope", device="cpu") + assert op.trace()["actual_backend"].startswith("rlkernel.reference.") + from rl_engine.backends.cuda.rope import multi_axis_rope as cuda_module + + monkeypatch.setattr(cuda_module, "native_available", lambda: False) + with pytest.raises(RuntimeError, match="unavailable"): + get_multi_axis_rope_op("cuda") + with pytest.raises(ValueError, match="unknown"): + get_multi_axis_rope_op("typo") + x, cos, sin = inputs("cpu") + assert_bits(multi_axis_rope(x, cos, sin), op(x, cos, sin)) + + +def test_cuda_registry_uses_native_when_present(): + from rl_engine.backends.cuda.rope.multi_axis_rope import native_available + from rl_engine.runtime.registry import KernelRegistry + + if not torch.cuda.is_available() or not native_available(): + pytest.skip("native CUDA multi_axis_rope extension required") + op = KernelRegistry().get_op("multi_axis_rope", device="cuda") + assert op.trace()["actual_backend"] == "rlkernel.cuda.qwenimage_multi_axis_rope.v1" + x, cos, sin = inputs("cuda") + torch.testing.assert_close( + op(x, cos, sin).double(), complex_reference(x, cos, sin), atol=8e-7, rtol=8e-7 + ) + + +def test_cuda_registry_falls_back_to_triton(monkeypatch): + from rl_engine.runtime.registry import KernelRegistry + + if not torch.cuda.is_available() or torch.version.hip is not None: + pytest.skip("NVIDIA CUDA required") + pytest.importorskip("triton") + from rl_engine.backends.cuda.rope import multi_axis_rope as cuda_module + + monkeypatch.setattr(cuda_module, "native_available", lambda: False) + op = KernelRegistry().get_op("multi_axis_rope", device="cuda") + assert op.trace()["actual_backend"] == "rlkernel.triton.qwenimage_multi_axis_rope.v1" + x, cos, sin = inputs("cuda") + torch.testing.assert_close( + op(x, cos, sin).double(), complex_reference(x, cos, sin), atol=8e-7, rtol=8e-7 + ) + + +@pytest.mark.parametrize( + "bad", + [torch.zeros(2, 3), torch.zeros(2, 4, dtype=torch.int64), torch.full((2, 3), 4096)], +) +def test_reject_invalid_positions(bad): + with pytest.raises((TypeError, ValueError)): + build_multi_axis_cos_sin(bad) + + +@pytest.mark.parametrize("dtype", [torch.float32, torch.float16, torch.bfloat16]) +def test_cuda_triton_byte_parity_and_graph(dtype): + if not torch.cuda.is_available() or torch.version.hip is not None: + pytest.skip("NVIDIA CUDA required") + from rl_engine.backends.cuda.rope.multi_axis_rope import native_available + + if not native_available(): + pytest.skip("native multi_axis_rope extension required") + pytest.importorskip("triton") + cuda, triton = get_multi_axis_rope_op("cuda"), get_multi_axis_rope_op("triton") + x, cos, sin = inputs("cuda", dtype, batched=True, seq=17, heads=24) + x.requires_grad_() + grad = torch.randn_like(x) + yc, yt = cuda(x, cos, sin), triton(x, cos, sin) + dc = torch.autograd.grad(yc, x, grad)[0] + dt = torch.autograd.grad(yt, x, grad)[0] + assert_bits(yc, yt) + assert_bits(dc, dt) + reference = get_multi_axis_rope_op("reference") + yr = reference(x, cos, sin) + dr = torch.autograd.grad(yr, x, grad)[0] + assert_bits(yc, yr) + assert_bits(dc, dr) + # Capture forward + autograd after warmup on a side stream. + for op in (cuda, triton): + stream = torch.cuda.Stream() + stream.wait_stream(torch.cuda.current_stream()) + with torch.cuda.stream(stream): + static_x = x.detach().clone().requires_grad_() + static_grad = grad.clone() + for _ in range(3): + torch.autograd.grad(op(static_x, cos, sin), static_x, static_grad) + torch.cuda.current_stream().wait_stream(stream) + graph = torch.cuda.CUDAGraph() + gc.collect() + was_enabled = gc.isenabled() + gc.disable() + try: + with torch.cuda.graph(graph, stream=stream): + result = op(static_x, cos, sin) + derivative = torch.autograd.grad(result, static_x, static_grad)[0] + finally: + if was_enabled: + gc.enable() + for _ in range(3): + graph.replay() + torch.cuda.synchronize() + assert_bits(result, yc) + assert_bits(derivative, dc) + # Replay must consume updated static buffers, not just retain captured outputs. + with torch.no_grad(): + static_x.copy_(x * 2) + static_grad.copy_(grad * 0.5) + graph.replay() + torch.cuda.synchronize() + expected = op(static_x, cos, sin) + expected_dx = torch.autograd.grad(expected, static_x, static_grad)[0] + assert_bits(result, expected) + assert_bits(derivative, expected_dx) + + +def test_native_entrypoint_checks_and_nondefault_stream(): + if not torch.cuda.is_available() or torch.version.hip is not None: + pytest.skip("NVIDIA CUDA required") + from rl_engine.backends.cuda.rope.multi_axis_rope import native_available + + if not native_available(): + pytest.skip("native multi_axis_rope extension required") + x, cos, sin = inputs("cuda") + native = torch.ops.rl_kernel_native.multi_axis_rope + for c, s in [(cos.half(), sin), (cos, sin[:-1]), (cos.cpu(), sin)]: + with pytest.raises(RuntimeError): + native(x, c, s, False) + stream = torch.cuda.Stream() + stream.wait_stream(torch.cuda.current_stream()) + with torch.cuda.stream(stream): + xx = x * 3 + y = native(xx, cos, sin, False) + torch.cuda.current_stream().wait_stream(stream) + torch.testing.assert_close(y.double(), complex_reference(xx, cos, sin), atol=2e-6, rtol=2e-6) + + +def test_device_guard(): + if torch.cuda.device_count() < 2 or torch.version.hip is not None: + pytest.skip("requires two NVIDIA GPUs") + from rl_engine.backends.cuda.rope.multi_axis_rope import native_available + + if not native_available(): + pytest.skip("native multi_axis_rope extension required") + with torch.cuda.device(0): + x, cos, sin = inputs("cuda:1") + for backend in ("cuda", "triton"): + y = get_multi_axis_rope_op(backend)(x, cos, sin) + assert y.device == x.device + assert torch.cuda.current_device() == 0 + torch.testing.assert_close( + y.double(), complex_reference(x, cos, sin), atol=8e-7, rtol=8e-7 + ) + + +@pytest.mark.parametrize("default_device", ["meta", "cuda"]) +def test_canonical_tables_ignore_default_device(default_device): + if default_device == "cuda" and not torch.cuda.is_available(): + pytest.skip("CUDA context requires CUDA") + positions, text = qwen_image_positions(3, 5, 7) + expected = build_multi_axis_cos_sin(positions) + with torch.device(default_device): + image_again, text_again = qwen_image_positions(3, 5, 7) + actual = build_multi_axis_cos_sin(positions, device="cpu") + assert image_again.device.type == "cpu" and text_again.device.type == "cpu" + assert_bits(image_again, positions) + assert_bits(text_again, text) + for a, b in zip(actual, expected, strict=True): + assert_bits(a, b) + + +def test_packed_qkv_and_independent_table_strides(implementation): + op, device = implementation + packed = torch.randn(2, 17, 3, 3, 256, device=device) + x = packed[:, :, 0, :, ::2].detach().requires_grad_() + cos_base = torch.randn(2, 17, 128, device=device) + cos = cos_base[..., ::2] + sin = torch.randn(17, 2, 64, device=device).transpose(0, 1) + grad = torch.randn_like(x) + result = op(x, cos, sin) + dx = torch.autograd.grad(result, x, grad)[0] + assert result.is_contiguous() + torch.testing.assert_close( + result.double(), complex_reference(x, cos, sin), atol=2e-6, rtol=2e-6 + ) + torch.testing.assert_close( + dx.double(), complex_reference(grad, cos, sin, transpose=True), atol=2e-6, rtol=2e-6 + ) + if device == "cuda": + with torch.profiler.profile(activities=[torch.profiler.ProfilerActivity.CPU]) as profile: + op(x, cos, sin) + copies = {event.key for event in profile.key_averages()} & { + "aten::contiguous", + "aten::clone", + "aten::copy_", + } + assert not copies, f"RoPE materialized its strided inputs: {copies}" + + +@pytest.mark.parametrize("backend", ["cuda", "triton"]) +def test_dispatcher_fake_autograd_and_fullgraph_compile(backend): + if not torch.cuda.is_available() or torch.version.hip is not None: + pytest.skip("NVIDIA CUDA required") + if backend == "cuda": + from rl_engine.backends.cuda.rope.multi_axis_rope import _rotation, native_available + + if not native_available(): + pytest.skip("build native RoPE library") + else: + pytest.importorskip("triton") + from rl_engine.backends.shared.triton.rope.multi_axis_rope import _rotation + op = get_multi_axis_rope_op(backend) + x, cos, sin = inputs("cuda", batched=True, seq=5, heads=2) + x.requires_grad_() + checks = torch.library.opcheck(_rotation, (x, cos, sin, False)) + assert all(value == "SUCCESS" for value in checks.values()) + compiled = torch.compile(op, fullgraph=True) + grad = torch.randn_like(x) + expected = op(x, cos, sin) + expected_dx = torch.autograd.grad(expected, x, grad)[0] + result = compiled(x, cos, sin) + dx = torch.autograd.grad(result, x, grad)[0] + assert_bits(result, expected) + assert_bits(dx, expected_dx) + + +@pytest.mark.parametrize("backend", ["cuda", "triton"]) +@pytest.mark.ws1 +def test_standard_ws1_four_judgments(backend): + if not torch.cuda.is_available() or torch.version.hip is not None: + pytest.skip("NVIDIA CUDA required") + if backend == "cuda": + from rl_engine.backends.cuda.rope.multi_axis_rope import native_available + + if not native_available(): + pytest.skip("build native RoPE library") + else: + pytest.importorskip("triton") + from rl_engine.validation.operators.qwen_image_rope import run_multi_axis_rope_ws1 + + result = run_multi_axis_rope_ws1(backend) + assert result["passed"] + assert result["trace"]["actual_backend"].startswith(f"rlkernel.{backend}.") + for phase in ("forward", "gradient"): + assert result[phase]["accuracy_reports"] + assert result[phase]["invariance_reports"] + assert all(report["passed"] for report in result[phase]["invariance_reports"]) diff --git a/tests/validation/operators/test_operator_inputs.py b/tests/validation/operators/test_operator_inputs.py index 9e5015bf4..b2c054b18 100644 --- a/tests/validation/operators/test_operator_inputs.py +++ b/tests/validation/operators/test_operator_inputs.py @@ -103,6 +103,24 @@ def test_cp_attention_operator_spec_registers_backward_grad_inputs(): assert candidate.name == "pytorch-cp_attention" +def test_multi_axis_rope_operator_spec_and_seeded_inputs(): + args = _args(op="multi_axis_rope", batch=2, seq=7, n_heads=3, input_mode="random") + case = make_operator_case(args, torch.float32, torch.device("cpu")) + inputs = make_operator_inputs("multi_axis_rope", args, torch.float32, torch.device("cpu")) + repeated = make_operator_inputs("multi_axis_rope", args, torch.float32, torch.device("cpu")) + assert case.op_class == "elementwise" + assert case.grad_input_names == ("x",) + assert inputs["x"].shape == (2, 7, 3, 128) + assert inputs["cos"].shape == inputs["sin"].shape == (7, 64) + for name in inputs: + assert torch.equal(inputs[name], repeated[name]) + assert operator_shape_name("multi_axis_rope", args) == "2x7x3x128" + with pytest.raises(ValueError, match="128"): + make_operator_inputs( + "multi_axis_rope", _args(head_dim=64), torch.float32, torch.device("cpu") + ) + + def test_constant_linear_logp_inputs_match_operator_contract(): args = _args(input_mode="constant", constant_value=0.5, token_value=3) inputs = make_operator_inputs("linear_logp", args, torch.float32, torch.device("cpu")) diff --git a/tools/validation/operators/check_multi_axis_rope_install.py b/tools/validation/operators/check_multi_axis_rope_install.py new file mode 100644 index 000000000..56cfe2ccb --- /dev/null +++ b/tools/validation/operators/check_multi_axis_rope_install.py @@ -0,0 +1,67 @@ +# SPDX-License-Identifier: Apache-2.0 +"""Run from outside the checkout after installing a RoPE-enabled wheel. + +Does not change sys.path or compile/load a JIT extension. --package-root makes +the expected installation location an explicit assertion in package validation. +""" + +import argparse +import importlib.util +import json +from pathlib import Path + +import torch + +import rl_engine +from rl_engine.ops.rope.multi_axis_rope import ( + build_multi_axis_cos_sin, + get_multi_axis_rope_op, + qwen_image_positions, +) + + +def main(): + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--package-root", type=Path, required=True) + parser.add_argument("--output", type=Path, required=True) + args = parser.parse_args() + package = Path(rl_engine.__file__).resolve().parent + assert package.is_relative_to(args.package_root.resolve()), package + license_path = package / "backends/shared/triton/rope/FLASH_ATTENTION_LICENSE" + assert "Redistribution and use" in license_path.read_text(encoding="utf-8") + positions, _ = qwen_image_positions(3, 5, 0) + cos, sin = build_multi_axis_cos_sin(positions, device="cuda") + packed = torch.randn(2, 15, 3, 3, 128, device="cuda", dtype=torch.bfloat16) + x = packed[:, :, 0].detach().requires_grad_() + grad = torch.randn_like(x) + gold = get_multi_axis_rope_op("reference")(x, cos, sin) + gold_dx = torch.autograd.grad(gold, x, grad)[0] + outcomes = {} + for backend in ("cuda", "triton"): + op = get_multi_axis_rope_op(backend, device=x.device) + for mode, fn in (("eager", op), ("compiled", torch.compile(op, fullgraph=True))): + result = fn(x, cos, sin) + dx = torch.autograd.grad(result, x, grad)[0] + assert torch.equal(result.view(torch.uint8), gold.contiguous().view(torch.uint8)) + assert torch.equal(dx.view(torch.uint8), gold_dx.contiguous().view(torch.uint8)) + outcomes[f"{backend}_{mode}"] = "passed" + outcomes[f"{backend}_trace"] = op.trace() + native = importlib.util.find_spec("rl_engine._rope_cuda") + report = { + "package": str(package), + "companion_library": native.origin if native else None, + "native_dispatcher_registered": hasattr(torch.ops.rl_kernel_native, "multi_axis_rope"), + "torch": torch.__version__, + "gpu": torch.cuda.get_device_name(), + "license_present": True, + "packed_qkv_stride": list(x.stride()), + "checks": outcomes, + "passed": True, + } + args.output.parent.mkdir(parents=True, exist_ok=True) + args.output.write_text(json.dumps(report, indent=2) + "\n", encoding="utf-8") + print(json.dumps(report, indent=2)) + + +if __name__ == "__main__": + main() diff --git a/tools/validation/operators/check_qwenimage_rope_upstream.py b/tools/validation/operators/check_qwenimage_rope_upstream.py new file mode 100644 index 000000000..153000904 --- /dev/null +++ b/tools/validation/operators/check_qwenimage_rope_upstream.py @@ -0,0 +1,99 @@ +# SPDX-License-Identifier: Apache-2.0 +"""Audit table and rotation semantics against a checksum-pinned upstream file. + +Supply transformer_qwenimage.py from Diffusers revision +40f63690da440e8c382ab53e67508835deb1ed2f. Only the two audited symbol definitions +are executed, after verifying the source hash. No Diffusers installation or model +checkpoint is required. This checks CPU semantics, not full-model integration. +""" + +from __future__ import annotations + +import argparse +import ast +import functools +import hashlib +import json +import logging +import sys +from pathlib import Path + +import torch + +ROOT = Path(__file__).resolve().parents[3] +sys.path.insert(0, str(ROOT)) +SOURCE_SHA256 = "0bad11901deae69f05aec01fd464e51e38582d455068f2851ee665036179321f" + + +def main(): + from rl_engine.ops.rope.multi_axis_rope import ( + build_multi_axis_cos_sin, + get_multi_axis_rope_op, + qwen_image_positions, + ) + + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--source", type=Path, required=True) + parser.add_argument("--output", type=Path, required=True) + args = parser.parse_args() + source = args.source.read_text(encoding="utf-8") + digest = hashlib.sha256(source.encode()).hexdigest() + if digest != SOURCE_SHA256: + raise ValueError("upstream source does not match the pinned revision checksum") + symbols = {"QwenEmbedRope", "apply_rotary_emb_qwen"} + nodes = [n for n in ast.parse(source).body if getattr(n, "name", None) in symbols] + if len(nodes) != 2: + raise ValueError("missing audited upstream definitions") + namespace = { + "torch": torch, + "nn": torch.nn, + "lru_cache_unless_export": functools.lru_cache, + "logger": logging.getLogger("upstream_audit"), + } + exec( + compile(ast.Module(body=nodes, type_ignores=[]), str(args.source), "exec"), + namespace, + ) + embed = namespace["QwenEmbedRope"](theta=10000, axes_dim=[16, 56, 56], scale_rope=True) + apply_upstream = namespace["apply_rotary_emb_qwen"] + op = get_multi_axis_rope_op("reference") + torch.manual_seed(386) + results = [] + for height, width in ((2, 3), (3, 5), (64, 64), (83, 83), (104, 58)): + positions = qwen_image_positions(height, width, 7) + upstream_tables = embed((1, height, width), device=torch.device("cpu"), max_txt_seq_len=7) + for stream, pos, upstream in zip( + ("image", "text"), positions, upstream_tables, strict=True + ): + cos, sin = build_multi_axis_cos_sin(pos) + assert torch.equal(cos, upstream.real) and torch.equal(sin, upstream.imag) + for dtype in (torch.float32, torch.float16, torch.bfloat16): + x = torch.randn(1, len(pos), 2, 128, dtype=dtype, requires_grad=True) + grad = torch.randn_like(x) + expected = apply_upstream(x, upstream, use_real=False) + actual = op(x, cos, sin) + eg = torch.autograd.grad(expected, x, grad)[0] + ag = torch.autograd.grad(actual, x, grad)[0] + tol = {torch.float32: 8e-7, torch.float16: 0.002, torch.bfloat16: 0.02}[dtype] + torch.testing.assert_close(actual, expected, atol=tol, rtol=tol) + torch.testing.assert_close(ag, eg, atol=tol, rtol=tol) + results.append( + { + "grid": [height, width], + "stream": stream, + "dtype": str(dtype), + "table_bitwise": True, + "forward_max_abs": (actual.float() - expected.float()).abs().max().item(), + "dx_max_abs": (ag.float() - eg.float()).abs().max().item(), + } + ) + args.output.parent.mkdir(parents=True, exist_ok=True) + args.output.write_text(json.dumps({"source_sha256": digest, "cases": results}, indent=2) + "\n") + print( + f"{len(results)} upstream forward/backward cases passed; " + "all image/text tables exactly equal" + ) + + +if __name__ == "__main__": + main() diff --git a/tools/validation/operators/validate_multi_axis_rope.py b/tools/validation/operators/validate_multi_axis_rope.py new file mode 100644 index 000000000..d18ef7cbb --- /dev/null +++ b/tools/validation/operators/validate_multi_axis_rope.py @@ -0,0 +1,85 @@ +# SPDX-License-Identifier: Apache-2.0 +"""Compile only the production RoPE CUDA source for focused GPU validation. + +Loads the production TORCH_LIBRARY registration directly. Normal package builds +include the same source in rl_engine._C. No generated binding or monkeypatch is +used. This validates the operator library, not every unrelated native operator. +""" + +from __future__ import annotations + +import argparse +import json +import os +import runpy +import sys +from pathlib import Path + +ROOT = Path(__file__).resolve().parents[3] +sys.path.insert(0, str(ROOT)) + + +def load_focused_extension(): + import torch + from torch.utils.cpp_extension import CUDA_HOME, load + + from rl_engine.backends.cuda.rope.multi_axis_rope import native_available + + if native_available(): + print("Using the installed production RoPE dispatcher library") + return + + build = ROOT / "artifacts" / "multi_axis_rope_native" + build.mkdir(parents=True, exist_ok=True) + # Conda CUDA toolkits use lib/, whereas the Windows SDK convention is lib/x64/. + link_flags = [] + if os.name == "nt" and CUDA_HOME and (Path(CUDA_HOME) / "lib/cudart.lib").is_file(): + link_flags.append(f'/LIBPATH:"{Path(CUDA_HOME) / "lib"}"') + os.environ["PATH"] = str(Path(sys.executable).parent) + os.pathsep + os.environ["PATH"] + load( + name="rlk_multi_axis_rope_native", + sources=[str(ROOT / "csrc/cuda/rope/multi_axis_rope.cu")], + build_directory=str(build), + extra_cflags=(["/O2", "/std:c++17"] if os.name == "nt" else ["-O3", "-std=c++17"]), + extra_cuda_cflags=["-O3", "-std=c++17", "--fmad=false"], + extra_ldflags=link_flags, + is_python_module=False, + verbose=True, + ) + assert hasattr(torch.ops.rl_kernel_native, "multi_axis_rope") + print( + f"Focused CUDA source build; torch={torch.__version__}, GPU={torch.cuda.get_device_name()}" + ) + + +def main(): + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--benchmark", action="store_true") + parser.add_argument("--ws1", choices=("cuda", "triton")) + parser.add_argument("--output", type=Path) + args, remaining = parser.parse_known_args() + load_focused_extension() + if args.ws1: + from rl_engine.validation.operators.qwen_image_rope import run_multi_axis_rope_ws1 + + result = run_multi_axis_rope_ws1(args.ws1) + if args.output: + args.output.parent.mkdir(parents=True, exist_ok=True) + args.output.write_text(json.dumps(result, indent=2) + "\n", encoding="utf-8") + print(f"WS1 four judgments backend={args.ws1}: passed={result['passed']}") + raise SystemExit(0 if result["passed"] else 1) + elif args.benchmark: + sys.argv = ["benchmark_multi_axis_rope", *remaining] + if args.output: + sys.argv.extend(["--output", str(args.output)]) + runpy.run_module("benchmarks.operators.rope.benchmark_multi_axis_rope", run_name="__main__") + else: + import pytest + + raise SystemExit( + pytest.main([str(ROOT / "tests/ops/rope/test_multi_axis_rope.py"), *remaining]) + ) + + +if __name__ == "__main__": + main() From a0f9276195b020681ea01c5ca47bc33fa260277a Mon Sep 17 00:00:00 2001 From: dandan20031031 <2425926169@qq.com> Date: Wed, 28 Oct 2026 21:49:14 +0800 Subject: [PATCH 2/5] test(rope): qualify the complete CUDA package on WSL2 Add an explicit full-extension GPU gate and CPU CI for test-qwenimage. Record WSL2 wheel/import, WS1, layout, compiler, CPU-only, memory-check and mature-baseline evidence; expose the loaded full native library in the installation probe. Refs #386. Signed-off-by: dandan20031031 <2425926169@qq.com> --- .github/workflows/qwenimage-rope.yml | 82 + ci/scripts/run_qwenimage_rope_ws1.sh | 89 + docs/operators/multi-axis-rope.md | 23 + .../README.md | 132 + .../benchmark-bfloat16.json | 2781 +++++++++++++++++ .../cpu-contracts.txt | 20 + .../cpu-environment.json | 7 + .../full-extension.json | 22 + .../full-inplace-build.txt | 20 + .../full-wheel-build.txt | 10 + .../gpu-environment.json | 15 + .../installed-wheel.json | 38 + .../layout.txt | 2 + .../manifest.json | 58 + .../memcheck.txt | 4 + .../operator-cuda.json | 70 + .../operator-tests.xml | 1 + .../operator-triton.json | 70 + .../package.json | 11 + .../precommit.txt | 7 + .../python-packages.json | 114 + .../quality.txt | 34 + .../regression.txt | 10 + .../tests-and-ws1.txt | 11 + .../upstream.json | 335 ++ .../ws1-cuda.json | 1962 ++++++++++++ .../ws1-triton.json | 1962 ++++++++++++ .../check_multi_axis_rope_install.py | 2 + 28 files changed, 7892 insertions(+) create mode 100644 .github/workflows/qwenimage-rope.yml create mode 100644 ci/scripts/run_qwenimage_rope_ws1.sh create mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/README.md create mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/benchmark-bfloat16.json create mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu-contracts.txt create mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu-environment.json create mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/full-extension.json create mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/full-inplace-build.txt create mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/full-wheel-build.txt create mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/gpu-environment.json create mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/installed-wheel.json create mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/layout.txt create mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/manifest.json create mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/memcheck.txt create mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/operator-cuda.json create mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/operator-tests.xml create mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/operator-triton.json create mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/package.json create mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/precommit.txt create mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/python-packages.json create mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/quality.txt create mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/regression.txt create mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/tests-and-ws1.txt create mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/upstream.json create mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/ws1-cuda.json create mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/ws1-triton.json diff --git a/.github/workflows/qwenimage-rope.yml b/.github/workflows/qwenimage-rope.yml new file mode 100644 index 000000000..7819acb9d --- /dev/null +++ b/.github/workflows/qwenimage-rope.yml @@ -0,0 +1,82 @@ +# SPDX-License-Identifier: Apache-2.0 +name: Qwen-Image RoPE + +on: + pull_request: + branches: [test-qwenimage] + types: [opened, synchronize, reopened, ready_for_review] + workflow_dispatch: + +permissions: + contents: read + +jobs: + cpu-contracts: + runs-on: ubuntu-latest + strategy: + fail-fast: false + matrix: + python-version: ['3.10', '3.12'] + steps: + - uses: actions/checkout@v4 + - uses: actions/setup-python@v5 + with: + python-version: ${{ matrix.python-version }} + cache: pip + - name: Install CPU reference dependencies + run: | + python -m pip install pip==25.3 setuptools==80.9.0 wheel==0.45.1 + python -m pip install torch==2.9.1 --index-url https://download.pytorch.org/whl/cpu + python -m pip install --no-build-isolation -e . pytest + - name: Run CPU contracts, input registration and build-profile checks + env: + OMP_NUM_THREADS: '1' + MKL_NUM_THREADS: '1' + run: | + python -m pytest -q -rs \ + tests/ops/rope/test_multi_axis_rope.py \ + tests/validation/operators/test_operator_inputs.py \ + tests/build/test_multi_axis_rope_build.py \ + --junitxml=artifacts/rope-cpu.xml + - uses: actions/upload-artifact@v4 + if: always() + with: + name: rope-cpu-${{ matrix.python-version }} + path: artifacts/rope-cpu.xml + + static-and-docs: + runs-on: ubuntu-latest + steps: + - uses: actions/checkout@v4 + with: + fetch-depth: 0 + - uses: actions/setup-python@v5 + with: + python-version: '3.12' + cache: pip + - run: python -m pip install pre-commit==4.3.0 -r requirements/docs.txt + - name: Check changed files + run: | + python - <<'PY' + import json + import os + import re + import subprocess + import sys + + with open(os.environ['GITHUB_EVENT_PATH'], encoding='utf-8') as stream: + event = json.load(stream) + base = event.get('pull_request', {}).get('base', {}).get('sha') + if base: + assert re.fullmatch(r'[0-9a-f]{40}', base) + else: + base = 'HEAD^' + changed = subprocess.check_output([ + 'git', 'diff', '--name-only', '--diff-filter=ACMR', '-z', base, 'HEAD' + ]).decode('utf-8').split('\0') + files = [('./' + path if path.startswith('-') else path) for path in changed if path] + if files: + subprocess.check_call([sys.executable, '-m', 'pre_commit', 'run', '--files', *files]) + PY + - name: Build documentation + run: mkdocs build --strict -f mkdocs.yaml diff --git a/ci/scripts/run_qwenimage_rope_ws1.sh b/ci/scripts/run_qwenimage_rope_ws1.sh new file mode 100644 index 000000000..447e93f14 --- /dev/null +++ b/ci/scripts/run_qwenimage_rope_ws1.sh @@ -0,0 +1,89 @@ +#!/usr/bin/env bash +# SPDX-License-Identifier: Apache-2.0 +# Operator-only gate: requires an ordinary, complete rl_engine._C installation. +set -euo pipefail + +ROPE_ROOT="$(cd "$(dirname "$0")/../.." && pwd)" +cd "$ROPE_ROOT" +ROPE_PYTHON="${ROPE_PYTHON:-python3}" +ROPE_WS1_OUTPUT="${ROPE_WS1_OUTPUT:-${TMPDIR:-/tmp}/rlk-qwenimage-rope}" +mkdir -p "$ROPE_WS1_OUTPUT" +export RL_KERNEL_REQUIRE_EXT=1 + +"$ROPE_PYTHON" - "$ROPE_WS1_OUTPUT/environment.json" <<'PY' +import hashlib +import json +import platform +import sys +from pathlib import Path + +import torch +import triton + +from build_tools.envs import env_flag +from rl_engine import _C + +if env_flag("RL_KERNEL_ROPE_ONLY"): + raise SystemExit("The complete _C build is required; unset RL_KERNEL_ROPE_ONLY") +if env_flag("KERNEL_ALIGN_USE_FAST_MATH"): + raise SystemExit("This numerical profile does not qualify fast-math builds") +if torch.version.hip is not None or not torch.cuda.is_available(): + raise SystemExit("A real NVIDIA CUDA device is required") +if not torch.cuda.is_bf16_supported(): + raise SystemExit("The WS1 BF16 profile requires BF16-capable hardware") +if not hasattr(torch.ops.rl_kernel_native, "multi_axis_rope"): + raise SystemExit("The imported _C library does not register multi_axis_rope") +library = Path(_C.__file__).resolve() +report = { + "scope": "Qwen-Image multi_axis_rope operator; ordinary full _C build", + "python": platform.python_version(), + "platform": platform.platform(), + "torch": torch.__version__, + "cuda": torch.version.cuda, + "triton": triton.__version__, + "gpu": torch.cuda.get_device_name(), + "capability": list(torch.cuda.get_device_capability()), + "native_library": str(library), + "native_library_sha256": hashlib.sha256(library.read_bytes()).hexdigest(), +} +Path(sys.argv[1]).write_text(json.dumps(report, indent=2) + "\n", encoding="utf-8") +print(json.dumps(report)) +PY + +"$ROPE_PYTHON" -m pytest tests/ops/rope/test_multi_axis_rope.py -q -rs \ + --junitxml="$ROPE_WS1_OUTPUT/operator-tests.xml" + +# GPU/backward/compile cases must execute. Only the two-device guard may skip +# on a single-GPU runner; missing CUDA/Triton support is not acceptance evidence. +"$ROPE_PYTHON" - "$ROPE_WS1_OUTPUT/operator-tests.xml" <<'PY' +import sys +import xml.etree.ElementTree as ET + +root = ET.parse(sys.argv[1]).getroot() +unexpected = [case.get("name") for case in root.iter("testcase") + if case.find("skipped") is not None and case.get("name") != "test_device_guard"] +if unexpected: + raise SystemExit(f"Required RoPE tests skipped: {unexpected}") +PY + +for backend in cuda triton; do + "$ROPE_PYTHON" tools/validation/operators/check_operator.py \ + --op multi_axis_rope --candidate "$backend" --device cuda --dtype bf16 \ + --batch 2 --seq 7 --n-heads 3 --check-grad --json \ + > "$ROPE_WS1_OUTPUT/operator-$backend.json" + "$ROPE_PYTHON" tools/validation/operators/validate_multi_axis_rope.py \ + --ws1 "$backend" --output "$ROPE_WS1_OUTPUT/ws1-$backend.json" +done + +"$ROPE_PYTHON" - "$ROPE_WS1_OUTPUT" <<'PY' +import json +import sys +from pathlib import Path + +root = Path(sys.argv[1]) +for backend in ("cuda", "triton"): + report = json.loads((root / f"ws1-{backend}.json").read_text(encoding="utf-8")) + if not report["passed"] or report["trace"]["actual_backend"] != f"rlkernel.{backend}.qwenimage_multi_axis_rope.v1": + raise SystemExit(f"WS1 failed or selected an unexpected backend: {backend}") +print("Qwen-Image RoPE full-extension operator gate passed") +PY diff --git a/docs/operators/multi-axis-rope.md b/docs/operators/multi-axis-rope.md index fa908a03e..16899a755 100644 --- a/docs/operators/multi-axis-rope.md +++ b/docs/operators/multi-axis-rope.md @@ -172,6 +172,29 @@ with C1's elementwise tolerance profile and C2's logical configuration transform This gate qualifies the Qwen-Image operator and does not certify the Dense Qwen3 chain, distributed execution, or model assembly. +For a complete native package installation on Linux, run the dedicated gate: + +```bash +export CUDA_HOME=/usr/local/cuda-12.8 +export PATH="$CUDA_HOME/bin:$PATH" +RL_KERNEL_REQUIRE_EXT=1 python -m pip install --no-build-isolation -e . +ROPE_PYTHON=python bash ci/scripts/run_qwenimage_rope_ws1.sh +``` + +The gate requires a real NVIDIA GPU, Triton and an ordinary `rl_engine._C` +containing the native RoPE registration. It rejects the operator-only profile +and fast-math mode, checks required tests actually execute, and permits only +the two-device guard to skip on a single-GPU host. Results include the loaded +binary fingerprint, ordinary operator CLI checks and both WS1 backend reports. +Set `ROPE_WS1_OUTPUT` to select the artifact directory. + +On WSL2, keep the checkout and Python environment in the Linux filesystem. +GNU linking of the package's CUDA driver calls may need +`LIBRARY_PATH=/usr/lib/wsl/lib`. Record the WSL kernel and physical GPU in the +evidence; a WSL2 profile does not qualify a bare-metal Linux server or another +GPU architecture. The `Qwen-Image RoPE` workflow runs CPU contracts and changed-file +checks for PRs targeting `test-qwenimage`; GPU qualification uses the gate above. + Tests cover independent forward/VJP accuracy, adjacent-pair layout, shared and batched positions, odd and rectangular grids, batch/padding/chunk invariance, non-contiguous tensors, packed QKV views without copies, default-device contexts, diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/README.md b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/README.md new file mode 100644 index 000000000..ddc879318 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/README.md @@ -0,0 +1,132 @@ +# Qwen-Image RoPE: full native package qualification on WSL2 + +Recorded 2026-10-28 for [PR #623](https://github.com/RL-Align/RL-Kernel/pull/623). +Numerical/native sources are based on `ac17ab87e8d93dabe93dfdc134e9853d47de7af1`; +this follow-up adds the reproducible gate, CPU CI routing and package provenance. +Source fingerprints are in `manifest.json`. Historical Windows measurements stay +in the separate `qwenimage_multi_axis_rope_v2_sm120` directory. + +## Environment and build + +| Item | Recorded value | +| --- | --- | +| Linux distribution | Ubuntu 24.04.5 LTS, x86_64, glibc 2.39 | +| Virtualization | WSL 3.0.3 on Windows 11 | +| Linux kernel | 6.18.54.1-microsoft-standard-WSL2 | +| GPU / driver | RTX 5080, SM120, 16 GB / Windows driver 617.42 | +| Python / PyTorch / Triton | 3.12.3 / 2.9.1+cu128 / 3.5.1 | +| CUDA compiler / host compiler | CUDA 12.8.93 / GCC 13.3.0 | +| Compute Sanitizer | 2025.1.0.0, CUDA 12.8 | +| Build profile | Ordinary complete `rl_engine._C`; 10 source translation units | +| CUDA architecture / parallelism | `TORCH_CUDA_ARCH_LIST=12.0`, `MAX_JOBS=2` | + +CUDA compiler, development libraries and command-line tools were installed from +NVIDIA's WSL-Ubuntu repository; the existing Windows GPU driver provides CUDA +driver access. The WSL driver directory `/usr/lib/wsl/lib` was added to +`LIBRARY_PATH` for linking the package's existing CUDA driver API calls. +`RL_KERNEL_ROPE_ONLY`, forced SM90 and fast-math flags were unset. + +The normal PEP 517 wheel build compiled all ten configured translation units, +including the existing logprob/GEMM/norm/activation/attention/collective code and +the new RoPE source. The complete `_C` loaded and registered the native RoPE +dispatcher. `full-extension.json` records sources, flags, binary path and hash; +`package.json` records the built wheel and installed library fingerprints. + +The checkout, virtual environments, build outputs and caches are inside the +Linux filesystem, backed by the D: virtual disk. A fresh Git clone preserved all +27 tracked symlinks. This is a WSL2 Linux userspace/compiler qualification on +SM120; bare-metal Linux and A100/H100 remain separate hardware profiles. + +## Results + +| Check | Result | Evidence | +| --- | --- | --- | +| Ordinary complete `_C` wheel and in-place build/import | Passed | `full-extension.json`, `full-wheel-build.txt`, `package.json` | +| Operator accuracy/invariance/autograd/graphs/opcheck/fullgraph | 81 passed, 1 two-GPU guard skip | `tests-and-ws1.txt`, `operator-tests.xml` | +| CUDA and Triton WS1 four judgments | Both passed with explicit provenance | `ws1-cuda.json`, `ws1-triton.json` | +| Standard operator CLI forward/input gradient | Both passed | `operator-cuda.json`, `operator-triton.json` | +| Related regressions | 195 passed, 7 Ascend skips | `regression.txt` | +| Repository layout with extension enforcement enabled | 37 passed, zero skips/failures | `layout.txt` | +| Fresh installed full wheel, outside source checkout | Eager/compiled forward and input-gradient bytes exact for both backends | `installed-wheel.json` | +| Actual CPU-only PyTorch wheel/reference installation | 57 passed, 52 expected GPU skips; reference forward/backward passed | `cpu-contracts.txt`, `cpu-environment.json` | +| Compute Sanitizer accuracy/stride/empty/stream subset | 34 passed; zero errors | `memcheck.txt` | +| Pinned actual Diffusers symbols | 30 forward/backward cases passed; image/text tables exact | `upstream.json` | +| Reference-grid benchmark correctness | 40 comparisons passed | `benchmark-bfloat16.json` | +| Shell/workflow syntax, changed-file checks, strict docs | Passed locally | `quality.txt` | + +The full-extension GPU gate rejects missing native/Triton execution. Its only +allowed operator-test skip is the two-device guard on a one-GPU runner. +The CPU-only result is separate: it uses `torch==2.9.1+cpu`, and its installed +package contains no native CUDA binary. CPU skips are not GPU acceptance evidence. +The 37 passing Linux layout tests contrast with the five Windows baseline +failures recorded earlier; those historical logs were not rewritten. + +The installed-wheel probe starts in a fresh process outside the checkout, +loads `_C` from site-packages and reports that exact path. It tests packed-QKV +forward/input-gradient equality in eager and `torch.compile(fullgraph=True)` +for both CUDA and Triton. The FlashAttention BSD license is present in the wheel. + +## Standalone benchmark + +BF16, batch 1, 24 heads, head dimension 128. CUDA Graph medians below are +**forward / forward+backward milliseconds per call**: + +| Grid / layout | Native CUDA | Adapted Triton | FlashAttention, no FMA | Diffusers complex | +| --- | --- | --- | --- | --- | +| 64x64 contiguous | 0.05780 / 0.12967 | 0.04929 / 0.12414 | 0.04898 / 0.12407 | 0.24446 / 0.51281 | +| 64x64 packed QKV | 0.05868 / 0.13079 | 0.04731 / 0.11841 | 0.04722 / 0.11847 | 0.24754 / 0.51583 | +| 83x83 contiguous | 0.11198 / 0.22392 | 0.10720 / 0.22538 | 0.10739 / 0.21420 | 0.52579 / 1.06721 | +| 83x83 packed QKV | 0.11233 / 0.22486 | 0.10745 / 0.21478 | 0.10744 / 0.21469 | 0.52820 / 1.06947 | +| 104x58 contiguous | 0.09797 / 0.19708 | 0.09366 / 0.18497 | 0.09362 / 0.18458 | 0.44992 / 0.90110 | +| 104x58 packed QKV | 0.09610 / 0.19612 | 0.09310 / 0.18828 | 0.09323 / 0.18817 | 0.44979 / 0.90276 | + +These are 20 calls per graph, five warm calls/replays and 30 replay measurements. +Three synchronized eager-wall trials of 30 calls are also in the JSON. Tables, +compilation and input/backing allocations are outside timing. Extra peak includes +the returned output or gradient. The short synthetic case is included in raw data. +At 4096 tokens, both GPU paths' forward extra peak is 24 MiB for either layout, +the returned BF16 output size. Packed inputs are read without materialization. + +Original FlashAttention and its autograd wrapper are checksum-pinned to +`98c1cee0ed59307833841cdbf567fe336ceec13a`. The no-FMA baseline changes only the +launch contraction option. CUDA, adapted Triton and that baseline have zero +differing forward/input-gradient bytes against the unfused FP32 gold in every +benchmark case. Default FlashAttention and complex arithmetic use declared tolerances. +The adapted Triton path is close to the mature baseline; no performance advantage +over FlashAttention is established. Timings are one-process measurements on a +shared Windows/WSL GPU without clock pinning, not a model-throughput result or +a controlled comparison against the earlier Windows process. + +## Reproduction and CI + +Use the matching CUDA 12.8 toolkit and Linux PyTorch 2.9.1+cu128 in a fresh +Linux virtual environment. Other Python package versions are recorded in +`python-packages.json`. From the Linux checkout: + +```bash +export CUDA_HOME=/usr/local/cuda-12.8 +export PATH="$CUDA_HOME/bin:$PATH" +export LIBRARY_PATH="/usr/lib/wsl/lib${LIBRARY_PATH:+:$LIBRARY_PATH}" +export MAX_JOBS=2 TORCH_CUDA_ARCH_LIST=12.0 RL_KERNEL_REQUIRE_EXT=1 +unset RL_KERNEL_ROPE_ONLY KERNEL_ALIGN_FORCE_SM90 KERNEL_ALIGN_USE_FAST_MATH +python -m pip wheel --no-build-isolation --no-deps . -w /tmp/rope-full-wheel +python -m pip install --force-reinstall --no-deps /tmp/rope-full-wheel/rl_kernel-*.whl +python setup.py build_ext --inplace +ROPE_PYTHON=python ROPE_WS1_OUTPUT=/tmp/rope-ws1 bash ci/scripts/run_qwenimage_rope_ws1.sh +python ci/run.py layout -q -rs +compute-sanitizer --tool memcheck --error-exitcode 1 python -m pytest \ + tests/ops/rope/test_multi_axis_rope.py -q \ + -k "forward_backward_accuracy or noncontiguous_input_tables_and_gradient or empty_forward_backward or packed_qkv_and_independent_table_strides or native_entrypoint_checks" +``` + +The new [operator gate](../../../ci/scripts/run_qwenimage_rope_ws1.sh) requires +the complete `_C` registration, records actual library/backend identity and +enforces the GPU skip policy. The [CPU workflow](../../../.github/workflows/qwenimage-rope.yml) +targets `test-qwenimage`, runs Python 3.10/3.12 CPU contracts, checks changed files +and builds documentation. This report's actual CPU-only run used Python 3.12; +hosted CI status is tracked on the PR. GPU runners use the explicit gate. + +Bare-metal Linux, A100/H100, two-GPU guards, QK-norm fusion, model/block assembly +and trajectories remain unqualified. Building the entire library does not +certify every native operator. No Dense Qwen3/distributed certification or +cross-device/version byte-equality claim is made. diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/benchmark-bfloat16.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/benchmark-bfloat16.json new file mode 100644 index 000000000..ca80105b0 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/benchmark-bfloat16.json @@ -0,0 +1,2781 @@ +{ + "gpu": "NVIDIA GeForce RTX 5080", + "capability": [ + 12, + 0 + ], + "os": "Linux-6.18.54.1-microsoft-standard-WSL2-x86_64-with-glibc2.39", + "python": "3.12.3 (main, Aug 31 2026, 10:18:26) [GCC 13.3.0]", + "torch": "2.9.1+cu128", + "cuda_runtime": "12.8", + "triton": "3.5.1", + "dtype": "bfloat16", + "base_commit": "ac17ab87e8d93dabe93dfdc134e9853d47de7af1", + "source_sha256": { + "csrc/cuda/rope/multi_axis_rope.cu": "f5db5cf5d4962053194455d9cde8cf91cf4e824bad7abf72e8410a404bcb0e44", + "rl_engine/backends/shared/triton/rope/_flash_attn_rotary.py": "cf0eebf315830a4178053a71a77e650b58aa85bba1d1a08aff7677350af755d9", + "rl_engine/backends/shared/triton/rope/multi_axis_rope.py": "f08ff90627c66e15e9c558f8a68f132434a741a76997b3800bfbf73dc8940e9e", + "rl_engine/backends/cuda/rope/multi_axis_rope.py": "d779d9b60ce5d1d6651925f723049930d898024fdf0548ce29a2019b84753d65", + "rl_engine/ops/rope/_registration.py": "add0375d1639444908091d01e8d29781f7a975beffec64b77bfaa81a5b9a0ae3", + "rl_engine/ops/rope/multi_axis_rope.py": "07bcff1160f22c1ff6730be731206b2f4a153e2fada26afc93fc884e780568cb", + "rl_engine/reference/rope/multi_axis_rope.py": "c9a94730047f51659f7aa0a443e88d6ce7d03d00d6813a9aed548cafbf6347bf", + "benchmarks/operators/rope/benchmark_multi_axis_rope.py": "6961a55803fbd4e566bd060558fb59e658e7f1f704e61ca97f912e5d16b55884" + }, + "timing": "CUDA graph ms/call and synchronized eager wall ms/call; warm allocation; excludes tables, compilation, and initial allocation of inputs", + "memory": "extra peak allocated bytes excluding inputs/tables/grad; includes returned output or gradient; packed QKV backing storage allocated before timing", + "iterations": 30, + "warmup": 5, + "trials": 3, + "graph_calls": 20, + "flash_baseline": { + "included": true, + "revision": "98c1cee0ed59307833841cdbf567fe336ceec13a", + "kernel_sha256_normalized": "d2037e623854c702fb075784fe4292d4a89d6b9ac1ae33cab3d6788f8f94f240", + "layer_sha256_normalized": "a79cb42dfc33e126227bdba27ae7a024d7d5cc5c8fd7f7a5f26c68d501024e0a", + "mode": "upstream rotary_kernel, apply_rotary and ApplyRotaryEmb, interleaved=True", + "no_fma_change": "only add enable_fp_fusion=False at the original kernel launch" + }, + "trace": { + "cuda": { + "contract": "qwenimage-interleaved-rope-fp32-v1", + "actual_backend": "rlkernel.cuda.qwenimage_multi_axis_rope.v1", + "kernel_id": "torch.ops.rl_kernel_native.multi_axis_rope", + "layout": "BSHD, adjacent channel pairs", + "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", + "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" + }, + "triton": { + "contract": "qwenimage-interleaved-rope-fp32-v1", + "actual_backend": "rlkernel.triton.qwenimage_multi_axis_rope.v1", + "kernel_id": "rl_engine.backends.shared.triton.rope._flash_attn_rotary.rotary_kernel", + "layout": "BSHD, adjacent channel pairs", + "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", + "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" + } + }, + "qualification": "Exact bytes vs unfused FP32 gold for CUDA, Triton and flash_no_fma; declared tolerance for upstream complex and default FlashAttention", + "correctness": [ + { + "backend": "upstream_complex", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "cuda", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "triton", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "flash_default", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "flash_no_fma", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "upstream_complex", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "cuda", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "triton", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "flash_default", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "flash_no_fma", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "upstream_complex", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "different_bytes_y_dx": [ + 130, + 139 + ] + }, + { + "backend": "cuda", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "triton", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "flash_default", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "different_bytes_y_dx": [ + 130, + 139 + ] + }, + { + "backend": "flash_no_fma", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "upstream_complex", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "different_bytes_y_dx": [ + 141, + 139 + ] + }, + { + "backend": "cuda", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "triton", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "flash_default", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "different_bytes_y_dx": [ + 141, + 139 + ] + }, + { + "backend": "flash_no_fma", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "upstream_complex", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "different_bytes_y_dx": [ + 235, + 224 + ] + }, + { + "backend": "cuda", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "triton", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "flash_default", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "different_bytes_y_dx": [ + 235, + 224 + ] + }, + { + "backend": "flash_no_fma", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "upstream_complex", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "different_bytes_y_dx": [ + 221, + 226 + ] + }, + { + "backend": "cuda", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "triton", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "flash_default", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "different_bytes_y_dx": [ + 221, + 226 + ] + }, + { + "backend": "flash_no_fma", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "upstream_complex", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "different_bytes_y_dx": [ + 195, + 199 + ] + }, + { + "backend": "cuda", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "triton", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "flash_default", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "different_bytes_y_dx": [ + 195, + 199 + ] + }, + { + "backend": "flash_no_fma", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "upstream_complex", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "different_bytes_y_dx": [ + 230, + 220 + ] + }, + { + "backend": "cuda", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "triton", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + }, + { + "backend": "flash_default", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "different_bytes_y_dx": [ + 230, + 220 + ] + }, + { + "backend": "flash_no_fma", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "different_bytes_y_dx": [ + 0, + 0 + ] + } + ], + "results": [ + { + "backend": "upstream_complex", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 18432, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.0034856000915169717, + "graph_min_ms": 0.0034543998539447785, + "graph_max_ms": 0.0035312000662088396, + "eager_wall_median_ms": 0.08385686669498682, + "eager_wall_trials_ms": [ + 0.09771363335554877, + 0.08163450002030004, + 0.08385686669498682 + ], + "extra_peak_bytes": 184320 + }, + { + "backend": "upstream_complex", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 18432, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.007651200145483017, + "graph_min_ms": 0.007601600140333176, + "graph_max_ms": 0.007692799717187881, + "eager_wall_median_ms": 0.19571573332844613, + "eager_wall_trials_ms": [ + 0.2535254000273805, + 0.19571573332844613, + 0.18590186670432254 + ], + "extra_peak_bytes": 187392 + }, + { + "backend": "cuda", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 18432, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.0012047999538481236, + "graph_min_ms": 0.0011695999652147292, + "graph_max_ms": 0.0012400000356137752, + "eager_wall_median_ms": 0.05762376664885475, + "eager_wall_trials_ms": [ + 0.07130193328824437, + 0.05762376664885475, + 0.057149166665719044 + ], + "extra_peak_bytes": 36864 + }, + { + "backend": "cuda", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 18432, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.002329600043594837, + "graph_min_ms": 0.002223999984562397, + "graph_max_ms": 0.0026383999735116958, + "eager_wall_median_ms": 0.17063139997238372, + "eager_wall_trials_ms": [ + 0.17063139997238372, + 0.16431806670880178, + 0.24000986668397672 + ], + "extra_peak_bytes": 73728 + }, + { + "backend": "triton", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 18432, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.0013432000298053027, + "graph_min_ms": 0.0010928000323474408, + "graph_max_ms": 0.006166400015354156, + "eager_wall_median_ms": 0.08490003328915918, + "eager_wall_trials_ms": [ + 0.08490003328915918, + 0.0852208666401566, + 0.08217023332690587 + ], + "extra_peak_bytes": 36864 + }, + { + "backend": "triton", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 18432, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.002125599980354309, + "graph_min_ms": 0.0020735999569296837, + "graph_max_ms": 0.0021679999306797983, + "eager_wall_median_ms": 0.3603733666144156, + "eager_wall_trials_ms": [ + 0.3862667332820517, + 0.31255590001819655, + 0.3603733666144156 + ], + "extra_peak_bytes": 73728 + }, + { + "backend": "flash_default", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 18432, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.0011024000123143196, + "graph_min_ms": 0.0010544000193476676, + "graph_max_ms": 0.01971679925918579, + "eager_wall_median_ms": 0.07505266664035541, + "eager_wall_trials_ms": [ + 0.09395336667997374, + 0.07483469998987857, + 0.07505266664035541 + ], + "extra_peak_bytes": 36864 + }, + { + "backend": "flash_default", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 18432, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.0020584000274538996, + "graph_min_ms": 0.002022399939596653, + "graph_max_ms": 0.002151999995112419, + "eager_wall_median_ms": 0.3825640333161573, + "eager_wall_trials_ms": [ + 0.3916477333405055, + 0.38042896667320747, + 0.3825640333161573 + ], + "extra_peak_bytes": 73728 + }, + { + "backend": "flash_no_fma", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 18432, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.0011288000270724297, + "graph_min_ms": 0.0010703999549150466, + "graph_max_ms": 0.0025040000677108766, + "eager_wall_median_ms": 0.07896646663236122, + "eager_wall_trials_ms": [ + 0.07896646663236122, + 0.08368196668016026, + 0.07650009999148703 + ], + "extra_peak_bytes": 36864 + }, + { + "backend": "flash_no_fma", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 18432, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.002057600021362305, + "graph_min_ms": 0.002022399939596653, + "graph_max_ms": 0.0021487999707460403, + "eager_wall_median_ms": 0.3408336999806731, + "eager_wall_trials_ms": [ + 0.3408336999806731, + 0.32808933331883355, + 0.35737479999321903 + ], + "extra_peak_bytes": 73728 + }, + { + "backend": "upstream_complex", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 55296, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.0037903999909758568, + "graph_min_ms": 0.003734400123357773, + "graph_max_ms": 0.0038384001702070234, + "eager_wall_median_ms": 0.026099166704322368, + "eager_wall_trials_ms": [ + 0.027518266688275617, + 0.02535739998468974, + 0.026099166704322368 + ], + "extra_peak_bytes": 184320 + }, + { + "backend": "upstream_complex", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 55296, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.008167200163006781, + "graph_min_ms": 0.00809599980711937, + "graph_max_ms": 0.008207999914884568, + "eager_wall_median_ms": 0.22357730000900725, + "eager_wall_trials_ms": [ + 0.255850233346185, + 0.22357730000900725, + 0.2018136666568656 + ], + "extra_peak_bytes": 187392 + }, + { + "backend": "cuda", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 55296, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.001232000021263957, + "graph_min_ms": 0.0011552000418305398, + "graph_max_ms": 0.0026464000344276427, + "eager_wall_median_ms": 0.05505463335187718, + "eager_wall_trials_ms": [ + 0.05505463335187718, + 0.052852566659566946, + 0.05531913332864254 + ], + "extra_peak_bytes": 36864 + }, + { + "backend": "cuda", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 55296, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.002279200032353401, + "graph_min_ms": 0.00220320001244545, + "graph_max_ms": 0.005976000055670738, + "eager_wall_median_ms": 0.2586098333267728, + "eager_wall_trials_ms": [ + 0.2587837333218583, + 0.25146820001585485, + 0.2586098333267728 + ], + "extra_peak_bytes": 73728 + }, + { + "backend": "triton", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 55296, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.0011359999887645244, + "graph_min_ms": 0.001071999967098236, + "graph_max_ms": 0.002603200078010559, + "eager_wall_median_ms": 0.03804060003555302, + "eager_wall_trials_ms": [ + 0.03804060003555302, + 0.04575726664673615, + 0.03741419999035619 + ], + "extra_peak_bytes": 36864 + }, + { + "backend": "triton", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 55296, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.0021271999925374986, + "graph_min_ms": 0.0020896000787615776, + "graph_max_ms": 0.016212800145149232, + "eager_wall_median_ms": 0.24443720000514682, + "eager_wall_trials_ms": [ + 0.281011633342132, + 0.24443720000514682, + 0.23313996668245332 + ], + "extra_peak_bytes": 73728 + }, + { + "backend": "flash_default", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 55296, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.001104000024497509, + "graph_min_ms": 0.001065600011497736, + "graph_max_ms": 0.0011983999982476234, + "eager_wall_median_ms": 0.078476833308135, + "eager_wall_trials_ms": [ + 0.07911106664929927, + 0.07519836666081876, + 0.078476833308135 + ], + "extra_peak_bytes": 36864 + }, + { + "backend": "flash_default", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 55296, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.0020527999848127367, + "graph_min_ms": 0.0020207999274134635, + "graph_max_ms": 0.0021023999899625776, + "eager_wall_median_ms": 0.30830893338134047, + "eager_wall_trials_ms": [ + 0.3673079666744646, + 0.29416493331761256, + 0.30830893338134047 + ], + "extra_peak_bytes": 73728 + }, + { + "backend": "flash_no_fma", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 55296, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.0011272000148892403, + "graph_min_ms": 0.0010511999949812888, + "graph_max_ms": 0.002239999920129776, + "eager_wall_median_ms": 0.07592803334167304, + "eager_wall_trials_ms": [ + 0.0775905666766145, + 0.07592803334167304, + 0.0742380666755101 + ], + "extra_peak_bytes": 36864 + }, + { + "backend": "flash_no_fma", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6, + 24, + 128 + ], + "input_stride": [ + 55296, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.0020535999909043314, + "graph_min_ms": 0.002014400064945221, + "graph_max_ms": 0.0020975999534130095, + "eager_wall_median_ms": 0.31871896662778454, + "eager_wall_trials_ms": [ + 0.31871896662778454, + 0.33167873334605247, + 0.2915063333299865 + ], + "extra_peak_bytes": 73728 + }, + { + "backend": "upstream_complex", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 12582912, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.24445840120315554, + "graph_min_ms": 0.24224960803985596, + "graph_max_ms": 0.29114880561828616, + "eager_wall_median_ms": 0.247559633377629, + "eager_wall_trials_ms": [ + 0.25793806665509084, + 0.2471523333345734, + 0.247559633377629 + ], + "extra_peak_bytes": 125829120 + }, + { + "backend": "upstream_complex", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 12582912, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.5128072023391723, + "graph_min_ms": 0.49546399116516116, + "graph_max_ms": 0.5571055889129639, + "eager_wall_median_ms": 0.5245368333514003, + "eager_wall_trials_ms": [ + 0.5245368333514003, + 0.518849266639639, + 0.5350402333230401 + ], + "extra_peak_bytes": 127926272 + }, + { + "backend": "cuda", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 12582912, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.05780480206012726, + "graph_min_ms": 0.057633602619171144, + "graph_max_ms": 0.07439360022544861, + "eager_wall_median_ms": 0.06806933330760027, + "eager_wall_trials_ms": [ + 0.07017140002668991, + 0.06806933330760027, + 0.06168783335548748 + ], + "extra_peak_bytes": 25165824 + }, + { + "backend": "cuda", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 12582912, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.1296720027923584, + "graph_min_ms": 0.12902239561080933, + "graph_max_ms": 0.15688639879226685, + "eager_wall_median_ms": 0.2333034999537631, + "eager_wall_trials_ms": [ + 0.23402639999403618, + 0.19631726663646987, + 0.2333034999537631 + ], + "extra_peak_bytes": 50331648 + }, + { + "backend": "triton", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 12582912, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.049288800358772276, + "graph_min_ms": 0.04815520048141479, + "graph_max_ms": 0.06759200096130372, + "eager_wall_median_ms": 0.08061226665934858, + "eager_wall_trials_ms": [ + 0.07100329997532147, + 0.0930295666573026, + 0.08061226665934858 + ], + "extra_peak_bytes": 25165824 + }, + { + "backend": "triton", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 12582912, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.12414479851722718, + "graph_min_ms": 0.12152320146560669, + "graph_max_ms": 0.16048799753189086, + "eager_wall_median_ms": 0.35214599999259616, + "eager_wall_trials_ms": [ + 0.35569400000288925, + 0.35214599999259616, + 0.32247826663175755 + ], + "extra_peak_bytes": 50331648 + }, + { + "backend": "flash_default", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 12582912, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.04924079924821853, + "graph_min_ms": 0.04797120094299316, + "graph_max_ms": 0.068995201587677, + "eager_wall_median_ms": 0.13288023334704727, + "eager_wall_trials_ms": [ + 0.14359453337723002, + 0.12000013339275029, + 0.13288023334704727 + ], + "extra_peak_bytes": 25165824 + }, + { + "backend": "flash_default", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 12582912, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.12360240221023559, + "graph_min_ms": 0.1215999960899353, + "graph_max_ms": 0.16038399934768677, + "eager_wall_median_ms": 0.35429976663484314, + "eager_wall_trials_ms": [ + 0.34954116666388774, + 0.3772011999899405, + 0.35429976663484314 + ], + "extra_peak_bytes": 50331648 + }, + { + "backend": "flash_no_fma", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 12582912, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.0489767998456955, + "graph_min_ms": 0.0479856014251709, + "graph_max_ms": 0.06645600199699402, + "eager_wall_median_ms": 0.1330274333061728, + "eager_wall_trials_ms": [ + 0.1330274333061728, + 0.1200082000044252, + 0.13366413331823424 + ], + "extra_peak_bytes": 25165824 + }, + { + "backend": "flash_no_fma", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 12582912, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.12407039999961852, + "graph_min_ms": 0.12193280458450317, + "graph_max_ms": 0.1605712056159973, + "eager_wall_median_ms": 0.36529683329717955, + "eager_wall_trials_ms": [ + 0.3895335000076254, + 0.36529683329717955, + 0.3503133667133322 + ], + "extra_peak_bytes": 50331648 + }, + { + "backend": "upstream_complex", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 37748736, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.2475391983985901, + "graph_min_ms": 0.24529919624328614, + "graph_max_ms": 0.27714240550994873, + "eager_wall_median_ms": 0.25804310001452296, + "eager_wall_trials_ms": [ + 0.25804310001452296, + 0.26510979999632883, + 0.2503190333300154 + ], + "extra_peak_bytes": 125829120 + }, + { + "backend": "upstream_complex", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 37748736, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.5158272027969361, + "graph_min_ms": 0.49741439819335936, + "graph_max_ms": 0.5595920085906982, + "eager_wall_median_ms": 0.5221092666639985, + "eager_wall_trials_ms": [ + 0.5362037999778598, + 0.5221092666639985, + 0.5111432000073061 + ], + "extra_peak_bytes": 127926272 + }, + { + "backend": "cuda", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 37748736, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.05868320167064667, + "graph_min_ms": 0.058520001173019406, + "graph_max_ms": 0.0754975974559784, + "eager_wall_median_ms": 0.06486223331497361, + "eager_wall_trials_ms": [ + 0.06433113333817649, + 0.08489736665069358, + 0.06486223331497361 + ], + "extra_peak_bytes": 25165824 + }, + { + "backend": "cuda", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 37748736, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.13079119920730592, + "graph_min_ms": 0.13011519908905028, + "graph_max_ms": 0.1711359977722168, + "eager_wall_median_ms": 0.21280260001124893, + "eager_wall_trials_ms": [ + 0.2440493333779159, + 0.20012976668416135, + 0.21280260001124893 + ], + "extra_peak_bytes": 50331648 + }, + { + "backend": "triton", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 37748736, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.04731200039386749, + "graph_min_ms": 0.04663040041923523, + "graph_max_ms": 0.06482719779014587, + "eager_wall_median_ms": 0.08378989999376547, + "eager_wall_trials_ms": [ + 0.08378989999376547, + 0.07920856666411662, + 0.09377416663483018 + ], + "extra_peak_bytes": 25165824 + }, + { + "backend": "triton", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 37748736, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.11841280460357666, + "graph_min_ms": 0.11643359661102295, + "graph_max_ms": 0.14446239471435546, + "eager_wall_median_ms": 0.3498463667104564, + "eager_wall_trials_ms": [ + 0.3498463667104564, + 0.35651856663510745, + 0.3451830999741408 + ], + "extra_peak_bytes": 50331648 + }, + { + "backend": "flash_default", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 37748736, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.047358399629592894, + "graph_min_ms": 0.04637599885463715, + "graph_max_ms": 0.06467199921607972, + "eager_wall_median_ms": 0.12952533331069085, + "eager_wall_trials_ms": [ + 0.12952533331069085, + 0.1346659000167468, + 0.12932516668418734 + ], + "extra_peak_bytes": 25165824 + }, + { + "backend": "flash_default", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 37748736, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.11804000139236451, + "graph_min_ms": 0.11632959842681885, + "graph_max_ms": 0.13619199991226197, + "eager_wall_median_ms": 0.3809861666619933, + "eager_wall_trials_ms": [ + 0.3809861666619933, + 0.39690620002753957, + 0.3660936666468236 + ], + "extra_peak_bytes": 50331648 + }, + { + "backend": "flash_no_fma", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 37748736, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.04721840023994446, + "graph_min_ms": 0.04616160094738007, + "graph_max_ms": 0.06441760063171387, + "eager_wall_median_ms": 0.1290158666355031, + "eager_wall_trials_ms": [ + 0.13406613331123177, + 0.1290158666355031, + 0.12222770001244498 + ], + "extra_peak_bytes": 25165824 + }, + { + "backend": "flash_no_fma", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 4096, + 24, + 128 + ], + "input_stride": [ + 37748736, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.11847440004348755, + "graph_min_ms": 0.11651359796524048, + "graph_max_ms": 0.1468176007270813, + "eager_wall_median_ms": 0.4179154666417162, + "eager_wall_trials_ms": [ + 0.4434371333142432, + 0.4179154666417162, + 0.4101795332947707 + ], + "extra_peak_bytes": 50331648 + }, + { + "backend": "upstream_complex", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 21163008, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.5257872104644776, + "graph_min_ms": 0.5080063819885254, + "graph_max_ms": 0.5525023937225342, + "eager_wall_median_ms": 0.5246295667044857, + "eager_wall_trials_ms": [ + 0.5097412666752158, + 0.528316466625256, + 0.5246295667044857 + ], + "extra_peak_bytes": 211630080 + }, + { + "backend": "upstream_complex", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 21163008, + 3072, + 128, + 1 + ], + "graph_median_ms": 1.067212772369385, + "graph_min_ms": 1.0389776229858398, + "graph_max_ms": 1.0976528167724608, + "eager_wall_median_ms": 1.0610819333426964, + "eager_wall_trials_ms": [ + 1.0702722999970622, + 1.0479776333644015, + 1.0610819333426964 + ], + "extra_peak_bytes": 215157248 + }, + { + "backend": "cuda", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 21163008, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.11197760105133056, + "graph_min_ms": 0.11102559566497802, + "graph_max_ms": 0.13988800048828126, + "eager_wall_median_ms": 0.12250546666715915, + "eager_wall_trials_ms": [ + 0.11638536664880424, + 0.13100786666958206, + 0.12250546666715915 + ], + "extra_peak_bytes": 42326016 + }, + { + "backend": "cuda", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 21163008, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.22392239570617678, + "graph_min_ms": 0.22292959690093994, + "graph_max_ms": 0.26030399799346926, + "eager_wall_median_ms": 0.24186149997452353, + "eager_wall_trials_ms": [ + 0.24518829995940905, + 0.23021603331775017, + 0.24186149997452353 + ], + "extra_peak_bytes": 84652032 + }, + { + "backend": "triton", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 21163008, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.10719599723815917, + "graph_min_ms": 0.10527039766311645, + "graph_max_ms": 0.123745596408844, + "eager_wall_median_ms": 0.10772686667526916, + "eager_wall_trials_ms": [ + 0.10711883332987782, + 0.10772686667526916, + 0.10832119999880281 + ], + "extra_peak_bytes": 42326016 + }, + { + "backend": "triton", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 21163008, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.2253783941268921, + "graph_min_ms": 0.21142399311065674, + "graph_max_ms": 0.25633440017700193, + "eager_wall_median_ms": 0.29183916667534504, + "eager_wall_trials_ms": [ + 0.29183916667534504, + 0.31016709999676095, + 0.27477109997562366 + ], + "extra_peak_bytes": 84652032 + }, + { + "backend": "flash_default", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 21163008, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.10732319951057434, + "graph_min_ms": 0.10543199777603149, + "graph_max_ms": 0.12232480049133301, + "eager_wall_median_ms": 0.11885433329249888, + "eager_wall_trials_ms": [ + 0.11160320003303544, + 0.11885433329249888, + 0.148421633275575 + ], + "extra_peak_bytes": 42326016 + }, + { + "backend": "flash_default", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 21163008, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.22528079748153684, + "graph_min_ms": 0.21176319122314452, + "graph_max_ms": 0.24576001167297362, + "eager_wall_median_ms": 0.41978080001475365, + "eager_wall_trials_ms": [ + 0.4051902666712219, + 0.42735956667456776, + 0.41978080001475365 + ], + "extra_peak_bytes": 84652032 + }, + { + "backend": "flash_no_fma", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 21163008, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.10738959908485413, + "graph_min_ms": 0.10557440519332886, + "graph_max_ms": 0.1310464024543762, + "eager_wall_median_ms": 0.1100779333379857, + "eager_wall_trials_ms": [ + 0.11928496666466042, + 0.10885563333431492, + 0.1100779333379857 + ], + "extra_peak_bytes": 42326016 + }, + { + "backend": "flash_no_fma", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 21163008, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.21419600248336793, + "graph_min_ms": 0.2123471975326538, + "graph_max_ms": 0.24199841022491456, + "eager_wall_median_ms": 0.42007340004298993, + "eager_wall_trials_ms": [ + 0.4409272999813159, + 0.42007340004298993, + 0.40205200002674246 + ], + "extra_peak_bytes": 84652032 + }, + { + "backend": "upstream_complex", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 63489024, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.5282040119171143, + "graph_min_ms": 0.5110288143157959, + "graph_max_ms": 0.5885712146759033, + "eager_wall_median_ms": 0.547415499992591, + "eager_wall_trials_ms": [ + 0.547415499992591, + 0.5160773666527044, + 0.5478402666388623 + ], + "extra_peak_bytes": 211630080 + }, + { + "backend": "upstream_complex", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 63489024, + 9216, + 128, + 1 + ], + "graph_median_ms": 1.0694655895233156, + "graph_min_ms": 1.0273263931274415, + "graph_max_ms": 1.0971407890319824, + "eager_wall_median_ms": 1.0825770000034634, + "eager_wall_trials_ms": [ + 1.0599153666286536, + 1.0825770000034634, + 1.0843000333503976 + ], + "extra_peak_bytes": 215157248 + }, + { + "backend": "cuda", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 63489024, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.11232720017433166, + "graph_min_ms": 0.11165920495986939, + "graph_max_ms": 0.13819040060043336, + "eager_wall_median_ms": 0.12610396667393312, + "eager_wall_trials_ms": [ + 0.12610396667393312, + 0.11783170005704353, + 0.1361641000282058 + ], + "extra_peak_bytes": 42326016 + }, + { + "backend": "cuda", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 63489024, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.22485519647598268, + "graph_min_ms": 0.2238991975784302, + "graph_max_ms": 0.2728208065032959, + "eager_wall_median_ms": 0.2468356666213367, + "eager_wall_trials_ms": [ + 0.2468356666213367, + 0.25081796666199807, + 0.23804076669572774 + ], + "extra_peak_bytes": 84652032 + }, + { + "backend": "triton", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 63489024, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.10744640231132507, + "graph_min_ms": 0.10558079481124878, + "graph_max_ms": 0.12357280254364014, + "eager_wall_median_ms": 0.10917900005248764, + "eager_wall_trials_ms": [ + 0.108484466666899, + 0.11029286670236615, + 0.10917900005248764 + ], + "extra_peak_bytes": 42326016 + }, + { + "backend": "triton", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 63489024, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.21478239297866822, + "graph_min_ms": 0.21232800483703612, + "graph_max_ms": 0.2458656072616577, + "eager_wall_median_ms": 0.3064794333719571, + "eager_wall_trials_ms": [ + 0.3064794333719571, + 0.29039823330094805, + 0.3166810000038822 + ], + "extra_peak_bytes": 84652032 + }, + { + "backend": "flash_default", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 63489024, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.10760959982872009, + "graph_min_ms": 0.10598560571670532, + "graph_max_ms": 0.13379520177841187, + "eager_wall_median_ms": 0.11580663334219328, + "eager_wall_trials_ms": [ + 0.11580663334219328, + 0.12037279999882837, + 0.10920460002428929 + ], + "extra_peak_bytes": 42326016 + }, + { + "backend": "flash_default", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 63489024, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.2147576093673706, + "graph_min_ms": 0.21242880821228027, + "graph_max_ms": 0.2415616035461426, + "eager_wall_median_ms": 0.42939333337320323, + "eager_wall_trials_ms": [ + 0.4258981332895928, + 0.45155456670424127, + 0.42939333337320323 + ], + "extra_peak_bytes": 84652032 + }, + { + "backend": "flash_no_fma", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 63489024, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.1074351966381073, + "graph_min_ms": 0.10590399503707885, + "graph_max_ms": 0.13916159868240358, + "eager_wall_median_ms": 0.11420323335187277, + "eager_wall_trials_ms": [ + 0.11669516667704254, + 0.11420323335187277, + 0.11166136664542137 + ], + "extra_peak_bytes": 42326016 + }, + { + "backend": "flash_no_fma", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6889, + 24, + 128 + ], + "input_stride": [ + 63489024, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.21468960046768187, + "graph_min_ms": 0.21285600662231446, + "graph_max_ms": 0.24125280380249023, + "eager_wall_median_ms": 0.4036813999846345, + "eager_wall_trials_ms": [ + 0.40494703334843507, + 0.4027272333284297, + 0.4036813999846345 + ], + "extra_peak_bytes": 84652032 + }, + { + "backend": "upstream_complex", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 18530304, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.44991838932037354, + "graph_min_ms": 0.4312863826751709, + "graph_max_ms": 0.4923136234283447, + "eager_wall_median_ms": 0.4565672999888193, + "eager_wall_trials_ms": [ + 0.46419443333434174, + 0.4565672999888193, + 0.43665636664324364 + ], + "extra_peak_bytes": 185303040 + }, + { + "backend": "upstream_complex", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 18530304, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.9010960102081299, + "graph_min_ms": 0.8724479675292969, + "graph_max_ms": 0.9627167701721191, + "eager_wall_median_ms": 0.9057802333093908, + "eager_wall_trials_ms": [ + 0.9057802333093908, + 0.9399917666087276, + 0.8861856666650662 + ], + "extra_peak_bytes": 188391424 + }, + { + "backend": "cuda", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 18530304, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.09797199964523315, + "graph_min_ms": 0.0971664011478424, + "graph_max_ms": 0.12116639614105225, + "eager_wall_median_ms": 0.10149109997049284, + "eager_wall_trials_ms": [ + 0.10149109997049284, + 0.10145610000715048, + 0.10150603332779913 + ], + "extra_peak_bytes": 37060608 + }, + { + "backend": "cuda", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 18530304, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.19708160161972046, + "graph_min_ms": 0.19586080312728882, + "graph_max_ms": 0.2340928077697754, + "eager_wall_median_ms": 0.23985833334639514, + "eager_wall_trials_ms": [ + 0.2599421666673152, + 0.21093993333731, + 0.23985833334639514 + ], + "extra_peak_bytes": 74809344 + }, + { + "backend": "triton", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 18530304, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.09365599751472473, + "graph_min_ms": 0.09192960262298584, + "graph_max_ms": 0.1098512053489685, + "eager_wall_median_ms": 0.09574563337082509, + "eager_wall_trials_ms": [ + 0.10015153335795428, + 0.09574563337082509, + 0.09447829997952795 + ], + "extra_peak_bytes": 37060608 + }, + { + "backend": "triton", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 18530304, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.18496879935264587, + "graph_min_ms": 0.18224799633026123, + "graph_max_ms": 0.20439839363098145, + "eager_wall_median_ms": 0.27221376664480584, + "eager_wall_trials_ms": [ + 0.30083316663270426, + 0.2585961666530541, + 0.27221376664480584 + ], + "extra_peak_bytes": 74809344 + }, + { + "backend": "flash_default", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 18530304, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.09344720244407653, + "graph_min_ms": 0.09223840236663819, + "graph_max_ms": 0.10959680080413818, + "eager_wall_median_ms": 0.10214689997762132, + "eager_wall_trials_ms": [ + 0.10214689997762132, + 0.09693080003974804, + 0.10462310001457809 + ], + "extra_peak_bytes": 37060608 + }, + { + "backend": "flash_default", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 18530304, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.18451839685440063, + "graph_min_ms": 0.1817072033882141, + "graph_max_ms": 0.2279695987701416, + "eager_wall_median_ms": 0.38619293330460397, + "eager_wall_trials_ms": [ + 0.4084413000479496, + 0.38619293330460397, + 0.3539249666573596 + ], + "extra_peak_bytes": 74809344 + }, + { + "backend": "flash_no_fma", + "phase": "forward", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 18530304, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.09362159967422486, + "graph_min_ms": 0.09215999841690063, + "graph_max_ms": 0.1108240008354187, + "eager_wall_median_ms": 0.10443516669814319, + "eager_wall_trials_ms": [ + 0.10080796667656007, + 0.10632786670612404, + 0.10443516669814319 + ], + "extra_peak_bytes": 37060608 + }, + { + "backend": "flash_no_fma", + "phase": "forward_backward", + "layout": "contiguous", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 18530304, + 3072, + 128, + 1 + ], + "graph_median_ms": 0.18458240032196044, + "graph_min_ms": 0.18273439407348632, + "graph_max_ms": 0.21199359893798828, + "eager_wall_median_ms": 0.3521715333287527, + "eager_wall_trials_ms": [ + 0.3573932999643148, + 0.32707516663019004, + 0.3521715333287527 + ], + "extra_peak_bytes": 74809344 + }, + { + "backend": "upstream_complex", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 55590912, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.44979281425476075, + "graph_min_ms": 0.4314896106719971, + "graph_max_ms": 0.48181757926940916, + "eager_wall_median_ms": 0.4543782333106113, + "eager_wall_trials_ms": [ + 0.4543782333106113, + 0.4393956000058097, + 0.4658290666460137 + ], + "extra_peak_bytes": 185303040 + }, + { + "backend": "upstream_complex", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 55590912, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.9027584075927735, + "graph_min_ms": 0.8707056045532227, + "graph_max_ms": 0.9478416442871094, + "eager_wall_median_ms": 0.9098065333091654, + "eager_wall_trials_ms": [ + 0.92959719998665, + 0.9058386999943954, + 0.9098065333091654 + ], + "extra_peak_bytes": 188391424 + }, + { + "backend": "cuda", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 55590912, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.0960999995470047, + "graph_min_ms": 0.09566079974174499, + "graph_max_ms": 0.1124384045600891, + "eager_wall_median_ms": 0.10085823329670045, + "eager_wall_trials_ms": [ + 0.10085823329670045, + 0.1027275666880693, + 0.10056819998377857 + ], + "extra_peak_bytes": 37060608 + }, + { + "backend": "cuda", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 55590912, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.19612080454826356, + "graph_min_ms": 0.19543039798736572, + "graph_max_ms": 0.22353920936584473, + "eager_wall_median_ms": 0.21260113335301867, + "eager_wall_trials_ms": [ + 0.21260113335301867, + 0.23863556671130937, + 0.20749836670195995 + ], + "extra_peak_bytes": 74809344 + }, + { + "backend": "triton", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 55590912, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.09310480058193207, + "graph_min_ms": 0.09154719710350037, + "graph_max_ms": 0.1167631983757019, + "eager_wall_median_ms": 0.10304526670855314, + "eager_wall_trials_ms": [ + 0.09351696662633913, + 0.10913006669094709, + 0.10304526670855314 + ], + "extra_peak_bytes": 37060608 + }, + { + "backend": "triton", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 55590912, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.18827759623527526, + "graph_min_ms": 0.18518240451812745, + "graph_max_ms": 0.21609120368957518, + "eager_wall_median_ms": 0.33319789999950444, + "eager_wall_trials_ms": [ + 0.33319789999950444, + 0.3456709999833644, + 0.3028138333320385 + ], + "extra_peak_bytes": 74809344 + }, + { + "backend": "flash_default", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 55590912, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.09317199885845184, + "graph_min_ms": 0.09170079827308655, + "graph_max_ms": 0.11182240247726441, + "eager_wall_median_ms": 0.13384176666780453, + "eager_wall_trials_ms": [ + 0.12049876671274735, + 0.15613870000379393, + 0.13384176666780453 + ], + "extra_peak_bytes": 37060608 + }, + { + "backend": "flash_default", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 55590912, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.18772000074386597, + "graph_min_ms": 0.18648799657821655, + "graph_max_ms": 0.2264847993850708, + "eager_wall_median_ms": 0.3818082333358082, + "eager_wall_trials_ms": [ + 0.44640976666414645, + 0.3674994000296768, + 0.3818082333358082 + ], + "extra_peak_bytes": 74809344 + }, + { + "backend": "flash_no_fma", + "phase": "forward", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 55590912, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.09322640001773834, + "graph_min_ms": 0.09153280258178711, + "graph_max_ms": 0.11082240343093872, + "eager_wall_median_ms": 0.13263150000663396, + "eager_wall_trials_ms": [ + 0.13263150000663396, + 0.13320910002221353, + 0.12726086667195582 + ], + "extra_peak_bytes": 37060608 + }, + { + "backend": "flash_no_fma", + "phase": "forward_backward", + "layout": "packed_qkv", + "shape": [ + 1, + 6032, + 24, + 128 + ], + "input_stride": [ + 55590912, + 9216, + 128, + 1 + ], + "graph_median_ms": 0.18817280530929564, + "graph_min_ms": 0.18564800024032593, + "graph_max_ms": 0.21453280448913575, + "eager_wall_median_ms": 0.3810788333794335, + "eager_wall_trials_ms": [ + 0.3668230666638313, + 0.3885829333739821, + 0.3810788333794335 + ], + "extra_peak_bytes": 74809344 + } + ] +} diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu-contracts.txt b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu-contracts.txt new file mode 100644 index 000000000..de9819e68 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu-contracts.txt @@ -0,0 +1,20 @@ +...........ssssssssssss.ss.ss.ss...ssssss.ss......ssssssssssss.ss...ssss [ 66%] +s.s.ssssss........................... [100%] +=========================== short test summary info ============================ +SKIPPED [12] tests/ops/rope/test_multi_axis_rope.py:79: NVIDIA CUDA required +SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:100: NVIDIA CUDA required +SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:110: NVIDIA CUDA required +SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:133: NVIDIA CUDA required +SKIPPED [6] tests/ops/rope/test_multi_axis_rope.py:148: NVIDIA CUDA required +SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:158: NVIDIA CUDA required +SKIPPED [12] tests/ops/rope/test_multi_axis_rope.py:174: NVIDIA CUDA required +SKIPPED [1] tests/ops/rope/test_multi_axis_rope.py:217: native CUDA multi_axis_rope extension required +SKIPPED [1] tests/ops/rope/test_multi_axis_rope.py:230: NVIDIA CUDA required +SKIPPED [3] tests/ops/rope/test_multi_axis_rope.py:255: NVIDIA CUDA required +SKIPPED [1] tests/ops/rope/test_multi_axis_rope.py:315: NVIDIA CUDA required +SKIPPED [1] tests/ops/rope/test_multi_axis_rope.py:336: requires two NVIDIA GPUs +SKIPPED [1] tests/ops/rope/test_multi_axis_rope.py:355: CUDA context requires CUDA +SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:368: NVIDIA CUDA required +SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:399: NVIDIA CUDA required +SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:427: NVIDIA CUDA required +57 passed, 52 skipped in 1.70s diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu-environment.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu-environment.json new file mode 100644 index 000000000..e41fd057f --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu-environment.json @@ -0,0 +1,7 @@ +{ + "torch": "2.9.1+cpu", + "cuda": null, + "package": "/home/dandan/venvs/rlk-rope-cpu/lib/python3.12/site-packages/rl_engine/__init__.py", + "reference_forward_backward": true, + "passed": true +} diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/full-extension.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/full-extension.json new file mode 100644 index 000000000..420aa5da8 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/full-extension.json @@ -0,0 +1,22 @@ +{ + "profile": "ordinary complete _C; no RL_KERNEL_ROPE_ONLY", + "module": "/home/dandan/src/RL-Kernel/rl_engine/_C.cpython-312-x86_64-linux-gnu.so", + "binary_sha256": "2661e7188984c7d78acae1ef5186a3e90ec4aa8bb6f457522b23435a1f24f94a", + "sources": [ + "csrc/bindings/ops.cpp", + "csrc/cuda/logprob/fused_logp_kernel.cu", + "csrc/cuda/logprob/deterministic_logp_kernel.cu", + "csrc/cuda/gemm/det_gemm_kernel.cu", + "csrc/cuda/norm/rmsnorm.cu", + "csrc/cuda/activation/activation.cu", + "csrc/cuda/attention/deterministic_attention.cu", + "csrc/cuda/rope/multi_axis_rope.cu", + "csrc/cuda/collectives/deterministic_collective.cu", + "csrc/cuda/attention/prefix_shared_attention.cu" + ], + "cuda_home": "/usr/local/cuda-12.8", + "driver_link_search": "/usr/lib/wsl/lib", + "torch": "2.9.1+cu128", + "architecture": "12.0", + "max_jobs": "2" +} diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/full-inplace-build.txt b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/full-inplace-build.txt new file mode 100644 index 000000000..d61fb7517 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/full-inplace-build.txt @@ -0,0 +1,20 @@ +/home/dandan/src/RL-Kernel/.venv/lib/python3.12/site-packages/setuptools/config/_apply_pyprojecttoml.py:82: SetuptoolsDeprecationWarning: `project.license` as a TOML table is deprecated +!! + + ******************************************************************************** + Please use a simple string containing a SPDX expression for `project.license`. You can also use `project.license-files`. (Both options available on setuptools>=77.0.0). + + This deprecation is overdue, please update your project and remove deprecated + calls to avoid build errors in the future. + + See https://packaging.python.org/en/latest/guides/writing-pyproject-toml/#license for details. + ******************************************************************************** + +!! + corresp(dist, value, root_dir) +running build_ext +W1028 18:37:59.459000 9791 .venv/lib/python3.12/site-packages/torch/utils/cpp_extension.py:531] There are no x86_64-linux-gnu-g++ version bounds defined for CUDA version 12.8 +building 'rl_engine._C' extension +ninja: no work to do. +x86_64-linux-gnu-g++ -fno-strict-overflow -Wsign-compare -DNDEBUG -g -O2 -Wall -shared -Wl,-O1 -Wl,-Bsymbolic-functions /home/dandan/src/RL-Kernel/build/temp.linux-x86_64-cpython-312/csrc/bindings/ops.o /home/dandan/src/RL-Kernel/build/temp.linux-x86_64-cpython-312/csrc/cuda/activation/activation.o /home/dandan/src/RL-Kernel/build/temp.linux-x86_64-cpython-312/csrc/cuda/attention/deterministic_attention.o /home/dandan/src/RL-Kernel/build/temp.linux-x86_64-cpython-312/csrc/cuda/attention/prefix_shared_attention.o /home/dandan/src/RL-Kernel/build/temp.linux-x86_64-cpython-312/csrc/cuda/collectives/deterministic_collective.o /home/dandan/src/RL-Kernel/build/temp.linux-x86_64-cpython-312/csrc/cuda/gemm/det_gemm_kernel.o /home/dandan/src/RL-Kernel/build/temp.linux-x86_64-cpython-312/csrc/cuda/logprob/deterministic_logp_kernel.o /home/dandan/src/RL-Kernel/build/temp.linux-x86_64-cpython-312/csrc/cuda/logprob/fused_logp_kernel.o /home/dandan/src/RL-Kernel/build/temp.linux-x86_64-cpython-312/csrc/cuda/norm/rmsnorm.o /home/dandan/src/RL-Kernel/build/temp.linux-x86_64-cpython-312/csrc/cuda/rope/multi_axis_rope.o -L/home/dandan/src/RL-Kernel/.venv/lib/python3.12/site-packages/torch/lib -L/usr/local/cuda-12.8/lib64 -L/usr/lib/x86_64-linux-gnu -lc10 -ltorch -ltorch_cpu -ltorch_python -lcudart -lc10_cuda -ltorch_cuda -o build/lib.linux-x86_64-cpython-312/rl_engine/_C.cpython-312-x86_64-linux-gnu.so -Wl,-rpath,$ORIGIN/../torch/lib -lcuda +copying build/lib.linux-x86_64-cpython-312/rl_engine/_C.cpython-312-x86_64-linux-gnu.so -> rl_engine diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/full-wheel-build.txt b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/full-wheel-build.txt new file mode 100644 index 000000000..592c254f2 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/full-wheel-build.txt @@ -0,0 +1,10 @@ +Processing /home/dandan/src/RL-Kernel + Preparing metadata (pyproject.toml): started + Preparing metadata (pyproject.toml): finished with status 'done' +Building wheels for collected packages: RL-Kernel + Building wheel for RL-Kernel (pyproject.toml): started + Building wheel for RL-Kernel (pyproject.toml): still running... + Building wheel for RL-Kernel (pyproject.toml): finished with status 'done' + Created wheel for RL-Kernel: filename=rl_kernel-0.1.0-cp312-cp312-linux_x86_64.whl size=5020592 sha256=32a39cd9cbb30b87816f6584096238fc2e4b0803eb387a2d3e1e80f3d1b40d8a + Stored in directory: /tmp/pip-ephem-wheel-cache-8gtts1_e/wheels/de/2c/1f/c224e9a9cc43d1b39ee96ba5703938f5f1d66da44d8695a43a +Successfully built RL-Kernel diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/gpu-environment.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/gpu-environment.json new file mode 100644 index 000000000..e40096583 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/gpu-environment.json @@ -0,0 +1,15 @@ +{ + "scope": "Qwen-Image multi_axis_rope operator; ordinary full _C build", + "python": "3.12.3", + "platform": "Linux-6.18.54.1-microsoft-standard-WSL2-x86_64-with-glibc2.39", + "torch": "2.9.1+cu128", + "cuda": "12.8", + "triton": "3.5.1", + "gpu": "NVIDIA GeForce RTX 5080", + "capability": [ + 12, + 0 + ], + "native_library": "/home/dandan/src/RL-Kernel/rl_engine/_C.cpython-312-x86_64-linux-gnu.so", + "native_library_sha256": "2661e7188984c7d78acae1ef5186a3e90ec4aa8bb6f457522b23435a1f24f94a" +} diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/installed-wheel.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/installed-wheel.json new file mode 100644 index 000000000..28d3066fc --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/installed-wheel.json @@ -0,0 +1,38 @@ +{ + "package": "/home/dandan/src/RL-Kernel/.venv/lib/python3.12/site-packages/rl_engine", + "companion_library": null, + "full_extension_library": "/home/dandan/src/RL-Kernel/.venv/lib/python3.12/site-packages/rl_engine/_C.cpython-312-x86_64-linux-gnu.so", + "native_dispatcher_registered": true, + "torch": "2.9.1+cu128", + "gpu": "NVIDIA GeForce RTX 5080", + "license_present": true, + "packed_qkv_stride": [ + 17280, + 1152, + 128, + 1 + ], + "checks": { + "cuda_eager": "passed", + "cuda_compiled": "passed", + "cuda_trace": { + "contract": "qwenimage-interleaved-rope-fp32-v1", + "actual_backend": "rlkernel.cuda.qwenimage_multi_axis_rope.v1", + "kernel_id": "torch.ops.rl_kernel_native.multi_axis_rope", + "layout": "BSHD, adjacent channel pairs", + "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", + "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" + }, + "triton_eager": "passed", + "triton_compiled": "passed", + "triton_trace": { + "contract": "qwenimage-interleaved-rope-fp32-v1", + "actual_backend": "rlkernel.triton.qwenimage_multi_axis_rope.v1", + "kernel_id": "rl_engine.backends.shared.triton.rope._flash_attn_rotary.rotary_kernel", + "layout": "BSHD, adjacent channel pairs", + "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", + "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" + } + }, + "passed": true +} diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/layout.txt b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/layout.txt new file mode 100644 index 000000000..8c81ecbf5 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/layout.txt @@ -0,0 +1,2 @@ +..................................... [100%] +37 passed in 3.81s diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/manifest.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/manifest.json new file mode 100644 index 000000000..3acb62a79 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/manifest.json @@ -0,0 +1,58 @@ +{ + "profile": "Ubuntu 24.04.5 / WSL2 / RTX 5080 SM120 / ordinary complete _C", + "numerical_source_commit": "ac17ab87e8d93dabe93dfdc134e9853d47de7af1", + "source_sha256": { + ".github/workflows/qwenimage-rope.yml": "39e7e02a1f1f75603482e6acce9eb99ba8b5c5260ec6043271714061843651c3", + "MANIFEST.in": "87e59e6c40c67e54ae9bcd057957ac7d7823067cb455cbc4dff45b1ec88bdb3c", + "benchmarks/operators/rope/benchmark_multi_axis_rope.py": "6961a55803fbd4e566bd060558fb59e658e7f1f704e61ca97f912e5d16b55884", + "build_tools/extensions.py": "3723f14a3fb9307889790af15b5e3f360c4c46e12b8fef22423070e9660f70c5", + "ci/scripts/run_qwenimage_rope_ws1.sh": "8e1e12c1aad9231f215505a172b3a5416c54e7cbfb8856f5ff00d225be7bfa95", + "csrc/bindings/ops.cpp": "10b76a5e236c0dc3715f9703e28c057f404f21f2a9802245d5221c4997b2c3f4", + "csrc/cuda/activation/activation.cu": "35ef0cefcaa6167aec7a901fd19a2f6c90a8b9f01c2e786187aacdeee10c4b63", + "csrc/cuda/attention/deterministic_attention.cu": "99b44addd61b999a5fb6a04ca227ed8814d454368f361a8fd032e83d4d83a561", + "csrc/cuda/attention/prefix_shared_attention.cu": "218ab9647cd274f7a4d8246f8df8e875c5cf431a5a73f252e946b10c74a2d63f", + "csrc/cuda/collectives/deterministic_collective.cu": "9e114d5aae7a73a912e4e4afaa23a74126bef5557ba3b54832da0c0cdf9d97a2", + "csrc/cuda/gemm/det_gemm_kernel.cu": "08d08b0a32900fd9590d47f04819b34a2d37cbf53996aa5ed1373d88eae29f33", + "csrc/cuda/logprob/deterministic_logp_kernel.cu": "271399040a0489a3cabd9e0c06b64fa8b1bac92ab08710fce8922bcb1ff2b04f", + "csrc/cuda/logprob/fused_logp_kernel.cu": "ea67bcd307f0c63753397225543d333f58cf3d66a0ae69b643499d1e6729d8e7", + "csrc/cuda/norm/rmsnorm.cu": "fbb765563dbe9c063b35f38354321eaf522cf62ba03addf327154395db9683b9", + "csrc/cuda/rope/multi_axis_rope.cu": "f5db5cf5d4962053194455d9cde8cf91cf4e824bad7abf72e8410a404bcb0e44", + "docs/operators/multi-axis-rope.md": "1fb3859a9a2639aae809ece52a785737e9966098c24a08626b9240dad0d6cbf6", + "pyproject.toml": "c2d948784b379a0a828bbf3d8cfd9e0e7398a1ee5694fe8298509514684b2709", + "rl_engine/backends/cuda/rope/multi_axis_rope.py": "d779d9b60ce5d1d6651925f723049930d898024fdf0548ce29a2019b84753d65", + "rl_engine/backends/shared/triton/rope/_flash_attn_rotary.py": "cf0eebf315830a4178053a71a77e650b58aa85bba1d1a08aff7677350af755d9", + "rl_engine/backends/shared/triton/rope/multi_axis_rope.py": "f08ff90627c66e15e9c558f8a68f132434a741a76997b3800bfbf73dc8940e9e", + "rl_engine/ops/rope/_registration.py": "add0375d1639444908091d01e8d29781f7a975beffec64b77bfaa81a5b9a0ae3", + "rl_engine/ops/rope/multi_axis_rope.py": "07bcff1160f22c1ff6730be731206b2f4a153e2fada26afc93fc884e780568cb", + "rl_engine/reference/rope/multi_axis_rope.py": "c9a94730047f51659f7aa0a443e88d6ce7d03d00d6813a9aed548cafbf6347bf", + "rl_engine/validation/operators/qwen_image_rope.py": "c3ddfc732f32de3c633de16064810d67202254ea78fd8107d558b471bb15a37f", + "tests/ops/rope/test_multi_axis_rope.py": "c9ff3eb6e6e013243fa741f853debb256fdac6c1d3a3820eb7de764a830e240c", + "tools/validation/operators/check_multi_axis_rope_install.py": "f295fc8ac5878ac13b6255c2f0be83b66d34ff2ad2f75111bb8b633508a865a7" + }, + "evidence_sha256": { + "benchmark-bfloat16.json": "e9c1e4b1085623d84800656c179e4a0e86d70a0ef7fc9ed92f9e6ccc50f0f491", + "cpu-contracts.txt": "9518f53a6efc603bca1c52f160e5430f1006bb2abc9a7ef57694dae6feac5d78", + "cpu-environment.json": "4dfc68f3d76af05a050d830840b4a6b8c8fa0a4c12c29d11d632a56a57a2d4c6", + "full-extension.json": "8fae782e3c5e6aae59de6fe6b4e7ae185729e49d2282dbd0cb3b9ca08c0274cd", + "full-inplace-build.txt": "04d5ef122650ab305847d58850a6c3b44438eacf97920492e1eb1c5580e9a21e", + "full-wheel-build.txt": "a7883dfb1211798c869ede1ee0b6f1dfb32edd7b27b359ee18c6e5ab9fe449d5", + "gpu-environment.json": "bf808ab535d90b9b52dfcfb86543809e7d412cb48418677923cad22783bb9b8e", + "installed-wheel.json": "c726717106f30336e0a541e4a38d1f1b3858ed5dd262516c86618d4511a43554", + "layout.txt": "daef0733010cdad31123cb6b8cb041fe6c626319d527a5d97c88d1b958479039", + "memcheck.txt": "1a0c572f8f0d18e3699c847dbb69b345e4fd19200cabda3a39385d588cff0f10", + "operator-cuda.json": "fd95922bac2c5921744a8e601e9a267e51b7b7d7d44d94745f53964c2da2feb1", + "operator-tests.xml": "0ed98ddf85e34ba87f41465be1db4f417997e82b50cc5f8f5a83bd00f1f1543e", + "operator-triton.json": "7791aa1dfa68652dd66e3c332d8f036e9e21851e1df7fd41a090d594e22b06ff", + "package.json": "3f75c06fdee58afb1cdd9a2a611b432c8be1960fa640d96263c86b87423fe7bd", + "precommit.txt": "7cf90e06cf73194271c176ebc2d147e64cf8f5ed522d8e0e03093a4647f63759", + "python-packages.json": "617f8cef5a24f3b2a683e70e0b491e47db6d24a7ca848e5b3e461e87993a7201", + "quality.txt": "8826f0e406946a923400d792dd60d0f545b32f9fcffa3a148268c969e30e14c0", + "README.md": "af365e658fe4875df643704b2495e1f115c6c889f5d449880233d49e511040ac", + "regression.txt": "039cd9813a2ea54bfc5b7750fc0555b54d2bf30b09dbca2eba4a0403578ab9b0", + "tests-and-ws1.txt": "68397e27144c9244743dbcff39ffdca8a77a30eaf1e8180b2ca84dfc1216d633", + "upstream.json": "04e5472a82477a305b996ccbb420d0f3cc7b39af4053c867d37916a9bf5de51a", + "ws1-cuda.json": "7302e11f798816178622393095858871e2dc53ab925509eb130196a40b3c18ad", + "ws1-triton.json": "b21f985fbdff7e9b14f689d5eac25c9df5d6d3b0189900085837d5fc400f9d38" + }, + "hash_mode": "SHA256 of recorded Linux file bytes; text is LF" +} diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/memcheck.txt b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/memcheck.txt new file mode 100644 index 000000000..915746141 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/memcheck.txt @@ -0,0 +1,4 @@ +========= COMPUTE-SANITIZER +.................................. [100%] +34 passed, 48 deselected in 3.49s +========= ERROR SUMMARY: 0 errors diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/operator-cuda.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/operator-cuda.json new file mode 100644 index 000000000..801312e4a --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/operator-cuda.json @@ -0,0 +1,70 @@ +{ + "suite_name": "multi_axis_rope", + "total_candidates": 1, + "passed_candidates": 1, + "pass_rate": 1.0, + "passed": true, + "candidates": [ + { + "candidate_name": "cuda-multi_axis_rope", + "backend": "cuda", + "total_outputs": 2, + "passed_outputs": 2, + "pass_rate": 1.0, + "passed": true, + "cases": [ + { + "case_name": "multi_axis_rope-torch.bfloat16-2x7x3x128", + "dtype": "torch.bfloat16", + "op_class": "elementwise", + "passed": true, + "outputs": [ + { + "output_index": 0, + "shape": [ + 2, + 7, + 3, + 128 + ], + "candidate_dtype": "torch.bfloat16", + "gold_dtype": "torch.float32", + "atol": 0.02, + "rtol": 0.016, + "max_abs_error": 0.007704973220825195, + "mean_abs_error": 0.0006018930580466986, + "max_rel_error": 0.003828622167930007, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference", + "message": "" + }, + { + "output_index": 1, + "shape": [ + 2, + 7, + 3, + 128 + ], + "candidate_dtype": "torch.bfloat16", + "gold_dtype": "torch.bfloat16", + "atol": 0.02, + "rtol": 0.016, + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference", + "message": "gradient:x" + } + ] + } + ], + "backend_provenance": null + } + ] +} diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/operator-tests.xml b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/operator-tests.xml new file mode 100644 index 000000000..9e26942b9 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/operator-tests.xml @@ -0,0 +1 @@ +/home/dandan/src/RL-Kernel/tests/ops/rope/test_multi_axis_rope.py:336: requires two NVIDIA GPUs diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/operator-triton.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/operator-triton.json new file mode 100644 index 000000000..e105fb723 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/operator-triton.json @@ -0,0 +1,70 @@ +{ + "suite_name": "multi_axis_rope", + "total_candidates": 1, + "passed_candidates": 1, + "pass_rate": 1.0, + "passed": true, + "candidates": [ + { + "candidate_name": "triton-multi_axis_rope", + "backend": "triton", + "total_outputs": 2, + "passed_outputs": 2, + "pass_rate": 1.0, + "passed": true, + "cases": [ + { + "case_name": "multi_axis_rope-torch.bfloat16-2x7x3x128", + "dtype": "torch.bfloat16", + "op_class": "elementwise", + "passed": true, + "outputs": [ + { + "output_index": 0, + "shape": [ + 2, + 7, + 3, + 128 + ], + "candidate_dtype": "torch.bfloat16", + "gold_dtype": "torch.float32", + "atol": 0.02, + "rtol": 0.016, + "max_abs_error": 0.007704973220825195, + "mean_abs_error": 0.0006018930580466986, + "max_rel_error": 0.003828622167930007, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference", + "message": "" + }, + { + "output_index": 1, + "shape": [ + 2, + 7, + 3, + 128 + ], + "candidate_dtype": "torch.bfloat16", + "gold_dtype": "torch.bfloat16", + "atol": 0.02, + "rtol": 0.016, + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference", + "message": "gradient:x" + } + ] + } + ], + "backend_provenance": null + } + ] +} diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/package.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/package.json new file mode 100644 index 000000000..1dce0c941 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/package.json @@ -0,0 +1,11 @@ +{ + "profile": "ordinary complete _C wheel; RL_KERNEL_ROPE_ONLY unset", + "wheel": "rl_kernel-0.1.0-cp312-cp312-linux_x86_64.whl", + "wheel_sha256": "32a39cd9cbb30b87816f6584096238fc2e4b0803eb387a2d3e1e80f3d1b40d8a", + "native_member": "rl_engine/_C.cpython-312-x86_64-linux-gnu.so", + "native_sha256": "2661e7188984c7d78acae1ef5186a3e90ec4aa8bb6f457522b23435a1f24f94a", + "license_present": true, + "installed_library": "/home/dandan/src/RL-Kernel/.venv/lib/python3.12/site-packages/rl_engine/_C.cpython-312-x86_64-linux-gnu.so", + "installed_library_sha256": "2661e7188984c7d78acae1ef5186a3e90ec4aa8bb6f457522b23435a1f24f94a", + "cpu_wheel_contains_no_cuda_binary": true +} diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/precommit.txt b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/precommit.txt new file mode 100644 index 000000000..475a8c0d8 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/precommit.txt @@ -0,0 +1,7 @@ +trim trailing whitespace.................................................Passed +fix end of files.........................................................Passed +check yaml...............................................................Passed +check for added large files..............................................Passed +black....................................................................Passed +isort....................................................................Passed +flake8...................................................................Passed diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/python-packages.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/python-packages.json new file mode 100644 index 000000000..41997f91d --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/python-packages.json @@ -0,0 +1,114 @@ +{ + "GitPython": "3.2.1", + "Jinja2": "3.1.6", + "Markdown": "3.11.2", + "MarkupSafe": "3.0.3", + "PyYAML": "6.0.3", + "Pygments": "2.21.0", + "RL-Kernel": "0.1.0", + "accelerate": "1.15.0", + "annotated-doc": "0.0.5", + "annotated-types": "0.8.0", + "anyio": "4.15.1", + "ast_serialize": "0.13.1", + "babel": "2.18.0", + "backrefs": "8.1", + "black": "24.4.2", + "bracex": "3.2", + "certifi": "2026.10.23", + "cfgv": "3.5.0", + "charset-normalizer": "3.5.2", + "click": "8.5.0", + "colorama": "0.4.6", + "distlib": "0.4.4", + "filelock": "3.32.3", + "flake8": "7.0.0", + "fsspec": "2026.7.0", + "ghp-import": "2.1.0", + "gitdb": "4.0.12", + "h11": "0.16.0", + "hf-xet": "1.7.0", + "httpcore": "1.0.9", + "httpx": "0.28.1", + "huggingface_hub": "1.33.0", + "identify": "2.6.20", + "idna": "3.20", + "iniconfig": "2.3.1", + "isort": "5.13.2", + "librt": "0.16.0", + "markdown-it-py": "4.2.0", + "mccabe": "0.7.0", + "mdurl": "0.1.2", + "mergedeep": "1.3.4", + "mkdocs": "1.6.1", + "mkdocs-awesome-nav": "3.3.0", + "mkdocs-get-deps": "0.2.2", + "mkdocs-git-revision-date-localized-plugin": "1.6.2", + "mkdocs-material": "9.8.0", + "mkdocs-material-extensions": "1.3.1", + "mpmath": "1.3.0", + "mypy": "2.4.0", + "mypy_extensions": "1.1.0", + "natsort": "8.4.0", + "networkx": "3.6.1", + "ninja": "1.13.2", + "nodeenv": "1.11.0", + "numpy": "2.5.4", + "nvidia-cublas-cu12": "12.8.4.1", + "nvidia-cuda-cupti-cu12": "12.8.90", + "nvidia-cuda-nvrtc-cu12": "12.8.93", + "nvidia-cuda-runtime-cu12": "12.8.90", + "nvidia-cudnn-cu12": "9.10.2.21", + "nvidia-cufft-cu12": "11.3.3.83", + "nvidia-cufile-cu12": "1.13.1.3", + "nvidia-curand-cu12": "10.3.9.90", + "nvidia-cusolver-cu12": "11.7.3.90", + "nvidia-cusparse-cu12": "12.5.8.93", + "nvidia-cusparselt-cu12": "0.7.1", + "nvidia-nccl-cu12": "2.27.5", + "nvidia-nvjitlink-cu12": "12.8.93", + "nvidia-nvshmem-cu12": "3.3.20", + "nvidia-nvtx-cu12": "12.8.90", + "packaging": "26.3", + "paginate": "0.5.7", + "pathspec": "1.1.1", + "pip": "25.3", + "platformdirs": "4.13.7", + "pluggy": "1.7.0", + "pre_commit": "4.7.0", + "psutil": "7.2.2", + "pycodestyle": "2.11.1", + "pydantic": "2.14.1", + "pydantic_core": "2.50.1", + "pyflakes": "3.2.0", + "pymdown-extensions": "12.3", + "pytest": "9.2.0", + "python-dateutil": "3.0.0", + "python-discovery": "1.6.7", + "pyyaml_env_tag": "1.1", + "regex": "2026.9.29", + "requests": "2.34.2", + "rich": "15.0.0", + "ruff": "0.17.3", + "safetensors": "0.8.0", + "setuptools": "80.9.0", + "shellingham": "1.5.4", + "six": "1.17.0", + "smmap": "5.0.3", + "sympy": "1.14.0", + "tabulate": "0.10.0", + "tokenizers": "0.22.2", + "torch": "2.9.1+cu128", + "tqdm": "4.70.1", + "transformers": "5.13.1", + "triton": "3.5.1", + "typer": "0.27.3", + "typing-inspection": "0.4.4", + "typing_extensions": "4.16.0", + "tzdata": "2026.6", + "urllib3": "2.8.0", + "virtualenv": "21.15.1", + "watchdog": "6.0.0", + "wcmatch": "11.2", + "wheel": "0.45.1" +} diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/quality.txt b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/quality.txt new file mode 100644 index 000000000..8e84b2473 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/quality.txt @@ -0,0 +1,34 @@ +Local verification; no CUDA numerical/kernel sources changed for this follow-up. + +workflow-validation.log +Workflow YAML, branch coverage and embedded Python syntax validated + +isort.log + + +black.log +All done! ✨ 🍰 ✨ +19 files would be left unchanged. + +flake8.log + + +ruff.log +All checks passed! + +docs.log +INFO - Cleaning site directory +INFO - Building documentation to directory: /home/dandan/src/RL-Kernel/site +INFO - The following pages exist in the docs directory, but are not included in the "nav" configuration: + - blog/2026-07-08-announcing-rl-kernel-linear-logp-for-vime-zh.md + - blog/2026-08-09-dsv4-flash-moe-consistency-roadmap.md + - operators/matmul.md + - operators/rope.md + - usage/cross-config-drift-report.md + - usage/cuda-rocm-consistency-audit.md + - usage/h100-configurable.md + - usage/h100-cp-gradient-validation.md + - usage/h100-matrix-validation.md + - usage/h100-pcp-validation.md + - usage/rocm-sparse-performance.md +INFO - Documentation built in 1.15 seconds diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/regression.txt b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/regression.txt new file mode 100644 index 000000000..95d27f812 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/regression.txt @@ -0,0 +1,10 @@ +........................................................................ [ 35%] +........................................................................ [ 71%] +...................................................sssssss [100%] +=========================== short test summary info ============================ +SKIPPED [3] tests/ops/rope/test_rope.py:390: rope_apply_ascend is not compiled (requires an Ascend NPU build). +SKIPPED [1] tests/ops/rope/test_rope.py:406: rope_apply_ascend is not compiled (requires an Ascend NPU build). +SKIPPED [1] tests/ops/rope/test_rope.py:413: rope_apply_ascend is not compiled (requires an Ascend NPU build). +SKIPPED [1] tests/ops/rope/test_rope.py:420: rope_apply_ascend is not compiled (requires an Ascend NPU build). +SKIPPED [1] tests/ops/rope/test_rope.py:433: rope_apply_ascend is not compiled (requires an Ascend NPU build). +195 passed, 7 skipped in 4.56s diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/tests-and-ws1.txt b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/tests-and-ws1.txt new file mode 100644 index 000000000..0b8ce46dd --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/tests-and-ws1.txt @@ -0,0 +1,11 @@ +{"scope": "Qwen-Image multi_axis_rope operator; ordinary full _C build", "python": "3.12.3", "platform": "Linux-6.18.54.1-microsoft-standard-WSL2-x86_64-with-glibc2.39", "torch": "2.9.1+cu128", "cuda": "12.8", "triton": "3.5.1", "gpu": "NVIDIA GeForce RTX 5080", "capability": [12, 0], "native_library": "/home/dandan/src/RL-Kernel/rl_engine/_C.cpython-312-x86_64-linux-gnu.so", "native_library_sha256": "2661e7188984c7d78acae1ef5186a3e90ec4aa8bb6f457522b23435a1f24f94a"} +........................................................................ [ 87%] +s......... [100%] +=========================== short test summary info ============================ +SKIPPED [1] tests/ops/rope/test_multi_axis_rope.py:336: requires two NVIDIA GPUs +81 passed, 1 skipped in 17.49s +Using the installed production RoPE dispatcher library +WS1 four judgments backend=cuda: passed=True +Using the installed production RoPE dispatcher library +WS1 four judgments backend=triton: passed=True +Qwen-Image RoPE full-extension operator gate passed diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/upstream.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/upstream.json new file mode 100644 index 000000000..032e23098 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/upstream.json @@ -0,0 +1,335 @@ +{ + "source_sha256": "0bad11901deae69f05aec01fd464e51e38582d455068f2851ee665036179321f", + "cases": [ + { + "grid": [ + 2, + 3 + ], + "stream": "image", + "dtype": "torch.float32", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 2, + 3 + ], + "stream": "image", + "dtype": "torch.float16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 2, + 3 + ], + "stream": "image", + "dtype": "torch.bfloat16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 2, + 3 + ], + "stream": "text", + "dtype": "torch.float32", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 2, + 3 + ], + "stream": "text", + "dtype": "torch.float16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 2, + 3 + ], + "stream": "text", + "dtype": "torch.bfloat16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 3, + 5 + ], + "stream": "image", + "dtype": "torch.float32", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 3, + 5 + ], + "stream": "image", + "dtype": "torch.float16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 3, + 5 + ], + "stream": "image", + "dtype": "torch.bfloat16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 3, + 5 + ], + "stream": "text", + "dtype": "torch.float32", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 3, + 5 + ], + "stream": "text", + "dtype": "torch.float16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 3, + 5 + ], + "stream": "text", + "dtype": "torch.bfloat16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 64, + 64 + ], + "stream": "image", + "dtype": "torch.float32", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 64, + 64 + ], + "stream": "image", + "dtype": "torch.float16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 64, + 64 + ], + "stream": "image", + "dtype": "torch.bfloat16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 64, + 64 + ], + "stream": "text", + "dtype": "torch.float32", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 64, + 64 + ], + "stream": "text", + "dtype": "torch.float16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 64, + 64 + ], + "stream": "text", + "dtype": "torch.bfloat16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 83, + 83 + ], + "stream": "image", + "dtype": "torch.float32", + "table_bitwise": true, + "forward_max_abs": 2.384185791015625e-07, + "dx_max_abs": 2.384185791015625e-07 + }, + { + "grid": [ + 83, + 83 + ], + "stream": "image", + "dtype": "torch.float16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 83, + 83 + ], + "stream": "image", + "dtype": "torch.bfloat16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 83, + 83 + ], + "stream": "text", + "dtype": "torch.float32", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 83, + 83 + ], + "stream": "text", + "dtype": "torch.float16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 83, + 83 + ], + "stream": "text", + "dtype": "torch.bfloat16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 104, + 58 + ], + "stream": "image", + "dtype": "torch.float32", + "table_bitwise": true, + "forward_max_abs": 2.384185791015625e-07, + "dx_max_abs": 2.384185791015625e-07 + }, + { + "grid": [ + 104, + 58 + ], + "stream": "image", + "dtype": "torch.float16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 104, + 58 + ], + "stream": "image", + "dtype": "torch.bfloat16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 104, + 58 + ], + "stream": "text", + "dtype": "torch.float32", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 104, + 58 + ], + "stream": "text", + "dtype": "torch.float16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + }, + { + "grid": [ + 104, + 58 + ], + "stream": "text", + "dtype": "torch.bfloat16", + "table_bitwise": true, + "forward_max_abs": 0.0, + "dx_max_abs": 0.0 + } + ] +} diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/ws1-cuda.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/ws1-cuda.json new file mode 100644 index 000000000..3a3e06daf --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/ws1-cuda.json @@ -0,0 +1,1962 @@ +{ + "scope": "Qwen-Image multi_axis_rope operator; C1 elementwise and C2 config transforms", + "trace": { + "contract": "qwenimage-interleaved-rope-fp32-v1", + "actual_backend": "rlkernel.cuda.qwenimage_multi_axis_rope.v1", + "kernel_id": "torch.ops.rl_kernel_native.multi_axis_rope", + "layout": "BSHD, adjacent channel pairs", + "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", + "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" + }, + "passed": true, + "forward": { + "op_name": "multi_axis_rope", + "backend_profile": "cuda_bf16", + "accuracy_reports": [ + { + "config_id": "BN/full", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.054684440139681e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/chunked", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/chunked", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.054684440139681e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s0", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/full/s0", + "fp32_reference" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004673302173614502, + "mean_abs_error": 6.917398422956467e-05, + "max_rel_error": 0.003753266530111432, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s1", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/full/s1", + "fp32_reference" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.00048516690731048584, + "mean_abs_error": 7.802648178767413e-05, + "max_rel_error": 0.0037793784867972136, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s2", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/full/s2", + "fp32_reference" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004864335060119629, + "mean_abs_error": 7.697803084738553e-05, + "max_rel_error": 0.003821910358965397, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s3", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/full/s3", + "fp32_reference" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.710371912457049e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s0", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/chunked/s0", + "fp32_reference" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004673302173614502, + "mean_abs_error": 6.917398422956467e-05, + "max_rel_error": 0.003753266530111432, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s1", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/chunked/s1", + "fp32_reference" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.00048516690731048584, + "mean_abs_error": 7.802648178767413e-05, + "max_rel_error": 0.0037793784867972136, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s2", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/chunked/s2", + "fp32_reference" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004864335060119629, + "mean_abs_error": 7.697803084738553e-05, + "max_rel_error": 0.003821910358965397, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s3", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/chunked/s3", + "fp32_reference" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.710371912457049e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/permuted", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/permuted", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.054684440139681e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/padded_right", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/padded_right", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.054684440139681e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/padded_left", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/padded_left", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.054684440139681e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + } + ], + "invariance_reports": [ + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/chunked", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "BN/chunked" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s0", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s0" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s1", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s1" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s2", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s2" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s3", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s3" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s0", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s0" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s1", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s1" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s2", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s2" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s3", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s3" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/permuted", + "transform_kind": "permutation", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "BN/permuted" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/padded_right", + "transform_kind": "padding", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "BN/padded_right" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/padded_left", + "transform_kind": "padding", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "BN/padded_left" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + } + ], + "logprob_smoke": null, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + }, + "candidate_id": "rlkernel.cuda.qwenimage_multi_axis_rope.v1", + "device": "cuda", + "compute_capability": "sm120", + "seed": 20260812, + "fallback_reason": null, + "passed": true, + "provenance_valid": true, + "metadata_valid": true, + "observed_kernel_id": "torch.ops.rl_kernel_native.multi_axis_rope" + }, + "gradient": { + "op_name": "multi_axis_rope", + "backend_profile": "cuda_bf16", + "accuracy_reports": [ + { + "config_id": "BN/full", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.699562850873917e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/chunked", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/chunked", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.699562850873917e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s0", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/full/s0", + "fp32_reference" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 2.7887988835573196e-05, + "mean_abs_error": 3.6821973026235355e-06, + "max_rel_error": 1.8164952993392944, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s1", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/full/s1", + "fp32_reference" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.7113780965446495e-06, + "max_rel_error": 0.4316565990447998, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s2", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/full/s2", + "fp32_reference" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 2.995925024151802e-05, + "mean_abs_error": 4.716822331829462e-06, + "max_rel_error": 0.7439075112342834, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s3", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/full/s3", + "fp32_reference" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.0211638659238815e-05, + "mean_abs_error": 4.9605891945248e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s0", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/chunked/s0", + "fp32_reference" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 2.7887988835573196e-05, + "mean_abs_error": 3.6821973026235355e-06, + "max_rel_error": 1.8164952993392944, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s1", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/chunked/s1", + "fp32_reference" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.7113780965446495e-06, + "max_rel_error": 0.4316565990447998, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s2", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/chunked/s2", + "fp32_reference" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 2.995925024151802e-05, + "mean_abs_error": 4.716822331829462e-06, + "max_rel_error": 0.7439075112342834, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s3", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/chunked/s3", + "fp32_reference" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.0211638659238815e-05, + "mean_abs_error": 4.9605891945248e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/permuted", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/permuted", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.699562850873917e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/padded_right", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/padded_right", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.699562850873917e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/padded_left", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/padded_left", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.699562850873917e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + } + ], + "invariance_reports": [ + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/chunked", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "BN/chunked" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s0", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s0" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s1", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s1" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s2", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s2" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s3", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s3" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s0", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s0" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s1", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s1" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s2", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s2" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s3", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s3" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/permuted", + "transform_kind": "permutation", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "BN/permuted" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/padded_right", + "transform_kind": "padding", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "BN/padded_right" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/padded_left", + "transform_kind": "padding", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "BN/padded_left" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + } + ], + "singleton_aggregate_reports": [], + "backend_provenance": { + "backend_profile": "cuda_bf16", + "requested_backend": "cuda", + "actual_backend": "cuda", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + }, + "candidate_id": "rlkernel.cuda.qwenimage_multi_axis_rope.v1", + "device": "cuda", + "compute_capability": "sm120", + "seed": 20260812, + "fallback_reason": null, + "passed": true, + "provenance_valid": true, + "metadata_valid": true, + "loss_reduction": "sum_over_active_tokens_then_optional_mean_by_active_count", + "active_token_denominator": 27, + "grad_tensor_names": [ + "dx" + ], + "first_failing_op": null, + "first_failing_tensor": null, + "first_failing_config_pair": null, + "observed_kernel_id": "torch.ops.rl_kernel_native.multi_axis_rope" + } +} diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/ws1-triton.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/ws1-triton.json new file mode 100644 index 000000000..628fcbf70 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/ws1-triton.json @@ -0,0 +1,1962 @@ +{ + "scope": "Qwen-Image multi_axis_rope operator; C1 elementwise and C2 config transforms", + "trace": { + "contract": "qwenimage-interleaved-rope-fp32-v1", + "actual_backend": "rlkernel.triton.qwenimage_multi_axis_rope.v1", + "kernel_id": "rl_engine.backends.shared.triton.rope._flash_attn_rotary.rotary_kernel", + "layout": "BSHD, adjacent channel pairs", + "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", + "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" + }, + "passed": true, + "forward": { + "op_name": "multi_axis_rope", + "backend_profile": "triton_cuda_bf16", + "accuracy_reports": [ + { + "config_id": "BN/full", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.054684440139681e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/chunked", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/chunked", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.054684440139681e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s0", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/full/s0", + "fp32_reference" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004673302173614502, + "mean_abs_error": 6.917398422956467e-05, + "max_rel_error": 0.003753266530111432, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s1", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/full/s1", + "fp32_reference" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.00048516690731048584, + "mean_abs_error": 7.802648178767413e-05, + "max_rel_error": 0.0037793784867972136, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s2", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/full/s2", + "fp32_reference" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004864335060119629, + "mean_abs_error": 7.697803084738553e-05, + "max_rel_error": 0.003821910358965397, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s3", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/full/s3", + "fp32_reference" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.710371912457049e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s0", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/chunked/s0", + "fp32_reference" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004673302173614502, + "mean_abs_error": 6.917398422956467e-05, + "max_rel_error": 0.003753266530111432, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s1", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/chunked/s1", + "fp32_reference" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.00048516690731048584, + "mean_abs_error": 7.802648178767413e-05, + "max_rel_error": 0.0037793784867972136, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s2", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/chunked/s2", + "fp32_reference" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004864335060119629, + "mean_abs_error": 7.697803084738553e-05, + "max_rel_error": 0.003821910358965397, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s3", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "B1-singleton_aggregate/chunked/s3", + "fp32_reference" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.710371912457049e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/permuted", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/permuted", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.054684440139681e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/padded_right", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/padded_right", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.054684440139681e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/padded_left", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/padded_left", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 0.0004878789186477661, + "mean_abs_error": 8.054684440139681e-05, + "max_rel_error": 0.0038475203327834606, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "forward_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + } + ], + "invariance_reports": [ + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/chunked", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "BN/chunked" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s0", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s0" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s1", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s1" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s2", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s2" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s3", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s3" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s0", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s0" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s1", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s1" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s2", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s2" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s3", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s3" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/permuted", + "transform_kind": "permutation", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "BN/permuted" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/padded_right", + "transform_kind": "padding", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "BN/padded_right" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/padded_left", + "transform_kind": "padding", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "output", + "config_pair": [ + "BN/full", + "BN/padded_left" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "forward_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + } + ], + "logprob_smoke": null, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + }, + "candidate_id": "rlkernel.triton.qwenimage_multi_axis_rope.v1", + "device": "cuda", + "compute_capability": "sm120", + "seed": 20260812, + "fallback_reason": null, + "passed": true, + "provenance_valid": true, + "metadata_valid": true, + "observed_kernel_id": "rl_engine.backends.shared.triton.rope._flash_attn_rotary.rotary_kernel" + }, + "gradient": { + "op_name": "multi_axis_rope", + "backend_profile": "triton_cuda_bf16", + "accuracy_reports": [ + { + "config_id": "BN/full", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.699562850873917e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/chunked", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/chunked", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.699562850873917e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s0", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/full/s0", + "fp32_reference" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 2.7887988835573196e-05, + "mean_abs_error": 3.6821973026235355e-06, + "max_rel_error": 1.8164952993392944, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s1", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/full/s1", + "fp32_reference" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.7113780965446495e-06, + "max_rel_error": 0.4316565990447998, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s2", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/full/s2", + "fp32_reference" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 2.995925024151802e-05, + "mean_abs_error": 4.716822331829462e-06, + "max_rel_error": 0.7439075112342834, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/full/s3", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/full/s3", + "fp32_reference" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.0211638659238815e-05, + "mean_abs_error": 4.9605891945248e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s0", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/chunked/s0", + "fp32_reference" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 2.7887988835573196e-05, + "mean_abs_error": 3.6821973026235355e-06, + "max_rel_error": 1.8164952993392944, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s1", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/chunked/s1", + "fp32_reference" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.7113780965446495e-06, + "max_rel_error": 0.4316565990447998, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s2", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/chunked/s2", + "fp32_reference" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 2.995925024151802e-05, + "mean_abs_error": 4.716822331829462e-06, + "max_rel_error": 0.7439075112342834, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "B1-singleton_aggregate/chunked/s3", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "B1-singleton_aggregate/chunked/s3", + "fp32_reference" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.0211638659238815e-05, + "mean_abs_error": 4.9605891945248e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/permuted", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/permuted", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.699562850873917e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/padded_right", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/padded_right", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.699562850873917e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + }, + { + "config_id": "BN/padded_left", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/padded_left", + "fp32_reference" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "float32", + "max_abs_error": 3.213249146938324e-05, + "mean_abs_error": 4.699562850873917e-06, + "max_rel_error": 5.814913272857666, + "atol": 0.02, + "rtol": 0.016, + "passed": true, + "judgment": "gradient_accuracy", + "comparison_lhs_role": "bf16_candidate", + "comparison_rhs_role": "fp32_reference" + } + ], + "passed": true, + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + } + } + ], + "invariance_reports": [ + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/chunked", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "BN/chunked" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s0", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s0" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s1", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s1" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s2", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s2" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/full/s3", + "transform_kind": "batch_size", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/full/s3" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s0", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s0" + ], + "shape": [ + 3, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s1", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s1" + ], + "shape": [ + 8, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s2", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s2" + ], + "shape": [ + 5, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "B1-singleton_aggregate/chunked/s3", + "transform_kind": "chunk", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "B1-singleton_aggregate/chunked/s3" + ], + "shape": [ + 11, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/permuted", + "transform_kind": "permutation", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "BN/permuted" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/padded_right", + "transform_kind": "padding", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "BN/padded_right" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + }, + { + "canonical_config_id": "BN/full", + "transformed_config_id": "BN/padded_left", + "transform_kind": "padding", + "op_class": "elementwise", + "dtype": "bfloat16", + "backend_profile": "triton_cuda_bf16", + "details": [ + { + "tensor_name": "dx", + "config_pair": [ + "BN/full", + "BN/padded_left" + ], + "shape": [ + 27, + 3, + 128 + ], + "dtype": "bfloat16", + "max_abs_error": 0.0, + "mean_abs_error": 0.0, + "max_rel_error": 0.0, + "atol": 0.0, + "rtol": 0.0, + "passed": true, + "judgment": "gradient_invariance", + "comparison_lhs_role": "transformed_config", + "comparison_rhs_role": "canonical_config" + } + ], + "passed": true + } + ], + "singleton_aggregate_reports": [], + "backend_provenance": { + "backend_profile": "triton_cuda_bf16", + "requested_backend": "triton", + "actual_backend": "triton", + "execution_dtype": "bfloat16", + "accumulation_dtype": "float32", + "output_dtype": "bfloat16", + "reference_dtype": "float32", + "candidate_tf32_enabled": false, + "reference_tf32_enabled": false + }, + "candidate_id": "rlkernel.triton.qwenimage_multi_axis_rope.v1", + "device": "cuda", + "compute_capability": "sm120", + "seed": 20260812, + "fallback_reason": null, + "passed": true, + "provenance_valid": true, + "metadata_valid": true, + "loss_reduction": "sum_over_active_tokens_then_optional_mean_by_active_count", + "active_token_denominator": 27, + "grad_tensor_names": [ + "dx" + ], + "first_failing_op": null, + "first_failing_tensor": null, + "first_failing_config_pair": null, + "observed_kernel_id": "rl_engine.backends.shared.triton.rope._flash_attn_rotary.rotary_kernel" + } +} diff --git a/tools/validation/operators/check_multi_axis_rope_install.py b/tools/validation/operators/check_multi_axis_rope_install.py index 56cfe2ccb..4fa42d15f 100644 --- a/tools/validation/operators/check_multi_axis_rope_install.py +++ b/tools/validation/operators/check_multi_axis_rope_install.py @@ -13,6 +13,7 @@ import torch import rl_engine +from rl_engine.backends import extension from rl_engine.ops.rope.multi_axis_rope import ( build_multi_axis_cos_sin, get_multi_axis_rope_op, @@ -50,6 +51,7 @@ def main(): report = { "package": str(package), "companion_library": native.origin if native else None, + "full_extension_library": getattr(extension._C, "__file__", None), "native_dispatcher_registered": hasattr(torch.ops.rl_kernel_native, "multi_axis_rope"), "torch": torch.__version__, "gpu": torch.cuda.get_device_name(), From 84751473006401c19c1c8c17844e8ef8b1c44375 Mon Sep 17 00:00:00 2001 From: dandan20031031 <2425926169@qq.com> Date: Wed, 28 Oct 2026 22:00:50 +0800 Subject: [PATCH 3/5] test(rope): record Python 3.10 CPU qualification Signed-off-by: dandan20031031 <2425926169@qq.com> --- .../README.md | 14 ++++++++----- .../cpu310-contracts.txt | 20 +++++++++++++++++++ .../cpu310-environment.json | 9 +++++++++ .../manifest.json | 4 +++- 4 files changed, 41 insertions(+), 6 deletions(-) create mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu310-contracts.txt create mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu310-environment.json diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/README.md b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/README.md index ddc879318..9ef00fdb1 100644 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/README.md +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/README.md @@ -48,7 +48,8 @@ SM120; bare-metal Linux and A100/H100 remain separate hardware profiles. | Related regressions | 195 passed, 7 Ascend skips | `regression.txt` | | Repository layout with extension enforcement enabled | 37 passed, zero skips/failures | `layout.txt` | | Fresh installed full wheel, outside source checkout | Eager/compiled forward and input-gradient bytes exact for both backends | `installed-wheel.json` | -| Actual CPU-only PyTorch wheel/reference installation | 57 passed, 52 expected GPU skips; reference forward/backward passed | `cpu-contracts.txt`, `cpu-environment.json` | +| Actual CPU-only PyTorch wheel/reference installation, Python 3.10.22 | 57 passed, 52 expected GPU skips; reference forward/backward passed | `cpu310-contracts.txt`, `cpu310-environment.json` | +| Actual CPU-only PyTorch wheel/reference installation, Python 3.12.3 | 57 passed, 52 expected GPU skips; reference forward/backward passed | `cpu-contracts.txt`, `cpu-environment.json` | | Compute Sanitizer accuracy/stride/empty/stream subset | 34 passed; zero errors | `memcheck.txt` | | Pinned actual Diffusers symbols | 30 forward/backward cases passed; image/text tables exact | `upstream.json` | | Reference-grid benchmark correctness | 40 comparisons passed | `benchmark-bfloat16.json` | @@ -56,8 +57,10 @@ SM120; bare-metal Linux and A100/H100 remain separate hardware profiles. The full-extension GPU gate rejects missing native/Triton execution. Its only allowed operator-test skip is the two-device guard on a one-GPU runner. -The CPU-only result is separate: it uses `torch==2.9.1+cpu`, and its installed -package contains no native CUDA binary. CPU skips are not GPU acceptance evidence. +The CPU-only results are separate: both Python versions use `torch==2.9.1+cpu`, +and their installed packages contain no native CUDA binary. CPU skips are not +GPU acceptance evidence. Both runs install the normal CPU package and exercise +the workflow's three test files locally; these are not hosted CI results. The 37 passing Linux layout tests contrast with the five Windows baseline failures recorded earlier; those historical logs were not rewritten. @@ -123,8 +126,9 @@ The new [operator gate](../../../ci/scripts/run_qwenimage_rope_ws1.sh) requires the complete `_C` registration, records actual library/backend identity and enforces the GPU skip policy. The [CPU workflow](../../../.github/workflows/qwenimage-rope.yml) targets `test-qwenimage`, runs Python 3.10/3.12 CPU contracts, checks changed files -and builds documentation. This report's actual CPU-only run used Python 3.12; -hosted CI status is tracked on the PR. GPU runners use the explicit gate. +and builds documentation. This report includes local CPU-only runs on both +Python 3.10 and 3.12; hosted CI status is tracked on the PR. GPU runners use the +explicit gate. Bare-metal Linux, A100/H100, two-GPU guards, QK-norm fusion, model/block assembly and trajectories remain unqualified. Building the entire library does not diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu310-contracts.txt b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu310-contracts.txt new file mode 100644 index 000000000..6f0703516 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu310-contracts.txt @@ -0,0 +1,20 @@ +...........ssssssssssss.ss.ss.ss...ssssss.ss......ssssssssssss.ss...ssss [ 66%] +s.s.ssssss........................... [100%] +=========================== short test summary info ============================ +SKIPPED [12] tests/ops/rope/test_multi_axis_rope.py:79: NVIDIA CUDA required +SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:100: NVIDIA CUDA required +SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:110: NVIDIA CUDA required +SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:133: NVIDIA CUDA required +SKIPPED [6] tests/ops/rope/test_multi_axis_rope.py:148: NVIDIA CUDA required +SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:158: NVIDIA CUDA required +SKIPPED [12] tests/ops/rope/test_multi_axis_rope.py:174: NVIDIA CUDA required +SKIPPED [1] tests/ops/rope/test_multi_axis_rope.py:217: native CUDA multi_axis_rope extension required +SKIPPED [1] tests/ops/rope/test_multi_axis_rope.py:230: NVIDIA CUDA required +SKIPPED [3] tests/ops/rope/test_multi_axis_rope.py:255: NVIDIA CUDA required +SKIPPED [1] tests/ops/rope/test_multi_axis_rope.py:315: NVIDIA CUDA required +SKIPPED [1] tests/ops/rope/test_multi_axis_rope.py:336: requires two NVIDIA GPUs +SKIPPED [1] tests/ops/rope/test_multi_axis_rope.py:355: CUDA context requires CUDA +SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:368: NVIDIA CUDA required +SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:399: NVIDIA CUDA required +SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:427: NVIDIA CUDA required +57 passed, 52 skipped in 1.72s diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu310-environment.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu310-environment.json new file mode 100644 index 000000000..833e69038 --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu310-environment.json @@ -0,0 +1,9 @@ +{ + "python": "3.10.22", + "torch": "2.9.1+cpu", + "cuda": null, + "package": "/home/dandan/venvs/rlk-rope-cpu310/lib/python3.10/site-packages/rl_engine", + "reference_forward_backward": true, + "cpu_wheel_contains_no_cuda_binary": true, + "passed": true +} diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/manifest.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/manifest.json index 3acb62a79..9db29dacc 100644 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/manifest.json +++ b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/manifest.json @@ -30,9 +30,12 @@ "tools/validation/operators/check_multi_axis_rope_install.py": "f295fc8ac5878ac13b6255c2f0be83b66d34ff2ad2f75111bb8b633508a865a7" }, "evidence_sha256": { + "README.md": "3ad6f8ef514b4fc61d51fb0ce8cb6d420d2180251c9ee4f84b8c931aa3a75094", "benchmark-bfloat16.json": "e9c1e4b1085623d84800656c179e4a0e86d70a0ef7fc9ed92f9e6ccc50f0f491", "cpu-contracts.txt": "9518f53a6efc603bca1c52f160e5430f1006bb2abc9a7ef57694dae6feac5d78", "cpu-environment.json": "4dfc68f3d76af05a050d830840b4a6b8c8fa0a4c12c29d11d632a56a57a2d4c6", + "cpu310-contracts.txt": "02f57808b9553cefe6f4dcec06d46a5163157b4290cc0c590b344a1616057852", + "cpu310-environment.json": "7130d4c164c98804284cc0f1e07ae6265c35e7498c2fb8285a5e6df87cb010b2", "full-extension.json": "8fae782e3c5e6aae59de6fe6b4e7ae185729e49d2282dbd0cb3b9ca08c0274cd", "full-inplace-build.txt": "04d5ef122650ab305847d58850a6c3b44438eacf97920492e1eb1c5580e9a21e", "full-wheel-build.txt": "a7883dfb1211798c869ede1ee0b6f1dfb32edd7b27b359ee18c6e5ab9fe449d5", @@ -47,7 +50,6 @@ "precommit.txt": "7cf90e06cf73194271c176ebc2d147e64cf8f5ed522d8e0e03093a4647f63759", "python-packages.json": "617f8cef5a24f3b2a683e70e0b491e47db6d24a7ca848e5b3e461e87993a7201", "quality.txt": "8826f0e406946a923400d792dd60d0f545b32f9fcffa3a148268c969e30e14c0", - "README.md": "af365e658fe4875df643704b2495e1f115c6c889f5d449880233d49e511040ac", "regression.txt": "039cd9813a2ea54bfc5b7750fc0555b54d2bf30b09dbca2eba4a0403578ab9b0", "tests-and-ws1.txt": "68397e27144c9244743dbcff39ffdca8a77a30eaf1e8180b2ca84dfc1216d633", "upstream.json": "04e5472a82477a305b996ccbb420d0f3cc7b39af4053c867d37916a9bf5de51a", From 5931ce7b6d95799a6c19e0d3ae3eac958a768acf Mon Sep 17 00:00:00 2001 From: dandan20031031 <2425926169@qq.com> Date: Wed, 28 Oct 2026 22:35:46 +0800 Subject: [PATCH 4/5] docs(rope): consolidate validation evidence for review Signed-off-by: dandan20031031 <2425926169@qq.com> --- .../qwenimage_multi_axis_rope_sm120.md | 122 + .../README.md | 152 - .../baseline-layout.txt | 138 - .../benchmark-bfloat16.json | 2781 ----------------- .../installed-wheel.json | 37 - .../installed-wheel.txt | 38 - .../layout.txt | 138 - .../manifest.json | 65 - .../memcheck.txt | 8 - .../operator-cuda.json | 70 - .../operator-triton.json | 70 - .../package-build.txt | 40 - .../package.json | 31 - .../precommit.txt | 7 - .../quality.txt | 58 - .../regression.txt | 10 - .../tests.txt | 5 - .../upstream.json | 335 -- .../ws1-cuda.json | 1962 ------------ .../ws1-cuda.txt | 3 - .../ws1-triton.json | 1962 ------------ .../ws1-triton.txt | 3 - .../README.md | 136 - .../benchmark-bfloat16.json | 2781 ----------------- .../cpu-contracts.txt | 20 - .../cpu-environment.json | 7 - .../cpu310-contracts.txt | 20 - .../cpu310-environment.json | 9 - .../full-extension.json | 22 - .../full-inplace-build.txt | 20 - .../full-wheel-build.txt | 10 - .../gpu-environment.json | 15 - .../installed-wheel.json | 38 - .../layout.txt | 2 - .../manifest.json | 60 - .../memcheck.txt | 4 - .../operator-cuda.json | 70 - .../operator-tests.xml | 1 - .../operator-triton.json | 70 - .../package.json | 11 - .../precommit.txt | 7 - .../python-packages.json | 114 - .../quality.txt | 34 - .../regression.txt | 10 - .../tests-and-ws1.txt | 11 - .../upstream.json | 335 -- .../ws1-cuda.json | 1962 ------------ .../ws1-triton.json | 1962 ------------ 48 files changed, 122 insertions(+), 15644 deletions(-) create mode 100644 reports/experiments/qwenimage_multi_axis_rope_sm120.md delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/README.md delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/baseline-layout.txt delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/benchmark-bfloat16.json delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/installed-wheel.json delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/installed-wheel.txt delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/layout.txt delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/manifest.json delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/memcheck.txt delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/operator-cuda.json delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/operator-triton.json delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/package-build.txt delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/package.json delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/precommit.txt delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/quality.txt delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/regression.txt delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/tests.txt delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/upstream.json delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-cuda.json delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-cuda.txt delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-triton.json delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-triton.txt delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/README.md delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/benchmark-bfloat16.json delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu-contracts.txt delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu-environment.json delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu310-contracts.txt delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu310-environment.json delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/full-extension.json delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/full-inplace-build.txt delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/full-wheel-build.txt delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/gpu-environment.json delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/installed-wheel.json delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/layout.txt delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/manifest.json delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/memcheck.txt delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/operator-cuda.json delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/operator-tests.xml delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/operator-triton.json delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/package.json delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/precommit.txt delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/python-packages.json delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/quality.txt delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/regression.txt delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/tests-and-ws1.txt delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/upstream.json delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/ws1-cuda.json delete mode 100644 reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/ws1-triton.json diff --git a/reports/experiments/qwenimage_multi_axis_rope_sm120.md b/reports/experiments/qwenimage_multi_axis_rope_sm120.md new file mode 100644 index 000000000..a7d82f65b --- /dev/null +++ b/reports/experiments/qwenimage_multi_axis_rope_sm120.md @@ -0,0 +1,122 @@ +# Qwen-Image multi-axis RoPE: SM120 validation + +Recorded 2026-10-28 for [PR #623](https://github.com/RL-Align/RL-Kernel/pull/623). +The numerical/native implementation is based on `ac17ab87e8d93dabe93dfdc134e9853d47de7af1`. +The [operator documentation](../../docs/operators/multi-axis-rope.md) defines the +contract, upstream reuse, API and detailed reproduction commands. + +## Environments + +Both profiles use one RTX 5080 (SM120, 16 GB), Windows driver 617.42, +PyTorch 2.9.1+cu128 and CUDA Toolkit 12.8. They are separate recorded runs, +not a controlled Windows-versus-WSL performance comparison. + +| Item | WSL2 Linux | Windows | +| --- | --- | --- | +| OS | Ubuntu 24.04.5, glibc 2.39; WSL 3.0.3 on Windows 11 | Windows | +| Kernel | 6.18.54.1-microsoft-standard-WSL2 | Native Windows | +| Python | 3.12.3 | 3.13.9 | +| Compiler | CUDA 12.8.93, GCC 13.3 | CUDA 12.8, MSVC 2022 x64 | +| Triton | 3.5.1 | triton-windows 3.5.1.post24 | +| Native build | Ordinary complete `rl_engine._C`, all 10 configured translation units | Explicit operator-only `rl_engine._rope_cuda` | + +The WSL checkout, environments and caches are in the Linux filesystem. +The normal full wheel was built with `RL_KERNEL_ROPE_ONLY`, forced SM90 and +fast-math flags unset. Linking existing CUDA driver calls used +`LIBRARY_PATH=/usr/lib/wsl/lib`; the Windows driver supplies CUDA driver access. + +## Recorded results + +| Check | WSL2 Linux | Windows | +| --- | --- | --- | +| Public operator suite | 81 passed, 1 two-GPU guard skip | 81 passed, 1 two-GPU guard skip | +| Related regressions | 195 passed, 7 Ascend skips | 195 passed, 7 Ascend skips | +| CUDA/Triton WS1 forward/gradient accuracy and invariance | Both passed with backend provenance | Both passed with backend provenance | +| Standard operator CLI forward/input gradient | Both backends passed | Both backends passed | +| Pinned actual Diffusers symbols | 30 cases passed; image/text tables exact | 30 cases passed; image/text tables exact | +| Compute Sanitizer | 34 selected cases passed, zero errors | 34 selected cases passed, zero errors | +| Fresh installed wheel outside checkout | Full wheel: both backends' eager/compiled forward and input-gradient bytes exact | Operator wheel: same checks passed | +| CPU-only PyTorch 2.9.1+cpu, Python 3.10.22 and 3.12.3 | Each: 57 passed, 52 expected GPU skips; installed reference forward/backward passed | Not recorded | +| Repository layout | 37 passed, zero skips/failures | 31 passed, 5 failures, 1 skip | +| Changed-file static checks and strict docs | Passed locally | Passed locally | + +The operator suite covers FP32/FP16/BF16, independent accuracy, byte invariance, +packed-QKV/no-copy and table strides, empty inputs, higher input VJPs, stream +selection, graph replay, opcheck and fullgraph forward/backward. The installed +wheels contain the FlashAttention BSD license; the CPU packages contain no CUDA +binary. Local CPU checks are separate from hosted CI and GPU acceptance. + +The five historical Windows layout failures also reproduce on its clean base +(27 passed, 5 failures, 1 skip): driver linking and Windows symlink privileges or +placeholders. The later Linux run preserved all 27 tracked symlinks and passed. +The earlier Windows logs remain unchanged in the archive. + +## Standalone benchmark + +WSL2 BF16, batch 1, 24 heads, D=128. CUDA Graph medians in ms/call; +each cell is **forward / forward+backward**. + +| Grid / layout | Native CUDA | Adapted Triton | FlashAttention, no FMA | Diffusers complex | +| --- | --- | --- | --- | --- | +| 64x64 contiguous | 0.05780 / 0.12967 | 0.04929 / 0.12414 | 0.04898 / 0.12407 | 0.24446 / 0.51281 | +| 64x64 packed QKV | 0.05868 / 0.13079 | 0.04731 / 0.11841 | 0.04722 / 0.11847 | 0.24754 / 0.51583 | +| 83x83 contiguous | 0.11198 / 0.22392 | 0.10720 / 0.22538 | 0.10739 / 0.21420 | 0.52579 / 1.06721 | +| 83x83 packed QKV | 0.11233 / 0.22486 | 0.10745 / 0.21478 | 0.10744 / 0.21469 | 0.52820 / 1.06947 | +| 104x58 contiguous | 0.09797 / 0.19708 | 0.09366 / 0.18497 | 0.09362 / 0.18458 | 0.44992 / 0.90110 | +| 104x58 packed QKV | 0.09610 / 0.19612 | 0.09310 / 0.18828 | 0.09323 / 0.18817 | 0.44979 / 0.90276 | + +All 40 pre-timing comparisons passed; CUDA, adapted Triton and the no-FMA +baseline had zero differing output/input-gradient bytes against the FP32 gold. +The archive includes 80 measurements, the short case, default FlashAttention, +eager wall timings and peak allocation. At 4096 tokens, either GPU path's forward +extra peak is 24 MiB, the returned output size, for both layouts. + +Each graph contains 20 calls, with five warm calls/replays and 30 measured +replays. Eager wall timings use three synchronized trials of 30 calls. +Tables, compilation and input/backing allocation are excluded; extra peak +includes the returned output or gradient. The upstream source/autograd wrapper +is pinned to `98c1cee0ed59307833841cdbf567fe336ceec13a`; the no-FMA variant changes +only contraction. These single-process runs use a shared GPU without clock +pinning and establish no advantage over FlashAttention or model speedup. + +## Reproduction + +Use a matching CUDA 12.8/PyTorch environment in a Linux checkout: + +```bash +export CUDA_HOME=/usr/local/cuda-12.8 +export PATH="$CUDA_HOME/bin:$PATH" +export LIBRARY_PATH="/usr/lib/wsl/lib${LIBRARY_PATH:+:$LIBRARY_PATH}" +export MAX_JOBS=2 TORCH_CUDA_ARCH_LIST=12.0 RL_KERNEL_REQUIRE_EXT=1 +unset RL_KERNEL_ROPE_ONLY KERNEL_ALIGN_FORCE_SM90 KERNEL_ALIGN_USE_FAST_MATH +python -m pip wheel --no-build-isolation --no-deps . -w /tmp/rope-full-wheel +python -m pip install --force-reinstall --no-deps /tmp/rope-full-wheel/rl_kernel-*.whl +python setup.py build_ext --inplace +ROPE_PYTHON=python ROPE_WS1_OUTPUT=/tmp/rope-ws1 bash ci/scripts/run_qwenimage_rope_ws1.sh +python ci/run.py layout -q -rs +compute-sanitizer --tool memcheck --error-exitcode 1 python -m pytest \ + tests/ops/rope/test_multi_axis_rope.py -q \ + -k "forward_backward_accuracy or noncontiguous_input_tables_and_gradient or empty_forward_backward or packed_qkv_and_independent_table_strides or native_entrypoint_checks" +``` + +The [GPU gate](../../ci/scripts/run_qwenimage_rope_ws1.sh) requires the full +native extension and Triton, records loaded binary/backend identity, and permits +only the two-GPU guard to skip. The [CPU workflow](../../.github/workflows/qwenimage-rope.yml) +contains the Python 3.10/3.12 contract commands. The operator documentation covers +installed-wheel probes, pinned upstream audits and benchmark reproduction. +Hosted CI status is tracked on the PR. + +## Archived evidence and qualification limits + +Detailed per-case JSON, JUnit, build/test logs, dependency versions, and +source/binary SHA256 manifests remain available at immutable commit +`84751473006401c19c1c8c17844e8ef8b1c44375`: + +- [WSL2 complete-extension evidence](https://github.com/dandan20031031/RL-Kernel/tree/84751473006401c19c1c8c17844e8ef8b1c44375/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120) +- [Historical Windows operator-profile evidence](https://github.com/dandan20031031/RL-Kernel/tree/84751473006401c19c1c8c17844e8ef8b1c44375/reports/experiments/qwenimage_multi_axis_rope_v2_sm120) + +This summary preserves those recorded results; consolidating reports does not +rerun the numerical suites. Bare-metal Linux, A100/H100, two-GPU guards, +QK-norm fusion, model/block assembly and trajectories remain unqualified. +Building the full library does not certify every native operator or Dense Qwen3. +No distributed or cross-version/device byte-equality claim is made. diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/README.md b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/README.md deleted file mode 100644 index 64232c731..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/README.md +++ /dev/null @@ -1,152 +0,0 @@ -# Qwen-Image multi-axis RoPE validation — 2026-10-28 - -Evidence for the CUDA `multi_axis_rope` row of -[RFC #386](https://github.com/RL-Align/RL-Kernel/issues/386), based on -`test-qwenimage@16b81175874ecf11d5ff05c00de3b69e9457591b`. -The source fingerprints in `manifest.json` identify the final implementation. -This is single-GPU operator evidence; it does not qualify a model or trajectory. - -## Implementation - -- Adjacent pairs, axes `[16, 56, 56]`, canonical centered image positions and - diagonal text offsets, pinned to Diffusers `40f63690da440e8c382ab53e67508835deb1ed2f`. -- Explicit CPU table factories remain CPU factories inside CUDA/meta default-device contexts. -- Native CUDA reads all input/table strides directly, including packed QKV views. - The shared Triton implementation adapts FlashAttention's existing rotary kernel - at `98c1cee0ed59307833841cdbf567fe336ceec13a`, retaining attribution and BSD-3 terms. -- Both GPU paths have dispatcher, fake and input-autograd registration. The - arithmetic is separate FP32 multiply/add followed by one output cast. -- OP_SPECS, the ordinary input builder and a gradient adapter expose the operator - to existing validation APIs. The operator gate runs all four forward/gradient - accuracy/invariance judgments with explicit backend/kernel provenance. -- Default CUDA builds include the native source in `_C`. The explicit - `RL_KERNEL_ROPE_ONLY=1` build profile also produces an installable companion - library, without compiling unrelated native operators. - -## Environment - -| Item | Value | -| --- | --- | -| GPU | NVIDIA GeForce RTX 5080, SM120, 16 GB | -| Driver | 617.42 | -| OS / Python | Windows / 3.13.9 | -| PyTorch / CUDA toolkit | 2.9.1+cu128 / 12.8 | -| Triton | triton-windows 3.5.1.post24; runtime 3.5.1 | -| Native compiler | MSVC 2022 x64, C++17; CUDA target 12.0 | -| Sanitizer | NVIDIA Compute Sanitizer 13.3, memcheck | -| Benchmark dtype | BF16, FP32 position tables | - -The installed wheel is CPython 3.13/Windows-specific. Toolkit/driver versions -have different roles; the native code was compiled with CUDA 12.8, not 13.3. - -## Validation - -| Check | Result | Evidence | -| --- | --- | --- | -| Public operator suite using the built companion library | 81 passed, 1 two-GPU skip | `tests.txt` | -| Operator/build/registry/validation regressions | 195 passed, 7 Ascend skips | `regression.txt` | -| CUDA WS1 four judgments | Passed | `ws1-cuda.json` | -| Triton WS1 four judgments | Passed | `ws1-triton.json` | -| Standard CLI, BF16 forward and input gradient, CUDA/Triton | Both passed | `operator-{cuda,triton}.json` | -| Pinned actual Diffusers symbols | 30 cases passed; image/text tables exact | `upstream.json` | -| Memory safety on selected accuracy/strided/empty/stream cases | 34 passed, zero memcheck errors | `memcheck.txt` | -| Source distribution + operator wheel build | Passed; native sources and BSD license present | `package-build.txt`, `package.json` | -| Fresh interpreter outside checkout after wheel installation | CUDA/Triton eager and fullgraph forward/input-gradient bytes exact | `installed-wheel.json` | -| Changed Python formatting/static checks; strict docs build | Passed | `quality.txt` | -| Repository layout check | 5 failures, 31 passed, 1 skipped | `layout.txt` | -| Same layout check on clean base | Same 5 failures, 27 passed, 1 skipped | `baseline-layout.txt` | - -Operator tests cover FP32/FP16/BF16, shared/per-sample tables, odd and rectangular -grids, batch/chunk/padding invariance, higher input VJPs, strided gradients, -independent table strides, empty tensors, invalid-input rejection, non-default -streams, graph replay with changed static buffers, `torch.library.opcheck`, and -`torch.compile(..., fullgraph=True)` through backward. A CPU profiler assertion -rejects `contiguous`, `clone` or `copy_` inside either GPU call for a packed view. - -The layout failures reproduce on the clean base: the Windows driver-link flag -assertion, symlink creation without Windows privilege, and three checks affected -by symlink-disabled JSON/source placeholders. The added build tests pass. This -report does not describe the entire repository test suite as green. - -## Mature baseline and performance - -`benchmark-bfloat16.json` contains 40 pre-timing correctness comparisons and 80 -measurements: four geometries, two layouts, five implementations, two phases. -The CUDA, adapted Triton and contraction-disabled upstream FlashAttention paths -all have zero differing output/gradient bytes against the unfused FP32 gold. -Original FlashAttention and Diffusers complex arithmetic use declared tolerances. - -The optional FlashAttention comparison loads the checksum-pinned original -`rotary_kernel`, `apply_rotary` and `ApplyRotaryEmb` class without installing -flash-attn. `flash_no_fma` adds only `enable_fp_fusion=False` to the original -launch; the default baseline retains its fused arithmetic. - -CUDA Graph medians in ms/call; each cell is **forward / forward+backward**. -Batch 1, 24 heads, D=128: - -| Grid / layout | Native CUDA | Adapted Triton | FlashAttention, no FMA | Diffusers complex | -| --- | --- | --- | --- | --- | -| 64×64 contiguous | 0.05771 / 0.12953 | 0.05191 / 0.12463 | 0.05202 / 0.12443 | 0.24277 / 0.50118 | -| 64×64 packed QKV | 0.05866 / 0.13089 | 0.04835 / 0.12314 | 0.04896 / 0.12374 | 0.24599 / 0.50083 | -| 83×83 contiguous | 0.11203 / 0.22478 | 0.10685 / 0.21339 | 0.10693 / 0.21351 | 0.51009 / 1.05054 | -| 83×83 packed QKV | 0.11280 / 0.22514 | 0.10758 / 0.21415 | 0.10746 / 0.21431 | 0.51117 / 1.05461 | -| 104×58 contiguous | 0.09809 / 0.19646 | 0.09356 / 0.18708 | 0.09362 / 0.18709 | 0.43411 / 0.89356 | -| 104×58 packed QKV | 0.09596 / 0.19664 | 0.09341 / 0.18776 | 0.09344 / 0.18786 | 0.43668 / 0.89561 | - -The adapted Triton path is close to its mature upstream baseline. The native -CUDA path avoids requiring Triton but is slower in this graph comparison. -No speed advantage over FlashAttention is established. At 4096 tokens, either -GPU path's forward extra peak is 24 MiB, exactly the returned BF16 output size, -for both layouts. There is no extra packed-input materialization. - -Five warm calls/replays, 20 calls per captured graph, median of 30 replay -measurements. Synchronized eager wall time is also recorded as three trials of -30 calls. It includes Python/dispatcher/autograd overhead and can differ -substantially from graph time. Tables, compilation, input allocation and packed -backing storage are outside timing. Extra peak allocation includes the returned -output or gradient. These are one-machine, one-process microbenchmarks, without -clock pinning; close differences do not establish a stable winner. - -## Reproduction - -Set a toolkit matching PyTorch, the target architecture and a matching host -compiler. On the recorded Windows host, run from an x64 MSVC environment with -`CUDA_HOME`/`CUDA_PATH` selecting CUDA 12.8 and `TORCH_CUDA_ARCH_LIST=12.0`. -Use `$env:RL_KERNEL_ROPE_ONLY='1'` for the operator package profile. - -```bash -python setup.py sdist bdist_wheel -python setup.py build_ext --inplace -python -m pytest tests/ops/rope/test_multi_axis_rope.py -q -rs -python tools/validation/operators/check_operator.py --op multi_axis_rope \ - --candidate cuda --device cuda --dtype bf16 --batch 2 --seq 7 --n-heads 3 --check-grad --json -python tools/validation/operators/validate_multi_axis_rope.py --ws1 cuda --output ws1-cuda.json -python tools/validation/operators/validate_multi_axis_rope.py --ws1 triton --output ws1-triton.json -python tools/validation/operators/validate_multi_axis_rope.py --benchmark \ - --flash-source /path/to/flash_attn/ops/triton/rotary.py \ - --flash-layer-source /path/to/flash_attn/layers/rotary.py --output benchmark.json -python tools/validation/operators/check_qwenimage_rope_upstream.py \ - --source /path/to/transformer_qwenimage.py --output upstream.json -compute-sanitizer --tool memcheck --error-exitcode 1 python \ - tools/validation/operators/validate_multi_axis_rope.py -q \ - -k "forward_backward_accuracy or noncontiguous_input_tables_and_gradient or empty_forward_backward or packed_qkv_and_independent_table_strides or native_entrypoint_checks" -mkdocs build --strict -f mkdocs.yaml -``` - -Install the built wheel, change to a directory outside the checkout, then run -`check_multi_axis_rope_install.py --package-root /path/to/site-packages --output installed-wheel.json`. -This probe neither changes sys.path nor invokes the focused JIT loader. -See `quality.txt` for the regression command and static-check scope. - -## Related work and remaining qualification - -The [operator page](../../../docs/operators/multi-axis-rope.md) documents -FlashAttention reuse and the related Ascend #411, SGLang and vLLM-Omni paths. -The live repository search on 2026-10-28 found #411 for the exact row name; -its head remains `285c4dccab7e9048956584caf431946d82405c09`. Its half-split and -uncentered reference differs from the pinned adjacent-pair/centered contract. - -Linux, A100/H100, multi-GPU guards, the complete monolithic `_C` build, QK-norm -fusion, model/block assembly and trajectories remain unqualified. There is no -Dense Qwen3 or distributed certification implied by the operator WS1 adapter. -Cross-version/device byte equality and fast-math builds are not claimed. diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/baseline-layout.txt b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/baseline-layout.txt deleted file mode 100644 index 16a5c97e8..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/baseline-layout.txt +++ /dev/null @@ -1,138 +0,0 @@ -.FF.............s........F.FF.... [100%] -================================== FAILURES =================================== -________________ test_cuda_build_keeps_existing_ipc_collective ________________ - -monkeypatch = <_pytest.monkeypatch.MonkeyPatch object at 0x0000020FF72168B0> - - def test_cuda_build_keeps_existing_ipc_collective(monkeypatch) -> None: - extension = _load_extension_config(monkeypatch, hip=None) - - assert "csrc/cuda/collectives/deterministic_collective.cu" in extension["sources"] - assert "-DKERNEL_ALIGN_WITH_CUDA" in extension["extra_compile_args"]["cxx"] - assert "-DKERNEL_ALIGN_WITH_ROCM" not in extension["extra_compile_args"]["cxx"] -> assert "-lcuda" in extension["extra_link_args"] -E AssertionError: assert '-lcuda' in ['-Wl,-rpath,$ORIGIN/../torch/lib'] - -tests\build\test_build_platform_collectives.py:61: AssertionError -____________ test_package_data_json_links_resolve_to_real_content _____________ - -tmp_path = WindowsPath('C:/Users/123/AppData/Local/Temp/pytest-of-123/pytest-4/test_package_data_json_links_r0') - - def test_package_data_json_links_resolve_to_real_content(tmp_path) -> None: - target = tmp_path / "configs" / "profile.json" - target.parent.mkdir() - target.write_text('{"name": "profile"}\n', encoding="utf-8") - source = tmp_path / "rl_engine" / "profile.json" - source.parent.mkdir() -> source.symlink_to("../configs/profile.json") - -tests\build\test_build_platform_collectives.py:70: -_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ - -self = WindowsPath('C:/Users/123/AppData/Local/Temp/pytest-of-123/pytest-4/test_package_data_json_links_r0/rl_engine/profile.json') -target = '../configs/profile.json', target_is_directory = False - - def symlink_to(self, target, target_is_directory=False): - """ - Make this path a symlink pointing to the target path. - Note the order of arguments (link, target) is the reverse of os.symlink. - """ -> os.symlink(target, self, target_is_directory) -E OSError: [WinError 1314] �ͻ���û���������Ȩ��: '../configs/profile.json' -> 'C:\\Users\\123\\AppData\\Local\\Temp\\pytest-of-123\\pytest-4\\test_package_data_json_links_r0\\rl_engine\\profile.json' - -D:\anaconda\Lib\pathlib\_local.py:789: OSError -_______ test_recorded_algorithm_sources_resolve_to_real_implementations _______ - - def test_recorded_algorithm_sources_resolve_to_real_implementations(): - from rl_engine.config.workload import load_manifest - - moves = json.loads((ROOT / "tools/migration/layout.json").read_text())["moves"] - for case in load_manifest().representative_cases: - source, symbol = case["provenance_evidence"]["algorithm_source"].rsplit(":", 1) - implementation = ROOT / moves.get(source, source) -> text = implementation.read_text() - ^^^^^^^^^^^^^^^^^^^^^^^^^^ - -tests\runtime\test_layout_compatibility.py:87: -_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ -D:\anaconda\Lib\pathlib\_local.py:546: in read_text - return PathBase.read_text(self, encoding, errors, newline) - ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ -_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ - -self = WindowsPath('D:/dandan123/dandan/issue/investigation-rl-kernel/baseline-validation/csrc/cuda/attention/deterministic_attention.cu') -encoding = 'locale', errors = None, newline = None - - def read_text(self, encoding=None, errors=None, newline=None): - """ - Open the file in text mode, read it, and close the file. - """ - with self.open(mode='r', encoding=encoding, errors=errors, newline=newline) as f: -> return f.read() - ^^^^^^^^ -E UnicodeDecodeError: 'gbk' codec can't decode byte 0x92 in position 173: illegal multibyte sequence - -D:\anaconda\Lib\pathlib\_abc.py:633: UnicodeDecodeError -_ test_packaged_contract_bytes_match_main[rl_engine.config-workloads/qwen3_8b.json-80395644b5e84ed1a76729ffc6f78eb11f41b9ab42fb36e2d00f19635d3df0f6] _ - -package = 'rl_engine.config', resource = 'workloads/qwen3_8b.json' -digest = '80395644b5e84ed1a76729ffc6f78eb11f41b9ab42fb36e2d00f19635d3df0f6' - - @pytest.mark.parametrize( - ("package", "resource", "digest"), - [ - ( - "rl_engine.config", - "workloads/qwen3_8b.json", - "80395644b5e84ed1a76729ffc6f78eb11f41b9ab42fb36e2d00f19635d3df0f6", - ), - ( - "rl_engine.contracts", - "profiles/precision/ws1.json", - "044167ec21e372017a2758dc09f805b368283c06b83f64061877628cbd101532", - ), - ], - ) - def test_packaged_contract_bytes_match_main(package, resource, digest): -> assert hashlib.sha256(files(package).joinpath(resource).read_bytes()).hexdigest() == digest -E AssertionError: assert 'c97e966fb075...75e111e052d7d' == '80395644b5e8...f19635d3df0f6' -E -E - 80395644b5e84ed1a76729ffc6f78eb11f41b9ab42fb36e2d00f19635d3df0f6 -E + c97e966fb075dad8ce7dcd7a23a58ffb856e7474dfb3c50510675e111e052d7d - -tests\runtime\test_layout_compatibility.py:118: AssertionError -_ test_packaged_contract_bytes_match_main[rl_engine.contracts-profiles/precision/ws1.json-044167ec21e372017a2758dc09f805b368283c06b83f64061877628cbd101532] _ - -package = 'rl_engine.contracts', resource = 'profiles/precision/ws1.json' -digest = '044167ec21e372017a2758dc09f805b368283c06b83f64061877628cbd101532' - - @pytest.mark.parametrize( - ("package", "resource", "digest"), - [ - ( - "rl_engine.config", - "workloads/qwen3_8b.json", - "80395644b5e84ed1a76729ffc6f78eb11f41b9ab42fb36e2d00f19635d3df0f6", - ), - ( - "rl_engine.contracts", - "profiles/precision/ws1.json", - "044167ec21e372017a2758dc09f805b368283c06b83f64061877628cbd101532", - ), - ], - ) - def test_packaged_contract_bytes_match_main(package, resource, digest): -> assert hashlib.sha256(files(package).joinpath(resource).read_bytes()).hexdigest() == digest -E AssertionError: assert '2020c02e7af4...313d6752cd34e' == '044167ec21e3...7628cbd101532' -E -E - 044167ec21e372017a2758dc09f805b368283c06b83f64061877628cbd101532 -E + 2020c02e7af4e2ca4163e1ecfb9c491baef07b73deba39b0f44313d6752cd34e - -tests\runtime\test_layout_compatibility.py:118: AssertionError -=========================== short test summary info =========================== -FAILED tests/build/test_build_platform_collectives.py::test_cuda_build_keeps_existing_ipc_collective -FAILED tests/build/test_build_platform_collectives.py::test_package_data_json_links_resolve_to_real_content -FAILED tests/runtime/test_layout_compatibility.py::test_recorded_algorithm_sources_resolve_to_real_implementations -FAILED tests/runtime/test_layout_compatibility.py::test_packaged_contract_bytes_match_main[rl_engine.config-workloads/qwen3_8b.json-80395644b5e84ed1a76729ffc6f78eb11f41b9ab42fb36e2d00f19635d3df0f6] -FAILED tests/runtime/test_layout_compatibility.py::test_packaged_contract_bytes_match_main[rl_engine.contracts-profiles/precision/ws1.json-044167ec21e372017a2758dc09f805b368283c06b83f64061877628cbd101532] -5 failed, 27 passed, 1 skipped in 5.70s diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/benchmark-bfloat16.json b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/benchmark-bfloat16.json deleted file mode 100644 index 72fb42bbe..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/benchmark-bfloat16.json +++ /dev/null @@ -1,2781 +0,0 @@ -{ - "gpu": "NVIDIA GeForce RTX 5080", - "capability": [ - 12, - 0 - ], - "os": "Windows-11-10.0.26200-SP0", - "python": "3.13.9 | packaged by Anaconda, Inc. | (main, Oct 21 2025, 19:09:58) [MSC v.1929 64 bit (AMD64)]", - "torch": "2.9.1+cu128", - "cuda_runtime": "12.8", - "triton": "3.5.1", - "dtype": "bfloat16", - "base_commit": "16b81175874ecf11d5ff05c00de3b69e9457591b", - "source_sha256": { - "csrc/cuda/rope/multi_axis_rope.cu": "f5db5cf5d4962053194455d9cde8cf91cf4e824bad7abf72e8410a404bcb0e44", - "rl_engine/backends/shared/triton/rope/_flash_attn_rotary.py": "5c6984da351e5228f7549629d4ff53e5772b9a6ef2a01ba2a38493e769b15ab5", - "rl_engine/backends/shared/triton/rope/multi_axis_rope.py": "5b989bd63efd77913267b48654d10204bc563768c6c5d4e1cde9fb55e1354547", - "rl_engine/backends/cuda/rope/multi_axis_rope.py": "d779d9b60ce5d1d6651925f723049930d898024fdf0548ce29a2019b84753d65", - "rl_engine/ops/rope/_registration.py": "add0375d1639444908091d01e8d29781f7a975beffec64b77bfaa81a5b9a0ae3", - "rl_engine/ops/rope/multi_axis_rope.py": "07bcff1160f22c1ff6730be731206b2f4a153e2fada26afc93fc884e780568cb", - "rl_engine/reference/rope/multi_axis_rope.py": "c9a94730047f51659f7aa0a443e88d6ce7d03d00d6813a9aed548cafbf6347bf", - "benchmarks/operators/rope/benchmark_multi_axis_rope.py": "6961a55803fbd4e566bd060558fb59e658e7f1f704e61ca97f912e5d16b55884" - }, - "timing": "CUDA graph ms/call and synchronized eager wall ms/call; warm allocation; excludes tables, compilation, and initial allocation of inputs", - "memory": "extra peak allocated bytes excluding inputs/tables/grad; includes returned output or gradient; packed QKV backing storage allocated before timing", - "iterations": 30, - "warmup": 5, - "trials": 3, - "graph_calls": 20, - "flash_baseline": { - "included": true, - "revision": "98c1cee0ed59307833841cdbf567fe336ceec13a", - "kernel_sha256_normalized": "d2037e623854c702fb075784fe4292d4a89d6b9ac1ae33cab3d6788f8f94f240", - "layer_sha256_normalized": "a79cb42dfc33e126227bdba27ae7a024d7d5cc5c8fd7f7a5f26c68d501024e0a", - "mode": "upstream rotary_kernel, apply_rotary and ApplyRotaryEmb, interleaved=True", - "no_fma_change": "only add enable_fp_fusion=False at the original kernel launch" - }, - "trace": { - "cuda": { - "contract": "qwenimage-interleaved-rope-fp32-v1", - "actual_backend": "rlkernel.cuda.qwenimage_multi_axis_rope.v1", - "kernel_id": "torch.ops.rl_kernel_native.multi_axis_rope", - "layout": "BSHD, adjacent channel pairs", - "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", - "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" - }, - "triton": { - "contract": "qwenimage-interleaved-rope-fp32-v1", - "actual_backend": "rlkernel.triton.qwenimage_multi_axis_rope.v1", - "kernel_id": "rl_engine.backends.shared.triton.rope._flash_attn_rotary.rotary_kernel", - "layout": "BSHD, adjacent channel pairs", - "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", - "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" - } - }, - "qualification": "Exact bytes vs unfused FP32 gold for CUDA, Triton and flash_no_fma; declared tolerance for upstream complex and default FlashAttention", - "correctness": [ - { - "backend": "upstream_complex", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "cuda", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "triton", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "flash_default", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "flash_no_fma", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "upstream_complex", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "cuda", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "triton", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "flash_default", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "flash_no_fma", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "upstream_complex", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "different_bytes_y_dx": [ - 130, - 139 - ] - }, - { - "backend": "cuda", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "triton", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "flash_default", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "different_bytes_y_dx": [ - 130, - 139 - ] - }, - { - "backend": "flash_no_fma", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "upstream_complex", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "different_bytes_y_dx": [ - 139, - 140 - ] - }, - { - "backend": "cuda", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "triton", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "flash_default", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "different_bytes_y_dx": [ - 139, - 140 - ] - }, - { - "backend": "flash_no_fma", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "upstream_complex", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "different_bytes_y_dx": [ - 234, - 225 - ] - }, - { - "backend": "cuda", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "triton", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "flash_default", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "different_bytes_y_dx": [ - 234, - 225 - ] - }, - { - "backend": "flash_no_fma", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "upstream_complex", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "different_bytes_y_dx": [ - 220, - 226 - ] - }, - { - "backend": "cuda", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "triton", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "flash_default", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "different_bytes_y_dx": [ - 220, - 226 - ] - }, - { - "backend": "flash_no_fma", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "upstream_complex", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "different_bytes_y_dx": [ - 196, - 203 - ] - }, - { - "backend": "cuda", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "triton", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "flash_default", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "different_bytes_y_dx": [ - 196, - 203 - ] - }, - { - "backend": "flash_no_fma", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "upstream_complex", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "different_bytes_y_dx": [ - 230, - 218 - ] - }, - { - "backend": "cuda", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "triton", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "flash_default", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "different_bytes_y_dx": [ - 230, - 218 - ] - }, - { - "backend": "flash_no_fma", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - } - ], - "results": [ - { - "backend": "upstream_complex", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 18432, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.003558399900794029, - "graph_min_ms": 0.0035264000296592714, - "graph_max_ms": 0.015438400208950043, - "eager_wall_median_ms": 0.0339666667059646, - "eager_wall_trials_ms": [ - 0.0356399999873247, - 0.0339666667059646, - 0.03332666662269427 - ], - "extra_peak_bytes": 184320 - }, - { - "backend": "upstream_complex", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 18432, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.007688799872994423, - "graph_min_ms": 0.007651200145483017, - "graph_max_ms": 0.007790400087833405, - "eager_wall_median_ms": 0.21265999997316007, - "eager_wall_trials_ms": [ - 0.21187666667780528, - 0.21265999997316007, - 0.22587333332921844 - ], - "extra_peak_bytes": 187392 - }, - { - "backend": "cuda", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 18432, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.0012016000226140022, - "graph_min_ms": 0.0011567999608814716, - "graph_max_ms": 0.0012384000234305858, - "eager_wall_median_ms": 0.02289333333465038, - "eager_wall_trials_ms": [ - 0.02289333333465038, - 0.02245666667780218, - 0.02458333328831941 - ], - "extra_peak_bytes": 36864 - }, - { - "backend": "cuda", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 18432, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.0022280000150203703, - "graph_min_ms": 0.0021888000890612604, - "graph_max_ms": 0.0022655999287962913, - "eager_wall_median_ms": 0.1764300000407578, - "eager_wall_trials_ms": [ - 0.20031666669334905, - 0.1764300000407578, - 0.15246333338533682 - ], - "extra_peak_bytes": 73728 - }, - { - "backend": "triton", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 18432, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.001142400037497282, - "graph_min_ms": 0.0010591999627649785, - "graph_max_ms": 0.005691200122237205, - "eager_wall_median_ms": 0.057766666698929235, - "eager_wall_trials_ms": [ - 0.07126333330234047, - 0.057766666698929235, - 0.05640333335274287 - ], - "extra_peak_bytes": 36864 - }, - { - "backend": "triton", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 18432, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.002125599980354309, - "graph_min_ms": 0.002091200090944767, - "graph_max_ms": 0.002160000056028366, - "eager_wall_median_ms": 0.21872666663208898, - "eager_wall_trials_ms": [ - 0.22702333335473668, - 0.21872666663208898, - 0.21611666670651175 - ], - "extra_peak_bytes": 73728 - }, - { - "backend": "flash_default", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 18432, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.0010935999918729067, - "graph_min_ms": 0.001033599954098463, - "graph_max_ms": 0.0011536000296473504, - "eager_wall_median_ms": 0.10024000002886169, - "eager_wall_trials_ms": [ - 0.10126666669142044, - 0.09858333329854456, - 0.10024000002886169 - ], - "extra_peak_bytes": 36864 - }, - { - "backend": "flash_default", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 18432, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.0020527999848127367, - "graph_min_ms": 0.002019199915230274, - "graph_max_ms": 0.0021023999899625776, - "eager_wall_median_ms": 0.3111766666431019, - "eager_wall_trials_ms": [ - 0.30222666664485587, - 0.3111766666431019, - 0.333893333360417 - ], - "extra_peak_bytes": 73728 - }, - { - "backend": "flash_no_fma", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 18432, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.0010991999879479408, - "graph_min_ms": 0.001033599954098463, - "graph_max_ms": 0.0011392000131309032, - "eager_wall_median_ms": 0.10421333336125826, - "eager_wall_trials_ms": [ - 0.10421333336125826, - 0.10106999998242827, - 0.12048333337588701 - ], - "extra_peak_bytes": 36864 - }, - { - "backend": "flash_no_fma", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 18432, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.0020503999665379526, - "graph_min_ms": 0.002001599967479706, - "graph_max_ms": 0.00886240005493164, - "eager_wall_median_ms": 0.31822000000829576, - "eager_wall_trials_ms": [ - 0.28550666665978497, - 0.31822000000829576, - 0.33777000001767493 - ], - "extra_peak_bytes": 73728 - }, - { - "backend": "upstream_complex", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 55296, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.00379519984126091, - "graph_min_ms": 0.003745599836111069, - "graph_max_ms": 0.003859199956059456, - "eager_wall_median_ms": 0.10653666668076767, - "eager_wall_trials_ms": [ - 0.10653666668076767, - 0.10879000001295935, - 0.09587666669782872 - ], - "extra_peak_bytes": 184320 - }, - { - "backend": "upstream_complex", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 55296, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.008064000308513642, - "graph_min_ms": 0.008011200278997422, - "graph_max_ms": 0.008102399855852127, - "eager_wall_median_ms": 0.2260500000071867, - "eager_wall_trials_ms": [ - 0.23646333332483968, - 0.2187999999781217, - 0.2260500000071867 - ], - "extra_peak_bytes": 187392 - }, - { - "backend": "cuda", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 55296, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.0012024000287055969, - "graph_min_ms": 0.0011392000131309032, - "graph_max_ms": 0.0025040000677108766, - "eager_wall_median_ms": 0.023706666615908034, - "eager_wall_trials_ms": [ - 0.023706666615908034, - 0.024339999981748406, - 0.023019999995691858 - ], - "extra_peak_bytes": 36864 - }, - { - "backend": "cuda", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 55296, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.002226400002837181, - "graph_min_ms": 0.002174399979412556, - "graph_max_ms": 0.0022671999409794807, - "eager_wall_median_ms": 0.157756666643157, - "eager_wall_trials_ms": [ - 0.165013333329019, - 0.15389999995628992, - 0.157756666643157 - ], - "extra_peak_bytes": 73728 - }, - { - "backend": "triton", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 55296, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.0011008000001311302, - "graph_min_ms": 0.0010688000358641148, - "graph_max_ms": 0.0011359999887645244, - "eager_wall_median_ms": 0.059236666675133165, - "eager_wall_trials_ms": [ - 0.053669999942940194, - 0.06250666665437166, - 0.059236666675133165 - ], - "extra_peak_bytes": 36864 - }, - { - "backend": "triton", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 55296, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.0021215999498963357, - "graph_min_ms": 0.002065600082278252, - "graph_max_ms": 0.0021679999306797983, - "eager_wall_median_ms": 0.25444666662224336, - "eager_wall_trials_ms": [ - 0.25444666662224336, - 0.27043333332888625, - 0.2240899999984928 - ], - "extra_peak_bytes": 73728 - }, - { - "backend": "flash_default", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 55296, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.001101600006222725, - "graph_min_ms": 0.001065600011497736, - "graph_max_ms": 0.0011376000009477138, - "eager_wall_median_ms": 0.12566333334689261, - "eager_wall_trials_ms": [ - 0.12566333334689261, - 0.12918333328949907, - 0.10983666664590903 - ], - "extra_peak_bytes": 36864 - }, - { - "backend": "flash_default", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 55296, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.0020527999848127367, - "graph_min_ms": 0.0020160000771284104, - "graph_max_ms": 0.0021023999899625776, - "eager_wall_median_ms": 0.28767333333235, - "eager_wall_trials_ms": [ - 0.2609133333559536, - 0.28767333333235, - 0.3289566666353494 - ], - "extra_peak_bytes": 73728 - }, - { - "backend": "flash_no_fma", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 55296, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.001101600006222725, - "graph_min_ms": 0.0010223999619483949, - "graph_max_ms": 0.0022975999861955644, - "eager_wall_median_ms": 0.0995366666757036, - "eager_wall_trials_ms": [ - 0.10132333336514421, - 0.0995366666757036, - 0.0990000000456348 - ], - "extra_peak_bytes": 36864 - }, - { - "backend": "flash_no_fma", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 55296, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.0020527999848127367, - "graph_min_ms": 0.0020128000527620316, - "graph_max_ms": 0.0021007999777793883, - "eager_wall_median_ms": 0.3198099999887442, - "eager_wall_trials_ms": [ - 0.33375666668386356, - 0.30899333335886087, - 0.3198099999887442 - ], - "extra_peak_bytes": 73728 - }, - { - "backend": "upstream_complex", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 12582912, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.2427680015563965, - "graph_min_ms": 0.24123198986053468, - "graph_max_ms": 0.3049216032028198, - "eager_wall_median_ms": 0.25121666670505266, - "eager_wall_trials_ms": [ - 0.24960999999166233, - 0.2669566666857766, - 0.25121666670505266 - ], - "extra_peak_bytes": 125829120 - }, - { - "backend": "upstream_complex", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 12582912, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.5011807918548584, - "graph_min_ms": 0.4943295955657959, - "graph_max_ms": 0.5467391967773437, - "eager_wall_median_ms": 0.5078799999561548, - "eager_wall_trials_ms": [ - 0.5254200000005463, - 0.5078799999561548, - 0.5070699999729792 - ], - "extra_peak_bytes": 127926272 - }, - { - "backend": "cuda", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 12582912, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.05770559906959534, - "graph_min_ms": 0.05762240290641785, - "graph_max_ms": 0.07225120067596436, - "eager_wall_median_ms": 0.061780000032740645, - "eager_wall_trials_ms": [ - 0.06511333334249987, - 0.06060333331940152, - 0.061780000032740645 - ], - "extra_peak_bytes": 25165824 - }, - { - "backend": "cuda", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 12582912, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.12953280210494994, - "graph_min_ms": 0.12918239831924438, - "graph_max_ms": 0.14451359510421752, - "eager_wall_median_ms": 0.15664999997776854, - "eager_wall_trials_ms": [ - 0.2333100000517637, - 0.15664999997776854, - 0.14973666669296412 - ], - "extra_peak_bytes": 50331648 - }, - { - "backend": "triton", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 12582912, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.05190559923648834, - "graph_min_ms": 0.05078719854354859, - "graph_max_ms": 0.06408159732818604, - "eager_wall_median_ms": 0.0845666666161075, - "eager_wall_trials_ms": [ - 0.08885000000494377, - 0.08201333330362104, - 0.0845666666161075 - ], - "extra_peak_bytes": 25165824 - }, - { - "backend": "triton", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 12582912, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.12463040351867677, - "graph_min_ms": 0.12331520318984986, - "graph_max_ms": 0.1529088020324707, - "eager_wall_median_ms": 0.2539966667124342, - "eager_wall_trials_ms": [ - 0.2757299999454214, - 0.25390333333537757, - 0.2539966667124342 - ], - "extra_peak_bytes": 50331648 - }, - { - "backend": "flash_default", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 12582912, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.05208719968795776, - "graph_min_ms": 0.05123199820518494, - "graph_max_ms": 0.06861280202865601, - "eager_wall_median_ms": 0.13019999999717888, - "eager_wall_trials_ms": [ - 0.1370399999965836, - 0.12465666665472479, - 0.13019999999717888 - ], - "extra_peak_bytes": 25165824 - }, - { - "backend": "flash_default", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 12582912, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.12441120147705079, - "graph_min_ms": 0.12314239740371705, - "graph_max_ms": 0.139137601852417, - "eager_wall_median_ms": 0.340176666698729, - "eager_wall_trials_ms": [ - 0.35652999995363643, - 0.340176666698729, - 0.33921999996285496 - ], - "extra_peak_bytes": 50331648 - }, - { - "backend": "flash_no_fma", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 12582912, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.052018401026725766, - "graph_min_ms": 0.05074719786643982, - "graph_max_ms": 0.07086880207061767, - "eager_wall_median_ms": 0.10117999997116083, - "eager_wall_trials_ms": [ - 0.10250666667464732, - 0.10117999997116083, - 0.10012000002461718 - ], - "extra_peak_bytes": 25165824 - }, - { - "backend": "flash_no_fma", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 12582912, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.12442959547042846, - "graph_min_ms": 0.12299360036849975, - "graph_max_ms": 0.14664160013198851, - "eager_wall_median_ms": 0.34945666666317265, - "eager_wall_trials_ms": [ - 0.36080999998375773, - 0.3163533333160255, - 0.34945666666317265 - ], - "extra_peak_bytes": 50331648 - }, - { - "backend": "upstream_complex", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 37748736, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.2459879994392395, - "graph_min_ms": 0.24467840194702148, - "graph_max_ms": 0.272545599937439, - "eager_wall_median_ms": 0.2536533333720096, - "eager_wall_trials_ms": [ - 0.27815666665749933, - 0.2529566666756485, - 0.2536533333720096 - ], - "extra_peak_bytes": 125829120 - }, - { - "backend": "upstream_complex", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 37748736, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.5008344173431396, - "graph_min_ms": 0.4996367931365967, - "graph_max_ms": 0.5373487949371338, - "eager_wall_median_ms": 0.5262199999682101, - "eager_wall_trials_ms": [ - 0.5269766666970099, - 0.5070566666593853, - 0.5262199999682101 - ], - "extra_peak_bytes": 127926272 - }, - { - "backend": "cuda", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 37748736, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.05865840017795562, - "graph_min_ms": 0.058550399541854856, - "graph_max_ms": 0.08192480206489564, - "eager_wall_median_ms": 0.06282000000889336, - "eager_wall_trials_ms": [ - 0.06282000000889336, - 0.06640333334265354, - 0.06277333335068154 - ], - "extra_peak_bytes": 25165824 - }, - { - "backend": "cuda", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 37748736, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.1308944046497345, - "graph_min_ms": 0.13030719757080078, - "graph_max_ms": 0.145305597782135, - "eager_wall_median_ms": 0.1688300000144712, - "eager_wall_trials_ms": [ - 0.1688300000144712, - 0.15922666667999388, - 0.17570666665657578 - ], - "extra_peak_bytes": 50331648 - }, - { - "backend": "triton", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 37748736, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.04834639877080918, - "graph_min_ms": 0.047383999824523924, - "graph_max_ms": 0.06508319973945617, - "eager_wall_median_ms": 0.08053666667062014, - "eager_wall_trials_ms": [ - 0.09511666667094687, - 0.07707999999790142, - 0.08053666667062014 - ], - "extra_peak_bytes": 25165824 - }, - { - "backend": "triton", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 37748736, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.1231384038925171, - "graph_min_ms": 0.12216000556945801, - "graph_max_ms": 0.1478927969932556, - "eager_wall_median_ms": 0.2741499999804849, - "eager_wall_trials_ms": [ - 0.28788666665301815, - 0.2741499999804849, - 0.24507333334137607 - ], - "extra_peak_bytes": 50331648 - }, - { - "backend": "flash_default", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 37748736, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.048105600476264956, - "graph_min_ms": 0.0474047988653183, - "graph_max_ms": 0.06945440173149109, - "eager_wall_median_ms": 0.12940666668631215, - "eager_wall_trials_ms": [ - 0.14179000002817096, - 0.12940666668631215, - 0.12683000004471978 - ], - "extra_peak_bytes": 25165824 - }, - { - "backend": "flash_default", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 37748736, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.12347760200500488, - "graph_min_ms": 0.12205439805984497, - "graph_max_ms": 0.14951200485229493, - "eager_wall_median_ms": 0.36926333335334977, - "eager_wall_trials_ms": [ - 0.36926333335334977, - 0.3322200000184239, - 0.36990666667406913 - ], - "extra_peak_bytes": 50331648 - }, - { - "backend": "flash_no_fma", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 37748736, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.04896320104598999, - "graph_min_ms": 0.047470399737358095, - "graph_max_ms": 0.06636160016059875, - "eager_wall_median_ms": 0.14505666664869446, - "eager_wall_trials_ms": [ - 0.14700000004571243, - 0.14505666664869446, - 0.13565333332129134 - ], - "extra_peak_bytes": 25165824 - }, - { - "backend": "flash_no_fma", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 37748736, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.12373999953269958, - "graph_min_ms": 0.1219599962234497, - "graph_max_ms": 0.1542415976524353, - "eager_wall_median_ms": 0.33643999995547347, - "eager_wall_trials_ms": [ - 0.33643999995547347, - 0.3219899999749032, - 0.3583933333478247 - ], - "extra_peak_bytes": 50331648 - }, - { - "backend": "upstream_complex", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 21163008, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.51009202003479, - "graph_min_ms": 0.5082543849945068, - "graph_max_ms": 0.5431327819824219, - "eager_wall_median_ms": 0.526480000007723, - "eager_wall_trials_ms": [ - 0.5092500000197712, - 0.5466633333223095, - 0.526480000007723 - ], - "extra_peak_bytes": 211630080 - }, - { - "backend": "upstream_complex", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 21163008, - 3072, - 128, - 1 - ], - "graph_median_ms": 1.0505383968353272, - "graph_min_ms": 1.0236895561218262, - "graph_max_ms": 1.129139232635498, - "eager_wall_median_ms": 1.044759999967937, - "eager_wall_trials_ms": [ - 1.044759999967937, - 1.0434166666527744, - 1.0474333333453008 - ], - "extra_peak_bytes": 215157248 - }, - { - "backend": "cuda", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 21163008, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.11203439831733704, - "graph_min_ms": 0.1114351987838745, - "graph_max_ms": 0.14540959596633912, - "eager_wall_median_ms": 0.11692666666931473, - "eager_wall_trials_ms": [ - 0.11580999998841435, - 0.11692666666931473, - 0.12043333335896023 - ], - "extra_peak_bytes": 42326016 - }, - { - "backend": "cuda", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 21163008, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.22477760314941406, - "graph_min_ms": 0.22440481185913086, - "graph_max_ms": 0.24935200214385986, - "eager_wall_median_ms": 0.23845333331943644, - "eager_wall_trials_ms": [ - 0.23845333331943644, - 0.2526333333056148, - 0.230336666633472 - ], - "extra_peak_bytes": 84652032 - }, - { - "backend": "triton", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 21163008, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.10685279965400696, - "graph_min_ms": 0.10539519786834717, - "graph_max_ms": 0.12008960247039795, - "eager_wall_median_ms": 0.11027333330275724, - "eager_wall_trials_ms": [ - 0.11027333330275724, - 0.10863333330538201, - 0.11599999997997656 - ], - "extra_peak_bytes": 42326016 - }, - { - "backend": "triton", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 21163008, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.21339280605316163, - "graph_min_ms": 0.2126607894897461, - "graph_max_ms": 0.23967199325561522, - "eager_wall_median_ms": 0.27515000001585577, - "eager_wall_trials_ms": [ - 0.2870199999961187, - 0.2669133333256468, - 0.27515000001585577 - ], - "extra_peak_bytes": 84652032 - }, - { - "backend": "flash_default", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 21163008, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.10692960023880005, - "graph_min_ms": 0.10526080131530761, - "graph_max_ms": 0.12183040380477905, - "eager_wall_median_ms": 0.10955999999472017, - "eager_wall_trials_ms": [ - 0.1204166666866513, - 0.10955999999472017, - 0.10893666670502473 - ], - "extra_peak_bytes": 42326016 - }, - { - "backend": "flash_default", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 21163008, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.21346559524536132, - "graph_min_ms": 0.21225121021270751, - "graph_max_ms": 0.2430032014846802, - "eager_wall_median_ms": 0.338726666692916, - "eager_wall_trials_ms": [ - 0.33888666663794237, - 0.3213533333716138, - 0.338726666692916 - ], - "extra_peak_bytes": 84652032 - }, - { - "backend": "flash_no_fma", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 21163008, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.10693199634552003, - "graph_min_ms": 0.10515999794006348, - "graph_max_ms": 0.12142080068588257, - "eager_wall_median_ms": 0.1100600000427221, - "eager_wall_trials_ms": [ - 0.11902333335456206, - 0.1100600000427221, - 0.10893666664439176 - ], - "extra_peak_bytes": 42326016 - }, - { - "backend": "flash_no_fma", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 21163008, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.21350640058517456, - "graph_min_ms": 0.2123823881149292, - "graph_max_ms": 0.23954401016235352, - "eager_wall_median_ms": 0.3223733333167426, - "eager_wall_trials_ms": [ - 0.29901333333934116, - 0.3223733333167426, - 0.326143333343983 - ], - "extra_peak_bytes": 84652032 - }, - { - "backend": "upstream_complex", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 63489024, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.5111671924591064, - "graph_min_ms": 0.5095888137817383, - "graph_max_ms": 0.5389599800109863, - "eager_wall_median_ms": 0.5120633333717706, - "eager_wall_trials_ms": [ - 0.5099033333560025, - 0.532273333374178, - 0.5120633333717706 - ], - "extra_peak_bytes": 211630080 - }, - { - "backend": "upstream_complex", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 63489024, - 9216, - 128, - 1 - ], - "graph_median_ms": 1.0546072006225586, - "graph_min_ms": 1.0272704124450684, - "graph_max_ms": 1.1091728210449219, - "eager_wall_median_ms": 1.0453999999905743, - "eager_wall_trials_ms": [ - 1.0565066666458733, - 1.0453999999905743, - 1.042836666662576 - ], - "extra_peak_bytes": 215157248 - }, - { - "backend": "cuda", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 63489024, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.11280480027198792, - "graph_min_ms": 0.11215200424194335, - "graph_max_ms": 0.13529759645462036, - "eager_wall_median_ms": 0.12308333334658528, - "eager_wall_trials_ms": [ - 0.11563333331044608, - 0.12308333334658528, - 0.12365000002318993 - ], - "extra_peak_bytes": 42326016 - }, - { - "backend": "cuda", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 63489024, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.2251416087150574, - "graph_min_ms": 0.22450559139251708, - "graph_max_ms": 0.2644239902496338, - "eager_wall_median_ms": 0.2328933333653064, - "eager_wall_trials_ms": [ - 0.23086333330866182, - 0.2328933333653064, - 0.25025666667109664 - ], - "extra_peak_bytes": 84652032 - }, - { - "backend": "triton", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 63489024, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.1075808048248291, - "graph_min_ms": 0.1061360001564026, - "graph_max_ms": 0.12249759435653687, - "eager_wall_median_ms": 0.10984000000462402, - "eager_wall_trials_ms": [ - 0.1298533332980393, - 0.10860666661756113, - 0.10984000000462402 - ], - "extra_peak_bytes": 42326016 - }, - { - "backend": "triton", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 63489024, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.21414560079574585, - "graph_min_ms": 0.21337599754333497, - "graph_max_ms": 0.23953759670257568, - "eager_wall_median_ms": 0.26501666670810664, - "eager_wall_trials_ms": [ - 0.2692166666747653, - 0.23041333333821967, - 0.26501666670810664 - ], - "extra_peak_bytes": 84652032 - }, - { - "backend": "flash_default", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 63489024, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.10766559839248657, - "graph_min_ms": 0.10606240034103394, - "graph_max_ms": 0.13348640203475953, - "eager_wall_median_ms": 0.11105333330003002, - "eager_wall_trials_ms": [ - 0.11928999999023897, - 0.11105333330003002, - 0.11004666666849516 - ], - "extra_peak_bytes": 42326016 - }, - { - "backend": "flash_default", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 63489024, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.21448720693588258, - "graph_min_ms": 0.21337599754333497, - "graph_max_ms": 0.25474560260772705, - "eager_wall_median_ms": 0.30118666666870314, - "eager_wall_trials_ms": [ - 0.32142000000021653, - 0.30118666666870314, - 0.2931233333583805 - ], - "extra_peak_bytes": 84652032 - }, - { - "backend": "flash_no_fma", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 63489024, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.1074567973613739, - "graph_min_ms": 0.10611200332641602, - "graph_max_ms": 0.1223423957824707, - "eager_wall_median_ms": 0.11078333336627111, - "eager_wall_trials_ms": [ - 0.12015666664713838, - 0.10989999997642978, - 0.11078333336627111 - ], - "extra_peak_bytes": 42326016 - }, - { - "backend": "flash_no_fma", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 63489024, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.21430879831314087, - "graph_min_ms": 0.2134079933166504, - "graph_max_ms": 0.23155200481414795, - "eager_wall_median_ms": 0.3170500000123866, - "eager_wall_trials_ms": [ - 0.3170500000123866, - 0.32710333331730607, - 0.31006333332091646 - ], - "extra_peak_bytes": 84652032 - }, - { - "backend": "upstream_complex", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 18530304, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.43411200046539306, - "graph_min_ms": 0.4321375846862793, - "graph_max_ms": 0.46126718521118165, - "eager_wall_median_ms": 0.4498766666680846, - "eager_wall_trials_ms": [ - 0.43220333336648764, - 0.4498766666680846, - 0.46697666663627996 - ], - "extra_peak_bytes": 185303040 - }, - { - "backend": "upstream_complex", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 18530304, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.8935583591461181, - "graph_min_ms": 0.870751953125, - "graph_max_ms": 0.9448016166687012, - "eager_wall_median_ms": 0.8892600000156866, - "eager_wall_trials_ms": [ - 0.8870666666553007, - 0.8892600000156866, - 0.8902699999453034 - ], - "extra_peak_bytes": 188391424 - }, - { - "backend": "cuda", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 18530304, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.09808879792690278, - "graph_min_ms": 0.09759039878845215, - "graph_max_ms": 0.11292320489883423, - "eager_wall_median_ms": 0.11404999998679462, - "eager_wall_trials_ms": [ - 0.11404999998679462, - 0.11419666667886001, - 0.102829999984048 - ], - "extra_peak_bytes": 37060608 - }, - { - "backend": "cuda", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 18530304, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.19646480083465576, - "graph_min_ms": 0.19599839448928832, - "graph_max_ms": 0.21091840267181397, - "eager_wall_median_ms": 0.20405333331533862, - "eager_wall_trials_ms": [ - 0.20405333331533862, - 0.20323666661473302, - 0.2081700000417186 - ], - "extra_peak_bytes": 74809344 - }, - { - "backend": "triton", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 18530304, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.09356399774551391, - "graph_min_ms": 0.09257439970970154, - "graph_max_ms": 0.10800800323486329, - "eager_wall_median_ms": 0.10138999999374694, - "eager_wall_trials_ms": [ - 0.10716000003109609, - 0.10138999999374694, - 0.10003000000627556 - ], - "extra_peak_bytes": 37060608 - }, - { - "backend": "triton", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 18530304, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.18708080053329468, - "graph_min_ms": 0.1862496018409729, - "graph_max_ms": 0.21909279823303224, - "eager_wall_median_ms": 0.26604333331003244, - "eager_wall_trials_ms": [ - 0.30756000002535683, - 0.2481433333135404, - 0.26604333331003244 - ], - "extra_peak_bytes": 74809344 - }, - { - "backend": "flash_default", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 18530304, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.09362079799175263, - "graph_min_ms": 0.09246399998664856, - "graph_max_ms": 0.10841599702835084, - "eager_wall_median_ms": 0.10133000002194119, - "eager_wall_trials_ms": [ - 0.14010333331195093, - 0.10133000002194119, - 0.10039666667580605 - ], - "extra_peak_bytes": 37060608 - }, - { - "backend": "flash_default", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 18530304, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.18697920441627502, - "graph_min_ms": 0.18585760593414308, - "graph_max_ms": 0.21307199001312255, - "eager_wall_median_ms": 0.3559766667118917, - "eager_wall_trials_ms": [ - 0.3559766667118917, - 0.3666866666511244, - 0.35564333332634607 - ], - "extra_peak_bytes": 74809344 - }, - { - "backend": "flash_no_fma", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 18530304, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.09362399876117705, - "graph_min_ms": 0.09233919978141784, - "graph_max_ms": 0.10795680284500123, - "eager_wall_median_ms": 0.1356800000091122, - "eager_wall_trials_ms": [ - 0.14346666666824603, - 0.1356800000091122, - 0.1269966666465431 - ], - "extra_peak_bytes": 37060608 - }, - { - "backend": "flash_no_fma", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 18530304, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.18709120154380798, - "graph_min_ms": 0.18642239570617675, - "graph_max_ms": 0.22162079811096191, - "eager_wall_median_ms": 0.3122899999652873, - "eager_wall_trials_ms": [ - 0.29407333331619157, - 0.3122899999652873, - 0.36742666664698237 - ], - "extra_peak_bytes": 74809344 - }, - { - "backend": "upstream_complex", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 55590912, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.43668241500854493, - "graph_min_ms": 0.43459200859069824, - "graph_max_ms": 0.4815072059631348, - "eager_wall_median_ms": 0.4541133333380761, - "eager_wall_trials_ms": [ - 0.4636933333434475, - 0.43558999999125564, - 0.4541133333380761 - ], - "extra_peak_bytes": 185303040 - }, - { - "backend": "upstream_complex", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 55590912, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.8956063747406007, - "graph_min_ms": 0.8733648300170899, - "graph_max_ms": 0.9540863990783691, - "eager_wall_median_ms": 0.8916500000244317, - "eager_wall_trials_ms": [ - 0.8923599999737538, - 0.8916500000244317, - 0.8882866666681366 - ], - "extra_peak_bytes": 188391424 - }, - { - "backend": "cuda", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 55590912, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.09595520198345184, - "graph_min_ms": 0.09553279876708984, - "graph_max_ms": 0.11051679849624634, - "eager_wall_median_ms": 0.1001133333678202, - "eager_wall_trials_ms": [ - 0.1001133333678202, - 0.09971333335367187, - 0.10664000001270324 - ], - "extra_peak_bytes": 37060608 - }, - { - "backend": "cuda", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 55590912, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.19663520455360411, - "graph_min_ms": 0.19614880084991454, - "graph_max_ms": 0.22351360321044922, - "eager_wall_median_ms": 0.20444999997077198, - "eager_wall_trials_ms": [ - 0.22424333334735516, - 0.20414666669239523, - 0.20444999997077198 - ], - "extra_peak_bytes": 74809344 - }, - { - "backend": "triton", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 55590912, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.0934136003255844, - "graph_min_ms": 0.0916208028793335, - "graph_max_ms": 0.10537919998168946, - "eager_wall_median_ms": 0.09456333336856915, - "eager_wall_trials_ms": [ - 0.09251333334153362, - 0.09570666667665743, - 0.09456333336856915 - ], - "extra_peak_bytes": 37060608 - }, - { - "backend": "triton", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 55590912, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.1877624034881592, - "graph_min_ms": 0.18668639659881592, - "graph_max_ms": 0.21486239433288573, - "eager_wall_median_ms": 0.24588000002646973, - "eager_wall_trials_ms": [ - 0.26271666671770316, - 0.24588000002646973, - 0.23198333334827717 - ], - "extra_peak_bytes": 74809344 - }, - { - "backend": "flash_default", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 55590912, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.09341759979724884, - "graph_min_ms": 0.0921504020690918, - "graph_max_ms": 0.11279360055923462, - "eager_wall_median_ms": 0.14909333337224476, - "eager_wall_trials_ms": [ - 0.15717333335487638, - 0.12570666664638944, - 0.14909333337224476 - ], - "extra_peak_bytes": 37060608 - }, - { - "backend": "flash_default", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 55590912, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.1880079984664917, - "graph_min_ms": 0.18690240383148193, - "graph_max_ms": 0.2143280029296875, - "eager_wall_median_ms": 0.3311766666835562, - "eager_wall_trials_ms": [ - 0.3571433333490859, - 0.3311766666835562, - 0.32030999997611315 - ], - "extra_peak_bytes": 74809344 - }, - { - "backend": "flash_no_fma", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 55590912, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.0934448003768921, - "graph_min_ms": 0.09198240041732789, - "graph_max_ms": 0.10836479663848878, - "eager_wall_median_ms": 0.132580000051045, - "eager_wall_trials_ms": [ - 0.132580000051045, - 0.14571666670235572, - 0.12959000002107737 - ], - "extra_peak_bytes": 37060608 - }, - { - "backend": "flash_no_fma", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 55590912, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.18786399960517886, - "graph_min_ms": 0.18675040006637572, - "graph_max_ms": 0.2141263961791992, - "eager_wall_median_ms": 0.34321999998307245, - "eager_wall_trials_ms": [ - 0.36083333331286366, - 0.3284899999925983, - 0.34321999998307245 - ], - "extra_peak_bytes": 74809344 - } - ] -} diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/installed-wheel.json b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/installed-wheel.json deleted file mode 100644 index 468ce07af..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/installed-wheel.json +++ /dev/null @@ -1,37 +0,0 @@ -{ - "package": "D:\\dandan123\\dandan\\issue\\investigation-rl-kernel\\RL-Kernel\\.venv\\Lib\\site-packages\\rl_engine", - "companion_library": "D:\\dandan123\\dandan\\issue\\investigation-rl-kernel\\RL-Kernel\\.venv\\Lib\\site-packages\\rl_engine\\_rope_cuda.cp313-win_amd64.pyd", - "native_dispatcher_registered": true, - "torch": "2.9.1+cu128", - "gpu": "NVIDIA GeForce RTX 5080", - "license_present": true, - "packed_qkv_stride": [ - 17280, - 1152, - 128, - 1 - ], - "checks": { - "cuda_eager": "passed", - "cuda_compiled": "passed", - "cuda_trace": { - "contract": "qwenimage-interleaved-rope-fp32-v1", - "actual_backend": "rlkernel.cuda.qwenimage_multi_axis_rope.v1", - "kernel_id": "torch.ops.rl_kernel_native.multi_axis_rope", - "layout": "BSHD, adjacent channel pairs", - "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", - "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" - }, - "triton_eager": "passed", - "triton_compiled": "passed", - "triton_trace": { - "contract": "qwenimage-interleaved-rope-fp32-v1", - "actual_backend": "rlkernel.triton.qwenimage_multi_axis_rope.v1", - "kernel_id": "rl_engine.backends.shared.triton.rope._flash_attn_rotary.rotary_kernel", - "layout": "BSHD, adjacent channel pairs", - "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", - "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" - } - }, - "passed": true -} diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/installed-wheel.txt b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/installed-wheel.txt deleted file mode 100644 index 01c1dd377..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/installed-wheel.txt +++ /dev/null @@ -1,38 +0,0 @@ -WARNING 10-28 17:08:06 [RL-Kernel]: Core binary extension (_C) unavailable: cannot import name '_C' from 'rl_engine' (D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\.venv\Lib\site-packages\rl_engine\__init__.py). Falling back to native code. -{ - "package": "D:\\dandan123\\dandan\\issue\\investigation-rl-kernel\\RL-Kernel\\.venv\\Lib\\site-packages\\rl_engine", - "companion_library": "D:\\dandan123\\dandan\\issue\\investigation-rl-kernel\\RL-Kernel\\.venv\\Lib\\site-packages\\rl_engine\\_rope_cuda.cp313-win_amd64.pyd", - "native_dispatcher_registered": true, - "torch": "2.9.1+cu128", - "gpu": "NVIDIA GeForce RTX 5080", - "license_present": true, - "packed_qkv_stride": [ - 17280, - 1152, - 128, - 1 - ], - "checks": { - "cuda_eager": "passed", - "cuda_compiled": "passed", - "cuda_trace": { - "contract": "qwenimage-interleaved-rope-fp32-v1", - "actual_backend": "rlkernel.cuda.qwenimage_multi_axis_rope.v1", - "kernel_id": "torch.ops.rl_kernel_native.multi_axis_rope", - "layout": "BSHD, adjacent channel pairs", - "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", - "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" - }, - "triton_eager": "passed", - "triton_compiled": "passed", - "triton_trace": { - "contract": "qwenimage-interleaved-rope-fp32-v1", - "actual_backend": "rlkernel.triton.qwenimage_multi_axis_rope.v1", - "kernel_id": "rl_engine.backends.shared.triton.rope._flash_attn_rotary.rotary_kernel", - "layout": "BSHD, adjacent channel pairs", - "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", - "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" - } - }, - "passed": true -} diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/layout.txt b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/layout.txt deleted file mode 100644 index 0086d019b..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/layout.txt +++ /dev/null @@ -1,138 +0,0 @@ -.FF.............s............F.FF.... [100%] -================================== FAILURES =================================== -________________ test_cuda_build_keeps_existing_ipc_collective ________________ - -monkeypatch = <_pytest.monkeypatch.MonkeyPatch object at 0x000001D4B0E668B0> - - def test_cuda_build_keeps_existing_ipc_collective(monkeypatch) -> None: - extension = _load_extension_config(monkeypatch, hip=None) - - assert "csrc/cuda/collectives/deterministic_collective.cu" in extension["sources"] - assert "-DKERNEL_ALIGN_WITH_CUDA" in extension["extra_compile_args"]["cxx"] - assert "-DKERNEL_ALIGN_WITH_ROCM" not in extension["extra_compile_args"]["cxx"] -> assert "-lcuda" in extension["extra_link_args"] -E AssertionError: assert '-lcuda' in ['-Wl,-rpath,$ORIGIN/../torch/lib'] - -tests\build\test_build_platform_collectives.py:61: AssertionError -____________ test_package_data_json_links_resolve_to_real_content _____________ - -tmp_path = WindowsPath('C:/Users/123/AppData/Local/Temp/pytest-of-123/pytest-3/test_package_data_json_links_r0') - - def test_package_data_json_links_resolve_to_real_content(tmp_path) -> None: - target = tmp_path / "configs" / "profile.json" - target.parent.mkdir() - target.write_text('{"name": "profile"}\n', encoding="utf-8") - source = tmp_path / "rl_engine" / "profile.json" - source.parent.mkdir() -> source.symlink_to("../configs/profile.json") - -tests\build\test_build_platform_collectives.py:70: -_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ - -self = WindowsPath('C:/Users/123/AppData/Local/Temp/pytest-of-123/pytest-3/test_package_data_json_links_r0/rl_engine/profile.json') -target = '../configs/profile.json', target_is_directory = False - - def symlink_to(self, target, target_is_directory=False): - """ - Make this path a symlink pointing to the target path. - Note the order of arguments (link, target) is the reverse of os.symlink. - """ -> os.symlink(target, self, target_is_directory) -E OSError: [WinError 1314] �ͻ���û���������Ȩ��: '../configs/profile.json' -> 'C:\\Users\\123\\AppData\\Local\\Temp\\pytest-of-123\\pytest-3\\test_package_data_json_links_r0\\rl_engine\\profile.json' - -D:\anaconda\Lib\pathlib\_local.py:789: OSError -_______ test_recorded_algorithm_sources_resolve_to_real_implementations _______ - - def test_recorded_algorithm_sources_resolve_to_real_implementations(): - from rl_engine.config.workload import load_manifest - - moves = json.loads((ROOT / "tools/migration/layout.json").read_text())["moves"] - for case in load_manifest().representative_cases: - source, symbol = case["provenance_evidence"]["algorithm_source"].rsplit(":", 1) - implementation = ROOT / moves.get(source, source) -> text = implementation.read_text() - ^^^^^^^^^^^^^^^^^^^^^^^^^^ - -tests\runtime\test_layout_compatibility.py:87: -_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ -D:\anaconda\Lib\pathlib\_local.py:546: in read_text - return PathBase.read_text(self, encoding, errors, newline) - ^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ -_ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ _ - -self = WindowsPath('D:/dandan123/dandan/issue/investigation-rl-kernel/RL-Kernel/csrc/cuda/attention/deterministic_attention.cu') -encoding = 'locale', errors = None, newline = None - - def read_text(self, encoding=None, errors=None, newline=None): - """ - Open the file in text mode, read it, and close the file. - """ - with self.open(mode='r', encoding=encoding, errors=errors, newline=newline) as f: -> return f.read() - ^^^^^^^^ -E UnicodeDecodeError: 'gbk' codec can't decode byte 0x92 in position 173: illegal multibyte sequence - -D:\anaconda\Lib\pathlib\_abc.py:633: UnicodeDecodeError -_ test_packaged_contract_bytes_match_main[rl_engine.config-workloads/qwen3_8b.json-80395644b5e84ed1a76729ffc6f78eb11f41b9ab42fb36e2d00f19635d3df0f6] _ - -package = 'rl_engine.config', resource = 'workloads/qwen3_8b.json' -digest = '80395644b5e84ed1a76729ffc6f78eb11f41b9ab42fb36e2d00f19635d3df0f6' - - @pytest.mark.parametrize( - ("package", "resource", "digest"), - [ - ( - "rl_engine.config", - "workloads/qwen3_8b.json", - "80395644b5e84ed1a76729ffc6f78eb11f41b9ab42fb36e2d00f19635d3df0f6", - ), - ( - "rl_engine.contracts", - "profiles/precision/ws1.json", - "044167ec21e372017a2758dc09f805b368283c06b83f64061877628cbd101532", - ), - ], - ) - def test_packaged_contract_bytes_match_main(package, resource, digest): -> assert hashlib.sha256(files(package).joinpath(resource).read_bytes()).hexdigest() == digest -E AssertionError: assert 'c97e966fb075...75e111e052d7d' == '80395644b5e8...f19635d3df0f6' -E -E - 80395644b5e84ed1a76729ffc6f78eb11f41b9ab42fb36e2d00f19635d3df0f6 -E + c97e966fb075dad8ce7dcd7a23a58ffb856e7474dfb3c50510675e111e052d7d - -tests\runtime\test_layout_compatibility.py:118: AssertionError -_ test_packaged_contract_bytes_match_main[rl_engine.contracts-profiles/precision/ws1.json-044167ec21e372017a2758dc09f805b368283c06b83f64061877628cbd101532] _ - -package = 'rl_engine.contracts', resource = 'profiles/precision/ws1.json' -digest = '044167ec21e372017a2758dc09f805b368283c06b83f64061877628cbd101532' - - @pytest.mark.parametrize( - ("package", "resource", "digest"), - [ - ( - "rl_engine.config", - "workloads/qwen3_8b.json", - "80395644b5e84ed1a76729ffc6f78eb11f41b9ab42fb36e2d00f19635d3df0f6", - ), - ( - "rl_engine.contracts", - "profiles/precision/ws1.json", - "044167ec21e372017a2758dc09f805b368283c06b83f64061877628cbd101532", - ), - ], - ) - def test_packaged_contract_bytes_match_main(package, resource, digest): -> assert hashlib.sha256(files(package).joinpath(resource).read_bytes()).hexdigest() == digest -E AssertionError: assert '2020c02e7af4...313d6752cd34e' == '044167ec21e3...7628cbd101532' -E -E - 044167ec21e372017a2758dc09f805b368283c06b83f64061877628cbd101532 -E + 2020c02e7af4e2ca4163e1ecfb9c491baef07b73deba39b0f44313d6752cd34e - -tests\runtime\test_layout_compatibility.py:118: AssertionError -=========================== short test summary info =========================== -FAILED tests/build/test_build_platform_collectives.py::test_cuda_build_keeps_existing_ipc_collective -FAILED tests/build/test_build_platform_collectives.py::test_package_data_json_links_resolve_to_real_content -FAILED tests/runtime/test_layout_compatibility.py::test_recorded_algorithm_sources_resolve_to_real_implementations -FAILED tests/runtime/test_layout_compatibility.py::test_packaged_contract_bytes_match_main[rl_engine.config-workloads/qwen3_8b.json-80395644b5e84ed1a76729ffc6f78eb11f41b9ab42fb36e2d00f19635d3df0f6] -FAILED tests/runtime/test_layout_compatibility.py::test_packaged_contract_bytes_match_main[rl_engine.contracts-profiles/precision/ws1.json-044167ec21e372017a2758dc09f805b368283c06b83f64061877628cbd101532] -5 failed, 31 passed, 1 skipped in 5.67s diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/manifest.json b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/manifest.json deleted file mode 100644 index 190cf7834..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/manifest.json +++ /dev/null @@ -1,65 +0,0 @@ -{ - "recorded_at_utc": "2026-10-28T09:12:00.361759+00:00", - "base_commit": "16b81175874ecf11d5ff05c00de3b69e9457591b", - "source_sha256": { - "MANIFEST.in": "706ec113547f7012cfa6d15d062817d5d2d9e0c0e18e0fb706eca5ba2335bcc5", - "benchmarks/operators/rope/benchmark_multi_axis_rope.py": "6961a55803fbd4e566bd060558fb59e658e7f1f704e61ca97f912e5d16b55884", - "build_tools/extensions.py": "cf07cd49cba0ba7b4da459c26546380b7b288232fcc032e7641863e09ae3014a", - "csrc/bindings/rope_module.cpp": "c69eb3f200de306ca5885fd7c4c7ccf49c822ed1f2f1f639223ec36f0c51b524", - "csrc/cuda/rope/multi_axis_rope.cu": "f5db5cf5d4962053194455d9cde8cf91cf4e824bad7abf72e8410a404bcb0e44", - "docs/.nav.yml": "814a6d60839465e8f9405a9a26306dbc203b0fcb6aa36db264fd7651624d7ef3", - "docs/operators/README.md": "ab926bf5a3fb719e4a57cb65f0784a52019f6ed4bd49c0b5f1b15a1765fd6d7a", - "docs/operators/multi-axis-rope.md": "79f94810ee15ead5f1ce24715456f65f12b383728f2a3571104614b1a33b47d8", - "pyproject.toml": "a132159ae2a2c077d9fe001194fa2b9ab382acbcd58a9eccdea9387286bfb258", - "rl_engine/backends/cuda/rope/multi_axis_rope.py": "d779d9b60ce5d1d6651925f723049930d898024fdf0548ce29a2019b84753d65", - "rl_engine/backends/shared/triton/rope/FLASH_ATTENTION_LICENSE": "8c9ccb96c065e706135b6cbad279b721da6156e51f3a5f27c6b3329af9416d73", - "rl_engine/backends/shared/triton/rope/_flash_attn_rotary.py": "5c6984da351e5228f7549629d4ff53e5772b9a6ef2a01ba2a38493e769b15ab5", - "rl_engine/backends/shared/triton/rope/multi_axis_rope.py": "5b989bd63efd77913267b48654d10204bc563768c6c5d4e1cde9fb55e1354547", - "rl_engine/ops/rope/_registration.py": "add0375d1639444908091d01e8d29781f7a975beffec64b77bfaa81a5b9a0ae3", - "rl_engine/ops/rope/multi_axis_rope.py": "07bcff1160f22c1ff6730be731206b2f4a153e2fada26afc93fc884e780568cb", - "rl_engine/reference/rope/multi_axis_rope.py": "c9a94730047f51659f7aa0a443e88d6ce7d03d00d6813a9aed548cafbf6347bf", - "rl_engine/runtime/registry.py": "9f9c69bb6c2e941b057a66ac2d388be021cbb4b55cb3c602976844aabc6106e1", - "rl_engine/validation/operators/gradient_adapters.py": "5ee2e9ad619a4cad6e81d8f4a6ccf81ecf45e63a7764d0fcb4c204206b37c16b", - "rl_engine/validation/operators/operator_inputs.py": "55c8d9e7cbde2070662dd1d0ebe9888aef106b45fb05902057aec032d29a9836", - "rl_engine/validation/operators/operator_specs.py": "cef2e5bed8075509d5d5a538b378477b09f1adfe5fd01829156357cefa99f983", - "rl_engine/validation/operators/qwen_image_rope.py": "c3ddfc732f32de3c633de16064810d67202254ea78fd8107d558b471bb15a37f", - "tests/build/test_multi_axis_rope_build.py": "90c3740b2d4946acc6206d952622c19afaa0089cca904919088563427555b956", - "tests/ops/rope/test_multi_axis_rope.py": "1a0950b09eec407e8ffd490b40dda4118f8d9b876f4db4af8aac648af1a76687", - "tests/validation/operators/test_operator_inputs.py": "36365f662194f158276b3c07f6447471a0cf17253ec420ab63c8b9149c511b49", - "tools/validation/operators/check_multi_axis_rope_install.py": "ff8198ff55779ea2ede88d4ecfe07caa627631b4b08bce38139ccfa147abf538", - "tools/validation/operators/check_qwenimage_rope_upstream.py": "cc368d5bebf0f0acb21b7455348592d30442b6a8d349de0f49d5a09c5f20ce4e", - "tools/validation/operators/validate_multi_axis_rope.py": "def4747da9ea9194c761bef24479017ac74ac10de46a3fbe1a9528ec6ee07731" - }, - "evidence_sha256": { - "baseline-layout.txt": "b050d798a5e3a6c2d43c7c22e4e1140529ebf05ebab982df17f501f5b69cff12", - "benchmark-bfloat16.json": "0679a422e8129a8a36ada938b024794334f784cdea28013574f99c70673e13a0", - "installed-wheel.json": "60f6a1a48d35fa25a639a1ff168ce0c37659e0258737422cafda3741d8ef0200", - "installed-wheel.txt": "2be6479a8e30938410095dbcc7724ad114cb61cb35ef86c7f86b8d507a3193f6", - "layout.txt": "83169d7995bdadc14744a0ed0fff69b334f2ed6ff0b92cda005a9cf0f733116d", - "memcheck.txt": "f6f05d4e5b772909f29baf2a1d6be04b4456660ca024e5067c3f64c7efa4b81f", - "operator-cuda.json": "1a231d26f2411a38cffe70fdb59b388227fc0e7da9625370b165d7d590e1e011", - "operator-triton.json": "af2be1f4df1bf9d7afb4802ca02c81826e64acacae891436f83ad6b3342109f5", - "package-build.txt": "34b11d626dc4988f3d862e4122599a0b1a949b1d799436981dd74f3ee5c1b07e", - "package.json": "e267d82707cdeb0c53d120928cd32a544a44190e8e41c43f3b6c1d4668275a7e", - "precommit.txt": "04f9db16780e0ad0b41e6450d362e6f0a7bdb07128596c8ea13f8ea8bb2cc674", - "quality.txt": "d4728355fe488102ca644f510795c474b54f8d42bb92ddfda34557ba521c6afd", - "README.md": "aa2cabc1e8423f4d737d0db90786142c4a919041dbf97335b17991475a4cbaca", - "regression.txt": "5d058ac93516d9ce1f486fbd4f35b84a254a32d8ddf26eedf41cf2388153e004", - "tests.txt": "f6af2230f6a0b5d993cb6bca2724a82f772c6b33df456edce6e728996f2a76d5", - "upstream.json": "8c087d60013f0266303c6d60237d87e5f1ff67b03e8e8c1900de4b20349e2e93", - "ws1-cuda.json": "531124223ae5fe43eca890fb1ccf80e09f33affe34bab132c7c17742081d85a3", - "ws1-cuda.txt": "b4cf09ef5165e535dba9ed19eeec08f15eccaf86b549bb1a253fd08e87be35ab", - "ws1-triton.json": "29eedee8e50b8750d0ce9be95deb39fc6aa9ab47b0867fe20ce2e7ad5cb47548", - "ws1-triton.txt": "a1daefba35f7c1a06d222750ad1ce9c4bdc97cc03410969e7b4f4ce011a6770a" - }, - "layout_failures_equal_clean_base": true, - "layout_failure_ids": [ - "tests/build/test_build_platform_collectives.py::test_cuda_build_keeps_existing_ipc_collective", - "tests/build/test_build_platform_collectives.py::test_package_data_json_links_resolve_to_real_content", - "tests/runtime/test_layout_compatibility.py::test_recorded_algorithm_sources_resolve_to_real_implementations", - "tests/runtime/test_layout_compatibility.py::test_packaged_contract_bytes_match_main[rl_engine.config-workloads/qwen3_8b.json-80395644b5e84ed1a76729ffc6f78eb11f41b9ab42fb36e2d00f19635d3df0f6]", - "tests/runtime/test_layout_compatibility.py::test_packaged_contract_bytes_match_main[rl_engine.contracts-profiles/precision/ws1.json-044167ec21e372017a2758dc09f805b368283c06b83f64061877628cbd101532]" - ], - "package_source_and_license_checks": "passed", - "hash_mode": "SHA256 of recorded working-tree bytes; upstream baselines separately normalize newlines" -} diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/memcheck.txt b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/memcheck.txt deleted file mode 100644 index dd59e40b0..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/memcheck.txt +++ /dev/null @@ -1,8 +0,0 @@ -========= COMPUTE-SANITIZER -WARNING 10-28 16:46:09 [RL-Kernel]: Core binary extension (_C) unavailable: cannot import name '_C' from 'rl_engine' (D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\rl_engine\__init__.py). Falling back to native code. -W1028 16:46:10.075000 16608 site-packages\torch\utils\cpp_extension.py:480] Error checking compiler version for cl: 'cp1' codec can't decode bytes in position 0--1: No mapping for the Unicode character exists in the target code page. -ninja: no work to do. -Focused CUDA source build; torch=2.9.1+cu128, GPU=NVIDIA GeForce RTX 5080 -.................................. [100%] -34 passed, 48 deselected in 3.96s -========= ERROR SUMMARY: 0 errors diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/operator-cuda.json b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/operator-cuda.json deleted file mode 100644 index 03d59bb4b..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/operator-cuda.json +++ /dev/null @@ -1,70 +0,0 @@ -{ - "suite_name": "multi_axis_rope", - "total_candidates": 1, - "passed_candidates": 1, - "pass_rate": 1.0, - "passed": true, - "candidates": [ - { - "candidate_name": "cuda-multi_axis_rope", - "backend": "cuda", - "total_outputs": 2, - "passed_outputs": 2, - "pass_rate": 1.0, - "passed": true, - "cases": [ - { - "case_name": "multi_axis_rope-torch.bfloat16-2x7x3x128", - "dtype": "torch.bfloat16", - "op_class": "elementwise", - "passed": true, - "outputs": [ - { - "output_index": 0, - "shape": [ - 2, - 7, - 3, - 128 - ], - "candidate_dtype": "torch.bfloat16", - "gold_dtype": "torch.float32", - "atol": 0.02, - "rtol": 0.016, - "max_abs_error": 0.007704973220825195, - "mean_abs_error": 0.0006018931162543595, - "max_rel_error": 0.003828622167930007, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference", - "message": "" - }, - { - "output_index": 1, - "shape": [ - 2, - 7, - 3, - 128 - ], - "candidate_dtype": "torch.bfloat16", - "gold_dtype": "torch.bfloat16", - "atol": 0.02, - "rtol": 0.016, - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference", - "message": "gradient:x" - } - ] - } - ], - "backend_provenance": null - } - ] -} diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/operator-triton.json b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/operator-triton.json deleted file mode 100644 index 6f086edac..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/operator-triton.json +++ /dev/null @@ -1,70 +0,0 @@ -{ - "suite_name": "multi_axis_rope", - "total_candidates": 1, - "passed_candidates": 1, - "pass_rate": 1.0, - "passed": true, - "candidates": [ - { - "candidate_name": "triton-multi_axis_rope", - "backend": "triton", - "total_outputs": 2, - "passed_outputs": 2, - "pass_rate": 1.0, - "passed": true, - "cases": [ - { - "case_name": "multi_axis_rope-torch.bfloat16-2x7x3x128", - "dtype": "torch.bfloat16", - "op_class": "elementwise", - "passed": true, - "outputs": [ - { - "output_index": 0, - "shape": [ - 2, - 7, - 3, - 128 - ], - "candidate_dtype": "torch.bfloat16", - "gold_dtype": "torch.float32", - "atol": 0.02, - "rtol": 0.016, - "max_abs_error": 0.007704973220825195, - "mean_abs_error": 0.0006018931162543595, - "max_rel_error": 0.003828622167930007, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference", - "message": "" - }, - { - "output_index": 1, - "shape": [ - 2, - 7, - 3, - 128 - ], - "candidate_dtype": "torch.bfloat16", - "gold_dtype": "torch.bfloat16", - "atol": 0.02, - "rtol": 0.016, - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference", - "message": "gradient:x" - } - ] - } - ], - "backend_provenance": null - } - ] -} diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/package-build.txt b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/package-build.txt deleted file mode 100644 index d63d79d52..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/package-build.txt +++ /dev/null @@ -1,40 +0,0 @@ -Selected lines from the completed sdist/wheel build log. -Full local log SHA256: bfbcd342efaf8caa36adcd208acee385db1ba6fbc4326fe8e0325f9523f199b0 -running sdist -copying rl_engine\backends\shared\triton\rope\FLASH_ATTENTION_LICENSE -> rl_kernel-0.1.0\rl_engine\backends\shared\triton\rope -running bdist_wheel -copying rl_engine\backends\shared\triton\rope\FLASH_ATTENTION_LICENSE -> build\lib.win-amd64-cpython-313\rl_engine\backends\shared\triton\rope -copying rl_engine\backends\shared\triton\rope\FLASH_ATTENTION_LICENSE -> build\lib.win-amd64-cpython-313\rl_engine\backends\shared\triton\rope -running build_ext -ninja: no work to do. -"C:\Program Files (x86)\Microsoft Visual Studio\2022\BuildTools\VC\Tools\MSVC\14.44.35207\bin\HostX64\x64\link.exe" /nologo /INCREMENTAL:NO /LTCG /DLL /MANIFEST:EMBED,ID=2 /MANIFESTUAC:NO /LIBPATH:D:\dandan123\dandan\myvllm\.cuda-12.8\Library\lib /LIBPATH:D:\anaconda\Lib\site-packages\torch\lib /LIBPATH:D:\dandan123\dandan\myvllm\.cuda-12.8\Library\lib\x64 /LIBPATH:D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\.venv\libs /LIBPATH:D:\anaconda\libs /LIBPATH:D:\anaconda /LIBPATH:D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\.venv\PCbuild\amd64 "/LIBPATH:C:\Program Files (x86)\Microsoft Visual Studio\2022\BuildTools\VC\Tools\MSVC\14.44.35207\lib\x64" "/LIBPATH:C:\Program Files (x86)\Windows Kits\10\lib\10.0.26100.0\ucrt\x64" "/LIBPATH:C:\Program Files (x86)\Windows Kits\10\\lib\10.0.26100.0\\um\x64" c10.lib torch.lib torch_cpu.lib torch_python.lib cudart.lib c10_cuda.lib torch_cuda.lib /EXPORT:PyInit__rope_cuda D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\build\temp.win-amd64-cpython-313\Release\csrc\bindings\rope_module.obj D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\build\temp.win-amd64-cpython-313\Release\csrc\cuda\rope\multi_axis_rope.obj /OUT:build\lib.win-amd64-cpython-313\rl_engine\_rope_cuda.cp313-win_amd64.pyd /IMPLIB:D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\build\temp.win-amd64-cpython-313\Release\csrc\bindings\_rope_cuda.cp313-win_amd64.lib -copying build\lib.win-amd64-cpython-313\rl_engine\backends\shared\triton\rope\FLASH_ATTENTION_LICENSE -> build\bdist.win-amd64\wheel\.\rl_engine\backends\shared\triton\rope -copying build\lib.win-amd64-cpython-313\rl_engine\_rope_cuda.cp313-win_amd64.pyd -> build\bdist.win-amd64\wheel\.\rl_engine -adding 'rl_engine/_rope_cuda.cp313-win_amd64.pyd' -adding 'rl_engine/backends/shared/triton/rope/FLASH_ATTENTION_LICENSE' - -In-place build log: -D:\anaconda\Lib\site-packages\setuptools\config\_apply_pyprojecttoml.py:82: SetuptoolsDeprecationWarning: `project.license` as a TOML table is deprecated -!! - - ******************************************************************************** - Please use a simple string containing a SPDX expression for `project.license`. You can also use `project.license-files`. (Both options available on setuptools>=77.0.0). - - By 2027-Feb-18, you need to update your project and remove deprecated calls - or your builds will no longer be supported. - - See https://packaging.python.org/en/latest/guides/writing-pyproject-toml/#license for details. - ******************************************************************************** - -!! - corresp(dist, value, root_dir) -running build_ext -W1028 16:52:54.107000 9384 site-packages\torch\utils\cpp_extension.py:480] Error checking compiler version for cl: 'cp1' codec can't decode bytes in position 0--1: No mapping for the Unicode character exists in the target code page. -building 'rl_engine._rope_cuda' extension -W1028 16:52:54.150000 9384 site-packages\torch\utils\cpp_extension.py:480] Error checking compiler version for cl: 'cp1' codec can't decode bytes in position 0--1: No mapping for the Unicode character exists in the target code page. -ninja: no work to do. -"C:\Program Files (x86)\Microsoft Visual Studio\2022\BuildTools\VC\Tools\MSVC\14.44.35207\bin\HostX64\x64\link.exe" /nologo /INCREMENTAL:NO /LTCG /DLL /MANIFEST:EMBED,ID=2 /MANIFESTUAC:NO /LIBPATH:D:\dandan123\dandan\myvllm\.cuda-12.8\Library\lib /LIBPATH:D:\anaconda\Lib\site-packages\torch\lib /LIBPATH:D:\dandan123\dandan\myvllm\.cuda-12.8\Library\lib\x64 /LIBPATH:D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\.venv\libs /LIBPATH:D:\anaconda\libs /LIBPATH:D:\anaconda /LIBPATH:D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\.venv\PCbuild\amd64 "/LIBPATH:C:\Program Files (x86)\Microsoft Visual Studio\2022\BuildTools\VC\Tools\MSVC\14.44.35207\lib\x64" "/LIBPATH:C:\Program Files (x86)\Windows Kits\10\lib\10.0.26100.0\ucrt\x64" "/LIBPATH:C:\Program Files (x86)\Windows Kits\10\\lib\10.0.26100.0\\um\x64" c10.lib torch.lib torch_cpu.lib torch_python.lib cudart.lib c10_cuda.lib torch_cuda.lib /EXPORT:PyInit__rope_cuda D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\build\temp.win-amd64-cpython-313\Release\csrc\bindings\rope_module.obj D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\build\temp.win-amd64-cpython-313\Release\csrc\cuda\rope\multi_axis_rope.obj /OUT:build\lib.win-amd64-cpython-313\rl_engine\_rope_cuda.cp313-win_amd64.pyd /IMPLIB:D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\build\temp.win-amd64-cpython-313\Release\csrc\bindings\_rope_cuda.cp313-win_amd64.lib - 正在创建库 D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\build\temp.win-amd64-cpython-313\Release\csrc\bindings\_rope_cuda.cp313-win_amd64.lib 和对象 D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\build\temp.win-amd64-cpython-313\Release\csrc\bindings\_rope_cuda.cp313-win_amd64.exp -正在生成代码 -已完成代码的生成 -copying build\lib.win-amd64-cpython-313\rl_engine\_rope_cuda.cp313-win_amd64.pyd -> rl_engine diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/package.json b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/package.json deleted file mode 100644 index 51e102428..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/package.json +++ /dev/null @@ -1,31 +0,0 @@ -{ - "wheel": "rl_kernel-0.1.0-cp313-cp313-win_amd64.whl", - "wheel_sha256": "bc5418bbc4be779ec6d5b4fa99d54b3244f2570d60a29adc3e6a1c38be25a864", - "sdist": "rl_kernel-0.1.0.tar.gz", - "sdist_sha256": "56026574fa3307e85e6e278298a09deee13b43d4ec5c61899698ac5e55163a12", - "native_library_member": "rl_engine/_rope_cuda.cp313-win_amd64.pyd", - "native_library_sha256": "c0878eeac05b6b400ed7c3935561de525ce3af8b2a050b5d8bfe5cdd66650940", - "inplace_library_sha256": "aab5bf0777da43609e067e6a41bc37c46fdf6b1dd0c60d1071c15bbc4f707242", - "binary_hash_scope": "Wheel and in-place libraries were linked separately; record both hashes", - "license_in_both_archives": true, - "license_sha256": "8c9ccb96c065e706135b6cbad279b721da6156e51f3a5f27c6b3329af9416d73", - "wheel_sources_match_checkout": [ - "rl_engine/backends/cuda/rope/multi_axis_rope.py", - "rl_engine/backends/shared/triton/rope/_flash_attn_rotary.py", - "rl_engine/backends/shared/triton/rope/multi_axis_rope.py", - "rl_engine/ops/rope/_registration.py", - "rl_engine/ops/rope/multi_axis_rope.py", - "rl_engine/reference/rope/multi_axis_rope.py", - "rl_engine/runtime/registry.py", - "rl_engine/validation/operators/gradient_adapters.py", - "rl_engine/validation/operators/operator_inputs.py", - "rl_engine/validation/operators/operator_specs.py", - "rl_engine/validation/operators/qwen_image_rope.py" - ], - "sdist_required_sources_match_checkout": [ - "csrc/bindings/rope_module.cpp", - "csrc/cuda/rope/multi_axis_rope.cu", - "rl_engine/backends/shared/triton/rope/FLASH_ATTENTION_LICENSE" - ], - "profile": "RL_KERNEL_ROPE_ONLY=1; full monolithic _C build not qualified" -} diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/precommit.txt b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/precommit.txt deleted file mode 100644 index 475a8c0d8..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/precommit.txt +++ /dev/null @@ -1,7 +0,0 @@ -trim trailing whitespace.................................................Passed -fix end of files.........................................................Passed -check yaml...............................................................Passed -check for added large files..............................................Passed -black....................................................................Passed -isort....................................................................Passed -flake8...................................................................Passed diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/quality.txt b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/quality.txt deleted file mode 100644 index 2756f4004..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/quality.txt +++ /dev/null @@ -1,58 +0,0 @@ -Changed-file checks; repository-pinned Black 24.4.2, isort 5.13.2 and Flake8 7.0.0. -Test commands are separate from the static checks below. -Regression command: python -m pytest tests/build/test_multi_axis_rope_build.py tests/validation/operators/test_operator_inputs.py tests/validation/operators/test_gradient_invariance.py tests/validation/operators/test_forward_invariance.py tests/validation/operators/test_elementwise_inventory.py tests/validation/operators/test_four_judgment_matrix.py tests/validation/operators/test_ws1_workload.py tests/runtime/test_kernel_registry.py tests/runtime/test_dispatch.py tests/ops/rope/test_rope.py -q -rs -Layout output copies have only trailing whitespace normalized by pre-commit. - -Command: python -m isort --check-only --profile black --line-length 100 benchmarks/operators/rope/benchmark_multi_axis_rope.py build_tools/extensions.py rl_engine/backends/cuda/rope/multi_axis_rope.py rl_engine/backends/shared/triton/rope/_flash_attn_rotary.py rl_engine/backends/shared/triton/rope/multi_axis_rope.py rl_engine/ops/rope/_registration.py rl_engine/ops/rope/multi_axis_rope.py rl_engine/reference/rope/multi_axis_rope.py rl_engine/runtime/registry.py rl_engine/validation/operators/gradient_adapters.py rl_engine/validation/operators/operator_inputs.py rl_engine/validation/operators/operator_specs.py rl_engine/validation/operators/qwen_image_rope.py tests/build/test_multi_axis_rope_build.py tests/ops/rope/test_multi_axis_rope.py tests/validation/operators/test_operator_inputs.py tools/validation/operators/check_multi_axis_rope_install.py tools/validation/operators/check_qwenimage_rope_upstream.py tools/validation/operators/validate_multi_axis_rope.py -Exit code: 0 - - -Command: python -m black --check --line-length 100 benchmarks/operators/rope/benchmark_multi_axis_rope.py build_tools/extensions.py rl_engine/backends/cuda/rope/multi_axis_rope.py rl_engine/backends/shared/triton/rope/_flash_attn_rotary.py rl_engine/backends/shared/triton/rope/multi_axis_rope.py rl_engine/ops/rope/_registration.py rl_engine/ops/rope/multi_axis_rope.py rl_engine/reference/rope/multi_axis_rope.py rl_engine/runtime/registry.py rl_engine/validation/operators/gradient_adapters.py rl_engine/validation/operators/operator_inputs.py rl_engine/validation/operators/operator_specs.py rl_engine/validation/operators/qwen_image_rope.py tests/build/test_multi_axis_rope_build.py tests/ops/rope/test_multi_axis_rope.py tests/validation/operators/test_operator_inputs.py tools/validation/operators/check_multi_axis_rope_install.py tools/validation/operators/check_qwenimage_rope_upstream.py tools/validation/operators/validate_multi_axis_rope.py -Exit code: 0 -All done! ✨ 🍰 ✨ - -19 files would be left unchanged. - -Command: python -m flake8 --max-line-length=100 --extend-ignore=E203,E704 benchmarks/operators/rope/benchmark_multi_axis_rope.py build_tools/extensions.py rl_engine/backends/cuda/rope/multi_axis_rope.py rl_engine/backends/shared/triton/rope/_flash_attn_rotary.py rl_engine/backends/shared/triton/rope/multi_axis_rope.py rl_engine/ops/rope/_registration.py rl_engine/ops/rope/multi_axis_rope.py rl_engine/reference/rope/multi_axis_rope.py rl_engine/runtime/registry.py rl_engine/validation/operators/gradient_adapters.py rl_engine/validation/operators/operator_inputs.py rl_engine/validation/operators/operator_specs.py rl_engine/validation/operators/qwen_image_rope.py tests/build/test_multi_axis_rope_build.py tests/ops/rope/test_multi_axis_rope.py tests/validation/operators/test_operator_inputs.py tools/validation/operators/check_multi_axis_rope_install.py tools/validation/operators/check_qwenimage_rope_upstream.py tools/validation/operators/validate_multi_axis_rope.py -Exit code: 0 - - -Command: python -m ruff check benchmarks/operators/rope/benchmark_multi_axis_rope.py build_tools/extensions.py rl_engine/backends/cuda/rope/multi_axis_rope.py rl_engine/backends/shared/triton/rope/_flash_attn_rotary.py rl_engine/backends/shared/triton/rope/multi_axis_rope.py rl_engine/ops/rope/_registration.py rl_engine/ops/rope/multi_axis_rope.py rl_engine/reference/rope/multi_axis_rope.py rl_engine/runtime/registry.py rl_engine/validation/operators/gradient_adapters.py rl_engine/validation/operators/operator_inputs.py rl_engine/validation/operators/operator_specs.py rl_engine/validation/operators/qwen_image_rope.py tests/build/test_multi_axis_rope_build.py tests/ops/rope/test_multi_axis_rope.py tests/validation/operators/test_operator_inputs.py tools/validation/operators/check_multi_axis_rope_install.py tools/validation/operators/check_qwenimage_rope_upstream.py tools/validation/operators/validate_multi_axis_rope.py -Exit code: 0 -All checks passed! - -Command: python -m mkdocs build --strict -f mkdocs.yaml -Exit code: 0 -INFO - Cleaning site directory - -INFO - Building documentation to directory: D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\site - -mkdocs_git_revision_date_localized_plugin: 'D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\docs\operators\multi-axis-rope.md' has no git logs, using current timestamp - -INFO - The following pages exist in the docs directory, but are not included in the "nav" configuration: - - - blog\2026-07-08-announcing-rl-kernel-linear-logp-for-vime-zh.md - - - blog\2026-08-09-dsv4-flash-moe-consistency-roadmap.md - - - operators\matmul.md - - - operators\rope.md - - - usage\cross-config-drift-report.md - - - usage\cuda-rocm-consistency-audit.md - - - usage\h100-configurable.md - - - usage\h100-cp-gradient-validation.md - - - usage\h100-matrix-validation.md - - - usage\h100-pcp-validation.md - - - usage\rocm-sparse-performance.md - -INFO - Documentation built in 4.01 seconds - -Repository pre-commit: passed; see precommit.txt. diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/regression.txt b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/regression.txt deleted file mode 100644 index 066e09cec..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/regression.txt +++ /dev/null @@ -1,10 +0,0 @@ -........................................................................ [ 35%] -........................................................................ [ 71%] -...................................................sssssss [100%] -=========================== short test summary info =========================== -SKIPPED [3] tests\ops\rope\test_rope.py:390: rope_apply_ascend is not compiled (requires an Ascend NPU build). -SKIPPED [1] tests\ops\rope\test_rope.py:406: rope_apply_ascend is not compiled (requires an Ascend NPU build). -SKIPPED [1] tests\ops\rope\test_rope.py:413: rope_apply_ascend is not compiled (requires an Ascend NPU build). -SKIPPED [1] tests\ops\rope\test_rope.py:420: rope_apply_ascend is not compiled (requires an Ascend NPU build). -SKIPPED [1] tests\ops\rope\test_rope.py:433: rope_apply_ascend is not compiled (requires an Ascend NPU build). -195 passed, 7 skipped in 4.31s diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/tests.txt b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/tests.txt deleted file mode 100644 index 38e8e21e6..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/tests.txt +++ /dev/null @@ -1,5 +0,0 @@ -........................................................................ [ 87%] -s......... [100%] -=========================== short test summary info =========================== -SKIPPED [1] tests\ops\rope\test_multi_axis_rope.py:336: requires two NVIDIA GPUs -81 passed, 1 skipped in 5.55s diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/upstream.json b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/upstream.json deleted file mode 100644 index 1f89880c1..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/upstream.json +++ /dev/null @@ -1,335 +0,0 @@ -{ - "source_sha256": "0bad11901deae69f05aec01fd464e51e38582d455068f2851ee665036179321f", - "cases": [ - { - "grid": [ - 2, - 3 - ], - "stream": "image", - "dtype": "torch.float32", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 2, - 3 - ], - "stream": "image", - "dtype": "torch.float16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 2, - 3 - ], - "stream": "image", - "dtype": "torch.bfloat16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 2, - 3 - ], - "stream": "text", - "dtype": "torch.float32", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 2, - 3 - ], - "stream": "text", - "dtype": "torch.float16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 2, - 3 - ], - "stream": "text", - "dtype": "torch.bfloat16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 3, - 5 - ], - "stream": "image", - "dtype": "torch.float32", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 3, - 5 - ], - "stream": "image", - "dtype": "torch.float16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 3, - 5 - ], - "stream": "image", - "dtype": "torch.bfloat16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 3, - 5 - ], - "stream": "text", - "dtype": "torch.float32", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 3, - 5 - ], - "stream": "text", - "dtype": "torch.float16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 3, - 5 - ], - "stream": "text", - "dtype": "torch.bfloat16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 64, - 64 - ], - "stream": "image", - "dtype": "torch.float32", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 64, - 64 - ], - "stream": "image", - "dtype": "torch.float16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 64, - 64 - ], - "stream": "image", - "dtype": "torch.bfloat16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 64, - 64 - ], - "stream": "text", - "dtype": "torch.float32", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 64, - 64 - ], - "stream": "text", - "dtype": "torch.float16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 64, - 64 - ], - "stream": "text", - "dtype": "torch.bfloat16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 83, - 83 - ], - "stream": "image", - "dtype": "torch.float32", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 83, - 83 - ], - "stream": "image", - "dtype": "torch.float16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 83, - 83 - ], - "stream": "image", - "dtype": "torch.bfloat16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 83, - 83 - ], - "stream": "text", - "dtype": "torch.float32", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 83, - 83 - ], - "stream": "text", - "dtype": "torch.float16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 83, - 83 - ], - "stream": "text", - "dtype": "torch.bfloat16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 104, - 58 - ], - "stream": "image", - "dtype": "torch.float32", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 104, - 58 - ], - "stream": "image", - "dtype": "torch.float16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 104, - 58 - ], - "stream": "image", - "dtype": "torch.bfloat16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 104, - 58 - ], - "stream": "text", - "dtype": "torch.float32", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 104, - 58 - ], - "stream": "text", - "dtype": "torch.float16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 104, - 58 - ], - "stream": "text", - "dtype": "torch.bfloat16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - } - ] -} diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-cuda.json b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-cuda.json deleted file mode 100644 index 20ad4dd78..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-cuda.json +++ /dev/null @@ -1,1962 +0,0 @@ -{ - "scope": "Qwen-Image multi_axis_rope operator; C1 elementwise and C2 config transforms", - "trace": { - "contract": "qwenimage-interleaved-rope-fp32-v1", - "actual_backend": "rlkernel.cuda.qwenimage_multi_axis_rope.v1", - "kernel_id": "torch.ops.rl_kernel_native.multi_axis_rope", - "layout": "BSHD, adjacent channel pairs", - "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", - "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" - }, - "passed": true, - "forward": { - "op_name": "multi_axis_rope", - "backend_profile": "cuda_bf16", - "accuracy_reports": [ - { - "config_id": "BN/full", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.054684440139681e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/chunked", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/chunked", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.054684440139681e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s0", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/full/s0", - "fp32_reference" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004673302173614502, - "mean_abs_error": 6.917401333339512e-05, - "max_rel_error": 0.003753266530111432, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s1", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/full/s1", - "fp32_reference" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.00048516690731048584, - "mean_abs_error": 7.802647451171651e-05, - "max_rel_error": 0.0037793784867972136, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s2", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/full/s2", - "fp32_reference" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004864335060119629, - "mean_abs_error": 7.697804539930075e-05, - "max_rel_error": 0.003821910358965397, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s3", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/full/s3", - "fp32_reference" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.710371912457049e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s0", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/chunked/s0", - "fp32_reference" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004673302173614502, - "mean_abs_error": 6.917401333339512e-05, - "max_rel_error": 0.003753266530111432, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s1", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/chunked/s1", - "fp32_reference" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.00048516690731048584, - "mean_abs_error": 7.802647451171651e-05, - "max_rel_error": 0.0037793784867972136, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s2", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/chunked/s2", - "fp32_reference" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004864335060119629, - "mean_abs_error": 7.697804539930075e-05, - "max_rel_error": 0.003821910358965397, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s3", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/chunked/s3", - "fp32_reference" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.710371912457049e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/permuted", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/permuted", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.054684440139681e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/padded_right", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/padded_right", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.054684440139681e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/padded_left", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/padded_left", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.054684440139681e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - } - ], - "invariance_reports": [ - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/chunked", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "BN/chunked" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s0", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s0" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s1", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s1" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s2", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s2" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s3", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s3" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s0", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s0" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s1", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s1" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s2", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s2" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s3", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s3" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/permuted", - "transform_kind": "permutation", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "BN/permuted" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/padded_right", - "transform_kind": "padding", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "BN/padded_right" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/padded_left", - "transform_kind": "padding", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "BN/padded_left" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - } - ], - "logprob_smoke": null, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - }, - "candidate_id": "rlkernel.cuda.qwenimage_multi_axis_rope.v1", - "device": "cuda", - "compute_capability": "sm120", - "seed": 20260812, - "fallback_reason": null, - "passed": true, - "provenance_valid": true, - "metadata_valid": true, - "observed_kernel_id": "torch.ops.rl_kernel_native.multi_axis_rope" - }, - "gradient": { - "op_name": "multi_axis_rope", - "backend_profile": "cuda_bf16", - "accuracy_reports": [ - { - "config_id": "BN/full", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.699562850873917e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/chunked", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/chunked", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.699562850873917e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s0", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/full/s0", - "fp32_reference" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 2.7887988835573196e-05, - "mean_abs_error": 3.6821968478761846e-06, - "max_rel_error": 1.8164952993392944, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s1", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/full/s1", - "fp32_reference" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.7113780965446495e-06, - "max_rel_error": 0.4316565990447998, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s2", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/full/s2", - "fp32_reference" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 2.995925024151802e-05, - "mean_abs_error": 4.7168209675874095e-06, - "max_rel_error": 0.7439075112342834, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s3", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/full/s3", - "fp32_reference" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.0211638659238815e-05, - "mean_abs_error": 4.9605891945248e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s0", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/chunked/s0", - "fp32_reference" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 2.7887988835573196e-05, - "mean_abs_error": 3.6821968478761846e-06, - "max_rel_error": 1.8164952993392944, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s1", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/chunked/s1", - "fp32_reference" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.7113780965446495e-06, - "max_rel_error": 0.4316565990447998, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s2", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/chunked/s2", - "fp32_reference" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 2.995925024151802e-05, - "mean_abs_error": 4.7168209675874095e-06, - "max_rel_error": 0.7439075112342834, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s3", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/chunked/s3", - "fp32_reference" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.0211638659238815e-05, - "mean_abs_error": 4.9605891945248e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/permuted", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/permuted", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.699562850873917e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/padded_right", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/padded_right", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.699562850873917e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/padded_left", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/padded_left", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.699562850873917e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - } - ], - "invariance_reports": [ - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/chunked", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "BN/chunked" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s0", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s0" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s1", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s1" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s2", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s2" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s3", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s3" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s0", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s0" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s1", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s1" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s2", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s2" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s3", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s3" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/permuted", - "transform_kind": "permutation", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "BN/permuted" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/padded_right", - "transform_kind": "padding", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "BN/padded_right" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/padded_left", - "transform_kind": "padding", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "BN/padded_left" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - } - ], - "singleton_aggregate_reports": [], - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - }, - "candidate_id": "rlkernel.cuda.qwenimage_multi_axis_rope.v1", - "device": "cuda", - "compute_capability": "sm120", - "seed": 20260812, - "fallback_reason": null, - "passed": true, - "provenance_valid": true, - "metadata_valid": true, - "loss_reduction": "sum_over_active_tokens_then_optional_mean_by_active_count", - "active_token_denominator": 27, - "grad_tensor_names": [ - "dx" - ], - "first_failing_op": null, - "first_failing_tensor": null, - "first_failing_config_pair": null, - "observed_kernel_id": "torch.ops.rl_kernel_native.multi_axis_rope" - } -} diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-cuda.txt b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-cuda.txt deleted file mode 100644 index e69d60986..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-cuda.txt +++ /dev/null @@ -1,3 +0,0 @@ -WARNING 10-28 16:54:40 [RL-Kernel]: Core binary extension (_C) unavailable: cannot import name '_C' from 'rl_engine' (D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\rl_engine\__init__.py). Falling back to native code. -Using the installed production RoPE dispatcher library -WS1 four judgments backend=cuda: passed=True diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-triton.json b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-triton.json deleted file mode 100644 index d81adfea1..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-triton.json +++ /dev/null @@ -1,1962 +0,0 @@ -{ - "scope": "Qwen-Image multi_axis_rope operator; C1 elementwise and C2 config transforms", - "trace": { - "contract": "qwenimage-interleaved-rope-fp32-v1", - "actual_backend": "rlkernel.triton.qwenimage_multi_axis_rope.v1", - "kernel_id": "rl_engine.backends.shared.triton.rope._flash_attn_rotary.rotary_kernel", - "layout": "BSHD, adjacent channel pairs", - "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", - "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" - }, - "passed": true, - "forward": { - "op_name": "multi_axis_rope", - "backend_profile": "triton_cuda_bf16", - "accuracy_reports": [ - { - "config_id": "BN/full", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.054684440139681e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/chunked", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/chunked", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.054684440139681e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s0", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/full/s0", - "fp32_reference" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004673302173614502, - "mean_abs_error": 6.917401333339512e-05, - "max_rel_error": 0.003753266530111432, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s1", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/full/s1", - "fp32_reference" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.00048516690731048584, - "mean_abs_error": 7.802647451171651e-05, - "max_rel_error": 0.0037793784867972136, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s2", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/full/s2", - "fp32_reference" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004864335060119629, - "mean_abs_error": 7.697804539930075e-05, - "max_rel_error": 0.003821910358965397, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s3", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/full/s3", - "fp32_reference" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.710371912457049e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s0", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/chunked/s0", - "fp32_reference" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004673302173614502, - "mean_abs_error": 6.917401333339512e-05, - "max_rel_error": 0.003753266530111432, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s1", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/chunked/s1", - "fp32_reference" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.00048516690731048584, - "mean_abs_error": 7.802647451171651e-05, - "max_rel_error": 0.0037793784867972136, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s2", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/chunked/s2", - "fp32_reference" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004864335060119629, - "mean_abs_error": 7.697804539930075e-05, - "max_rel_error": 0.003821910358965397, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s3", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/chunked/s3", - "fp32_reference" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.710371912457049e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/permuted", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/permuted", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.054684440139681e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/padded_right", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/padded_right", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.054684440139681e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/padded_left", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/padded_left", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.054684440139681e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - } - ], - "invariance_reports": [ - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/chunked", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "BN/chunked" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s0", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s0" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s1", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s1" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s2", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s2" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s3", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s3" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s0", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s0" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s1", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s1" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s2", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s2" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s3", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s3" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/permuted", - "transform_kind": "permutation", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "BN/permuted" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/padded_right", - "transform_kind": "padding", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "BN/padded_right" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/padded_left", - "transform_kind": "padding", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "BN/padded_left" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - } - ], - "logprob_smoke": null, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - }, - "candidate_id": "rlkernel.triton.qwenimage_multi_axis_rope.v1", - "device": "cuda", - "compute_capability": "sm120", - "seed": 20260812, - "fallback_reason": null, - "passed": true, - "provenance_valid": true, - "metadata_valid": true, - "observed_kernel_id": "rl_engine.backends.shared.triton.rope._flash_attn_rotary.rotary_kernel" - }, - "gradient": { - "op_name": "multi_axis_rope", - "backend_profile": "triton_cuda_bf16", - "accuracy_reports": [ - { - "config_id": "BN/full", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.699562850873917e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/chunked", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/chunked", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.699562850873917e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s0", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/full/s0", - "fp32_reference" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 2.7887988835573196e-05, - "mean_abs_error": 3.6821968478761846e-06, - "max_rel_error": 1.8164952993392944, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s1", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/full/s1", - "fp32_reference" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.7113780965446495e-06, - "max_rel_error": 0.4316565990447998, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s2", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/full/s2", - "fp32_reference" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 2.995925024151802e-05, - "mean_abs_error": 4.7168209675874095e-06, - "max_rel_error": 0.7439075112342834, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s3", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/full/s3", - "fp32_reference" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.0211638659238815e-05, - "mean_abs_error": 4.9605891945248e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s0", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/chunked/s0", - "fp32_reference" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 2.7887988835573196e-05, - "mean_abs_error": 3.6821968478761846e-06, - "max_rel_error": 1.8164952993392944, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s1", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/chunked/s1", - "fp32_reference" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.7113780965446495e-06, - "max_rel_error": 0.4316565990447998, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s2", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/chunked/s2", - "fp32_reference" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 2.995925024151802e-05, - "mean_abs_error": 4.7168209675874095e-06, - "max_rel_error": 0.7439075112342834, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s3", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/chunked/s3", - "fp32_reference" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.0211638659238815e-05, - "mean_abs_error": 4.9605891945248e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/permuted", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/permuted", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.699562850873917e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/padded_right", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/padded_right", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.699562850873917e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/padded_left", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/padded_left", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.699562850873917e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - } - ], - "invariance_reports": [ - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/chunked", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "BN/chunked" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s0", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s0" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s1", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s1" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s2", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s2" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s3", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s3" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s0", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s0" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s1", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s1" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s2", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s2" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s3", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s3" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/permuted", - "transform_kind": "permutation", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "BN/permuted" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/padded_right", - "transform_kind": "padding", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "BN/padded_right" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/padded_left", - "transform_kind": "padding", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "BN/padded_left" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - } - ], - "singleton_aggregate_reports": [], - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - }, - "candidate_id": "rlkernel.triton.qwenimage_multi_axis_rope.v1", - "device": "cuda", - "compute_capability": "sm120", - "seed": 20260812, - "fallback_reason": null, - "passed": true, - "provenance_valid": true, - "metadata_valid": true, - "loss_reduction": "sum_over_active_tokens_then_optional_mean_by_active_count", - "active_token_denominator": 27, - "grad_tensor_names": [ - "dx" - ], - "first_failing_op": null, - "first_failing_tensor": null, - "first_failing_config_pair": null, - "observed_kernel_id": "rl_engine.backends.shared.triton.rope._flash_attn_rotary.rotary_kernel" - } -} diff --git a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-triton.txt b/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-triton.txt deleted file mode 100644 index 283d1933a..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_v2_sm120/ws1-triton.txt +++ /dev/null @@ -1,3 +0,0 @@ -WARNING 10-28 16:54:47 [RL-Kernel]: Core binary extension (_C) unavailable: cannot import name '_C' from 'rl_engine' (D:\dandan123\dandan\issue\investigation-rl-kernel\RL-Kernel\rl_engine\__init__.py). Falling back to native code. -Using the installed production RoPE dispatcher library -WS1 four judgments backend=triton: passed=True diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/README.md b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/README.md deleted file mode 100644 index 9ef00fdb1..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/README.md +++ /dev/null @@ -1,136 +0,0 @@ -# Qwen-Image RoPE: full native package qualification on WSL2 - -Recorded 2026-10-28 for [PR #623](https://github.com/RL-Align/RL-Kernel/pull/623). -Numerical/native sources are based on `ac17ab87e8d93dabe93dfdc134e9853d47de7af1`; -this follow-up adds the reproducible gate, CPU CI routing and package provenance. -Source fingerprints are in `manifest.json`. Historical Windows measurements stay -in the separate `qwenimage_multi_axis_rope_v2_sm120` directory. - -## Environment and build - -| Item | Recorded value | -| --- | --- | -| Linux distribution | Ubuntu 24.04.5 LTS, x86_64, glibc 2.39 | -| Virtualization | WSL 3.0.3 on Windows 11 | -| Linux kernel | 6.18.54.1-microsoft-standard-WSL2 | -| GPU / driver | RTX 5080, SM120, 16 GB / Windows driver 617.42 | -| Python / PyTorch / Triton | 3.12.3 / 2.9.1+cu128 / 3.5.1 | -| CUDA compiler / host compiler | CUDA 12.8.93 / GCC 13.3.0 | -| Compute Sanitizer | 2025.1.0.0, CUDA 12.8 | -| Build profile | Ordinary complete `rl_engine._C`; 10 source translation units | -| CUDA architecture / parallelism | `TORCH_CUDA_ARCH_LIST=12.0`, `MAX_JOBS=2` | - -CUDA compiler, development libraries and command-line tools were installed from -NVIDIA's WSL-Ubuntu repository; the existing Windows GPU driver provides CUDA -driver access. The WSL driver directory `/usr/lib/wsl/lib` was added to -`LIBRARY_PATH` for linking the package's existing CUDA driver API calls. -`RL_KERNEL_ROPE_ONLY`, forced SM90 and fast-math flags were unset. - -The normal PEP 517 wheel build compiled all ten configured translation units, -including the existing logprob/GEMM/norm/activation/attention/collective code and -the new RoPE source. The complete `_C` loaded and registered the native RoPE -dispatcher. `full-extension.json` records sources, flags, binary path and hash; -`package.json` records the built wheel and installed library fingerprints. - -The checkout, virtual environments, build outputs and caches are inside the -Linux filesystem, backed by the D: virtual disk. A fresh Git clone preserved all -27 tracked symlinks. This is a WSL2 Linux userspace/compiler qualification on -SM120; bare-metal Linux and A100/H100 remain separate hardware profiles. - -## Results - -| Check | Result | Evidence | -| --- | --- | --- | -| Ordinary complete `_C` wheel and in-place build/import | Passed | `full-extension.json`, `full-wheel-build.txt`, `package.json` | -| Operator accuracy/invariance/autograd/graphs/opcheck/fullgraph | 81 passed, 1 two-GPU guard skip | `tests-and-ws1.txt`, `operator-tests.xml` | -| CUDA and Triton WS1 four judgments | Both passed with explicit provenance | `ws1-cuda.json`, `ws1-triton.json` | -| Standard operator CLI forward/input gradient | Both passed | `operator-cuda.json`, `operator-triton.json` | -| Related regressions | 195 passed, 7 Ascend skips | `regression.txt` | -| Repository layout with extension enforcement enabled | 37 passed, zero skips/failures | `layout.txt` | -| Fresh installed full wheel, outside source checkout | Eager/compiled forward and input-gradient bytes exact for both backends | `installed-wheel.json` | -| Actual CPU-only PyTorch wheel/reference installation, Python 3.10.22 | 57 passed, 52 expected GPU skips; reference forward/backward passed | `cpu310-contracts.txt`, `cpu310-environment.json` | -| Actual CPU-only PyTorch wheel/reference installation, Python 3.12.3 | 57 passed, 52 expected GPU skips; reference forward/backward passed | `cpu-contracts.txt`, `cpu-environment.json` | -| Compute Sanitizer accuracy/stride/empty/stream subset | 34 passed; zero errors | `memcheck.txt` | -| Pinned actual Diffusers symbols | 30 forward/backward cases passed; image/text tables exact | `upstream.json` | -| Reference-grid benchmark correctness | 40 comparisons passed | `benchmark-bfloat16.json` | -| Shell/workflow syntax, changed-file checks, strict docs | Passed locally | `quality.txt` | - -The full-extension GPU gate rejects missing native/Triton execution. Its only -allowed operator-test skip is the two-device guard on a one-GPU runner. -The CPU-only results are separate: both Python versions use `torch==2.9.1+cpu`, -and their installed packages contain no native CUDA binary. CPU skips are not -GPU acceptance evidence. Both runs install the normal CPU package and exercise -the workflow's three test files locally; these are not hosted CI results. -The 37 passing Linux layout tests contrast with the five Windows baseline -failures recorded earlier; those historical logs were not rewritten. - -The installed-wheel probe starts in a fresh process outside the checkout, -loads `_C` from site-packages and reports that exact path. It tests packed-QKV -forward/input-gradient equality in eager and `torch.compile(fullgraph=True)` -for both CUDA and Triton. The FlashAttention BSD license is present in the wheel. - -## Standalone benchmark - -BF16, batch 1, 24 heads, head dimension 128. CUDA Graph medians below are -**forward / forward+backward milliseconds per call**: - -| Grid / layout | Native CUDA | Adapted Triton | FlashAttention, no FMA | Diffusers complex | -| --- | --- | --- | --- | --- | -| 64x64 contiguous | 0.05780 / 0.12967 | 0.04929 / 0.12414 | 0.04898 / 0.12407 | 0.24446 / 0.51281 | -| 64x64 packed QKV | 0.05868 / 0.13079 | 0.04731 / 0.11841 | 0.04722 / 0.11847 | 0.24754 / 0.51583 | -| 83x83 contiguous | 0.11198 / 0.22392 | 0.10720 / 0.22538 | 0.10739 / 0.21420 | 0.52579 / 1.06721 | -| 83x83 packed QKV | 0.11233 / 0.22486 | 0.10745 / 0.21478 | 0.10744 / 0.21469 | 0.52820 / 1.06947 | -| 104x58 contiguous | 0.09797 / 0.19708 | 0.09366 / 0.18497 | 0.09362 / 0.18458 | 0.44992 / 0.90110 | -| 104x58 packed QKV | 0.09610 / 0.19612 | 0.09310 / 0.18828 | 0.09323 / 0.18817 | 0.44979 / 0.90276 | - -These are 20 calls per graph, five warm calls/replays and 30 replay measurements. -Three synchronized eager-wall trials of 30 calls are also in the JSON. Tables, -compilation and input/backing allocations are outside timing. Extra peak includes -the returned output or gradient. The short synthetic case is included in raw data. -At 4096 tokens, both GPU paths' forward extra peak is 24 MiB for either layout, -the returned BF16 output size. Packed inputs are read without materialization. - -Original FlashAttention and its autograd wrapper are checksum-pinned to -`98c1cee0ed59307833841cdbf567fe336ceec13a`. The no-FMA baseline changes only the -launch contraction option. CUDA, adapted Triton and that baseline have zero -differing forward/input-gradient bytes against the unfused FP32 gold in every -benchmark case. Default FlashAttention and complex arithmetic use declared tolerances. -The adapted Triton path is close to the mature baseline; no performance advantage -over FlashAttention is established. Timings are one-process measurements on a -shared Windows/WSL GPU without clock pinning, not a model-throughput result or -a controlled comparison against the earlier Windows process. - -## Reproduction and CI - -Use the matching CUDA 12.8 toolkit and Linux PyTorch 2.9.1+cu128 in a fresh -Linux virtual environment. Other Python package versions are recorded in -`python-packages.json`. From the Linux checkout: - -```bash -export CUDA_HOME=/usr/local/cuda-12.8 -export PATH="$CUDA_HOME/bin:$PATH" -export LIBRARY_PATH="/usr/lib/wsl/lib${LIBRARY_PATH:+:$LIBRARY_PATH}" -export MAX_JOBS=2 TORCH_CUDA_ARCH_LIST=12.0 RL_KERNEL_REQUIRE_EXT=1 -unset RL_KERNEL_ROPE_ONLY KERNEL_ALIGN_FORCE_SM90 KERNEL_ALIGN_USE_FAST_MATH -python -m pip wheel --no-build-isolation --no-deps . -w /tmp/rope-full-wheel -python -m pip install --force-reinstall --no-deps /tmp/rope-full-wheel/rl_kernel-*.whl -python setup.py build_ext --inplace -ROPE_PYTHON=python ROPE_WS1_OUTPUT=/tmp/rope-ws1 bash ci/scripts/run_qwenimage_rope_ws1.sh -python ci/run.py layout -q -rs -compute-sanitizer --tool memcheck --error-exitcode 1 python -m pytest \ - tests/ops/rope/test_multi_axis_rope.py -q \ - -k "forward_backward_accuracy or noncontiguous_input_tables_and_gradient or empty_forward_backward or packed_qkv_and_independent_table_strides or native_entrypoint_checks" -``` - -The new [operator gate](../../../ci/scripts/run_qwenimage_rope_ws1.sh) requires -the complete `_C` registration, records actual library/backend identity and -enforces the GPU skip policy. The [CPU workflow](../../../.github/workflows/qwenimage-rope.yml) -targets `test-qwenimage`, runs Python 3.10/3.12 CPU contracts, checks changed files -and builds documentation. This report includes local CPU-only runs on both -Python 3.10 and 3.12; hosted CI status is tracked on the PR. GPU runners use the -explicit gate. - -Bare-metal Linux, A100/H100, two-GPU guards, QK-norm fusion, model/block assembly -and trajectories remain unqualified. Building the entire library does not -certify every native operator. No Dense Qwen3/distributed certification or -cross-device/version byte-equality claim is made. diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/benchmark-bfloat16.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/benchmark-bfloat16.json deleted file mode 100644 index ca80105b0..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/benchmark-bfloat16.json +++ /dev/null @@ -1,2781 +0,0 @@ -{ - "gpu": "NVIDIA GeForce RTX 5080", - "capability": [ - 12, - 0 - ], - "os": "Linux-6.18.54.1-microsoft-standard-WSL2-x86_64-with-glibc2.39", - "python": "3.12.3 (main, Aug 31 2026, 10:18:26) [GCC 13.3.0]", - "torch": "2.9.1+cu128", - "cuda_runtime": "12.8", - "triton": "3.5.1", - "dtype": "bfloat16", - "base_commit": "ac17ab87e8d93dabe93dfdc134e9853d47de7af1", - "source_sha256": { - "csrc/cuda/rope/multi_axis_rope.cu": "f5db5cf5d4962053194455d9cde8cf91cf4e824bad7abf72e8410a404bcb0e44", - "rl_engine/backends/shared/triton/rope/_flash_attn_rotary.py": "cf0eebf315830a4178053a71a77e650b58aa85bba1d1a08aff7677350af755d9", - "rl_engine/backends/shared/triton/rope/multi_axis_rope.py": "f08ff90627c66e15e9c558f8a68f132434a741a76997b3800bfbf73dc8940e9e", - "rl_engine/backends/cuda/rope/multi_axis_rope.py": "d779d9b60ce5d1d6651925f723049930d898024fdf0548ce29a2019b84753d65", - "rl_engine/ops/rope/_registration.py": "add0375d1639444908091d01e8d29781f7a975beffec64b77bfaa81a5b9a0ae3", - "rl_engine/ops/rope/multi_axis_rope.py": "07bcff1160f22c1ff6730be731206b2f4a153e2fada26afc93fc884e780568cb", - "rl_engine/reference/rope/multi_axis_rope.py": "c9a94730047f51659f7aa0a443e88d6ce7d03d00d6813a9aed548cafbf6347bf", - "benchmarks/operators/rope/benchmark_multi_axis_rope.py": "6961a55803fbd4e566bd060558fb59e658e7f1f704e61ca97f912e5d16b55884" - }, - "timing": "CUDA graph ms/call and synchronized eager wall ms/call; warm allocation; excludes tables, compilation, and initial allocation of inputs", - "memory": "extra peak allocated bytes excluding inputs/tables/grad; includes returned output or gradient; packed QKV backing storage allocated before timing", - "iterations": 30, - "warmup": 5, - "trials": 3, - "graph_calls": 20, - "flash_baseline": { - "included": true, - "revision": "98c1cee0ed59307833841cdbf567fe336ceec13a", - "kernel_sha256_normalized": "d2037e623854c702fb075784fe4292d4a89d6b9ac1ae33cab3d6788f8f94f240", - "layer_sha256_normalized": "a79cb42dfc33e126227bdba27ae7a024d7d5cc5c8fd7f7a5f26c68d501024e0a", - "mode": "upstream rotary_kernel, apply_rotary and ApplyRotaryEmb, interleaved=True", - "no_fma_change": "only add enable_fp_fusion=False at the original kernel launch" - }, - "trace": { - "cuda": { - "contract": "qwenimage-interleaved-rope-fp32-v1", - "actual_backend": "rlkernel.cuda.qwenimage_multi_axis_rope.v1", - "kernel_id": "torch.ops.rl_kernel_native.multi_axis_rope", - "layout": "BSHD, adjacent channel pairs", - "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", - "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" - }, - "triton": { - "contract": "qwenimage-interleaved-rope-fp32-v1", - "actual_backend": "rlkernel.triton.qwenimage_multi_axis_rope.v1", - "kernel_id": "rl_engine.backends.shared.triton.rope._flash_attn_rotary.rotary_kernel", - "layout": "BSHD, adjacent channel pairs", - "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", - "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" - } - }, - "qualification": "Exact bytes vs unfused FP32 gold for CUDA, Triton and flash_no_fma; declared tolerance for upstream complex and default FlashAttention", - "correctness": [ - { - "backend": "upstream_complex", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "cuda", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "triton", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "flash_default", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "flash_no_fma", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "upstream_complex", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "cuda", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "triton", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "flash_default", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "flash_no_fma", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "upstream_complex", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "different_bytes_y_dx": [ - 130, - 139 - ] - }, - { - "backend": "cuda", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "triton", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "flash_default", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "different_bytes_y_dx": [ - 130, - 139 - ] - }, - { - "backend": "flash_no_fma", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "upstream_complex", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "different_bytes_y_dx": [ - 141, - 139 - ] - }, - { - "backend": "cuda", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "triton", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "flash_default", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "different_bytes_y_dx": [ - 141, - 139 - ] - }, - { - "backend": "flash_no_fma", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "upstream_complex", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "different_bytes_y_dx": [ - 235, - 224 - ] - }, - { - "backend": "cuda", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "triton", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "flash_default", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "different_bytes_y_dx": [ - 235, - 224 - ] - }, - { - "backend": "flash_no_fma", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "upstream_complex", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "different_bytes_y_dx": [ - 221, - 226 - ] - }, - { - "backend": "cuda", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "triton", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "flash_default", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "different_bytes_y_dx": [ - 221, - 226 - ] - }, - { - "backend": "flash_no_fma", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "upstream_complex", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "different_bytes_y_dx": [ - 195, - 199 - ] - }, - { - "backend": "cuda", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "triton", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "flash_default", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "different_bytes_y_dx": [ - 195, - 199 - ] - }, - { - "backend": "flash_no_fma", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "upstream_complex", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "different_bytes_y_dx": [ - 230, - 220 - ] - }, - { - "backend": "cuda", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "triton", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - }, - { - "backend": "flash_default", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "different_bytes_y_dx": [ - 230, - 220 - ] - }, - { - "backend": "flash_no_fma", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "different_bytes_y_dx": [ - 0, - 0 - ] - } - ], - "results": [ - { - "backend": "upstream_complex", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 18432, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.0034856000915169717, - "graph_min_ms": 0.0034543998539447785, - "graph_max_ms": 0.0035312000662088396, - "eager_wall_median_ms": 0.08385686669498682, - "eager_wall_trials_ms": [ - 0.09771363335554877, - 0.08163450002030004, - 0.08385686669498682 - ], - "extra_peak_bytes": 184320 - }, - { - "backend": "upstream_complex", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 18432, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.007651200145483017, - "graph_min_ms": 0.007601600140333176, - "graph_max_ms": 0.007692799717187881, - "eager_wall_median_ms": 0.19571573332844613, - "eager_wall_trials_ms": [ - 0.2535254000273805, - 0.19571573332844613, - 0.18590186670432254 - ], - "extra_peak_bytes": 187392 - }, - { - "backend": "cuda", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 18432, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.0012047999538481236, - "graph_min_ms": 0.0011695999652147292, - "graph_max_ms": 0.0012400000356137752, - "eager_wall_median_ms": 0.05762376664885475, - "eager_wall_trials_ms": [ - 0.07130193328824437, - 0.05762376664885475, - 0.057149166665719044 - ], - "extra_peak_bytes": 36864 - }, - { - "backend": "cuda", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 18432, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.002329600043594837, - "graph_min_ms": 0.002223999984562397, - "graph_max_ms": 0.0026383999735116958, - "eager_wall_median_ms": 0.17063139997238372, - "eager_wall_trials_ms": [ - 0.17063139997238372, - 0.16431806670880178, - 0.24000986668397672 - ], - "extra_peak_bytes": 73728 - }, - { - "backend": "triton", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 18432, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.0013432000298053027, - "graph_min_ms": 0.0010928000323474408, - "graph_max_ms": 0.006166400015354156, - "eager_wall_median_ms": 0.08490003328915918, - "eager_wall_trials_ms": [ - 0.08490003328915918, - 0.0852208666401566, - 0.08217023332690587 - ], - "extra_peak_bytes": 36864 - }, - { - "backend": "triton", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 18432, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.002125599980354309, - "graph_min_ms": 0.0020735999569296837, - "graph_max_ms": 0.0021679999306797983, - "eager_wall_median_ms": 0.3603733666144156, - "eager_wall_trials_ms": [ - 0.3862667332820517, - 0.31255590001819655, - 0.3603733666144156 - ], - "extra_peak_bytes": 73728 - }, - { - "backend": "flash_default", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 18432, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.0011024000123143196, - "graph_min_ms": 0.0010544000193476676, - "graph_max_ms": 0.01971679925918579, - "eager_wall_median_ms": 0.07505266664035541, - "eager_wall_trials_ms": [ - 0.09395336667997374, - 0.07483469998987857, - 0.07505266664035541 - ], - "extra_peak_bytes": 36864 - }, - { - "backend": "flash_default", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 18432, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.0020584000274538996, - "graph_min_ms": 0.002022399939596653, - "graph_max_ms": 0.002151999995112419, - "eager_wall_median_ms": 0.3825640333161573, - "eager_wall_trials_ms": [ - 0.3916477333405055, - 0.38042896667320747, - 0.3825640333161573 - ], - "extra_peak_bytes": 73728 - }, - { - "backend": "flash_no_fma", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 18432, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.0011288000270724297, - "graph_min_ms": 0.0010703999549150466, - "graph_max_ms": 0.0025040000677108766, - "eager_wall_median_ms": 0.07896646663236122, - "eager_wall_trials_ms": [ - 0.07896646663236122, - 0.08368196668016026, - 0.07650009999148703 - ], - "extra_peak_bytes": 36864 - }, - { - "backend": "flash_no_fma", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 18432, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.002057600021362305, - "graph_min_ms": 0.002022399939596653, - "graph_max_ms": 0.0021487999707460403, - "eager_wall_median_ms": 0.3408336999806731, - "eager_wall_trials_ms": [ - 0.3408336999806731, - 0.32808933331883355, - 0.35737479999321903 - ], - "extra_peak_bytes": 73728 - }, - { - "backend": "upstream_complex", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 55296, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.0037903999909758568, - "graph_min_ms": 0.003734400123357773, - "graph_max_ms": 0.0038384001702070234, - "eager_wall_median_ms": 0.026099166704322368, - "eager_wall_trials_ms": [ - 0.027518266688275617, - 0.02535739998468974, - 0.026099166704322368 - ], - "extra_peak_bytes": 184320 - }, - { - "backend": "upstream_complex", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 55296, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.008167200163006781, - "graph_min_ms": 0.00809599980711937, - "graph_max_ms": 0.008207999914884568, - "eager_wall_median_ms": 0.22357730000900725, - "eager_wall_trials_ms": [ - 0.255850233346185, - 0.22357730000900725, - 0.2018136666568656 - ], - "extra_peak_bytes": 187392 - }, - { - "backend": "cuda", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 55296, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.001232000021263957, - "graph_min_ms": 0.0011552000418305398, - "graph_max_ms": 0.0026464000344276427, - "eager_wall_median_ms": 0.05505463335187718, - "eager_wall_trials_ms": [ - 0.05505463335187718, - 0.052852566659566946, - 0.05531913332864254 - ], - "extra_peak_bytes": 36864 - }, - { - "backend": "cuda", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 55296, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.002279200032353401, - "graph_min_ms": 0.00220320001244545, - "graph_max_ms": 0.005976000055670738, - "eager_wall_median_ms": 0.2586098333267728, - "eager_wall_trials_ms": [ - 0.2587837333218583, - 0.25146820001585485, - 0.2586098333267728 - ], - "extra_peak_bytes": 73728 - }, - { - "backend": "triton", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 55296, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.0011359999887645244, - "graph_min_ms": 0.001071999967098236, - "graph_max_ms": 0.002603200078010559, - "eager_wall_median_ms": 0.03804060003555302, - "eager_wall_trials_ms": [ - 0.03804060003555302, - 0.04575726664673615, - 0.03741419999035619 - ], - "extra_peak_bytes": 36864 - }, - { - "backend": "triton", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 55296, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.0021271999925374986, - "graph_min_ms": 0.0020896000787615776, - "graph_max_ms": 0.016212800145149232, - "eager_wall_median_ms": 0.24443720000514682, - "eager_wall_trials_ms": [ - 0.281011633342132, - 0.24443720000514682, - 0.23313996668245332 - ], - "extra_peak_bytes": 73728 - }, - { - "backend": "flash_default", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 55296, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.001104000024497509, - "graph_min_ms": 0.001065600011497736, - "graph_max_ms": 0.0011983999982476234, - "eager_wall_median_ms": 0.078476833308135, - "eager_wall_trials_ms": [ - 0.07911106664929927, - 0.07519836666081876, - 0.078476833308135 - ], - "extra_peak_bytes": 36864 - }, - { - "backend": "flash_default", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 55296, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.0020527999848127367, - "graph_min_ms": 0.0020207999274134635, - "graph_max_ms": 0.0021023999899625776, - "eager_wall_median_ms": 0.30830893338134047, - "eager_wall_trials_ms": [ - 0.3673079666744646, - 0.29416493331761256, - 0.30830893338134047 - ], - "extra_peak_bytes": 73728 - }, - { - "backend": "flash_no_fma", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 55296, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.0011272000148892403, - "graph_min_ms": 0.0010511999949812888, - "graph_max_ms": 0.002239999920129776, - "eager_wall_median_ms": 0.07592803334167304, - "eager_wall_trials_ms": [ - 0.0775905666766145, - 0.07592803334167304, - 0.0742380666755101 - ], - "extra_peak_bytes": 36864 - }, - { - "backend": "flash_no_fma", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6, - 24, - 128 - ], - "input_stride": [ - 55296, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.0020535999909043314, - "graph_min_ms": 0.002014400064945221, - "graph_max_ms": 0.0020975999534130095, - "eager_wall_median_ms": 0.31871896662778454, - "eager_wall_trials_ms": [ - 0.31871896662778454, - 0.33167873334605247, - 0.2915063333299865 - ], - "extra_peak_bytes": 73728 - }, - { - "backend": "upstream_complex", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 12582912, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.24445840120315554, - "graph_min_ms": 0.24224960803985596, - "graph_max_ms": 0.29114880561828616, - "eager_wall_median_ms": 0.247559633377629, - "eager_wall_trials_ms": [ - 0.25793806665509084, - 0.2471523333345734, - 0.247559633377629 - ], - "extra_peak_bytes": 125829120 - }, - { - "backend": "upstream_complex", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 12582912, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.5128072023391723, - "graph_min_ms": 0.49546399116516116, - "graph_max_ms": 0.5571055889129639, - "eager_wall_median_ms": 0.5245368333514003, - "eager_wall_trials_ms": [ - 0.5245368333514003, - 0.518849266639639, - 0.5350402333230401 - ], - "extra_peak_bytes": 127926272 - }, - { - "backend": "cuda", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 12582912, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.05780480206012726, - "graph_min_ms": 0.057633602619171144, - "graph_max_ms": 0.07439360022544861, - "eager_wall_median_ms": 0.06806933330760027, - "eager_wall_trials_ms": [ - 0.07017140002668991, - 0.06806933330760027, - 0.06168783335548748 - ], - "extra_peak_bytes": 25165824 - }, - { - "backend": "cuda", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 12582912, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.1296720027923584, - "graph_min_ms": 0.12902239561080933, - "graph_max_ms": 0.15688639879226685, - "eager_wall_median_ms": 0.2333034999537631, - "eager_wall_trials_ms": [ - 0.23402639999403618, - 0.19631726663646987, - 0.2333034999537631 - ], - "extra_peak_bytes": 50331648 - }, - { - "backend": "triton", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 12582912, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.049288800358772276, - "graph_min_ms": 0.04815520048141479, - "graph_max_ms": 0.06759200096130372, - "eager_wall_median_ms": 0.08061226665934858, - "eager_wall_trials_ms": [ - 0.07100329997532147, - 0.0930295666573026, - 0.08061226665934858 - ], - "extra_peak_bytes": 25165824 - }, - { - "backend": "triton", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 12582912, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.12414479851722718, - "graph_min_ms": 0.12152320146560669, - "graph_max_ms": 0.16048799753189086, - "eager_wall_median_ms": 0.35214599999259616, - "eager_wall_trials_ms": [ - 0.35569400000288925, - 0.35214599999259616, - 0.32247826663175755 - ], - "extra_peak_bytes": 50331648 - }, - { - "backend": "flash_default", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 12582912, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.04924079924821853, - "graph_min_ms": 0.04797120094299316, - "graph_max_ms": 0.068995201587677, - "eager_wall_median_ms": 0.13288023334704727, - "eager_wall_trials_ms": [ - 0.14359453337723002, - 0.12000013339275029, - 0.13288023334704727 - ], - "extra_peak_bytes": 25165824 - }, - { - "backend": "flash_default", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 12582912, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.12360240221023559, - "graph_min_ms": 0.1215999960899353, - "graph_max_ms": 0.16038399934768677, - "eager_wall_median_ms": 0.35429976663484314, - "eager_wall_trials_ms": [ - 0.34954116666388774, - 0.3772011999899405, - 0.35429976663484314 - ], - "extra_peak_bytes": 50331648 - }, - { - "backend": "flash_no_fma", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 12582912, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.0489767998456955, - "graph_min_ms": 0.0479856014251709, - "graph_max_ms": 0.06645600199699402, - "eager_wall_median_ms": 0.1330274333061728, - "eager_wall_trials_ms": [ - 0.1330274333061728, - 0.1200082000044252, - 0.13366413331823424 - ], - "extra_peak_bytes": 25165824 - }, - { - "backend": "flash_no_fma", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 12582912, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.12407039999961852, - "graph_min_ms": 0.12193280458450317, - "graph_max_ms": 0.1605712056159973, - "eager_wall_median_ms": 0.36529683329717955, - "eager_wall_trials_ms": [ - 0.3895335000076254, - 0.36529683329717955, - 0.3503133667133322 - ], - "extra_peak_bytes": 50331648 - }, - { - "backend": "upstream_complex", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 37748736, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.2475391983985901, - "graph_min_ms": 0.24529919624328614, - "graph_max_ms": 0.27714240550994873, - "eager_wall_median_ms": 0.25804310001452296, - "eager_wall_trials_ms": [ - 0.25804310001452296, - 0.26510979999632883, - 0.2503190333300154 - ], - "extra_peak_bytes": 125829120 - }, - { - "backend": "upstream_complex", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 37748736, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.5158272027969361, - "graph_min_ms": 0.49741439819335936, - "graph_max_ms": 0.5595920085906982, - "eager_wall_median_ms": 0.5221092666639985, - "eager_wall_trials_ms": [ - 0.5362037999778598, - 0.5221092666639985, - 0.5111432000073061 - ], - "extra_peak_bytes": 127926272 - }, - { - "backend": "cuda", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 37748736, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.05868320167064667, - "graph_min_ms": 0.058520001173019406, - "graph_max_ms": 0.0754975974559784, - "eager_wall_median_ms": 0.06486223331497361, - "eager_wall_trials_ms": [ - 0.06433113333817649, - 0.08489736665069358, - 0.06486223331497361 - ], - "extra_peak_bytes": 25165824 - }, - { - "backend": "cuda", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 37748736, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.13079119920730592, - "graph_min_ms": 0.13011519908905028, - "graph_max_ms": 0.1711359977722168, - "eager_wall_median_ms": 0.21280260001124893, - "eager_wall_trials_ms": [ - 0.2440493333779159, - 0.20012976668416135, - 0.21280260001124893 - ], - "extra_peak_bytes": 50331648 - }, - { - "backend": "triton", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 37748736, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.04731200039386749, - "graph_min_ms": 0.04663040041923523, - "graph_max_ms": 0.06482719779014587, - "eager_wall_median_ms": 0.08378989999376547, - "eager_wall_trials_ms": [ - 0.08378989999376547, - 0.07920856666411662, - 0.09377416663483018 - ], - "extra_peak_bytes": 25165824 - }, - { - "backend": "triton", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 37748736, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.11841280460357666, - "graph_min_ms": 0.11643359661102295, - "graph_max_ms": 0.14446239471435546, - "eager_wall_median_ms": 0.3498463667104564, - "eager_wall_trials_ms": [ - 0.3498463667104564, - 0.35651856663510745, - 0.3451830999741408 - ], - "extra_peak_bytes": 50331648 - }, - { - "backend": "flash_default", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 37748736, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.047358399629592894, - "graph_min_ms": 0.04637599885463715, - "graph_max_ms": 0.06467199921607972, - "eager_wall_median_ms": 0.12952533331069085, - "eager_wall_trials_ms": [ - 0.12952533331069085, - 0.1346659000167468, - 0.12932516668418734 - ], - "extra_peak_bytes": 25165824 - }, - { - "backend": "flash_default", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 37748736, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.11804000139236451, - "graph_min_ms": 0.11632959842681885, - "graph_max_ms": 0.13619199991226197, - "eager_wall_median_ms": 0.3809861666619933, - "eager_wall_trials_ms": [ - 0.3809861666619933, - 0.39690620002753957, - 0.3660936666468236 - ], - "extra_peak_bytes": 50331648 - }, - { - "backend": "flash_no_fma", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 37748736, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.04721840023994446, - "graph_min_ms": 0.04616160094738007, - "graph_max_ms": 0.06441760063171387, - "eager_wall_median_ms": 0.1290158666355031, - "eager_wall_trials_ms": [ - 0.13406613331123177, - 0.1290158666355031, - 0.12222770001244498 - ], - "extra_peak_bytes": 25165824 - }, - { - "backend": "flash_no_fma", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 4096, - 24, - 128 - ], - "input_stride": [ - 37748736, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.11847440004348755, - "graph_min_ms": 0.11651359796524048, - "graph_max_ms": 0.1468176007270813, - "eager_wall_median_ms": 0.4179154666417162, - "eager_wall_trials_ms": [ - 0.4434371333142432, - 0.4179154666417162, - 0.4101795332947707 - ], - "extra_peak_bytes": 50331648 - }, - { - "backend": "upstream_complex", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 21163008, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.5257872104644776, - "graph_min_ms": 0.5080063819885254, - "graph_max_ms": 0.5525023937225342, - "eager_wall_median_ms": 0.5246295667044857, - "eager_wall_trials_ms": [ - 0.5097412666752158, - 0.528316466625256, - 0.5246295667044857 - ], - "extra_peak_bytes": 211630080 - }, - { - "backend": "upstream_complex", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 21163008, - 3072, - 128, - 1 - ], - "graph_median_ms": 1.067212772369385, - "graph_min_ms": 1.0389776229858398, - "graph_max_ms": 1.0976528167724608, - "eager_wall_median_ms": 1.0610819333426964, - "eager_wall_trials_ms": [ - 1.0702722999970622, - 1.0479776333644015, - 1.0610819333426964 - ], - "extra_peak_bytes": 215157248 - }, - { - "backend": "cuda", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 21163008, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.11197760105133056, - "graph_min_ms": 0.11102559566497802, - "graph_max_ms": 0.13988800048828126, - "eager_wall_median_ms": 0.12250546666715915, - "eager_wall_trials_ms": [ - 0.11638536664880424, - 0.13100786666958206, - 0.12250546666715915 - ], - "extra_peak_bytes": 42326016 - }, - { - "backend": "cuda", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 21163008, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.22392239570617678, - "graph_min_ms": 0.22292959690093994, - "graph_max_ms": 0.26030399799346926, - "eager_wall_median_ms": 0.24186149997452353, - "eager_wall_trials_ms": [ - 0.24518829995940905, - 0.23021603331775017, - 0.24186149997452353 - ], - "extra_peak_bytes": 84652032 - }, - { - "backend": "triton", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 21163008, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.10719599723815917, - "graph_min_ms": 0.10527039766311645, - "graph_max_ms": 0.123745596408844, - "eager_wall_median_ms": 0.10772686667526916, - "eager_wall_trials_ms": [ - 0.10711883332987782, - 0.10772686667526916, - 0.10832119999880281 - ], - "extra_peak_bytes": 42326016 - }, - { - "backend": "triton", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 21163008, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.2253783941268921, - "graph_min_ms": 0.21142399311065674, - "graph_max_ms": 0.25633440017700193, - "eager_wall_median_ms": 0.29183916667534504, - "eager_wall_trials_ms": [ - 0.29183916667534504, - 0.31016709999676095, - 0.27477109997562366 - ], - "extra_peak_bytes": 84652032 - }, - { - "backend": "flash_default", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 21163008, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.10732319951057434, - "graph_min_ms": 0.10543199777603149, - "graph_max_ms": 0.12232480049133301, - "eager_wall_median_ms": 0.11885433329249888, - "eager_wall_trials_ms": [ - 0.11160320003303544, - 0.11885433329249888, - 0.148421633275575 - ], - "extra_peak_bytes": 42326016 - }, - { - "backend": "flash_default", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 21163008, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.22528079748153684, - "graph_min_ms": 0.21176319122314452, - "graph_max_ms": 0.24576001167297362, - "eager_wall_median_ms": 0.41978080001475365, - "eager_wall_trials_ms": [ - 0.4051902666712219, - 0.42735956667456776, - 0.41978080001475365 - ], - "extra_peak_bytes": 84652032 - }, - { - "backend": "flash_no_fma", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 21163008, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.10738959908485413, - "graph_min_ms": 0.10557440519332886, - "graph_max_ms": 0.1310464024543762, - "eager_wall_median_ms": 0.1100779333379857, - "eager_wall_trials_ms": [ - 0.11928496666466042, - 0.10885563333431492, - 0.1100779333379857 - ], - "extra_peak_bytes": 42326016 - }, - { - "backend": "flash_no_fma", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 21163008, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.21419600248336793, - "graph_min_ms": 0.2123471975326538, - "graph_max_ms": 0.24199841022491456, - "eager_wall_median_ms": 0.42007340004298993, - "eager_wall_trials_ms": [ - 0.4409272999813159, - 0.42007340004298993, - 0.40205200002674246 - ], - "extra_peak_bytes": 84652032 - }, - { - "backend": "upstream_complex", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 63489024, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.5282040119171143, - "graph_min_ms": 0.5110288143157959, - "graph_max_ms": 0.5885712146759033, - "eager_wall_median_ms": 0.547415499992591, - "eager_wall_trials_ms": [ - 0.547415499992591, - 0.5160773666527044, - 0.5478402666388623 - ], - "extra_peak_bytes": 211630080 - }, - { - "backend": "upstream_complex", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 63489024, - 9216, - 128, - 1 - ], - "graph_median_ms": 1.0694655895233156, - "graph_min_ms": 1.0273263931274415, - "graph_max_ms": 1.0971407890319824, - "eager_wall_median_ms": 1.0825770000034634, - "eager_wall_trials_ms": [ - 1.0599153666286536, - 1.0825770000034634, - 1.0843000333503976 - ], - "extra_peak_bytes": 215157248 - }, - { - "backend": "cuda", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 63489024, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.11232720017433166, - "graph_min_ms": 0.11165920495986939, - "graph_max_ms": 0.13819040060043336, - "eager_wall_median_ms": 0.12610396667393312, - "eager_wall_trials_ms": [ - 0.12610396667393312, - 0.11783170005704353, - 0.1361641000282058 - ], - "extra_peak_bytes": 42326016 - }, - { - "backend": "cuda", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 63489024, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.22485519647598268, - "graph_min_ms": 0.2238991975784302, - "graph_max_ms": 0.2728208065032959, - "eager_wall_median_ms": 0.2468356666213367, - "eager_wall_trials_ms": [ - 0.2468356666213367, - 0.25081796666199807, - 0.23804076669572774 - ], - "extra_peak_bytes": 84652032 - }, - { - "backend": "triton", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 63489024, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.10744640231132507, - "graph_min_ms": 0.10558079481124878, - "graph_max_ms": 0.12357280254364014, - "eager_wall_median_ms": 0.10917900005248764, - "eager_wall_trials_ms": [ - 0.108484466666899, - 0.11029286670236615, - 0.10917900005248764 - ], - "extra_peak_bytes": 42326016 - }, - { - "backend": "triton", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 63489024, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.21478239297866822, - "graph_min_ms": 0.21232800483703612, - "graph_max_ms": 0.2458656072616577, - "eager_wall_median_ms": 0.3064794333719571, - "eager_wall_trials_ms": [ - 0.3064794333719571, - 0.29039823330094805, - 0.3166810000038822 - ], - "extra_peak_bytes": 84652032 - }, - { - "backend": "flash_default", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 63489024, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.10760959982872009, - "graph_min_ms": 0.10598560571670532, - "graph_max_ms": 0.13379520177841187, - "eager_wall_median_ms": 0.11580663334219328, - "eager_wall_trials_ms": [ - 0.11580663334219328, - 0.12037279999882837, - 0.10920460002428929 - ], - "extra_peak_bytes": 42326016 - }, - { - "backend": "flash_default", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 63489024, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.2147576093673706, - "graph_min_ms": 0.21242880821228027, - "graph_max_ms": 0.2415616035461426, - "eager_wall_median_ms": 0.42939333337320323, - "eager_wall_trials_ms": [ - 0.4258981332895928, - 0.45155456670424127, - 0.42939333337320323 - ], - "extra_peak_bytes": 84652032 - }, - { - "backend": "flash_no_fma", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 63489024, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.1074351966381073, - "graph_min_ms": 0.10590399503707885, - "graph_max_ms": 0.13916159868240358, - "eager_wall_median_ms": 0.11420323335187277, - "eager_wall_trials_ms": [ - 0.11669516667704254, - 0.11420323335187277, - 0.11166136664542137 - ], - "extra_peak_bytes": 42326016 - }, - { - "backend": "flash_no_fma", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6889, - 24, - 128 - ], - "input_stride": [ - 63489024, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.21468960046768187, - "graph_min_ms": 0.21285600662231446, - "graph_max_ms": 0.24125280380249023, - "eager_wall_median_ms": 0.4036813999846345, - "eager_wall_trials_ms": [ - 0.40494703334843507, - 0.4027272333284297, - 0.4036813999846345 - ], - "extra_peak_bytes": 84652032 - }, - { - "backend": "upstream_complex", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 18530304, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.44991838932037354, - "graph_min_ms": 0.4312863826751709, - "graph_max_ms": 0.4923136234283447, - "eager_wall_median_ms": 0.4565672999888193, - "eager_wall_trials_ms": [ - 0.46419443333434174, - 0.4565672999888193, - 0.43665636664324364 - ], - "extra_peak_bytes": 185303040 - }, - { - "backend": "upstream_complex", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 18530304, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.9010960102081299, - "graph_min_ms": 0.8724479675292969, - "graph_max_ms": 0.9627167701721191, - "eager_wall_median_ms": 0.9057802333093908, - "eager_wall_trials_ms": [ - 0.9057802333093908, - 0.9399917666087276, - 0.8861856666650662 - ], - "extra_peak_bytes": 188391424 - }, - { - "backend": "cuda", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 18530304, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.09797199964523315, - "graph_min_ms": 0.0971664011478424, - "graph_max_ms": 0.12116639614105225, - "eager_wall_median_ms": 0.10149109997049284, - "eager_wall_trials_ms": [ - 0.10149109997049284, - 0.10145610000715048, - 0.10150603332779913 - ], - "extra_peak_bytes": 37060608 - }, - { - "backend": "cuda", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 18530304, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.19708160161972046, - "graph_min_ms": 0.19586080312728882, - "graph_max_ms": 0.2340928077697754, - "eager_wall_median_ms": 0.23985833334639514, - "eager_wall_trials_ms": [ - 0.2599421666673152, - 0.21093993333731, - 0.23985833334639514 - ], - "extra_peak_bytes": 74809344 - }, - { - "backend": "triton", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 18530304, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.09365599751472473, - "graph_min_ms": 0.09192960262298584, - "graph_max_ms": 0.1098512053489685, - "eager_wall_median_ms": 0.09574563337082509, - "eager_wall_trials_ms": [ - 0.10015153335795428, - 0.09574563337082509, - 0.09447829997952795 - ], - "extra_peak_bytes": 37060608 - }, - { - "backend": "triton", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 18530304, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.18496879935264587, - "graph_min_ms": 0.18224799633026123, - "graph_max_ms": 0.20439839363098145, - "eager_wall_median_ms": 0.27221376664480584, - "eager_wall_trials_ms": [ - 0.30083316663270426, - 0.2585961666530541, - 0.27221376664480584 - ], - "extra_peak_bytes": 74809344 - }, - { - "backend": "flash_default", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 18530304, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.09344720244407653, - "graph_min_ms": 0.09223840236663819, - "graph_max_ms": 0.10959680080413818, - "eager_wall_median_ms": 0.10214689997762132, - "eager_wall_trials_ms": [ - 0.10214689997762132, - 0.09693080003974804, - 0.10462310001457809 - ], - "extra_peak_bytes": 37060608 - }, - { - "backend": "flash_default", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 18530304, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.18451839685440063, - "graph_min_ms": 0.1817072033882141, - "graph_max_ms": 0.2279695987701416, - "eager_wall_median_ms": 0.38619293330460397, - "eager_wall_trials_ms": [ - 0.4084413000479496, - 0.38619293330460397, - 0.3539249666573596 - ], - "extra_peak_bytes": 74809344 - }, - { - "backend": "flash_no_fma", - "phase": "forward", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 18530304, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.09362159967422486, - "graph_min_ms": 0.09215999841690063, - "graph_max_ms": 0.1108240008354187, - "eager_wall_median_ms": 0.10443516669814319, - "eager_wall_trials_ms": [ - 0.10080796667656007, - 0.10632786670612404, - 0.10443516669814319 - ], - "extra_peak_bytes": 37060608 - }, - { - "backend": "flash_no_fma", - "phase": "forward_backward", - "layout": "contiguous", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 18530304, - 3072, - 128, - 1 - ], - "graph_median_ms": 0.18458240032196044, - "graph_min_ms": 0.18273439407348632, - "graph_max_ms": 0.21199359893798828, - "eager_wall_median_ms": 0.3521715333287527, - "eager_wall_trials_ms": [ - 0.3573932999643148, - 0.32707516663019004, - 0.3521715333287527 - ], - "extra_peak_bytes": 74809344 - }, - { - "backend": "upstream_complex", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 55590912, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.44979281425476075, - "graph_min_ms": 0.4314896106719971, - "graph_max_ms": 0.48181757926940916, - "eager_wall_median_ms": 0.4543782333106113, - "eager_wall_trials_ms": [ - 0.4543782333106113, - 0.4393956000058097, - 0.4658290666460137 - ], - "extra_peak_bytes": 185303040 - }, - { - "backend": "upstream_complex", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 55590912, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.9027584075927735, - "graph_min_ms": 0.8707056045532227, - "graph_max_ms": 0.9478416442871094, - "eager_wall_median_ms": 0.9098065333091654, - "eager_wall_trials_ms": [ - 0.92959719998665, - 0.9058386999943954, - 0.9098065333091654 - ], - "extra_peak_bytes": 188391424 - }, - { - "backend": "cuda", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 55590912, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.0960999995470047, - "graph_min_ms": 0.09566079974174499, - "graph_max_ms": 0.1124384045600891, - "eager_wall_median_ms": 0.10085823329670045, - "eager_wall_trials_ms": [ - 0.10085823329670045, - 0.1027275666880693, - 0.10056819998377857 - ], - "extra_peak_bytes": 37060608 - }, - { - "backend": "cuda", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 55590912, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.19612080454826356, - "graph_min_ms": 0.19543039798736572, - "graph_max_ms": 0.22353920936584473, - "eager_wall_median_ms": 0.21260113335301867, - "eager_wall_trials_ms": [ - 0.21260113335301867, - 0.23863556671130937, - 0.20749836670195995 - ], - "extra_peak_bytes": 74809344 - }, - { - "backend": "triton", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 55590912, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.09310480058193207, - "graph_min_ms": 0.09154719710350037, - "graph_max_ms": 0.1167631983757019, - "eager_wall_median_ms": 0.10304526670855314, - "eager_wall_trials_ms": [ - 0.09351696662633913, - 0.10913006669094709, - 0.10304526670855314 - ], - "extra_peak_bytes": 37060608 - }, - { - "backend": "triton", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 55590912, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.18827759623527526, - "graph_min_ms": 0.18518240451812745, - "graph_max_ms": 0.21609120368957518, - "eager_wall_median_ms": 0.33319789999950444, - "eager_wall_trials_ms": [ - 0.33319789999950444, - 0.3456709999833644, - 0.3028138333320385 - ], - "extra_peak_bytes": 74809344 - }, - { - "backend": "flash_default", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 55590912, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.09317199885845184, - "graph_min_ms": 0.09170079827308655, - "graph_max_ms": 0.11182240247726441, - "eager_wall_median_ms": 0.13384176666780453, - "eager_wall_trials_ms": [ - 0.12049876671274735, - 0.15613870000379393, - 0.13384176666780453 - ], - "extra_peak_bytes": 37060608 - }, - { - "backend": "flash_default", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 55590912, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.18772000074386597, - "graph_min_ms": 0.18648799657821655, - "graph_max_ms": 0.2264847993850708, - "eager_wall_median_ms": 0.3818082333358082, - "eager_wall_trials_ms": [ - 0.44640976666414645, - 0.3674994000296768, - 0.3818082333358082 - ], - "extra_peak_bytes": 74809344 - }, - { - "backend": "flash_no_fma", - "phase": "forward", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 55590912, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.09322640001773834, - "graph_min_ms": 0.09153280258178711, - "graph_max_ms": 0.11082240343093872, - "eager_wall_median_ms": 0.13263150000663396, - "eager_wall_trials_ms": [ - 0.13263150000663396, - 0.13320910002221353, - 0.12726086667195582 - ], - "extra_peak_bytes": 37060608 - }, - { - "backend": "flash_no_fma", - "phase": "forward_backward", - "layout": "packed_qkv", - "shape": [ - 1, - 6032, - 24, - 128 - ], - "input_stride": [ - 55590912, - 9216, - 128, - 1 - ], - "graph_median_ms": 0.18817280530929564, - "graph_min_ms": 0.18564800024032593, - "graph_max_ms": 0.21453280448913575, - "eager_wall_median_ms": 0.3810788333794335, - "eager_wall_trials_ms": [ - 0.3668230666638313, - 0.3885829333739821, - 0.3810788333794335 - ], - "extra_peak_bytes": 74809344 - } - ] -} diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu-contracts.txt b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu-contracts.txt deleted file mode 100644 index de9819e68..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu-contracts.txt +++ /dev/null @@ -1,20 +0,0 @@ -...........ssssssssssss.ss.ss.ss...ssssss.ss......ssssssssssss.ss...ssss [ 66%] -s.s.ssssss........................... [100%] -=========================== short test summary info ============================ -SKIPPED [12] tests/ops/rope/test_multi_axis_rope.py:79: NVIDIA CUDA required -SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:100: NVIDIA CUDA required -SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:110: NVIDIA CUDA required -SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:133: NVIDIA CUDA required -SKIPPED [6] tests/ops/rope/test_multi_axis_rope.py:148: NVIDIA CUDA required -SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:158: NVIDIA CUDA required -SKIPPED [12] tests/ops/rope/test_multi_axis_rope.py:174: NVIDIA CUDA required -SKIPPED [1] tests/ops/rope/test_multi_axis_rope.py:217: native CUDA multi_axis_rope extension required -SKIPPED [1] tests/ops/rope/test_multi_axis_rope.py:230: NVIDIA CUDA required -SKIPPED [3] tests/ops/rope/test_multi_axis_rope.py:255: NVIDIA CUDA required -SKIPPED [1] tests/ops/rope/test_multi_axis_rope.py:315: NVIDIA CUDA required -SKIPPED [1] tests/ops/rope/test_multi_axis_rope.py:336: requires two NVIDIA GPUs -SKIPPED [1] tests/ops/rope/test_multi_axis_rope.py:355: CUDA context requires CUDA -SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:368: NVIDIA CUDA required -SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:399: NVIDIA CUDA required -SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:427: NVIDIA CUDA required -57 passed, 52 skipped in 1.70s diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu-environment.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu-environment.json deleted file mode 100644 index e41fd057f..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu-environment.json +++ /dev/null @@ -1,7 +0,0 @@ -{ - "torch": "2.9.1+cpu", - "cuda": null, - "package": "/home/dandan/venvs/rlk-rope-cpu/lib/python3.12/site-packages/rl_engine/__init__.py", - "reference_forward_backward": true, - "passed": true -} diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu310-contracts.txt b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu310-contracts.txt deleted file mode 100644 index 6f0703516..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu310-contracts.txt +++ /dev/null @@ -1,20 +0,0 @@ -...........ssssssssssss.ss.ss.ss...ssssss.ss......ssssssssssss.ss...ssss [ 66%] -s.s.ssssss........................... [100%] -=========================== short test summary info ============================ -SKIPPED [12] tests/ops/rope/test_multi_axis_rope.py:79: NVIDIA CUDA required -SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:100: NVIDIA CUDA required -SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:110: NVIDIA CUDA required -SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:133: NVIDIA CUDA required -SKIPPED [6] tests/ops/rope/test_multi_axis_rope.py:148: NVIDIA CUDA required -SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:158: NVIDIA CUDA required -SKIPPED [12] tests/ops/rope/test_multi_axis_rope.py:174: NVIDIA CUDA required -SKIPPED [1] tests/ops/rope/test_multi_axis_rope.py:217: native CUDA multi_axis_rope extension required -SKIPPED [1] tests/ops/rope/test_multi_axis_rope.py:230: NVIDIA CUDA required -SKIPPED [3] tests/ops/rope/test_multi_axis_rope.py:255: NVIDIA CUDA required -SKIPPED [1] tests/ops/rope/test_multi_axis_rope.py:315: NVIDIA CUDA required -SKIPPED [1] tests/ops/rope/test_multi_axis_rope.py:336: requires two NVIDIA GPUs -SKIPPED [1] tests/ops/rope/test_multi_axis_rope.py:355: CUDA context requires CUDA -SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:368: NVIDIA CUDA required -SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:399: NVIDIA CUDA required -SKIPPED [2] tests/ops/rope/test_multi_axis_rope.py:427: NVIDIA CUDA required -57 passed, 52 skipped in 1.72s diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu310-environment.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu310-environment.json deleted file mode 100644 index 833e69038..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/cpu310-environment.json +++ /dev/null @@ -1,9 +0,0 @@ -{ - "python": "3.10.22", - "torch": "2.9.1+cpu", - "cuda": null, - "package": "/home/dandan/venvs/rlk-rope-cpu310/lib/python3.10/site-packages/rl_engine", - "reference_forward_backward": true, - "cpu_wheel_contains_no_cuda_binary": true, - "passed": true -} diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/full-extension.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/full-extension.json deleted file mode 100644 index 420aa5da8..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/full-extension.json +++ /dev/null @@ -1,22 +0,0 @@ -{ - "profile": "ordinary complete _C; no RL_KERNEL_ROPE_ONLY", - "module": "/home/dandan/src/RL-Kernel/rl_engine/_C.cpython-312-x86_64-linux-gnu.so", - "binary_sha256": "2661e7188984c7d78acae1ef5186a3e90ec4aa8bb6f457522b23435a1f24f94a", - "sources": [ - "csrc/bindings/ops.cpp", - "csrc/cuda/logprob/fused_logp_kernel.cu", - "csrc/cuda/logprob/deterministic_logp_kernel.cu", - "csrc/cuda/gemm/det_gemm_kernel.cu", - "csrc/cuda/norm/rmsnorm.cu", - "csrc/cuda/activation/activation.cu", - "csrc/cuda/attention/deterministic_attention.cu", - "csrc/cuda/rope/multi_axis_rope.cu", - "csrc/cuda/collectives/deterministic_collective.cu", - "csrc/cuda/attention/prefix_shared_attention.cu" - ], - "cuda_home": "/usr/local/cuda-12.8", - "driver_link_search": "/usr/lib/wsl/lib", - "torch": "2.9.1+cu128", - "architecture": "12.0", - "max_jobs": "2" -} diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/full-inplace-build.txt b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/full-inplace-build.txt deleted file mode 100644 index d61fb7517..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/full-inplace-build.txt +++ /dev/null @@ -1,20 +0,0 @@ -/home/dandan/src/RL-Kernel/.venv/lib/python3.12/site-packages/setuptools/config/_apply_pyprojecttoml.py:82: SetuptoolsDeprecationWarning: `project.license` as a TOML table is deprecated -!! - - ******************************************************************************** - Please use a simple string containing a SPDX expression for `project.license`. You can also use `project.license-files`. (Both options available on setuptools>=77.0.0). - - This deprecation is overdue, please update your project and remove deprecated - calls to avoid build errors in the future. - - See https://packaging.python.org/en/latest/guides/writing-pyproject-toml/#license for details. - ******************************************************************************** - -!! - corresp(dist, value, root_dir) -running build_ext -W1028 18:37:59.459000 9791 .venv/lib/python3.12/site-packages/torch/utils/cpp_extension.py:531] There are no x86_64-linux-gnu-g++ version bounds defined for CUDA version 12.8 -building 'rl_engine._C' extension -ninja: no work to do. -x86_64-linux-gnu-g++ -fno-strict-overflow -Wsign-compare -DNDEBUG -g -O2 -Wall -shared -Wl,-O1 -Wl,-Bsymbolic-functions /home/dandan/src/RL-Kernel/build/temp.linux-x86_64-cpython-312/csrc/bindings/ops.o /home/dandan/src/RL-Kernel/build/temp.linux-x86_64-cpython-312/csrc/cuda/activation/activation.o /home/dandan/src/RL-Kernel/build/temp.linux-x86_64-cpython-312/csrc/cuda/attention/deterministic_attention.o /home/dandan/src/RL-Kernel/build/temp.linux-x86_64-cpython-312/csrc/cuda/attention/prefix_shared_attention.o /home/dandan/src/RL-Kernel/build/temp.linux-x86_64-cpython-312/csrc/cuda/collectives/deterministic_collective.o /home/dandan/src/RL-Kernel/build/temp.linux-x86_64-cpython-312/csrc/cuda/gemm/det_gemm_kernel.o /home/dandan/src/RL-Kernel/build/temp.linux-x86_64-cpython-312/csrc/cuda/logprob/deterministic_logp_kernel.o /home/dandan/src/RL-Kernel/build/temp.linux-x86_64-cpython-312/csrc/cuda/logprob/fused_logp_kernel.o /home/dandan/src/RL-Kernel/build/temp.linux-x86_64-cpython-312/csrc/cuda/norm/rmsnorm.o /home/dandan/src/RL-Kernel/build/temp.linux-x86_64-cpython-312/csrc/cuda/rope/multi_axis_rope.o -L/home/dandan/src/RL-Kernel/.venv/lib/python3.12/site-packages/torch/lib -L/usr/local/cuda-12.8/lib64 -L/usr/lib/x86_64-linux-gnu -lc10 -ltorch -ltorch_cpu -ltorch_python -lcudart -lc10_cuda -ltorch_cuda -o build/lib.linux-x86_64-cpython-312/rl_engine/_C.cpython-312-x86_64-linux-gnu.so -Wl,-rpath,$ORIGIN/../torch/lib -lcuda -copying build/lib.linux-x86_64-cpython-312/rl_engine/_C.cpython-312-x86_64-linux-gnu.so -> rl_engine diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/full-wheel-build.txt b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/full-wheel-build.txt deleted file mode 100644 index 592c254f2..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/full-wheel-build.txt +++ /dev/null @@ -1,10 +0,0 @@ -Processing /home/dandan/src/RL-Kernel - Preparing metadata (pyproject.toml): started - Preparing metadata (pyproject.toml): finished with status 'done' -Building wheels for collected packages: RL-Kernel - Building wheel for RL-Kernel (pyproject.toml): started - Building wheel for RL-Kernel (pyproject.toml): still running... - Building wheel for RL-Kernel (pyproject.toml): finished with status 'done' - Created wheel for RL-Kernel: filename=rl_kernel-0.1.0-cp312-cp312-linux_x86_64.whl size=5020592 sha256=32a39cd9cbb30b87816f6584096238fc2e4b0803eb387a2d3e1e80f3d1b40d8a - Stored in directory: /tmp/pip-ephem-wheel-cache-8gtts1_e/wheels/de/2c/1f/c224e9a9cc43d1b39ee96ba5703938f5f1d66da44d8695a43a -Successfully built RL-Kernel diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/gpu-environment.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/gpu-environment.json deleted file mode 100644 index e40096583..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/gpu-environment.json +++ /dev/null @@ -1,15 +0,0 @@ -{ - "scope": "Qwen-Image multi_axis_rope operator; ordinary full _C build", - "python": "3.12.3", - "platform": "Linux-6.18.54.1-microsoft-standard-WSL2-x86_64-with-glibc2.39", - "torch": "2.9.1+cu128", - "cuda": "12.8", - "triton": "3.5.1", - "gpu": "NVIDIA GeForce RTX 5080", - "capability": [ - 12, - 0 - ], - "native_library": "/home/dandan/src/RL-Kernel/rl_engine/_C.cpython-312-x86_64-linux-gnu.so", - "native_library_sha256": "2661e7188984c7d78acae1ef5186a3e90ec4aa8bb6f457522b23435a1f24f94a" -} diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/installed-wheel.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/installed-wheel.json deleted file mode 100644 index 28d3066fc..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/installed-wheel.json +++ /dev/null @@ -1,38 +0,0 @@ -{ - "package": "/home/dandan/src/RL-Kernel/.venv/lib/python3.12/site-packages/rl_engine", - "companion_library": null, - "full_extension_library": "/home/dandan/src/RL-Kernel/.venv/lib/python3.12/site-packages/rl_engine/_C.cpython-312-x86_64-linux-gnu.so", - "native_dispatcher_registered": true, - "torch": "2.9.1+cu128", - "gpu": "NVIDIA GeForce RTX 5080", - "license_present": true, - "packed_qkv_stride": [ - 17280, - 1152, - 128, - 1 - ], - "checks": { - "cuda_eager": "passed", - "cuda_compiled": "passed", - "cuda_trace": { - "contract": "qwenimage-interleaved-rope-fp32-v1", - "actual_backend": "rlkernel.cuda.qwenimage_multi_axis_rope.v1", - "kernel_id": "torch.ops.rl_kernel_native.multi_axis_rope", - "layout": "BSHD, adjacent channel pairs", - "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", - "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" - }, - "triton_eager": "passed", - "triton_compiled": "passed", - "triton_trace": { - "contract": "qwenimage-interleaved-rope-fp32-v1", - "actual_backend": "rlkernel.triton.qwenimage_multi_axis_rope.v1", - "kernel_id": "rl_engine.backends.shared.triton.rope._flash_attn_rotary.rotary_kernel", - "layout": "BSHD, adjacent channel pairs", - "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", - "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" - } - }, - "passed": true -} diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/layout.txt b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/layout.txt deleted file mode 100644 index 8c81ecbf5..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/layout.txt +++ /dev/null @@ -1,2 +0,0 @@ -..................................... [100%] -37 passed in 3.81s diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/manifest.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/manifest.json deleted file mode 100644 index 9db29dacc..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/manifest.json +++ /dev/null @@ -1,60 +0,0 @@ -{ - "profile": "Ubuntu 24.04.5 / WSL2 / RTX 5080 SM120 / ordinary complete _C", - "numerical_source_commit": "ac17ab87e8d93dabe93dfdc134e9853d47de7af1", - "source_sha256": { - ".github/workflows/qwenimage-rope.yml": "39e7e02a1f1f75603482e6acce9eb99ba8b5c5260ec6043271714061843651c3", - "MANIFEST.in": "87e59e6c40c67e54ae9bcd057957ac7d7823067cb455cbc4dff45b1ec88bdb3c", - "benchmarks/operators/rope/benchmark_multi_axis_rope.py": "6961a55803fbd4e566bd060558fb59e658e7f1f704e61ca97f912e5d16b55884", - "build_tools/extensions.py": "3723f14a3fb9307889790af15b5e3f360c4c46e12b8fef22423070e9660f70c5", - "ci/scripts/run_qwenimage_rope_ws1.sh": "8e1e12c1aad9231f215505a172b3a5416c54e7cbfb8856f5ff00d225be7bfa95", - "csrc/bindings/ops.cpp": "10b76a5e236c0dc3715f9703e28c057f404f21f2a9802245d5221c4997b2c3f4", - "csrc/cuda/activation/activation.cu": "35ef0cefcaa6167aec7a901fd19a2f6c90a8b9f01c2e786187aacdeee10c4b63", - "csrc/cuda/attention/deterministic_attention.cu": "99b44addd61b999a5fb6a04ca227ed8814d454368f361a8fd032e83d4d83a561", - "csrc/cuda/attention/prefix_shared_attention.cu": "218ab9647cd274f7a4d8246f8df8e875c5cf431a5a73f252e946b10c74a2d63f", - "csrc/cuda/collectives/deterministic_collective.cu": "9e114d5aae7a73a912e4e4afaa23a74126bef5557ba3b54832da0c0cdf9d97a2", - "csrc/cuda/gemm/det_gemm_kernel.cu": "08d08b0a32900fd9590d47f04819b34a2d37cbf53996aa5ed1373d88eae29f33", - "csrc/cuda/logprob/deterministic_logp_kernel.cu": "271399040a0489a3cabd9e0c06b64fa8b1bac92ab08710fce8922bcb1ff2b04f", - "csrc/cuda/logprob/fused_logp_kernel.cu": "ea67bcd307f0c63753397225543d333f58cf3d66a0ae69b643499d1e6729d8e7", - "csrc/cuda/norm/rmsnorm.cu": "fbb765563dbe9c063b35f38354321eaf522cf62ba03addf327154395db9683b9", - "csrc/cuda/rope/multi_axis_rope.cu": "f5db5cf5d4962053194455d9cde8cf91cf4e824bad7abf72e8410a404bcb0e44", - "docs/operators/multi-axis-rope.md": "1fb3859a9a2639aae809ece52a785737e9966098c24a08626b9240dad0d6cbf6", - "pyproject.toml": "c2d948784b379a0a828bbf3d8cfd9e0e7398a1ee5694fe8298509514684b2709", - "rl_engine/backends/cuda/rope/multi_axis_rope.py": "d779d9b60ce5d1d6651925f723049930d898024fdf0548ce29a2019b84753d65", - "rl_engine/backends/shared/triton/rope/_flash_attn_rotary.py": "cf0eebf315830a4178053a71a77e650b58aa85bba1d1a08aff7677350af755d9", - "rl_engine/backends/shared/triton/rope/multi_axis_rope.py": "f08ff90627c66e15e9c558f8a68f132434a741a76997b3800bfbf73dc8940e9e", - "rl_engine/ops/rope/_registration.py": "add0375d1639444908091d01e8d29781f7a975beffec64b77bfaa81a5b9a0ae3", - "rl_engine/ops/rope/multi_axis_rope.py": "07bcff1160f22c1ff6730be731206b2f4a153e2fada26afc93fc884e780568cb", - "rl_engine/reference/rope/multi_axis_rope.py": "c9a94730047f51659f7aa0a443e88d6ce7d03d00d6813a9aed548cafbf6347bf", - "rl_engine/validation/operators/qwen_image_rope.py": "c3ddfc732f32de3c633de16064810d67202254ea78fd8107d558b471bb15a37f", - "tests/ops/rope/test_multi_axis_rope.py": "c9ff3eb6e6e013243fa741f853debb256fdac6c1d3a3820eb7de764a830e240c", - "tools/validation/operators/check_multi_axis_rope_install.py": "f295fc8ac5878ac13b6255c2f0be83b66d34ff2ad2f75111bb8b633508a865a7" - }, - "evidence_sha256": { - "README.md": "3ad6f8ef514b4fc61d51fb0ce8cb6d420d2180251c9ee4f84b8c931aa3a75094", - "benchmark-bfloat16.json": "e9c1e4b1085623d84800656c179e4a0e86d70a0ef7fc9ed92f9e6ccc50f0f491", - "cpu-contracts.txt": "9518f53a6efc603bca1c52f160e5430f1006bb2abc9a7ef57694dae6feac5d78", - "cpu-environment.json": "4dfc68f3d76af05a050d830840b4a6b8c8fa0a4c12c29d11d632a56a57a2d4c6", - "cpu310-contracts.txt": "02f57808b9553cefe6f4dcec06d46a5163157b4290cc0c590b344a1616057852", - "cpu310-environment.json": "7130d4c164c98804284cc0f1e07ae6265c35e7498c2fb8285a5e6df87cb010b2", - "full-extension.json": "8fae782e3c5e6aae59de6fe6b4e7ae185729e49d2282dbd0cb3b9ca08c0274cd", - "full-inplace-build.txt": "04d5ef122650ab305847d58850a6c3b44438eacf97920492e1eb1c5580e9a21e", - "full-wheel-build.txt": "a7883dfb1211798c869ede1ee0b6f1dfb32edd7b27b359ee18c6e5ab9fe449d5", - "gpu-environment.json": "bf808ab535d90b9b52dfcfb86543809e7d412cb48418677923cad22783bb9b8e", - "installed-wheel.json": "c726717106f30336e0a541e4a38d1f1b3858ed5dd262516c86618d4511a43554", - "layout.txt": "daef0733010cdad31123cb6b8cb041fe6c626319d527a5d97c88d1b958479039", - "memcheck.txt": "1a0c572f8f0d18e3699c847dbb69b345e4fd19200cabda3a39385d588cff0f10", - "operator-cuda.json": "fd95922bac2c5921744a8e601e9a267e51b7b7d7d44d94745f53964c2da2feb1", - "operator-tests.xml": "0ed98ddf85e34ba87f41465be1db4f417997e82b50cc5f8f5a83bd00f1f1543e", - "operator-triton.json": "7791aa1dfa68652dd66e3c332d8f036e9e21851e1df7fd41a090d594e22b06ff", - "package.json": "3f75c06fdee58afb1cdd9a2a611b432c8be1960fa640d96263c86b87423fe7bd", - "precommit.txt": "7cf90e06cf73194271c176ebc2d147e64cf8f5ed522d8e0e03093a4647f63759", - "python-packages.json": "617f8cef5a24f3b2a683e70e0b491e47db6d24a7ca848e5b3e461e87993a7201", - "quality.txt": "8826f0e406946a923400d792dd60d0f545b32f9fcffa3a148268c969e30e14c0", - "regression.txt": "039cd9813a2ea54bfc5b7750fc0555b54d2bf30b09dbca2eba4a0403578ab9b0", - "tests-and-ws1.txt": "68397e27144c9244743dbcff39ffdca8a77a30eaf1e8180b2ca84dfc1216d633", - "upstream.json": "04e5472a82477a305b996ccbb420d0f3cc7b39af4053c867d37916a9bf5de51a", - "ws1-cuda.json": "7302e11f798816178622393095858871e2dc53ab925509eb130196a40b3c18ad", - "ws1-triton.json": "b21f985fbdff7e9b14f689d5eac25c9df5d6d3b0189900085837d5fc400f9d38" - }, - "hash_mode": "SHA256 of recorded Linux file bytes; text is LF" -} diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/memcheck.txt b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/memcheck.txt deleted file mode 100644 index 915746141..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/memcheck.txt +++ /dev/null @@ -1,4 +0,0 @@ -========= COMPUTE-SANITIZER -.................................. [100%] -34 passed, 48 deselected in 3.49s -========= ERROR SUMMARY: 0 errors diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/operator-cuda.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/operator-cuda.json deleted file mode 100644 index 801312e4a..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/operator-cuda.json +++ /dev/null @@ -1,70 +0,0 @@ -{ - "suite_name": "multi_axis_rope", - "total_candidates": 1, - "passed_candidates": 1, - "pass_rate": 1.0, - "passed": true, - "candidates": [ - { - "candidate_name": "cuda-multi_axis_rope", - "backend": "cuda", - "total_outputs": 2, - "passed_outputs": 2, - "pass_rate": 1.0, - "passed": true, - "cases": [ - { - "case_name": "multi_axis_rope-torch.bfloat16-2x7x3x128", - "dtype": "torch.bfloat16", - "op_class": "elementwise", - "passed": true, - "outputs": [ - { - "output_index": 0, - "shape": [ - 2, - 7, - 3, - 128 - ], - "candidate_dtype": "torch.bfloat16", - "gold_dtype": "torch.float32", - "atol": 0.02, - "rtol": 0.016, - "max_abs_error": 0.007704973220825195, - "mean_abs_error": 0.0006018930580466986, - "max_rel_error": 0.003828622167930007, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference", - "message": "" - }, - { - "output_index": 1, - "shape": [ - 2, - 7, - 3, - 128 - ], - "candidate_dtype": "torch.bfloat16", - "gold_dtype": "torch.bfloat16", - "atol": 0.02, - "rtol": 0.016, - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference", - "message": "gradient:x" - } - ] - } - ], - "backend_provenance": null - } - ] -} diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/operator-tests.xml b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/operator-tests.xml deleted file mode 100644 index 9e26942b9..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/operator-tests.xml +++ /dev/null @@ -1 +0,0 @@ -/home/dandan/src/RL-Kernel/tests/ops/rope/test_multi_axis_rope.py:336: requires two NVIDIA GPUs diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/operator-triton.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/operator-triton.json deleted file mode 100644 index e105fb723..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/operator-triton.json +++ /dev/null @@ -1,70 +0,0 @@ -{ - "suite_name": "multi_axis_rope", - "total_candidates": 1, - "passed_candidates": 1, - "pass_rate": 1.0, - "passed": true, - "candidates": [ - { - "candidate_name": "triton-multi_axis_rope", - "backend": "triton", - "total_outputs": 2, - "passed_outputs": 2, - "pass_rate": 1.0, - "passed": true, - "cases": [ - { - "case_name": "multi_axis_rope-torch.bfloat16-2x7x3x128", - "dtype": "torch.bfloat16", - "op_class": "elementwise", - "passed": true, - "outputs": [ - { - "output_index": 0, - "shape": [ - 2, - 7, - 3, - 128 - ], - "candidate_dtype": "torch.bfloat16", - "gold_dtype": "torch.float32", - "atol": 0.02, - "rtol": 0.016, - "max_abs_error": 0.007704973220825195, - "mean_abs_error": 0.0006018930580466986, - "max_rel_error": 0.003828622167930007, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference", - "message": "" - }, - { - "output_index": 1, - "shape": [ - 2, - 7, - 3, - 128 - ], - "candidate_dtype": "torch.bfloat16", - "gold_dtype": "torch.bfloat16", - "atol": 0.02, - "rtol": 0.016, - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference", - "message": "gradient:x" - } - ] - } - ], - "backend_provenance": null - } - ] -} diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/package.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/package.json deleted file mode 100644 index 1dce0c941..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/package.json +++ /dev/null @@ -1,11 +0,0 @@ -{ - "profile": "ordinary complete _C wheel; RL_KERNEL_ROPE_ONLY unset", - "wheel": "rl_kernel-0.1.0-cp312-cp312-linux_x86_64.whl", - "wheel_sha256": "32a39cd9cbb30b87816f6584096238fc2e4b0803eb387a2d3e1e80f3d1b40d8a", - "native_member": "rl_engine/_C.cpython-312-x86_64-linux-gnu.so", - "native_sha256": "2661e7188984c7d78acae1ef5186a3e90ec4aa8bb6f457522b23435a1f24f94a", - "license_present": true, - "installed_library": "/home/dandan/src/RL-Kernel/.venv/lib/python3.12/site-packages/rl_engine/_C.cpython-312-x86_64-linux-gnu.so", - "installed_library_sha256": "2661e7188984c7d78acae1ef5186a3e90ec4aa8bb6f457522b23435a1f24f94a", - "cpu_wheel_contains_no_cuda_binary": true -} diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/precommit.txt b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/precommit.txt deleted file mode 100644 index 475a8c0d8..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/precommit.txt +++ /dev/null @@ -1,7 +0,0 @@ -trim trailing whitespace.................................................Passed -fix end of files.........................................................Passed -check yaml...............................................................Passed -check for added large files..............................................Passed -black....................................................................Passed -isort....................................................................Passed -flake8...................................................................Passed diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/python-packages.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/python-packages.json deleted file mode 100644 index 41997f91d..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/python-packages.json +++ /dev/null @@ -1,114 +0,0 @@ -{ - "GitPython": "3.2.1", - "Jinja2": "3.1.6", - "Markdown": "3.11.2", - "MarkupSafe": "3.0.3", - "PyYAML": "6.0.3", - "Pygments": "2.21.0", - "RL-Kernel": "0.1.0", - "accelerate": "1.15.0", - "annotated-doc": "0.0.5", - "annotated-types": "0.8.0", - "anyio": "4.15.1", - "ast_serialize": "0.13.1", - "babel": "2.18.0", - "backrefs": "8.1", - "black": "24.4.2", - "bracex": "3.2", - "certifi": "2026.10.23", - "cfgv": "3.5.0", - "charset-normalizer": "3.5.2", - "click": "8.5.0", - "colorama": "0.4.6", - "distlib": "0.4.4", - "filelock": "3.32.3", - "flake8": "7.0.0", - "fsspec": "2026.7.0", - "ghp-import": "2.1.0", - "gitdb": "4.0.12", - "h11": "0.16.0", - "hf-xet": "1.7.0", - "httpcore": "1.0.9", - "httpx": "0.28.1", - "huggingface_hub": "1.33.0", - "identify": "2.6.20", - "idna": "3.20", - "iniconfig": "2.3.1", - "isort": "5.13.2", - "librt": "0.16.0", - "markdown-it-py": "4.2.0", - "mccabe": "0.7.0", - "mdurl": "0.1.2", - "mergedeep": "1.3.4", - "mkdocs": "1.6.1", - "mkdocs-awesome-nav": "3.3.0", - "mkdocs-get-deps": "0.2.2", - "mkdocs-git-revision-date-localized-plugin": "1.6.2", - "mkdocs-material": "9.8.0", - "mkdocs-material-extensions": "1.3.1", - "mpmath": "1.3.0", - "mypy": "2.4.0", - "mypy_extensions": "1.1.0", - "natsort": "8.4.0", - "networkx": "3.6.1", - "ninja": "1.13.2", - "nodeenv": "1.11.0", - "numpy": "2.5.4", - "nvidia-cublas-cu12": "12.8.4.1", - "nvidia-cuda-cupti-cu12": "12.8.90", - "nvidia-cuda-nvrtc-cu12": "12.8.93", - "nvidia-cuda-runtime-cu12": "12.8.90", - "nvidia-cudnn-cu12": "9.10.2.21", - "nvidia-cufft-cu12": "11.3.3.83", - "nvidia-cufile-cu12": "1.13.1.3", - "nvidia-curand-cu12": "10.3.9.90", - "nvidia-cusolver-cu12": "11.7.3.90", - "nvidia-cusparse-cu12": "12.5.8.93", - "nvidia-cusparselt-cu12": "0.7.1", - "nvidia-nccl-cu12": "2.27.5", - "nvidia-nvjitlink-cu12": "12.8.93", - "nvidia-nvshmem-cu12": "3.3.20", - "nvidia-nvtx-cu12": "12.8.90", - "packaging": "26.3", - "paginate": "0.5.7", - "pathspec": "1.1.1", - "pip": "25.3", - "platformdirs": "4.13.7", - "pluggy": "1.7.0", - "pre_commit": "4.7.0", - "psutil": "7.2.2", - "pycodestyle": "2.11.1", - "pydantic": "2.14.1", - "pydantic_core": "2.50.1", - "pyflakes": "3.2.0", - "pymdown-extensions": "12.3", - "pytest": "9.2.0", - "python-dateutil": "3.0.0", - "python-discovery": "1.6.7", - "pyyaml_env_tag": "1.1", - "regex": "2026.9.29", - "requests": "2.34.2", - "rich": "15.0.0", - "ruff": "0.17.3", - "safetensors": "0.8.0", - "setuptools": "80.9.0", - "shellingham": "1.5.4", - "six": "1.17.0", - "smmap": "5.0.3", - "sympy": "1.14.0", - "tabulate": "0.10.0", - "tokenizers": "0.22.2", - "torch": "2.9.1+cu128", - "tqdm": "4.70.1", - "transformers": "5.13.1", - "triton": "3.5.1", - "typer": "0.27.3", - "typing-inspection": "0.4.4", - "typing_extensions": "4.16.0", - "tzdata": "2026.6", - "urllib3": "2.8.0", - "virtualenv": "21.15.1", - "watchdog": "6.0.0", - "wcmatch": "11.2", - "wheel": "0.45.1" -} diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/quality.txt b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/quality.txt deleted file mode 100644 index 8e84b2473..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/quality.txt +++ /dev/null @@ -1,34 +0,0 @@ -Local verification; no CUDA numerical/kernel sources changed for this follow-up. - -workflow-validation.log -Workflow YAML, branch coverage and embedded Python syntax validated - -isort.log - - -black.log -All done! ✨ 🍰 ✨ -19 files would be left unchanged. - -flake8.log - - -ruff.log -All checks passed! - -docs.log -INFO - Cleaning site directory -INFO - Building documentation to directory: /home/dandan/src/RL-Kernel/site -INFO - The following pages exist in the docs directory, but are not included in the "nav" configuration: - - blog/2026-07-08-announcing-rl-kernel-linear-logp-for-vime-zh.md - - blog/2026-08-09-dsv4-flash-moe-consistency-roadmap.md - - operators/matmul.md - - operators/rope.md - - usage/cross-config-drift-report.md - - usage/cuda-rocm-consistency-audit.md - - usage/h100-configurable.md - - usage/h100-cp-gradient-validation.md - - usage/h100-matrix-validation.md - - usage/h100-pcp-validation.md - - usage/rocm-sparse-performance.md -INFO - Documentation built in 1.15 seconds diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/regression.txt b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/regression.txt deleted file mode 100644 index 95d27f812..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/regression.txt +++ /dev/null @@ -1,10 +0,0 @@ -........................................................................ [ 35%] -........................................................................ [ 71%] -...................................................sssssss [100%] -=========================== short test summary info ============================ -SKIPPED [3] tests/ops/rope/test_rope.py:390: rope_apply_ascend is not compiled (requires an Ascend NPU build). -SKIPPED [1] tests/ops/rope/test_rope.py:406: rope_apply_ascend is not compiled (requires an Ascend NPU build). -SKIPPED [1] tests/ops/rope/test_rope.py:413: rope_apply_ascend is not compiled (requires an Ascend NPU build). -SKIPPED [1] tests/ops/rope/test_rope.py:420: rope_apply_ascend is not compiled (requires an Ascend NPU build). -SKIPPED [1] tests/ops/rope/test_rope.py:433: rope_apply_ascend is not compiled (requires an Ascend NPU build). -195 passed, 7 skipped in 4.56s diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/tests-and-ws1.txt b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/tests-and-ws1.txt deleted file mode 100644 index 0b8ce46dd..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/tests-and-ws1.txt +++ /dev/null @@ -1,11 +0,0 @@ -{"scope": "Qwen-Image multi_axis_rope operator; ordinary full _C build", "python": "3.12.3", "platform": "Linux-6.18.54.1-microsoft-standard-WSL2-x86_64-with-glibc2.39", "torch": "2.9.1+cu128", "cuda": "12.8", "triton": "3.5.1", "gpu": "NVIDIA GeForce RTX 5080", "capability": [12, 0], "native_library": "/home/dandan/src/RL-Kernel/rl_engine/_C.cpython-312-x86_64-linux-gnu.so", "native_library_sha256": "2661e7188984c7d78acae1ef5186a3e90ec4aa8bb6f457522b23435a1f24f94a"} -........................................................................ [ 87%] -s......... [100%] -=========================== short test summary info ============================ -SKIPPED [1] tests/ops/rope/test_multi_axis_rope.py:336: requires two NVIDIA GPUs -81 passed, 1 skipped in 17.49s -Using the installed production RoPE dispatcher library -WS1 four judgments backend=cuda: passed=True -Using the installed production RoPE dispatcher library -WS1 four judgments backend=triton: passed=True -Qwen-Image RoPE full-extension operator gate passed diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/upstream.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/upstream.json deleted file mode 100644 index 032e23098..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/upstream.json +++ /dev/null @@ -1,335 +0,0 @@ -{ - "source_sha256": "0bad11901deae69f05aec01fd464e51e38582d455068f2851ee665036179321f", - "cases": [ - { - "grid": [ - 2, - 3 - ], - "stream": "image", - "dtype": "torch.float32", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 2, - 3 - ], - "stream": "image", - "dtype": "torch.float16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 2, - 3 - ], - "stream": "image", - "dtype": "torch.bfloat16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 2, - 3 - ], - "stream": "text", - "dtype": "torch.float32", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 2, - 3 - ], - "stream": "text", - "dtype": "torch.float16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 2, - 3 - ], - "stream": "text", - "dtype": "torch.bfloat16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 3, - 5 - ], - "stream": "image", - "dtype": "torch.float32", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 3, - 5 - ], - "stream": "image", - "dtype": "torch.float16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 3, - 5 - ], - "stream": "image", - "dtype": "torch.bfloat16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 3, - 5 - ], - "stream": "text", - "dtype": "torch.float32", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 3, - 5 - ], - "stream": "text", - "dtype": "torch.float16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 3, - 5 - ], - "stream": "text", - "dtype": "torch.bfloat16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 64, - 64 - ], - "stream": "image", - "dtype": "torch.float32", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 64, - 64 - ], - "stream": "image", - "dtype": "torch.float16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 64, - 64 - ], - "stream": "image", - "dtype": "torch.bfloat16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 64, - 64 - ], - "stream": "text", - "dtype": "torch.float32", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 64, - 64 - ], - "stream": "text", - "dtype": "torch.float16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 64, - 64 - ], - "stream": "text", - "dtype": "torch.bfloat16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 83, - 83 - ], - "stream": "image", - "dtype": "torch.float32", - "table_bitwise": true, - "forward_max_abs": 2.384185791015625e-07, - "dx_max_abs": 2.384185791015625e-07 - }, - { - "grid": [ - 83, - 83 - ], - "stream": "image", - "dtype": "torch.float16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 83, - 83 - ], - "stream": "image", - "dtype": "torch.bfloat16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 83, - 83 - ], - "stream": "text", - "dtype": "torch.float32", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 83, - 83 - ], - "stream": "text", - "dtype": "torch.float16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 83, - 83 - ], - "stream": "text", - "dtype": "torch.bfloat16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 104, - 58 - ], - "stream": "image", - "dtype": "torch.float32", - "table_bitwise": true, - "forward_max_abs": 2.384185791015625e-07, - "dx_max_abs": 2.384185791015625e-07 - }, - { - "grid": [ - 104, - 58 - ], - "stream": "image", - "dtype": "torch.float16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 104, - 58 - ], - "stream": "image", - "dtype": "torch.bfloat16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 104, - 58 - ], - "stream": "text", - "dtype": "torch.float32", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 104, - 58 - ], - "stream": "text", - "dtype": "torch.float16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - }, - { - "grid": [ - 104, - 58 - ], - "stream": "text", - "dtype": "torch.bfloat16", - "table_bitwise": true, - "forward_max_abs": 0.0, - "dx_max_abs": 0.0 - } - ] -} diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/ws1-cuda.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/ws1-cuda.json deleted file mode 100644 index 3a3e06daf..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/ws1-cuda.json +++ /dev/null @@ -1,1962 +0,0 @@ -{ - "scope": "Qwen-Image multi_axis_rope operator; C1 elementwise and C2 config transforms", - "trace": { - "contract": "qwenimage-interleaved-rope-fp32-v1", - "actual_backend": "rlkernel.cuda.qwenimage_multi_axis_rope.v1", - "kernel_id": "torch.ops.rl_kernel_native.multi_axis_rope", - "layout": "BSHD, adjacent channel pairs", - "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", - "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" - }, - "passed": true, - "forward": { - "op_name": "multi_axis_rope", - "backend_profile": "cuda_bf16", - "accuracy_reports": [ - { - "config_id": "BN/full", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.054684440139681e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/chunked", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/chunked", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.054684440139681e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s0", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/full/s0", - "fp32_reference" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004673302173614502, - "mean_abs_error": 6.917398422956467e-05, - "max_rel_error": 0.003753266530111432, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s1", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/full/s1", - "fp32_reference" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.00048516690731048584, - "mean_abs_error": 7.802648178767413e-05, - "max_rel_error": 0.0037793784867972136, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s2", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/full/s2", - "fp32_reference" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004864335060119629, - "mean_abs_error": 7.697803084738553e-05, - "max_rel_error": 0.003821910358965397, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s3", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/full/s3", - "fp32_reference" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.710371912457049e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s0", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/chunked/s0", - "fp32_reference" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004673302173614502, - "mean_abs_error": 6.917398422956467e-05, - "max_rel_error": 0.003753266530111432, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s1", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/chunked/s1", - "fp32_reference" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.00048516690731048584, - "mean_abs_error": 7.802648178767413e-05, - "max_rel_error": 0.0037793784867972136, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s2", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/chunked/s2", - "fp32_reference" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004864335060119629, - "mean_abs_error": 7.697803084738553e-05, - "max_rel_error": 0.003821910358965397, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s3", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/chunked/s3", - "fp32_reference" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.710371912457049e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/permuted", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/permuted", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.054684440139681e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/padded_right", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/padded_right", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.054684440139681e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/padded_left", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/padded_left", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.054684440139681e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - } - ], - "invariance_reports": [ - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/chunked", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "BN/chunked" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s0", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s0" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s1", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s1" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s2", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s2" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s3", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s3" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s0", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s0" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s1", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s1" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s2", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s2" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s3", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s3" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/permuted", - "transform_kind": "permutation", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "BN/permuted" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/padded_right", - "transform_kind": "padding", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "BN/padded_right" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/padded_left", - "transform_kind": "padding", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "BN/padded_left" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - } - ], - "logprob_smoke": null, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - }, - "candidate_id": "rlkernel.cuda.qwenimage_multi_axis_rope.v1", - "device": "cuda", - "compute_capability": "sm120", - "seed": 20260812, - "fallback_reason": null, - "passed": true, - "provenance_valid": true, - "metadata_valid": true, - "observed_kernel_id": "torch.ops.rl_kernel_native.multi_axis_rope" - }, - "gradient": { - "op_name": "multi_axis_rope", - "backend_profile": "cuda_bf16", - "accuracy_reports": [ - { - "config_id": "BN/full", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.699562850873917e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/chunked", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/chunked", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.699562850873917e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s0", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/full/s0", - "fp32_reference" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 2.7887988835573196e-05, - "mean_abs_error": 3.6821973026235355e-06, - "max_rel_error": 1.8164952993392944, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s1", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/full/s1", - "fp32_reference" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.7113780965446495e-06, - "max_rel_error": 0.4316565990447998, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s2", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/full/s2", - "fp32_reference" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 2.995925024151802e-05, - "mean_abs_error": 4.716822331829462e-06, - "max_rel_error": 0.7439075112342834, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s3", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/full/s3", - "fp32_reference" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.0211638659238815e-05, - "mean_abs_error": 4.9605891945248e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s0", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/chunked/s0", - "fp32_reference" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 2.7887988835573196e-05, - "mean_abs_error": 3.6821973026235355e-06, - "max_rel_error": 1.8164952993392944, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s1", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/chunked/s1", - "fp32_reference" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.7113780965446495e-06, - "max_rel_error": 0.4316565990447998, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s2", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/chunked/s2", - "fp32_reference" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 2.995925024151802e-05, - "mean_abs_error": 4.716822331829462e-06, - "max_rel_error": 0.7439075112342834, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s3", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/chunked/s3", - "fp32_reference" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.0211638659238815e-05, - "mean_abs_error": 4.9605891945248e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/permuted", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/permuted", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.699562850873917e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/padded_right", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/padded_right", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.699562850873917e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/padded_left", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/padded_left", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.699562850873917e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - } - ], - "invariance_reports": [ - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/chunked", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "BN/chunked" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s0", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s0" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s1", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s1" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s2", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s2" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s3", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s3" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s0", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s0" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s1", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s1" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s2", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s2" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s3", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s3" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/permuted", - "transform_kind": "permutation", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "BN/permuted" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/padded_right", - "transform_kind": "padding", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "BN/padded_right" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/padded_left", - "transform_kind": "padding", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "BN/padded_left" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - } - ], - "singleton_aggregate_reports": [], - "backend_provenance": { - "backend_profile": "cuda_bf16", - "requested_backend": "cuda", - "actual_backend": "cuda", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - }, - "candidate_id": "rlkernel.cuda.qwenimage_multi_axis_rope.v1", - "device": "cuda", - "compute_capability": "sm120", - "seed": 20260812, - "fallback_reason": null, - "passed": true, - "provenance_valid": true, - "metadata_valid": true, - "loss_reduction": "sum_over_active_tokens_then_optional_mean_by_active_count", - "active_token_denominator": 27, - "grad_tensor_names": [ - "dx" - ], - "first_failing_op": null, - "first_failing_tensor": null, - "first_failing_config_pair": null, - "observed_kernel_id": "torch.ops.rl_kernel_native.multi_axis_rope" - } -} diff --git a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/ws1-triton.json b/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/ws1-triton.json deleted file mode 100644 index 628fcbf70..000000000 --- a/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120/ws1-triton.json +++ /dev/null @@ -1,1962 +0,0 @@ -{ - "scope": "Qwen-Image multi_axis_rope operator; C1 elementwise and C2 config transforms", - "trace": { - "contract": "qwenimage-interleaved-rope-fp32-v1", - "actual_backend": "rlkernel.triton.qwenimage_multi_axis_rope.v1", - "kernel_id": "rl_engine.backends.shared.triton.rope._flash_attn_rotary.rotary_kernel", - "layout": "BSHD, adjacent channel pairs", - "arithmetic": "fp32 separate multiply/add, one output cast, no reduction", - "table_revision": "40f63690da440e8c382ab53e67508835deb1ed2f" - }, - "passed": true, - "forward": { - "op_name": "multi_axis_rope", - "backend_profile": "triton_cuda_bf16", - "accuracy_reports": [ - { - "config_id": "BN/full", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.054684440139681e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/chunked", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/chunked", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.054684440139681e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s0", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/full/s0", - "fp32_reference" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004673302173614502, - "mean_abs_error": 6.917398422956467e-05, - "max_rel_error": 0.003753266530111432, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s1", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/full/s1", - "fp32_reference" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.00048516690731048584, - "mean_abs_error": 7.802648178767413e-05, - "max_rel_error": 0.0037793784867972136, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s2", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/full/s2", - "fp32_reference" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004864335060119629, - "mean_abs_error": 7.697803084738553e-05, - "max_rel_error": 0.003821910358965397, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s3", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/full/s3", - "fp32_reference" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.710371912457049e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s0", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/chunked/s0", - "fp32_reference" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004673302173614502, - "mean_abs_error": 6.917398422956467e-05, - "max_rel_error": 0.003753266530111432, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s1", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/chunked/s1", - "fp32_reference" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.00048516690731048584, - "mean_abs_error": 7.802648178767413e-05, - "max_rel_error": 0.0037793784867972136, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s2", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/chunked/s2", - "fp32_reference" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004864335060119629, - "mean_abs_error": 7.697803084738553e-05, - "max_rel_error": 0.003821910358965397, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s3", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "B1-singleton_aggregate/chunked/s3", - "fp32_reference" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.710371912457049e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/permuted", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/permuted", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.054684440139681e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/padded_right", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/padded_right", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.054684440139681e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/padded_left", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/padded_left", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 0.0004878789186477661, - "mean_abs_error": 8.054684440139681e-05, - "max_rel_error": 0.0038475203327834606, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "forward_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - } - ], - "invariance_reports": [ - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/chunked", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "BN/chunked" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s0", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s0" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s1", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s1" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s2", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s2" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s3", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s3" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s0", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s0" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s1", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s1" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s2", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s2" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s3", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s3" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/permuted", - "transform_kind": "permutation", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "BN/permuted" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/padded_right", - "transform_kind": "padding", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "BN/padded_right" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/padded_left", - "transform_kind": "padding", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "output", - "config_pair": [ - "BN/full", - "BN/padded_left" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "forward_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - } - ], - "logprob_smoke": null, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - }, - "candidate_id": "rlkernel.triton.qwenimage_multi_axis_rope.v1", - "device": "cuda", - "compute_capability": "sm120", - "seed": 20260812, - "fallback_reason": null, - "passed": true, - "provenance_valid": true, - "metadata_valid": true, - "observed_kernel_id": "rl_engine.backends.shared.triton.rope._flash_attn_rotary.rotary_kernel" - }, - "gradient": { - "op_name": "multi_axis_rope", - "backend_profile": "triton_cuda_bf16", - "accuracy_reports": [ - { - "config_id": "BN/full", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.699562850873917e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/chunked", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/chunked", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.699562850873917e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s0", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/full/s0", - "fp32_reference" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 2.7887988835573196e-05, - "mean_abs_error": 3.6821973026235355e-06, - "max_rel_error": 1.8164952993392944, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s1", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/full/s1", - "fp32_reference" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.7113780965446495e-06, - "max_rel_error": 0.4316565990447998, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s2", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/full/s2", - "fp32_reference" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 2.995925024151802e-05, - "mean_abs_error": 4.716822331829462e-06, - "max_rel_error": 0.7439075112342834, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/full/s3", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/full/s3", - "fp32_reference" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.0211638659238815e-05, - "mean_abs_error": 4.9605891945248e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s0", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/chunked/s0", - "fp32_reference" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 2.7887988835573196e-05, - "mean_abs_error": 3.6821973026235355e-06, - "max_rel_error": 1.8164952993392944, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s1", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/chunked/s1", - "fp32_reference" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.7113780965446495e-06, - "max_rel_error": 0.4316565990447998, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s2", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/chunked/s2", - "fp32_reference" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 2.995925024151802e-05, - "mean_abs_error": 4.716822331829462e-06, - "max_rel_error": 0.7439075112342834, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "B1-singleton_aggregate/chunked/s3", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "B1-singleton_aggregate/chunked/s3", - "fp32_reference" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.0211638659238815e-05, - "mean_abs_error": 4.9605891945248e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/permuted", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/permuted", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.699562850873917e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/padded_right", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/padded_right", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.699562850873917e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - }, - { - "config_id": "BN/padded_left", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/padded_left", - "fp32_reference" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "float32", - "max_abs_error": 3.213249146938324e-05, - "mean_abs_error": 4.699562850873917e-06, - "max_rel_error": 5.814913272857666, - "atol": 0.02, - "rtol": 0.016, - "passed": true, - "judgment": "gradient_accuracy", - "comparison_lhs_role": "bf16_candidate", - "comparison_rhs_role": "fp32_reference" - } - ], - "passed": true, - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - } - } - ], - "invariance_reports": [ - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/chunked", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "BN/chunked" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s0", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s0" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s1", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s1" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s2", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s2" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/full/s3", - "transform_kind": "batch_size", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/full/s3" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s0", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s0" - ], - "shape": [ - 3, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s1", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s1" - ], - "shape": [ - 8, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s2", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s2" - ], - "shape": [ - 5, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "B1-singleton_aggregate/chunked/s3", - "transform_kind": "chunk", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "B1-singleton_aggregate/chunked/s3" - ], - "shape": [ - 11, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/permuted", - "transform_kind": "permutation", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "BN/permuted" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/padded_right", - "transform_kind": "padding", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "BN/padded_right" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - }, - { - "canonical_config_id": "BN/full", - "transformed_config_id": "BN/padded_left", - "transform_kind": "padding", - "op_class": "elementwise", - "dtype": "bfloat16", - "backend_profile": "triton_cuda_bf16", - "details": [ - { - "tensor_name": "dx", - "config_pair": [ - "BN/full", - "BN/padded_left" - ], - "shape": [ - 27, - 3, - 128 - ], - "dtype": "bfloat16", - "max_abs_error": 0.0, - "mean_abs_error": 0.0, - "max_rel_error": 0.0, - "atol": 0.0, - "rtol": 0.0, - "passed": true, - "judgment": "gradient_invariance", - "comparison_lhs_role": "transformed_config", - "comparison_rhs_role": "canonical_config" - } - ], - "passed": true - } - ], - "singleton_aggregate_reports": [], - "backend_provenance": { - "backend_profile": "triton_cuda_bf16", - "requested_backend": "triton", - "actual_backend": "triton", - "execution_dtype": "bfloat16", - "accumulation_dtype": "float32", - "output_dtype": "bfloat16", - "reference_dtype": "float32", - "candidate_tf32_enabled": false, - "reference_tf32_enabled": false - }, - "candidate_id": "rlkernel.triton.qwenimage_multi_axis_rope.v1", - "device": "cuda", - "compute_capability": "sm120", - "seed": 20260812, - "fallback_reason": null, - "passed": true, - "provenance_valid": true, - "metadata_valid": true, - "loss_reduction": "sum_over_active_tokens_then_optional_mean_by_active_count", - "active_token_denominator": 27, - "grad_tensor_names": [ - "dx" - ], - "first_failing_op": null, - "first_failing_tensor": null, - "first_failing_config_pair": null, - "observed_kernel_id": "rl_engine.backends.shared.triton.rope._flash_attn_rotary.rotary_kernel" - } -} From 9e2e50fcd58e426d5fb67f1e1a2b79ef2ceaa5ef Mon Sep 17 00:00:00 2001 From: dandan20031031 <2425926169@qq.com> Date: Wed, 28 Oct 2026 23:35:02 +0800 Subject: [PATCH 5/5] test(rope): align validation fixtures and cover review findings Signed-off-by: dandan20031031 <2425926169@qq.com> --- csrc/cuda/rope/multi_axis_rope.cu | 13 ++-- docs/operators/multi-axis-rope.md | 11 +++ .../qwenimage_multi_axis_rope_sm120.md | 39 +++++++++- rl_engine/ops/rope/_registration.py | 3 + .../validation/operators/gradient_adapters.py | 16 +++- .../validation/operators/operator_inputs.py | 5 +- rl_engine/validation/operators/rope_inputs.py | 21 +++++ tests/ops/rope/test_multi_axis_rope.py | 78 ++++++++++++++++++- .../operators/test_operator_inputs.py | 65 ++++++++++++++++ 9 files changed, 236 insertions(+), 15 deletions(-) create mode 100644 rl_engine/validation/operators/rope_inputs.py diff --git a/csrc/cuda/rope/multi_axis_rope.cu b/csrc/cuda/rope/multi_axis_rope.cu index 3de49caba..27459c5fb 100644 --- a/csrc/cuda/rope/multi_axis_rope.cu +++ b/csrc/cuda/rope/multi_axis_rope.cu @@ -9,6 +9,9 @@ #include namespace { +constexpr int64_t kHeadDim = 128; +constexpr int64_t kPairsPerHead = kHeadDim / 2; + struct Strides { int64_t xb, xs, xh, xd; int64_t cb, cs, cd, sb, ss, sd; @@ -21,8 +24,8 @@ __global__ void rotate_pairs(const scalar_t* x, const float* cos, const float* s const int64_t step = static_cast(blockDim.x) * gridDim.x; for (int64_t i = static_cast(blockIdx.x) * blockDim.x + threadIdx.x; i < n; i += step) { - const int64_t pair = i % 64; - const int64_t row = i / 64; + const int64_t pair = i % kPairsPerHead; + const int64_t row = i / kPairsPerHead; const int64_t head = row % heads; const int64_t token = row / heads; const int64_t batch = token / seq; @@ -43,15 +46,15 @@ at::Tensor multi_axis_rope_apply(const at::Tensor& x, const at::Tensor& cos, const at::Tensor& sin, bool transpose) { TORCH_CHECK(x.is_cuda() && cos.is_cuda() && sin.is_cuda(), "multi_axis_rope: CUDA tensors required"); TORCH_CHECK(x.device() == cos.device() && x.device() == sin.device(), "multi_axis_rope: device mismatch"); - TORCH_CHECK(x.dim() == 4 && x.size(3) == 128 && x.size(2) > 0, + TORCH_CHECK(x.dim() == 4 && x.size(3) == kHeadDim && x.size(2) > 0, "multi_axis_rope: x must be [B, S, H, 128], H > 0"); TORCH_CHECK(x.scalar_type() == at::kFloat || x.scalar_type() == at::kHalf || x.scalar_type() == at::kBFloat16, "multi_axis_rope: unsupported input dtype"); TORCH_CHECK(cos.scalar_type() == at::kFloat && sin.scalar_type() == at::kFloat, "multi_axis_rope: tables must be fp32"); TORCH_CHECK(cos.sizes() == sin.sizes(), "multi_axis_rope: table shape mismatch"); - TORCH_CHECK((cos.dim() == 2 && cos.size(0) == x.size(1) && cos.size(1) == 64) || - (cos.dim() == 3 && cos.size(0) == x.size(0) && cos.size(1) == x.size(1) && cos.size(2) == 64), + TORCH_CHECK((cos.dim() == 2 && cos.size(0) == x.size(1) && cos.size(1) == kPairsPerHead) || + (cos.dim() == 3 && cos.size(0) == x.size(0) && cos.size(1) == x.size(1) && cos.size(2) == kPairsPerHead), "multi_axis_rope: tables must be [S, 64] or [B, S, 64]"); TORCH_CHECK(!cos.requires_grad() && !sin.requires_grad(), "multi_axis_rope: table gradients unsupported"); const c10::cuda::CUDAGuard guard(x.device()); diff --git a/docs/operators/multi-axis-rope.md b/docs/operators/multi-axis-rope.md index 16899a755..a7c8f6495 100644 --- a/docs/operators/multi-axis-rope.md +++ b/docs/operators/multi-axis-rope.md @@ -172,6 +172,17 @@ with C1's elementwise tolerance profile and C2's logical configuration transform This gate qualifies the Qwen-Image operator and does not certify the Dense Qwen3 chain, distributed execution, or model assembly. +The standard CLI and WS1 forward/gradient adapters share a centered image-grid +prefix from `qwen_image_positions`. The square grid contains the full logical +sample; a nonsquare sequence length uses its row-major prefix. WS1 selects rows +by `(sample_id, token_position)` from that sample's grid, so physical chunking, +reordering and padding cannot change active-token coordinates. Padding uses +zero coordinates. These are operator fixtures, not a model's image/text assembly. + +Backward provenance is recorded inside the opaque custom-op implementation, +outside fake/AOT tracing. Its Python `execution_count` is not a CUDA Graph replay +counter: replay executes the captured GPU kernels without re-entering Python. + For a complete native package installation on Linux, run the dedicated gate: ```bash diff --git a/reports/experiments/qwenimage_multi_axis_rope_sm120.md b/reports/experiments/qwenimage_multi_axis_rope_sm120.md index a7d82f65b..a827bb42e 100644 --- a/reports/experiments/qwenimage_multi_axis_rope_sm120.md +++ b/reports/experiments/qwenimage_multi_axis_rope_sm120.md @@ -1,7 +1,7 @@ # Qwen-Image multi-axis RoPE: SM120 validation Recorded 2026-10-28 for [PR #623](https://github.com/RL-Align/RL-Kernel/pull/623). -The numerical/native implementation is based on `ac17ab87e8d93dabe93dfdc134e9853d47de7af1`. +The rotation arithmetic is based on `ac17ab87e8d93dabe93dfdc134e9853d47de7af1`. The [operator documentation](../../docs/operators/multi-axis-rope.md) defines the contract, upstream reuse, API and detailed reproduction commands. @@ -25,7 +25,7 @@ The normal full wheel was built with `RL_KERNEL_ROPE_ONLY`, forced SM90 and fast-math flags unset. Linking existing CUDA driver calls used `LIBRARY_PATH=/usr/lib/wsl/lib`; the Windows driver supplies CUDA driver access. -## Recorded results +## Initial qualification | Check | WSL2 Linux | Windows | | --- | --- | --- | @@ -51,6 +51,37 @@ The five historical Windows layout failures also reproduce on its clean base placeholders. The later Linux run preserved all 27 tracked symlinks and passed. The earlier Windows logs remain unchanged in the archive. +## Automated-review follow-up + +The standard CLI and WS1 adapters now share a centered image-grid prefix, +selected by full logical sample length and token identity. Eleven CPU regression +cases cover empty/square/nonsquare lengths and packed/permuted/chunked/padded +layouts. Before the change, ten of those cases failed; all eleven now pass. +Additional tests pin non-CUDA reference routing, library-load error propagation, +shared-table broadcasting, and head-dimension stride 2 or outermost storage. +Native shape checks and indexing share named dimension constants; the rotation +arithmetic is unchanged. Backward telemetry documents its graph-replay limits. + +The following are new local WSL2 runs of the updated implementation and fixtures: + +| Check | Result | +| --- | --- | +| Rebuilt ordinary complete CUDA wheel; fresh installed-package eager/fullgraph forward/input gradient | Passed for both GPU backends | +| Public operator suite | 105 passed, 1 two-GPU guard skip | +| CUDA/Triton WS1 accuracy/invariance and ordinary CLI checks | Both passed with explicit backend provenance | +| Related regressions, including the eleven new input cases | 206 passed, 7 Ascend skips | +| Repository layout with native extension enforcement | 37 passed | +| CPU-only PyTorch 2.9.1+cpu, Python 3.10.22 and 3.12.3 | Each: 80 passed, 64 expected GPU skips; installed reference forward/backward passed | +| Expanded Compute Sanitizer subset, including the new stride cases | 52 passed, zero errors | +| Changed-file checks and strict docs | Passed locally | + +The full wheel contains the shared validation helper and the FlashAttention BSD +license. Wheel SHA256: +`9feb8200b948021c268c933f93efcd424b6b68d6c55bc24c3004d6336deb3a13`. +Non-CUDA routing tests exercise candidate selection with CPU tensors; they do +not certify NPU/ROCm/MUSA hardware. The benchmark below retains its initial-run +measurements; no new timing comparison was performed for this follow-up. + ## Standalone benchmark WSL2 BF16, batch 1, 24 heads, D=128. CUDA Graph medians in ms/call; @@ -115,8 +146,8 @@ source/binary SHA256 manifests remain available at immutable commit - [WSL2 complete-extension evidence](https://github.com/dandan20031031/RL-Kernel/tree/84751473006401c19c1c8c17844e8ef8b1c44375/reports/experiments/qwenimage_multi_axis_rope_wsl2_sm120) - [Historical Windows operator-profile evidence](https://github.com/dandan20031031/RL-Kernel/tree/84751473006401c19c1c8c17844e8ef8b1c44375/reports/experiments/qwenimage_multi_axis_rope_v2_sm120) -This summary preserves those recorded results; consolidating reports does not -rerun the numerical suites. Bare-metal Linux, A100/H100, two-GPU guards, +Those archives identify the initial runs; the follow-up table records subsequent +reruns. Bare-metal Linux, A100/H100, two-GPU guards, QK-norm fusion, model/block assembly and trajectories remain unqualified. Building the full library does not certify every native operator or Dense Qwen3. No distributed or cross-version/device byte-equality claim is made. diff --git a/rl_engine/ops/rope/_registration.py b/rl_engine/ops/rope/_registration.py index 343d41a9c..33798a03e 100644 --- a/rl_engine/ops/rope/_registration.py +++ b/rl_engine/ops/rope/_registration.py @@ -15,6 +15,9 @@ def rotation( validate_inputs(x, cos, sin) result = implementation(x, cos, sin, transpose) if transpose: + # Keep telemetry inside the opaque custom op; fake/AOT tracing + # must not execute this Python lock. CUDA Graph replay runs only + # captured kernels, so execution_count is not a replay counter. record_backward( "multi_axis_rope", kernel_id=kernel_id, diff --git a/rl_engine/validation/operators/gradient_adapters.py b/rl_engine/validation/operators/gradient_adapters.py index 6bc6e8104..48da77045 100644 --- a/rl_engine/validation/operators/gradient_adapters.py +++ b/rl_engine/validation/operators/gradient_adapters.py @@ -32,6 +32,7 @@ MissingBackwardError, ) from rl_engine.validation.operators.operator_specs import OP_SPECS, _load_object +from rl_engine.validation.operators.rope_inputs import qwen_image_grid_prefix REPO_ROOT = Path(__file__).resolve().parents[3] @@ -733,8 +734,19 @@ def _row_inputs( if head_dim != 128: raise ValueError("multi_axis_rope gradient adapter requires head_dim=128") rows = _stack_rows(keys, leading, (n_heads, head_dim), device=device, dtype=dtype) - positions = _row_positions(keys, device=torch.device("cpu")) - coords = torch.stack((positions % 7, positions % 127 - 63, positions % 61 - 30), dim=-1) + # Use the complete logical sample, not this physical call's rows, to + # keep grid centers stable under chunks, permutations and padding. + lengths: dict[str, int] = {} + for sample_id, position in tokens: + lengths[sample_id] = max(lengths.get(sample_id, 0), position + 1) + grids: dict[str, torch.Tensor] = {} + coords = torch.zeros((len(keys), 3), device="cpu", dtype=torch.int64) + for row, key in enumerate(keys): + if key is not None: + sample_id, position = key + if sample_id not in grids: + grids[sample_id] = qwen_image_grid_prefix(lengths[sample_id]) + coords[row] = grids[sample_id][position] cos, sin = build_multi_axis_cos_sin(coords, device=device) return {"x": rows.unsqueeze(0), "cos": cos, "sin": sin} if op_name == "rope": diff --git a/rl_engine/validation/operators/operator_inputs.py b/rl_engine/validation/operators/operator_inputs.py index 137d8d69b..0306a683d 100644 --- a/rl_engine/validation/operators/operator_inputs.py +++ b/rl_engine/validation/operators/operator_inputs.py @@ -8,6 +8,8 @@ import torch +from rl_engine.validation.operators.rope_inputs import qwen_image_grid_prefix + DEFAULT_HIDDEN = 4096 DEFAULT_N_HEADS = 32 DEFAULT_N_KV_HEADS = 8 @@ -259,8 +261,7 @@ def _make_multi_axis_rope_inputs(args, dtype, device): if _arg_int(args, "head_dim", 128) != 128: raise ValueError("multi_axis_rope requires head_dim=128") heads = _arg_int(args, "n_heads", 24) - row = torch.arange(seq, device="cpu", dtype=torch.int64) - positions = torch.stack((row % 7, (row // 127) % 4096, row % 127 - 63), dim=-1) + positions = qwen_image_grid_prefix(seq) cos, sin = build_multi_axis_cos_sin(positions, device=device) return { "x": _floating_tensor((batch, seq, heads, 128), args, dtype, device, offset=0), diff --git a/rl_engine/validation/operators/rope_inputs.py b/rl_engine/validation/operators/rope_inputs.py new file mode 100644 index 000000000..5e22f977e --- /dev/null +++ b/rl_engine/validation/operators/rope_inputs.py @@ -0,0 +1,21 @@ +# SPDX-License-Identifier: Apache-2.0 +"""Shared Qwen-Image coordinates for operator and logical-layout validation.""" + +import math + +import torch + +from rl_engine.ops.rope.multi_axis_rope import qwen_image_positions + + +def qwen_image_grid_prefix(sequence_length: int) -> torch.Tensor: + """Use a centered square-grid prefix for a complete logical sequence. + + CLI sequence lengths need not be square. Grid size depends on the full + sample length, never the physical batch, padding or current chunk length. + """ + if sequence_length < 0: + raise ValueError("sequence_length must be nonnegative") + side = math.isqrt(max(0, sequence_length - 1)) + 1 + image, _ = qwen_image_positions(side, side, 0) + return image[:sequence_length] diff --git a/tests/ops/rope/test_multi_axis_rope.py b/tests/ops/rope/test_multi_axis_rope.py index 6cbcb2b85..9f30c35b0 100644 --- a/tests/ops/rope/test_multi_axis_rope.py +++ b/tests/ops/rope/test_multi_axis_rope.py @@ -3,6 +3,7 @@ import gc import math +from types import SimpleNamespace import pytest import torch @@ -86,12 +87,15 @@ def test_forward_backward_accuracy(implementation, dtype, batched): result = op(x, cos, sin) dx = torch.autograd.grad(result, x, grad)[0] tolerance = {torch.float32: 8e-7, torch.float16: 0.002, torch.bfloat16: 0.02}[dtype] + # Make the oracle's batch axis explicit when the candidate reads shared tables. + c = cos if batched else cos.unsqueeze(0).expand(x.shape[0], -1, -1) + s = sin if batched else sin.unsqueeze(0).expand(x.shape[0], -1, -1) torch.testing.assert_close( - result.double(), complex_reference(x, cos, sin), atol=tolerance, rtol=tolerance + result.double(), complex_reference(x, c, s), atol=tolerance, rtol=tolerance ) torch.testing.assert_close( dx.double(), - complex_reference(grad, cos, sin, transpose=True), + complex_reference(grad, c, s, transpose=True), atol=tolerance, rtol=tolerance, ) @@ -145,6 +149,33 @@ def test_noncontiguous_input_tables_and_gradient(implementation): assert_bits(dx, dd) +@pytest.mark.parametrize("dtype", [torch.float32, torch.float16, torch.bfloat16]) +@pytest.mark.parametrize("layout", ["head_dim_stride_2", "head_dim_outermost"]) +def test_arbitrary_head_dimension_stride_and_gradient(implementation, dtype, layout): + op, device = implementation + _, cos, sin = inputs(device, dtype) + if layout == "head_dim_stride_2": + x = torch.randn(2, 7, 3, 256, device=device, dtype=dtype)[..., ::2] + assert x.stride(-1) == 2 + else: + x = torch.randn(128, 2, 7, 3, device=device, dtype=dtype).permute(1, 2, 3, 0) + assert x.stride(-1) > max(x.stride()[:-1]) + x.requires_grad_() + grad = torch.randn(2, 7, 3, 256, device=device, dtype=dtype)[..., ::2] + result = op(x, cos, sin) + dx = torch.autograd.grad(result, x, grad)[0] + tolerance = {torch.float32: 8e-7, torch.float16: 0.002, torch.bfloat16: 0.02}[dtype] + torch.testing.assert_close( + result.double(), complex_reference(x, cos, sin), atol=tolerance, rtol=tolerance + ) + torch.testing.assert_close( + dx.double(), + complex_reference(grad, cos, sin, transpose=True), + atol=tolerance, + rtol=tolerance, + ) + + @pytest.mark.parametrize("batch,seq", [(0, 7), (2, 0), (0, 0)]) def test_empty_forward_backward(implementation, batch, seq): op, device = implementation @@ -209,6 +240,49 @@ def test_explicit_backend_and_reference_dispatch(monkeypatch): assert_bits(multi_axis_rope(x, cos, sin), op(x, cos, sin)) +@pytest.mark.parametrize("platform", ["cpu", "npu", "rocm", "musa"]) +def test_non_cuda_registry_keeps_multi_axis_reference(platform, monkeypatch): + """Exercise candidate selection on CPU; no accelerator execution is certified.""" + from rl_engine.runtime.registry import KernelRegistry, OpBackend + + registry = KernelRegistry() + assert registry._priority_map[platform]["multi_axis_rope"] == [ + OpBackend.PYTORCH_MULTI_AXIS_ROPE + ] + monkeypatch.setattr(registry, "_platform", lambda: platform) + op = registry.get_op("multi_axis_rope") + assert op.trace()["actual_backend"] == "rlkernel.reference.qwenimage_multi_axis_rope.v1" + x, cos, sin = inputs("cpu") + assert_bits(op(x, cos, sin), get_multi_axis_rope_op("reference")(x, cos, sin)) + + +@pytest.mark.parametrize("exception_type", [OSError, RuntimeError]) +def test_operator_wheel_load_error_propagates(exception_type, monkeypatch): + from rl_engine.backends.cuda.rope import multi_axis_rope as cuda_module + + error = exception_type("probe library load failure") + + def fail_load(path): + raise error + + monkeypatch.setattr( + cuda_module, + "torch", + SimpleNamespace( + version=SimpleNamespace(hip=None), + ops=SimpleNamespace(rl_kernel_native=SimpleNamespace(), load_library=fail_load), + ), + ) + monkeypatch.setattr( + cuda_module.importlib.util, + "find_spec", + lambda name: SimpleNamespace(origin="probe.so"), + ) + with pytest.raises(exception_type, match="probe library load failure") as caught: + cuda_module.native_available() + assert caught.value is error + + def test_cuda_registry_uses_native_when_present(): from rl_engine.backends.cuda.rope.multi_axis_rope import native_available from rl_engine.runtime.registry import KernelRegistry diff --git a/tests/validation/operators/test_operator_inputs.py b/tests/validation/operators/test_operator_inputs.py index b2c054b18..0416d1540 100644 --- a/tests/validation/operators/test_operator_inputs.py +++ b/tests/validation/operators/test_operator_inputs.py @@ -8,6 +8,9 @@ import pytest import torch +from rl_engine.config.workload import LogicalSample +from rl_engine.ops.rope.multi_axis_rope import build_multi_axis_cos_sin, qwen_image_positions +from rl_engine.validation.operators.gradient_adapters import _row_inputs from rl_engine.validation.operators.operator_inputs import make_operator_inputs, operator_shape_name from rl_engine.validation.operators.operator_specs import ( make_candidate, @@ -121,6 +124,68 @@ def test_multi_axis_rope_operator_spec_and_seeded_inputs(): ) +@pytest.mark.parametrize( + "seq,side", [(0, 1), (1, 1), (7, 3), (15, 4), (16, 4), (4096, 64), (6889, 83)] +) +def test_multi_axis_rope_validation_uses_centered_image_grid_prefix(seq, side): + inputs = make_operator_inputs( + "multi_axis_rope", _args(seq=seq, n_heads=1), torch.float32, torch.device("cpu") + ) + image, _ = qwen_image_positions(side, side, 0) + cos, sin = build_multi_axis_cos_sin(image[:seq]) + assert torch.equal(inputs["cos"], cos) + assert torch.equal(inputs["sin"], sin) + + +@pytest.mark.parametrize("layout", ["packed", "permuted", "chunked", "padded"]) +def test_multi_axis_rope_validation_preserves_logical_tables_across_layouts(layout): + samples = ( + LogicalSample("short", tuple(range(7)), prompt_len=1, seq_len=7), + LogicalSample("long", tuple(range(15)), prompt_len=2, seq_len=15), + ) + tokens = { + (token.sample_id, token.token_position): token + for sample in samples + for token in sample.tokens() + } + expected = { + sample.sample_id: make_operator_inputs( + "multi_axis_rope", + _args(seq=sample.seq_len, n_heads=1), + torch.float32, + torch.device("cpu"), + ) + for sample in samples + } + keys = list(tokens) + if layout == "permuted": + keys.reverse() + elif layout == "chunked": + keys = [("short", 5), ("short", 6), ("long", 12)] + elif layout == "padded": + keys = [None, ("long", 4), None, ("short", 6), ("short", 0), None] + inputs = _row_inputs( + "multi_axis_rope", + keys, + tokens, + {}, + device=torch.device("cpu"), + dtype=torch.float32, + hidden=1, + vocab_size=17, + n_heads=1, + head_dim=128, + ) + for row, key in enumerate(keys): + if key is None: + assert torch.equal(inputs["cos"][row], torch.ones(64)) + assert torch.equal(inputs["sin"][row], torch.zeros(64)) + else: + sample_id, position = key + for name in ("cos", "sin"): + assert torch.equal(inputs[name][row], expected[sample_id][name][position]) + + def test_constant_linear_logp_inputs_match_operator_contract(): args = _args(input_mode="constant", constant_value=0.5, token_value=3) inputs = make_operator_inputs("linear_logp", args, torch.float32, torch.device("cpu"))