DeepEP Normal vs Low-Latency 双模式深度分析：从 Blog 第一性原理到代码实现

deep_ep/buffers/legacy.py
  └── Buffer(low_latency_mode=False)
        ├── notify_dispatch<kLowLatencyMode=false>   ──┐
        ├── dispatch<kLowLatencyMode=false>            ├── internode.cu
        ├── cached_notify<kLowLatencyMode=false>       │
        └── combine<kLowLatencyMode=false>             ──┘
        └── (intranode.cu: NVLink only 路径)

# deep_ep/buffers/legacy.py:37-93
class Buffer:
    def __init__(self,
                 group,
                 num_nvl_bytes: int = 0,
                 num_rdma_bytes: int = 0,
                 low_latency_mode: bool = False,  # ← False = Normal 模式
                 num_qps_per_rank: int = 24,
                 allow_nvlink_for_low_latency_mode: bool = True,
                 ...):
        self.low_latency_mode = low_latency_mode
        self.runtime = _C.Buffer(self.rank, self.group_size,
                                 num_nvl_bytes, num_rdma_bytes,
                                 low_latency_mode, ...)  # 传入 C++ 运行时

// csrc/kernels/legacy/internode.cu:487
enum class WarpRole {
    kRDMASender,              // RDMA 发送者：从源 GPU 读取 Dispatch Buffer，写入 RDMA send buffer
    kRDMASenderCoordinator,   // RDMA 发送协调者：轮询 RDMA tail，按 Chunk 发起 put_nbi
    kRDMAAndNVLForwarder,     // RDMA+NVL 转发者：从 RDMA buffer 读取，通过 NVLink 转发到目标 GPU
    kForwarderCoordinator,    // 转发协调者：收集各 Forwarder 的 head 最小值，更新远端 RDMA head
    kNVLReceivers             // NVLink 接收者：从 NVL 读取，写入 recv_x（最终 Receive Buffer）
};

// csrc/kernels/legacy/internode.cu:499-516
const auto num_channels = num_sms / 2, channel_id = sm_id / 2;
const bool is_forwarder = sm_id % 2 == 0;  // 偶数 SM 是 Forwarder，奇数 SM 是 Sender/Receiver

const auto role_meta = [=]() -> std::pair<WarpRole, int> {
    if (is_forwarder) {
        if (warp_id < LEGACY_NUM_MAX_NVL_PEERS) {
            return {WarpRole::kRDMAAndNVLForwarder, (warp_id + channel_id) % LEGACY_NUM_MAX_NVL_PEERS};
        } else {
            return {WarpRole::kForwarderCoordinator, warp_id - LEGACY_NUM_MAX_NVL_PEERS};
        }
    } else if (warp_id < kNumDispatchRDMASenderWarps) {
        return {WarpRole::kRDMASender, -1};
    } else if (warp_id == kNumDispatchRDMASenderWarps) {
        return {WarpRole::kRDMASenderCoordinator, -1};
    } else {
        return {WarpRole::kNVLReceivers, (warp_id + channel_id - kNumDispatchRDMASenderWarps) % LEGACY_NUM_MAX_NVL_PEERS};
    }
};

flowchart TB
    subgraph Stage1["Stage 1: IB Sending (kRDMASender)"]
        S1["读取 Dispatch Buffer<br/>写入 RDMA send buffer<br/>(SymBuffer)"]
        SC["kRDMASenderCoordinator<br/>轮询 tail, 按 Chunk 发起 put_nbi"]
    end

    subgraph Stage2["Stage 2: IB-to-NVLink Forwarding (kRDMAAndNVLForwarder)"]
        F1["等待 RDMA meta 到达<br/>(rdma_channel_meta)"]
        F2["轮询 src_rdma_rank<br/>从 RDMA buffer 读取 Token"]
        F3["TMA load → mbarrier_wait<br/>TMA store → NVL buffer"]
        FC["kForwarderCoordinator<br/>收集 min_head, 更新远端 head"]
    end

    subgraph Stage3["Stage 3: NVLink Receiving (kNVLReceivers)"]
        R1["等待 NVL channel prefix<br/>(nvl_channel_prefix_start/end)"]
        R2["从 NVL buffer 读取 Token<br/>TMA load + TMA store"]
        R3["写入 recv_x<br/>(最终 Receive Buffer)"]
    end

    Stage1 -->|"nvshmemi_ibgda_put_nbi_warp"| Stage2
    Stage2 -->|"st_release_sys_global(nvl_channel_tail)"| Stage3

