pub fn cublas_fp8_proj(
gpu: &dyn GpuBackend,
ptg_quant_k: KernelHandle,
act_bf16: DevicePtr,
act_fp8_scratch: DevicePtr,
act_scale_scratch: DevicePtr,
fp8w: &Fp8Weight,
out: DevicePtr,
m: u32,
n: u32,
k: u32,
stream: u64,
) -> Result<()>Expand description
Route a projection through native-FP8 cuBLASLt block-scaled matmul: quantize
the activation to FP8 + per-[token,128-of-K] VEC128 scales (the existing
per_token_group_quant_fp8 kernel), feed the FP8 weight + its per-128×128
block scales directly (zero dequant, zero extra weight memory). Both operands
128-block-scaled (cuBLASLt requires it). ~1.8× the bf16 path (152 vs 85 TF).
act_fp8_scratch/act_scale_scratch must hold the padded extents (the
buffers.fp8_act/fp8_act_scale arena buffers, sized for max_batch_tokens).