pub fn moe_grouped_gate_up_cutlass(
gpu: &dyn GpuBackend,
host: &MoeCutlassHostTables,
a: DevicePtr,
sorted_token_ids: DevicePtr,
c_gate: DevicePtr,
c_up: DevicePtr,
expert_offsets: DevicePtr,
inter: u32,
hidden: u32,
stream: u64,
) -> Result<Vec<i32>>Expand description
2026-09-25: Single-launch CUTLASS grouped NVFP4 gate_up GEMM through
metrale_gpu_runtime::cutlass::nvfp4_grouped_gate_up_fused, with the weights from host.
expert_offsets is the device i32 [num_experts + 1] prefix sum; the host entry needs it on
the host, so it is copied and the stream synchronised first. Returns that host copy, which
moe_grouped_down_cutlass accepts to skip its own copy and synchronise.