// csrc/kernels/legacy/internode.cu:527-529
// RDMA symmetric layout: Chunk 是通信粒度
auto rdma_channel_data = SymBuffer<uint8_t>(
    rdma_buffer_ptr,
    num_max_rdma_chunked_recv_tokens * num_bytes_per_token,  // ← Chunk 大小
    kNumRDMARanks, channel_id, num_channels);

// csrc/kernels/legacy/internode.cu:809-817
// 等待 Sender 完成一定数量的 Token 写入
auto num_tokens_processed = processed_tail - synced_last_issued_tail;
if (num_tokens_processed != synced_num_tokens_to_send
    and num_tokens_processed < num_max_rdma_chunked_send_tokens)
    continue;  // 未达到 Chunk 阈值，继续等待

// 按 Chunk 粒度发送
auto num_tokens_to_issue = min(num_tokens_processed, num_max_rdma_chunked_send_tokens);
nvshmemi_ibgda_put_nbi_warp<true>(dst_ptr, src_ptr, num_bytes_per_msg, ...);

// csrc/kernels/legacy/internode.cu:909-1013
} else if (warp_role == WarpRole::kRDMAAndNVLForwarder) {
    // 1. 等待 RDMA meta 到达（包含各 NVL rank 的 token 数前缀和）
    while (true) {
        auto meta_0 = ld_volatile_global(rdma_channel_meta.recv_buffer(lane_id) + dst_nvl_rank);
        auto meta_1 = ld_volatile_global(rdma_channel_meta.recv_buffer(lane_id) + LEGACY_NUM_MAX_NVL_PEERS + dst_nvl_rank);
        // meta_0, meta_1 < 0 表示数据已到达（编码为 -value - 1）
        if (meta_0 < 0 and meta_1 < 0 and meta_2 < 0 and meta_3 < 0) {
            int start_sum = -meta_0 - 1, end_sum = -meta_1 - 1;
            // 通知 NVL ranks
            st_relaxed_sys_global(nvl_channel_prefix_start.buffer() + lane_id, -start_sum - 1);
            st_relaxed_sys_global(nvl_channel_prefix_end.buffer() + lane_id, -end_sum - 1);
            break;
        }
    }

    // 2. 轮询 src_rdma_rank（round-robin），从 RDMA buffer 读取 Token
    while (__any_sync(0xffffffff, num_tokens_to_recv_from_rdma > 0)) {
        // 找到下一个有数据的 src_rdma_rank
        src_rdma_rank = (src_rdma_rank + 1) % kNumRDMARanks;
        if (__shfl_sync(0xffffffff, num_tokens_to_recv_from_rdma, src_rdma_rank) > 0) {
            if (lane_id == src_rdma_rank and cached_rdma_channel_head == cached_rdma_channel_tail)
                cached_rdma_channel_tail = ld_acquire_sys_global(rdma_channel_tail.buffer(src_rdma_rank));
            ...
        }

        // 3. 遍历每个 Token，通过 TMA 转发到 NVL buffer
        for (int i = src_rdma_head; i < src_rdma_tail; ++i) {
            auto rdma_slot_idx = i % num_max_rdma_chunked_recv_tokens;
            auto shifted = rdma_channel_data.recv_buffer(src_rdma_rank) + rdma_slot_idx * num_bytes_per_token;
            auto src_meta = ld_nc_global(reinterpret_cast<SourceMeta*>(shifted + hidden_bytes + scale_bytes));
            if (not is_in_dst_nvl_rank) continue;

            // TMA load from RDMA buffer, TMA store to NVL buffer
            int dst_slot_idx = (cached_nvl_channel_tail++) % num_max_nvl_chunked_recv_tokens;
            tma_load_1d(tma_buffer, shifted, tma_mbarrier, num_bytes_per_token, false);
            mbarrier_arrive_and_expect_tx(tma_mbarrier, num_bytes_per_token);
            mbarrier_wait(tma_mbarrier, tma_phase);
            tma_store_1d(tma_buffer, dst_shifted, num_bytes_per_token);
        }
    }
}

// csrc/kernels/legacy/internode.cu:1746
enum class WarpRole { kNVLSender, kNVLAndRDMAForwarder, kRDMAReceiver, kCoordinator };

