深度解读DeepSeek：开源周（Open Source Week）技术解读

深度解读DeepSeek：开源周（Open Source Week）技术解读
深度解读DeepSeek：源码解读 DeepSeek-V3
深度解读DeepSeek：技术原理
深度解读DeepSeek：发展历程

2025年2月24日至28日，DeepSeek通过连续五天的“开源周”活动，向全球开源了8个核心项目，覆盖AI训练、推理、并行计算等多个关键领域。

一、开源内容概览

DeepSeek 开源周共发布 五个核心项目 及多个辅助工具，涵盖 AI 开发的三大核心领域：计算优化、通信效率、存储加速。以下是各项目的核心价值：

Day1：FlashMLA

功能：针对 NVIDIA Hopper GPU 优化的多头线性注意力解码内核，支持可变长度序列处理。

突破：在 H800 GPU 上实现 580 TFLOPS 计算性能 和 3000 GB/s 内存带宽，推理效率提升 2-3 倍，适用于实时翻译、长文本处理等场景18。

意义：打破大厂对高效推理工具的垄断，降低开发者使用门槛，推动边缘设备部署。

Day2：DeepEP

功能：专为混合专家模型（MoE）设计的通信库，优化节点间数据分发与合并。

突破：通过低延迟内核和通信-计算重叠技术，实现 训练速度提升 3 倍、延迟降低 5 倍，支持 FP8 低精度通信。

意义：挑战英伟达 NCCL 生态，打破硬件与软件耦合的技术壁垒。

Day3：DeepGEMM

功能：基于 FP8 的高效矩阵乘法库，专为 MoE 模型优化。

突破：代码仅 300 行，通过即时编译（JIT）和 CUDA 核心双层累加技术，实现 1.1-2.7 倍加速，最高性能达 1350 TFLOPS。

意义：推动低精度计算普及，降低千亿参数模型部署成本，成为“AI 工业革命的基石”。

Day4：DualPipe & EPLB

功能：创新双向流水线并行算法（DualPipe）与动态负载均衡工具（EPLB）。

突破：通过任务交叉排布和专家模型动态复制，减少 GPU 空闲时间，优化资源利用率。

意义：类比“泰勒管理制”和“福特流水线”，重构 AI 训练流程，提升工业级效率。

Day5：3FS

功能：高性能分布式文件系统，支持 RDMA 网络和 SSD 存储。

突破：实现 6.6 TB/s 读取速度，加速海量数据训练与推理阶段的向量搜索。

意义：补全 AI 基础设施的最后一块拼图，解决存储瓶颈问题。

二、技术突破与创新

DeepSeek 开源周的核心技术突破体现在以下三方面：

• 硬件性能压榨

  • GPU 极限优化：如 FlashMLA 将 H800 GPU 的内存带宽利用率提升至理论极限的 90%，DeepGEMM 通过直接编写机器指令绕过 CUDA 生态限制。

  • 低精度计算革命：FP8 的广泛应用（如 DeepGEMM）在保证精度损失 <0.5% 的前提下，将存储和算力需求降低至 FP32 的 1/4。

• 并行计算重构

  • 通信与计算重叠：DeepEP 通过钩子机制实现通信与计算并行，减少 GPU 空闲时间。
  • 动态负载均衡：EPLB 根据专家模型调用频率动态调整任务分配，避免 GPU 资源浪费。

• 开源生态挑战

  • 对抗英伟达垄断：DeepGEMM 和 DeepEP 直接挑战 CUDA 和 NCCL 生态，推动国产软硬件适配。
  • 透明化技术黑箱：公开训练框架分析数据，推动行业技术共享与协作。

三、技术原理

DeepEP

MoE介绍：深度解读DeepSeek
网络知识扫盲：GPU通信互联技术介绍
传统的基于NVSwitch的All-to-All通信结构：

让我们再看看DeepSeek-V3的论文，它在里面提到的相关优化点：

3.2.2. EfficientImplementationofCross-NodeAll-to-AllCommunication

技术术语说明：

• “SM”：中文全称"流式多处理器"
• “IB”：InfiniBand
• “token”：在AI上下文统一译为"token"不译
• “warp”：采用计算机体系结构标准译法"线程束"
• “PTX”： Parallel Thread Execution

翻译一下：

