Function moe_grouped_gate_up_cutlass 

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

Single-launch CUTLASS grouped NVFP4 fused gate_up GEMM (Phase-2).

Bridges the load-time host snapshot of the per-expert pointer/scale tables (MoeCutlassHostTables) to the host-side spark_runtime::cutlass::nvfp4_grouped_gate_up_fused entry. a is the expert-contiguous bf16 activation [total_expanded, k]; expert_offsets is the device i32 [num_experts+1] prefix sum — the only per-call table, so the only one copied and the only reason for the synchronize (the C entry indexes offsets on the host before it can launch). Returns the host copy of expert_offsets so the paired down call can reuse it instead of repeating the D2H + synchronize. The two calls share the same offsets (both are driven by one moe_sort_by_expert), and each sync blocks the host until the GPU drains — halving them halves that stall.