昨天刚说TileKernels，说到大概率会在昇腾，没想到今天这么快就发了。读DeepSeek的技术报告总是能学到很多东西，包含了大模型从0到稳定部署运行的全面技术过程，没有哪一个做AI的机构能够做到这种程度，敬佩！
报告内容非常庞杂，本文率先深入解析一下技术报告中第3章的基础设施部分3.1和3.3，以及个人对DeepSeek V4目前为什么要同时支持GPU和NPU两种芯片的看法，后续再分析attention部分，这里也很精彩
EP all2all部分的细粒度计算通信重叠
通信与计算重叠本人之前提到了很多关于TP场景下Gemm+RS和AG+GEMM，可以进本人主页搜一搜，这里也放几个传送门
大模型计算/通信overlapped kernel(一)--动机和概念
大模型计算/通信overlapped kernel(二)--AllGather+Gemm/Gemm+ReduceScatter
ParallelKittens: 使用ThunderKittens也可以轻松开发计算/通信overlapped kernel啦
DeepSeek-V4 聚焦MoE 专家并行（EP）的通信计算重叠优化，核心是解决多 GPU 部署下 MoE 模型 All-to-All 通信瓶颈，实现极致效率的计算通信协同，这一点在V3里面也反复提到All-to-All的通信开销非常大，尤其是只有H800的情况下，DeepEP尽了它的全部能力
核心问题在于MoE Expert Parallelism 包含 Dispatch、Linear + activation + Linear、Combine（结果回收）三个阶段，传统实现中通信与计算串行执行，All-to-All 通信耗时严重拖慢端到端性能。为此 DeepSeek-V4 提出细粒度的通信 - 计算重叠方案，将专家拆分为多个wave，构建三级流水线：当前 Wave 执行 Linear 计算、下一个 Wave 进行 Dispatch 通信、上一个 Wave 执行 Combine 回收，让通信 latency 完全隐藏在计算之下。类似套路之前已经出现，在Flux和comet和Flashoverlap中已经见到过，其中最相似的是Flashoverlap
它们之间的共识是以Wave作为overlap粒度，wave的意思可简单理解为一群block tiles，传统的张量切分太粗，而细化到Tile又太细，只不过小区别是 FlashOverlap 中按矩阵 Tile 划分的 Wave，还是 DeepSeek 中按专家划分的 Wave，它们都将 Wave 作为了流水线并发的最优粒度。
但是它们也有很大不同，FlashOverlap并不是一个MegaKernel，没有把通信和计算写在一个 Kernel 里。它让计算 Kernel 保持纯粹，算完一个wave后，发一个极轻量的信号给另一个专门负责通信的流，通信任务直接调用NCCL 。为了让乱序算完的数据能在内存里连续拼好发给 NCCL，它还在算子首尾插入了内存Reordering，为的是通用性好，毕竟NCCL那肯定通用，支持所有通信原语，TP里的all reduce，reducescatter，allgather小意思。然而DeepSeek-V4完全相反，它并没有用标准的 NCCL ，而是将通信和计算融合在了一起，是一个 MegaMoE kernel，在这个kernel里面我可以随意重排指令，这是在flashoverlap的信号量多流方式无法做到
此外论文提到DeepSeek-V4采用Pull-based，这种通信方式对于all2all这种token数量多的且通信数据量小的场景非常合适，延迟很低，与之对立的叫Push-based，这种适合粗粒度的，或者一对多通信，比如broadcast，如果在all2all场景用Push based，想象一下，rank0上的token发到各个rank你都要通知一遍，那延迟可高了去了
高性能的批次无关和确定性算子
很多朋友可能无法意识到这个的重要性，但这个对训练精度，推理精度，以及从GPU porting到NPU是非常关键的一步。
在拥有万亿参数的模型训练中，经常玄学报错，实现 Batch-Invariant批次无关和 Deterministic确定性可以一定程度避免这种问题以及易于debug。
Batch-Invariant指同一个 Token，无论所在 Batch 大小如何，算出来的浮点数必须比特级一致（Bitwise Identical），在 GPU的大规模并行计算中，由于浮点数的加法不满足结合律即 (a+b)+c =! a+(b+c)，如果动态改变了多线程的任务分配和累加顺序（例如Split-K reduce），就会导致同一个 Token 在不同 Batch 下算出来的结果尾数有一点点微小的差异。DeepSeek 为此开发了dual kernel, 保证了模型固定 Batch 训练时的表现，与线上部署时应对并发请求剧烈波动时的表现绝对一致。
而Deterministic指相同输入在多次反向传播中梯度完全一致。由于 GPU 多线程在执行 atomicAdd等原子操作是乱序抢占的，这个做过实验，确实如此，这会导致极其微小的浮点数误差，在万亿参数的网络中可能被放大为 Loss Spikes，如果系统是非确定性的，当你想要复现这个 Spike 去查找是哪一层、哪个矩阵的数据溢出时，你重新跑一次，它可能就不崩了，你永远也找不到 Bug 在哪。所以DeepSeek 去掉了所有全局无序的原子加法，改用预先排序、分配独立隔离缓存、最后按固定顺序归约求和的方式，保证了确定性，由此你就可以复现问题
同时，更加深层次的原因，本人个人认为和昇腾NPU的架构也有着紧密联系，谷歌的TPU 和 昇腾NPU 都是极其依赖编译器做静态调度（Static Scheduling，如 XLA）的硬件，它们并不擅长处理动态场景，因为它们没有分支预测、乱序执行等等，所以批不变性和确定性是 NPU/TPU 这种基于Systolic Array/达芬奇的硬件最喜欢的运行方式，有利于静态图的生成。而GPU 由于其天然的动态性，它无所谓，所以DeepSeek 改写反向传播逻辑，给每个单元分配独立缓存再归约，这估计也是为了make ascend happy
而且，为什么deepseek的报告里面提到同时支持了NV GPU和HW NPU？个人认为，你在NPU上做pretrain-posttrain-inference，你总得有个baseline吧，然后慢慢从GPU porting到NPU对伐，一下子登天是不可能的事，所以NV GPU我估计是拿来做精度和性能的baseline的，这对国产芯片真的是一个非常大的突破。当 NPU 算出的结果跟 GPU 不一致时，如果你做不到NPU/GPU都有确定性支持，就无法判断是 NPU 编译器写错了、精度溢出了，还是仅仅因为浮点数乱序带来的正常波动。
本人没用过昇腾，以上只是猜测，说错请理性讨论，欢迎关注我