Function moe_expert_gate_up_shared_batch2 

pub fn moe_expert_gate_up_shared_batch2(
    gpu: &dyn GpuBackend,
    kernel: KernelHandle,
    input: DevicePtr,
    gate_packed_ptrs: DevicePtr,
    gate_scale_ptrs: DevicePtr,
    gate_scale2_vals: DevicePtr,
    gate_out: DevicePtr,
    up_packed_ptrs: DevicePtr,
    up_scale_ptrs: DevicePtr,
    up_scale2_vals: DevicePtr,
    up_out: DevicePtr,
    expert_indices: DevicePtr,
    sh_gate: &QuantizedWeight,
    sh_gate_out: DevicePtr,
    sh_up: &QuantizedWeight,
    sh_up_out: DevicePtr,
    n: u32,
    k: u32,
    top_k: u32,
    block_size: u32,
    stream: u64,
) -> Result<()>

2026-09-25: Gate and up GEMVs of the routed and shared experts for two tokens in one launch (row layout in the module doc). The shared expert uses direct weight pointers, the routed experts the pointer tables. A block_size of 256 selects the kernel’s wide decomposition (blockDim.x == 256); any other width runs the 128-thread one.