pub fn cublas_fp8_proj(
gpu: &dyn GpuBackend,
ptg_quant_k: Fp8ActQuant,
scale_kmajor_k: KernelHandle,
act_bf16: DevicePtr,
act_fp8_scratch: DevicePtr,
act_scale_scratch: DevicePtr,
act_scale_kmajor_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] scales (the existing
per_token_group_quant_fp8 kernel), adapt those scales to the layout
cuBLASLt documents, then 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/act_scale_kmajor_scratch must hold
the padded extents (the buffers.fp8_act/fp8_act_scale arena buffers,
sized for max_batch_tokens).