为确保DualPipe架构的计算性能，我们定制了高效的跨节点全互联通信内核（包括分发与聚合操作），以优化通信专用流式多处理器（SM）的资源占用。该内核实现与MoE门控算法及集群网络拓扑深度协同设计：在我们的集群中，跨节点GPU通过InfiniBand（IB）全互联，节点内则采用NVLink通信。NVLink提供160GB/s带宽，约为IB（50GB/s）的3.2倍。为充分发挥异构带宽优势，我们限制每个token最多分发至4个节点，从而降低IB流量。当token路由决策完成后，首先通过IB传输至目标节点中与本节点同索引的GPU，到达后立即通过NVLink转发至托管目标专家的特定GPU，避免被后续到达的token阻塞。这种设计使得IB与NVLink通信完全重叠（overlapped），每个token可高效选择每节点平均3.2个专家，且不产生额外NVLink开销。这意味着DeepSeek-V3实际仅路由8个专家时，在保持通信成本不变的前提下，理论上可扩展至最多13个专家（4节点×3.2专家/节点）。

DeepSeek还采用了warp（线程束）专用化技术：整套通信策略仅需占用20个SM即可完全榨取IB与NVLink的带宽潜力。具体实现中，我们采用线程束专业化技术（Bauer et al., 2014），将20个SM划分为10条通信通道。分发过程中：（1）IB发送、（2）IB至NVLink转发、（3）NVLink接收分别由专用线程束处理，各通信任务的线程束数量根据SM实际负载动态调节。聚合过程同理：（1）NVLink发送、（2）NVLink至IB转发与累加、（3）IB接收与累加也采用动态线程束分配。此外，由于分发与聚合内核均与计算流重叠执行，我们通过定制PTX指令集和自动调优通信分块大小，显著降低L2缓存占用，减少对其他SM计算内核的干扰。

下面正式看看DeepEP是如何实现的：

DeepEP是由DeepSeek团队推出的首个专为混合专家模型（MoE）设计的专家并行（EP）通信库。主要解决MoE模型在分布式训练和推理中的通信瓶颈问题，通过优化数据传输和资源调度，实现“降本增效”。

• 高效的全对全通信（All-to-All）：支持节点内（NVLink）和节点间（RDMA）的高带宽通信，优化数据在不同专家子网络间的快速交换。

• 动态资源调控：基于群组限制门控算法（group-limited gating），动态分配GPU计算单元（SM）数量，任务多时增加资源，任务少时降低功耗，减少资源浪费。支持低精度运算：原生支持FP8格式，减少内存占用并加速计算，适用于大规模分布式训练

之前的MoE普遍使用NCCL的p2p通信进行all-to-all，结果纷纷吐槽all-to-all性能差，带宽利用率低。但是，很少有人真的去分析all-to-all性能差的原因，并尝试去改进。而DeepEP的出现，可以说彻底解决了all-to-all打不满带宽的问题。DeepEP直接放弃了NCCL，转而使用更底层的NVSHMEM进行通信，结果基本用满了NVLink和IB的带宽。

ps：NCCL（NVIDIA Collective Communications Library）是英伟达（NVIDIA）开发的一个专为多GPU和多节点通信优化的库，主要用于加速分布式深度学习训练和科学计算中的集体通信操作。，支持以下关键操作：

• All-Reduce：将多个GPU上的数据汇总（如梯度求和）并分发回所有GPU。
• Broadcast：将单个GPU的数据广播到其他所有GPU。
• All-Gather：将每个GPU的数据收集到所有GPU。
• Reduce-Scatter：将数据汇总后按块分发给各个GPU。
• Point-to-Point：单GPU到单GPU的直接数据传输。

下面我们沿着README.md里的例子，分析一下DeepEP究竟是如何实现的。

https://github.com/deepseek-ai/DeepEP

DeepEP主要实现了EP通信过程（dispatch分发+combine聚合），涉及两种实现模式：

• Normal kernel模式：Nvlink+RDMA混合模式，用于训练/推理prefill阶段
• Low-latency kernel模式：IBGDA+纯RDMA模式，用于推理decode阶段

性能测试结果：

初始化：

_buffer = Buffer(group, num_nvl_bytes, num_rdma_bytes)

Buffer是DeepEP的核心数据结构，位于deep_ep/buffer.py。

