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<()>

cublas_fp8_proj for an activation that is ALREADY quantized — the caller ran per_token_group_quant_fp8 itself.

WHY the split (#917/#928): the dense FFN’s gate and up projections consume the SAME [M, K] input, so quantizing inside the GEMM helper would pay the per-token quant twice per layer. The FFN quantizes once and calls this for both, then quantizes the post-SiLU intermediate once for down.

⚠ SCALE LAYOUT. cuBLASLt reads the VEC128 B-scale tensor with the TOKEN index contiguous ([K/128, ceil16(M)]), not the [M, K/128] the quantizer writes — cuBLAS “Scaling factors layouts”, and the reason this helper needs act_scale_kmajor at all. Handing the quantizer’s buffer over directly is what the 2026-09-11 H100 run measured at rel_rms 7.7e-2 / ~33 000 BF16 ULP against the in-tree kernel on identical FP8 bytes; ATLAS_CUBLAS_SCALE_LAYOUT =rowmajor reproduces that reading deliberately.

⚠ PADDED-M EXTENTS. cuBLASLt is handed ceil16(M), so:

• out must hold ceil16(M) * N BF16 elements — the phantom rows are WRITTEN (well-defined: their activation scales are zeroed below).
• act_fp8 must hold ceil16(M) * K bytes, act_scale M * (K/128) f32 and act_scale_kmajor ceil16(M) * (K/128) f32 — the phantom rows are READ.

The arena sizes that headroom in; see the sizing notes in spark_runtime::buffers::sizes (fp8_act, ffn_act_a, ffn_act_scale, ffn_act_scale_kmajor, expert_gate_out, moe_output).