pub fn moe_weighted_sum_blend_fp8_grouped(
gpu: &dyn GpuBackend,
kernel: KernelHandle,
output: DevicePtr,
expert_out: DevicePtr,
expert_weights: DevicePtr,
token_to_perm: DevicePtr,
shared_out: DevicePtr,
input: DevicePtr,
gate_weight: DevicePtr,
hidden: u32,
top_k: u32,
k: u32,
num_tokens: u32,
stream: u64,
) -> Result<()>Expand description
2026-09-25: Grouped blend, one blockIdx.y per token. expert_out rows
are sorted positions; token_to_perm[token * top_k + k] maps a slot to its
row. gate_weight may be NULL (ungated shared expert).