benchmarks/benchmark_native_parity.py

#!/usr/bin/env python3
"""World-model Conv benchmark with native, wrapper, compile, and cuDNN paths."""

from __future__ import annotations

import argparse
import json
import os
import sys
from dataclasses import asdict, dataclass
from pathlib import Path

import torch
import torch.nn.functional as F


PACKAGE = Path(__file__).resolve().parents[1]
sys.path.insert(0, str(PACKAGE / "tests"))
from test_world_model_conv import (  # noqa: E402
    dequantize_linear_nvfp4,
    load_installed_ops,
    load_source_ops,
    quantize_conv_tensor,
)


CONV3D_SHAPES = {
    "causal-c32": (1, 2, 4, 16, 16, 32, 32),
    "causal-small": (1, 2, 4, 16, 16, 64, 64),
}
NVFP4_CONV3D_SHAPES = {
    "nvfp4-c64": (1, 2, 4, 16, 16, 64, 64),
    "nvfp4-c128": (1, 2, 4, 16, 16, 128, 128),
    "nvfp4-c512": (1, 2, 4, 16, 16, 512, 512),
}
CONV2D_SHAPES = {
    "resample-c64": (4, 32, 32, 64, 64),
    "resample-c320": (17, 32, 32, 320, 320),
}


@dataclass
class Result:
    workload: str
    shape: str
    native_us: float
    wrapper_us: float
    wrapper_native: float
    eager_cudnn_us: float
    compile_cudnn_us: float
    diagnostic_predequant_cudnn_us: float | None
    diagnostic_predequant_compile_us: float | None
    max_abs: float
    mean_abs: float
    p99_abs: float
    cosine: float
    accepted: bool


def bench(fn, warmup, iters):
    for _ in range(warmup):
        fn()
    torch.cuda.synchronize()
    start = torch.cuda.Event(enable_timing=True)
    end = torch.cuda.Event(enable_timing=True)
    start.record()
    for _ in range(iters):
        fn()
    end.record()
    torch.cuda.synchronize()
    return start.elapsed_time(end) * 1000.0 / iters


def build_native():
    from torch.utils.cpp_extension import load

    os.environ.setdefault("TORCH_CUDA_ARCH_LIST", "12.0a")
    return load(
        name="world_model_conv_raw_native",
        sources=[
            str(PACKAGE / "benchmarks/native_binding.cpp"),
            str(PACKAGE / "csrc/fp8_conv3d_sm120_v18.cu"),
            str(PACKAGE / "csrc/fp8_causal_conv3d_sm120.cu"),
            str(PACKAGE / "csrc/fp8_conv2d_3x3_sm120.cu"),
            str(PACKAGE / "csrc/nvfp4_causal_conv3d_sm120.cu"),
            str(PACKAGE / "csrc/nvfp4_causal_conv3d_residual_sm120.cu"),
            str(PACKAGE / "csrc/nvfp4_causal_conv3d_residual_k128_sm120.cu"),
        ],
        extra_include_paths=[str(PACKAGE / "csrc")],
        extra_cflags=["-O3"],
        extra_cuda_cflags=["-O3"],
        verbose=False,
    )


def metrics(got, ref):
    diff = (got.float() - ref.float()).abs().flatten()
    cosine = F.cosine_similarity(
        got.float().flatten(), ref.float().flatten(), dim=0
    ).item()
    return (
        diff.max().item(),
        diff.mean().item(),
        torch.quantile(diff, 0.99).item(),
        cosine,
    )


def source_call(ops, name, *args, out):
    if hasattr(ops, "_ops"):
        getattr(ops._ops, name)(*args, out)
    else:
        getattr(ops, name)(*args, out=out)