// csrc/kernels/legacy/internode.cu:87-95
template <bool kLowLatencyMode>
__forceinline__ __device__ int translate_dst_rdma_rank(const int dst_rdma_rank, const int nvl_rank) {
    // Low-Latency: 每个 GPU 独立 RDMA rank (dst_rdma_rank * 8 + nvl_rank)
    // Normal: 同一 node 的 GPU 共享 RDMA rank (dst_rdma_rank)
    return kLowLatencyMode ? (dst_rdma_rank * LEGACY_NUM_MAX_NVL_PEERS + nvl_rank) : dst_rdma_rank;
}

template <bool kLowLatencyMode>
__forceinline__ __device__ void nvshmem_sync_with_same_gpu_idx(const nvshmem_team_t& rdma_team) {
    // Low-Latency: 仅同步 RDMA team（减少同步开销）
    // Normal: 全局同步
    kLowLatencyMode ? void(nvshmem_sync(rdma_team)) : nvshmem_sync_all();
}

deep_ep/buffers/legacy.py
  └── Buffer(low_latency_mode=True)
        └── internode_ll.cu
              ├── clean_low_latency_buffer
              ├── dispatch (发送+接收一体化)
              ├── combine  (发送+接收一体化)
              ├── query_mask_buffer / update_mask_buffer / clean_mask_buffer

// csrc/kernels/legacy/internode_ll.cu:128-155
template <bool kUseFP8, bool kUseUE8M0, int kHidden>
__global__ __launch_bounds__(1024, 1) void dispatch(...) {
    const auto sm_id = static_cast<int>(blockIdx.x);
    const auto warp_id = thread_id / 32, lane_id = get_lane_id();
    const auto warp_group_id = warp_id / num_warps_per_group;
    const auto sub_warp_id = warp_id % num_warps_per_group;
    const auto responsible_expert_idx = sm_id * num_warp_groups + warp_group_id;

    // 只有 2 种 Warp 角色:
    // 1. warp_id < num_warps - 1: FP8 cast + IBGDA send
    // 2. warp_id == num_warps - 1: 读取 topk_idx, 统计 per-expert 计数
    if (warp_id < num_warps - 1) {
        // 发送阶段
    } else if (warp_id == num_warps - 1) {
        // 计数阶段
    }
}

// csrc/kernels/legacy/internode_ll.cu:202-278
for (int token_idx = sm_id; token_idx < num_tokens; token_idx += num_sms) {
    const auto x_int4 = static_cast<const int4*>(x) + token_idx * hidden_bf16_int4;
    const auto rdma_x_src_idx = reinterpret_cast<int*>(static_cast<uint8_t*>(rdma_x) + token_idx * num_bytes_per_msg);

    // FP8 cast（如果启用）
    for (int i = thread_id; i < hidden_bf16_int4; i += num_threads) {
        auto int4_value = __ldg(x_int4 + i);
        if constexpr (kUseFP8) {
            // 计算 local amax, scale
            calculate_fp8_scales(amax, scale, scale_inv, round_scale);
            // Cast into send buffer
            rdma_x_vec[i] = int2_value;
        }
    }
    asm volatile("bar.sync 1, %0;" ::"r"(num_threads));

    // 直接通过 IBGDA 发送，无 Forwarding，无 Chunk 聚合
    if (dst_expert_idx >= 0) {
        int slot_idx = lane_id == 0 ? atomicAdd(atomic_counter_per_expert + dst_expert_idx, 1) : 0;
        slot_idx = __shfl_sync(0xffffffff, slot_idx, 0);
        const auto dst_rank = dst_expert_idx / num_local_experts;
        const auto dst_ptr = reinterpret_cast<uint64_t>(rdma_recv_x) + ...;
        const auto dst_p2p_ptr = nvshmemi_get_p2p_ptr(dst_ptr, rank, dst_rank);
        if (not is_rank_masked<true>(mask_buffer_ptr, dst_rank)) {
            if (dst_p2p_ptr == 0) {
                nvshmemi_ibgda_put_nbi_warp(dst_ptr, src_ptr, num_bytes_per_msg, dst_rank, ...);
            } else {
                // P2P 直接写入
                UNROLLED_WARP_COPY(8, lane_id, num_int4_per_msg, dst_int4_ptr, src_int4_ptr, ld_nc_global, st_na_global);
            }
        }
    }
}

// csrc/kernels/legacy/internode_ll.cu:362-370
// 接收端按 Expert 组织，而非按 Rank（与 Normal 不同）
const auto rdma_recv_x_uint8 = static_cast<uint8_t*>(rdma_recv_x) +
    local_expert_idx * num_ranks * num_max_dispatch_tokens_per_rank * num_bytes_per_msg +
    src_rank * num_max_dispatch_tokens_per_rank * num_bytes_per_msg;
