Function w4a16_gemm_n128_ldb 

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

W4A16 GEMM with N_TILE=128: same kernel signature, wider N tile.

Grid: (ceil(N/128), ceil(M/64), 1) Block: (128, 1, 1) w4a16_gemm_n128 with an explicit transposed-B ROW STRIDE.

Needed when N is not a multiple of 16: the kernel’s B loads are 16-byte cp.async, which requires 16-byte-aligned sources, and row r sits at r * ldb. lm_head is the motivating case — its N is the vocab size, 248077 on this checkpoint, which is ODD and made 15 of every 16 k-rows fault with CUDA_ERROR_MISALIGNED_ADDRESS (the campaign’s long-standing “716”). Pass ldb = align_up(n, 128) with the pad columns zero-filled.