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
2026-09-25: Quantize the activation with per_token_group_quant_fp8 (FP8
plus one scale per token and 128 of K), then run
cublas_fp8_proj_prequant. The FP8 weight and its 128x128 block scales
go to cuBLASLt as they are, with no dequant.
The scratch buffers need the extents listed on
cublas_fp8_proj_prequant.