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.