Initialize legacy kernels compatibility mirror

README.md

@@ -1,135 +1,9 @@
-# fp8-gemm
+# flashrt/fp8-gemm
 
-FlashRT native CUDA FP8 GEMV/GEMM kernels for low-latency transformer and
-diffuser linear layers on NVIDIA Ada SM89 and Blackwell SM110/SM120 GPUs.
+This repository is a compatibility mirror for older `kernels` clients
+that resolve repositories through the default Hugging Face model repo API.
 
-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.
+Canonical Kernel Hub repo: https://huggingface.co/kernels/flashrt/fp8-gemm
 
-## 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.
+Do not edit this mirror by hand. It is generated from the Kernel Hub
+`vN` branches and contains the same `build/**` artifacts.