Function w4a16_gemm_n128_m128 

pub fn w4a16_gemm_n128_m128(
    gpu: &dyn GpuBackend,
    kernel: KernelHandle,
    input: DevicePtr,
    weight: &QuantizedWeight,
    output: DevicePtr,
    m: u32,
    n: u32,
    k: u32,
    stream: u64,
) -> Result<()>

2026-09-25: W4A16 GEMM with a 128-row CTA (two 64-row chunks): each CTA loads a B tile once for both chunks.

Grid contract: N is the fast axis (blockIdx.x = N block, blockIdx.y = M block), as every w4a16_gemm_t_m128 in kernels/ reads it. Several layers share this launcher (qwen3_attention, dense_ffn, qwen3_ssm, nemotron_mamba2, nemotron_moe), so swapping the axes here would mis-map every CTA of the other kernels with no error. An M-fast order needs its own kernel and launcher, as w4a4_gemm_mfast and fp8_gemm_t_m128_mfast have.