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.

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

@@ -169,6 +169,8 @@ def select_fp8_linear_tile(m: int, n: int, k: int, variant: int = 0) -> str:
         raise RuntimeError("small-M dispatcher currently supports variant=0 only")
     if k % 32:
         raise RuntimeError("SM120 requires k divisible by 32")
+    if m > 64:
+        return "cublaslt_fp8_large_m"
     if m <= 16:
         if k % 256 == 0:
             return "ld_fp8_gemm_16x128x256_w4" if n % 128 == 0 else "ld_fp8_gemm_16x64x256_w4"
@@ -193,7 +195,7 @@ def select_fp8_linear_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("only M=1 decode or 2 <= M <= 64 small-M rows are supported")
+    raise RuntimeError("M must be positive")
 
 
 def fp8_linear_bf16(
@@ -209,7 +211,7 @@ def fp8_linear_bf16(
     ``(M, K)`` and ``(N, K)``. ``alpha`` is a host float, normally the product
     of static per-tensor input and weight scales. SM110 uses the production
     CUTLASS Sq/T1/Wide dispatcher over full model row counts; SM120 uses the
-    hand-tuned M<=64 path.
+    hand-tuned M<=64 path and cuBLASLt for larger row counts.
     """
 
     if out is None:

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

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

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

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

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

@@ -1,6 +1,6 @@
 {
   "name": "fp8-gemm",
-  "id": "_fp8_gemm_cuda_8fba1d7",
+  "id": "_fp8_gemm_cuda_9ac1ace",
   "version": 1,
   "license": "Apache-2.0",
   "python-depends": [],
@@ -15,9 +15,9 @@
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "__init__.py": "HsFfsGY7AK/6Fziv4r2VQ+R2kDjS6hxKKgWqlA+zrKY=",
-      "_fp8_gemm_cuda_8fba1d7.abi3.so": "Q80ftellpduL5nWRUWDisZJ4F4m9gIxhyhUhwiWYyHQ=",
-      "_ops.py": "ptfSaRMzscUfi27OPI9uPdaK9uAgUUrImeBJSSDv3kY="
+      "__init__.py": "7UODiuG0STVFP2F4ZTuu+2z8Ih/u8572gTNX+uLBP/s=",
+      "_fp8_gemm_cuda_9ac1ace.abi3.so": "oxpa8yOZfsax7sENyF0apU3A/dMsTkGqP7JiwmCcvZM=",
+      "_ops.py": "8623JQCjzAtJO6ZqPnLQaFzDwRCUC0pkQWHUlwQbtAc="
     }
   },
   "provenance": {
@@ -27,7 +27,7 @@
       "dirty": false
     },
     "kernel": {
-      "sha": "8fba1d7f2e4e627b1a2c3c8018f01c0ca4f430ee",
+      "sha": "9ac1ace56531e1dee1125f5dd2f1cafc77b890fa",
       "dirty": false
     }
   }

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

@@ -169,6 +169,8 @@ def select_fp8_linear_tile(m: int, n: int, k: int, variant: int = 0) -> str:
         raise RuntimeError("small-M dispatcher currently supports variant=0 only")
     if k % 32:
         raise RuntimeError("SM120 requires k divisible by 32")
+    if m > 64:
+        return "cublaslt_fp8_large_m"
     if m <= 16:
         if k % 256 == 0:
             return "ld_fp8_gemm_16x128x256_w4" if n % 128 == 0 else "ld_fp8_gemm_16x64x256_w4"
@@ -193,7 +195,7 @@ def select_fp8_linear_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("only M=1 decode or 2 <= M <= 64 small-M rows are supported")
+    raise RuntimeError("M must be positive")
 
 
 def fp8_linear_bf16(
@@ -209,7 +211,7 @@ def fp8_linear_bf16(
     ``(M, K)`` and ``(N, K)``. ``alpha`` is a host float, normally the product
     of static per-tensor input and weight scales. SM110 uses the production
     CUTLASS Sq/T1/Wide dispatcher over full model row counts; SM120 uses the
-    hand-tuned M<=64 path.
+    hand-tuned M<=64 path and cuBLASLt for larger row counts.
     """
 
     if out is None:

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

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

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

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

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

@@ -1,6 +1,6 @@
 {
   "name": "fp8-gemm",
-  "id": "_fp8_gemm_cuda_8fba1d7",
+  "id": "_fp8_gemm_cuda_9ac1ace",
   "version": 1,
   "license": "Apache-2.0",
   "python-depends": [],
@@ -15,9 +15,9 @@
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "__init__.py": "HsFfsGY7AK/6Fziv4r2VQ+R2kDjS6hxKKgWqlA+zrKY=",
-      "_fp8_gemm_cuda_8fba1d7.abi3.so": "bzopByL+RJnJVmRH1mJ3YSe+6PtIhekscncdHEn1I2Q=",
-      "_ops.py": "ptfSaRMzscUfi27OPI9uPdaK9uAgUUrImeBJSSDv3kY="
+      "__init__.py": "7UODiuG0STVFP2F4ZTuu+2z8Ih/u8572gTNX+uLBP/s=",
+      "_fp8_gemm_cuda_9ac1ace.abi3.so": "MRxLWP8Hf/aHMMhpDRpEz6v0gVyfJ32jhZ2CJu7Y9Hs=",
+      "_ops.py": "8623JQCjzAtJO6ZqPnLQaFzDwRCUC0pkQWHUlwQbtAc="
     }
   },
   "provenance": {
@@ -27,7 +27,7 @@
       "dirty": false
     },
     "kernel": {
-      "sha": "8fba1d7f2e4e627b1a2c3c8018f01c0ca4f430ee",
+      "sha": "9ac1ace56531e1dee1125f5dd2f1cafc77b890fa",
       "dirty": false
     }
   }

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

