pub fn gdn_decode_f16_strided_norm(
gpu: &dyn GpuBackend,
kernel: KernelHandle,
h_state: DevicePtr,
query: DevicePtr,
key: DevicePtr,
value: DevicePtr,
gate: DevicePtr,
beta: DevicePtr,
z_gate: DevicePtr,
norm_weight: DevicePtr,
output: DevicePtr,
batch_size: u32,
num_k_heads: u32,
num_v_heads: u32,
k_dim: u32,
v_dim: u32,
qk_stride: u32,
v_stride: u32,
gb_stride: u32,
z_stride: u32,
out_stride: u32,
h_seq_stride: u64,
eps: f32,
stream: u64,
) -> Result<()>Expand description
2026-09-25: FP16 h-state twin of gdn_decode_f32_strided_norm, used when
ssm_h_fp16_enabled() is set. It adds one argument, h_seq_stride: the
per-sequence slot pitch of the h-state pool in __half elements. The pitch
depends on how the pool was sized, not on the head dims, so the caller
passes it.