moe_grouped_gate_up_cutlass

Function moe_grouped_gate_up_cutlass 

Source
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.