const auto recv_x_int4 = static_cast<int4*>(packed_recv_x) +
    local_expert_idx * num_ranks * num_max_dispatch_tokens_per_rank * hidden_int4;

// csrc/kernels/legacy/internode_ll.cu:383-408
if (sub_warp_id == 1 and lane_id == 0) {
    auto start_time = clock64();
    uint64_t wait_recv_cost = 0;
    if (not is_rank_masked(mask_buffer_ptr, src_rank)) {
        // 等待 Token 到达，带超时检测
        while ((num_recv_tokens = ld_acquire_sys_global(rdma_recv_count + ...)) == 0
               && (wait_recv_cost = clock64() - start_time) <= LEGACY_NUM_TIMEOUT_CYCLES)
            ;
    }
    // 超时则 mask 该 rank（低延迟场景的容错机制）
    if (wait_recv_cost > LEGACY_NUM_TIMEOUT_CYCLES) {
        printf("Warning: DeepEP timeout for dispatch receive, rank %d, ...", rank, ...);
        if (mask_buffer_ptr == nullptr) trap();
        atomicExch(mask_buffer_ptr + src_rank, 1);
    }
}

deep_ep/buffers/elastic.py
  └── ElasticBuffer
        ├── dispatch_impl<kIsScaleupNVLink=false>  → hybrid_dispatch.cuh (多节点 = Normal 等价)
        ├── dispatch_impl<kIsScaleupNVLink=true>   → dispatch.cuh (单节点 NVLink only = Low-Latency 等价)
        └── (num_scaleout_ranks == 1 时自动选择 direct)

// csrc/kernels/elastic/dispatch.hpp:53-78
static std::string generate_impl(const Args& args) {
    std::string header_name, func_name;
    if (args.num_scaleout_ranks == 1) {
        // 单节点模式 (类似 Low-Latency 场景)
        header_name = "dispatch";
        func_name = fmt::format("dispatch_impl<{}, {}, ...>",
            args.is_scaleup_nvlink,  // ← true = NVLink only
            ...);
    } else {
        // 跨节点模式 (类似 Normal 场景)
        header_name = "hybrid_dispatch";
        func_name = fmt::format("hybrid_dispatch_impl<{}, {}, ...>",
            args.num_scaleout_warps, args.num_forward_warps,  // ← 有 forwarding warps
            ...);
    }
}

// deep_ep/include/deep_ep/impls/dispatch.cuh:17-30
template <bool kIsScaleupNVLink,      // ← 核心模式开关
          bool kDoCPUSync,
          bool kReuseSlotIndices,
          int kNumSMs,
          int kNumNotifyWarps, int kNumDispatchWarps,
          int kNumRanks,
          int kNumHiddenBytes, int kNumSFPacks,
          int kNumMaxTokensPerRank,
          int kNumExperts, int kNumTopk, int kExpertAlignment,
          int kNumQPs, int64_t kNumTimeoutCycles,
          ...>
__global__ void dispatch_impl(...)

// deep_ep/include/deep_ep/impls/hybrid_dispatch.cuh:14-32
template <..., int kNumScaleoutWarps, int kNumForwardWarps,  // ← Forward warps
          int kNumScaleoutRanks, int kNumScaleupRanks, ...>
__global__ void hybrid_dispatch_impl(...) {
    // Warp 角色:
    // 1. Notify warps: 元数据交换
    // 2. Scaleout warps: RDMA 发送
    // 3. Forward warps: NVLink 转发 ← 对应 Normal 的 Forwarding
}

// deep_ep/include/deep_ep/impls/hybrid_dispatch.cuh:88-97
auto scaleup_buffer = layout::BufferLayout<false>(
    token_layout, kNumScaleupRanks, kNumScaleoutRanks * kNumMaxTokensPerRank, buffer);
auto scaleout_send_buffer = layout::BufferLayout<false>(
    token_layout, 1, kNumMaxTokensPerRank, scaleup_buffer.get_buffer_end_ptr());
auto scaleout_recv_buffer = layout::BufferLayout<false>(
    token_layout, kNumScaleoutRanks, kNumChannels * kNumMaxTokensPerChannel, ...);

// deep_ep/include/deep_ep/impls/hybrid_dispatch.cuh:24-27
int kNumChannelsPerSM = kNumScaleoutWarps,
int kNumChannels = kNumScaleoutWarps * kNumSMs,
int kNumMaxTokensPerChannel = math::constexpr_ceil_div(kNumMaxTokensPerRank, kNumChannels),
int kScaleoutUpdateInterval = 6,           // ← Chunk 更新间隔
int kNumSlotsPerForwardChunk = kScaleoutUpdateInterval,  // ← Chunk 大小

