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

W4A16 GEMM: C = A @ B with 2-M-chunk CTA (M_TILE2=128).

Halves weight re-reads vs w4a16_gemm_n128 for large M (ISL > 128): each CTA covers 128 rows of A, loading B once for both 64-row halves. ~2× speedup on qkvz (K=2048, N=12288) at ISL=1016.

Grid: (ceil(N/128), ceil(M/128), 1) Block: (128, 1, 1) SMEM: ~29.8 KB → 3 blocks/SM (vs 5 for m64 at ~19.6 KB).

GRID CONTRACT — N is the FAST axis (blockIdx.x = N-block, blockIdx.y = M-block). Every w4a16_gemm_t_m128 kernel across all model dirs reads it this way. This launcher is SHARED (qwen3_attention, dense_ffn, qwen3_ssm, nemotron_*), so the axes must NOT be swapped here to suit one model: doing so silently mis-maps every CTA for the other 18 kernels and produces garbage output with no error. If a model wants the m-fast (L2-friendly) order, add a SEPARATELY NAMED kernel + launcher (see w4a4_gemm_mfast / fp8_gemm_t_m128_mfast) rather than mutating this one.