class Buffer:"""专家并行(EP)通信核心缓冲区，专为混合专家(MoE)模型设计，支持以下通信模式：1. 节点内高吞吐量all-to-all通信(dispatch和combine操作，使用NVLink)2. 跨节点高吞吐量all-to-all通信(使用RDMA，不支持原子操作AR)3. 低延迟all-to-all通信(使用RDMA，支持原子操作AR)主要功能：- 管理节点内(NVLink)和跨节点(RDMA)的通信缓冲区- 提供高效的dispatch(分发)和combine(聚合)操作- 支持通信-计算重叠优化属性说明:num_sms: 用于高性能计算内核的流式多处理器(Streaming Multiprocessors)数量(SM是GPU的计算单元，影响并行计算能力)rank: 当前进程在通信组中的唯一标识(本地rank号)，范围[0, group_size-1]例如：在4机8卡训练中，rank 0-7对应机器1的8个GPUgroup_size: 通信组中的总进程数(即world size)group: PyTorch的进程组(ProcessGroup)对象，管理分布式通信上下文num_nvl_bytes: 节点内通过NVLink通信的缓冲区大小(字节)NVLink是NVIDIA GPU间的高速互联技术num_rdma_bytes: 跨节点RDMA通信的缓冲区大小(字节)(在低延迟模式下也用于节点内通信)RDMA是远程直接内存访问技术，可绕过CPUruntime: 底层C++运行时对象，实际执行通信操作的核心组件"""

Dispatch（分发）:

• dispatch_forward：前向传播，将输入数据（x）根据topk_idx和topk_weights分配到不同的专家（可能位于不同GPU/进程）。
• dispatch_backward：反向传播，通过combine操作合并来自各节点的梯度（反向传播的Dispatch对应正向Combine的逆过程）。

Combine（合并）:

• combine_forward：前向传播，聚合各专家处理后的结果，生成最终输出。
• combine_backward：反向传播，通过dispatch操作将梯度分发回各节点（反向传播的Combine对应正向Dispatch的逆过程）。

接下来我们看下dispatch_forward 过程：

def dispatch_forward(x: Union[torch.Tensor, Tuple[torch.Tensor, torch.Tensor]],topk_idx: torch.Tensor, topk_weights: torch.Tensor,num_experts: int, previous_event: Optional[EventOverlap] = None
) -> Tuple[Union[torch.Tensor, Tuple[torch.Tensor, torch.Tensor]],  # recv_xtorch.Tensor,  # recv_topk_idxtorch.Tensor,  # recv_topk_weightsList,          # num_recv_tokens_per_expert_listTuple,         # handle (通信句柄，用于后续操作)EventOverlap    # event (CUDA事件，用于同步)
]:"""MoE（混合专家）模型的前向分发（dispatch）操作。将输入数据根据topk索引和权重分发到不同的专家（可能位于不同GPU/进程）。Args:x: 输入数据，可以是单个张量或包含数据和门控值的元组。topk_idx: [num_tokens, top_k] 形状的张量，表示每个token选择的topk专家索引。topk_weights: [num_tokens, top_k] 形状的张量，对应topk专家的权重。num_experts: 专家总数（用于确定分发目标范围）。previous_event: 可选的CUDA事件，用于建立操作间的依赖关系（实现通信-计算重叠）。Returns:recv_x: 接收到的数据（来自其他进程的专家结果）。recv_topk_idx: 接收到的专家索引（可能与本地索引不同）。recv_topk_weights: 接收到的专家权重。num_recv_tokens_per_expert_list: 每个专家接收到的token数量统计列表。handle: 通信句柄，用于后续反向传播（combine操作）。event: CUDA事件，用于同步通信完成状态。"""global _buffer  # 使用全局通信缓冲区# === 步骤1: 计算分发布局 ===# 确定每个rank/专家需要处理的token数量及通信策略# 返回值:#   - num_tokens_per_rank: 每个进程接收的token总数#   - num_tokens_per_rdma_rank: 通过RDMA通信的token数（跨节点）#   - num_tokens_per_expert: 每个专家处理的token数#   - is_token_in_rank: 标记token是否应由当前rank处理#   - updated_event: 更新后的CUDA事件（包含新依赖）num_tokens_per_rank, num_tokens_per_rdma_rank, num_tokens_per_expert, is_token_in_rank, previous_event = \_buffer.get_dispatch_layout(topk_idx, num_experts,previous_event=previous_event,  # 继承事件依赖async_finish=True,              # 异步执行allocate_on_comm_stream=previous_event is not None  # 内存分配绑定到通信流)# === 步骤2: 执行MoE分发 ===# 实际的数据分发操作（可能涉及跨GPU/NVLink或跨节点/RDMA通信）# 注意:#   - 此操作会阻塞CPU直到GPU完成，因此不兼容CUDA Graph#   - 使用async_finish=True实现异步通信recv_x, recv_topk_idx, recv_topk_weights, num_recv_tokens_per_expert_list, handle, event = \_buffer.dispatch(x,topk_idx=topk_idx,topk_weights=topk_weights,num_tokens_per_rank=num_tokens_per_rank,          # 布局计算结果num_tokens_per_rdma_rank=num_tokens_per_rdma_rank,is_token_in_rank=is_token_in_rank,num_tokens_per_expert=num_tokens_per_expert,previous_event=previous_event,  # 确保在依赖事件完成后执行async_finish=True,              # 异步通信allocate_on_comm_stream=True    # 内存分配在通信流执行（避免阻塞计算流）)# 返回结果及通信元信息（用于后续操作和反向传播）return recv_x, recv_topk_idx, recv_topk_weights, num_recv_tokens_per_expert_list, handle, event

