Function cublas_fp8_proj_prequant 

pub fn cublas_fp8_proj_prequant(
    gpu: &dyn GpuBackend,
    scale_kmajor_k: KernelHandle,
    act_fp8: DevicePtr,
    act_scale: DevicePtr,
    act_scale_kmajor: DevicePtr,
    fp8w: &Fp8Weight,
    out: DevicePtr,
    m: u32,
    n: u32,
    k: u32,
    stream: u64,
) -> Result<()>

2026-09-25: cublas_fp8_proj for an activation the caller already ran through per_token_group_quant_fp8, so one quantization can feed several projections (the dense FFN’s gate and up share one).

With cublas_scale_layout_kmajor on, the scales are first transposed into act_scale_kmajor; a zero kernel handle or scratch pointer is an error.

cuBLASLt is handed ceil16(M) rows (cublas_fp8_m_pad), so:

• out must hold ceil16(M) * N BF16 elements; the pad rows are written.
• act_fp8 must hold ceil16(M) * K bytes; the pad rows are zeroed here.
• act_scale_kmajor must hold ceil16(M) * (K/128) f32, and act_scale M * (K/128) f32, or ceil16(M) * (K/128) under rowmajor, where the pad rows are zeroed in place.