File size: 3,764 Bytes
9022070 | 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31 32 33 34 35 36 37 38 39 40 41 42 43 44 45 46 47 48 49 50 51 52 53 54 55 56 57 58 59 60 61 62 63 64 65 66 67 68 69 70 71 72 73 74 75 76 77 78 79 80 81 | # Source Sync
- Upstream FlashRT source: `../official/FlashRT`
- Initial package date: June 20, 2026
- SM89 source revision: `70b8eac4b05e9193bd99631cf872c5a971b59f5d`
- SM110 sync revision: `132049d7c3a3534fb7d35676cd726f39408b1af6`
Copied source files:
- `csrc/gemm/fp8_gemv_m1_sm120.cu`
- `csrc/gemm/fp8_gemv_m1_sm120.cuh`
- `csrc/gemm/fp8_smallM_handtuned_sm120.cu`
- `csrc/gemm/fp8_smallM_handtuned_sm120.cuh`
- `csrc/gemm/fp8_smallM_handtuned_ldmatrix_sm120.cu`
- `csrc/gemm/fp8_smallM_handtuned_ldmatrix_sm120.cuh`
- `csrc/gemm/cutlass_sm120_block128_fp8_gemm.cu`
- `csrc/gemm/cutlass_sm120_block128_fp8_gemm.cuh`
- `csrc/gemm/fp8_block128_gemm_mma_sm89.cu`
- `csrc/gemm/fp8_block128_gemm_mma_sm89.cuh`
- `csrc/gemm/fp8_bs_gemm_device.cuh`
- `csrc/gemm/fp8_gemv_m1_sm89.cu`
- `csrc/gemm/fp8_gemv_m1_sm89.cuh`
- `csrc/gemm/gemm_types_sm100.h`
- `csrc/gemm/cutlass_sm100.cu`
The SM110 copies are package-local as `csrc/gemm_types_sm110.h` and
`csrc/cutlass_sm110_fp8_gemm.cu`. The C declarations in
`csrc/cutlass_sm110_fp8_gemm.cuh` are packaging glue; the upstream pointer API
declares them in its aggregate binding instead.
Local packaging edits:
- Added Tensor-facing PyTorch custom ops in `torch-ext/torch_binding.cpp`.
- Added Python wrappers and fake registrations in `torch-ext/fp8_gemm`.
- Kept public APIs model-agnostic; no raw pointer or stream arguments.
- Bound the upstream measured `32x128-w4-s1` fused SwiGLU producer without
changing its CUDA tile or arithmetic.
- Added a Tensor-facing SM110 dispatcher over the upstream BF16-output Sq, T1,
and Wide tactics. The public dispatcher and diagnostic variants do not alter
the copied GEMM templates or arithmetic.
- Renamed SM100 source filenames locally to make their SM110 package role
explicit; CUTLASS still uses the SM100-family architecture templates when
compiling for `sm_110a`.
- The SM110 build uses `-O3`, `--expt-relaxed-constexpr`, and
`--use_fast_math`, matching the validated native path.
Architecture assumptions:
- CUDA 12.8+ for SM89/SM120; CUDA 13.0+ for SM110.
- NVIDIA Ada SM89 for block-128 scaled GEMM/GEMV.
- NVIDIA Blackwell SM110a for per-tensor Sq/T1/Wide FP8 GEMM with BF16 output.
- NVIDIA Blackwell SM120a for the original public APIs. The per-tensor FP8 MMA path uses
`.kind::f8f6f4` instructions and must be compiled for `sm_120a`, not plain
`sm_120`.
- The SM110 kernel depends on the builder-provided CUTLASS 4.5 package. The
package flake is pinned to a builder revision that exports `cutlass_4_5`.
- The release flake temporarily pins
`LiangSu8899/kernels@d720fa9`, based on
`huggingface/kernels@e9152aa`. The fork preserves the upstream builder
sources and changes only the stale CUTLASS 4.5.2 fixed-output hash from the
specified value to the value returned by the upstream archive. Return to an
upstream revision after that hash correction lands.
Runtime constraints:
- Inputs are FP8 E4M3 tensors with layout `input[M, K]` and `weight[N, K]`.
- Output is BF16 `out[M, N]`.
- `K` must be divisible by 32.
- On SM120, `M` must be `1` or in `2..64`. M=128 remains an internal tuning
item because the validated correct SM120 tile is not performance-positive
enough for public release.
- On SM110, `N` and `K` must be divisible by 16. The full-row Sq/T1/Wide path
has been validated on `M` from 1 through 1024 across PI0.5, GROOT,
Cosmos Edge, and LingBot projection families.
- `alpha` is a host float scale multiplier, normally
`input_scale * weight_scale`.
- The blockwise path consumes FP32 scales with layouts `(M, K/128)` and
`(N/128, K/128)`. It is the same CUTLASS kernel and schedule dispatcher used
by the upstream FlashRT pointer API.
- Blockwise scaling is not exposed on SM110 in this increment.
|