DeepSeek 的硬件优化基因

一、DeepSeek 再度登场,但这次冲击的性质不同

2026 年 4 月 24 日,DeepSeek 公开了 V4 预览版。1.6 万亿参数的 V4 Pro 与 2,840 亿参数的 V4 Flash 同时发布,两款模型都标榜 100 万 token 上下文窗口与原生多模态能力。如果只是一次例行的后续模型发布,它无法重现一年前 R1 让美国大型科技公司市值蒸发约 1 万亿美元的那种冲击。市场早已学会了一个事实——中国 AI 既便宜又强大。

因此,V4 此次的冲击点不在模型本身,而在别处。据报道,DeepSeek 优先在华为昇腾(Ascend)950 系列与寒武纪(Cambricon)芯片上对 V4 进行了优化,并刻意没有向英伟达和 AMD 提供同样的早期优化接入窗口。

一年前还在证明"用被阉割的 H800 也能做出前沿模型"这一命题的公司,如今已经更进一步,宣告:完全可以不需要英伟达。

DeepSeek 的真正资产,不是模型权重,而是榨干被制裁硬件的工程文化本身。而这种文化的根,不在 AI 实验室,而在对冲基金的交易室里。

二、HFT 的基因——为什么 DeepSeek 是一家"连内核都要削"的公司

DeepSeek 的母公司是杭州的量化对冲基金幻方(High-Flyer)。其资产管理规模据称约为 80 亿美元,创始人梁文锋同时担任两家公司的 CEO。这一出身背景不仅仅意味着"资金来源",其意义更深——这家公司的 DNA 就是不一样的。

高频交易(HFT)与量化交易,本质上是"以微秒为单位削减延迟的产业"。算法的优雅程度,远不如能否削掉最后一个时钟周期重要。网卡固件、操作系统内核的中断处理、NUMA 节点之间的内存访问模式,甚至数据中心内部的网线长度,全都被视作变量。HFT 交易员问的第一个问题不是"创意",而是"每瓦特能输出多少"。

这种文化原封不动地被移植到 LLM 团队,是 DeepSeek 与其他大型实验室的决定性差异。一般 AI 研究室的出发点是"再多买点 GPU 就行了",而 DeepSeek 的出发点是"如何从手上已有的 GPU 再榨出最后 1%"。即便用同样的 H800 集群,HFT 出身的团队会下沉到 PTX/SASS 级别的指令调度,刻意使用虽然非确定性但更快的 PTX 指令,甚至下沉到 warp 级别的流水线。这不只是"工程做得好"那种程度的故事——其成本函数本身就不一样。

对 DeepSeek 来说,"削硬件"不是美德,而是默认值。这是本文要强调的第一个命题。

三、连内核都要削的团队功力——开源基础设施全家桶

DeepSeek 以 2025 年的"开源周"为起点,陆续公开了自家训练与推理基础设施的核心组件。值得注意的是,这套组合呈现出来的不是"只把模型做好的公司",而是"把整个训练系统垂直整合的公司"。下表整理了核心组件。

项目	功能领域	核心价值
DeepEP	MoE all-to-all 通信库	在 NVLink/RDMA 不对称带宽环境下消除专家并行的通信瓶颈。支持 FP8 dispatch 与不占用 SM 的 hook 式通信-计算重叠。
DeepGEMM	FP8 GEMM 内核	同时支持 Hopper 与 Blackwell 的简洁 BLAS 内核。通过细粒度缩放与 JIT 编译,把效率推到同类库的极限。
FlashMLA	Multi-head Latent Attention 内核	在 H800 SXM5 上把内存带宽推到接近理论极限的注意力内核。把推理成本压缩了一个数量级。
DualPipe	双向流水线并行	通过把前向/反向计算与通信交错排列消除流水线气泡。让大型 MoE 训练可以不依赖张量并行的核心算法。
3FS	分布式文件系统	专为 AI 训练与推理工作负载设计的高性能分布式存储。把"不让数据流水线本身成为瓶颈"做到了基础设施层。
TileKernels / EPLB	基于 tile 的内核与专家负载均衡	基于 tilelang 的内核库,以及 MoE 专家负载均衡算法。一家公司通过垂直整合承担整条训练流水线。

