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

2026-09-25: w4a16_gemm_n128 with an explicit row stride ldb for the transposed B, so row r starts at r * ldb.

The tile kernels load B in 16-byte cp.async chunks, which need 16-byte aligned sources, so an N that is not a multiple of 16 cannot be the stride. The lm_head’s N is the vocab size, which can be odd; its padded twin is built at align_up(vocab, 128) with the pad columns zeroed (impl_a1.rs, transpose.rs), and callers pass that as ldb.