pub fn cublas_fp8_rowwise_proj(
gpu: &dyn GpuBackend,
derived: &DerivedWeights,
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 ROW-WISE native-FP8 cuBLASLt (the fp8 path GB10
supports). Weight is re-quantized once to per-row fp8 (cached); the activation
is quantized per-token each call. ~1.8× the bf16 path (152 vs 85 TF), and
frees the bf16-dequant memory the bf16 path holds.
act_fp8_scratch ≥ m*k fp8 bytes; act_scale_scratch ≥ m f32 (e.g. the
buffers.fp8_act / fp8_act_scale arena buffers).