Promote latest kernel artifacts to main

README.md

@@ -1,9 +1,135 @@
-# flashrt/fp8-gemm
+# fp8-gemm
 
-This repository is a compatibility mirror for older `kernels` clients
-that resolve repositories through the default Hugging Face model repo API.
+FlashRT native CUDA FP8 GEMV/GEMM kernels for low-latency transformer and
+diffuser linear layers on NVIDIA Ada SM89 and Blackwell SM110/SM120 GPUs.
 
-Canonical Kernel Hub repo: https://huggingface.co/kernels/flashrt/fp8-gemm
+This package exposes the hand-tuned FP8 E4M3 decode and small-M kernels as
+Tensor APIs for Hugging Face Kernel Hub. It is intended for model runtimes that
+already hold activations and weights in FP8 and want a low-overhead BF16 output
+linear path.
 
-Do not edit this mirror by hand. It is generated from the Kernel Hub
-`vN` branches and contains the same `build/**` artifacts.
+## Available Functions
+
+- `fp8_linear_bf16(input, weight, alpha=1.0, out=None, variant=0)`
+- `fp8_linear_residual_bf16(input, weight, residual, alpha=1.0, variant=0)`
+- `fp8_linear_bias_bf16(input, weight, bias, alpha=1.0, out=None)`
+- `fp8_linear_bias_residual_bf16(input, weight, bias, residual, alpha=1.0)`
+- `fp8_linear_bias_gelu_bf16(input, weight, bias, alpha=1.0, out=None)`
+- `fp8_blockwise_linear_bf16(input, weight, input_scale, weight_scale, out=None)`
+- `fp8_blockwise_swiglu_quantize_fp8(input, gate_up_weight, input_scale, gate_up_weight_scale, output=None, output_scale=None)`
+- `select_fp8_linear_tile(m, n, k, variant=0)`
+
+Tensor contract:
+
+- `input`: `torch.float8_e4m3fn`, shape `(M, K)`, contiguous CUDA tensor.
+- `weight`: `torch.float8_e4m3fn`, shape `(N, K)`, contiguous CUDA tensor.
+- `out`: `torch.bfloat16`, shape `(M, N)`.
+- `residual`: `torch.bfloat16`, shape `(1, N)` or `(N,)`, only supported for
+  the `M=1` decode GEMV path.
+- `K % 16 == 0`; SM120 additionally requires `K % 32 == 0`.
+- On SM120, `M == 1` uses dedicated GEMV and `2 <= M <= 64` uses small-M
+  GEMM tiles.
+- On SM110 (Jetson AGX Thor), the per-tensor API uses the production FlashRT
+  CUTLASS Sq/T1/Wide family and supports the validated model-shape matrix from
+  decode through large vision/backbone rows. The large-M production band is
+  validated from `M=65` through `M=1024`, including PI0.5 prefill QKV, O,
+  gate/up, and down projections at `M=712..970`. `N` and `K` must be divisible
+  by 16.
+- The three BF16 bias APIs are SM110-only. They accept BF16 `(N,)` bias and
+  preserve the same row-major FP8 `(M,K)` input and `(N,K)` weight contract.
+  The residual API updates a BF16 `(M,N)` tensor in place. The GELU API uses
+  the tanh approximation.
+- SM110 `variant=0` is the production auto dispatcher. Diagnostic variants are
+  `1=Sq`, `2=T1`, and `3=Wide`; they are correctness-tested but should not be
+  pinned by model integrations without a shape-specific benchmark.
+- The per-tensor kernels use Blackwell FP8 MMA instructions and are not valid
+  for SM89. SM89 support is provided by the blockwise API below.
+- `alpha` is a host float. For per-tensor FP8 quantization, pass
+  `float(input_scale * weight_scale)` from your static calibration metadata.
+
+The blockwise API uses a separate contract:
+
+- `input`: FP8 E4M3 `(M, K)`.
+- `weight`: FP8 E4M3 `(N, K)`.
+- `input_scale`: FP32 `(M, K / 128)`.
+- `weight_scale`: FP32 `(N / 128, K / 128)`.
+- `N` and `K` must be divisible by 128; `M` is unrestricted.
+- Output is BF16 `(M, N)`.
+- On SM89, the blockwise API dispatches to the production FlashRT native
+  `mma.sync.aligned.m16n8k32` GEMM/GEMV implementation.
+- On SM120, it dispatches to the production FlashRT CUTLASS block-scaled
+  implementation.
+- SM110 is intentionally not claimed by the blockwise API; use the per-tensor
+  static-scale path there. Other architectures are rejected explicitly.
+
+The fused SM89 producer accepts FP8 `(M,K)` input, FP8 `(2*N,K)` gate/up
+weight, block-128 FP32 scales, and returns FP8 `(M,N)` plus FP32 `(M,N/128)`
+output scales. Its public range is `1 <= M <= 256` with `N` and `K` divisible
+by 128. It is rejected explicitly on non-SM89 GPUs.
+
+## Minimal Usage
+
+```python
+from kernels import get_kernel
+import torch
+
+ops = get_kernel("flashrt/fp8-gemm", version=1, trust_remote_code=True)
+
+x = torch.randn((16, 4096), device="cuda", dtype=torch.bfloat16).to(torch.float8_e4m3fn)
+w = torch.randn((8192, 4096), device="cuda", dtype=torch.bfloat16).to(torch.float8_e4m3fn)
+
+y = ops.fp8_linear_bf16(x, w, alpha=1.0)
+```
+
+SM110 bias epilogues:
+
+```python
+bias = torch.randn((8192,), device="cuda", dtype=torch.bfloat16)
+residual = torch.randn((16, 8192), device="cuda", dtype=torch.bfloat16)
+
+y = ops.fp8_linear_bias_bf16(x, w, bias, alpha=1.0)
+ops.fp8_linear_bias_residual_bf16(x, w, bias, residual, alpha=1.0)
+y_gelu = ops.fp8_linear_bias_gelu_bf16(x, w, bias, alpha=1.0)
+```
+
+Warm each distinct SM110 bias shape once before CUDA Graph capture. The
+cuBLASLt fallback lazily creates and caches its descriptor, algorithm, and
+workspace on the first call; replay itself performs no allocation.
+
+Decode residual path:
+
+```python
+x = torch.randn((1, 4096), device="cuda", dtype=torch.bfloat16).to(torch.float8_e4m3fn)
+w = torch.randn((4096, 4096), device="cuda", dtype=torch.bfloat16).to(torch.float8_e4m3fn)
+residual = torch.zeros((1, 4096), device="cuda", dtype=torch.bfloat16)
+
+ops.fp8_linear_residual_bf16(x, w, residual, alpha=1.0)
+```
+
+Block-128 scaling:
+
+```python
+m, k, n = 51, 1536, 1536
+x = torch.randn((m, k), device="cuda").to(torch.float8_e4m3fn)
+w = torch.randn((n, k), device="cuda").to(torch.float8_e4m3fn)
+x_scale = torch.ones((m, k // 128), device="cuda", dtype=torch.float32)
+w_scale = torch.ones((n // 128, k // 128), device="cuda", dtype=torch.float32)
+
+y = ops.fp8_blockwise_linear_bf16(x, w, x_scale, w_scale)
+```
+
+## Validation
+
+```bash
+python fp8-gemm/tests/test_fp8_gemm.py --backend source --mode full
+python fp8-gemm/benchmarks/benchmark.py --backend source --mode headline
+python fp8-gemm/benchmarks/benchmark.py --backend source --mode pi05-prefill
+python fp8-gemm/benchmarks/benchmark_bias.py --backend source
+```
+
+The SM110 full sweep covers PI0.5, GROOT N1.6/N1.7, Cosmos Edge, and LingBot
+VLA projection families, plus decode, generic small-M, the `M=65` large-M
+boundary, and the three SigLIP bias epilogues. Public
+benchmark tables are only updated after source correctness, installed artifact
+correctness, shape/tile sweeps, `torch.compile(fullgraph=True)`, CUDA Graph
+replay, and parity against the original FlashRT native pointer entry pass.