这套组合传达的信息很清楚:DeepSeek 把通信库(DeepEP)、矩阵运算内核(DeepGEMM)、注意力内核(FlashMLA)、流水线算法(DualPipe)、分布式文件系统(3FS),全部由一家公司同时运营。即使是美国的科技巨头,在这些层面也仍大量依赖外部库(NCCL、cuBLAS、FlashAttention 等),与之形成鲜明对比。借用半导体行业的说法,DeepSeek 更像是"从晶圆厂到封装一体化运营的 AI 公司"——一家 AI 界的 IDM。

四、把算法和硬件放在一起设计——把约束当作武器

DeepSeek 的真正差异化,不是某一两项单独的技术,而是它在"同一间办公室里"设计模型架构和硬件约束。Multi-head Latent Attention(MLA)就是典型——通过把 KV 缓存压到 5%~13% 的水平,使 H800 80GB 的 HBM 上能跑更大的批量。从 V3.2 开始引入的 DeepSeek Sparse Attention(DSA),使用一个名为"闪电索引器(Lightning Indexer)"的轻量打分器先筛选 token,再只对 top-K 真正执行注意力。它把随序列长度二次膨胀的注意力开销,压到了接近线性的水平。

FP8 训练也是同一个脉络。DeepSeek 承认 H800 张量核心的 FP8 累加精度只有约 14 位,然后通过每隔一定区间(约 128 个元素)将部分和提升到 CUDA 核心的 FP32 寄存器中累加,绕开了这一限制。

所有这些决定的共同点很清楚——当其他大型实验室还在按"先做模型,再去买硬件"的顺序行事时,DeepSeek 从相反的方向出发:把模型的形状贴合我们已有的芯片。这源于对硬件的深刻理解,以及 HFT 行业把性命押在硬件调优上的传统所积累出来的"常识"——一种把成本与性能比压到最优的常识。

五、【深度解析】DeepEP 解剖——它如何解决了 MoE 时代的通信瓶颈?

从 CUDA 一直追到 NVSHMEM——一个库所传递的信息分量

上一章我们用一句话总结了 DeepEP——"MoE all-to-all 通信库"。要理解这一句话为什么对整个产业具有破坏力,我们必须反过来,从 CUDA 开始一层一层往下追。本章在技术层面有些深入,但结论可以浓缩为一句话:

只发布模型的时代结束了。我们已经进入连基础设施内核都要做最优化的时代。

5.1 重新审视 GPU 通信——CUDA、NVLink、RDMA

CUDA 既是写运行在 NVIDIA GPU 上代码的语言,也是中间件。其语法类似 C/C++,但增加了直接驱动 GPU 上数千核心的扩展。为什么非要用 CUDA?原因很简单——普通 CPU 代码无法同时驱动 GPU 上的数千个核心。我们日常使用 PyTorch、TensorFlow,但拆开它们,底层调用的归根结底就是 CUDA 内核。矩阵乘法、注意力、归一化——都是某个人写的 CUDA 内核。

模型小,一张 GPU 就够。但今天的 LLM 早已破坏了这个假设。几百 GB 大小的模型放不进单张 GPU 的 HBM(高带宽显存)。最终必须在多张 GPU 之间拆分,而一旦拆分,GPU 之间就需要交换数据。这就是 GPU 通信。

GPU 通信分两个层级:同一台服务器内部用 NVLink,服务器之间用 InfiniBand RDMA。打个比方:NVLink 是市内电话,RDMA 是长途电话。两者都很快,但 NVLink 远比 RDMA 快。而这种通信本身,最终也是由 CUDA 内核来处理的。换句话说,"好的通信内核"直接等同于"快速的训练和推理"。通信内核占用多少 GPU 的 SM(流式多处理器)资源,就直接对应模型训练的速度。

