Function moe_expert_gate_up_shared_bf16_batch2 

pub fn moe_expert_gate_up_shared_bf16_batch2(
    gpu: &dyn GpuBackend,
    kernel: KernelHandle,
    input: DevicePtr,
    gate_weight_ptrs: DevicePtr,
    gate_out: DevicePtr,
    up_weight_ptrs: DevicePtr,
    up_out: DevicePtr,
    expert_indices: DevicePtr,
    sh_gate_weight: DevicePtr,
    sh_gate_out: DevicePtr,
    sh_up_weight: DevicePtr,
    sh_up_out: DevicePtr,
    n: u32,
    k: u32,
    top_k: u32,
    stream: u64,
) -> Result<()>

2026-09-25: moe_expert_gate_up_shared_bf16 for two tokens in one launch (called from MoeLayer::forward_k2). blockIdx.y in [0, 2*top_k) is a routed expert (token y / top_k, slot y % top_k, output row token * top_k + slot); blockIdx.y == 2*top_k is the shared expert, computed for both tokens into rows 0 and 1.