benchmarks/benchmark.py

@@ -149,13 +149,20 @@ def load_source_ops() -> SourceOps:
             str(cutlass_include.parent / "tools" / "util" / "include"),
         ]
     else:
+        if not (cutlass_include / "cutlass" / "cutlass.h").is_file():
+            raise RuntimeError("set CUTLASS_INCLUDE for the SM120 source benchmark")
         cuda_sources = [
             str(PACKAGE / "csrc" / "fp8_gemv_m1_sm120.cu"),
             str(PACKAGE / "csrc" / "fp8_smallM_handtuned_sm120.cu"),
             str(PACKAGE / "csrc" / "fp8_smallM_handtuned_ldmatrix_sm120.cu"),
+            str(PACKAGE / "csrc" / "cutlass_sm120_block128_fp8_gemm.cu"),
+            str(PACKAGE / "csrc" / "cublaslt_fp8_bias_sm110.cu"),
         ]
         source_define = "-DFLASHRT_FP8_GEMM_SOURCE_SM120_ONLY"
-        extra_includes = []
+        extra_includes = [
+            str(cutlass_include),
+            str(cutlass_include.parent / "tools" / "util" / "include"),
+        ]
     load(
         name=namespace,
         sources=[str(PACKAGE / "torch-ext" / "torch_binding.cpp"), *cuda_sources],
@@ -250,7 +257,7 @@ def select_tile(m: int, n: int, k: int, variant: int = 0) -> str:
         if n % 128 == 0:
             return "ld_fp8_gemm_64x128x128_w4"
         return "ld_fp8_gemm_64x64x128_w4"
-    raise RuntimeError("unsupported M")
+    return "cublaslt_fp8_large_m"
 
 
 def make_inputs(m: int, k: int, n: int, seed: int):

benchmarks/benchmark_bias.py

@@ -1,5 +1,5 @@
 #!/usr/bin/env python3
-"""Benchmark SM110 BF16-output FP8 GEMM epilogues."""
+"""Benchmark SM110/SM120 BF16-output FP8 GEMM epilogues."""
 
 from __future__ import annotations
 
@@ -66,29 +66,59 @@ class SourceOps:
 def load_source_ops() -> SourceOps:
     from torch.utils.cpp_extension import load
 
-    cutlass = Path(os.environ["CUTLASS_INCLUDE"])
-    os.environ.setdefault("TORCH_CUDA_ARCH_LIST", "11.0a")
+    capability = torch.cuda.get_device_capability()
+    if capability == (11, 0):
+        cutlass = Path(os.environ["CUTLASS_INCLUDE"])
+        os.environ.setdefault("TORCH_CUDA_ARCH_LIST", "11.0a")
+        cuda_sources = [
+            str(PACKAGE / "csrc" / "cutlass_sm110_fp8_gemm.cu"),
+            str(PACKAGE / "csrc" / "cublaslt_fp8_bias_sm110.cu"),
+        ]
+        source_define = "-DFLASHRT_FP8_GEMM_SOURCE_SM110_ONLY"
+        extra_includes = [
+            str(cutlass),
+            str(cutlass.parent / "tools" / "util" / "include"),
+        ]
+    elif capability[0] >= 12:
+        cutlass = Path(os.environ["CUTLASS_INCLUDE"])
+        os.environ.setdefault("TORCH_CUDA_ARCH_LIST", "12.0a")
+        cuda_sources = [
+            str(PACKAGE / "csrc" / "fp8_gemv_m1_sm120.cu"),
+            str(PACKAGE / "csrc" / "fp8_smallM_handtuned_sm120.cu"),
+            str(PACKAGE / "csrc" / "fp8_smallM_handtuned_ldmatrix_sm120.cu"),
+            str(PACKAGE / "csrc" / "cutlass_sm120_block128_fp8_gemm.cu"),
+            str(PACKAGE / "csrc" / "cublaslt_fp8_bias_sm110.cu"),
+        ]
+        source_define = "-DFLASHRT_FP8_GEMM_SOURCE_SM120_ONLY"
+        extra_includes = [
+            str(cutlass),
+            str(cutlass.parent / "tools" / "util" / "include"),
+        ]
+    else:
+        raise RuntimeError(f"SM110 or SM120 is required, got {capability}")
     namespace = "fp8_gemm_bias_source_bench"
     load(
         name=namespace,
         sources=[
             str(PACKAGE / "torch-ext" / "torch_binding.cpp"),
-            str(PACKAGE / "csrc" / "cutlass_sm110_fp8_gemm.cu"),
-            str(PACKAGE / "csrc" / "cublaslt_fp8_bias_sm110.cu"),
+            *cuda_sources,
         ],
         extra_include_paths=[
             str(PACKAGE / "csrc"),
             str(REGISTRATION_INCLUDE),
-            str(cutlass),
-            str(cutlass.parent / "tools" / "util" / "include"),
+            *extra_includes,
         ],
         extra_cflags=[
             "-O3", "-DNDEBUG", "-DCUDA_KERNEL",
-            "-DFLASHRT_FP8_GEMM_SOURCE_SM110_ONLY",
+            source_define,
         ],
         extra_cuda_cflags=[
             "-O3", "-DNDEBUG", "--expt-relaxed-constexpr", "--use_fast_math",
-            "-DCUDA_KERNEL", "-DFLASHRT_FP8_GEMM_SOURCE_SM110_ONLY",
+            "-U__CUDA_NO_HALF_OPERATORS__",
+            "-U__CUDA_NO_HALF_CONVERSIONS__",
+            "-U__CUDA_NO_BFLOAT16_CONVERSIONS__",
+            "-U__CUDA_NO_HALF2_OPERATORS__",
+            "-DCUDA_KERNEL", source_define,
         ],
         verbose=False,
     )
@@ -154,8 +184,9 @@ def main() -> int:
     parser.add_argument("--iterations", type=int, default=64)
     parser.add_argument("--json-out", default=None)
     args = parser.parse_args()
-    if torch.cuda.get_device_capability() != (11, 0):
-        raise SystemExit("SM110 is required")
+    capability = torch.cuda.get_device_capability()
+    if capability != (11, 0) and capability[0] < 12:
+        raise SystemExit(f"SM110 or SM120 is required, got {capability}")
     ops = load_source_ops() if args.backend == "source" else load_installed_ops(args.artifact)
     native = load_native()
     native_stream = int(torch.cuda.current_stream().cuda_stream)

build/torch212-cxx11-cu130-x86_64-linux/__init__.py

@@ -37,6 +37,34 @@ def _fp8_linear_residual_bf16_fake(
     return None
 
 
+def _check_bias_linear_shapes(input, weight, bias, out) -> None:
+    if input.dim() != 2 or weight.dim() != 2:
+        raise RuntimeError("input and weight must be rank-2 tensors")
+    if input.shape[1] != weight.shape[1]:
+        raise RuntimeError("input and weight K dimensions must match")
+    if bias.shape != (weight.shape[0],):
+        raise RuntimeError("bias must have shape (weight.shape[0],)")
+    if out.shape != (input.shape[0], weight.shape[0]):
+        raise RuntimeError("out must have shape (input.shape[0], weight.shape[0])")
+
+
+@torch.library.register_fake(add_op_namespace_prefix("fp8_linear_bias_bf16"))
+def _fp8_linear_bias_bf16_fake(input, weight, bias, alpha: float, out) -> None:
+    _check_bias_linear_shapes(input, weight, bias, out)
+
+
+@torch.library.register_fake(add_op_namespace_prefix("fp8_linear_bias_residual_bf16"))
+def _fp8_linear_bias_residual_bf16_fake(
+    input, weight, bias, alpha: float, residual
+) -> None:
+    _check_bias_linear_shapes(input, weight, bias, residual)
+
+
+@torch.library.register_fake(add_op_namespace_prefix("fp8_linear_bias_gelu_bf16"))
+def _fp8_linear_bias_gelu_bf16_fake(input, weight, bias, alpha: float, out) -> None:
+    _check_bias_linear_shapes(input, weight, bias, out)
+
+
 @torch.library.register_fake(add_op_namespace_prefix("fp8_blockwise_linear_bf16"))
 def _fp8_blockwise_linear_bf16_fake(
     input: torch.Tensor,
@@ -95,8 +123,8 @@ def select_fp8_linear_tile(m: int, n: int, k: int, variant: int = 0) -> str:
     variant = int(variant)
     if m <= 0 or n <= 0 or k <= 0:
         raise RuntimeError("m, n, and k must be positive")
-    if k % 32 != 0:
-        raise RuntimeError("k must be divisible by 32")
+    if k % 16 != 0:
+        raise RuntimeError("k must be divisible by 16")
     capability = torch.cuda.get_device_capability() if torch.cuda.is_available() else None
     if capability == (11, 0):
         forced = {1: "sm110_sq_bf16", 2: "sm110_t1_bf16", 3: "sm110_wide_bf16"}
@@ -106,6 +134,12 @@ def select_fp8_linear_tile(m: int, n: int, k: int, variant: int = 0) -> str:
             raise RuntimeError("SM110 requires n and k divisible by 16")
         if variant:
             return forced[variant]
+        if m >= 512 and k == 2048 and 2048 <= n <= 2560:
+            return "sm110_sq_bf16"
+        if m >= 512 and n >= 16 * k:
+            return "sm110_t1_bf16"
+        if m >= 512 and k >= 4 * n:
+            return "sm110_wide_bf16"
         if n >= 8 * k:
             return "sm110_wide_bf16"
         if m >= 128 and k >= 4 * n:
@@ -116,6 +150,8 @@ def select_fp8_linear_tile(m: int, n: int, k: int, variant: int = 0) -> str:
             return "sm110_wide_bf16"
         return "sm110_t1_bf16"
     if m == 1:
+        if k % 32:
+            raise RuntimeError("SM120 requires k divisible by 32")
         if variant == 4:
             return "gemv_fp8_m1_w4"
         if variant == 8:
@@ -131,6 +167,8 @@ def select_fp8_linear_tile(m: int, n: int, k: int, variant: int = 0) -> str:
         return "gemv_fp8_m1_w16"
     if variant != 0:
         raise RuntimeError("small-M dispatcher currently supports variant=0 only")
+    if k % 32:
+        raise RuntimeError("SM120 requires k divisible by 32")
     if m <= 16:
         if k % 256 == 0:
             return "ld_fp8_gemm_16x128x256_w4" if n % 128 == 0 else "ld_fp8_gemm_16x64x256_w4"
@@ -197,6 +235,56 @@ def fp8_linear_residual_bf16(
     return residual
 
 
+def fp8_linear_bias_bf16(
+    input: torch.Tensor,
+    weight: torch.Tensor,
+    bias: torch.Tensor,
+    alpha: float = 1.0,
+    out: torch.Tensor | None = None,
+) -> torch.Tensor:
+    """SM110 FP8 linear with fused BF16 bias and BF16 output."""
+    if out is None:
+        out = torch.empty(
+            (input.shape[0], weight.shape[0]),
+            device=input.device,
+            dtype=torch.bfloat16,
+        )
+    ops.fp8_linear_bias_bf16(input, weight, bias, float(alpha), out)
+    return out
+
+
+def fp8_linear_bias_residual_bf16(
+    input: torch.Tensor,
+    weight: torch.Tensor,
+    bias: torch.Tensor,
+    residual: torch.Tensor,
+    alpha: float = 1.0,
+) -> torch.Tensor:
+    """SM110 fused ``residual += alpha * input @ weight.T + bias``."""
+    ops.fp8_linear_bias_residual_bf16(
+        input, weight, bias, float(alpha), residual
+    )
+    return residual
+
+
+def fp8_linear_bias_gelu_bf16(
+    input: torch.Tensor,
+    weight: torch.Tensor,
+    bias: torch.Tensor,
+    alpha: float = 1.0,
+    out: torch.Tensor | None = None,
+) -> torch.Tensor:
+    """SM110 FP8 linear with fused BF16 bias and GELU epilogue."""
+    if out is None:
+        out = torch.empty(
+            (input.shape[0], weight.shape[0]),
+            device=input.device,
+            dtype=torch.bfloat16,
+        )
+    ops.fp8_linear_bias_gelu_bf16(input, weight, bias, float(alpha), out)
+    return out
+
+
 def fp8_blockwise_linear_bf16(
     input: torch.Tensor,
     weight: torch.Tensor,
@@ -248,6 +336,9 @@ def fp8_blockwise_swiglu_quantize_fp8(
 __all__ = [
     "fp8_linear_bf16",
     "fp8_linear_residual_bf16",
+    "fp8_linear_bias_bf16",
+    "fp8_linear_bias_residual_bf16",
+    "fp8_linear_bias_gelu_bf16",
     "fp8_blockwise_linear_bf16",
     "fp8_blockwise_swiglu_quantize_fp8",
     "select_fp8_linear_tile",

build/torch212-cxx11-cu130-x86_64-linux/{_fp8_gemm_cuda_bf15968.abi3.so → _fp8_gemm_cuda_8fba1d7.abi3.so}

@@ -1,3 +1,3 @@
 version https://git-lfs.github.com/spec/v1
-oid sha256:d8c90f60cb56efb79b25a6db2abe5f58336c7327902a0bcddfb7e4d3bde41dbb
-size 5390512
+oid sha256:3341b44bcf4bcbc602b01c1e0f409d63b8a3b4b0ef97ff5a9bccbc1b97bd74da
+size 5665360

build/torch212-cxx11-cu130-x86_64-linux/_ops.py

@@ -1,9 +1,9 @@
 import torch
-from . import _fp8_gemm_cuda_bf15968
-ops = torch.ops._fp8_gemm_cuda_bf15968
+from . import _fp8_gemm_cuda_8fba1d7
+ops = torch.ops._fp8_gemm_cuda_8fba1d7
 
 def add_op_namespace_prefix(op_name: str):
     """
     Prefix op by namespace.
     """
-    return f"_fp8_gemm_cuda_bf15968::{op_name}"
+    return f"_fp8_gemm_cuda_8fba1d7::{op_name}"

build/torch212-cxx11-cu130-x86_64-linux/metadata.json

@@ -1,6 +1,6 @@
 {
   "name": "fp8-gemm",
-  "id": "_fp8_gemm_cuda_bf15968",
+  "id": "_fp8_gemm_cuda_8fba1d7",
   "version": 1,
   "license": "Apache-2.0",
   "python-depends": [],
@@ -15,9 +15,9 @@
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "__init__.py": "hKsdFER2ZmYz36/zz9LYE7ddsK/R1dprHJFx+9yBzks=",
-      "_fp8_gemm_cuda_bf15968.abi3.so": "2MkPYMtW77ebJabbKr5fWDNscyeQKgvN37fk073kHbs=",
-      "_ops.py": "wDFSPD8vh2K2qjj/9mdUYv6LrXuGz6MQCb2B/e8nqWs="
+      "__init__.py": "HsFfsGY7AK/6Fziv4r2VQ+R2kDjS6hxKKgWqlA+zrKY=",
+      "_fp8_gemm_cuda_8fba1d7.abi3.so": "M0G0S89Ly8YCsBweD0CdY7ijtLDvl/9am8y8G5e9dNo=",
+      "_ops.py": "ptfSaRMzscUfi27OPI9uPdaK9uAgUUrImeBJSSDv3kY="
     }
   },
   "provenance": {
@@ -27,7 +27,7 @@
       "dirty": false
     },
     "kernel": {
-      "sha": "bf15968fd9c281a7ca107e7696300fe091f588a2",
+      "sha": "8fba1d7f2e4e627b1a2c3c8018f01c0ca4f430ee",
       "dirty": false
     }
   }

build/torch212-cxx11-cu132-x86_64-linux/__init__.py

@@ -37,6 +37,34 @@ def _fp8_linear_residual_bf16_fake(
     return None
 
 
+def _check_bias_linear_shapes(input, weight, bias, out) -> None:
+    if input.dim() != 2 or weight.dim() != 2:
+        raise RuntimeError("input and weight must be rank-2 tensors")
+    if input.shape[1] != weight.shape[1]:
+        raise RuntimeError("input and weight K dimensions must match")
+    if bias.shape != (weight.shape[0],):
+        raise RuntimeError("bias must have shape (weight.shape[0],)")
+    if out.shape != (input.shape[0], weight.shape[0]):
+        raise RuntimeError("out must have shape (input.shape[0], weight.shape[0])")
+
+
+@torch.library.register_fake(add_op_namespace_prefix("fp8_linear_bias_bf16"))
+def _fp8_linear_bias_bf16_fake(input, weight, bias, alpha: float, out) -> None:
+    _check_bias_linear_shapes(input, weight, bias, out)
+
+
+@torch.library.register_fake(add_op_namespace_prefix("fp8_linear_bias_residual_bf16"))
+def _fp8_linear_bias_residual_bf16_fake(
+    input, weight, bias, alpha: float, residual
+) -> None:
+    _check_bias_linear_shapes(input, weight, bias, residual)
+
+
+@torch.library.register_fake(add_op_namespace_prefix("fp8_linear_bias_gelu_bf16"))
+def _fp8_linear_bias_gelu_bf16_fake(input, weight, bias, alpha: float, out) -> None:
+    _check_bias_linear_shapes(input, weight, bias, out)
+
+
 @torch.library.register_fake(add_op_namespace_prefix("fp8_blockwise_linear_bf16"))
 def _fp8_blockwise_linear_bf16_fake(
     input: torch.Tensor,
@@ -95,8 +123,8 @@ def select_fp8_linear_tile(m: int, n: int, k: int, variant: int = 0) -> str:
     variant = int(variant)
     if m <= 0 or n <= 0 or k <= 0:
         raise RuntimeError("m, n, and k must be positive")
-    if k % 32 != 0:
-        raise RuntimeError("k must be divisible by 32")
+    if k % 16 != 0:
+        raise RuntimeError("k must be divisible by 16")
     capability = torch.cuda.get_device_capability() if torch.cuda.is_available() else None
     if capability == (11, 0):
         forced = {1: "sm110_sq_bf16", 2: "sm110_t1_bf16", 3: "sm110_wide_bf16"}
@@ -106,6 +134,12 @@ def select_fp8_linear_tile(m: int, n: int, k: int, variant: int = 0) -> str:
             raise RuntimeError("SM110 requires n and k divisible by 16")
         if variant:
             return forced[variant]
+        if m >= 512 and k == 2048 and 2048 <= n <= 2560:
+            return "sm110_sq_bf16"
+        if m >= 512 and n >= 16 * k:
+            return "sm110_t1_bf16"
+        if m >= 512 and k >= 4 * n:
+            return "sm110_wide_bf16"
         if n >= 8 * k:
             return "sm110_wide_bf16"
         if m >= 128 and k >= 4 * n:
@@ -116,6 +150,8 @@ def select_fp8_linear_tile(m: int, n: int, k: int, variant: int = 0) -> str:
             return "sm110_wide_bf16"
         return "sm110_t1_bf16"
     if m == 1:
+        if k % 32:
+            raise RuntimeError("SM120 requires k divisible by 32")
         if variant == 4:
             return "gemv_fp8_m1_w4"
         if variant == 8:
@@ -131,6 +167,8 @@ def select_fp8_linear_tile(m: int, n: int, k: int, variant: int = 0) -> str:
         return "gemv_fp8_m1_w16"
     if variant != 0:
         raise RuntimeError("small-M dispatcher currently supports variant=0 only")
+    if k % 32:
+        raise RuntimeError("SM120 requires k divisible by 32")
     if m <= 16:
         if k % 256 == 0:
             return "ld_fp8_gemm_16x128x256_w4" if n % 128 == 0 else "ld_fp8_gemm_16x64x256_w4"
@@ -197,6 +235,56 @@ def fp8_linear_residual_bf16(
     return residual
 
 
+def fp8_linear_bias_bf16(
+    input: torch.Tensor,
+    weight: torch.Tensor,
+    bias: torch.Tensor,
+    alpha: float = 1.0,
+    out: torch.Tensor | None = None,
+) -> torch.Tensor:
+    """SM110 FP8 linear with fused BF16 bias and BF16 output."""
+    if out is None:
+        out = torch.empty(
+            (input.shape[0], weight.shape[0]),
+            device=input.device,
+            dtype=torch.bfloat16,
+        )
+    ops.fp8_linear_bias_bf16(input, weight, bias, float(alpha), out)
+    return out
+
+
+def fp8_linear_bias_residual_bf16(
+    input: torch.Tensor,
+    weight: torch.Tensor,
+    bias: torch.Tensor,
+    residual: torch.Tensor,
+    alpha: float = 1.0,
+) -> torch.Tensor:
+    """SM110 fused ``residual += alpha * input @ weight.T + bias``."""
+    ops.fp8_linear_bias_residual_bf16(
+        input, weight, bias, float(alpha), residual
+    )
+    return residual
+
+
+def fp8_linear_bias_gelu_bf16(
+    input: torch.Tensor,
+    weight: torch.Tensor,
+    bias: torch.Tensor,
+    alpha: float = 1.0,
+    out: torch.Tensor | None = None,
+) -> torch.Tensor:
+    """SM110 FP8 linear with fused BF16 bias and GELU epilogue."""
+    if out is None:
+        out = torch.empty(
+            (input.shape[0], weight.shape[0]),
+            device=input.device,
+            dtype=torch.bfloat16,
+        )
+    ops.fp8_linear_bias_gelu_bf16(input, weight, bias, float(alpha), out)
+    return out
+
+
 def fp8_blockwise_linear_bf16(
     input: torch.Tensor,
     weight: torch.Tensor,
@@ -248,6 +336,9 @@ def fp8_blockwise_swiglu_quantize_fp8(
 __all__ = [
     "fp8_linear_bf16",
     "fp8_linear_residual_bf16",
+    "fp8_linear_bias_bf16",
+    "fp8_linear_bias_residual_bf16",
+    "fp8_linear_bias_gelu_bf16",
     "fp8_blockwise_linear_bf16",
     "fp8_blockwise_swiglu_quantize_fp8",
     "select_fp8_linear_tile",

build/torch212-cxx11-cu132-x86_64-linux/{_fp8_gemm_cuda_bf15968.abi3.so → _fp8_gemm_cuda_8fba1d7.abi3.so}

@@ -1,3 +1,3 @@
 version https://git-lfs.github.com/spec/v1
-oid sha256:9b025dcd2da72fc44d87e0bb8a68168e56e2972bca5417631ff1f534615c31ae
-size 5431528
+oid sha256:43cd1fb5e965a5db8be675915160e2b192781789bd808c61ca1521c22598c874
+size 5710472

build/torch212-cxx11-cu132-x86_64-linux/_ops.py

@@ -1,9 +1,9 @@
 import torch
-from . import _fp8_gemm_cuda_bf15968
-ops = torch.ops._fp8_gemm_cuda_bf15968
+from . import _fp8_gemm_cuda_8fba1d7
+ops = torch.ops._fp8_gemm_cuda_8fba1d7
 
 def add_op_namespace_prefix(op_name: str):
     """
     Prefix op by namespace.
     """
-    return f"_fp8_gemm_cuda_bf15968::{op_name}"
+    return f"_fp8_gemm_cuda_8fba1d7::{op_name}"

build/torch212-cxx11-cu132-x86_64-linux/metadata.json

@@ -1,6 +1,6 @@
 {
   "name": "fp8-gemm",
-  "id": "_fp8_gemm_cuda_bf15968",
+  "id": "_fp8_gemm_cuda_8fba1d7",
   "version": 1,
   "license": "Apache-2.0",
   "python-depends": [],
@@ -15,9 +15,9 @@
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "__init__.py": "hKsdFER2ZmYz36/zz9LYE7ddsK/R1dprHJFx+9yBzks=",
-      "_fp8_gemm_cuda_bf15968.abi3.so": "mwJdzS2nL8RNh+C7imgWjlbilyvKVBdjH/H1NGFcMa4=",
-      "_ops.py": "wDFSPD8vh2K2qjj/9mdUYv6LrXuGz6MQCb2B/e8nqWs="
+      "__init__.py": "HsFfsGY7AK/6Fziv4r2VQ+R2kDjS6hxKKgWqlA+zrKY=",
+      "_fp8_gemm_cuda_8fba1d7.abi3.so": "Q80ftellpduL5nWRUWDisZJ4F4m9gIxhyhUhwiWYyHQ=",
+      "_ops.py": "ptfSaRMzscUfi27OPI9uPdaK9uAgUUrImeBJSSDv3kY="
     }
   },
   "provenance": {
@@ -27,7 +27,7 @@
       "dirty": false
     },
     "kernel": {
-      "sha": "bf15968fd9c281a7ca107e7696300fe091f588a2",
+      "sha": "8fba1d7f2e4e627b1a2c3c8018f01c0ca4f430ee",
       "dirty": false
     }
   }

build/torch213-cxx11-cu130-x86_64-linux/__init__.py

@@ -37,6 +37,34 @@ def _fp8_linear_residual_bf16_fake(
     return None
 
 
+def _check_bias_linear_shapes(input, weight, bias, out) -> None:
+    if input.dim() != 2 or weight.dim() != 2:
+        raise RuntimeError("input and weight must be rank-2 tensors")
+    if input.shape[1] != weight.shape[1]:
+        raise RuntimeError("input and weight K dimensions must match")
+    if bias.shape != (weight.shape[0],):
+        raise RuntimeError("bias must have shape (weight.shape[0],)")
+    if out.shape != (input.shape[0], weight.shape[0]):
+        raise RuntimeError("out must have shape (input.shape[0], weight.shape[0])")
+
+
+@torch.library.register_fake(add_op_namespace_prefix("fp8_linear_bias_bf16"))
+def _fp8_linear_bias_bf16_fake(input, weight, bias, alpha: float, out) -> None:
+    _check_bias_linear_shapes(input, weight, bias, out)
+
+
+@torch.library.register_fake(add_op_namespace_prefix("fp8_linear_bias_residual_bf16"))
+def _fp8_linear_bias_residual_bf16_fake(
+    input, weight, bias, alpha: float, residual
+) -> None:
+    _check_bias_linear_shapes(input, weight, bias, residual)
+
+
+@torch.library.register_fake(add_op_namespace_prefix("fp8_linear_bias_gelu_bf16"))
+def _fp8_linear_bias_gelu_bf16_fake(input, weight, bias, alpha: float, out) -> None:
+    _check_bias_linear_shapes(input, weight, bias, out)
+
+
 @torch.library.register_fake(add_op_namespace_prefix("fp8_blockwise_linear_bf16"))
 def _fp8_blockwise_linear_bf16_fake(
     input: torch.Tensor,
@@ -95,8 +123,8 @@ def select_fp8_linear_tile(m: int, n: int, k: int, variant: int = 0) -> str:
     variant = int(variant)
     if m <= 0 or n <= 0 or k <= 0:
         raise RuntimeError("m, n, and k must be positive")
-    if k % 32 != 0:
-        raise RuntimeError("k must be divisible by 32")
+    if k % 16 != 0:
+        raise RuntimeError("k must be divisible by 16")
     capability = torch.cuda.get_device_capability() if torch.cuda.is_available() else None
     if capability == (11, 0):
         forced = {1: "sm110_sq_bf16", 2: "sm110_t1_bf16", 3: "sm110_wide_bf16"}
@@ -106,6 +134,12 @@ def select_fp8_linear_tile(m: int, n: int, k: int, variant: int = 0) -> str:
             raise RuntimeError("SM110 requires n and k divisible by 16")
         if variant:
             return forced[variant]
+        if m >= 512 and k == 2048 and 2048 <= n <= 2560:
+            return "sm110_sq_bf16"
+        if m >= 512 and n >= 16 * k:
+            return "sm110_t1_bf16"
+        if m >= 512 and k >= 4 * n:
+            return "sm110_wide_bf16"
         if n >= 8 * k:
             return "sm110_wide_bf16"
         if m >= 128 and k >= 4 * n:
@@ -116,6 +150,8 @@ def select_fp8_linear_tile(m: int, n: int, k: int, variant: int = 0) -> str:
             return "sm110_wide_bf16"
         return "sm110_t1_bf16"
     if m == 1:
+        if k % 32:
+            raise RuntimeError("SM120 requires k divisible by 32")
         if variant == 4:
             return "gemv_fp8_m1_w4"
         if variant == 8:
@@ -131,6 +167,8 @@ def select_fp8_linear_tile(m: int, n: int, k: int, variant: int = 0) -> str:
         return "gemv_fp8_m1_w16"
     if variant != 0:
         raise RuntimeError("small-M dispatcher currently supports variant=0 only")
+    if k % 32:
+        raise RuntimeError("SM120 requires k divisible by 32")
     if m <= 16:
         if k % 256 == 0:
             return "ld_fp8_gemm_16x128x256_w4" if n % 128 == 0 else "ld_fp8_gemm_16x64x256_w4"
@@ -197,6 +235,56 @@ def fp8_linear_residual_bf16(
     return residual
 
 
+def fp8_linear_bias_bf16(
+    input: torch.Tensor,
+    weight: torch.Tensor,
+    bias: torch.Tensor,
+    alpha: float = 1.0,
+    out: torch.Tensor | None = None,
+) -> torch.Tensor:
+    """SM110 FP8 linear with fused BF16 bias and BF16 output."""
+    if out is None:
+        out = torch.empty(
+            (input.shape[0], weight.shape[0]),
+            device=input.device,
+            dtype=torch.bfloat16,
+        )
+    ops.fp8_linear_bias_bf16(input, weight, bias, float(alpha), out)
+    return out
+
+
+def fp8_linear_bias_residual_bf16(
+    input: torch.Tensor,
+    weight: torch.Tensor,
+    bias: torch.Tensor,
+    residual: torch.Tensor,
+    alpha: float = 1.0,
+) -> torch.Tensor:
+    """SM110 fused ``residual += alpha * input @ weight.T + bias``."""
+    ops.fp8_linear_bias_residual_bf16(
+        input, weight, bias, float(alpha), residual
+    )
+    return residual
+
+
+def fp8_linear_bias_gelu_bf16(
+    input: torch.Tensor,
+    weight: torch.Tensor,
+    bias: torch.Tensor,
+    alpha: float = 1.0,
+    out: torch.Tensor | None = None,
+) -> torch.Tensor:
+    """SM110 FP8 linear with fused BF16 bias and GELU epilogue."""
+    if out is None:
+        out = torch.empty(
+            (input.shape[0], weight.shape[0]),
+            device=input.device,
+            dtype=torch.bfloat16,
+        )
+    ops.fp8_linear_bias_gelu_bf16(input, weight, bias, float(alpha), out)
+    return out
+
+
 def fp8_blockwise_linear_bf16(
     input: torch.Tensor,
     weight: torch.Tensor,
@@ -248,6 +336,9 @@ def fp8_blockwise_swiglu_quantize_fp8(
 __all__ = [
     "fp8_linear_bf16",
     "fp8_linear_residual_bf16",
+    "fp8_linear_bias_bf16",
+    "fp8_linear_bias_residual_bf16",
+    "fp8_linear_bias_gelu_bf16",
     "fp8_blockwise_linear_bf16",
     "fp8_blockwise_swiglu_quantize_fp8",
     "select_fp8_linear_tile",

build/torch213-cxx11-cu130-x86_64-linux/{_fp8_gemm_cuda_bf15968.abi3.so → _fp8_gemm_cuda_8fba1d7.abi3.so}

@@ -1,3 +1,3 @@
 version https://git-lfs.github.com/spec/v1
-oid sha256:90adef3d22bb98b7dc1b2b48d860f0fa2fe682125acb31b1c554f889e0f171ce
-size 5390352
+oid sha256:6f3a290722fe4499c9566447d662776127bee8fb4885e92c72771d1c49f52364
+size 5665200

build/torch213-cxx11-cu130-x86_64-linux/_ops.py

@@ -1,9 +1,9 @@
 import torch
-from . import _fp8_gemm_cuda_bf15968
-ops = torch.ops._fp8_gemm_cuda_bf15968
+from . import _fp8_gemm_cuda_8fba1d7
+ops = torch.ops._fp8_gemm_cuda_8fba1d7
 
 def add_op_namespace_prefix(op_name: str):
     """
     Prefix op by namespace.
     """
-    return f"_fp8_gemm_cuda_bf15968::{op_name}"
+    return f"_fp8_gemm_cuda_8fba1d7::{op_name}"

build/torch213-cxx11-cu130-x86_64-linux/metadata.json

@@ -1,6 +1,6 @@
 {
   "name": "fp8-gemm",
-  "id": "_fp8_gemm_cuda_bf15968",
+  "id": "_fp8_gemm_cuda_8fba1d7",
   "version": 1,
   "license": "Apache-2.0",
   "python-depends": [],
@@ -15,9 +15,9 @@
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "__init__.py": "hKsdFER2ZmYz36/zz9LYE7ddsK/R1dprHJFx+9yBzks=",
-      "_fp8_gemm_cuda_bf15968.abi3.so": "kK3vPSK7mLfcGytI2GDw+i/mghJayzGxxVT4ieDxcc4=",
-      "_ops.py": "wDFSPD8vh2K2qjj/9mdUYv6LrXuGz6MQCb2B/e8nqWs="
+      "__init__.py": "HsFfsGY7AK/6Fziv4r2VQ+R2kDjS6hxKKgWqlA+zrKY=",
+      "_fp8_gemm_cuda_8fba1d7.abi3.so": "bzopByL+RJnJVmRH1mJ3YSe+6PtIhekscncdHEn1I2Q=",
+      "_ops.py": "ptfSaRMzscUfi27OPI9uPdaK9uAgUUrImeBJSSDv3kY="
     }
   },
   "provenance": {
@@ -27,7 +27,7 @@
       "dirty": false
     },
     "kernel": {
-      "sha": "bf15968fd9c281a7ca107e7696300fe091f588a2",
+      "sha": "8fba1d7f2e4e627b1a2c3c8018f01c0ca4f430ee",
       "dirty": false
     }
   }

build/torch213-cxx11-cu132-x86_64-linux/__init__.py

@@ -37,6 +37,34 @@ def _fp8_linear_residual_bf16_fake(
     return None
 
 
+def _check_bias_linear_shapes(input, weight, bias, out) -> None:
+    if input.dim() != 2 or weight.dim() != 2:
+        raise RuntimeError("input and weight must be rank-2 tensors")
+    if input.shape[1] != weight.shape[1]:
+        raise RuntimeError("input and weight K dimensions must match")
+    if bias.shape != (weight.shape[0],):
+        raise RuntimeError("bias must have shape (weight.shape[0],)")
+    if out.shape != (input.shape[0], weight.shape[0]):
+        raise RuntimeError("out must have shape (input.shape[0], weight.shape[0])")
+
+
+@torch.library.register_fake(add_op_namespace_prefix("fp8_linear_bias_bf16"))
+def _fp8_linear_bias_bf16_fake(input, weight, bias, alpha: float, out) -> None:
+    _check_bias_linear_shapes(input, weight, bias, out)
+
+
+@torch.library.register_fake(add_op_namespace_prefix("fp8_linear_bias_residual_bf16"))
+def _fp8_linear_bias_residual_bf16_fake(
+    input, weight, bias, alpha: float, residual
+) -> None:
+    _check_bias_linear_shapes(input, weight, bias, residual)
+
+
+@torch.library.register_fake(add_op_namespace_prefix("fp8_linear_bias_gelu_bf16"))
+def _fp8_linear_bias_gelu_bf16_fake(input, weight, bias, alpha: float, out) -> None:
+    _check_bias_linear_shapes(input, weight, bias, out)
+
+
 @torch.library.register_fake(add_op_namespace_prefix("fp8_blockwise_linear_bf16"))
 def _fp8_blockwise_linear_bf16_fake(
     input: torch.Tensor,
@@ -95,8 +123,8 @@ def select_fp8_linear_tile(m: int, n: int, k: int, variant: int = 0) -> str:
     variant = int(variant)
     if m <= 0 or n <= 0 or k <= 0:
         raise RuntimeError("m, n, and k must be positive")
-    if k % 32 != 0:
-        raise RuntimeError("k must be divisible by 32")
+    if k % 16 != 0:
+        raise RuntimeError("k must be divisible by 16")
     capability = torch.cuda.get_device_capability() if torch.cuda.is_available() else None
     if capability == (11, 0):
         forced = {1: "sm110_sq_bf16", 2: "sm110_t1_bf16", 3: "sm110_wide_bf16"}
@@ -106,6 +134,12 @@ def select_fp8_linear_tile(m: int, n: int, k: int, variant: int = 0) -> str:
             raise RuntimeError("SM110 requires n and k divisible by 16")
         if variant:
             return forced[variant]
+        if m >= 512 and k == 2048 and 2048 <= n <= 2560:
+            return "sm110_sq_bf16"
+        if m >= 512 and n >= 16 * k:
+            return "sm110_t1_bf16"
+        if m >= 512 and k >= 4 * n:
+            return "sm110_wide_bf16"
         if n >= 8 * k:
             return "sm110_wide_bf16"
         if m >= 128 and k >= 4 * n:
@@ -116,6 +150,8 @@ def select_fp8_linear_tile(m: int, n: int, k: int, variant: int = 0) -> str:
             return "sm110_wide_bf16"
         return "sm110_t1_bf16"
     if m == 1:
+        if k % 32:
+            raise RuntimeError("SM120 requires k divisible by 32")
         if variant == 4:
             return "gemv_fp8_m1_w4"
         if variant == 8:
@@ -131,6 +167,8 @@ def select_fp8_linear_tile(m: int, n: int, k: int, variant: int = 0) -> str:
         return "gemv_fp8_m1_w16"
     if variant != 0:
         raise RuntimeError("small-M dispatcher currently supports variant=0 only")
+    if k % 32:
+        raise RuntimeError("SM120 requires k divisible by 32")
     if m <= 16:
         if k % 256 == 0:
             return "ld_fp8_gemm_16x128x256_w4" if n % 128 == 0 else "ld_fp8_gemm_16x64x256_w4"
@@ -197,6 +235,56 @@ def fp8_linear_residual_bf16(
     return residual
 
 
+def fp8_linear_bias_bf16(
+    input: torch.Tensor,
+    weight: torch.Tensor,
+    bias: torch.Tensor,
+    alpha: float = 1.0,
+    out: torch.Tensor | None = None,
+) -> torch.Tensor:
+    """SM110 FP8 linear with fused BF16 bias and BF16 output."""
+    if out is None:
+        out = torch.empty(
+            (input.shape[0], weight.shape[0]),
+            device=input.device,
+            dtype=torch.bfloat16,
+        )
+    ops.fp8_linear_bias_bf16(input, weight, bias, float(alpha), out)
+    return out
+
+
+def fp8_linear_bias_residual_bf16(
+    input: torch.Tensor,
+    weight: torch.Tensor,
+    bias: torch.Tensor,
+    residual: torch.Tensor,
+    alpha: float = 1.0,
+) -> torch.Tensor:
+    """SM110 fused ``residual += alpha * input @ weight.T + bias``."""
+    ops.fp8_linear_bias_residual_bf16(
+        input, weight, bias, float(alpha), residual
+    )
+    return residual
+
+
+def fp8_linear_bias_gelu_bf16(
+    input: torch.Tensor,
+    weight: torch.Tensor,
+    bias: torch.Tensor,
+    alpha: float = 1.0,
+    out: torch.Tensor | None = None,
+) -> torch.Tensor:
+    """SM110 FP8 linear with fused BF16 bias and GELU epilogue."""
+    if out is None:
+        out = torch.empty(
+            (input.shape[0], weight.shape[0]),
+            device=input.device,
+            dtype=torch.bfloat16,
+        )
+    ops.fp8_linear_bias_gelu_bf16(input, weight, bias, float(alpha), out)
+    return out
+
+
 def fp8_blockwise_linear_bf16(
     input: torch.Tensor,
     weight: torch.Tensor,
@@ -248,6 +336,9 @@ def fp8_blockwise_swiglu_quantize_fp8(
 __all__ = [
     "fp8_linear_bf16",
     "fp8_linear_residual_bf16",
+    "fp8_linear_bias_bf16",
+    "fp8_linear_bias_residual_bf16",
+    "fp8_linear_bias_gelu_bf16",
     "fp8_blockwise_linear_bf16",
     "fp8_blockwise_swiglu_quantize_fp8",
     "select_fp8_linear_tile",

build/torch213-cxx11-cu132-x86_64-linux/{_fp8_gemm_cuda_bf15968.abi3.so → _fp8_gemm_cuda_8fba1d7.abi3.so}

@@ -1,3 +1,3 @@
 version https://git-lfs.github.com/spec/v1
-oid sha256:e5395ce05668ccf4bec5f12109f073fb6ce743810f486ea98d67de77574a17b7
-size 5431368
+oid sha256:357edcc14c6ab989908a21219b641179346b797eb5ea9d3b302aef7b3bee3988
+size 5710312

build/torch213-cxx11-cu132-x86_64-linux/_ops.py

@@ -1,9 +1,9 @@
 import torch
-from . import _fp8_gemm_cuda_bf15968
-ops = torch.ops._fp8_gemm_cuda_bf15968
+from . import _fp8_gemm_cuda_8fba1d7
+ops = torch.ops._fp8_gemm_cuda_8fba1d7
 
 def add_op_namespace_prefix(op_name: str):
     """
     Prefix op by namespace.
     """
-    return f"_fp8_gemm_cuda_bf15968::{op_name}"
+    return f"_fp8_gemm_cuda_8fba1d7::{op_name}"

build/torch213-cxx11-cu132-x86_64-linux/metadata.json

@@ -1,6 +1,6 @@
 {
   "name": "fp8-gemm",
-  "id": "_fp8_gemm_cuda_bf15968",
+  "id": "_fp8_gemm_cuda_8fba1d7",
   "version": 1,
   "license": "Apache-2.0",
   "python-depends": [],
@@ -15,9 +15,9 @@
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "__init__.py": "hKsdFER2ZmYz36/zz9LYE7ddsK/R1dprHJFx+9yBzks=",
-      "_fp8_gemm_cuda_bf15968.abi3.so": "5Tlc4FZozPS+xfEhCfBz+2znQ4EPSG6pjWfed1dKF7c=",
-      "_ops.py": "wDFSPD8vh2K2qjj/9mdUYv6LrXuGz6MQCb2B/e8nqWs="
+      "__init__.py": "HsFfsGY7AK/6Fziv4r2VQ+R2kDjS6hxKKgWqlA+zrKY=",
+      "_fp8_gemm_cuda_8fba1d7.abi3.so": "NX7cwUxquYmQiiEhm2QReTRreX616p07MCrvezvuOYg=",
+      "_ops.py": "ptfSaRMzscUfi27OPI9uPdaK9uAgUUrImeBJSSDv3kY="
     }
   },
   "provenance": {
@@ -27,7 +27,7 @@
       "dirty": false
     },
     "kernel": {
-      "sha": "bf15968fd9c281a7ca107e7696300fe091f588a2",
+      "sha": "8fba1d7f2e4e627b1a2c3c8018f01c0ca4f430ee",
       "dirty": false
     }
   }