Function w8a16_gemm_m16_strided 

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

Strided sibling of w8a16_gemm_m16: a_row_stride / c_row_stride are the A and C row pitches in ELEMENTS, for callers whose rows are not contiguous. The multi-seq decode QKV buffer is [n, per_seq_qkv] with Q at offset 0, K after Q and V after K inside every row, so one launch per projection writes all m rows straight into their slots — the same reason w8a16_gemv_batch16_strided exists, and the same argument order.

a_row_stride must keep each activation row 16-byte aligned (a multiple of 8 BF16): the kernel stages A with 16-byte cp.async chunks.