DeepEP V1Legacy实战指南基于 NVSHMEM 的专家并行 MoE 通信库全解析【免费下载链接】DeepEPDeepEP: an efficient expert-parallel communication library项目地址: https://gitcode.com/GitHub_Trending/de/DeepEP本文是 DeepEP V1NVSHMEM 后端版本的存档技术文档围绕 MoE 专家并行EP中最为核心的dispatch / combine 全对全通信内核展开涵盖高吞吐普通内核与低延迟内核两套方案。文中所有接口、参数、命令行与性能数据均以当前仓库 docs/legacy.md 为骨架并辅以 deep_ep/buffers/legacy.py、csrc/legacy/config.hpp、setup.py 等源码级证据。读完本文你将掌握如何在训练/预填充阶段使用支持 SM 数量控制与 FP8 的高吞吐 dispatch/combine如何在推理解码阶段使用纯 RDMA 低延迟内核并借助接收 hook 实现零 SM 占用的通信-计算重叠以及集群网络VL、自适应路由与自动调优的最佳实践。1. 背景与定位DeepEP V1 是什么DeepEPDeepEveryParallelV1 是面向现代机器学习的高性能通信库聚焦于专家并行EP场景。它提供被称为MoE dispatch 和 combine的高吞吐、低延迟全对全all-to-allGPU 内核并支持 FP8 等低精度操作。对于训练与推理预填充prefilling任务V1 提供了一组针对非对称域带宽转发如从 NVLink 域转发到 RDMA 域优化的内核——这与 DeepSeek-V3 论文中提出的分组受限门控group-limited gating算法相对应这些内核具备高吞吐特性且支持 SMStreaming Multiprocessor数量控制。对于延迟敏感的推理解码decodingV1 则提供了一套纯 RDMA 的低延迟内核以最小化时延并引入了基于 hook 的通信-计算重叠方法该方法不占用任何 SM 资源。需要特别说明的是V1 基于 NVSHMEM 后端实现属于当前仓库中的存档Legacy版本最新的 V2 已在 README.md 中有完整说明——V2 完成了对专家并行的彻底重构改用更轻量的 NCCL Gin 后端并将高吞吐与低延迟 API 统一到单个ElasticBuffer接口中。V1 文档仍保留在 docs/legacy.md 供历史版本用户参考本文即围绕该 V1 体系展开。同时请注意本库的实现可能与 DeepSeek-V3 论文存在细微差异。从源码结构上看V1 的实现分为三部分Python 封装层deep_ep/buffers/legacy.py 定义了Buffer类负责缓冲区分配、布局计算、dispatch/combine 调用以及低延迟模式管理C 运行时层csrc/legacy/buffer.hpp 实现了通信缓冲区的底层管理NVLink IPC 句柄同步、NVSHMEM 唯一 ID 同步、通信流、barrier 信号与主机侧 MoE 接收计数等内核与配置层csrc/kernels/legacy/ 下的intranode.cu、internode.cu、internode_ll.cu、layout.cu对应机内/机间/低延迟三类内核csrc/legacy/config.hpp 定义了调优配置结构与缓冲区大小估算。2. 性能基准V1 的性能数据在 docs/legacy.md 中给出测试环境为 H800 GPUNVLink 最大带宽约 160 GB/s每卡连接一块 CX7 InfiniBand 400 Gb/s RDMA 网卡最大带宽约 50 GB/s。2.1 普通内核NVLink RDMA 转发测试遵循 DeepSeek-V3/R1 预训练配置每批 4096 token、hidden 维度 7168、top-4 组、top-8 专家、FP8 dispatch BF16 combine。类型Dispatch #EP瓶颈带宽Combine #EP瓶颈带宽机内Intranode8153 GB/s (NVLink)8158 GB/s (NVLink)机间Internode1643 GB/s (RDMA)1643 GB/s (RDMA)机间Internode3258 GB/s (RDMA)3257 GB/s (RDMA)机间Internode6451 GB/s (RDMA)6450 GB/s (RDMA)可以看出机内场景下 NVLink 带宽几乎被饱和利用153~158 GB/s vs 理论约 160 GB/s机间场景受 RDMA 带宽限制43~58 GB/s vs 理论约 50 GB/s。2.2 低延迟内核纯 RDMA测试遵循典型的 DeepSeek-V3/R1 生产配置每批 128 token、hidden 维度 7168、top-8 专家、FP8 dispatch BF16 combine。Dispatch #EP延迟RDMA 带宽Combine #EP延迟RDMA 带宽877 us98 GB/s8114 us127 GB/s16118 us63 GB/s16195 us74 GB/s32155 us48 GB/s32273 us53 GB/s64173 us43 GB/s64314 us46 GB/s128192 us39 GB/s128369 us39 GB/s256194 us39 GB/s256360 us40 GB/s低延迟内核在 #EP8 时可将 dispatch 延迟压到 77 us且 RDMA 带宽利用率极高。需要说明的是以上均为官方文档在特定硬件与配置下测得的数据如 README 所述V2 相较 V1 在峰值性能上有进一步提升同时 SM 占用显著减少但本文仅陈述 V1 文档内的原始数据不做跨版本对比推断。3. 安装与构建3.1 环境需求根据 docs/legacy.mdV1 运行环境要求如下GPUAmpereSM80、HopperSM90GPU或其他支持 SM90 PTX ISA 的架构Python3.8 及以上CUDA 版本SM80 GPU 需要 CUDA 11.0 及以上SM90 GPU 需要 CUDA 12.3 及以上PyTorch2.1 及以上网络机内通信需要 NVLink机间通信需要 RDMA 网络3.2 安装 NVSHMEM 依赖DeepEP V1 依赖 NVSHMEMV2 中 NVSHMEM 仅用于 legacy 方法。完整的安装指引见仓库内文档 docs/nvshmem.md要点如下NVSHMEM 版本要求3.3.9 及以上DeepEP 与上游 NVSHMEM 3.3.9 兼容硬件前提节点内 GPU 需通过 NVLink 互联跨节点需 RDMA 设备并支持 GPUDirect RDMA低延迟内核还需InfiniBand GPUDirect AsyncIBGDA支持可通过 tarballx86_64/aarch64、RPM/deb、conda-forge 或 pipnvidia-nvshmem-cu12安装启用 IBGDA 有两种方式其一修改/etc/modprobe.d/nvidia.conf配置options nvidia NVreg_EnableStreamMemOPs1 NVreg_RegistryDwordsPeerMappingOverride1;后sudo update-initramfs -u并重启其二安装 GDRCopy 并加载gdrdrv内核模块带少量性能开销但在无法修改驱动 regkey 时可用非 RPM/deb 安装时需在 shell 中设置NVSHMEM_DIR、LD_LIBRARY_PATH与PATH验证命令nvshmem-info -a。3.3 开发模式构建# 构建并为 SO 文件建立符号链接 NVSHMEM_DIR/path/to/installed/nvshmem python setup.py build # 可根据自身平台修改具体的 SO 名称 ln -s build/lib.linux-x86_64-cpython-38/deep_ep_cpp.cpython-38-x86_64-linux-gnu.so # 运行测试用例 # 注意可根据集群实际情况修改 init_dist 函数 # 当前仓库中该函数位于 deep_ep/utils/envs.py测试脚本中通过 from deep_ep.utils.envs import init_dist 导入 # 并需在多节点上启动 python tests/test_intranode.py python tests/test_internode.py python tests/test_low_latency.py说明legacy 文档撰写时测试文件位于tests/目录根下当前仓库已将其整理至 tests/legacy/test_intranode.py、test_internode.py、test_low_latency.pyinit_dist的实现也迁移到了 deep_ep/utils/envs.py。构建流程本身不变。从 setup.py 的源码可以看到仅当指定了NVSHMEM_DIR时构建系统才会把csrc/kernels/legacy/internode.cu、csrc/kernels/legacy/internode_ll.cu、csrc/kernels/backend/nvshmem.cu加入编译源文件并链接libnvshmem_device静态库与nvshmem_host动态库这也印证了文档中不指定NVSHMEM_DIR将禁用所有机间与低延迟特性的说明。3.4 安装NVSHMEM_DIR/path/to/installed/nvshmem python setup.py install3.5 安装环境变量环境变量取值含义NVSHMEM_DIR路径NVSHMEM 安装目录不指定则禁用所有机间与低延迟特性DISABLE_SM90_FEATURES0 或 1是否禁用 SM90 特性SM90 设备或 CUDA 11 环境下必须设置TORCH_CUDA_ARCH_LIST架构列表目标架构列表例如TORCH_CUDA_ARCH_LIST9.0DISABLE_AGGRESSIVE_PTX_INSTRS0 或 1是否禁用激进的 load/store 指令详见下文Undefined-behavior PTX 用法源码层面对这些变量的处理位于 setup.py当DISABLE_SM90_FEATURES1时默认把TORCH_CUDA_ARCH_LIST设为8.0优先 A100并向编译命令注入-DDISABLE_SM90_FEATURES禁用 FP8、部分 launch 方式与 TMA否则默认架构为9.0优先 H800 系列并追加-rdctrue --ptxas-options--register-usage-level10。DISABLE_AGGRESSIVE_PTX_INSTRS默认值为 1当前仓库 setup.py 的读取默认即为1且当目标架构不是9.0时会被强制置 1——因为.L1::no_allocate这类指令仅在 SM90 上验证过。4. 网络配置DeepEP 已在 InfiniBand 网络上经过完整测试理论上同样兼容 RoCERDMA over Converged Ethernet。4.1 流量隔离InfiniBand 通过虚拟通道Virtual Lanes, VL支持流量隔离。为防止不同类型流量相互干扰官方建议将以下负载划分到不同虚拟通道使用普通内核的负载使用低延迟内核的负载其他负载对于 DeepEP V1可通过设置NVSHMEM_IB_SL环境变量控制虚拟通道分配。从 deep_ep/buffers/legacy.py 的Buffer.__init__可以看到低延迟模式下库还会主动设置一批 NVSHMEM 环境变量以适配通信需求包括NVSHMEM_IB_ENABLE_IBGDA1强制启用 IBGDANVSHMEM_IBGDA_NUM_RC_PER_PE设置每个对端的 RC可靠连接QP 数量即构造参数num_qps_per_rankNVSHMEM_QP_DEPTH默认 1024保证 QP 深度始终大于在途 WR 数量从而跳过 WQ 槽位检查NVSHMEM_MAX_TEAMS7与NVSHMEM_DISABLE_NVLS1减少 GPU 显存占用、禁用 NVLink SHARPNVSHMEM_CUMEM_GRANULARITY2^29NVSHMEM 初始化至少需要 256 MiB 的分配粒度。4.2 自适应路由自适应路由adaptive routing是 InfiniBand 交换机提供的进阶路由特性可将流量均匀分布到多条路径上。启用它可以完全消除路由冲突导致的网络拥塞但会引入额外延迟。官方建议重负载环境启用自适应路由轻负载环境使用静态路由4.3 拥塞控制官方在生产环境中未观察到显著拥塞因此拥塞控制保持禁用。5. 接口实战一模型训练与推理预填充——普通内核普通内核适用于模型训练或推理预填充阶段不含反向部分其核心接口是 deep_ep/buffers/legacy.py 中的Buffer类。下面给出文档中的完整示例。5.1 完整示例代码import torch import torch.distributed as dist from typing import List, Tuple, Optional, Union from deep_ep import Buffer, EventOverlap # 通信缓冲区运行时分配 _buffer: Optional[Buffer] None # 设置使用的 SM 数量 # 注意这是一个静态变量 Buffer.set_num_sms(24) # 可在框架初始化时调用该函数 def get_buffer(group: dist.ProcessGroup, hidden_bytes: int) - Buffer: global _buffer # 注意也可以用所有测试得到的自动调优结果替换 get_*_config num_nvl_bytes, num_rdma_bytes 0, 0 for config in (Buffer.get_dispatch_config(group.size()), Buffer.get_combine_config(group.size())): num_nvl_bytes max(config.get_nvl_buffer_size_hint(hidden_bytes, group.size()), num_nvl_bytes) num_rdma_bytes max(config.get_rdma_buffer_size_hint(hidden_bytes, group.size()), num_rdma_bytes) # 若缓冲区不存在或大小不足则重新分配 if _buffer is None or _buffer.group ! group or _buffer.num_nvl_bytes num_nvl_bytes or _buffer.num_rdma_bytes num_rdma_bytes: _buffer Buffer(group, num_nvl_bytes, num_rdma_bytes) return _buffer def get_hidden_bytes(x: torch.Tensor) - int: t x[0] if isinstance(x, tuple) else x return t.size(1) * max(t.element_size(), 2) 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]], torch.Tensor, torch.Tensor, List, Tuple, EventOverlap]: # 注意可选的 previous_event 是希望作为 dispatch 内核依赖的 CUDA 事件 # 可用于通信-计算重叠更多信息请参考 Buffer.dispatch 的文档 global _buffer # 在实际 dispatch 前计算布局 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_eventprevious_event, async_finishTrue, allocate_on_comm_streamprevious_event is not None) # 执行 MoE dispatch # 注意CPU 会等待 GPU 的信号到达因此此方式与 CUDA graph 不兼容 # 除非指定 num_worst_tokens但该标志仅适用于机内场景 recv_x, recv_topk_idx, recv_topk_weights, num_recv_tokens_per_expert_list, handle, event \ _buffer.dispatch(x, topk_idxtopk_idx, topk_weightstopk_weights, num_tokens_per_ranknum_tokens_per_rank, num_tokens_per_rdma_ranknum_tokens_per_rdma_rank, is_token_in_rankis_token_in_rank, num_tokens_per_expertnum_tokens_per_expert, previous_eventprevious_event, async_finishTrue, allocate_on_comm_streamTrue) # 事件管理请参考 EventOverlap 类的文档 return recv_x, recv_topk_idx, recv_topk_weights, num_recv_tokens_per_expert_list, handle, event def dispatch_backward(grad_recv_x: torch.Tensor, grad_recv_topk_weights: torch.Tensor, handle: Tuple) - \ Tuple[torch.Tensor, torch.Tensor, EventOverlap]: global _buffer # MoE dispatch 的反向过程实际上是一次 combine combined_grad_x, combined_grad_recv_topk_weights, event \ _buffer.combine(grad_recv_x, handle, topk_weightsgrad_recv_topk_weights, async_finishTrue) return combined_grad_x, combined_grad_recv_topk_weights, event def combine_forward(x: torch.Tensor, handle: Tuple, previous_event: Optional[EventOverlap] None) - \ Tuple[torch.Tensor, EventOverlap]: global _buffer # 执行 MoE combine combined_x, _, event _buffer.combine(x, handle, async_finishTrue, previous_eventprevious_event, allocate_on_comm_streamprevious_event is not None) return combined_x, event def combine_backward(grad_combined_x: Union[torch.Tensor, Tuple[torch.Tensor, torch.Tensor]], handle: Tuple, previous_event: Optional[EventOverlap] None) - \ Tuple[Union[torch.Tensor, Tuple[torch.Tensor, torch.Tensor]], EventOverlap]: global _buffer # MoE combine 的反向过程实际上是一次 dispatch grad_x, _, _, _, _, event _buffer.dispatch(grad_combined_x, handlehandle, async_finishTrue, previous_eventprevious_event, allocate_on_comm_streamprevious_event is not None) return grad_x, event此外在 dispatch 函数内部我们可能无法预先知道当前 rank 需要接收多少 token因此会涉及一次隐式的 CPU 等待 GPU 接收计数信号的过程如下图所示。5.2 关键机制与源码解读1布局计算先行。get_dispatch_layout在真正通信之前根据topk_idx计算路由元数据返回四类张量num_tokens_per_rank发往各 rank 的 token 数、num_tokens_per_rdma_rank发往各 RDMA rank 的 token 数机内场景为None、num_tokens_per_expert发往各专家的 token 数、is_token_in_ranktoken 是否发往某 rank 的布尔掩码。对应实现在 deep_ep/buffers/legacy.py测试 tests/legacy/test_intranode.py 中会逐一校验这些元数据与手工计算一致。2SM 数量为静态变量且必须为偶数。示例中Buffer.set_num_sms(24)将高吞吐内核使用的 SM 数设为 24源码中该静态变量的默认值为 20且set_num_sms断言新值必须能被 2 整除deep_ep/buffers/legacy.py。原因是内核按通道组织通道数 num_sms / 2见 csrc/legacy/config.hpp。3Config 决定吞吐。Buffer.get_dispatch_config(group.size())与get_combine_config(group.size())针对不同 EP rank 数返回内置调优配置。以 dispatch 为例deep_ep/buffers/legacy.py 中的配置表覆盖 2~160 个 rank如 8 rank 时为Config(20, 6, 256, 6, 128)128 rank 时为Config(20, 20, 560, 12, 128)。这些数值对应 csrc/legacy/config.hpp 中Config结构的五个字段num_sms、NVLink 分块发送/接收 token 数、RDMA 分块发送/接收 token 数。get_nvl_buffer_size_hint与get_rdma_buffer_size_hint则按通道数、对端数与 token 上限估算通信缓冲区大小。4handle 复用与 CUDA graph。dispatch返回的handle记录了前缀矩阵、源索引等路由元数据再次调用时传入handle即可跳过布局重算测试中的cached dispatch路径。dispatch的num_worst_tokens参数若指定了最坏接收 token 数则不再需要 CPU 同步从而兼容 CUDA graph但该参数仅对机内内核有效——机间路径会直接断言num_worst_tokens 0见 deep_ep/buffers/legacy.py。5事件与重叠。EventOverlap是 CUDA 事件的轻量封装deep_ep/utils/event.pycapture()捕获当前流事件current_stream_wait()让当前计算流等待通信完成配合async_finishTrue与allocate_on_comm_streamTrue可实现通信与计算的重叠示例代码最后用event.current_stream_wait()消费结果。该封装同时用extra_tensors模拟record_stream语义避免与 CUDA graph 冲突。6FP8 支持。dispatch的输入x可以是torch.bfloat16张量也可以是(fp8_e4m3fn 数据, float 缩放因子)元组缩放因子形状为[num_tokens, hidden // 128]combine 则以 BF16 进行加权归约。测试脚本中用per_token_cast_to_fp8/per_token_cast_back构造并还原 FP8 数据以验证正确性。6. 接口实战二推理解码——低延迟内核低延迟内核适用于推理解码阶段基于 IBGDA 纯 RDMA 实现示例代码如下。6.1 完整示例代码import torch import torch.distributed as dist from typing import Tuple, Optional from deep_ep import Buffer # 通信缓冲区运行时分配 # 注意低延迟内核没有 SM 控制 API _buffer: Optional[Buffer] None # 可在框架初始化时调用该函数 def get_buffer(group: dist.ProcessGroup, num_max_dispatch_tokens_per_rank: int, hidden: int, num_experts: int) - Buffer: # 注意低延迟模式比普通模式消耗更多显存 # 因此建议 num_max_dispatch_tokens_per_rank解码引擎的实际 batch size小于 256 global _buffer num_rdma_bytes Buffer.get_low_latency_rdma_size_hint(num_max_dispatch_tokens_per_rank, hidden, group.size(), num_experts) # 若缓冲区不存在或大小不足则重新分配 if _buffer is None or _buffer.group ! group or not _buffer.low_latency_mode or _buffer.num_rdma_bytes num_rdma_bytes: # 注意为了获得最佳性能QP 数量**必须**等于本地专家数 assert num_experts % group.size() 0 _buffer Buffer(group, 0, num_rdma_bytes, low_latency_modeTrue, num_qps_per_ranknum_experts // group.size()) return _buffer def low_latency_dispatch(hidden_states: torch.Tensor, topk_idx: torch.Tensor, num_max_dispatch_tokens_per_rank: int, num_experts: int): global _buffer # 执行 MoE dispatch兼容 CUDA graph但重放后可能需要恢复部分缓冲区状态 recv_hidden_states, recv_expert_count, handle, event, hook \ _buffer.low_latency_dispatch(hidden_states, topk_idx, num_max_dispatch_tokens_per_rank, num_experts, async_finishFalse, return_recv_hookTrue) # 注意只有调用 hook() 后实际张量才会真正到达 # 这对双 batch 重叠非常有用且**不占用任何 SM** # 如果不想重叠请设置 return_recv_hookFalse # 之后可以使用我们的 GEMM 库按该特定格式进行计算 return recv_hidden_states, recv_expert_count, handle, event, hook def low_latency_combine(hidden_states: torch.Tensor, topk_idx: torch.Tensor, topk_weights: torch.Tensor, handle: Tuple): global _buffer # 执行 MoE combine兼容 CUDA graph但重放后可能需要恢复部分缓冲区状态 combined_hidden_states, event_overlap, hook \ _buffer.low_latency_combine(hidden_states, topk_idx, topk_weights, handle, async_finishFalse, return_recv_hookTrue) # 注意行为与 dispatch 内核中描述的一致 return combined_hidden_states, event_overlap, hook对于双微批two-micro-batch重叠可参考下图。借助接收 hook 接口RDMA 网络流量在后台进行不消耗计算部分任何 GPU SM。但需要注意重叠的各段可以调整——attention / dispatch / MoE / combine 四段并不一定执行时间完全相同你可以根据负载调整阶段设置。6.2 IBGDA 与缓冲区布局低延迟内核要求所有 rank无论机内还是机间都通过 RDMA 可见必须启用 IBGDA。在Buffer.__init__中只要num_rdma_ranks 1或处于低延迟模式就会自动设置NVSHMEM_IB_ENABLE_IBGDA1、NVSHMEM_IBGDA_NUM_RC_PER_PEnum_qps_per_rank等环境变量deep_ep/buffers/legacy.py。缓冲区布局由 csrc/legacy/config.hpp 的LowLatencyLayout计算RDMA 缓冲区被划分为两套对称的 odd/even 结构对应双缓冲每套包含发送缓冲区、接收缓冲区与信令缓冲区消息大小按 dispatch 消息int4控制头 hidden 数据或 FP8 数据缩放因子与 combine 消息每 128 通道的 min/max 头 hidden 数据分别计算。Buffer.get_low_latency_rdma_size_hint即据此给出最低 RDMA 缓冲区需求。由于只有两套缓冲同一时刻最多只能持有 2 个低延迟内核的结果张量否则结果会被覆盖。6.3 接收 hook 与零 SM 重叠return_recv_hookTrue时低延迟内核只发出 RDMA 请求而不真正接收数据调用返回的hook()才保证数据到达。这一设计使得网络传输可以在后台进行计算核如 GEMM无需让出 SM。注意低延迟内核没有 SM 控制 API这是与普通内核的重要区别。低延迟 dispatch 还支持以下进阶参数源码见 deep_ep/buffers/legacy.pyuse_fp8启用 FP8 转换接收数据为(FP8 张量, 缩放因子)元组形状为[num_local_experts, num_max_dispatch_tokens_per_rank * num_ranks, hidden]round_scale/use_ue8m0将缩放因子舍入为 2 的幂 / 使用 UE8M0 打包格式需与round_scaleTrue搭配cumulative_local_expert_recv_stats与dispatch_wait_recv_cost_stats累计统计张量前者用于在线服务 EP 负载均衡监控后者用于定位慢 rank 异常topk_idx中-1表示该位置未选择任何专家是合法输入。6.4 缓冲区清理与容错mask/shrink由于低延迟内核要求部分缓冲区零初始化当缓冲区变脏时例如先运行了普通 dispatch/combine必须在执行低延迟内核前调用buffer.clean_low_latency_buffer(num_max_dispatch_tokens_per_rank, hidden, num_experts)清理deep_ep/buffers/legacy.py。另外构造Buffer(..., enable_shrinkTrue)可开启 shrink收缩模式配合low_latency_update_mask_buffer(rank, mask)、low_latency_query_mask_buffer、low_latency_clean_mask_buffer可动态屏蔽/恢复某个 rank 的通信用于在线容错。测试 tests/legacy/test_low_latency.py 的--shrink-test模式会模拟 dispatch/combine/clean 阶段分别有 rank 掉线并校验 mask 缓冲状态与结果正确性。7. 测试与自动调优7.1 测试脚本与命令行参数V1 的三个测试脚本位于 tests/legacy/均基于torch.multiprocessing.spawn启动多进程分布式环境test_intranode.py机内内核主要参数参数默认值含义--num-processes8进程数默认 8 卡--num-tokens4096token 数--hidden7168hidden 维度--num-topk8top-k 专家数--num-experts256专家总数--allow-mnnvl-启用 MNNVL 支持test_internode.py机间内核额外提供--num-topk-groups默认min(num_nodes, 4)对应分组受限门控的组数、--pressure-test-mode0/1/2不做压力测试/做压力测试但不基准/做压力测试且带基准与--test-ll-compatibility与低延迟内核兼容性测试。机间测试要求num_local_ranks 8且num_ranks 8即每节点 8 卡并跨多节点。test_low_latency.py低延迟内核参数默认--num-tokens 128、--num-experts 288并提供--use-logfmt测试 LogFMT combine、--disable-nvlink禁用 NVLink、--pressure-test、--shrink-test等开关。7.2 调优流程解读测试脚本不仅验证正确性还内置了自动调优逻辑以test_internode.py为例dispatch 阶段在 NVL chunk4~44 步长 4与 RDMA chunk4~32 步长 4的网格上穷举Config用 Kineto 记录真实耗时并选出最优组合再由 rank 0 通过all_gather广播最优配置combine 阶段类似地在 NVL chunk 1~7、RDMA chunk 网格上搜索。文档建议为在你的集群上获得更优性能应运行全部测试并采用调优出的最佳配置——内置默认配置是针对 DeepSeek 内部集群优化的不一定适配你的网络拓扑。8. Roadmap 与注意事项V18.1 RoadmapV1 路线图勾选表示已完成如下AR 支持重构低延迟模式 AR 代码A100 支持仅机内低延迟 dispatch 内核支持 BF16机内低延迟内核支持 NVLink 协议使用 TMA 拷贝替代 LD/ST机内内核机间内核低延迟内核SM-free 内核与重构完全移除未定义行为的 PTX 指令8.2 更简单的潜在整体设计V1 实现使用队列管理通信缓冲区这能节省显存但引入了复杂性与潜在死锁。如果你要基于 DeepEP V1 自行实现文档建议考虑使用按最大容量分配的固定大小缓冲区换取简单性与更优性能该替代方案的详细讨论见 DeepEP 仓库 issue 39。8.3 Undefined-behavior PTX 用法为追求极致性能V1 使用了一处未定义行为的 PTX 用法用只读 PTX 指令ld.global.nc.L1::no_allocate.L2::256B读取易变volatile数据。PTX 修饰符.nc表示使用非相干缓存。虽然在 Hopper 架构上配合.L1::no_allocate实测正确性有保障且性能显著更好但这是通过实测验证而非架构规范保证的。作者推测原因非相干缓存与 L1 是统一的而.L1::no_allocate不只是提示而是强选项因此 L1 中无脏数据可保证正确性。最初由于 NVCC 无法自动展开 volatile 读取的 PTX作者尝试了__ldg即ld.nc。即使与手工展开的 volatile 读取相比__ldg也明显更快可能得益于额外的编译器优化但结果可能错误或读到脏数据。查阅 PTX 文档后作者发现 Hopper 架构上 L1 与非相干缓存统一推测.L1::no_allocate可解决问题最终促成这一发现。如果你在其他平台上发现内核无法正常工作可以在setup.py构建时设置DISABLE_AGGRESSIVE_PTX_INSTRS1禁用该特性或提交 issue 反馈。8.4 集群自动调优建议为在自有集群上获得更好性能建议运行 tests/legacy/ 下全部测试并采用最优自动调优配置内置默认配置针对 DeepSeek 内部集群优化直接使用未必最优。此外机内与低延迟内核混用时务必在切换前调用clean_low_latency_buffer低延迟模式建议将解码引擎 batch sizenum_max_dispatch_tokens_per_rank控制在 256 以内以控制显存开销。结语DeepEP V1 以 NVSHMEM 为后端为 MoE 专家并行提供了机内/机间的高吞吐 dispatch/combine 与纯 RDMA 的低延迟内核两套方案并通过 SM 控制、FP8 低精度、事件重叠与接收 hook 等机制覆盖了训练、预填充与解码全链路。本文以 docs/legacy.md 为主体结合 deep_ep/buffers/legacy.py、csrc/legacy/config.hpp、csrc/legacy/buffer.hpp、setup.py 与 tests/legacy/ 测试代码对接口语义与底层原理做了展开。V2 用户请以 README.md 为准而依赖 V1NVSHMEM的存量系统仍可参考本文完成安装、调优与集成。【免费下载链接】DeepEPDeepEP: an efficient expert-parallel communication library项目地址: https://gitcode.com/GitHub_Trending/de/DeepEP创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考