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
2026-09-25: Route a projection through row-wise FP8 cuBLASLt. The weight
pair comes from requant_weight_rowwise_fp8_cached; the activation is
quantized per token on every call.
cuBLASLt is handed ceil16(m) rows, so act_fp8_scratch must hold
ceil16(m) * k bytes, act_scale_scratch ceil16(m) f32 and out
ceil16(m) * n BF16 elements. The pad rows of both scratch buffers are
zeroed here, and the pad output rows are written.