DeepSeek DeepEP探索：低延迟分发（low latency dispatch）的架构设计与实现

1. 低延迟分发的核心挑战与设计理念

在分布式AI推理场景中，数据分发的延迟直接影响整体系统性能。传统分发方案通常采用两阶段处理：先通过机间网络传输元数据，再通过NVLink转发实际数据。这种设计虽然节省显存，但不可避免地增加了通信延迟。DeepSeek DeepEP的low latency dispatch方案直击这一痛点，其设计哲学可概括为"用显存换速度"——通过预分配充足缓冲区、消除中间通知步骤，实现端到端的直接数据传输。

实测数据显示，在典型MoE（混合专家）模型推理场景中，该方案能将分发延迟降低40-60%。这得益于三个关键设计：首先是卡间直接RDMA通信，绕过传统的网络协议栈；其次是省略notify步骤，将元数据与数据合并传输；最后是创新的warp角色分配机制，实现计算与通信的精细重叠。这些优化使得系统在7B参数规模的模型上，单个token的分发时间能控制在20微秒以内。

2. 架构实现的关键技术细节

2.1 RDMA通信的极致优化

传统分布式训练中，GPU间通信往往需要CPU参与协调。DeepEP通过NVSHMEM的IBGDA（InfiniBand GPU Direct Async）模式，实现了GPU显存到显存的直接读写。代码中通过设置特定环境变量开启这一特性：

os.environ['NVSHMEM_DISABLE_P2P'] = '1' os.environ['NVSHMEM_IB_ENABLE_IBGDA'] = '1' os.environ['NVSHMEM_IBGDA_NIC_HANDLER'] = 'gpu'

这种设计带来两个显著优势：一是完全绕过CPU的参与，减少上下文切换开销；二是支持RDMA write with immediate特性，使得接收方能即时感知数据到达。实测表明，在100Gbps的InfiniBand网络下，这种直接通信方式比传统方案降低约30%的通信延迟。

2.2 显存管理的权衡策略

低延迟模式需要预先分配大量显存缓冲区。通过get_low_latency_rdma_size_hint函数可以计算所需缓冲区大小：

size_t get_low_latency_rdma_size_hint(int num_max_dispatch_tokens_per_rank, int hidden, int num_ranks, int num_experts) { auto num_bytes = LowLatencyLayout(nullptr, num_max_dispatch_tokens_per_rank, hidden, num_ranks, num_experts).total_bytes; return ((num_bytes + NUM_BUFFER_ALIGNMENT_BYTES) / NUM_BUFFER_ALIGNMENT_BYTES) * NUM_BUFFER_ALIGNMENT_BYTES; }

缓冲区设计采用"最坏情况"原则：发送缓冲区按最大可能token数分配，而接收缓冲区则假设所有token都可能汇聚到单个expert。虽然这会增加显存占用（通常比普通模式多2-3倍），但换来了确定性的内存访问模式和零拷贝的数据传输。

3. Warp级并行化设计

3.1 精细化的线程角色分配

DeepEP创新性地将SM内的warp划分为不同功能组。如下图所示，每个SM包含多个warp group，每个group对应一个expert处理：

SM架构： ┌──────────────┐ │ Warp Group 0 │→ Expert 0 ├──────────────┤ │ Warp Group 1 │→ Expert 1 ├──────────────┤ │ ... │ ├──────────────┤ │ Warp Group N │→ Expert N └──────────────┘

具体角色分配通过以下计算实现：

const auto warp_group_id = warp_id / kNumWarpsPerGroup; const auto sub_warp_id = warp_id % kNumWarpsPerGroup; const auto responsible_expert_idx = sm_id * kNumWarpGroups + warp_group_id;

这种设计带来两个好处：一是保证每个expert有专属计算资源，避免竞争；二是通过warp级别的任务划分，天然支持处理不同大小的工作负载。

3.2 数据发送的流水线优化

数据发送过程采用双缓冲设计，关键代码如下：

if (dst_rank != rank) { nvshmemi_ibgda_put_nbi_warp(dst_ptr, src_ptr, num_bytes_per_msg, dst_rank, dst_expert_local_idx, lane_id, slot_idx); } else { UNROLLED_WARP_COPY(8, lane_id, num_int4_per_msg, dst_int4_ptr, src_int4_ptr, ld_nc_global, st_na_global); }

对于跨节点通信使用RDMA write，而同节点通信则直接内存拷贝。通过UNROLLED_WARP_COPY宏展开循环，实现8倍指令级并行。实测显示，这种混合传输策略相比纯RDMA方案，在同节点通信场景下能提升15%的带宽利用率。

4. 同步机制的创新设计

4.1 无锁化的进度跟踪

系统采用原子计数器实现跨SM的进度同步：

atomic_add_release_global(atomic_finish_counter_per_expert + dst_expert_idx, 1);

每个完成数据发送的warp会递增计数器，而负责统计的warp则会添加补偿值。通过精心设计的FINISHED_SUM_TAG机制（通常设置为远大于最大可能token数的2的幂次方），实现无锁化的完成状态检测：

while (ld_acquire_global(atomic_finish_counter_per_expert + responsible_expert_idx) != FINISHED_SUM_TAG * 2);

这种设计完美解决了分布式系统中常见的"最后一个包"问题，避免了显式的全局同步操作。

4.2 接收端的自适应处理

接收端采用事件驱动模型，通过轮询完成队列(CQ)来感知数据到达：

nvshmemi_ibgda_poll_recv(src_rank, local_expert_idx); num_recv_tokens = ld_acquire_global(rdma_recv_count + local_expert_idx * num_ranks + src_rank);

为了提高吞吐，接收处理采用双阶段流水线：warp 0负责通知处理，warp 1并行执行数据拷贝。通过__syncwarp()指令保证组内线程同步，同时使用共享内存减少全局内存访问：

__shared__ int shared_num_recv_tokens[kNumWarpGroups]; __shared__ int shared_recv_token_begin_idx[kNumWarpGroups];

在实际部署中，这种设计使得接收端能在数据到达后1-2微秒内开始处理，极大缩短了端到端延迟。