pub fn gdn_verify_fused_conv_kn_batched(
gpu: &dyn GpuBackend,
kernel: KernelHandle,
conv_state: DevicePtr,
new_input: DevicePtr,
weight: &DenseWeight,
output: DevicePtr,
conv_state_inter: DevicePtr,
num_tokens: u32,
d_inner: u32,
d_conv: u32,
qk_channels: u32,
head_dim: u32,
input_stride: u32,
output_stride: u32,
inter_stride: u32,
l2_eps: f32,
n_seq: u32,
conv_state_seq_stride: u32,
input_seq_stride: u32,
output_seq_stride: u32,
inter_seq_stride: u32,
stream: u64,
) -> Result<()>Expand description
2026-09-25: gdn_verify_fused_conv_kn for n_seq sequences in one launch,
one grid row (blockIdx.y) per sequence.
The kernel offsets conv_state, new_input, output and
conv_state_inter by seq times their *_seq_stride, then runs the same
per-sequence body as gdn_verify_fused_conv_kn, so each sequence gets the
result a separate call would give it.
Grid: (ceil(d_inner/256), n_seq, 1) Block: (256, 1, 1)