@@ -169,6 +169,8 @@ def select_fp8_linear_tile(m: int, n: int, k: int, variant: int = 0) -> str:
         raise RuntimeError("small-M dispatcher currently supports variant=0 only")
     if k % 32:
         raise RuntimeError("SM120 requires k divisible by 32")
+    if m > 64:
+        return "cublaslt_fp8_large_m"
     if m <= 16:
         if k % 256 == 0:
             return "ld_fp8_gemm_16x128x256_w4" if n % 128 == 0 else "ld_fp8_gemm_16x64x256_w4"
@@ -193,7 +195,7 @@ def select_fp8_linear_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("only M=1 decode or 2 <= M <= 64 small-M rows are supported")
+    raise RuntimeError("M must be positive")
 
 
 def fp8_linear_bf16(
@@ -209,7 +211,7 @@ def fp8_linear_bf16(
     ``(M, K)`` and ``(N, K)``. ``alpha`` is a host float, normally the product
     of static per-tensor input and weight scales. SM110 uses the production
     CUTLASS Sq/T1/Wide dispatcher over full model row counts; SM120 uses the
-    hand-tuned M<=64 path.
+    hand-tuned M<=64 path and cuBLASLt for larger row counts.
     """
 
     if out is None:

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

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

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

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

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

@@ -1,6 +1,6 @@
 {
   "name": "fp8-gemm",
-  "id": "_fp8_gemm_cuda_8fba1d7",
+  "id": "_fp8_gemm_cuda_9ac1ace",
   "version": 1,
   "license": "Apache-2.0",
   "python-depends": [],
@@ -15,9 +15,9 @@
   "digest": {
     "algorithm": "sha256",
     "files": {
-      "__init__.py": "HsFfsGY7AK/6Fziv4r2VQ+R2kDjS6hxKKgWqlA+zrKY=",
-      "_fp8_gemm_cuda_8fba1d7.abi3.so": "NX7cwUxquYmQiiEhm2QReTRreX616p07MCrvezvuOYg=",
-      "_ops.py": "ptfSaRMzscUfi27OPI9uPdaK9uAgUUrImeBJSSDv3kY="
+      "__init__.py": "7UODiuG0STVFP2F4ZTuu+2z8Ih/u8572gTNX+uLBP/s=",
+      "_fp8_gemm_cuda_9ac1ace.abi3.so": "GcubdTSg4wXYG7K5yOFUAHKcjC4+Lt5IGmo2kv+7o7s=",
+      "_ops.py": "8623JQCjzAtJO6ZqPnLQaFzDwRCUC0pkQWHUlwQbtAc="
     }
   },
   "provenance": {
@@ -27,7 +27,7 @@
       "dirty": false
     },
     "kernel": {
-      "sha": "8fba1d7f2e4e627b1a2c3c8018f01c0ca4f430ee",
+      "sha": "9ac1ace56531e1dee1125f5dd2f1cafc77b890fa",
       "dirty": false
     }
   }