def run_conv3d(ops, native, label, shape, args):
    n, tc, tn, h, w, ci, co = shape
    cache = (torch.randn((n, tc, h, w, ci), device="cuda") * 0.1).to(
        torch.float8_e4m3fn
    )
    new = (torch.randn((n, tn, h, w, ci), device="cuda") * 0.1).to(
        torch.float8_e4m3fn
    )
    weight = (torch.randn((co, 3, 3, 3, ci), device="cuda") * 0.1).to(
        torch.float8_e4m3fn
    )
    bias = (torch.randn(co, device="cuda") * 0.01).to(torch.bfloat16)
    out = torch.empty((n, tn, h, w, co), device="cuda", dtype=torch.bfloat16)
    alpha = 0.75

    wrapper = lambda: source_call(
        ops,
        "fp8_causal_conv3d_ndhwc_bf16",
        cache,
        new,
        weight,
        bias,
        alpha,
        out=out,
    )
    raw = lambda: native.causal_conv3d(
        cache, new, weight, bias, alpha, out
    )

    def cudnn_ref():
        x = torch.cat((cache, new), dim=1).float().permute(0, 4, 1, 2, 3)
        wt = weight.float().permute(0, 4, 1, 2, 3)
        y = F.conv3d(x, wt, padding=(0, 1, 1))
        y = y.mul(alpha).add(bias.float().view(1, -1, 1, 1, 1))
        out.copy_(y[:, :, :tn].permute(0, 2, 3, 4, 1).to(torch.bfloat16))

    compiled = torch.compile(cudnn_ref, fullgraph=True)
    wrapper()
    got = out.clone()
    cudnn_ref()
    ref = out.clone()
    max_abs, mean_abs, p99_abs, cosine = metrics(got, ref)
    native_us = bench(raw, args.warmup, args.iters)
    wrapper_us = bench(wrapper, args.warmup, args.iters)
    eager_us = bench(cudnn_ref, args.warmup, args.iters)
    compile_us = bench(compiled, args.warmup, args.iters)
    return Result(
        label,
        str(shape),
        native_us,
        wrapper_us,
        wrapper_us / native_us,
        eager_us,
        compile_us,
        None,
        None,
        max_abs,
        mean_abs,
        p99_abs,
        cosine,
        wrapper_us - native_us <= max(0.5, native_us * 0.05)
        and wrapper_us <= min(eager_us, compile_us) * 0.98
        and cosine >= 0.999
        and mean_abs <= 0.01,
    )


def run_conv3d_residual(ops, native, label, shape, args):
    n, tc, tn, h, w, ci, co = shape
    cache = (torch.randn((n, tc, h, w, ci), device="cuda") * 0.1).to(
        torch.float8_e4m3fn
    )
    new = (torch.randn((n, tn, h, w, ci), device="cuda") * 0.1).to(
        torch.float8_e4m3fn
    )
    weight = (torch.randn((co, 3, 3, 3, ci), device="cuda") * 0.1).to(
        torch.float8_e4m3fn
    )
    bias = (torch.randn(co, device="cuda") * 0.01).to(torch.bfloat16)
    residual = torch.randn(
        (n, co, tn, h, w), device="cuda", dtype=torch.bfloat16
    )
    out = torch.empty_like(residual)
    alpha = 0.75
    wrapper = lambda: source_call(
        ops,
        "fp8_conv3d_v18_ncdhw_res_bf16out",
        cache,
        new,
        weight,
        bias,
        residual,
        alpha,
        out=out,
    )
    raw = lambda: native.causal_conv3d_residual(
        cache, new, weight, bias, residual, alpha, out
    )

    def cudnn_ref():
        x = torch.cat((cache, new), dim=1).float().permute(0, 4, 1, 2, 3)
        wt = weight.float().permute(0, 4, 1, 2, 3)
        y = F.conv3d(x, wt, padding=(0, 1, 1))
        y = y.mul(alpha).add(bias.float().view(1, -1, 1, 1, 1))
        y = (
            y[:, :, :tn].to(torch.bfloat16).float()
            + residual.float()
        ).to(torch.bfloat16)
        out.copy_(y)

    compiled = torch.compile(cudnn_ref, fullgraph=True)
    wrapper()
    got = out.clone()
    cudnn_ref()
    ref = out.clone()
    max_abs, mean_abs, p99_abs, cosine = metrics(got, ref)
    native_us = bench(raw, args.warmup, args.iters)
    wrapper_us = bench(wrapper, args.warmup, args.iters)
    eager_us = bench(cudnn_ref, args.warmup, args.iters)
    compile_us = bench(compiled, args.warmup, args.iters)
    return Result(
        f"{label}-residual",
        str(shape),
        native_us,
        wrapper_us,
        wrapper_us / native_us,
        eager_us,
        compile_us,
        None,
        None,
        max_abs,
        mean_abs,
        p99_abs,
        cosine,
        wrapper_us - native_us <= max(0.5, native_us * 0.05)
        and wrapper_us <= min(eager_us, compile_us) * 0.98
        and cosine >= 0.999
        and mean_abs <= 0.01,
    )