dispatch_forward函数步骤：

• 异步通信：通过async_finish=True和EventOverlap实现通信与计算重叠，提升效率。
• 布局计算：调用_buffer.get_dispatch_layout，确定分发策略（如每个进程处理的Token数量）：根据本地的topk_idx，来计算本地要发往每个rank和每个expert的token数量。其内部使用了GPU来加速计算，具体的kernel代码我们略过。
• 数据分发：调用_buffer.dispatch，将输入数据分发到目标进程。如果需要机间通信，则调用internode_dispatch。
• 返回值：
  • recv_x：接收到的数据（来自其他进程）。
  • recv_topk_idx/recv_topk_weights：接收到的索引和权重。
  • event：CUDA事件，用于后续同步。

因为DeepEP主要就是对机间通信进行了很大的优化，因此我们来看其内部是怎么实现的。
Buffer的 self.internode_dispatch直接调用了self.runtime.internode_dispatch，代码位于csrc/deep_ep.cpp

std::tuple<torch::Tensor, ...>
Buffer::internode_dispatch(const torch::Tensor& x, ...) {bool cached_mode = cached_rank_prefix_matrix.has_value();// 1个channel对应2个SMconst int num_channels = config.num_sms / 2;// 设置comm_stream// Allocate all tensors on comm stream if set// NOTES: do not allocate tensors upfront!auto compute_stream = at::cuda::getCurrentCUDAStream();if (allocate_on_comm_stream) {EP_HOST_ASSERT(previous_event.has_value() and async);at::cuda::setCurrentCUDAStream(comm_stream);}// 等待前置任务完成// Wait previous tasks to be finishedif (previous_event.has_value()) {stream_wait(comm_stream, previous_event.value());} else {stream_wait(comm_stream, compute_stream);}if (cached_mode) {// 如果之前进行过dispatch，则可以重用之前的结果internode::cached_notify(...);}else {// 否则，需要进行计算internode::notify_dispatch(...);}// 等待notify_dispatch完成，正式进行dispatchinternode::dispatch(...);
}

notify_dispatch使用NVSHMEM在所有rank之间进行通信，计算互相发送的token数量以及负责的token区域，具体包括如下内容：

rdma_channel_prefix_matrix：形状(num_rdma_ranks, num_channels)，每个channel要发往每个RDMA节点token数量的前缀和
recv_rdma_rank_prefix_sum：形状(num_rdma_ranks)，每个RDMA节点要接收的token数量
gbl_channel_prefix_matrix：形状(num_ranks, num_channels)，每个channel要发往每个GPU的token数量的前缀和
recv_gbl_rank_prefix_sum：形状(num_ranks)，每个GPU要接收的token数量
moe_recv_counter：int，总共要接收的token数量
moe_recv_expert_counter：int[NUM_MAX_LOCAL_EXPERTS]，每个本地的expert要接收的token数量

然后，创建接收数据的tensor。最后，正式进行dispatch。dispatch的核心代码位于csrc/kernels/internode.cu，由于代码非常长，我们这里用文字来讲解它的流程。

在MoE里，一个token可能会发往多个GPU，这些GPU可能位于多个节点上（Deepseek-V3规定了一个token最多发往4个节点）。对于一个token，它首先经过rdma channel，从本地传输到所有的远端节点上编号相同的GPU。然后再经过nvl_channel，传输远端节点中所有的目标GPU上。

DeepEP使用多个channel发送数据。DeepEP将每个GPU上的数据划分为num_channels个连续的段，每一段用一个channel发送。其中每个channel包含一个rdma_channel和一个nvl_channel。rdma_channel和nvl_channel本质上都是环形队列。

下面这张图展示了整体的工作流程，注意：为了方便，这里只展示了一个token发往一个目标GPU的过程。实际上，每个token至多发往4个dst rdma rank，8个dst nvl rank。图中的黄框代表GPU，实线代表数据流经的路径，虚线代表控制信息。

参考：
https://blog.csdn.net/bit_mike/article/details/145566512
https://www.cnblogs.com/CQzhangyu/p/18741625