深度解读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

网络知识扫盲：GPU通信互联技术介绍
MoE介绍：深度解读DeepSeek

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

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

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

传统的基于NVSwitch的All-to-All通信结构：

通信内核（通信SM控制代码）的实现与MoE门控算法和集群网络拓扑是按照软硬件协同的思路来进行设计的。具体来说，在集群中，跨节点GPU与IB完全互连，节点内（单台服务器内）通信通过NVLink完成。NVLink 提供 160 GB/s 的带宽，约是IB的 3.2 倍 （50 GB/s）。

之前的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的直接数据传输。

为了有效利用IB和NVLink的不同带宽，DeepSeek将每个Token（词元）的分发限制为最多4 个节点，从而减少IB流量限制的影响。 对于每个Token，在做节点间路由决策时，先通过IB传输到目标节点上具有相同节点内索引的GPU；到达目标节点后，再通过NVLink 转发到托管目标专家的特定GPU。通过这种方式，通过IB和NVLink的通信重叠，平均每个Token可以在每个节点选择3.2名专家，而不会产生额外的NVLink开销。

实际算法中，V3/R1只通过路由选择了8个专家，但在保持相同通信成本的情况下，该架构可以扩展到最多13 个专家（4 个节点x3.2个专家/节点）。

DeepSeek还采用了warp（线程束）专用化技术，将20个SM划分为10个通信信道。

1）在调度过程中，（a）IB 发送、（b）IB 到NVLink 转发、（c） NVLink 接收由相应的warp处理。分配给每个通信任务的warp数量会根据所有SM的实际工作负载动态调整。

2）在合并过程中，（a） NVLink 发送、（b）NVLink到IB的转发和累积、（c）IB接收和累积也由动态调整的warp处理。

3）dispatching 和combining kernel都与计算流重叠，采用定制的PTX（Parallel Thread Execution）指令以自动调整通信块大小，减少了对L2缓存的使用和对其他 SM 的干扰。

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

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

性能测试结果：

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

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

初始化：

_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

我们来看其中两个具体的函数：

• get_dispatch_layout：
  布局计算，确定分发策略（如每个进程处理的Token数量）
  根据本地的topk_idx，来计算本地要发往每个rank和每个expert的token数量。其内部使用了GPU来加速计算，具体的kernel代码我们略过。
• dispatch：
  数据分发，将输入数据分发到目标进程
  如果需要机间通信，则调用internode_dispatch。因为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