def run_conv2d(ops, native, label, shape, args):
    n, h, w, ci, co = shape
    input = (torch.randn((n, h, w, ci), device="cuda") * 0.1).to(
        torch.float8_e4m3fn
    )
    weight = (torch.randn((co, 3, 3, ci), device="cuda") * 0.1).to(
        torch.float8_e4m3fn
    )
    bias = (torch.randn(co, device="cuda") * 0.01).to(torch.bfloat16)
    out = torch.empty((n, h, w, co), device="cuda", dtype=torch.bfloat16)
    alpha = 0.75
    wrapper = lambda: source_call(
        ops,
        "fp8_conv2d_3x3_nhwc_bf16",
        input,
        weight,
        bias,
        alpha,
        out=out,
    )
    raw = lambda: native.conv2d(input, weight, bias, alpha, out)

    def cudnn_ref():
        x = input.float().permute(0, 3, 1, 2)
        wt = weight.float().permute(0, 3, 1, 2)
        y = F.conv2d(x, wt, padding=1).mul(alpha)
        y = y.add(bias.float().view(1, -1, 1, 1))
        out.copy_(y.permute(0, 2, 3, 1).to(torch.bfloat16))

    compiled = torch.compile(cudnn_ref, fullgraph=True)
    wrapper()
    got = out.clone()
    cudnn_ref()
    ref = out.clone()
    max_abs, mean_abs, p99_abs, cosine = metrics(got, ref)
    native_us = bench(raw, args.warmup, args.iters)
    wrapper_us = bench(wrapper, args.warmup, args.iters)
    eager_us = bench(cudnn_ref, args.warmup, args.iters)
    compile_us = bench(compiled, args.warmup, args.iters)
    return Result(
        label,
        str(shape),
        native_us,
        wrapper_us,
        wrapper_us / native_us,
        eager_us,
        compile_us,
        None,
        None,
        max_abs,
        mean_abs,
        p99_abs,
        cosine,
        wrapper_us - native_us <= max(0.5, native_us * 0.05)
        and wrapper_us <= min(eager_us, compile_us) * 0.98
        and cosine >= 0.999
        and mean_abs <= 0.01,
    )