// deep_ep/include/deep_ep/impls/dispatch.cuh:362-393
// Issue TMA NVLink stores (直接写入，无转发)
const auto dst_ptr = stored_dst_slot_idx >= 0 ?
    gin.get_sym_ptr<team_t>(recv_buffer.get_token_buffer(stored_dst_slot_idx).get_base_ptr(), stored_dst_rank_idx) :
    nullptr;
if (dst_ptr != nullptr)
    ptx::tma_store_1d(dst_ptr, tma_buffer.get_base_ptr(), tma_buffer.get_num_bytes<false>());

// Issue RDMA put (仅对非 NVLink 可达的 rank)
if constexpr (not kIsScaleupNVLink) {
    if (stored_dst_slot_idx >= 0 and dst_ptr == nullptr) {
        gin.put<team_t>(recv_buffer.get_token_buffer(stored_dst_slot_idx).get_base_ptr(),
                        send_buffer_ptr, tma_buffer.get_num_bytes<false>(), stored_dst_rank_idx);
    }
}

// deep_ep/include/deep_ep/impls/dispatch.cuh:278-292
// 每个 Token 独立处理，无 Chunk 聚合
const auto token_start = dispatch_warp_idx * kNumSMs + sm_idx;
const auto token_stride = kNumDispatchWarps * kNumSMs;
for (int token_idx = token_start; token_idx < num_tokens; token_idx += token_stride) {
    // Issue data TMA (单 Token)
    if (ptx::elect_one_sync()) {
        ptx::tma_load_1d(tma_buffer.get_hidden_ptr(),
                         math::advance_ptr(x, token_i64_idx * kNumHiddenBytes),
                         mbarrier_ptr, kNumHiddenBytes);
    }
    // ...
}

