w8a16_gemm_m16_strided

Function w8a16_gemm_m16_strided 

Source
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<()>
Expand description

2026-09-25: Strided 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 (qkv_fp8_batch.rs) is [n, per_seq_qkv] with Q at offset 0, K after Q and V after K in every row, so one launch per projection writes all m rows into their slots. Same argument order as w8a16_gemv_batch16_strided.

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