前阵子帮朋友在实验室的 AMD MI250（CDNA2）机器上部署 DeepSeek-V4-Flash，模型倒是跑起来了，GPU 利用率也看着正常，可生成出来的内容总透着一股“不对劲”——有些句子读不通，有些甚至出现重复 token 的死循环。一开始我以为是采样参数没调好，直到把 logits 导出来跟另一台 NV 机器对比，才发现输出分布早就偏了。后来看到这篇论文的标题：先修正确性，再谈性能，一下就把我戳中了。
这篇论文讲的是怎么在 AMD MI250（CDNA2）上把 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 跟当时主流的数据中心卡做直观对比，大概是这种感觉：
| 项目 | MI250X（CDNA2） | A100 80GB | H100 80GB |
|---|---|---|---|
| 显存容量 | 128GB HBM2e | 80GB HBM2e | 80GB HBM3 |
| 显存带宽 | 约 3.2TB/s（双 GCD 合计） | 约 2TB/s | 约 3.35TB/s |
| FP16/BF16 矩阵算力 | 约 95.7 TFLOPS（带稀疏则更高） | 约 312 TFLOPS（带稀疏） | 约 990 TFLOPS（带稀疏） |
| FP64 算力 | 约 47.9 TFLOPS（极强） | 约 9.7 TFLOPS | 约 34 TFLOPS |
| 软件生态 | ROCm / HIP | CUDA | CUDA |
所以 MI250 跑大模型，真正的优势是大显存和高带宽，而不是绝对算力。DeepSeek-V4-Flash 这种模型，它的注意力机制经过压缩之后 KV cache 占用更小，长上下文推理时反而能把 MI250 的显存优势吃得比较透。论文里也是从这个角度切入的：要在 MI250 上验证的其实不是“能不能塞进去”，而是“塞进去之后 FlashAttention 这些算子能不能在 ROCm 上给出正确结果”。
1.2 DeepSeek-V4-Flash 给 MI250 出了哪些难题
DeepSeek 系列模型的架构特点，简单说就是两种核心设计：MLA（Multi-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 的分组 GEMM（Grouped 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 输入，但累加过程用 FP32；softmax 的统计量（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 矩阵乘 | FP16 | FP16 输入 + FP32 累加 | 约 10 倍 |
| Softmax 统计量 | FP16 | FP32 | 约 5 倍 |
| PV 加权求和 | FP16 输入 + FP16 累加 | FP16 输入 + FP32 累加 | 约 8 倍 |
| MoE 路由 top-k | FP16 | FP32 比较 + 固定 tie-break 规则 | 完全消除随机差异 |
| 最终输出 logits 层 | FP16 | BF16 + 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_dim=128 时 block_M/block_N 的不同组合，而不是沿用 CUDA 版本里那组默认值。
第三，双 GCD 的利用方式。MI250 最独特的地方就是两个 GCD。论文讨论了两种模式：MCM（Memory 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 上结果不对”的坑。先把正确性钉死在某个具体输入上，再去思考性能和扩展，你会发现自己排查问题的效率翻了一倍不止。
