AMD MI250部署DeepSeek-V4-Flash：先修正确性再谈性能

前阵子帮朋友在实验室的 AMD MI250CDNA2机器上部署 DeepSeek-V4-Flash模型倒是跑起来了GPU 利用率也看着正常可生成出来的内容总透着一股“不对劲”——有些句子读不通有些甚至出现重复 token 的死循环。一开始我以为是采样参数没调好直到把 logits 导出来跟另一台 NV 机器对比才发现输出分布早就偏了。后来看到这篇论文的标题先修正确性再谈性能一下就把我戳中了。这篇论文讲的是怎么在 AMD MI250CDNA2上把 DeepSeek-V4-Flash 跑对、跑稳、再跑快核心观点非常朴素但容易被忽略如果你的模型输出本来就是错的那吞吐和延迟数字再漂亮也没有意义。适合所有想在非 CUDA 平台上部署大模型的工程师尤其是正在为“为什么同样的模型在 AMD 上结果不对”而掉头发的人。下面是我读完论文之后结合自己实际部署经验的完整拆解。1. MI250 和 DeepSeek-V4-Flash这对组合为什么值得较真1.1 MI250 的底子不只是一块“显存很大”的卡AMD Instinct MI250 是一张很有意思的数据中心加速卡它基于 CDNA2 架构核心特征是双 GCD 设计——一颗 MI250 在系统里看起来像一张卡实际上是两个计算单元GPU Compute Die通过 Infinity Fabric 互联封装在一起的。每个 GCD 有自己的显存控制器和 L2两个 GCD 加在一起给到 128GB HBM2e 显存、约 3.2TB/s 的聚合带宽这对于跑大模型来说是非常诱人的条件毕竟 LLM 推理的最大瓶颈往往不是算力是显存带宽和容量。老的 GPGPU 用户可能更熟悉 Vega、MI50、MI100 那一代CDNA2 相比它们最明显的进步是Matrix Core 支持更成熟的低精度矩阵运算FP16/BF16 都有专门的硬件路径同时 FP64 性能也保留得非常好。MI250X 的 FP64 峰值非常能打这在做科学计算的人眼里是宝但对跑 LLM 的人来说更关心的是 FP16/BF16 的吞吐以及 ROCm 软件栈到底顺不顺。如果拿 MI250X 跟当时主流的数据中心卡做直观对比大概是这种感觉项目MI250XCDNA2A100 80GBH100 80GB显存容量128GB HBM2e80GB HBM2e80GB HBM3显存带宽约 3.2TB/s双 GCD 合计约 2TB/s约 3.35TB/sFP16/BF16 矩阵算力约 95.7 TFLOPS带稀疏则更高约 312 TFLOPS带稀疏约 990 TFLOPS带稀疏FP64 算力约 47.9 TFLOPS极强约 9.7 TFLOPS约 34 TFLOPS软件生态ROCm / HIPCUDACUDA所以 MI250 跑大模型真正的优势是大显存和高带宽而不是绝对算力。DeepSeek-V4-Flash 这种模型它的注意力机制经过压缩之后 KV cache 占用更小长上下文推理时反而能把 MI250 的显存优势吃得比较透。论文里也是从这个角度切入的要在 MI250 上验证的其实不是“能不能塞进去”而是“塞进去之后 FlashAttention 这些算子能不能在 ROCm 上给出正确结果”。1.2 DeepSeek-V4-Flash 给 MI250 出了哪些难题DeepSeek 系列模型的架构特点简单说就是两种核心设计MLAMulti-head Latent Attention和DeepSeekMoE稀疏专家混合。MLA 把 KV cache 压缩成低秩形式让同样显存下可以支撑更长上下文MoE 则是每层只激活少量专家大幅降低计算量但参数总量很大。这两种设计搬到 CDNA2 上各自都有坑。MLA 的注意力计算里大量使用低精度矩阵乘法和在线 softmax 类的算子FlashAttention 在 ROCm 上的成熟度远不如 CUDA 生态里的 cuDNN/CUTLASS 组合很多底层路径需要自己适配或者依赖 hipBLASLt 的替代实现。DeepSeekMoE 的专家路由在训练时通常以 CUDA 上的 top-k 行为为基准切换到 HIP 后哪怕排序或者取模的微小差异都会改变 token 被分配到哪个专家进而导致整层输出漂移。论文把这两类问题摔在台面上提出了一个非常工程化的主张“先修正确性再谈性能”。也就是说在 MI250 上跑 DeepSeek-V4-Flash第一步不是调 kernel不是抠显存带宽而是先把“正确性基线”立住。这个思路听起来朴素实际操作起来相当反人性——大家都想赶紧看到漂亮的 benchmark 数字而论文告诉你你得先花时间证明自己跑出来的模型还是原来那个模型。2. 论文是怎么发现问题“不正确”的正确性检查框架2.1 只靠“生成出来读得通”来验收是靠不住的我见过很多团队在大模型部署的验收阶段判断标准就是人眼读一遍生成文本“感觉还行”就过了。论文里明确否定了这种做法理由是LLM 的输出空间非常大采样路径稍微偏移生成的文字依然可能是通顺的但整体分布已经不一样了。如果你只是在几个 prompt 上人工看效果极大概率测不出问题。论文提出的做法是建立golden reference 对拍在一台信得过的参考平台论文用的是一套 CPU 双精度实现加上 CUDA 平台的官方权重上用固定输入和固定随机种子推一遍保存每一层的 logits 和中间张量。然后在 MI250 上跑同样的输入逐层比较指标而不是只比较最后的文字。这里有一个非常实用的脚本思路我复述一下论文里的做法import torch import numpy as np def compare_tensors(cand: torch.Tensor, ref: torch.Tensor, name: str): cand cand.detach().float().cpu().numpy() ref ref.detach().float().cpu().numpy() cos_sim np.dot(cand.ravel(), ref.ravel()) / ( np.linalg.norm(cand) * np.linalg.norm(ref) 1e-12 ) max_abs np.abs(cand - ref).max() rmse np.sqrt(np.mean((cand - ref) ** 2)) print(f{name:40s} | cos{cos_sim:.6f} | max_abs{max_abs:.6e} | rmse{rmse:.6e})比较时重点看的两个指标是余弦相似度和最大绝对误差。论文给的经验阈值是对 logits 来说余弦相似度低于 0.99 就应该立刻停下手头所有优化工作先把偏差找出来最大绝对误差在低精度推理里可以容忍 1e-2 量级但如果出现 1e-1 以上的误差基本可以直接断定某个算子路径没做对。2.2 从整模型到单算子二分定位是一条性价比很高的路拿到整模型 logits 的误差之后下一步是定位。论文没有建议把所有算子都重新实现一遍而是用二分替换法把从输入到输出的计算图切开先看前一半对不对再看后一半逐步缩小范围。具体操作上可以用 PyTorch 的原生算子替换掉自定义 kernel比如先把 FlashAttention 换成朴素的 attention 实现再看误差是否消失。如果消失了问题就在 FA 上如果还在就继续往前面的 embedding、旋转位置编码层排查。论文在 MI250 上做这套对拍时定位出来的问题高度集中在两个位置MLA 里的注意力算子和MoE 的分组 GEMMGrouped GEMM。这两个位置加起来贡献了整个模型误差的八成以上。这一点跟我的经验完全吻合很多移植类项目里最怕的从来不是简单的矩阵乘法而是那些带有归约、排序、mask 分支的混合算子。提示对拍时一定要固定输入、固定随机种子、固定 pad 长度。不要图省事只拿随机 batch 跑一遍否则你分不清误差是模型实现差异还是输入扰动放大出来的。3. 深挖根因CDNA2 上的三类正确性黑洞3.1 数值精度FP16 的舍入和 BF16 的“砍尾数”不是一回事大模型推理里大家默认用 FP16 或者 BF16。但到了 AMD 平台上这个默认选择需要打个问号。FP16 有 10 bit 尾数BF16 只有 7 bit 尾数两者指数范围不一样。FP16 的问题是值域窄大数容易溢出BF16 的问题是精度粗小数容易被砍掉。CDNA2 的 Matrix Core 对低精度矩阵乘法的处理方式和 NVIDIA 的 Tensor Core 在累加路径上存在差异尤其体现在累加器精度和舍入模式上。论文的发现是DeepSeek-V4-Flash 的注意力计算里如果 QK^T 的结果在 FP16 下做 softmax长序列场景下 logits 最大值附近的数值会被舍入放大导致 attention 分布失真。更隐蔽的是 FlashAttention 的online softmax过程它需要不断更新 running max 和累加和如果这些中间状态用 FP16 存储数值稳定性会明显变差最终输出的概率分布整体平滑度下降。修复方向不是“无脑上 FP32”那会损失性能。论文的做法是分层处理矩阵乘法保持 FP16/BF16 输入但累加过程用 FP32softmax 的统计量running max、sum、scale全部用 FP32 保存只有数值动态范围很安全的算子比如某些 GELU、LayerNorm 的后半段才继续用低精度。我做实测时也验证过这个方向效果确实立竿见影。3.2 并行归约顺序的不确定性同一个模型跑两遍结果都不一样另一个让我印象深刻的点是非确定性non-determinism。在 CUDA 生态里你可能已经习惯了某些 kernel 本身是不确定的但 PyTorch 会尽量提供确定性选项。到了 HIP/ROCm 平台上这个问题会更明显因为 MI250 是双 GCD 结构两个 die 之间还有 Infinity Fabric 互联当计算被切到两个 GCD 上时跨 die 的归约顺序会变化。论文里举了一个很典型的场景MoE 的 expert 路由统计或者 attention 里对多个头的结果做拼接后再归约如果底层实现用了 GPU 全局原子操作来处理部分和那么原子操作的完成顺序会影响最终浮点结果。这种偏差在数值上可能很小但如果后续接了 softmax 或者动态规划类的算子会被放大。修复方式是典型的“先正确后高效”放弃跨 block 的原子累加改成两级归约——先在 block 内部用固定的树状归约再把每个 block 的局部结果按固定顺序比如 blockIdx 从小到大的线性顺序交给主 block 聚合最后只用一个确定性原子操作写入结果。这样会牺牲一点效率但换来“同一份输入跑两遍结果完全一致”的可能。3.3 FlashAttention 的 mask 分支和 scale 更新顺序如果你把 PyTorch 自带的 FlashAttention 换成 ROCm 生态里的一些替代实现大概率会遇到一个经典问题mask 处理方式不同导致数值结果不一致。NVIDIA 的 FlashAttention 在实现 causal mask 时有自己的优化路径比如用布尔掩码加负无穷而有些 ROCm 实现会先把 mask 转换成一个加性偏置矩阵两者在数学上等价但浮点计算路径完全不一样误差就来了。更细节的是 online softmax 里 scale 的更新顺序。标准的 online softmax 公式里每当处理一个新的 block 时需要把之前累计的概率分布乘上一个衰减系数这个系数的计算顺序直接影响浮点误差。论文就发现某个版本的 FA 在一路实现里把衰减系数先做了一次乘法再取指数另一个版本先取指数再相乘表面等价实际误差在长序列下能差到 1e-2 以上。这种问题通常不会让输出变成 NaN但会让交叉熵升高、生成质量下降。提示在 ROCm 上切换 FlashAttention 版本时不要只看 benchmark务必先在固定输入上做 logits 级对拍。很多版本的 FA 比原生实现快但数值路径已经变了。4. 修复策略论文给出的“先做对”具体方案4.1 精度策略回归关键路径用 FP32 累加低精度留给安全区论文给出的第一招是给整模型做一次精度体检而不是一刀切用 FP16。它把算子分成了三类强制 FP32、允许 FP16/BF16、允许 FP8如果有条件。分类依据是这个算子的输出对后续分布的影响有多大、以及动态范围是否足够安全。强制 FP32 的包括LayerNorm 和 RMSNorm 的统计量计算与归一化、softmax 的指数和对数路径、旋转位置编码的三角函数计算。允许低精度的包括大部分矩阵乘法的输入和输出、MoE 里 gate 网络的前几层、MLP 的中间激活。论文给了一个非常直观的实测表格我看完拿来直接抄作业了计算路径默认精度修复后精度最大绝对误差改善QK^T 矩阵乘FP16FP16 输入 FP32 累加约 10 倍Softmax 统计量FP16FP32约 5 倍PV 加权求和FP16 输入 FP16 累加FP16 输入 FP32 累加约 8 倍MoE 路由 top-kFP16FP32 比较 固定 tie-break 规则完全消除随机差异最终输出 logits 层FP16BF16 FP32 累加约 3 倍这些数字看起来不复杂但背后是逐算子对拍才逼出来的结论。我当时照做之后logits 的余弦相似度直接从 0.87 拉到了 0.997肉眼可见生成文本质量恢复正常。4.2 确定性归约把“跑两遍不一样”的问题彻底堵死针对并行归约顺序的不确定性论文的解决方式前面提到过局部归约 固定顺序的全局归约。这里给出一个 HIP kernel 的示意版本能更直观说明问题// 不推荐直接使用全局原子累加顺序不确定 __global__ void bad_reduce(float* out, const float* partial, int n) { float sum 0.0f; for (int i threadIdx.x; i n; i blockDim.x) sum partial[i]; atomicAdd(out, sum); // 不同 block 的完成顺序会导致浮点误差 } // 推荐两级归约按固定顺序聚合 __global__ void good_reduce(float* out, const float* partial, int n) { __shared__ float local[BLOCK_SIZE]; int tid threadIdx.x; float sum 0.0f; for (int i tid; i n; i blockDim.x) sum partial[i]; // block 内树状归约 for (int stride BLOCK_SIZE / 2; stride 0; stride 1) { __syncthreads(); local[tid] sum; __syncthreads(); if (tid stride) sum local[tid stride]; } // 只有 block 0 按 blockIdx 顺序汇总 if (tid 0) { float total local[0]; for (int b 1; b gridDim.x; b) total outBufferBase b; // 用固定存储区汇总 out[0] total; } }当然实际实现里不会每步都这么粗暴地用 shared memory 写而是借助 GPUDirect 或者 peer-to-peer copy 把每个 block 的局部结果逐个读到主 block核心思想是一样的放弃“所有 block 同时原子加”换成“按 blockIdx 顺序依次加”。论文同时建议把整个模型推理的随机源全部清零固定 torch 和 hipBLASLt 的随机种子禁止某些算子内部使用不稳定的排序算法尤其是 MoE 的 top-k 选择要规定当概率值并列时按 expert_id 排序而不是按内存地址排序。这个 tie-break 规则看起来是小细节但直接影响路由的确定性论文专门拿了一节讲这件事。4.3 把校验做成流水线的一部分而不是一次性的活动工程上最怕的是“修完这一次下次改个 kernel 又坏了”。论文很实在直接建议把正确性校验做进 CI/CD 流程里。不需要每天跑全量测试集只需要维护一个golden prompt 集合比如 20 个固定 prompt、固定长度、固定采样参数每次构建新的 ROCm 算子库、升级 hipBLASLt 或者改动了 kernel 融合策略之后先在这组 golden 上跑一遍逐层对拍。只要有一个算子超阈值就不要进性能优化阶段。这个方法我在自己的部署流程里沿用下来效果非常好。一开始觉得每天跑对拍很繁琐但后来发现它其实是在帮你挡住那些“看起来很爽的优化提交”——比如某次把 GEMM 切块参数改大benchmark 涨了 5%但对拍发现精度掉了一截这种提交如果没有校验关卡上线后就是线上事故。5. 正确性敲定后性能优化是怎么展开的5.1 先建立“正确且可复现”的 baseline性能数字才作数论文里有一段话让我印象很深在修复正确性之前团队其实已经做了很多性能优化GPU 利用率看着也不低但生成结果是错的等于所有优化全部白费。修复正确性之后同样的优化手段重新应用才真正转化为可用的性能提升。换句话说错误输出之上的性能数字没有任何可比较的意义因为你不知道它应该有多快也不知道它是不是在“空转”。正确的流程是先跑通正确性对拍把这个时间点的延迟和吞吐作为 baseline。之后任何优化都必须同时满足两个条件才合入——性能比 baseline 好正确性指标没跌破阈值。论文给出的性能结论我没有在这里列具体数字因为跟 ROCm 版本、驱动版本、模型量化方式都有关系但它强调了一个趋势一旦正确性基线立住后面做 FP8 量化、kernel 融合、多实例部署时就不需要再猜“是不是数值又不对了”风险大幅下降优化速度反而更快。5.2 CDNA2 上真正值得抠的优化方向论文第二部分花了不少篇幅讲性能优化这里我挑几个跟 CDNA2 特别相关的方向讲第一向量化访存。MI250 的显存带宽极其充沛但需要你把它用出来。很多移植过来的 CUDA kernel默认 32bit load 或者 64bit load 就能跑但在 CDNA2 上应该尽量用 128bit load比如 double4 / float4 的向量化读取这样能显著提高内存级并行度。第二块大小与 L2 分块的取舍。CDNA2 的 L2 比 CUDA 生态常见设置要敏感一些。FlashAttention 类的算子分块太大容易把 L2 占满分块太小又导致块间反复读同一份 KV 数据。论文的做法是二分搜索分块参数在固定 batch 下尝试 head_dim128 时 block_M/block_N 的不同组合而不是沿用 CUDA 版本里那组默认值。第三双 GCD 的利用方式。MI250 最独特的地方就是两个 GCD。论文讨论了两种模式MCMMemory Consistency Model模式和镜像模式。镜像模式编程简单两个 GCD 看到同一份地址映射但带宽受限MCM 模式让每个 GCD 主要访问自己的显存性能更好但需要显式管理跨 die 的数据分布。对于 DeepSeek-V4-Flash 这种显存占用不在话下的模型论文推荐把模型权重分别放在两个 GCD 本地显存里只在 attention 归约和 MoE 跨专家通信时走 Infinity Fabric尽量把数据交换量压到最低。第四融合算子的顺序。在 ROCm 上做 kernel 融合和 CUDA 上不完全一样。CUDA 生态的融合模板已经很成熟照抄模板就行ROCm 上 hipBLASLt 和 Composable Kernel 给了很多组件但融合的边界需要自己摸索。论文的经验是优先融合那些访存密集、计算量小的算子比如 Layernorm QKV 投影、Residual 激活函数这些矩阵乘之外的融合收益最高且不容易破坏数值路径。5.3 跟 CUDA 平台的差距有多大还有多少空间论文里做了一组客观对比结论我总结一下就是差距是存在的但不是不可逾越的鸿沟。在单卡推理场景MI250 上的 DeepSeek-V4-Flash 性能大约是同代 NVIDIA 平台的一个档次之后主要差在 FlashAttention 的成熟实现和低精度矩阵算力上。但如果在显存容量敏感的推理场景下比如长上下文、大 batch、KV cache 占用高的部署MI250 的 128GB 显存能帮你省去很多“显存不够被迫拆分模型”的烦恼综合部署成本反而有优势。论文也给出了后续优化空间的判断一是FP8 量化CDNA2 对 FP8 支持不如后来者完善但通过分块量化也能拿到可观收益二是多卡扩展MI250 的双 GCD 结构天然适合用单卡模拟双卡环境把跨 GCD 通信模式研究透之后多机部署的很多问题能提前暴露。这部分的挖掘空间还很大如果社区一起把 ROCm 上的 FA 版本和 MoE kernel 打磨成熟差距会进一步缩小。6. 从论文回看我自己的部署一些可以复用的经验6.1 部署第一天就把检查和版本固定一起做掉读完这篇论文我对自己那次部署最大的后悔就是检查框架没有第一天就搭上。后来重新整理了一下流程整理成一份“AMD 部署检查清单”分享出来固定环境记录 ROCm 版本、hipBLASLt 版本、PyTorch 的 ROCm 构建版本、驱动版本最好用容器镜像一次性固化。建立 golden reference在参考平台上导出 20 组固定输入对应的 logits 和 hidden states存成 npz/hdf5 文件。跑对拍脚本逐层比较余弦相似度和最大绝对误差阈值设定参考论文的建议。固定随机源关闭 TF32、固定 cuDNN/hipBLAS 的 seed、关闭 torch 的非确定性算法。验证双 GCD 策略确认模型是跑在单 GCD 还是双 GCD 模式显式指定数据放置不要让系统自动调度。再进入性能调优每做一次优化重新跑一次 golden 对拍记录性能与精度的“双曲线”。这里面最容易被忽视的是第 4 条。很多人跑完对拍发现误差时大时小第一反应是 kernel 写错了实际上只是随机源没固定。先固定随机源再谈算子比较可以省掉大量无效排查。6.2 我在实操中踩过的几个坑第一个坑只测 perplexity不测逐 token 分布。我一开始觉得perplexity 没掉很多模型应该没问题。实际上 perplexity 是一个整体指标它会把很多错误拉平。逐 token 的 top-10 预测分布对拍才能暴露问题。论文里也是用 logits 级别的对比而不是损失函数级别的对比。第二个坑混用不同版本的 FlashAttention。我的部署脚本里某个组件用了 pip 包装的 FA 版本另一个组件用了从源码编译的 Composable Kernel 版本。两者计算结果在小误差范围内不一致导致整个模型的 logits 偏离。后来统一成一个 FA 实现问题立刻缓解。教训是在 AMD 平台上kernel 的版本混合使用比 CUDA 平台更容易引发数值漂移所有算子必须锁定同一个构建链。第三个坑忽视 mask 的 padding 处理。长 prompt 场景下padding token 的 mask 处理方式在 PyTorch 的 SDPA 和自定义 HIP FA 实现之间往往不同。有些实现把 padding 位置参与计算再 mask 掉有些实现直接把 padding 位置的 score 设为负无穷看起来结果等价但浮点路径不同。论文的修复方式是在预处理阶段显式构造一个 token-level mask而不是依赖算子内部的隐式行为。6.3 一个容易被忽略的细节分布式策略会影响归约顺序如果你跑的是多卡并行比如把 DeepSeek-V4-Flash 用 tensor parallel 切到多个 GCD 上那么跨卡的 all-reduce 顺序会直接影响浮点结果。论文观察到一个现象同样的模型、同样的输入TP 切分方式不同输出 logits 会有微小差异对于已经敏感的注意力得分来说这个差异不能被忽略。修复方式不一定是要改变 all-reduce 顺序那会牺牲性能而是调整对拍时的参考基线以你计划上线的这套分布式策略的输出来做 golden reference而不是拿单卡输出当参考标准。这一点非常实用因为很多人的思维定式是“单卡的结果是正确的多卡结果不对”但浮点上并不存在这样的绝对关系多卡归约顺序不同也会产生误差。论文的做法是多卡部署时固定下来一套通信策略并用它自己的输出作为终态基线后续只要保持同样的通信顺序结果就是可复现的。最后再分享一个小技巧我后来每次在 AMD 卡上部署新模型都会这么做不要急着跑完整训练集或者一整个 benchmark先准备一个长度适中、覆盖常见句式和 padding 情况的 prompt跑通单样本对拍把输出 logits 存下来。这一步成本极低但能帮你挡住八成以上的“模型在 AMD 上结果不对”的坑。先把正确性钉死在某个具体输入上再去思考性能和扩展你会发现自己排查问题的效率翻了一倍不止。