Uploaded using `kernel-builder`.
Browse files- benchmarks/benchmark_native_parity.py +520 -0
- build/torch211-cxx11-cu128-x86_64-linux/__init__.py +331 -1
- build/torch211-cxx11-cu128-x86_64-linux/_ops.py +3 -3
- build/torch211-cxx11-cu128-x86_64-linux/{_world_model_conv_cuda_f14c443.abi3.so → _world_model_conv_cuda_7781728.abi3.so} +2 -2
- build/torch211-cxx11-cu128-x86_64-linux/metadata.json +1 -1
- build/torch211-cxx11-cu130-x86_64-linux/__init__.py +331 -1
- build/torch211-cxx11-cu130-x86_64-linux/_ops.py +3 -3
- build/torch211-cxx11-cu130-x86_64-linux/{_world_model_conv_cuda_f14c443.abi3.so → _world_model_conv_cuda_7781728.abi3.so} +2 -2
- build/torch211-cxx11-cu130-x86_64-linux/metadata.json +1 -1
- build/torch212-cxx11-cu130-x86_64-linux/__init__.py +331 -1
- build/torch212-cxx11-cu130-x86_64-linux/_ops.py +3 -3
- build/torch212-cxx11-cu130-x86_64-linux/{_world_model_conv_cuda_f14c443.abi3.so → _world_model_conv_cuda_7781728.abi3.so} +2 -2
- build/torch212-cxx11-cu130-x86_64-linux/metadata.json +1 -1
- build/torch212-cxx11-cu132-x86_64-linux/__init__.py +331 -1
- build/torch212-cxx11-cu132-x86_64-linux/_ops.py +3 -3
- build/torch212-cxx11-cu132-x86_64-linux/{_world_model_conv_cuda_f14c443.abi3.so → _world_model_conv_cuda_7781728.abi3.so} +2 -2
- build/torch212-cxx11-cu132-x86_64-linux/metadata.json +1 -1
benchmarks/benchmark_native_parity.py
ADDED
|
@@ -0,0 +1,520 @@
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 1 |
+
#!/usr/bin/env python3
|
| 2 |
+
"""World-model Conv benchmark with native, wrapper, compile, and cuDNN paths."""
|
| 3 |
+
|
| 4 |
+
from __future__ import annotations
|
| 5 |
+
|
| 6 |
+
import argparse
|
| 7 |
+
import json
|
| 8 |
+
import os
|
| 9 |
+
import sys
|
| 10 |
+
from dataclasses import asdict, dataclass
|
| 11 |
+
from pathlib import Path
|
| 12 |
+
|
| 13 |
+
import torch
|
| 14 |
+
import torch.nn.functional as F
|
| 15 |
+
|
| 16 |
+
|
| 17 |
+
PACKAGE = Path(__file__).resolve().parents[1]
|
| 18 |
+
sys.path.insert(0, str(PACKAGE / "tests"))
|
| 19 |
+
from test_world_model_conv import ( # noqa: E402
|
| 20 |
+
dequantize_linear_nvfp4,
|
| 21 |
+
load_installed_ops,
|
| 22 |
+
load_source_ops,
|
| 23 |
+
quantize_conv_tensor,
|
| 24 |
+
)
|
| 25 |
+
|
| 26 |
+
|
| 27 |
+
CONV3D_SHAPES = {
|
| 28 |
+
"causal-c32": (1, 2, 4, 16, 16, 32, 32),
|
| 29 |
+
"causal-small": (1, 2, 4, 16, 16, 64, 64),
|
| 30 |
+
}
|
| 31 |
+
NVFP4_CONV3D_SHAPES = {
|
| 32 |
+
"nvfp4-c64": (1, 2, 4, 16, 16, 64, 64),
|
| 33 |
+
"nvfp4-c128": (1, 2, 4, 16, 16, 128, 128),
|
| 34 |
+
"nvfp4-c512": (1, 2, 4, 16, 16, 512, 512),
|
| 35 |
+
}
|
| 36 |
+
CONV2D_SHAPES = {
|
| 37 |
+
"resample-c64": (4, 32, 32, 64, 64),
|
| 38 |
+
"resample-c320": (17, 32, 32, 320, 320),
|
| 39 |
+
}
|
| 40 |
+
|
| 41 |
+
|
| 42 |
+
@dataclass
|
| 43 |
+
class Result:
|
| 44 |
+
workload: str
|
| 45 |
+
shape: str
|
| 46 |
+
native_us: float
|
| 47 |
+
wrapper_us: float
|
| 48 |
+
wrapper_native: float
|
| 49 |
+
eager_cudnn_us: float
|
| 50 |
+
compile_cudnn_us: float
|
| 51 |
+
diagnostic_predequant_cudnn_us: float | None
|
| 52 |
+
diagnostic_predequant_compile_us: float | None
|
| 53 |
+
max_abs: float
|
| 54 |
+
mean_abs: float
|
| 55 |
+
p99_abs: float
|
| 56 |
+
cosine: float
|
| 57 |
+
accepted: bool
|
| 58 |
+
|
| 59 |
+
|
| 60 |
+
def bench(fn, warmup, iters):
|
| 61 |
+
for _ in range(warmup):
|
| 62 |
+
fn()
|
| 63 |
+
torch.cuda.synchronize()
|
| 64 |
+
start = torch.cuda.Event(enable_timing=True)
|
| 65 |
+
end = torch.cuda.Event(enable_timing=True)
|
| 66 |
+
start.record()
|
| 67 |
+
for _ in range(iters):
|
| 68 |
+
fn()
|
| 69 |
+
end.record()
|
| 70 |
+
torch.cuda.synchronize()
|
| 71 |
+
return start.elapsed_time(end) * 1000.0 / iters
|
| 72 |
+
|
| 73 |
+
|
| 74 |
+
def build_native():
|
| 75 |
+
from torch.utils.cpp_extension import load
|
| 76 |
+
|
| 77 |
+
os.environ.setdefault("TORCH_CUDA_ARCH_LIST", "12.0a")
|
| 78 |
+
return load(
|
| 79 |
+
name="world_model_conv_raw_native",
|
| 80 |
+
sources=[
|
| 81 |
+
str(PACKAGE / "benchmarks/native_binding.cpp"),
|
| 82 |
+
str(PACKAGE / "csrc/fp8_conv3d_sm120_v18.cu"),
|
| 83 |
+
str(PACKAGE / "csrc/fp8_causal_conv3d_sm120.cu"),
|
| 84 |
+
str(PACKAGE / "csrc/fp8_conv2d_3x3_sm120.cu"),
|
| 85 |
+
str(PACKAGE / "csrc/nvfp4_causal_conv3d_sm120.cu"),
|
| 86 |
+
str(PACKAGE / "csrc/nvfp4_causal_conv3d_residual_sm120.cu"),
|
| 87 |
+
str(PACKAGE / "csrc/nvfp4_causal_conv3d_residual_k128_sm120.cu"),
|
| 88 |
+
],
|
| 89 |
+
extra_include_paths=[str(PACKAGE / "csrc")],
|
| 90 |
+
extra_cflags=["-O3"],
|
| 91 |
+
extra_cuda_cflags=["-O3"],
|
| 92 |
+
verbose=False,
|
| 93 |
+
)
|
| 94 |
+
|
| 95 |
+
|
| 96 |
+
def metrics(got, ref):
|
| 97 |
+
diff = (got.float() - ref.float()).abs().flatten()
|
| 98 |
+
cosine = F.cosine_similarity(
|
| 99 |
+
got.float().flatten(), ref.float().flatten(), dim=0
|
| 100 |
+
).item()
|
| 101 |
+
return (
|
| 102 |
+
diff.max().item(),
|
| 103 |
+
diff.mean().item(),
|
| 104 |
+
torch.quantile(diff, 0.99).item(),
|
| 105 |
+
cosine,
|
| 106 |
+
)
|
| 107 |
+
|
| 108 |
+
|
| 109 |
+
def source_call(ops, name, *args, out):
|
| 110 |
+
if hasattr(ops, "_ops"):
|
| 111 |
+
getattr(ops._ops, name)(*args, out)
|
| 112 |
+
else:
|
| 113 |
+
getattr(ops, name)(*args, out=out)
|
| 114 |
+
|
| 115 |
+
|
| 116 |
+
def run_conv3d(ops, native, label, shape, args):
|
| 117 |
+
n, tc, tn, h, w, ci, co = shape
|
| 118 |
+
cache = (torch.randn((n, tc, h, w, ci), device="cuda") * 0.1).to(
|
| 119 |
+
torch.float8_e4m3fn
|
| 120 |
+
)
|
| 121 |
+
new = (torch.randn((n, tn, h, w, ci), device="cuda") * 0.1).to(
|
| 122 |
+
torch.float8_e4m3fn
|
| 123 |
+
)
|
| 124 |
+
weight = (torch.randn((co, 3, 3, 3, ci), device="cuda") * 0.1).to(
|
| 125 |
+
torch.float8_e4m3fn
|
| 126 |
+
)
|
| 127 |
+
bias = (torch.randn(co, device="cuda") * 0.01).to(torch.bfloat16)
|
| 128 |
+
out = torch.empty((n, tn, h, w, co), device="cuda", dtype=torch.bfloat16)
|
| 129 |
+
alpha = 0.75
|
| 130 |
+
|
| 131 |
+
wrapper = lambda: source_call(
|
| 132 |
+
ops,
|
| 133 |
+
"fp8_causal_conv3d_ndhwc_bf16",
|
| 134 |
+
cache,
|
| 135 |
+
new,
|
| 136 |
+
weight,
|
| 137 |
+
bias,
|
| 138 |
+
alpha,
|
| 139 |
+
out=out,
|
| 140 |
+
)
|
| 141 |
+
raw = lambda: native.causal_conv3d(
|
| 142 |
+
cache, new, weight, bias, alpha, out
|
| 143 |
+
)
|
| 144 |
+
|
| 145 |
+
def cudnn_ref():
|
| 146 |
+
x = torch.cat((cache, new), dim=1).float().permute(0, 4, 1, 2, 3)
|
| 147 |
+
wt = weight.float().permute(0, 4, 1, 2, 3)
|
| 148 |
+
y = F.conv3d(x, wt, padding=(0, 1, 1))
|
| 149 |
+
y = y.mul(alpha).add(bias.float().view(1, -1, 1, 1, 1))
|
| 150 |
+
out.copy_(y[:, :, :tn].permute(0, 2, 3, 4, 1).to(torch.bfloat16))
|
| 151 |
+
|
| 152 |
+
compiled = torch.compile(cudnn_ref, fullgraph=True)
|
| 153 |
+
wrapper()
|
| 154 |
+
got = out.clone()
|
| 155 |
+
cudnn_ref()
|
| 156 |
+
ref = out.clone()
|
| 157 |
+
max_abs, mean_abs, p99_abs, cosine = metrics(got, ref)
|
| 158 |
+
native_us = bench(raw, args.warmup, args.iters)
|
| 159 |
+
wrapper_us = bench(wrapper, args.warmup, args.iters)
|
| 160 |
+
eager_us = bench(cudnn_ref, args.warmup, args.iters)
|
| 161 |
+
compile_us = bench(compiled, args.warmup, args.iters)
|
| 162 |
+
return Result(
|
| 163 |
+
label,
|
| 164 |
+
str(shape),
|
| 165 |
+
native_us,
|
| 166 |
+
wrapper_us,
|
| 167 |
+
wrapper_us / native_us,
|
| 168 |
+
eager_us,
|
| 169 |
+
compile_us,
|
| 170 |
+
None,
|
| 171 |
+
None,
|
| 172 |
+
max_abs,
|
| 173 |
+
mean_abs,
|
| 174 |
+
p99_abs,
|
| 175 |
+
cosine,
|
| 176 |
+
wrapper_us - native_us <= max(0.5, native_us * 0.05)
|
| 177 |
+
and wrapper_us <= min(eager_us, compile_us) * 0.98
|
| 178 |
+
and cosine >= 0.999
|
| 179 |
+
and mean_abs <= 0.01,
|
| 180 |
+
)
|
| 181 |
+
|
| 182 |
+
|
| 183 |
+
def run_conv3d_residual(ops, native, label, shape, args):
|
| 184 |
+
n, tc, tn, h, w, ci, co = shape
|
| 185 |
+
cache = (torch.randn((n, tc, h, w, ci), device="cuda") * 0.1).to(
|
| 186 |
+
torch.float8_e4m3fn
|
| 187 |
+
)
|
| 188 |
+
new = (torch.randn((n, tn, h, w, ci), device="cuda") * 0.1).to(
|
| 189 |
+
torch.float8_e4m3fn
|
| 190 |
+
)
|
| 191 |
+
weight = (torch.randn((co, 3, 3, 3, ci), device="cuda") * 0.1).to(
|
| 192 |
+
torch.float8_e4m3fn
|
| 193 |
+
)
|
| 194 |
+
bias = (torch.randn(co, device="cuda") * 0.01).to(torch.bfloat16)
|
| 195 |
+
residual = torch.randn(
|
| 196 |
+
(n, co, tn, h, w), device="cuda", dtype=torch.bfloat16
|
| 197 |
+
)
|
| 198 |
+
out = torch.empty_like(residual)
|
| 199 |
+
alpha = 0.75
|
| 200 |
+
wrapper = lambda: source_call(
|
| 201 |
+
ops,
|
| 202 |
+
"fp8_conv3d_v18_ncdhw_res_bf16out",
|
| 203 |
+
cache,
|
| 204 |
+
new,
|
| 205 |
+
weight,
|
| 206 |
+
bias,
|
| 207 |
+
residual,
|
| 208 |
+
alpha,
|
| 209 |
+
out=out,
|
| 210 |
+
)
|
| 211 |
+
raw = lambda: native.causal_conv3d_residual(
|
| 212 |
+
cache, new, weight, bias, residual, alpha, out
|
| 213 |
+
)
|
| 214 |
+
|
| 215 |
+
def cudnn_ref():
|
| 216 |
+
x = torch.cat((cache, new), dim=1).float().permute(0, 4, 1, 2, 3)
|
| 217 |
+
wt = weight.float().permute(0, 4, 1, 2, 3)
|
| 218 |
+
y = F.conv3d(x, wt, padding=(0, 1, 1))
|
| 219 |
+
y = y.mul(alpha).add(bias.float().view(1, -1, 1, 1, 1))
|
| 220 |
+
y = (
|
| 221 |
+
y[:, :, :tn].to(torch.bfloat16).float()
|
| 222 |
+
+ residual.float()
|
| 223 |
+
).to(torch.bfloat16)
|
| 224 |
+
out.copy_(y)
|
| 225 |
+
|
| 226 |
+
compiled = torch.compile(cudnn_ref, fullgraph=True)
|
| 227 |
+
wrapper()
|
| 228 |
+
got = out.clone()
|
| 229 |
+
cudnn_ref()
|
| 230 |
+
ref = out.clone()
|
| 231 |
+
max_abs, mean_abs, p99_abs, cosine = metrics(got, ref)
|
| 232 |
+
native_us = bench(raw, args.warmup, args.iters)
|
| 233 |
+
wrapper_us = bench(wrapper, args.warmup, args.iters)
|
| 234 |
+
eager_us = bench(cudnn_ref, args.warmup, args.iters)
|
| 235 |
+
compile_us = bench(compiled, args.warmup, args.iters)
|
| 236 |
+
return Result(
|
| 237 |
+
f"{label}-residual",
|
| 238 |
+
str(shape),
|
| 239 |
+
native_us,
|
| 240 |
+
wrapper_us,
|
| 241 |
+
wrapper_us / native_us,
|
| 242 |
+
eager_us,
|
| 243 |
+
compile_us,
|
| 244 |
+
None,
|
| 245 |
+
None,
|
| 246 |
+
max_abs,
|
| 247 |
+
mean_abs,
|
| 248 |
+
p99_abs,
|
| 249 |
+
cosine,
|
| 250 |
+
wrapper_us - native_us <= max(0.5, native_us * 0.05)
|
| 251 |
+
and wrapper_us <= min(eager_us, compile_us) * 0.98
|
| 252 |
+
and cosine >= 0.999
|
| 253 |
+
and mean_abs <= 0.01,
|
| 254 |
+
)
|
| 255 |
+
|
| 256 |
+
|
| 257 |
+
def run_conv2d(ops, native, label, shape, args):
|
| 258 |
+
n, h, w, ci, co = shape
|
| 259 |
+
input = (torch.randn((n, h, w, ci), device="cuda") * 0.1).to(
|
| 260 |
+
torch.float8_e4m3fn
|
| 261 |
+
)
|
| 262 |
+
weight = (torch.randn((co, 3, 3, ci), device="cuda") * 0.1).to(
|
| 263 |
+
torch.float8_e4m3fn
|
| 264 |
+
)
|
| 265 |
+
bias = (torch.randn(co, device="cuda") * 0.01).to(torch.bfloat16)
|
| 266 |
+
out = torch.empty((n, h, w, co), device="cuda", dtype=torch.bfloat16)
|
| 267 |
+
alpha = 0.75
|
| 268 |
+
wrapper = lambda: source_call(
|
| 269 |
+
ops,
|
| 270 |
+
"fp8_conv2d_3x3_nhwc_bf16",
|
| 271 |
+
input,
|
| 272 |
+
weight,
|
| 273 |
+
bias,
|
| 274 |
+
alpha,
|
| 275 |
+
out=out,
|
| 276 |
+
)
|
| 277 |
+
raw = lambda: native.conv2d(input, weight, bias, alpha, out)
|
| 278 |
+
|
| 279 |
+
def cudnn_ref():
|
| 280 |
+
x = input.float().permute(0, 3, 1, 2)
|
| 281 |
+
wt = weight.float().permute(0, 3, 1, 2)
|
| 282 |
+
y = F.conv2d(x, wt, padding=1).mul(alpha)
|
| 283 |
+
y = y.add(bias.float().view(1, -1, 1, 1))
|
| 284 |
+
out.copy_(y.permute(0, 2, 3, 1).to(torch.bfloat16))
|
| 285 |
+
|
| 286 |
+
compiled = torch.compile(cudnn_ref, fullgraph=True)
|
| 287 |
+
wrapper()
|
| 288 |
+
got = out.clone()
|
| 289 |
+
cudnn_ref()
|
| 290 |
+
ref = out.clone()
|
| 291 |
+
max_abs, mean_abs, p99_abs, cosine = metrics(got, ref)
|
| 292 |
+
native_us = bench(raw, args.warmup, args.iters)
|
| 293 |
+
wrapper_us = bench(wrapper, args.warmup, args.iters)
|
| 294 |
+
eager_us = bench(cudnn_ref, args.warmup, args.iters)
|
| 295 |
+
compile_us = bench(compiled, args.warmup, args.iters)
|
| 296 |
+
return Result(
|
| 297 |
+
label,
|
| 298 |
+
str(shape),
|
| 299 |
+
native_us,
|
| 300 |
+
wrapper_us,
|
| 301 |
+
wrapper_us / native_us,
|
| 302 |
+
eager_us,
|
| 303 |
+
compile_us,
|
| 304 |
+
None,
|
| 305 |
+
None,
|
| 306 |
+
max_abs,
|
| 307 |
+
mean_abs,
|
| 308 |
+
p99_abs,
|
| 309 |
+
cosine,
|
| 310 |
+
wrapper_us - native_us <= max(0.5, native_us * 0.05)
|
| 311 |
+
and wrapper_us <= min(eager_us, compile_us) * 0.98
|
| 312 |
+
and cosine >= 0.999
|
| 313 |
+
and mean_abs <= 0.01,
|
| 314 |
+
)
|
| 315 |
+
|
| 316 |
+
|
| 317 |
+
def run_nvfp4_conv3d(ops, native, label, shape, args, residual_path):
|
| 318 |
+
n, tc, tn, h, w, ci, co = shape
|
| 319 |
+
cache_bf16 = (
|
| 320 |
+
torch.randn((n, tc, h, w, ci), device="cuda") * 0.1
|
| 321 |
+
).to(torch.bfloat16)
|
| 322 |
+
input_bf16 = (
|
| 323 |
+
torch.randn((n, tn, h, w, ci), device="cuda") * 0.1
|
| 324 |
+
).to(torch.bfloat16)
|
| 325 |
+
weight_bf16 = (
|
| 326 |
+
torch.randn((co, 3, 3, 3, ci), device="cuda") * 0.1
|
| 327 |
+
).to(torch.bfloat16)
|
| 328 |
+
cache, cache_sf = quantize_conv_tensor(cache_bf16)
|
| 329 |
+
input, input_sf = quantize_conv_tensor(input_bf16)
|
| 330 |
+
weight, weight_sf = quantize_conv_tensor(weight_bf16)
|
| 331 |
+
cache_dequant = dequantize_linear_nvfp4(
|
| 332 |
+
cache.reshape(-1, ci // 2), cache_sf.reshape(-1, ci // 16)
|
| 333 |
+
).reshape_as(cache_bf16)
|
| 334 |
+
input_dequant = dequantize_linear_nvfp4(
|
| 335 |
+
input.reshape(-1, ci // 2), input_sf.reshape(-1, ci // 16)
|
| 336 |
+
).reshape_as(input_bf16)
|
| 337 |
+
weight_dequant = dequantize_linear_nvfp4(
|
| 338 |
+
weight.reshape(-1, ci // 2), weight_sf.reshape(-1, ci // 16)
|
| 339 |
+
).reshape_as(weight_bf16)
|
| 340 |
+
bias = (torch.randn(co, device="cuda") * 0.01).to(torch.bfloat16)
|
| 341 |
+
alpha = 0.75
|
| 342 |
+
|
| 343 |
+
if residual_path:
|
| 344 |
+
residual = torch.randn(
|
| 345 |
+
(n, co, tn, h, w), device="cuda", dtype=torch.bfloat16
|
| 346 |
+
)
|
| 347 |
+
out = torch.empty_like(residual)
|
| 348 |
+
wrapper = lambda: source_call(
|
| 349 |
+
ops,
|
| 350 |
+
"nvfp4_causal_conv3d_residual_ncdhw_bf16",
|
| 351 |
+
cache, input, weight, cache_sf, input_sf, weight_sf, bias,
|
| 352 |
+
residual, None, alpha, out=out,
|
| 353 |
+
)
|
| 354 |
+
raw = lambda: native.nvfp4_causal_conv3d_residual(
|
| 355 |
+
cache, input, weight, cache_sf, input_sf, weight_sf, bias,
|
| 356 |
+
residual, alpha, out,
|
| 357 |
+
)
|
| 358 |
+
else:
|
| 359 |
+
residual = None
|
| 360 |
+
out = torch.empty(
|
| 361 |
+
(n, tn, h, w, co), device="cuda", dtype=torch.bfloat16
|
| 362 |
+
)
|
| 363 |
+
wrapper = lambda: source_call(
|
| 364 |
+
ops,
|
| 365 |
+
"nvfp4_causal_conv3d_ndhwc_bf16",
|
| 366 |
+
cache, input, weight, cache_sf, input_sf, weight_sf, bias,
|
| 367 |
+
None, alpha, out=out,
|
| 368 |
+
)
|
| 369 |
+
raw = lambda: native.nvfp4_causal_conv3d(
|
| 370 |
+
cache, input, weight, cache_sf, input_sf, weight_sf, bias,
|
| 371 |
+
alpha, out,
|
| 372 |
+
)
|
| 373 |
+
|
| 374 |
+
def store_cudnn_result(cache_value, input_value):
|
| 375 |
+
x = torch.cat((cache_value, input_value), dim=1).permute(
|
| 376 |
+
0, 4, 1, 2, 3
|
| 377 |
+
)
|
| 378 |
+
wt = weight_dequant.permute(0, 4, 1, 2, 3)
|
| 379 |
+
y = F.conv3d(x, wt, padding=(0, 1, 1))[:, :, :tn]
|
| 380 |
+
y = y.mul(alpha).add(bias.float().view(1, -1, 1, 1, 1))
|
| 381 |
+
if residual_path:
|
| 382 |
+
out.copy_(
|
| 383 |
+
(y.to(torch.bfloat16).float() + residual.float()).to(
|
| 384 |
+
torch.bfloat16
|
| 385 |
+
)
|
| 386 |
+
)
|
| 387 |
+
else:
|
| 388 |
+
out.copy_(y.permute(0, 2, 3, 4, 1).to(torch.bfloat16))
|
| 389 |
+
|
| 390 |
+
def predequant_cudnn_ref():
|
| 391 |
+
store_cudnn_result(cache_dequant, input_dequant)
|
| 392 |
+
|
| 393 |
+
magnitude = torch.tensor(
|
| 394 |
+
[0.0, 0.5, 1.0, 1.5, 2.0, 3.0, 4.0, 6.0], device="cuda"
|
| 395 |
+
)
|
| 396 |
+
scale_values, scale_bytes = [], []
|
| 397 |
+
for byte in list(range(0x78)) + [0xFE]:
|
| 398 |
+
exponent = (byte >> 3) & 0xF
|
| 399 |
+
mantissa = byte & 0x7
|
| 400 |
+
value = (
|
| 401 |
+
(mantissa / 8.0) * (2.0 ** -6)
|
| 402 |
+
if exponent == 0
|
| 403 |
+
else (1.0 + mantissa / 8.0) * (2.0 ** (exponent - 7))
|
| 404 |
+
)
|
| 405 |
+
scale_values.append(value)
|
| 406 |
+
scale_bytes.append(byte)
|
| 407 |
+
scale_lookup = torch.zeros(256, device="cuda")
|
| 408 |
+
scale_lookup[
|
| 409 |
+
torch.tensor(scale_bytes, device="cuda", dtype=torch.long)
|
| 410 |
+
] = torch.tensor(scale_values, device="cuda")
|
| 411 |
+
|
| 412 |
+
def unpack(packed_value, scale_value):
|
| 413 |
+
low = packed_value & 0xF
|
| 414 |
+
high = packed_value >> 4
|
| 415 |
+
low_value = magnitude[(low & 0x7).long()] * torch.where(
|
| 416 |
+
low & 0x8 != 0, -1.0, 1.0
|
| 417 |
+
)
|
| 418 |
+
high_value = magnitude[(high & 0x7).long()] * torch.where(
|
| 419 |
+
high & 0x8 != 0, -1.0, 1.0
|
| 420 |
+
)
|
| 421 |
+
values = torch.stack((low_value, high_value), dim=-1).flatten(-2)
|
| 422 |
+
scales_value = scale_lookup[scale_value.long()].repeat_interleave(
|
| 423 |
+
16, dim=-1
|
| 424 |
+
)
|
| 425 |
+
return values * scales_value
|
| 426 |
+
|
| 427 |
+
def cudnn_ref():
|
| 428 |
+
cache_value = unpack(cache, cache_sf).reshape_as(cache_bf16)
|
| 429 |
+
input_value = unpack(input, input_sf).reshape_as(input_bf16)
|
| 430 |
+
store_cudnn_result(cache_value, input_value)
|
| 431 |
+
|
| 432 |
+
compiled = torch.compile(cudnn_ref, fullgraph=True)
|
| 433 |
+
predequant_compiled = torch.compile(predequant_cudnn_ref, fullgraph=True)
|
| 434 |
+
wrapper()
|
| 435 |
+
got = out.clone()
|
| 436 |
+
cudnn_ref()
|
| 437 |
+
ref = out.clone()
|
| 438 |
+
max_abs, mean_abs, p99_abs, cosine = metrics(got, ref)
|
| 439 |
+
native_us = bench(raw, args.warmup, args.iters)
|
| 440 |
+
wrapper_us = bench(wrapper, args.warmup, args.iters)
|
| 441 |
+
eager_us = bench(cudnn_ref, args.warmup, args.iters)
|
| 442 |
+
compile_us = bench(compiled, args.warmup, args.iters)
|
| 443 |
+
diagnostic_eager_us = bench(
|
| 444 |
+
predequant_cudnn_ref, args.warmup, args.iters
|
| 445 |
+
)
|
| 446 |
+
diagnostic_compile_us = bench(
|
| 447 |
+
predequant_compiled, args.warmup, args.iters
|
| 448 |
+
)
|
| 449 |
+
return Result(
|
| 450 |
+
f"{label}{'-residual' if residual_path else ''}",
|
| 451 |
+
str(shape),
|
| 452 |
+
native_us,
|
| 453 |
+
wrapper_us,
|
| 454 |
+
wrapper_us / native_us,
|
| 455 |
+
eager_us,
|
| 456 |
+
compile_us,
|
| 457 |
+
diagnostic_eager_us,
|
| 458 |
+
diagnostic_compile_us,
|
| 459 |
+
max_abs,
|
| 460 |
+
mean_abs,
|
| 461 |
+
p99_abs,
|
| 462 |
+
cosine,
|
| 463 |
+
wrapper_us - native_us <= max(0.5, native_us * 0.05)
|
| 464 |
+
and wrapper_us <= min(eager_us, compile_us) * 0.98
|
| 465 |
+
and cosine >= 0.998
|
| 466 |
+
and mean_abs <= 0.02,
|
| 467 |
+
)
|
| 468 |
+
|
| 469 |
+
|
| 470 |
+
def main():
|
| 471 |
+
parser = argparse.ArgumentParser()
|
| 472 |
+
parser.add_argument("--backend", choices=["source", "installed"], default="source")
|
| 473 |
+
parser.add_argument("--artifact")
|
| 474 |
+
parser.add_argument("--warmup", type=int, default=10)
|
| 475 |
+
parser.add_argument("--iters", type=int, default=30)
|
| 476 |
+
parser.add_argument("--output")
|
| 477 |
+
args = parser.parse_args()
|
| 478 |
+
ops = (
|
| 479 |
+
load_source_ops()
|
| 480 |
+
if args.backend == "source"
|
| 481 |
+
else load_installed_ops(args.artifact)
|
| 482 |
+
)
|
| 483 |
+
native = build_native()
|
| 484 |
+
rows = [
|
| 485 |
+
*(run_conv3d(ops, native, name, shape, args)
|
| 486 |
+
for name, shape in CONV3D_SHAPES.items()),
|
| 487 |
+
*(run_conv3d_residual(ops, native, name, shape, args)
|
| 488 |
+
for name, shape in CONV3D_SHAPES.items()
|
| 489 |
+
if shape[-1] % 8 == 0),
|
| 490 |
+
*(run_conv2d(ops, native, name, shape, args)
|
| 491 |
+
for name, shape in CONV2D_SHAPES.items()),
|
| 492 |
+
*(run_nvfp4_conv3d(ops, native, name, shape, args, False)
|
| 493 |
+
for name, shape in NVFP4_CONV3D_SHAPES.items()),
|
| 494 |
+
*(run_nvfp4_conv3d(ops, native, name, shape, args, True)
|
| 495 |
+
for name, shape in NVFP4_CONV3D_SHAPES.items()),
|
| 496 |
+
]
|
| 497 |
+
for row in rows:
|
| 498 |
+
print(
|
| 499 |
+
f"{row.workload}: native={row.native_us:.3f}us "
|
| 500 |
+
f"wrapper={row.wrapper_us:.3f}us ({row.wrapper_native:.3f}) "
|
| 501 |
+
f"cuDNN-eager={row.eager_cudnn_us:.3f}us "
|
| 502 |
+
f"cuDNN-compile={row.compile_cudnn_us:.3f}us "
|
| 503 |
+
+ (
|
| 504 |
+
f"predequant-compile="
|
| 505 |
+
f"{row.diagnostic_predequant_compile_us:.3f}us "
|
| 506 |
+
if row.diagnostic_predequant_compile_us is not None
|
| 507 |
+
else ""
|
| 508 |
+
)
|
| 509 |
+
+ f"cos={row.cosine:.7f} accepted={row.accepted}"
|
| 510 |
+
)
|
| 511 |
+
if args.output:
|
| 512 |
+
path = Path(args.output)
|
| 513 |
+
path.parent.mkdir(parents=True, exist_ok=True)
|
| 514 |
+
path.write_text(json.dumps([asdict(row) for row in rows], indent=2) + "\n")
|
| 515 |
+
if not all(row.accepted for row in rows):
|
| 516 |
+
raise SystemExit("world-model Conv acceptance failed")
|
| 517 |
+
|
| 518 |
+
|
| 519 |
+
if __name__ == "__main__":
|
| 520 |
+
main()
|
build/torch211-cxx11-cu128-x86_64-linux/__init__.py
CHANGED
|
@@ -27,6 +27,11 @@ def _fp8_conv3d_fake(
|
|
| 27 |
raise RuntimeError("cache_x must have shape (N,2,H,W,Ci)")
|
| 28 |
if weight.shape != (co, 3, 3, 3, ci):
|
| 29 |
raise RuntimeError("weight must have shape (Co,3,3,3,Ci)")
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 30 |
if residual.shape != (n, co, t_new, h, w) or out.shape != residual.shape:
|
| 31 |
raise RuntimeError("residual/out must be NCDHW")
|
| 32 |
if bias.shape != (co,):
|
|
@@ -34,6 +39,179 @@ def _fp8_conv3d_fake(
|
|
| 34 |
return None
|
| 35 |
|
| 36 |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 37 |
def fp8_conv3d_v18_ncdhw_res_bf16out(
|
| 38 |
cache_x: torch.Tensor,
|
| 39 |
new_x: torch.Tensor,
|
|
@@ -54,4 +232,156 @@ def fp8_conv3d_v18_ncdhw_res_bf16out(
|
|
| 54 |
return out
|
| 55 |
|
| 56 |
|
| 57 |
-
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 27 |
raise RuntimeError("cache_x must have shape (N,2,H,W,Ci)")
|
| 28 |
if weight.shape != (co, 3, 3, 3, ci):
|
| 29 |
raise RuntimeError("weight must have shape (Co,3,3,3,Ci)")
|
| 30 |
+
if ci not in (32, 64):
|
| 31 |
+
raise RuntimeError(
|
| 32 |
+
"FP8 Conv3D is accepted only for Ci=32/64; use the "
|
| 33 |
+
"strong-library or NVFP4 path for larger channels"
|
| 34 |
+
)
|
| 35 |
if residual.shape != (n, co, t_new, h, w) or out.shape != residual.shape:
|
| 36 |
raise RuntimeError("residual/out must be NCDHW")
|
| 37 |
if bias.shape != (co,):
|
|
|
|
| 39 |
return None
|
| 40 |
|
| 41 |
|
| 42 |
+
@torch.library.register_fake(
|
| 43 |
+
add_op_namespace_prefix("fp8_causal_conv3d_ndhwc_bf16")
|
| 44 |
+
)
|
| 45 |
+
def _fp8_causal_conv3d_fake(
|
| 46 |
+
cache_x: torch.Tensor,
|
| 47 |
+
new_x: torch.Tensor,
|
| 48 |
+
weight: torch.Tensor,
|
| 49 |
+
bias: torch.Tensor,
|
| 50 |
+
alpha: float,
|
| 51 |
+
out: torch.Tensor,
|
| 52 |
+
) -> None:
|
| 53 |
+
if cache_x.dim() != 5 or new_x.dim() != 5:
|
| 54 |
+
raise RuntimeError("cache_x/new_x must be NDHWC")
|
| 55 |
+
n, t, h, w, ci = new_x.shape
|
| 56 |
+
co = weight.shape[0]
|
| 57 |
+
if cache_x.shape != (n, 2, h, w, ci):
|
| 58 |
+
raise RuntimeError("cache_x must have shape (N,2,H,W,Ci)")
|
| 59 |
+
if weight.shape != (co, 3, 3, 3, ci):
|
| 60 |
+
raise RuntimeError("weight must have shape (Co,3,3,3,Ci)")
|
| 61 |
+
if ci not in (32, 64):
|
| 62 |
+
raise RuntimeError(
|
| 63 |
+
"FP8 Conv3D is accepted only for Ci=32/64; use the "
|
| 64 |
+
"strong-library or NVFP4 path for larger channels"
|
| 65 |
+
)
|
| 66 |
+
if bias.shape != (co,) or out.shape != (n, t, h, w, co):
|
| 67 |
+
raise RuntimeError("bias/out shape mismatch")
|
| 68 |
+
return None
|
| 69 |
+
|
| 70 |
+
|
| 71 |
+
@torch.library.register_fake(
|
| 72 |
+
add_op_namespace_prefix("fp8_conv2d_3x3_nhwc_bf16")
|
| 73 |
+
)
|
| 74 |
+
def _fp8_conv2d_fake(
|
| 75 |
+
input: torch.Tensor,
|
| 76 |
+
weight: torch.Tensor,
|
| 77 |
+
bias: torch.Tensor,
|
| 78 |
+
alpha: float,
|
| 79 |
+
out: torch.Tensor,
|
| 80 |
+
) -> None:
|
| 81 |
+
if input.dim() != 4:
|
| 82 |
+
raise RuntimeError("input must have shape (N,H,W,Ci)")
|
| 83 |
+
n, h, w, ci = input.shape
|
| 84 |
+
co = weight.shape[0]
|
| 85 |
+
if weight.shape != (co, 3, 3, ci):
|
| 86 |
+
raise RuntimeError("weight must have shape (Co,3,3,Ci)")
|
| 87 |
+
if bias.shape != (co,) or out.shape != (n, h, w, co):
|
| 88 |
+
raise RuntimeError("bias/out shape mismatch")
|
| 89 |
+
return None
|
| 90 |
+
|
| 91 |
+
|
| 92 |
+
@torch.library.register_fake(
|
| 93 |
+
add_op_namespace_prefix("fp8_conv2d_3x3_ncdhw_bf16")
|
| 94 |
+
)
|
| 95 |
+
def _fp8_conv2d_ncdhw_fake(
|
| 96 |
+
input: torch.Tensor,
|
| 97 |
+
weight: torch.Tensor,
|
| 98 |
+
bias: torch.Tensor,
|
| 99 |
+
alpha: float,
|
| 100 |
+
out: torch.Tensor,
|
| 101 |
+
) -> None:
|
| 102 |
+
if input.dim() != 5:
|
| 103 |
+
raise RuntimeError("input must have shape (B,T,H,W,Ci)")
|
| 104 |
+
b, t, h, w, ci = input.shape
|
| 105 |
+
co = weight.shape[0]
|
| 106 |
+
if weight.shape != (co, 3, 3, ci):
|
| 107 |
+
raise RuntimeError("weight must have shape (Co,3,3,Ci)")
|
| 108 |
+
if bias.shape != (co,) or out.shape != (b, co, t, h, w):
|
| 109 |
+
raise RuntimeError("bias/out shape mismatch")
|
| 110 |
+
return None
|
| 111 |
+
|
| 112 |
+
|
| 113 |
+
def _check_nvfp4_conv_shapes(
|
| 114 |
+
cache_packed: torch.Tensor,
|
| 115 |
+
new_packed: torch.Tensor,
|
| 116 |
+
weight_packed: torch.Tensor,
|
| 117 |
+
cache_sf: torch.Tensor,
|
| 118 |
+
new_sf: torch.Tensor,
|
| 119 |
+
weight_sf: torch.Tensor,
|
| 120 |
+
bias: torch.Tensor,
|
| 121 |
+
outer_weight: torch.Tensor | None,
|
| 122 |
+
) -> tuple[int, int, int, int, int]:
|
| 123 |
+
if (
|
| 124 |
+
cache_packed.dim() != 5
|
| 125 |
+
or new_packed.dim() != 5
|
| 126 |
+
or weight_packed.dim() != 5
|
| 127 |
+
):
|
| 128 |
+
raise RuntimeError("packed inputs must be five-dimensional")
|
| 129 |
+
n, t, h, w, ci_half = new_packed.shape
|
| 130 |
+
ci = ci_half * 2
|
| 131 |
+
co = weight_packed.shape[0]
|
| 132 |
+
if (
|
| 133 |
+
(ci != 64 and ci % 128 != 0)
|
| 134 |
+
or co % 8 != 0
|
| 135 |
+
or cache_packed.shape != (n, 2, h, w, ci // 2)
|
| 136 |
+
or weight_packed.shape != (co, 3, 3, 3, ci // 2)
|
| 137 |
+
or cache_sf.shape != (n, 2, h, w, ci // 16)
|
| 138 |
+
or new_sf.shape != (n, t, h, w, ci // 16)
|
| 139 |
+
or weight_sf.shape != (co, 3, 3, 3, ci // 16)
|
| 140 |
+
or bias.shape != (co,)
|
| 141 |
+
or (outer_weight is not None and outer_weight.shape != (co,))
|
| 142 |
+
):
|
| 143 |
+
raise RuntimeError(
|
| 144 |
+
"NVFP4 Conv3D accepts Ci=64 or multiples of 128; "
|
| 145 |
+
"other shapes must use the strong-library path"
|
| 146 |
+
)
|
| 147 |
+
return n, t, h, w, co
|
| 148 |
+
|
| 149 |
+
|
| 150 |
+
@torch.library.register_fake(
|
| 151 |
+
add_op_namespace_prefix("nvfp4_causal_conv3d_ndhwc_bf16")
|
| 152 |
+
)
|
| 153 |
+
def _nvfp4_causal_conv3d_fake(
|
| 154 |
+
cache_packed: torch.Tensor,
|
| 155 |
+
new_packed: torch.Tensor,
|
| 156 |
+
weight_packed: torch.Tensor,
|
| 157 |
+
cache_sf: torch.Tensor,
|
| 158 |
+
new_sf: torch.Tensor,
|
| 159 |
+
weight_sf: torch.Tensor,
|
| 160 |
+
bias: torch.Tensor,
|
| 161 |
+
outer_weight: torch.Tensor | None,
|
| 162 |
+
alpha: float,
|
| 163 |
+
out: torch.Tensor,
|
| 164 |
+
) -> None:
|
| 165 |
+
del alpha
|
| 166 |
+
n, t, h, w, co = _check_nvfp4_conv_shapes(
|
| 167 |
+
cache_packed,
|
| 168 |
+
new_packed,
|
| 169 |
+
weight_packed,
|
| 170 |
+
cache_sf,
|
| 171 |
+
new_sf,
|
| 172 |
+
weight_sf,
|
| 173 |
+
bias,
|
| 174 |
+
outer_weight,
|
| 175 |
+
)
|
| 176 |
+
if out.shape != (n, t, h, w, co):
|
| 177 |
+
raise RuntimeError("out must have shape (N,T,H,W,Co)")
|
| 178 |
+
return None
|
| 179 |
+
|
| 180 |
+
|
| 181 |
+
@torch.library.register_fake(
|
| 182 |
+
add_op_namespace_prefix(
|
| 183 |
+
"nvfp4_causal_conv3d_residual_ncdhw_bf16"
|
| 184 |
+
)
|
| 185 |
+
)
|
| 186 |
+
def _nvfp4_causal_conv3d_residual_fake(
|
| 187 |
+
cache_packed: torch.Tensor,
|
| 188 |
+
new_packed: torch.Tensor,
|
| 189 |
+
weight_packed: torch.Tensor,
|
| 190 |
+
cache_sf: torch.Tensor,
|
| 191 |
+
new_sf: torch.Tensor,
|
| 192 |
+
weight_sf: torch.Tensor,
|
| 193 |
+
bias: torch.Tensor,
|
| 194 |
+
residual: torch.Tensor,
|
| 195 |
+
outer_weight: torch.Tensor | None,
|
| 196 |
+
alpha: float,
|
| 197 |
+
out: torch.Tensor,
|
| 198 |
+
) -> None:
|
| 199 |
+
del alpha
|
| 200 |
+
n, t, h, w, co = _check_nvfp4_conv_shapes(
|
| 201 |
+
cache_packed,
|
| 202 |
+
new_packed,
|
| 203 |
+
weight_packed,
|
| 204 |
+
cache_sf,
|
| 205 |
+
new_sf,
|
| 206 |
+
weight_sf,
|
| 207 |
+
bias,
|
| 208 |
+
outer_weight,
|
| 209 |
+
)
|
| 210 |
+
if residual.shape != (n, co, t, h, w) or out.shape != residual.shape:
|
| 211 |
+
raise RuntimeError("residual/out must have shape (N,Co,T,H,W)")
|
| 212 |
+
return None
|
| 213 |
+
|
| 214 |
+
|
| 215 |
def fp8_conv3d_v18_ncdhw_res_bf16out(
|
| 216 |
cache_x: torch.Tensor,
|
| 217 |
new_x: torch.Tensor,
|
|
|
|
| 232 |
return out
|
| 233 |
|
| 234 |
|
| 235 |
+
def fp8_causal_conv3d_ndhwc_bf16(
|
| 236 |
+
cache_x: torch.Tensor,
|
| 237 |
+
new_x: torch.Tensor,
|
| 238 |
+
weight: torch.Tensor,
|
| 239 |
+
bias: torch.Tensor,
|
| 240 |
+
alpha: float = 1.0,
|
| 241 |
+
*,
|
| 242 |
+
out: Optional[torch.Tensor] = None,
|
| 243 |
+
) -> torch.Tensor:
|
| 244 |
+
"""FP8 causal 3D convolution with virtual two-frame cache concat."""
|
| 245 |
+
|
| 246 |
+
n, t, h, w, _ = new_x.shape
|
| 247 |
+
if out is None:
|
| 248 |
+
out = torch.empty(
|
| 249 |
+
(n, t, h, w, weight.shape[0]),
|
| 250 |
+
device=new_x.device,
|
| 251 |
+
dtype=torch.bfloat16,
|
| 252 |
+
)
|
| 253 |
+
ops.fp8_causal_conv3d_ndhwc_bf16(
|
| 254 |
+
cache_x, new_x, weight, bias, float(alpha), out
|
| 255 |
+
)
|
| 256 |
+
return out
|
| 257 |
+
|
| 258 |
+
|
| 259 |
+
def fp8_conv2d_3x3_nhwc_bf16(
|
| 260 |
+
input: torch.Tensor,
|
| 261 |
+
weight: torch.Tensor,
|
| 262 |
+
bias: torch.Tensor,
|
| 263 |
+
alpha: float = 1.0,
|
| 264 |
+
*,
|
| 265 |
+
out: Optional[torch.Tensor] = None,
|
| 266 |
+
) -> torch.Tensor:
|
| 267 |
+
"""FP8 3x3 Conv2D with NHWC input/output and BF16 epilogue."""
|
| 268 |
+
|
| 269 |
+
if out is None:
|
| 270 |
+
out = torch.empty(
|
| 271 |
+
(*input.shape[:3], weight.shape[0]),
|
| 272 |
+
device=input.device,
|
| 273 |
+
dtype=torch.bfloat16,
|
| 274 |
+
)
|
| 275 |
+
ops.fp8_conv2d_3x3_nhwc_bf16(
|
| 276 |
+
input, weight, bias, float(alpha), out
|
| 277 |
+
)
|
| 278 |
+
return out
|
| 279 |
+
|
| 280 |
+
|
| 281 |
+
def fp8_conv2d_3x3_ncdhw_bf16(
|
| 282 |
+
input: torch.Tensor,
|
| 283 |
+
weight: torch.Tensor,
|
| 284 |
+
bias: torch.Tensor,
|
| 285 |
+
alpha: float = 1.0,
|
| 286 |
+
*,
|
| 287 |
+
out: Optional[torch.Tensor] = None,
|
| 288 |
+
) -> torch.Tensor:
|
| 289 |
+
"""FP8 3x3 Conv2D over B*T frames with direct BF16 NCDHW output."""
|
| 290 |
+
|
| 291 |
+
if out is None:
|
| 292 |
+
out = torch.empty(
|
| 293 |
+
(
|
| 294 |
+
input.shape[0],
|
| 295 |
+
weight.shape[0],
|
| 296 |
+
input.shape[1],
|
| 297 |
+
input.shape[2],
|
| 298 |
+
input.shape[3],
|
| 299 |
+
),
|
| 300 |
+
device=input.device,
|
| 301 |
+
dtype=torch.bfloat16,
|
| 302 |
+
)
|
| 303 |
+
ops.fp8_conv2d_3x3_ncdhw_bf16(
|
| 304 |
+
input, weight, bias, float(alpha), out
|
| 305 |
+
)
|
| 306 |
+
return out
|
| 307 |
+
|
| 308 |
+
|
| 309 |
+
def nvfp4_causal_conv3d_ndhwc_bf16(
|
| 310 |
+
cache_packed: torch.Tensor,
|
| 311 |
+
new_packed: torch.Tensor,
|
| 312 |
+
weight_packed: torch.Tensor,
|
| 313 |
+
cache_sf: torch.Tensor,
|
| 314 |
+
new_sf: torch.Tensor,
|
| 315 |
+
weight_sf: torch.Tensor,
|
| 316 |
+
bias: torch.Tensor,
|
| 317 |
+
outer_weight: torch.Tensor | None = None,
|
| 318 |
+
alpha: float = 1.0,
|
| 319 |
+
*,
|
| 320 |
+
out: Optional[torch.Tensor] = None,
|
| 321 |
+
) -> torch.Tensor:
|
| 322 |
+
"""NVFP4 causal Conv3D with linear UE4M3 scale factors and BF16 NDHWC output."""
|
| 323 |
+
|
| 324 |
+
n, t, h, w, _ = new_packed.shape
|
| 325 |
+
if out is None:
|
| 326 |
+
out = torch.empty(
|
| 327 |
+
(n, t, h, w, weight_packed.shape[0]),
|
| 328 |
+
device=new_packed.device,
|
| 329 |
+
dtype=torch.bfloat16,
|
| 330 |
+
)
|
| 331 |
+
ops.nvfp4_causal_conv3d_ndhwc_bf16(
|
| 332 |
+
cache_packed,
|
| 333 |
+
new_packed,
|
| 334 |
+
weight_packed,
|
| 335 |
+
cache_sf,
|
| 336 |
+
new_sf,
|
| 337 |
+
weight_sf,
|
| 338 |
+
bias,
|
| 339 |
+
outer_weight,
|
| 340 |
+
float(alpha),
|
| 341 |
+
out,
|
| 342 |
+
)
|
| 343 |
+
return out
|
| 344 |
+
|
| 345 |
+
|
| 346 |
+
def nvfp4_causal_conv3d_residual_ncdhw_bf16(
|
| 347 |
+
cache_packed: torch.Tensor,
|
| 348 |
+
new_packed: torch.Tensor,
|
| 349 |
+
weight_packed: torch.Tensor,
|
| 350 |
+
cache_sf: torch.Tensor,
|
| 351 |
+
new_sf: torch.Tensor,
|
| 352 |
+
weight_sf: torch.Tensor,
|
| 353 |
+
bias: torch.Tensor,
|
| 354 |
+
residual: torch.Tensor,
|
| 355 |
+
outer_weight: torch.Tensor | None = None,
|
| 356 |
+
alpha: float = 1.0,
|
| 357 |
+
*,
|
| 358 |
+
out: Optional[torch.Tensor] = None,
|
| 359 |
+
) -> torch.Tensor:
|
| 360 |
+
"""NVFP4 causal Conv3D with fused bias, residual, and BF16 NCDHW output."""
|
| 361 |
+
|
| 362 |
+
if out is None:
|
| 363 |
+
out = torch.empty_like(residual)
|
| 364 |
+
ops.nvfp4_causal_conv3d_residual_ncdhw_bf16(
|
| 365 |
+
cache_packed,
|
| 366 |
+
new_packed,
|
| 367 |
+
weight_packed,
|
| 368 |
+
cache_sf,
|
| 369 |
+
new_sf,
|
| 370 |
+
weight_sf,
|
| 371 |
+
bias,
|
| 372 |
+
residual,
|
| 373 |
+
outer_weight,
|
| 374 |
+
float(alpha),
|
| 375 |
+
out,
|
| 376 |
+
)
|
| 377 |
+
return out
|
| 378 |
+
|
| 379 |
+
|
| 380 |
+
__all__ = [
|
| 381 |
+
"fp8_conv3d_v18_ncdhw_res_bf16out",
|
| 382 |
+
"fp8_causal_conv3d_ndhwc_bf16",
|
| 383 |
+
"fp8_conv2d_3x3_nhwc_bf16",
|
| 384 |
+
"fp8_conv2d_3x3_ncdhw_bf16",
|
| 385 |
+
"nvfp4_causal_conv3d_ndhwc_bf16",
|
| 386 |
+
"nvfp4_causal_conv3d_residual_ncdhw_bf16",
|
| 387 |
+
]
|
build/torch211-cxx11-cu128-x86_64-linux/_ops.py
CHANGED
|
@@ -1,9 +1,9 @@
|
|
| 1 |
import torch
|
| 2 |
-
from . import
|
| 3 |
-
ops = torch.ops.
|
| 4 |
|
| 5 |
def add_op_namespace_prefix(op_name: str):
|
| 6 |
"""
|
| 7 |
Prefix op by namespace.
|
| 8 |
"""
|
| 9 |
-
return f"
|
|
|
|
| 1 |
import torch
|
| 2 |
+
from . import _world_model_conv_cuda_7781728
|
| 3 |
+
ops = torch.ops._world_model_conv_cuda_7781728
|
| 4 |
|
| 5 |
def add_op_namespace_prefix(op_name: str):
|
| 6 |
"""
|
| 7 |
Prefix op by namespace.
|
| 8 |
"""
|
| 9 |
+
return f"_world_model_conv_cuda_7781728::{op_name}"
|
build/torch211-cxx11-cu128-x86_64-linux/{_world_model_conv_cuda_f14c443.abi3.so → _world_model_conv_cuda_7781728.abi3.so}
RENAMED
|
@@ -1,3 +1,3 @@
|
|
| 1 |
version https://git-lfs.github.com/spec/v1
|
| 2 |
-
oid sha256:
|
| 3 |
-
size
|
|
|
|
| 1 |
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:d963296bec8c011fc81e22ef3d66ae0df353d36550f0844bbdf5541a4f23cd2a
|
| 3 |
+
size 1271616
|
build/torch211-cxx11-cu128-x86_64-linux/metadata.json
CHANGED
|
@@ -1,6 +1,6 @@
|
|
| 1 |
{
|
| 2 |
"name": "world-model-conv",
|
| 3 |
-
"id": "
|
| 4 |
"version": 1,
|
| 5 |
"license": "Apache-2.0",
|
| 6 |
"python-depends": [],
|
|
|
|
| 1 |
{
|
| 2 |
"name": "world-model-conv",
|
| 3 |
+
"id": "_world_model_conv_cuda_7781728",
|
| 4 |
"version": 1,
|
| 5 |
"license": "Apache-2.0",
|
| 6 |
"python-depends": [],
|
build/torch211-cxx11-cu130-x86_64-linux/__init__.py
CHANGED
|
@@ -27,6 +27,11 @@ def _fp8_conv3d_fake(
|
|
| 27 |
raise RuntimeError("cache_x must have shape (N,2,H,W,Ci)")
|
| 28 |
if weight.shape != (co, 3, 3, 3, ci):
|
| 29 |
raise RuntimeError("weight must have shape (Co,3,3,3,Ci)")
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 30 |
if residual.shape != (n, co, t_new, h, w) or out.shape != residual.shape:
|
| 31 |
raise RuntimeError("residual/out must be NCDHW")
|
| 32 |
if bias.shape != (co,):
|
|
@@ -34,6 +39,179 @@ def _fp8_conv3d_fake(
|
|
| 34 |
return None
|
| 35 |
|
| 36 |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 37 |
def fp8_conv3d_v18_ncdhw_res_bf16out(
|
| 38 |
cache_x: torch.Tensor,
|
| 39 |
new_x: torch.Tensor,
|
|
@@ -54,4 +232,156 @@ def fp8_conv3d_v18_ncdhw_res_bf16out(
|
|
| 54 |
return out
|
| 55 |
|
| 56 |
|
| 57 |
-
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 27 |
raise RuntimeError("cache_x must have shape (N,2,H,W,Ci)")
|
| 28 |
if weight.shape != (co, 3, 3, 3, ci):
|
| 29 |
raise RuntimeError("weight must have shape (Co,3,3,3,Ci)")
|
| 30 |
+
if ci not in (32, 64):
|
| 31 |
+
raise RuntimeError(
|
| 32 |
+
"FP8 Conv3D is accepted only for Ci=32/64; use the "
|
| 33 |
+
"strong-library or NVFP4 path for larger channels"
|
| 34 |
+
)
|
| 35 |
if residual.shape != (n, co, t_new, h, w) or out.shape != residual.shape:
|
| 36 |
raise RuntimeError("residual/out must be NCDHW")
|
| 37 |
if bias.shape != (co,):
|
|
|
|
| 39 |
return None
|
| 40 |
|
| 41 |
|
| 42 |
+
@torch.library.register_fake(
|
| 43 |
+
add_op_namespace_prefix("fp8_causal_conv3d_ndhwc_bf16")
|
| 44 |
+
)
|
| 45 |
+
def _fp8_causal_conv3d_fake(
|
| 46 |
+
cache_x: torch.Tensor,
|
| 47 |
+
new_x: torch.Tensor,
|
| 48 |
+
weight: torch.Tensor,
|
| 49 |
+
bias: torch.Tensor,
|
| 50 |
+
alpha: float,
|
| 51 |
+
out: torch.Tensor,
|
| 52 |
+
) -> None:
|
| 53 |
+
if cache_x.dim() != 5 or new_x.dim() != 5:
|
| 54 |
+
raise RuntimeError("cache_x/new_x must be NDHWC")
|
| 55 |
+
n, t, h, w, ci = new_x.shape
|
| 56 |
+
co = weight.shape[0]
|
| 57 |
+
if cache_x.shape != (n, 2, h, w, ci):
|
| 58 |
+
raise RuntimeError("cache_x must have shape (N,2,H,W,Ci)")
|
| 59 |
+
if weight.shape != (co, 3, 3, 3, ci):
|
| 60 |
+
raise RuntimeError("weight must have shape (Co,3,3,3,Ci)")
|
| 61 |
+
if ci not in (32, 64):
|
| 62 |
+
raise RuntimeError(
|
| 63 |
+
"FP8 Conv3D is accepted only for Ci=32/64; use the "
|
| 64 |
+
"strong-library or NVFP4 path for larger channels"
|
| 65 |
+
)
|
| 66 |
+
if bias.shape != (co,) or out.shape != (n, t, h, w, co):
|
| 67 |
+
raise RuntimeError("bias/out shape mismatch")
|
| 68 |
+
return None
|
| 69 |
+
|
| 70 |
+
|
| 71 |
+
@torch.library.register_fake(
|
| 72 |
+
add_op_namespace_prefix("fp8_conv2d_3x3_nhwc_bf16")
|
| 73 |
+
)
|
| 74 |
+
def _fp8_conv2d_fake(
|
| 75 |
+
input: torch.Tensor,
|
| 76 |
+
weight: torch.Tensor,
|
| 77 |
+
bias: torch.Tensor,
|
| 78 |
+
alpha: float,
|
| 79 |
+
out: torch.Tensor,
|
| 80 |
+
) -> None:
|
| 81 |
+
if input.dim() != 4:
|
| 82 |
+
raise RuntimeError("input must have shape (N,H,W,Ci)")
|
| 83 |
+
n, h, w, ci = input.shape
|
| 84 |
+
co = weight.shape[0]
|
| 85 |
+
if weight.shape != (co, 3, 3, ci):
|
| 86 |
+
raise RuntimeError("weight must have shape (Co,3,3,Ci)")
|
| 87 |
+
if bias.shape != (co,) or out.shape != (n, h, w, co):
|
| 88 |
+
raise RuntimeError("bias/out shape mismatch")
|
| 89 |
+
return None
|
| 90 |
+
|
| 91 |
+
|
| 92 |
+
@torch.library.register_fake(
|
| 93 |
+
add_op_namespace_prefix("fp8_conv2d_3x3_ncdhw_bf16")
|
| 94 |
+
)
|
| 95 |
+
def _fp8_conv2d_ncdhw_fake(
|
| 96 |
+
input: torch.Tensor,
|
| 97 |
+
weight: torch.Tensor,
|
| 98 |
+
bias: torch.Tensor,
|
| 99 |
+
alpha: float,
|
| 100 |
+
out: torch.Tensor,
|
| 101 |
+
) -> None:
|
| 102 |
+
if input.dim() != 5:
|
| 103 |
+
raise RuntimeError("input must have shape (B,T,H,W,Ci)")
|
| 104 |
+
b, t, h, w, ci = input.shape
|
| 105 |
+
co = weight.shape[0]
|
| 106 |
+
if weight.shape != (co, 3, 3, ci):
|
| 107 |
+
raise RuntimeError("weight must have shape (Co,3,3,Ci)")
|
| 108 |
+
if bias.shape != (co,) or out.shape != (b, co, t, h, w):
|
| 109 |
+
raise RuntimeError("bias/out shape mismatch")
|
| 110 |
+
return None
|
| 111 |
+
|
| 112 |
+
|
| 113 |
+
def _check_nvfp4_conv_shapes(
|
| 114 |
+
cache_packed: torch.Tensor,
|
| 115 |
+
new_packed: torch.Tensor,
|
| 116 |
+
weight_packed: torch.Tensor,
|
| 117 |
+
cache_sf: torch.Tensor,
|
| 118 |
+
new_sf: torch.Tensor,
|
| 119 |
+
weight_sf: torch.Tensor,
|
| 120 |
+
bias: torch.Tensor,
|
| 121 |
+
outer_weight: torch.Tensor | None,
|
| 122 |
+
) -> tuple[int, int, int, int, int]:
|
| 123 |
+
if (
|
| 124 |
+
cache_packed.dim() != 5
|
| 125 |
+
or new_packed.dim() != 5
|
| 126 |
+
or weight_packed.dim() != 5
|
| 127 |
+
):
|
| 128 |
+
raise RuntimeError("packed inputs must be five-dimensional")
|
| 129 |
+
n, t, h, w, ci_half = new_packed.shape
|
| 130 |
+
ci = ci_half * 2
|
| 131 |
+
co = weight_packed.shape[0]
|
| 132 |
+
if (
|
| 133 |
+
(ci != 64 and ci % 128 != 0)
|
| 134 |
+
or co % 8 != 0
|
| 135 |
+
or cache_packed.shape != (n, 2, h, w, ci // 2)
|
| 136 |
+
or weight_packed.shape != (co, 3, 3, 3, ci // 2)
|
| 137 |
+
or cache_sf.shape != (n, 2, h, w, ci // 16)
|
| 138 |
+
or new_sf.shape != (n, t, h, w, ci // 16)
|
| 139 |
+
or weight_sf.shape != (co, 3, 3, 3, ci // 16)
|
| 140 |
+
or bias.shape != (co,)
|
| 141 |
+
or (outer_weight is not None and outer_weight.shape != (co,))
|
| 142 |
+
):
|
| 143 |
+
raise RuntimeError(
|
| 144 |
+
"NVFP4 Conv3D accepts Ci=64 or multiples of 128; "
|
| 145 |
+
"other shapes must use the strong-library path"
|
| 146 |
+
)
|
| 147 |
+
return n, t, h, w, co
|
| 148 |
+
|
| 149 |
+
|
| 150 |
+
@torch.library.register_fake(
|
| 151 |
+
add_op_namespace_prefix("nvfp4_causal_conv3d_ndhwc_bf16")
|
| 152 |
+
)
|
| 153 |
+
def _nvfp4_causal_conv3d_fake(
|
| 154 |
+
cache_packed: torch.Tensor,
|
| 155 |
+
new_packed: torch.Tensor,
|
| 156 |
+
weight_packed: torch.Tensor,
|
| 157 |
+
cache_sf: torch.Tensor,
|
| 158 |
+
new_sf: torch.Tensor,
|
| 159 |
+
weight_sf: torch.Tensor,
|
| 160 |
+
bias: torch.Tensor,
|
| 161 |
+
outer_weight: torch.Tensor | None,
|
| 162 |
+
alpha: float,
|
| 163 |
+
out: torch.Tensor,
|
| 164 |
+
) -> None:
|
| 165 |
+
del alpha
|
| 166 |
+
n, t, h, w, co = _check_nvfp4_conv_shapes(
|
| 167 |
+
cache_packed,
|
| 168 |
+
new_packed,
|
| 169 |
+
weight_packed,
|
| 170 |
+
cache_sf,
|
| 171 |
+
new_sf,
|
| 172 |
+
weight_sf,
|
| 173 |
+
bias,
|
| 174 |
+
outer_weight,
|
| 175 |
+
)
|
| 176 |
+
if out.shape != (n, t, h, w, co):
|
| 177 |
+
raise RuntimeError("out must have shape (N,T,H,W,Co)")
|
| 178 |
+
return None
|
| 179 |
+
|
| 180 |
+
|
| 181 |
+
@torch.library.register_fake(
|
| 182 |
+
add_op_namespace_prefix(
|
| 183 |
+
"nvfp4_causal_conv3d_residual_ncdhw_bf16"
|
| 184 |
+
)
|
| 185 |
+
)
|
| 186 |
+
def _nvfp4_causal_conv3d_residual_fake(
|
| 187 |
+
cache_packed: torch.Tensor,
|
| 188 |
+
new_packed: torch.Tensor,
|
| 189 |
+
weight_packed: torch.Tensor,
|
| 190 |
+
cache_sf: torch.Tensor,
|
| 191 |
+
new_sf: torch.Tensor,
|
| 192 |
+
weight_sf: torch.Tensor,
|
| 193 |
+
bias: torch.Tensor,
|
| 194 |
+
residual: torch.Tensor,
|
| 195 |
+
outer_weight: torch.Tensor | None,
|
| 196 |
+
alpha: float,
|
| 197 |
+
out: torch.Tensor,
|
| 198 |
+
) -> None:
|
| 199 |
+
del alpha
|
| 200 |
+
n, t, h, w, co = _check_nvfp4_conv_shapes(
|
| 201 |
+
cache_packed,
|
| 202 |
+
new_packed,
|
| 203 |
+
weight_packed,
|
| 204 |
+
cache_sf,
|
| 205 |
+
new_sf,
|
| 206 |
+
weight_sf,
|
| 207 |
+
bias,
|
| 208 |
+
outer_weight,
|
| 209 |
+
)
|
| 210 |
+
if residual.shape != (n, co, t, h, w) or out.shape != residual.shape:
|
| 211 |
+
raise RuntimeError("residual/out must have shape (N,Co,T,H,W)")
|
| 212 |
+
return None
|
| 213 |
+
|
| 214 |
+
|
| 215 |
def fp8_conv3d_v18_ncdhw_res_bf16out(
|
| 216 |
cache_x: torch.Tensor,
|
| 217 |
new_x: torch.Tensor,
|
|
|
|
| 232 |
return out
|
| 233 |
|
| 234 |
|
| 235 |
+
def fp8_causal_conv3d_ndhwc_bf16(
|
| 236 |
+
cache_x: torch.Tensor,
|
| 237 |
+
new_x: torch.Tensor,
|
| 238 |
+
weight: torch.Tensor,
|
| 239 |
+
bias: torch.Tensor,
|
| 240 |
+
alpha: float = 1.0,
|
| 241 |
+
*,
|
| 242 |
+
out: Optional[torch.Tensor] = None,
|
| 243 |
+
) -> torch.Tensor:
|
| 244 |
+
"""FP8 causal 3D convolution with virtual two-frame cache concat."""
|
| 245 |
+
|
| 246 |
+
n, t, h, w, _ = new_x.shape
|
| 247 |
+
if out is None:
|
| 248 |
+
out = torch.empty(
|
| 249 |
+
(n, t, h, w, weight.shape[0]),
|
| 250 |
+
device=new_x.device,
|
| 251 |
+
dtype=torch.bfloat16,
|
| 252 |
+
)
|
| 253 |
+
ops.fp8_causal_conv3d_ndhwc_bf16(
|
| 254 |
+
cache_x, new_x, weight, bias, float(alpha), out
|
| 255 |
+
)
|
| 256 |
+
return out
|
| 257 |
+
|
| 258 |
+
|
| 259 |
+
def fp8_conv2d_3x3_nhwc_bf16(
|
| 260 |
+
input: torch.Tensor,
|
| 261 |
+
weight: torch.Tensor,
|
| 262 |
+
bias: torch.Tensor,
|
| 263 |
+
alpha: float = 1.0,
|
| 264 |
+
*,
|
| 265 |
+
out: Optional[torch.Tensor] = None,
|
| 266 |
+
) -> torch.Tensor:
|
| 267 |
+
"""FP8 3x3 Conv2D with NHWC input/output and BF16 epilogue."""
|
| 268 |
+
|
| 269 |
+
if out is None:
|
| 270 |
+
out = torch.empty(
|
| 271 |
+
(*input.shape[:3], weight.shape[0]),
|
| 272 |
+
device=input.device,
|
| 273 |
+
dtype=torch.bfloat16,
|
| 274 |
+
)
|
| 275 |
+
ops.fp8_conv2d_3x3_nhwc_bf16(
|
| 276 |
+
input, weight, bias, float(alpha), out
|
| 277 |
+
)
|
| 278 |
+
return out
|
| 279 |
+
|
| 280 |
+
|
| 281 |
+
def fp8_conv2d_3x3_ncdhw_bf16(
|
| 282 |
+
input: torch.Tensor,
|
| 283 |
+
weight: torch.Tensor,
|
| 284 |
+
bias: torch.Tensor,
|
| 285 |
+
alpha: float = 1.0,
|
| 286 |
+
*,
|
| 287 |
+
out: Optional[torch.Tensor] = None,
|
| 288 |
+
) -> torch.Tensor:
|
| 289 |
+
"""FP8 3x3 Conv2D over B*T frames with direct BF16 NCDHW output."""
|
| 290 |
+
|
| 291 |
+
if out is None:
|
| 292 |
+
out = torch.empty(
|
| 293 |
+
(
|
| 294 |
+
input.shape[0],
|
| 295 |
+
weight.shape[0],
|
| 296 |
+
input.shape[1],
|
| 297 |
+
input.shape[2],
|
| 298 |
+
input.shape[3],
|
| 299 |
+
),
|
| 300 |
+
device=input.device,
|
| 301 |
+
dtype=torch.bfloat16,
|
| 302 |
+
)
|
| 303 |
+
ops.fp8_conv2d_3x3_ncdhw_bf16(
|
| 304 |
+
input, weight, bias, float(alpha), out
|
| 305 |
+
)
|
| 306 |
+
return out
|
| 307 |
+
|
| 308 |
+
|
| 309 |
+
def nvfp4_causal_conv3d_ndhwc_bf16(
|
| 310 |
+
cache_packed: torch.Tensor,
|
| 311 |
+
new_packed: torch.Tensor,
|
| 312 |
+
weight_packed: torch.Tensor,
|
| 313 |
+
cache_sf: torch.Tensor,
|
| 314 |
+
new_sf: torch.Tensor,
|
| 315 |
+
weight_sf: torch.Tensor,
|
| 316 |
+
bias: torch.Tensor,
|
| 317 |
+
outer_weight: torch.Tensor | None = None,
|
| 318 |
+
alpha: float = 1.0,
|
| 319 |
+
*,
|
| 320 |
+
out: Optional[torch.Tensor] = None,
|
| 321 |
+
) -> torch.Tensor:
|
| 322 |
+
"""NVFP4 causal Conv3D with linear UE4M3 scale factors and BF16 NDHWC output."""
|
| 323 |
+
|
| 324 |
+
n, t, h, w, _ = new_packed.shape
|
| 325 |
+
if out is None:
|
| 326 |
+
out = torch.empty(
|
| 327 |
+
(n, t, h, w, weight_packed.shape[0]),
|
| 328 |
+
device=new_packed.device,
|
| 329 |
+
dtype=torch.bfloat16,
|
| 330 |
+
)
|
| 331 |
+
ops.nvfp4_causal_conv3d_ndhwc_bf16(
|
| 332 |
+
cache_packed,
|
| 333 |
+
new_packed,
|
| 334 |
+
weight_packed,
|
| 335 |
+
cache_sf,
|
| 336 |
+
new_sf,
|
| 337 |
+
weight_sf,
|
| 338 |
+
bias,
|
| 339 |
+
outer_weight,
|
| 340 |
+
float(alpha),
|
| 341 |
+
out,
|
| 342 |
+
)
|
| 343 |
+
return out
|
| 344 |
+
|
| 345 |
+
|
| 346 |
+
def nvfp4_causal_conv3d_residual_ncdhw_bf16(
|
| 347 |
+
cache_packed: torch.Tensor,
|
| 348 |
+
new_packed: torch.Tensor,
|
| 349 |
+
weight_packed: torch.Tensor,
|
| 350 |
+
cache_sf: torch.Tensor,
|
| 351 |
+
new_sf: torch.Tensor,
|
| 352 |
+
weight_sf: torch.Tensor,
|
| 353 |
+
bias: torch.Tensor,
|
| 354 |
+
residual: torch.Tensor,
|
| 355 |
+
outer_weight: torch.Tensor | None = None,
|
| 356 |
+
alpha: float = 1.0,
|
| 357 |
+
*,
|
| 358 |
+
out: Optional[torch.Tensor] = None,
|
| 359 |
+
) -> torch.Tensor:
|
| 360 |
+
"""NVFP4 causal Conv3D with fused bias, residual, and BF16 NCDHW output."""
|
| 361 |
+
|
| 362 |
+
if out is None:
|
| 363 |
+
out = torch.empty_like(residual)
|
| 364 |
+
ops.nvfp4_causal_conv3d_residual_ncdhw_bf16(
|
| 365 |
+
cache_packed,
|
| 366 |
+
new_packed,
|
| 367 |
+
weight_packed,
|
| 368 |
+
cache_sf,
|
| 369 |
+
new_sf,
|
| 370 |
+
weight_sf,
|
| 371 |
+
bias,
|
| 372 |
+
residual,
|
| 373 |
+
outer_weight,
|
| 374 |
+
float(alpha),
|
| 375 |
+
out,
|
| 376 |
+
)
|
| 377 |
+
return out
|
| 378 |
+
|
| 379 |
+
|
| 380 |
+
__all__ = [
|
| 381 |
+
"fp8_conv3d_v18_ncdhw_res_bf16out",
|
| 382 |
+
"fp8_causal_conv3d_ndhwc_bf16",
|
| 383 |
+
"fp8_conv2d_3x3_nhwc_bf16",
|
| 384 |
+
"fp8_conv2d_3x3_ncdhw_bf16",
|
| 385 |
+
"nvfp4_causal_conv3d_ndhwc_bf16",
|
| 386 |
+
"nvfp4_causal_conv3d_residual_ncdhw_bf16",
|
| 387 |
+
]
|
build/torch211-cxx11-cu130-x86_64-linux/_ops.py
CHANGED
|
@@ -1,9 +1,9 @@
|
|
| 1 |
import torch
|
| 2 |
-
from . import
|
| 3 |
-
ops = torch.ops.
|
| 4 |
|
| 5 |
def add_op_namespace_prefix(op_name: str):
|
| 6 |
"""
|
| 7 |
Prefix op by namespace.
|
| 8 |
"""
|
| 9 |
-
return f"
|
|
|
|
| 1 |
import torch
|
| 2 |
+
from . import _world_model_conv_cuda_7781728
|
| 3 |
+
ops = torch.ops._world_model_conv_cuda_7781728
|
| 4 |
|
| 5 |
def add_op_namespace_prefix(op_name: str):
|
| 6 |
"""
|
| 7 |
Prefix op by namespace.
|
| 8 |
"""
|
| 9 |
+
return f"_world_model_conv_cuda_7781728::{op_name}"
|
build/torch211-cxx11-cu130-x86_64-linux/{_world_model_conv_cuda_f14c443.abi3.so → _world_model_conv_cuda_7781728.abi3.so}
RENAMED
|
@@ -1,3 +1,3 @@
|
|
| 1 |
version https://git-lfs.github.com/spec/v1
|
| 2 |
-
oid sha256:
|
| 3 |
-
size
|
|
|
|
| 1 |
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:8dfb73bc6c3dd31b73b180a9781a09a736607d33fb7846cf3b5d6d9dd99804cd
|
| 3 |
+
size 1242552
|
build/torch211-cxx11-cu130-x86_64-linux/metadata.json
CHANGED
|
@@ -1,6 +1,6 @@
|
|
| 1 |
{
|
| 2 |
"name": "world-model-conv",
|
| 3 |
-
"id": "
|
| 4 |
"version": 1,
|
| 5 |
"license": "Apache-2.0",
|
| 6 |
"python-depends": [],
|
|
|
|
| 1 |
{
|
| 2 |
"name": "world-model-conv",
|
| 3 |
+
"id": "_world_model_conv_cuda_7781728",
|
| 4 |
"version": 1,
|
| 5 |
"license": "Apache-2.0",
|
| 6 |
"python-depends": [],
|
build/torch212-cxx11-cu130-x86_64-linux/__init__.py
CHANGED
|
@@ -27,6 +27,11 @@ def _fp8_conv3d_fake(
|
|
| 27 |
raise RuntimeError("cache_x must have shape (N,2,H,W,Ci)")
|
| 28 |
if weight.shape != (co, 3, 3, 3, ci):
|
| 29 |
raise RuntimeError("weight must have shape (Co,3,3,3,Ci)")
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 30 |
if residual.shape != (n, co, t_new, h, w) or out.shape != residual.shape:
|
| 31 |
raise RuntimeError("residual/out must be NCDHW")
|
| 32 |
if bias.shape != (co,):
|
|
@@ -34,6 +39,179 @@ def _fp8_conv3d_fake(
|
|
| 34 |
return None
|
| 35 |
|
| 36 |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 37 |
def fp8_conv3d_v18_ncdhw_res_bf16out(
|
| 38 |
cache_x: torch.Tensor,
|
| 39 |
new_x: torch.Tensor,
|
|
@@ -54,4 +232,156 @@ def fp8_conv3d_v18_ncdhw_res_bf16out(
|
|
| 54 |
return out
|
| 55 |
|
| 56 |
|
| 57 |
-
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 27 |
raise RuntimeError("cache_x must have shape (N,2,H,W,Ci)")
|
| 28 |
if weight.shape != (co, 3, 3, 3, ci):
|
| 29 |
raise RuntimeError("weight must have shape (Co,3,3,3,Ci)")
|
| 30 |
+
if ci not in (32, 64):
|
| 31 |
+
raise RuntimeError(
|
| 32 |
+
"FP8 Conv3D is accepted only for Ci=32/64; use the "
|
| 33 |
+
"strong-library or NVFP4 path for larger channels"
|
| 34 |
+
)
|
| 35 |
if residual.shape != (n, co, t_new, h, w) or out.shape != residual.shape:
|
| 36 |
raise RuntimeError("residual/out must be NCDHW")
|
| 37 |
if bias.shape != (co,):
|
|
|
|
| 39 |
return None
|
| 40 |
|
| 41 |
|
| 42 |
+
@torch.library.register_fake(
|
| 43 |
+
add_op_namespace_prefix("fp8_causal_conv3d_ndhwc_bf16")
|
| 44 |
+
)
|
| 45 |
+
def _fp8_causal_conv3d_fake(
|
| 46 |
+
cache_x: torch.Tensor,
|
| 47 |
+
new_x: torch.Tensor,
|
| 48 |
+
weight: torch.Tensor,
|
| 49 |
+
bias: torch.Tensor,
|
| 50 |
+
alpha: float,
|
| 51 |
+
out: torch.Tensor,
|
| 52 |
+
) -> None:
|
| 53 |
+
if cache_x.dim() != 5 or new_x.dim() != 5:
|
| 54 |
+
raise RuntimeError("cache_x/new_x must be NDHWC")
|
| 55 |
+
n, t, h, w, ci = new_x.shape
|
| 56 |
+
co = weight.shape[0]
|
| 57 |
+
if cache_x.shape != (n, 2, h, w, ci):
|
| 58 |
+
raise RuntimeError("cache_x must have shape (N,2,H,W,Ci)")
|
| 59 |
+
if weight.shape != (co, 3, 3, 3, ci):
|
| 60 |
+
raise RuntimeError("weight must have shape (Co,3,3,3,Ci)")
|
| 61 |
+
if ci not in (32, 64):
|
| 62 |
+
raise RuntimeError(
|
| 63 |
+
"FP8 Conv3D is accepted only for Ci=32/64; use the "
|
| 64 |
+
"strong-library or NVFP4 path for larger channels"
|
| 65 |
+
)
|
| 66 |
+
if bias.shape != (co,) or out.shape != (n, t, h, w, co):
|
| 67 |
+
raise RuntimeError("bias/out shape mismatch")
|
| 68 |
+
return None
|
| 69 |
+
|
| 70 |
+
|
| 71 |
+
@torch.library.register_fake(
|
| 72 |
+
add_op_namespace_prefix("fp8_conv2d_3x3_nhwc_bf16")
|
| 73 |
+
)
|
| 74 |
+
def _fp8_conv2d_fake(
|
| 75 |
+
input: torch.Tensor,
|
| 76 |
+
weight: torch.Tensor,
|
| 77 |
+
bias: torch.Tensor,
|
| 78 |
+
alpha: float,
|
| 79 |
+
out: torch.Tensor,
|
| 80 |
+
) -> None:
|
| 81 |
+
if input.dim() != 4:
|
| 82 |
+
raise RuntimeError("input must have shape (N,H,W,Ci)")
|
| 83 |
+
n, h, w, ci = input.shape
|
| 84 |
+
co = weight.shape[0]
|
| 85 |
+
if weight.shape != (co, 3, 3, ci):
|
| 86 |
+
raise RuntimeError("weight must have shape (Co,3,3,Ci)")
|
| 87 |
+
if bias.shape != (co,) or out.shape != (n, h, w, co):
|
| 88 |
+
raise RuntimeError("bias/out shape mismatch")
|
| 89 |
+
return None
|
| 90 |
+
|
| 91 |
+
|
| 92 |
+
@torch.library.register_fake(
|
| 93 |
+
add_op_namespace_prefix("fp8_conv2d_3x3_ncdhw_bf16")
|
| 94 |
+
)
|
| 95 |
+
def _fp8_conv2d_ncdhw_fake(
|
| 96 |
+
input: torch.Tensor,
|
| 97 |
+
weight: torch.Tensor,
|
| 98 |
+
bias: torch.Tensor,
|
| 99 |
+
alpha: float,
|
| 100 |
+
out: torch.Tensor,
|
| 101 |
+
) -> None:
|
| 102 |
+
if input.dim() != 5:
|
| 103 |
+
raise RuntimeError("input must have shape (B,T,H,W,Ci)")
|
| 104 |
+
b, t, h, w, ci = input.shape
|
| 105 |
+
co = weight.shape[0]
|
| 106 |
+
if weight.shape != (co, 3, 3, ci):
|
| 107 |
+
raise RuntimeError("weight must have shape (Co,3,3,Ci)")
|
| 108 |
+
if bias.shape != (co,) or out.shape != (b, co, t, h, w):
|
| 109 |
+
raise RuntimeError("bias/out shape mismatch")
|
| 110 |
+
return None
|
| 111 |
+
|
| 112 |
+
|
| 113 |
+
def _check_nvfp4_conv_shapes(
|
| 114 |
+
cache_packed: torch.Tensor,
|
| 115 |
+
new_packed: torch.Tensor,
|
| 116 |
+
weight_packed: torch.Tensor,
|
| 117 |
+
cache_sf: torch.Tensor,
|
| 118 |
+
new_sf: torch.Tensor,
|
| 119 |
+
weight_sf: torch.Tensor,
|
| 120 |
+
bias: torch.Tensor,
|
| 121 |
+
outer_weight: torch.Tensor | None,
|
| 122 |
+
) -> tuple[int, int, int, int, int]:
|
| 123 |
+
if (
|
| 124 |
+
cache_packed.dim() != 5
|
| 125 |
+
or new_packed.dim() != 5
|
| 126 |
+
or weight_packed.dim() != 5
|
| 127 |
+
):
|
| 128 |
+
raise RuntimeError("packed inputs must be five-dimensional")
|
| 129 |
+
n, t, h, w, ci_half = new_packed.shape
|
| 130 |
+
ci = ci_half * 2
|
| 131 |
+
co = weight_packed.shape[0]
|
| 132 |
+
if (
|
| 133 |
+
(ci != 64 and ci % 128 != 0)
|
| 134 |
+
or co % 8 != 0
|
| 135 |
+
or cache_packed.shape != (n, 2, h, w, ci // 2)
|
| 136 |
+
or weight_packed.shape != (co, 3, 3, 3, ci // 2)
|
| 137 |
+
or cache_sf.shape != (n, 2, h, w, ci // 16)
|
| 138 |
+
or new_sf.shape != (n, t, h, w, ci // 16)
|
| 139 |
+
or weight_sf.shape != (co, 3, 3, 3, ci // 16)
|
| 140 |
+
or bias.shape != (co,)
|
| 141 |
+
or (outer_weight is not None and outer_weight.shape != (co,))
|
| 142 |
+
):
|
| 143 |
+
raise RuntimeError(
|
| 144 |
+
"NVFP4 Conv3D accepts Ci=64 or multiples of 128; "
|
| 145 |
+
"other shapes must use the strong-library path"
|
| 146 |
+
)
|
| 147 |
+
return n, t, h, w, co
|
| 148 |
+
|
| 149 |
+
|
| 150 |
+
@torch.library.register_fake(
|
| 151 |
+
add_op_namespace_prefix("nvfp4_causal_conv3d_ndhwc_bf16")
|
| 152 |
+
)
|
| 153 |
+
def _nvfp4_causal_conv3d_fake(
|
| 154 |
+
cache_packed: torch.Tensor,
|
| 155 |
+
new_packed: torch.Tensor,
|
| 156 |
+
weight_packed: torch.Tensor,
|
| 157 |
+
cache_sf: torch.Tensor,
|
| 158 |
+
new_sf: torch.Tensor,
|
| 159 |
+
weight_sf: torch.Tensor,
|
| 160 |
+
bias: torch.Tensor,
|
| 161 |
+
outer_weight: torch.Tensor | None,
|
| 162 |
+
alpha: float,
|
| 163 |
+
out: torch.Tensor,
|
| 164 |
+
) -> None:
|
| 165 |
+
del alpha
|
| 166 |
+
n, t, h, w, co = _check_nvfp4_conv_shapes(
|
| 167 |
+
cache_packed,
|
| 168 |
+
new_packed,
|
| 169 |
+
weight_packed,
|
| 170 |
+
cache_sf,
|
| 171 |
+
new_sf,
|
| 172 |
+
weight_sf,
|
| 173 |
+
bias,
|
| 174 |
+
outer_weight,
|
| 175 |
+
)
|
| 176 |
+
if out.shape != (n, t, h, w, co):
|
| 177 |
+
raise RuntimeError("out must have shape (N,T,H,W,Co)")
|
| 178 |
+
return None
|
| 179 |
+
|
| 180 |
+
|
| 181 |
+
@torch.library.register_fake(
|
| 182 |
+
add_op_namespace_prefix(
|
| 183 |
+
"nvfp4_causal_conv3d_residual_ncdhw_bf16"
|
| 184 |
+
)
|
| 185 |
+
)
|
| 186 |
+
def _nvfp4_causal_conv3d_residual_fake(
|
| 187 |
+
cache_packed: torch.Tensor,
|
| 188 |
+
new_packed: torch.Tensor,
|
| 189 |
+
weight_packed: torch.Tensor,
|
| 190 |
+
cache_sf: torch.Tensor,
|
| 191 |
+
new_sf: torch.Tensor,
|
| 192 |
+
weight_sf: torch.Tensor,
|
| 193 |
+
bias: torch.Tensor,
|
| 194 |
+
residual: torch.Tensor,
|
| 195 |
+
outer_weight: torch.Tensor | None,
|
| 196 |
+
alpha: float,
|
| 197 |
+
out: torch.Tensor,
|
| 198 |
+
) -> None:
|
| 199 |
+
del alpha
|
| 200 |
+
n, t, h, w, co = _check_nvfp4_conv_shapes(
|
| 201 |
+
cache_packed,
|
| 202 |
+
new_packed,
|
| 203 |
+
weight_packed,
|
| 204 |
+
cache_sf,
|
| 205 |
+
new_sf,
|
| 206 |
+
weight_sf,
|
| 207 |
+
bias,
|
| 208 |
+
outer_weight,
|
| 209 |
+
)
|
| 210 |
+
if residual.shape != (n, co, t, h, w) or out.shape != residual.shape:
|
| 211 |
+
raise RuntimeError("residual/out must have shape (N,Co,T,H,W)")
|
| 212 |
+
return None
|
| 213 |
+
|
| 214 |
+
|
| 215 |
def fp8_conv3d_v18_ncdhw_res_bf16out(
|
| 216 |
cache_x: torch.Tensor,
|
| 217 |
new_x: torch.Tensor,
|
|
|
|
| 232 |
return out
|
| 233 |
|
| 234 |
|
| 235 |
+
def fp8_causal_conv3d_ndhwc_bf16(
|
| 236 |
+
cache_x: torch.Tensor,
|
| 237 |
+
new_x: torch.Tensor,
|
| 238 |
+
weight: torch.Tensor,
|
| 239 |
+
bias: torch.Tensor,
|
| 240 |
+
alpha: float = 1.0,
|
| 241 |
+
*,
|
| 242 |
+
out: Optional[torch.Tensor] = None,
|
| 243 |
+
) -> torch.Tensor:
|
| 244 |
+
"""FP8 causal 3D convolution with virtual two-frame cache concat."""
|
| 245 |
+
|
| 246 |
+
n, t, h, w, _ = new_x.shape
|
| 247 |
+
if out is None:
|
| 248 |
+
out = torch.empty(
|
| 249 |
+
(n, t, h, w, weight.shape[0]),
|
| 250 |
+
device=new_x.device,
|
| 251 |
+
dtype=torch.bfloat16,
|
| 252 |
+
)
|
| 253 |
+
ops.fp8_causal_conv3d_ndhwc_bf16(
|
| 254 |
+
cache_x, new_x, weight, bias, float(alpha), out
|
| 255 |
+
)
|
| 256 |
+
return out
|
| 257 |
+
|
| 258 |
+
|
| 259 |
+
def fp8_conv2d_3x3_nhwc_bf16(
|
| 260 |
+
input: torch.Tensor,
|
| 261 |
+
weight: torch.Tensor,
|
| 262 |
+
bias: torch.Tensor,
|
| 263 |
+
alpha: float = 1.0,
|
| 264 |
+
*,
|
| 265 |
+
out: Optional[torch.Tensor] = None,
|
| 266 |
+
) -> torch.Tensor:
|
| 267 |
+
"""FP8 3x3 Conv2D with NHWC input/output and BF16 epilogue."""
|
| 268 |
+
|
| 269 |
+
if out is None:
|
| 270 |
+
out = torch.empty(
|
| 271 |
+
(*input.shape[:3], weight.shape[0]),
|
| 272 |
+
device=input.device,
|
| 273 |
+
dtype=torch.bfloat16,
|
| 274 |
+
)
|
| 275 |
+
ops.fp8_conv2d_3x3_nhwc_bf16(
|
| 276 |
+
input, weight, bias, float(alpha), out
|
| 277 |
+
)
|
| 278 |
+
return out
|
| 279 |
+
|
| 280 |
+
|
| 281 |
+
def fp8_conv2d_3x3_ncdhw_bf16(
|
| 282 |
+
input: torch.Tensor,
|
| 283 |
+
weight: torch.Tensor,
|
| 284 |
+
bias: torch.Tensor,
|
| 285 |
+
alpha: float = 1.0,
|
| 286 |
+
*,
|
| 287 |
+
out: Optional[torch.Tensor] = None,
|
| 288 |
+
) -> torch.Tensor:
|
| 289 |
+
"""FP8 3x3 Conv2D over B*T frames with direct BF16 NCDHW output."""
|
| 290 |
+
|
| 291 |
+
if out is None:
|
| 292 |
+
out = torch.empty(
|
| 293 |
+
(
|
| 294 |
+
input.shape[0],
|
| 295 |
+
weight.shape[0],
|
| 296 |
+
input.shape[1],
|
| 297 |
+
input.shape[2],
|
| 298 |
+
input.shape[3],
|
| 299 |
+
),
|
| 300 |
+
device=input.device,
|
| 301 |
+
dtype=torch.bfloat16,
|
| 302 |
+
)
|
| 303 |
+
ops.fp8_conv2d_3x3_ncdhw_bf16(
|
| 304 |
+
input, weight, bias, float(alpha), out
|
| 305 |
+
)
|
| 306 |
+
return out
|
| 307 |
+
|
| 308 |
+
|
| 309 |
+
def nvfp4_causal_conv3d_ndhwc_bf16(
|
| 310 |
+
cache_packed: torch.Tensor,
|
| 311 |
+
new_packed: torch.Tensor,
|
| 312 |
+
weight_packed: torch.Tensor,
|
| 313 |
+
cache_sf: torch.Tensor,
|
| 314 |
+
new_sf: torch.Tensor,
|
| 315 |
+
weight_sf: torch.Tensor,
|
| 316 |
+
bias: torch.Tensor,
|
| 317 |
+
outer_weight: torch.Tensor | None = None,
|
| 318 |
+
alpha: float = 1.0,
|
| 319 |
+
*,
|
| 320 |
+
out: Optional[torch.Tensor] = None,
|
| 321 |
+
) -> torch.Tensor:
|
| 322 |
+
"""NVFP4 causal Conv3D with linear UE4M3 scale factors and BF16 NDHWC output."""
|
| 323 |
+
|
| 324 |
+
n, t, h, w, _ = new_packed.shape
|
| 325 |
+
if out is None:
|
| 326 |
+
out = torch.empty(
|
| 327 |
+
(n, t, h, w, weight_packed.shape[0]),
|
| 328 |
+
device=new_packed.device,
|
| 329 |
+
dtype=torch.bfloat16,
|
| 330 |
+
)
|
| 331 |
+
ops.nvfp4_causal_conv3d_ndhwc_bf16(
|
| 332 |
+
cache_packed,
|
| 333 |
+
new_packed,
|
| 334 |
+
weight_packed,
|
| 335 |
+
cache_sf,
|
| 336 |
+
new_sf,
|
| 337 |
+
weight_sf,
|
| 338 |
+
bias,
|
| 339 |
+
outer_weight,
|
| 340 |
+
float(alpha),
|
| 341 |
+
out,
|
| 342 |
+
)
|
| 343 |
+
return out
|
| 344 |
+
|
| 345 |
+
|
| 346 |
+
def nvfp4_causal_conv3d_residual_ncdhw_bf16(
|
| 347 |
+
cache_packed: torch.Tensor,
|
| 348 |
+
new_packed: torch.Tensor,
|
| 349 |
+
weight_packed: torch.Tensor,
|
| 350 |
+
cache_sf: torch.Tensor,
|
| 351 |
+
new_sf: torch.Tensor,
|
| 352 |
+
weight_sf: torch.Tensor,
|
| 353 |
+
bias: torch.Tensor,
|
| 354 |
+
residual: torch.Tensor,
|
| 355 |
+
outer_weight: torch.Tensor | None = None,
|
| 356 |
+
alpha: float = 1.0,
|
| 357 |
+
*,
|
| 358 |
+
out: Optional[torch.Tensor] = None,
|
| 359 |
+
) -> torch.Tensor:
|
| 360 |
+
"""NVFP4 causal Conv3D with fused bias, residual, and BF16 NCDHW output."""
|
| 361 |
+
|
| 362 |
+
if out is None:
|
| 363 |
+
out = torch.empty_like(residual)
|
| 364 |
+
ops.nvfp4_causal_conv3d_residual_ncdhw_bf16(
|
| 365 |
+
cache_packed,
|
| 366 |
+
new_packed,
|
| 367 |
+
weight_packed,
|
| 368 |
+
cache_sf,
|
| 369 |
+
new_sf,
|
| 370 |
+
weight_sf,
|
| 371 |
+
bias,
|
| 372 |
+
residual,
|
| 373 |
+
outer_weight,
|
| 374 |
+
float(alpha),
|
| 375 |
+
out,
|
| 376 |
+
)
|
| 377 |
+
return out
|
| 378 |
+
|
| 379 |
+
|
| 380 |
+
__all__ = [
|
| 381 |
+
"fp8_conv3d_v18_ncdhw_res_bf16out",
|
| 382 |
+
"fp8_causal_conv3d_ndhwc_bf16",
|
| 383 |
+
"fp8_conv2d_3x3_nhwc_bf16",
|
| 384 |
+
"fp8_conv2d_3x3_ncdhw_bf16",
|
| 385 |
+
"nvfp4_causal_conv3d_ndhwc_bf16",
|
| 386 |
+
"nvfp4_causal_conv3d_residual_ncdhw_bf16",
|
| 387 |
+
]
|
build/torch212-cxx11-cu130-x86_64-linux/_ops.py
CHANGED
|
@@ -1,9 +1,9 @@
|
|
| 1 |
import torch
|
| 2 |
-
from . import
|
| 3 |
-
ops = torch.ops.
|
| 4 |
|
| 5 |
def add_op_namespace_prefix(op_name: str):
|
| 6 |
"""
|
| 7 |
Prefix op by namespace.
|
| 8 |
"""
|
| 9 |
-
return f"
|
|
|
|
| 1 |
import torch
|
| 2 |
+
from . import _world_model_conv_cuda_7781728
|
| 3 |
+
ops = torch.ops._world_model_conv_cuda_7781728
|
| 4 |
|
| 5 |
def add_op_namespace_prefix(op_name: str):
|
| 6 |
"""
|
| 7 |
Prefix op by namespace.
|
| 8 |
"""
|
| 9 |
+
return f"_world_model_conv_cuda_7781728::{op_name}"
|
build/torch212-cxx11-cu130-x86_64-linux/{_world_model_conv_cuda_f14c443.abi3.so → _world_model_conv_cuda_7781728.abi3.so}
RENAMED
|
@@ -1,3 +1,3 @@
|
|
| 1 |
version https://git-lfs.github.com/spec/v1
|
| 2 |
-
oid sha256:
|
| 3 |
-
size
|
|
|
|
| 1 |
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:21963c9f0459cc516906960624490bd34898ac669436bf3a4e6cfd4f6f2bc9c7
|
| 3 |
+
size 1248816
|
build/torch212-cxx11-cu130-x86_64-linux/metadata.json
CHANGED
|
@@ -1,6 +1,6 @@
|
|
| 1 |
{
|
| 2 |
"name": "world-model-conv",
|
| 3 |
-
"id": "
|
| 4 |
"version": 1,
|
| 5 |
"license": "Apache-2.0",
|
| 6 |
"python-depends": [],
|
|
|
|
| 1 |
{
|
| 2 |
"name": "world-model-conv",
|
| 3 |
+
"id": "_world_model_conv_cuda_7781728",
|
| 4 |
"version": 1,
|
| 5 |
"license": "Apache-2.0",
|
| 6 |
"python-depends": [],
|
build/torch212-cxx11-cu132-x86_64-linux/__init__.py
CHANGED
|
@@ -27,6 +27,11 @@ def _fp8_conv3d_fake(
|
|
| 27 |
raise RuntimeError("cache_x must have shape (N,2,H,W,Ci)")
|
| 28 |
if weight.shape != (co, 3, 3, 3, ci):
|
| 29 |
raise RuntimeError("weight must have shape (Co,3,3,3,Ci)")
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 30 |
if residual.shape != (n, co, t_new, h, w) or out.shape != residual.shape:
|
| 31 |
raise RuntimeError("residual/out must be NCDHW")
|
| 32 |
if bias.shape != (co,):
|
|
@@ -34,6 +39,179 @@ def _fp8_conv3d_fake(
|
|
| 34 |
return None
|
| 35 |
|
| 36 |
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 37 |
def fp8_conv3d_v18_ncdhw_res_bf16out(
|
| 38 |
cache_x: torch.Tensor,
|
| 39 |
new_x: torch.Tensor,
|
|
@@ -54,4 +232,156 @@ def fp8_conv3d_v18_ncdhw_res_bf16out(
|
|
| 54 |
return out
|
| 55 |
|
| 56 |
|
| 57 |
-
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
|
| 27 |
raise RuntimeError("cache_x must have shape (N,2,H,W,Ci)")
|
| 28 |
if weight.shape != (co, 3, 3, 3, ci):
|
| 29 |
raise RuntimeError("weight must have shape (Co,3,3,3,Ci)")
|
| 30 |
+
if ci not in (32, 64):
|
| 31 |
+
raise RuntimeError(
|
| 32 |
+
"FP8 Conv3D is accepted only for Ci=32/64; use the "
|
| 33 |
+
"strong-library or NVFP4 path for larger channels"
|
| 34 |
+
)
|
| 35 |
if residual.shape != (n, co, t_new, h, w) or out.shape != residual.shape:
|
| 36 |
raise RuntimeError("residual/out must be NCDHW")
|
| 37 |
if bias.shape != (co,):
|
|
|
|
| 39 |
return None
|
| 40 |
|
| 41 |
|
| 42 |
+
@torch.library.register_fake(
|
| 43 |
+
add_op_namespace_prefix("fp8_causal_conv3d_ndhwc_bf16")
|
| 44 |
+
)
|
| 45 |
+
def _fp8_causal_conv3d_fake(
|
| 46 |
+
cache_x: torch.Tensor,
|
| 47 |
+
new_x: torch.Tensor,
|
| 48 |
+
weight: torch.Tensor,
|
| 49 |
+
bias: torch.Tensor,
|
| 50 |
+
alpha: float,
|
| 51 |
+
out: torch.Tensor,
|
| 52 |
+
) -> None:
|
| 53 |
+
if cache_x.dim() != 5 or new_x.dim() != 5:
|
| 54 |
+
raise RuntimeError("cache_x/new_x must be NDHWC")
|
| 55 |
+
n, t, h, w, ci = new_x.shape
|
| 56 |
+
co = weight.shape[0]
|
| 57 |
+
if cache_x.shape != (n, 2, h, w, ci):
|
| 58 |
+
raise RuntimeError("cache_x must have shape (N,2,H,W,Ci)")
|
| 59 |
+
if weight.shape != (co, 3, 3, 3, ci):
|
| 60 |
+
raise RuntimeError("weight must have shape (Co,3,3,3,Ci)")
|
| 61 |
+
if ci not in (32, 64):
|
| 62 |
+
raise RuntimeError(
|
| 63 |
+
"FP8 Conv3D is accepted only for Ci=32/64; use the "
|
| 64 |
+
"strong-library or NVFP4 path for larger channels"
|
| 65 |
+
)
|
| 66 |
+
if bias.shape != (co,) or out.shape != (n, t, h, w, co):
|
| 67 |
+
raise RuntimeError("bias/out shape mismatch")
|
| 68 |
+
return None
|
| 69 |
+
|
| 70 |
+
|
| 71 |
+
@torch.library.register_fake(
|
| 72 |
+
add_op_namespace_prefix("fp8_conv2d_3x3_nhwc_bf16")
|
| 73 |
+
)
|
| 74 |
+
def _fp8_conv2d_fake(
|
| 75 |
+
input: torch.Tensor,
|
| 76 |
+
weight: torch.Tensor,
|
| 77 |
+
bias: torch.Tensor,
|
| 78 |
+
alpha: float,
|
| 79 |
+
out: torch.Tensor,
|
| 80 |
+
) -> None:
|
| 81 |
+
if input.dim() != 4:
|
| 82 |
+
raise RuntimeError("input must have shape (N,H,W,Ci)")
|
| 83 |
+
n, h, w, ci = input.shape
|
| 84 |
+
co = weight.shape[0]
|
| 85 |
+
if weight.shape != (co, 3, 3, ci):
|
| 86 |
+
raise RuntimeError("weight must have shape (Co,3,3,Ci)")
|
| 87 |
+
if bias.shape != (co,) or out.shape != (n, h, w, co):
|
| 88 |
+
raise RuntimeError("bias/out shape mismatch")
|
| 89 |
+
return None
|
| 90 |
+
|
| 91 |
+
|
| 92 |
+
@torch.library.register_fake(
|
| 93 |
+
add_op_namespace_prefix("fp8_conv2d_3x3_ncdhw_bf16")
|
| 94 |
+
)
|
| 95 |
+
def _fp8_conv2d_ncdhw_fake(
|
| 96 |
+
input: torch.Tensor,
|
| 97 |
+
weight: torch.Tensor,
|
| 98 |
+
bias: torch.Tensor,
|
| 99 |
+
alpha: float,
|
| 100 |
+
out: torch.Tensor,
|
| 101 |
+
) -> None:
|
| 102 |
+
if input.dim() != 5:
|
| 103 |
+
raise RuntimeError("input must have shape (B,T,H,W,Ci)")
|
| 104 |
+
b, t, h, w, ci = input.shape
|
| 105 |
+
co = weight.shape[0]
|
| 106 |
+
if weight.shape != (co, 3, 3, ci):
|
| 107 |
+
raise RuntimeError("weight must have shape (Co,3,3,Ci)")
|
| 108 |
+
if bias.shape != (co,) or out.shape != (b, co, t, h, w):
|
| 109 |
+
raise RuntimeError("bias/out shape mismatch")
|
| 110 |
+
return None
|
| 111 |
+
|
| 112 |
+
|
| 113 |
+
def _check_nvfp4_conv_shapes(
|
| 114 |
+
cache_packed: torch.Tensor,
|
| 115 |
+
new_packed: torch.Tensor,
|
| 116 |
+
weight_packed: torch.Tensor,
|
| 117 |
+
cache_sf: torch.Tensor,
|
| 118 |
+
new_sf: torch.Tensor,
|
| 119 |
+
weight_sf: torch.Tensor,
|
| 120 |
+
bias: torch.Tensor,
|
| 121 |
+
outer_weight: torch.Tensor | None,
|
| 122 |
+
) -> tuple[int, int, int, int, int]:
|
| 123 |
+
if (
|
| 124 |
+
cache_packed.dim() != 5
|
| 125 |
+
or new_packed.dim() != 5
|
| 126 |
+
or weight_packed.dim() != 5
|
| 127 |
+
):
|
| 128 |
+
raise RuntimeError("packed inputs must be five-dimensional")
|
| 129 |
+
n, t, h, w, ci_half = new_packed.shape
|
| 130 |
+
ci = ci_half * 2
|
| 131 |
+
co = weight_packed.shape[0]
|
| 132 |
+
if (
|
| 133 |
+
(ci != 64 and ci % 128 != 0)
|
| 134 |
+
or co % 8 != 0
|
| 135 |
+
or cache_packed.shape != (n, 2, h, w, ci // 2)
|
| 136 |
+
or weight_packed.shape != (co, 3, 3, 3, ci // 2)
|
| 137 |
+
or cache_sf.shape != (n, 2, h, w, ci // 16)
|
| 138 |
+
or new_sf.shape != (n, t, h, w, ci // 16)
|
| 139 |
+
or weight_sf.shape != (co, 3, 3, 3, ci // 16)
|
| 140 |
+
or bias.shape != (co,)
|
| 141 |
+
or (outer_weight is not None and outer_weight.shape != (co,))
|
| 142 |
+
):
|
| 143 |
+
raise RuntimeError(
|
| 144 |
+
"NVFP4 Conv3D accepts Ci=64 or multiples of 128; "
|
| 145 |
+
"other shapes must use the strong-library path"
|
| 146 |
+
)
|
| 147 |
+
return n, t, h, w, co
|
| 148 |
+
|
| 149 |
+
|
| 150 |
+
@torch.library.register_fake(
|
| 151 |
+
add_op_namespace_prefix("nvfp4_causal_conv3d_ndhwc_bf16")
|
| 152 |
+
)
|
| 153 |
+
def _nvfp4_causal_conv3d_fake(
|
| 154 |
+
cache_packed: torch.Tensor,
|
| 155 |
+
new_packed: torch.Tensor,
|
| 156 |
+
weight_packed: torch.Tensor,
|
| 157 |
+
cache_sf: torch.Tensor,
|
| 158 |
+
new_sf: torch.Tensor,
|
| 159 |
+
weight_sf: torch.Tensor,
|
| 160 |
+
bias: torch.Tensor,
|
| 161 |
+
outer_weight: torch.Tensor | None,
|
| 162 |
+
alpha: float,
|
| 163 |
+
out: torch.Tensor,
|
| 164 |
+
) -> None:
|
| 165 |
+
del alpha
|
| 166 |
+
n, t, h, w, co = _check_nvfp4_conv_shapes(
|
| 167 |
+
cache_packed,
|
| 168 |
+
new_packed,
|
| 169 |
+
weight_packed,
|
| 170 |
+
cache_sf,
|
| 171 |
+
new_sf,
|
| 172 |
+
weight_sf,
|
| 173 |
+
bias,
|
| 174 |
+
outer_weight,
|
| 175 |
+
)
|
| 176 |
+
if out.shape != (n, t, h, w, co):
|
| 177 |
+
raise RuntimeError("out must have shape (N,T,H,W,Co)")
|
| 178 |
+
return None
|
| 179 |
+
|
| 180 |
+
|
| 181 |
+
@torch.library.register_fake(
|
| 182 |
+
add_op_namespace_prefix(
|
| 183 |
+
"nvfp4_causal_conv3d_residual_ncdhw_bf16"
|
| 184 |
+
)
|
| 185 |
+
)
|
| 186 |
+
def _nvfp4_causal_conv3d_residual_fake(
|
| 187 |
+
cache_packed: torch.Tensor,
|
| 188 |
+
new_packed: torch.Tensor,
|
| 189 |
+
weight_packed: torch.Tensor,
|
| 190 |
+
cache_sf: torch.Tensor,
|
| 191 |
+
new_sf: torch.Tensor,
|
| 192 |
+
weight_sf: torch.Tensor,
|
| 193 |
+
bias: torch.Tensor,
|
| 194 |
+
residual: torch.Tensor,
|
| 195 |
+
outer_weight: torch.Tensor | None,
|
| 196 |
+
alpha: float,
|
| 197 |
+
out: torch.Tensor,
|
| 198 |
+
) -> None:
|
| 199 |
+
del alpha
|
| 200 |
+
n, t, h, w, co = _check_nvfp4_conv_shapes(
|
| 201 |
+
cache_packed,
|
| 202 |
+
new_packed,
|
| 203 |
+
weight_packed,
|
| 204 |
+
cache_sf,
|
| 205 |
+
new_sf,
|
| 206 |
+
weight_sf,
|
| 207 |
+
bias,
|
| 208 |
+
outer_weight,
|
| 209 |
+
)
|
| 210 |
+
if residual.shape != (n, co, t, h, w) or out.shape != residual.shape:
|
| 211 |
+
raise RuntimeError("residual/out must have shape (N,Co,T,H,W)")
|
| 212 |
+
return None
|
| 213 |
+
|
| 214 |
+
|
| 215 |
def fp8_conv3d_v18_ncdhw_res_bf16out(
|
| 216 |
cache_x: torch.Tensor,
|
| 217 |
new_x: torch.Tensor,
|
|
|
|
| 232 |
return out
|
| 233 |
|
| 234 |
|
| 235 |
+
def fp8_causal_conv3d_ndhwc_bf16(
|
| 236 |
+
cache_x: torch.Tensor,
|
| 237 |
+
new_x: torch.Tensor,
|
| 238 |
+
weight: torch.Tensor,
|
| 239 |
+
bias: torch.Tensor,
|
| 240 |
+
alpha: float = 1.0,
|
| 241 |
+
*,
|
| 242 |
+
out: Optional[torch.Tensor] = None,
|
| 243 |
+
) -> torch.Tensor:
|
| 244 |
+
"""FP8 causal 3D convolution with virtual two-frame cache concat."""
|
| 245 |
+
|
| 246 |
+
n, t, h, w, _ = new_x.shape
|
| 247 |
+
if out is None:
|
| 248 |
+
out = torch.empty(
|
| 249 |
+
(n, t, h, w, weight.shape[0]),
|
| 250 |
+
device=new_x.device,
|
| 251 |
+
dtype=torch.bfloat16,
|
| 252 |
+
)
|
| 253 |
+
ops.fp8_causal_conv3d_ndhwc_bf16(
|
| 254 |
+
cache_x, new_x, weight, bias, float(alpha), out
|
| 255 |
+
)
|
| 256 |
+
return out
|
| 257 |
+
|
| 258 |
+
|
| 259 |
+
def fp8_conv2d_3x3_nhwc_bf16(
|
| 260 |
+
input: torch.Tensor,
|
| 261 |
+
weight: torch.Tensor,
|
| 262 |
+
bias: torch.Tensor,
|
| 263 |
+
alpha: float = 1.0,
|
| 264 |
+
*,
|
| 265 |
+
out: Optional[torch.Tensor] = None,
|
| 266 |
+
) -> torch.Tensor:
|
| 267 |
+
"""FP8 3x3 Conv2D with NHWC input/output and BF16 epilogue."""
|
| 268 |
+
|
| 269 |
+
if out is None:
|
| 270 |
+
out = torch.empty(
|
| 271 |
+
(*input.shape[:3], weight.shape[0]),
|
| 272 |
+
device=input.device,
|
| 273 |
+
dtype=torch.bfloat16,
|
| 274 |
+
)
|
| 275 |
+
ops.fp8_conv2d_3x3_nhwc_bf16(
|
| 276 |
+
input, weight, bias, float(alpha), out
|
| 277 |
+
)
|
| 278 |
+
return out
|
| 279 |
+
|
| 280 |
+
|
| 281 |
+
def fp8_conv2d_3x3_ncdhw_bf16(
|
| 282 |
+
input: torch.Tensor,
|
| 283 |
+
weight: torch.Tensor,
|
| 284 |
+
bias: torch.Tensor,
|
| 285 |
+
alpha: float = 1.0,
|
| 286 |
+
*,
|
| 287 |
+
out: Optional[torch.Tensor] = None,
|
| 288 |
+
) -> torch.Tensor:
|
| 289 |
+
"""FP8 3x3 Conv2D over B*T frames with direct BF16 NCDHW output."""
|
| 290 |
+
|
| 291 |
+
if out is None:
|
| 292 |
+
out = torch.empty(
|
| 293 |
+
(
|
| 294 |
+
input.shape[0],
|
| 295 |
+
weight.shape[0],
|
| 296 |
+
input.shape[1],
|
| 297 |
+
input.shape[2],
|
| 298 |
+
input.shape[3],
|
| 299 |
+
),
|
| 300 |
+
device=input.device,
|
| 301 |
+
dtype=torch.bfloat16,
|
| 302 |
+
)
|
| 303 |
+
ops.fp8_conv2d_3x3_ncdhw_bf16(
|
| 304 |
+
input, weight, bias, float(alpha), out
|
| 305 |
+
)
|
| 306 |
+
return out
|
| 307 |
+
|
| 308 |
+
|
| 309 |
+
def nvfp4_causal_conv3d_ndhwc_bf16(
|
| 310 |
+
cache_packed: torch.Tensor,
|
| 311 |
+
new_packed: torch.Tensor,
|
| 312 |
+
weight_packed: torch.Tensor,
|
| 313 |
+
cache_sf: torch.Tensor,
|
| 314 |
+
new_sf: torch.Tensor,
|
| 315 |
+
weight_sf: torch.Tensor,
|
| 316 |
+
bias: torch.Tensor,
|
| 317 |
+
outer_weight: torch.Tensor | None = None,
|
| 318 |
+
alpha: float = 1.0,
|
| 319 |
+
*,
|
| 320 |
+
out: Optional[torch.Tensor] = None,
|
| 321 |
+
) -> torch.Tensor:
|
| 322 |
+
"""NVFP4 causal Conv3D with linear UE4M3 scale factors and BF16 NDHWC output."""
|
| 323 |
+
|
| 324 |
+
n, t, h, w, _ = new_packed.shape
|
| 325 |
+
if out is None:
|
| 326 |
+
out = torch.empty(
|
| 327 |
+
(n, t, h, w, weight_packed.shape[0]),
|
| 328 |
+
device=new_packed.device,
|
| 329 |
+
dtype=torch.bfloat16,
|
| 330 |
+
)
|
| 331 |
+
ops.nvfp4_causal_conv3d_ndhwc_bf16(
|
| 332 |
+
cache_packed,
|
| 333 |
+
new_packed,
|
| 334 |
+
weight_packed,
|
| 335 |
+
cache_sf,
|
| 336 |
+
new_sf,
|
| 337 |
+
weight_sf,
|
| 338 |
+
bias,
|
| 339 |
+
outer_weight,
|
| 340 |
+
float(alpha),
|
| 341 |
+
out,
|
| 342 |
+
)
|
| 343 |
+
return out
|
| 344 |
+
|
| 345 |
+
|
| 346 |
+
def nvfp4_causal_conv3d_residual_ncdhw_bf16(
|
| 347 |
+
cache_packed: torch.Tensor,
|
| 348 |
+
new_packed: torch.Tensor,
|
| 349 |
+
weight_packed: torch.Tensor,
|
| 350 |
+
cache_sf: torch.Tensor,
|
| 351 |
+
new_sf: torch.Tensor,
|
| 352 |
+
weight_sf: torch.Tensor,
|
| 353 |
+
bias: torch.Tensor,
|
| 354 |
+
residual: torch.Tensor,
|
| 355 |
+
outer_weight: torch.Tensor | None = None,
|
| 356 |
+
alpha: float = 1.0,
|
| 357 |
+
*,
|
| 358 |
+
out: Optional[torch.Tensor] = None,
|
| 359 |
+
) -> torch.Tensor:
|
| 360 |
+
"""NVFP4 causal Conv3D with fused bias, residual, and BF16 NCDHW output."""
|
| 361 |
+
|
| 362 |
+
if out is None:
|
| 363 |
+
out = torch.empty_like(residual)
|
| 364 |
+
ops.nvfp4_causal_conv3d_residual_ncdhw_bf16(
|
| 365 |
+
cache_packed,
|
| 366 |
+
new_packed,
|
| 367 |
+
weight_packed,
|
| 368 |
+
cache_sf,
|
| 369 |
+
new_sf,
|
| 370 |
+
weight_sf,
|
| 371 |
+
bias,
|
| 372 |
+
residual,
|
| 373 |
+
outer_weight,
|
| 374 |
+
float(alpha),
|
| 375 |
+
out,
|
| 376 |
+
)
|
| 377 |
+
return out
|
| 378 |
+
|
| 379 |
+
|
| 380 |
+
__all__ = [
|
| 381 |
+
"fp8_conv3d_v18_ncdhw_res_bf16out",
|
| 382 |
+
"fp8_causal_conv3d_ndhwc_bf16",
|
| 383 |
+
"fp8_conv2d_3x3_nhwc_bf16",
|
| 384 |
+
"fp8_conv2d_3x3_ncdhw_bf16",
|
| 385 |
+
"nvfp4_causal_conv3d_ndhwc_bf16",
|
| 386 |
+
"nvfp4_causal_conv3d_residual_ncdhw_bf16",
|
| 387 |
+
]
|
build/torch212-cxx11-cu132-x86_64-linux/_ops.py
CHANGED
|
@@ -1,9 +1,9 @@
|
|
| 1 |
import torch
|
| 2 |
-
from . import
|
| 3 |
-
ops = torch.ops.
|
| 4 |
|
| 5 |
def add_op_namespace_prefix(op_name: str):
|
| 6 |
"""
|
| 7 |
Prefix op by namespace.
|
| 8 |
"""
|
| 9 |
-
return f"
|
|
|
|
| 1 |
import torch
|
| 2 |
+
from . import _world_model_conv_cuda_7781728
|
| 3 |
+
ops = torch.ops._world_model_conv_cuda_7781728
|
| 4 |
|
| 5 |
def add_op_namespace_prefix(op_name: str):
|
| 6 |
"""
|
| 7 |
Prefix op by namespace.
|
| 8 |
"""
|
| 9 |
+
return f"_world_model_conv_cuda_7781728::{op_name}"
|
build/torch212-cxx11-cu132-x86_64-linux/{_world_model_conv_cuda_f14c443.abi3.so → _world_model_conv_cuda_7781728.abi3.so}
RENAMED
|
@@ -1,3 +1,3 @@
|
|
| 1 |
version https://git-lfs.github.com/spec/v1
|
| 2 |
-
oid sha256:
|
| 3 |
-
size
|
|
|
|
| 1 |
version https://git-lfs.github.com/spec/v1
|
| 2 |
+
oid sha256:cf595248e19065ecdb2fe638be1a705647fad52912c88d0aa9c614195d09bf8f
|
| 3 |
+
size 1285680
|
build/torch212-cxx11-cu132-x86_64-linux/metadata.json
CHANGED
|
@@ -1,6 +1,6 @@
|
|
| 1 |
{
|
| 2 |
"name": "world-model-conv",
|
| 3 |
-
"id": "
|
| 4 |
"version": 1,
|
| 5 |
"license": "Apache-2.0",
|
| 6 |
"python-depends": [],
|
|
|
|
| 1 |
{
|
| 2 |
"name": "world-model-conv",
|
| 3 |
+
"id": "_world_model_conv_cuda_7781728",
|
| 4 |
"version": 1,
|
| 5 |
"license": "Apache-2.0",
|
| 6 |
"python-depends": [],
|