pub fn mla_kv_assemble_batched(
gpu: &dyn GpuBackend,
kernel: KernelHandle,
kv_expanded: DevicePtr,
k_rope_buf: DevicePtr,
k_out: DevicePtr,
v_out: DevicePtr,
num_tokens: u32,
nkv: u32,
nope: u32,
v_dim: u32,
rope: u32,
hd: u32,
kv_expanded_stride: u32,
stream: u64,
) -> Result<()>Expand description
2026-09-25: Per token, build K as [k_nope | k_rope] for each KV head (k_rope
shared by all heads) and extract V from kv_expanded. blockIdx.y 0 builds
K, 1 extracts V.