def run_nvfp4_conv3d(ops, native, label, shape, args, residual_path):
    n, tc, tn, h, w, ci, co = shape
    cache_bf16 = (
        torch.randn((n, tc, h, w, ci), device="cuda") * 0.1
    ).to(torch.bfloat16)
    input_bf16 = (
        torch.randn((n, tn, h, w, ci), device="cuda") * 0.1
    ).to(torch.bfloat16)
    weight_bf16 = (
        torch.randn((co, 3, 3, 3, ci), device="cuda") * 0.1
    ).to(torch.bfloat16)
    cache, cache_sf = quantize_conv_tensor(cache_bf16)
    input, input_sf = quantize_conv_tensor(input_bf16)
    weight, weight_sf = quantize_conv_tensor(weight_bf16)
    cache_dequant = dequantize_linear_nvfp4(
        cache.reshape(-1, ci // 2), cache_sf.reshape(-1, ci // 16)
    ).reshape_as(cache_bf16)
    input_dequant = dequantize_linear_nvfp4(
        input.reshape(-1, ci // 2), input_sf.reshape(-1, ci // 16)
    ).reshape_as(input_bf16)
    weight_dequant = dequantize_linear_nvfp4(
        weight.reshape(-1, ci // 2), weight_sf.reshape(-1, ci // 16)
    ).reshape_as(weight_bf16)
    bias = (torch.randn(co, device="cuda") * 0.01).to(torch.bfloat16)
    alpha = 0.75

    if residual_path:
        residual = torch.randn(
            (n, co, tn, h, w), device="cuda", dtype=torch.bfloat16
        )
        out = torch.empty_like(residual)
        wrapper = lambda: source_call(
            ops,
            "nvfp4_causal_conv3d_residual_ncdhw_bf16",
            cache, input, weight, cache_sf, input_sf, weight_sf, bias,
            residual, None, alpha, out=out,
        )
        raw = lambda: native.nvfp4_causal_conv3d_residual(
            cache, input, weight, cache_sf, input_sf, weight_sf, bias,
            residual, alpha, out,
        )
    else:
        residual = None
        out = torch.empty(
            (n, tn, h, w, co), device="cuda", dtype=torch.bfloat16
        )
        wrapper = lambda: source_call(
            ops,
            "nvfp4_causal_conv3d_ndhwc_bf16",
            cache, input, weight, cache_sf, input_sf, weight_sf, bias,
            None, alpha, out=out,
        )
        raw = lambda: native.nvfp4_causal_conv3d(
            cache, input, weight, cache_sf, input_sf, weight_sf, bias,
            alpha, out,
        )

    def store_cudnn_result(cache_value, input_value):
        x = torch.cat((cache_value, input_value), dim=1).permute(
            0, 4, 1, 2, 3
        )
        wt = weight_dequant.permute(0, 4, 1, 2, 3)
        y = F.conv3d(x, wt, padding=(0, 1, 1))[:, :, :tn]
        y = y.mul(alpha).add(bias.float().view(1, -1, 1, 1, 1))
        if residual_path:
            out.copy_(
                (y.to(torch.bfloat16).float() + residual.float()).to(
                    torch.bfloat16
                )
            )
        else:
            out.copy_(y.permute(0, 2, 3, 4, 1).to(torch.bfloat16))

    def predequant_cudnn_ref():
        store_cudnn_result(cache_dequant, input_dequant)

    magnitude = torch.tensor(
        [0.0, 0.5, 1.0, 1.5, 2.0, 3.0, 4.0, 6.0], device="cuda"
    )
    scale_values, scale_bytes = [], []
    for byte in list(range(0x78)) + [0xFE]:
        exponent = (byte >> 3) & 0xF
        mantissa = byte & 0x7
        value = (
            (mantissa / 8.0) * (2.0 ** -6)
            if exponent == 0
            else (1.0 + mantissa / 8.0) * (2.0 ** (exponent - 7))
        )
        scale_values.append(value)
        scale_bytes.append(byte)
    scale_lookup = torch.zeros(256, device="cuda")
    scale_lookup[
        torch.tensor(scale_bytes, device="cuda", dtype=torch.long)
    ] = torch.tensor(scale_values, device="cuda")

    def unpack(packed_value, scale_value):
        low = packed_value & 0xF
        high = packed_value >> 4
        low_value = magnitude[(low & 0x7).long()] * torch.where(
            low & 0x8 != 0, -1.0, 1.0
        )
        high_value = magnitude[(high & 0x7).long()] * torch.where(
            high & 0x8 != 0, -1.0, 1.0
        )
        values = torch.stack((low_value, high_value), dim=-1).flatten(-2)
        scales_value = scale_lookup[scale_value.long()].repeat_interleave(
            16, dim=-1
        )
        return values * scales_value

    def cudnn_ref():
        cache_value = unpack(cache, cache_sf).reshape_as(cache_bf16)
        input_value = unpack(input, input_sf).reshape_as(input_bf16)
        store_cudnn_result(cache_value, input_value)

    compiled = torch.compile(cudnn_ref, fullgraph=True)
    predequant_compiled = torch.compile(predequant_cudnn_ref, fullgraph=True)
    wrapper()
    got = out.clone()
    cudnn_ref()
    ref = out.clone()
    max_abs, mean_abs, p99_abs, cosine = metrics(got, ref)
    native_us = bench(raw, args.warmup, args.iters)
    wrapper_us = bench(wrapper, args.warmup, args.iters)
    eager_us = bench(cudnn_ref, args.warmup, args.iters)
    compile_us = bench(compiled, args.warmup, args.iters)
    diagnostic_eager_us = bench(
        predequant_cudnn_ref, args.warmup, args.iters
    )
    diagnostic_compile_us = bench(
        predequant_compiled, args.warmup, args.iters
    )
    return Result(
        f"{label}{'-residual' if residual_path else ''}",
        str(shape),
        native_us,
        wrapper_us,
        wrapper_us / native_us,
        eager_us,
        compile_us,
        diagnostic_eager_us,
        diagnostic_compile_us,
        max_abs,
        mean_abs,
        p99_abs,
        cosine,
        wrapper_us - native_us <= max(0.5, native_us * 0.05)
        and wrapper_us <= min(eager_us, compile_us) * 0.98
        and cosine >= 0.998
        and mean_abs <= 0.02,
    )


def main():
    parser = argparse.ArgumentParser()
    parser.add_argument("--backend", choices=["source", "installed"], default="source")
    parser.add_argument("--artifact")
    parser.add_argument("--warmup", type=int, default=10)
    parser.add_argument("--iters", type=int, default=30)
    parser.add_argument("--output")
    args = parser.parse_args()
    ops = (
        load_source_ops()
        if args.backend == "source"
        else load_installed_ops(args.artifact)
    )
    native = build_native()
    rows = [
        *(run_conv3d(ops, native, name, shape, args)
          for name, shape in CONV3D_SHAPES.items()),
        *(run_conv3d_residual(ops, native, name, shape, args)
          for name, shape in CONV3D_SHAPES.items()
          if shape[-1] % 8 == 0),
        *(run_conv2d(ops, native, name, shape, args)
          for name, shape in CONV2D_SHAPES.items()),
        *(run_nvfp4_conv3d(ops, native, name, shape, args, False)
          for name, shape in NVFP4_CONV3D_SHAPES.items()),
        *(run_nvfp4_conv3d(ops, native, name, shape, args, True)
          for name, shape in NVFP4_CONV3D_SHAPES.items()),
    ]
    for row in rows:
        print(
            f"{row.workload}: native={row.native_us:.3f}us "
            f"wrapper={row.wrapper_us:.3f}us ({row.wrapper_native:.3f}) "
            f"cuDNN-eager={row.eager_cudnn_us:.3f}us "
            f"cuDNN-compile={row.compile_cudnn_us:.3f}us "
            + (
                f"predequant-compile="
                f"{row.diagnostic_predequant_compile_us:.3f}us "
                if row.diagnostic_predequant_compile_us is not None
                else ""
            )
            + f"cos={row.cosine:.7f} accepted={row.accepted}"
        )
    if args.output:
        path = Path(args.output)
        path.parent.mkdir(parents=True, exist_ok=True)
        path.write_text(json.dumps([asdict(row) for row in rows], indent=2) + "\n")
    if not all(row.accepted for row in rows):
        raise SystemExit("world-model Conv acceptance failed")


if __name__ == "__main__":
    main()