w8a16_gemv_dual

spark_model::layers::ops

Function w8a16_gemv_dual 

Source 
pub fn w8a16_gemv_dual(
    gpu: &dyn GpuBackend,
    kernel: KernelHandle,
    input: DevicePtr,
    weight1: DevicePtr,
    row_scale1: DevicePtr,
    output1: DevicePtr,
    weight2: DevicePtr,
    row_scale2: DevicePtr,
    output2: DevicePtr,
    n: u32,
    k: u32,
    stream: u64,
) -> Result<()>
Expand description

W8A16 (FP8 E4M3) dual GEMV: two projections sharing the same BF16 input, one launch. blockIdx.z selects projection 0 (gate) vs 1 (up). Both N must be equal. Mirrors w4a16_gemv_dual but takes RAW DevicePtrs (FP8 weights are fp8w.weight / fp8w.row_scale, no QuantizedWeight wrapper, no scale2 f32).

Grid: (ceil(N/4), 1, 2) Block: (256, 1, 1)