5.2 MoE 真正的瓶颈——all-to-all 通信

MoE(Mixture of Experts)是只激活模型一部分"专家"的结构。每个 token,路由器都要决定"这个 token 应该送到哪个专家"。

问题在于,这个路由发生在每一对 GPU 之间。

8 张 GPU 就是每一步都触发 8 对 8 通信,64 张就是 64 对 64。这就是 all-to-all 通信。在 MoE 中,它分两个阶段:把 token 发出去的 dispatch,以及把结果收回来的 combine。

随着 MoE 模型变大,这种通信就变成了真正的瓶颈。不是算力不够——是通信卡住了你的脖子。在像 H800 这种 NVLink 带宽被砍到 H100 一半以下的"合规规格"芯片上,这个瓶颈是决定性的。DeepSeek 在训练自家 V3 和 R1 模型的过程中,直接对这一层做了优化,并把那份成果原封不动地开源了出来——这就是 DeepEP。等于把自家模型训练用的代码原封不动地公开了。

DeepEP 才是 DeepSeek 团队真正掏出来的"武林秘籍"。

一句话总结,DeepEP 的身份就是:MoE 专用 GPU 通信内核合集——用 CUDA 直接写的 dispatch/combine 内核。

5.3 两种内核——Normal 和 Low-Latency

DeepEP 分为两类内核:用于训练和推理 prefill 阶段的 Normal Kernel,以及用于推理 decoding 阶段的 Low-Latency Kernel。两种工作负载的特性完全相反,因此同一个库里实际上塞着两台不同的通信引擎。

【Normal Kernel】 训练和 prefill 一次处理一个大批量(官方基准是 4,096 token)。这时关键是同时驱动 NVLink 和 RDMA。DeepEP 把节点内通信走 NVLink、节点间通信走 RDMA,把两种带宽不对称地合在一起用。在 H800 + CX7 InfiniBand 400 Gb/s 环境下,官方测得的数据如下:

通信类型	Dispatch #EP	瓶颈带宽 (Dispatch)	Combine #EP	瓶颈带宽 (Combine)
Intranode	8	153 GB/s (NVLink)	8	158 GB/s (NVLink)
Internode	16	43 GB/s (RDMA)	16	43 GB/s (RDMA)
Internode	32	58 GB/s (RDMA)	32	57 GB/s (RDMA)
Internode	64	51 GB/s (RDMA)	64	50 GB/s (RDMA)

基准环境为 H800(NVLink 最大带宽约 160 GB/s),CX7 InfiniBand 400 Gb/s NIC(最大带宽约 50 GB/s),遵循 DeepSeek-V3/R1 预训练设置(每批 4,096 token,hidden 7,168,top-4 组 / top-8 专家,FP8 dispatch + BF16 combine)。

节点内 8EP 上 NVLink 153 GB/s——相当于理论值 160 GB/s 的 96%,可以说"已经贴在理论极限附近"。节点间 32EP 的 RDMA 58 GB/s 也超过了 50 GB/s 的标称上限,这是因为它把 NVLink 域的流量通过 RDMA 转发出去,以不对称方式合并了两种带宽。

