feat(VT-FP8-W8A8-CPU-ARM): register the static fp8 W8A8 path on the CPU backend (#468)

include/vt/fused_recipe.h

@@ -44,7 +44,9 @@ namespace vt {
 //      kSigmoidGate -> vt::SigmoidGateBf16  (a·sigmoid(b))
 //      kRmsNormGated-> vt::RmsNormGated     (gated rms-normalize)
 //      kRope        -> vt::RopeFromCache    (partial NeoX RoPE from a cos/sin cache)
-//      kQuantFp8    -> vt::QuantFp8Static   (static per-tensor fp8 terminal; CUDA-only)
+//      kQuantFp8    -> vt::QuantFp8Static   (static per-tensor fp8 terminal; CUDA + CPU
+//                                            since #468 — the fp8-terminal recipes now
+//                                            realize END-TO-END on CPU, not negotiated)
 //      kQuantFp4    -> vt::ScaledFp4Quant   (dynamic per-group fp4 terminal)
 //      kAttnQkNormRopeGate -> vt::AttnQkNormRopeGate (fused full-attention preamble)
 //

include/vt/ops.h

@@ -1453,10 +1453,19 @@ void MatmulNvfp4Cutlass(Queue& q, Tensor& out, const Tensor& a_packed, const Ten
 // per-tensor weight_scale + a f32 per-tensor input_scale (both applied directly:
 // dequant(w)=f8(w)*weight_scale, dequant(a)=f8(a)*input_scale).
 
-// QuantFp8Static (mirror vLLM static_scaled_fp8_quant, is_scale_inverted=False):
-//   out_fp8[i] = fp8_e4m3( clamp(x[i] / input_scale, -448, 448) )   // RNE hw cvt
+// QuantFp8Static (mirror vLLM static_scaled_fp8_quant):
+//   inv = 1/input_scale;  out_fp8[i] = fp8_e4m3( clamp(x[i] * inv, -448, 448) )
+// RNE convert. The scale is applied as a RECIPROCAL MULTIPLY, not a divide, and
+// the reciprocal is formed ONCE outside the elementwise loop — that is what
+// upstream ships: `x = val * scale` under is_scale_inverted=true
+// (csrc/quantization/w8a8/fp8/common.cuh:62, clamp at :68) with the inverse formed
+// by the caller (csrc/libtorch_stable/quantization/w8a8/fp8/common.cu:31,
+// `1.0f / scale[...]`). DO NOT "correct" the kernels to a divide to match a
+// prose formula: `x/s` and `x*(1/s)` differ by up to one f32 ulp, and near an
+// e4m3 tie that ulp changes the emitted byte on a default-ON 35B path.
 // Static per-tensor scale (NOT dynamic/per-token). x [M,K] f32/bf16, out [M,K]
-// i8 (raw fp8-e4m3fn bytes). CUDA only (the 35B W8A8 path is CUDA-resident).
+// i8 (raw fp8-e4m3fn bytes). CUDA + CPU (the CPU arm is the portable reference
+// that makes the fp8 seam testable without a GPU, #468).
 void QuantFp8Static(Queue& q, Tensor& out_fp8, const Tensor& x, float input_scale);
 
 // RmsNormQuantFp8 (fused fp8 RMSNorm -> static per-tensor activation quant). One
@@ -1507,7 +1516,16 @@ void RmsNormGatedQuantFp8(Queue& q, Tensor& out_fp8, const Tensor& x, const Tens
 // sequential scale_a·(scale_b·acc) — within fp8 tolerance, ported deviation).
 // a_fp8 [M,K] (= QuantFp8Static output), b_fp8 [N,K] the on-disk raw fp8-e4m3fn
 // weight (K contiguous). out [M,N] bf16 (cutlass epilogue) or f32 (via cast).
-// K,N multiples of 16 (128-bit fp8 alignment). CUDA-only (sm120a).
+// K,N multiples of 16 (128-bit fp8 alignment). CUDA (sm120a) + a CPU CORRECTNESS
+// REFERENCE (f32 accumulate, naive triple loop — no speed claim, and no
+// production model routes through it; it exists so the fp8 seam resolves on a
+// CPU queue and can be gated without a GPU, #468). The CPU arm is EXPECTED to
+// agree with the CUDA kernel to fp8/bf16 tolerance and NOT byte-for-byte, because
+// the CUDA arm reduces K in tensor-core order and rounds its epilogue through
+// bf16 — but that agreement is DECLARED AND OWED, not measured. No committed run
+// has compared the two arms: gate G2 of .agents/specs/vt-fp8-w8a8-cpu-arm.md is
+// PENDING for want of a GPU. Treat the tolerance above as the claim to be tested,
+// not as a result.
 void MatmulFp8Cutlass(Queue& q, Tensor& out, const Tensor& a_fp8, const Tensor& b_fp8,
                       float alpha);
 

src/vt/cpu/cpu_ops.cpp

@@ -532,6 +532,72 @@ uint8_t F32ToFp8(float f) {
                               static_cast<uint8_t>(mant));
 }
 
+// --- Static per-tensor FP8 W8A8 (VT-FP8-W8A8-CPU-ARM, #468). The CPU arm of the
+// path vLLM's ModelOptFp8LinearMethod runs: a static per-tensor activation quant
+// followed by a per-tensor fp8 GEMM. It exists so the fp8 seam is reachable, and
+// therefore testable, without a GPU.
+
+// QuantFp8Static CPU kernel — mirror of vLLM's static_scaled_fp8_quant
+// (csrc/quantization/w8a8/fp8/common.cuh:58-77 `scaled_fp8_conversion`):
+//   x = val * scale;  r = fmaxf(-448, fminf(x, 448));  hardware RNE convert
+// with the RECIPROCAL formed ONCE outside the loop, exactly as upstream forms it
+// (csrc/libtorch_stable/quantization/w8a8/fp8/common.cu:31 `1.0f / scale[...]`)
+// and exactly as our CUDA kernel does (cuda_matmul_fp8_cutlass.cu: `const float
+// inv = 1.0f / input_scale;` then `LoadIn(x, i) * inv`). It is a MULTIPLY BY THE
+// RECIPROCAL, not a divide: the two differ by up to one f32 ulp before the fp8
+// round, and near an e4m3 tie that ulp changes the emitted byte.
+//
+// F32ToFp8 supplies both remaining halves: it saturates (|a| >= 448 -> 0x7E,
+// which IS the encoding of 448, so clamp-then-convert and saturating-convert
+// coincide because 448 is the largest finite e4m3fn value) and it rounds to
+// nearest-even. The scale is per-TENSOR — upstream collapses the per-shard
+// input_scale to one scalar with `.max()` (modelopt.py:528) and then treats
+// `scale.numel() == 1` as a single group spanning the whole tensor
+// (common.cu:204-210). LoadF32 widens a bf16 x to f32 BEFORE the multiply, as
+// the CUDA kernel's LoadIn overload does, so both backends round at one point.
+void QuantFp8StaticKernel(Queue&, Tensor& out_fp8, const Tensor& x, float input_scale) {
+  const int64_t n = x.shape[0] * x.shape[1];
+  const float inv_scale = 1.0F / input_scale;
+  uint8_t* op = out_fp8.Ptr<uint8_t>();
+  ForRows(n, [&](int64_t r0, int64_t r1) {
+    for (int64_t i = r0; i < r1; ++i) op[i] = F32ToFp8(LoadF32(x, i) * inv_scale);
+  });
+}
+
+// MatmulFp8Cutlass CPU kernel: out[m,n] = alpha * Sum_k f8val(a[m,k])*f8val(b[n,k]),
+// f32 accumulate, ONE folded alpha (= input_scale*weight_scale — our recorded
+// deviation from upstream's two epilogue scalars, see include/vt/ops.h).
+//
+// A CORRECTNESS REFERENCE, NOT A PERFORMANCE PATH. It is a naive triple loop; it
+// makes no speed claim and nothing routes a production model through it. Its
+// purpose is that the fp8 GEMM seam resolves on a CPU queue so the surrounding
+// wiring can be gated without a GPU (#468).
+//
+// It is deliberately NOT a bit-mirror of the CUDA GEMM and does not claim to be:
+// the CUDA arm reduces K in tensor-core order and rounds its epilogue through
+// bf16, so the two agree to fp8/bf16 tolerance. Only the QUANT half above carries
+// a bit-exactness claim. Shaped like MatmulNvfp4Fp4Kernel: the A row is decoded
+// once per M and reused across N.
+void MatmulFp8CutlassKernel(Queue&, Tensor& out, const Tensor& a_fp8, const Tensor& b_fp8,
+                            float alpha) {
+  const int64_t m = a_fp8.shape[0], k = a_fp8.shape[1], n = b_fp8.shape[0];
+  const auto* ap = a_fp8.Ptr<uint8_t>();
+  const auto* bp = b_fp8.Ptr<uint8_t>();
+  ForRows(m, [&](int64_t r0, int64_t r1) {
+    std::vector<float> arow(static_cast<size_t>(k));
+    for (int64_t i = r0; i < r1; ++i) {
+      for (int64_t kk = 0; kk < k; ++kk)
+        arow[static_cast<size_t>(kk)] = Fp8ToF32(ap[i * k + kk]);
+      for (int64_t col = 0; col < n; ++col) {
+        float acc = 0.0F;
+        for (int64_t kk = 0; kk < k; ++kk)
+          acc += arow[static_cast<size_t>(kk)] * Fp8ToF32(bp[col * k + kk]);
+        StoreF32(out, i * n + col, alpha * acc);
+      }
+    }
+  });
+}
+
 // Fused fp8 RMSNorm -> static per-tensor quant (mirror vLLM Inductor
 // fused_add_rms_norm_static_fp8_quant, rms_quant_fusion.py:124). Same reduction
 // order as RmsNormKernel; the fp8 is taken from the SAME bf16-rounded normed value
@@ -3120,6 +3186,10 @@ struct Registrar {
                reinterpret_cast<void*>(static_cast<RmsNormFn>(&RmsNormKernel)));
     RegisterOp(OpId::kRmsNormQuantFp8, DeviceType::kCPU,
                reinterpret_cast<void*>(static_cast<RmsNormQuantFp8Fn>(&RmsNormQuantFp8Kernel)));
+    RegisterOp(OpId::kQuantFp8Static, DeviceType::kCPU,
+               reinterpret_cast<void*>(static_cast<QuantFp8StaticFn>(&QuantFp8StaticKernel)));
+    RegisterOp(OpId::kMatmulFp8Cutlass, DeviceType::kCPU,
+               reinterpret_cast<void*>(static_cast<MatmulFp8CutlassFn>(&MatmulFp8CutlassKernel)));
     RegisterOp(OpId::kSiluAndMul, DeviceType::kCPU,
                reinterpret_cast<void*>(static_cast<SiluAndMulFn>(&SiluAndMulKernel)));
     RegisterOp(OpId::kGeluAndMul, DeviceType::kCPU,

src/vt/cuda/cuda_matmul_fp8_cutlass.cu

@@ -317,8 +317,25 @@ void MatmulFp8CutlassKernelCuda(Queue& q, Tensor& out, const Tensor& a_fp8, cons
 }
 
 // ---- Static per-tensor fp8 activation quant (vLLM static_scaled_fp8_quant) ---
-// out_fp8[i] = fp8_e4m3(clamp(x[i]/input_scale, -448, 448)). __NV_SATFINITE cvt
-// saturates == clamp-then-cvt; RNE == vLLM's hardware cvt. Tin f32/bf16.
+// inv = 1/input_scale; out_fp8[i] = fp8_e4m3(clamp(x[i]*inv, -448, 448)).
+// A RECIPROCAL MULTIPLY, not a divide, and the reciprocal is hoisted out of the
+// loop — that is upstream's shipped form (`x = val * scale` with the inverse
+// formed by the caller: csrc/quantization/w8a8/fp8/common.cuh:62 and
+// csrc/libtorch_stable/quantization/w8a8/fp8/common.cu:31). The code below is
+// RIGHT; do not "fix" it into `x / input_scale` to match a prose formula. The two
+// differ by up to one f32 ulp before the fp8 round, and near an e4m3 tie that
+// ulp changes the emitted byte on a default-ON 35B path.
+// __NV_SATFINITE cvt saturates == clamp-then-cvt; RNE == vLLM's hardware cvt.
+// Tin f32/bf16.
+//
+// The CPU arm (src/vt/cpu/cpu_ops.cpp QuantFp8StaticKernel) is INTENDED to be the
+// byte-for-byte mirror of this kernel, and that equivalence is DECLARED AND OWED,
+// not measured. It is gate G2 of .agents/specs/vt-fp8-w8a8-cpu-arm.md, which is
+// PENDING for want of a GPU (#468). What IS measured is weaker and lives on the
+// CPU side: G1 proves the CPU kernel matches an independent e4m3 reference derived
+// from the format. Two implementations each matching a reference is not the same
+// claim as the two matching each other, so do not cite this comment as evidence
+// that they agree. Run tests/vt/test_ops_fp8_cpu.cpp on a CUDA host to close it.
 __device__ __forceinline__ uint8_t F32ToFp8Dev(float f) {
   return static_cast<uint8_t>(__nv_cvt_float_to_fp8(f, __NV_SATFINITE, __NV_E4M3));
 }

tests/CMakeLists.txt

@@ -1537,6 +1537,10 @@ if(VLLM_CPP_CUTLASS)
       "VT_FP4_TEST_PERSISTENT_SAVE=1;VT_FP4_PERSISTENT_CACHE=1;VT_FP4_AUTOTUNE_CACHE_READONLY=0;VT_FP4_AUTOTUNE_CACHE_PATH=${CMAKE_CURRENT_BINARY_DIR}/nvfp4-runtime-save-lifecycle.json;VT_FP4_AUTOTUNE_DELAY_US=0;VT_FP4_PLAN_CACHE=1;VT_FP4_AUTOTUNE=1;VT_FP4_FULL_TACTICS=1")
 endif()
 vllm_cpp_add_test(test_ops_fp8_cutlass vt/test_ops_fp8_cutlass.cpp)
+# VT-FP8-W8A8-CPU-ARM (#468): the CPU arm of the static fp8 W8A8 path. Runs on a
+# box with NO GPU by construction — that is the point of the row — and carries a
+# CUDA-gated arm (G2) for CPU-vs-CUDA byte agreement wherever a device exists.
+vllm_cpp_add_test(test_ops_fp8_cpu vt/test_ops_fp8_cpu.cpp)
 # Opt-in arm: run the fp8 plan-cache byte-exact case with the cache ENABLED
 # (VT_FP8_PLAN_CACHE=1 -> first MatmulFp8CublasLt call builds the plan fresh,
 # later calls hit the cache). Proves the cached-plan GEMM is BYTE-identical to the