Function w8a16_gemm_pipelined_m32_strided 

pub fn w8a16_gemm_pipelined_m32_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<()>

2026-09-25: Strided launch of w8a16_gemm_pipelined_m32.

input [M, a_row_stride] BF16 (first k of each row read), weight [N, K] FP8 E4M3 with block_scale [N/128, K/128] FP32, output [M, c_row_stride] BF16 (first n of each row written). Strides in elements. The A rows feed 16-byte cp.async chunks, so a_row_stride must keep them 16 B-aligned (multiple of 8).

Returns an error for m == 0, for k not a positive multiple of 128, for pitches shorter than k / n, and for an a_row_stride that is not a multiple of 8. Same argument order as w8a16_gemv_batch{4,16}_strided, so the multi-seq Q/K/V tier holds it as one more StridedBatchGemv arm.

Grid: (ceil(N/32), ceil(M/32), 1) Block: (256, 1, 1); 2026-09-27: at 1..=16 rows with K <= W8A16_M16_MAX_K, (ceil(N / 8), 1, 1).