📌 2025 年 4 月 22 日,腾讯网络平台部的优化补丁(PR #130)合入主分支,据报告在某些配置下带来高达 30% 的额外性能提升。一家科技巨头的通信团队主动给另一家公司的库提交补丁,这件事本身就说明 DeepEP 已经在事实上成为中国国内的标准。

【Low-Latency Kernel】 推理 decoding 是相反的工作负载。一次处理的批量很小(官方设置是 128 token),响应时间就是用户体验本身。有趣的是,这个模式完全不用 NVLink,只用纯 RDMA。原因是:NVLink 的建立开销,比 decoding 一个 token 的处理时间还大。当市内电话的接通时间比通话本身更长,直接打长途反而更快——这是个有点反直觉的悖论。官方基准如下:

Dispatch #EP	Latency	RDMA 带宽	Combine #EP	Latency	带宽
8	77 μs	98 GB/s	8	114 μs	127 GB/s
16	118 μs	63 GB/s	16	195 μs	74 GB/s
32	155 μs	48 GB/s	32	273 μs	53 GB/s
64	173 μs	43 GB/s	64	314 μs	46 GB/s
128	192 μs	39 GB/s	128	369 μs	39 GB/s
256	194 μs	39 GB/s	256	360 μs	40 GB/s

8EP 时 dispatch 77 μs、combine 114 μs 的数字意义不容小觑。即便扩展到 256EP,延迟也没有爆炸——dispatch 仍然只有 194 μs、combine 360 μs。对推理服务系统而言,这意味着一个承诺:专家数量再多,用户感知到的延迟几乎是恒定的。 2025 年 6 月 5 日的补丁(PR #173)进一步改进了低延迟内核,使其尽可能利用 NVLink。

5.4 基于 Hook 的通信-计算重叠与 FP8 dispatch

DeepEP 中最有意思的设计,是基于 hook 的通信-计算重叠。通常的通信库为了处理通信,会占用 GPU 的一部分 SM——能用于计算的 SM 就相应减少。DeepEP 的低延迟内核把 RDMA 流量挤到后台,把 SM 占用率压到零。它通过 hook 接口(return_recv_hook=True)让调用方直接控制"接收完成"的时点。

通过这一机制,把 Attention → Dispatch → MoE 计算 → Combine 这四段在两个 micro-batch 之间交错排布,GPU 的空闲气泡(idle bubble)实际上就收敛到了零。更重要的副作用是 CUDA Graph 兼容性。在推理服务里,CUDA Graph 是把每个 token 的开销削到个位数微秒级别的关键功能,而绝大多数通信库不支持它。DeepEP 的低延迟模式支持。

还有一点,DeepEP 原生支持 FP8 dispatch:发送 token 时压成 FP8,合并结果时用 BF16。精度保持在 BF16 水平的同时,通信量几乎砍掉一半。 DeepSeek-V3 论文中提出的 group-limited gating 算法也原封不动地集成在内。可以清楚地看到,他们几乎没改动自家 V3/R1 训练代码就直接公开了出来。

5.5 使用环境、依赖与门槛

DeepEP 的环境要求如下。支持的 GPU 是 Ampere(SM80)与 Hopper(SM90),官方测试基于 H800。软件层面,SM80 需 CUDA 11+,SM90 需 CUDA 12.3+,PyTorch 2.1+,Python 3.8+。

节点内通信前提是 NVLink,节点间通信前提是 RDMA。推荐网络是 InfiniBand CX7 400 Gb/s,RoCE(RDMA over Converged Ethernet)在文档里被列为"理论上兼容"——但这个"理论上"的尾巴值得注意。

门槛非常清晰:NVSHMEM 依赖。 NVSHMEM 是 NVIDIA 用于节点间通信的 SHMEM 库,安装它不是一行 pip install 能解决的,需要按照单独的构建指南来。DeepEP 仓库甚至专门附带了一份 NVSHMEM 安装指南文档。许可证是 MIT,但部分引用 NVSHMEM 的文件(csrc/kernels/ibgda_device.cuh、third-party/nvshmem.patch)受 NVIDIA 的 NVSHMEM SLA 约束。引入前需要做许可证审查。

实际使用方式出乎意料地简单:创建 Buffer 对象,调用 dispatch() 把 token 发给专家,跑 MoE 计算,再用 combine() 收回来。低延迟模式在创建 Buffer 时通过 low_latency_mode=True 标志分支,加入后台 hook。接口本身对 PyTorch 用户来说,适应起来不算难。

5.6 DeepEP 真正传递的信息

从这里开始,才是本文真正要谈的。DeepEP 的意义,远不止"快的通信库"这一层。

第一, MoE 训练已经成为离开通信优化就无法运转的领域。光靠普通的 PyTorch DDP 或 DeepSpeed,已经无法高效驱动几百张 GPU 规模的 MoE。专用内核事实上已成必需。DeepEP 正在拿下事实上的标准位置。
第二,DeepSeek 已经从"只发布模型的公司",升级到连基础设施开销库都一起发布的公司。这一举动,既能压低自家训练成本,也能诱导其他团队在做类似 MoE 模型时直接构建在自家代码基础之上——这是一次标准制定权的争夺。
第三, AMD ROCm 分支 Mori-EP 已经作为官方分支存在;UCCL 项目的 uccl-ep 分支支持异构 GPU(NVIDIA、AMD)与多种 NIC(EFA、Broadcom、CX7);蚂蚁集团在另一个分支上运营 SM-Free 优化系列;NVIDIA 自己也通过 Hybrid-EP 分支在试验 TMA 指令与 NVFP4 数据类型支持。
第四,Infrawaves 增加了双端口 NIC 与多 QP 支持,蚂蚁集团的 DeepXTrace 是诊断慢节点的分析工具。一个库周围,小型生态已经开始成形。这是"有人正在拆解 NVIDIA 锁定"的一个明确信号。让我们把这幅图带回韩国的现实。如果是在国内做 MoE 模型训练的团队,现在最先要评估的库就是DeepEP。尤其是那些拥有 H100/H200 集群的单位,引入成本对比训练效率的改进幅度,是当下就可量化的。如果是做推理服务的团队,光是把 Low-Latency Kernel 单独拎出来评估,就有它的价值——它是降低 decoding 延迟最快的路径之一。

缺点也很清楚:NVSHMEM 依赖、H800 之外环境的验证不足,以及在普通数据中心以太网上 RoCE 的工作情况需要单独验证。但这套代码本身的构成是 CUDA 58.9%、Python 20.3%、C++ 19.2%——足够"开放",可以直接读、直接改。如果韩国团队想尝试给自家 NPU(Rebellions、FuriosaAI 等)接后端,Mori-EP 或 UCCL-EP 这样的分支会是一个不错的起点。

只发布模型的时代结束了。现在是连基础设施内核都开源的时代。而推动这股潮流的,不是美国科技巨头,而是从被制裁的 H800 起步的杭州对冲基金团队的 DNA——这才是本文的核心要点。

六、美国制裁与 DeepSeek 的回应——三阶段战略

美国对华 AI 芯片出口管制,经过 2022 年 10 月的第一轮、2023 年 10 月的强化,以及 2024–2025 年的多轮追加,已逐渐成形。其核心是阻止前沿 GPU(H100、B200 系列)和高带宽互连出口到中国。H800 原本是为绕开监管而刻意降低 NVLink 带宽的"中国专供版",但 2023 年 10 月的扩大管制,最终连这条路也封死了。DeepSeek 的应对,可以总结为以下三个阶段。

6.1 第一阶段——用算法补硬件差距

V3、R1 阶段属于这一层。前面分析过的 FP8 训练、DualPipe、MLA,以及刚刚剖析完的 DeepEP,都是在"用受限硬件做出非受限硬件级别性能"这一明确目标下设计的。尤其是 DeepEP "把 H800 狭窄的 NVLink 与相对充裕的 RDMA 不对称合并"的设计,是把"由制裁塑造出来的硬件形态"本身作为算法输入变量的最清晰案例。这一阶段传达给美国决策者的第一个信号是:监管无法永远维持压倒性的硬件代差。

6.2 第二阶段——向中国国产芯片的"刻意"迁移

V4 中政治意义最重大的决定,是优先在华为昇腾 950 系列与寒武纪芯片上进行优化。据报道,DeepSeek 把 V4 的预优化接入窗口优先开放给中国芯片厂商,而对英伟达和 AMD 关闭了同样的窗口。中国媒体援引英伟达 CEO 黄仁勋的说法,称这件事 "可能会成为美国的灾难"。技术上仍存在差距。分析师把昇腾 950 评估在 H100 与 H200 之间,而中芯国际(SMIC)的 7 纳米级产能本身就是瓶颈。但政治信号很明确:一旦我们不需要美国芯片也能完成训练 这件事被证明,哪怕只是一次,美国基于芯片的封锁,其有效性与象征意义就会同时崩塌。

而 DeepEP 的 Mori-EP / UCCL-EP 分支已经支持 ROCm 和异构 NIC——这意味着"不依赖 NVIDIA 的 MoE 训练栈"的一块块组件,正在悄无声息地被一件件拼齐。

6.3 第三阶段——通过开源进行生态锁定

DeepSeek 的 V3、R1、V3.2、V4 都以 MIT 或与之等价的许可证开放权重。DeepEP、DeepGEMM、FlashMLA、DualPipe、3FS 也全部是 MIT。这不是慈善,而是一种明确的战略。当模型与基础设施都开源,全世界的云服务商、研究机构、企业,会自发地去分发、修改、改进"在中国芯片上优化的模型"和"中国公司做的通信库"。这些 know-how 在全球社区层面累积,然后再回流进中国硬件生态。

美国之所以已经开始考虑把模型权重本身也纳入出口管制对象,原因正在于此。原本以为只要锁住芯片就够了,结果一次开源,就能把全世界的开发者动员成"中国硬件的自愿优化劳动力"。DeepSeek 精准地理解并利用了这一点。

七、对韩国 AI 与软件产业的启示

DeepSeek 这个案例,向韩国的 AI、区块链、网络安全产业,提出了三个问题。

第一,"我们公司的成本瓶颈到底是什么?" 在充裕 GPU 上做"创意模型"的公司,与在稀缺 GPU 上榨"每瓦特 token"的公司,最终会做出不同的结果。韩国 AI 创业公司,需要扪心自问自己更接近哪一边。
第二,"我们的垂直整合做到了多深?"DeepSeek 把通信库、GEMM 内核、分布式文件系统都自己运营。如果韩国企业不能降低对 NCCL、cuBLAS、FlashAttention 这类外部库的依赖度,输的就不是成本之争,而是控制权之争。DeepEP 传达出来的最务实的信息是:门开着的时候,进去把代码读了。-第三,"制裁对我们意味着什么?" 美国对华制裁,对韩国是双刃剑。一方面,它给韩国半导体与 AI 企业带来短期的转移需求;另一方面,"韩国对美国芯片生态的依存度"会原封不动地转化为风险。DeepSeek 向华为昇腾迁移的过程,对韩国是一面镜子——映出我们自己的问题:如何拼出自家 NPU、AI 加速器生态。

在我自己运营的公司,以及我从事的区块链产业一线,同样的模式不断重复。深度依赖外部基础设施的项目,无法把成本结构握在自己手上。DeepSeek 的教训不局限于 AI 领域。基础设施栈的哪一层你能自己"削到内核",最终决定了在监管与冲击的时代,哪些公司能活下来。

八、结论——"最后一个时钟周期"的竞争

DeepSeek V4 真正的信息,不是 1.6 万亿参数这个数字。它真正展示的是:"受限硬件 + 深度系统工程 + 开源战略"这一组合,可以与美式的"充裕硬件 + 闭源模型"模式正面竞争。从 HFT 交易室里诞生的那个问题——"如何削掉最后一个时钟周期"——已经扩展为整个 AI 基础设施的设计哲学。

DeepEP 是这一哲学最浓缩的产物之一。一个 58.9% 由 CUDA 写成的库,既解开了 MoE 训练的通信瓶颈,又把推理 decoding 的延迟以微秒为单位削减,与此同时,它还以 MIT 许可证开放,把自己的标准传播给全世界的开发者。只发布模型的时代结束了。现在,是连基础设施内核都要做最优化的时代。

美国制裁在短期内是有效的。但当对手是一家"把约束本身当作成本函数输入变量"的公司时,这种制裁会被中和得比预期更快。DeepSeek 正是在最大的舞台上跑这个假设的公司,而这次实验的结果,将在未来五年里,作为决定全球 AI 产业格局的最重要变量之一保留下来。韩国产业,也不能把这个变量当作"别人家的事"。这个问题,很快也会送到我们门口。