# deep_ep/buffers/elastic.py:729-834
@weak_lru(maxsize=None)
def get_theoretical_num_sms(self, num_experts: int, num_topk: int,
                            num_scaleout_topk: int = 0,
                            rdma_gbs: float = 0, nvlink_gbs: float = 0,
                            sm_read_gbs: float = 200, sm_write_gbs: float = 50) -> int:
    """
    基于带宽模型估算 dispatch/combine kernel 的最优 SM 数量。
    假设 balanced gate distribution。
    """
    # 1. 获取带宽
    if rdma_gbs == 0 and self.num_rdma_ranks > 1:
        rdma_gbs = get_rdma_gbs()
    if nvlink_gbs == 0:
        nvlink_gbs = get_nvlink_gbs()

    # 2. 计算期望 top-k scale-out/scale-up ranks
    def get_expected_topk(num_groups: int) -> float:
        return num_groups * (1 - math.comb(num_experts - num_experts // num_groups, num_topk)
                                / math.comb(num_experts, num_topk))

    num_expected_scaleout_topk = get_expected_topk(self.num_scaleout_ranks) if self.num_scaleout_ranks > 1 else 0
    num_expected_topk = get_expected_topk(self.num_ranks)

    # 3. 计算各类型流量
    sm_read, sm_write = 0, 0
    rdma_traffic, nvlink_traffic = 0, 0

    sm_read += 1 / num_expected_topk  # 读取 tokens

    if self.num_scaleout_ranks > 1:
        # Hybrid 模式 (Normal 等价)
        sm_write += 1 / num_expected_topk
        sm_write += (1 / num_expected_topk) * (num_expected_scaleout_topk / self.num_scaleout_ranks)
        rdma_traffic += (1 / num_expected_topk) * (num_expected_scaleout_topk * (1 - 1 / self.num_scaleout_ranks))
        sm_read += num_expected_scaleout_topk / num_expected_topk  # Forward warps
        sm_write += 1  # Issue scaleup
        nvlink_traffic += 1 - (1 / self.num_scaleup_ranks)
    else:
        # Direct 模式 (Low-Latency 等价)
        if self.num_rdma_ranks > 1:
            sm_write += 1 / num_expected_topk
        sm_write += self.num_nvlink_ranks / self.num_ranks
        nvlink_traffic += self.num_nvlink_ranks / self.num_ranks * (1 - 1 / self.num_nvlink_ranks)
        rdma_traffic += (self.num_ranks - self.num_nvlink_ranks) / self.num_ranks

    # 4. 找到瓶颈
    if self.num_scaleout_ranks > 1 and (rdma_traffic / rdma_gbs) > (nvlink_traffic / nvlink_gbs):
        bounded_traffic, bounded_gbs = rdma_traffic, rdma_gbs
    else:
        bounded_traffic, bounded_gbs = nvlink_traffic, nvlink_gbs

    # 5. 计算 SM 数量
    num_device_sms = torch.cuda.get_device_properties('cuda').multi_processor_count
    num_sms = num_device_sms  # 无流量时默认全部 SM
    if bounded_traffic > 0:
        num_sms = max(
            bounded_gbs / bounded_traffic * sm_read / sm_read_gbs,
            bounded_gbs / bounded_traffic * sm_write / sm_write_gbs,
        )
    num_sms = align(max(4, math.ceil(num_sms * 1.25)), 2)
    num_sms = num_sms if self.prefer_overlap_with_compute else max(num_sms, 64)
    num_sms = min(num_sms, num_device_sms)
    return num_sms

所需 SM 数 = 瓶颈带宽 / 单 Token 流量 × 每 SM 读写需求 / 每 SM 带宽能力

# tests/legacy/test_internode.py:324-330 (Normal 模式测试)
num_sms = 24
num_qps_per_rank = max(num_sms, ll_num_experts // num_ranks if args.test_ll_compatibility else 0)
buffer = deep_ep.Buffer(group,
                        num_nvl_bytes=nvl_buffer_size,
                        num_rdma_bytes=rdma_buffer_size,
                        low_latency_mode=args.test_ll_compatibility,  # ← False = Normal
                        num_qps_per_rank=num_qps_per_rank,
                        ...)

# tests/legacy/test_low_latency.py:263-270 (Low-Latency 模式测试)
buffer = deep_ep.Buffer(group,
                        num_rdma_bytes=num_rdma_bytes,
                        low_latency_mode=True,  # ← True = Low-Latency
                        num_qps_per_rank=num_experts // num_ranks,
                        allow_nvlink_for_low_latency_mode=not args.disable_nvlink,
                        ...)

# README.md:148-162
_buffer = ElasticBuffer(
    group,
    num_max_tokens_per_rank=num_max_tokens_per_rank,
    hidden=hidden,
    num_topk=num_topk,
    use_fp8_dispatch=use_fp8_dispatch,
)

# V2 analytically calculates the optimal SM count — no more auto-tuning needed
_num_comm_sms = _buffer.get_theoretical_num_sms(num_experts, num_topk)

// csrc/kernels/elastic/dispatch.hpp:53-78
if (args.num_scaleout_ranks == 1) {
    header_name = "dispatch";          // ← 单节点，类似 Low-Latency
} else {
    header_name = "hybrid_dispatch";   // ← 多节点，类似 Normal
}

flowchart LR
    subgraph V1["V1: 物理分离"]
        V1N["Normal Kernel<br/>(internode.cu)"]
        V1L["Low-Latency Kernel<br/>(internode_ll.cu)"]
        V1N ---|"完全独立代码"| V1L
    end

    subgraph V2["V2: 逻辑统一"]
        V2J["JIT 代码生成<br/>(dispatch.hpp)"]
        V2N["hybrid_dispatch_impl<br/>(scaleout > 1)"]
        V2L["dispatch_impl<br/>(scaleout = 1)"]
        V2J -->|"kIsScaleupNVLink=false"| V2N
        V2J -->|"kIsScaleupNVLink=true"| V2L
    end

    V1 -->|"统一化"| V2

// csrc/kernels/legacy/compiled.cuh:5-18
#define LEGACY_NUM_MAX_NVL_PEERS 8
#define LEGACY_NUM_MAX_RDMA_PEERS 20
#define LEGACY_NUM_WORKSPACE_BYTES (32 * 1024 * 1024)
#define LEGACY_NUM_MAX_LOCAL_EXPERTS 1024
#define LEGACY_NUM_BUFFER_ALIGNMENT_BYTES 128
#define LEGACY_LOW_LATENCY_SEND_PHASE 1
#define LEGACY_LOW_LATENCY_RECV_PHASE 2
#define LEGACY_FINISHED_SUM_TAG 1024
#define LEGACY_NUM_WAIT_NANOSECONDS 500
#define LEGACY_NUM_CPU_TIMEOUT_SECS 100
#define LEGACY_NUM_TIMEOUT_CYCLES 200000000000ull  // 200G cycles ~= 100s