필사 모드: AMD 与 NVIDIA 有何不同 — 问题不在硬件，而在软件栈

• 引言 — HIPIFY 搬走 90% 之后，剩下的那 10%
• 硬件 — SM 与 CU，Tensor Core 与 Matrix Core
• warp 32 对 wavefront 64 — 最昂贵的一行代码
• 软件栈 — CUDA、ROCm 与 HIP
  • 库对应表
• 移植路径 — HIPIFY 能搞定的和搞不定的
  • 自动搬过去的部分
  • 必须手动修改的部分
  • 移植后必须做的验证
• 现实中真正的障碍是什么
  • 1. 内核生态的积累量
  • 2. 框架支持的滞后与支持矩阵
  • 3. 出问题时的信息量
• 什么样的负载下 AMD 是合理选择
• 结语 — 差距不在硅片，而在积累下来的内核
• 参考资料

引言 — HIPIFY 搬走 90% 之后，剩下的那 10%

把 CUDA 代码库移植到 AMD，大致会经历这样的过程。跑一遍 hipify-perl，大部分文件几秒钟就转换完了。cudaMalloc 变成 hipMalloc，__global__ 原封不动，内核启动的语法也一样。构建能过，小测试也能过，到这里花了半天。

接下来要花两周。追查一个结果有微妙错误的内核，发现 __shfl_xor 后面硬编码了一个 32；卡在一个性能只有一半的内核上，才发现共享内存的 tile 大小是按 warp 32 定下来的。有一个文件里带内联 PTX，只能整个重写。

这篇文章要讲的，正是这 10% 究竟是什么，以及比它重要得多的生态差距是什么。不摇旗呐喊，也不冷嘲热讽，只看结构。核对的基准是 ROCm 7.14.0（2026 年 7 月 16 日发布）、vLLM v0.26.0 文档、AMD ROCm 官方文档。

硬件 — SM 与 CU，Tensor Core 与 Matrix Core

先把术语对应上。概念层面几乎是一一对应的。

构建时需要指定目标架构，这一点也一样。就像 NVIDIA 上要给 -arch=sm_90，ROCm 上要给类似 gfx942 这样的值。自己设备上的值用 rocminfo 确认。

rocminfo | grep gfx        # 例如: gfx942 (MI300 系列), gfx950 (MI350 系列)
rocm-smi                   # 相当于 nvidia-smi。查看利用率、温度、功耗

把现在市面上会遇到的几代产品整理一下，是这样的。以下硬件数据是综合厂商公告与二手报道得出的，我本人并未实测。做采购决策时，请务必核对厂商官方的规格表。

这里要看的不是具体数字，而是两家公司的硬件在同一代产品上，大体达到了相近级别的内存容量和带宽这一事实。有的世代 AMD 在容量上领先，也有的世代 NVIDIA 追了上来。在这个维度上没有决定性差距。

所以实际工作中的差距来自别的地方。别的地方，就是这篇文章剩下的内容。

warp 32 对 wavefront 64 — 最昂贵的一行代码

在硬件差异里，唯一会直接伤到代码的就是这一条。NVIDIA 的 warp 是 32 个线程。AMD CDNA 系列数据中心 GPU 的 wavefront 是 64 个。RDNA 系列则是 32 个。

HIP 文档专门对这一点提出了警告。

Code should not assume a warp size of 32 or 64, as AMD GPU architectures have different warp sizes. The warpSize built-in should be used in device code.

这一行代码在实际中会引发三类 bug。

第一，shuffle 归约的迭代次数。回头看一下之前文章里写过的 warp 归约。

// 假设跑在 NVIDIA 上写的代码。在 AMD 上只会归约一半。
__inline__ __device__ float warpSum(float v) {
  for (int off = 16; off > 0; off >>= 1)          // 16 是假设 32 而来的
    v += __shfl_xor_sync(0xffffffffu, v, off);
  return v;
}

如果 wavefront 是 64，off 就得从 32 开始。从 16 开始，会导致前 32 个 lane 和后 32 个 lane 分别独立求和，最终结果只有一半是对的。能编译，不会崩溃，就是算错了。这是移植工作中最容易卡住最久的一类 bug。

可移植的写法是这样的。

// HIP。在 AMD 和 NVIDIA 上都能正确运行。
__device__ float waveSum(float v) {
  for (int off = warpSize / 2; off > 0; off >>= 1)
    v += __shfl_xor(v, off, warpSize);
  return v;
}

第二，lane mask 的位宽。在 NVIDIA 上，活跃 lane 掩码是 32 位，所以 0xffffffff 用起来很自然。到了 wavefront 64，就需要 64 位。HIP 文档在这里还写了一个具体的坑：在 64 宽的 wavefront 上，把一个 32 位整数移位超过 31 位，寄存器会被清零。解决办法是给 lane mask 用 uint64_t。

第三，tile 大小和 block 大小里藏着的隐含假设。如果把 block 大小 256 想成"8 个 warp"，并据此把共享内存里的部分和数组开了 8 格，那么在 AMD 上，wavefront 只有 4 个，只会用到 4 格，逻辑就错位了。反过来，只按 32 的倍数对齐的 block 大小，在 AMD 上未必是 64 的倍数，可能导致最后一个 wavefront 有一半空转。

原则只有一条。不要把 32 或 64 当成常量硬写进源码。要么用 warpSize，如果确实需要编译期常量，就集中放在一个地方，按架构分别替换。

软件栈 — CUDA、ROCm 与 HIP

把各层并排放在一起，是这样的。

NVIDIA                          AMD
──────────────────────          ──────────────────────
PyTorch / JAX / vLLM            PyTorch / vLLM
   │                               │
cuBLAS, cuDNN, NCCL             rocBLAS, MIOpen, RCCL
   │                               │
CUDA Runtime API                HIP Runtime API
   │                               │
CUDA Driver                     ROCr 运行时 + ROCk 内核驱动
   │                               │
NVCC → PTX → SASS               hipcc(LLVM) → AMDGCN ISA

HIP 在设计上刻意做成了和 CUDA 几乎一样的形状。函数名里的 cuda 前缀换成 hip，内核定义和启动语法基本相同。这不是巧合，而是设计目标本身。而且 HIP 代码在 NVIDIA GPU 上也能编译。也就是说，HIP 可以当作 CUDA 之上的一层薄封装来用，这也是需要同时支持两家厂商的库，会选择把 HIP 作为单一代码源的原因。

这里有件事值得说一下 ROCm 7.14.0。版本号从 7.2.4 跳到 7.14.0，看起来像打字错误，但确实如此。AMD 的发行说明明确写了"从 7.9.0 预览版开始的版本号不连续"，同时说明 ROCm 正在转向一套叫 TheRock 的模块化构建与发布系统。做法是让核心 SDK 更轻量，AI、数据科学、HPC 用的领域 SDK 则按需选择安装。

这件事本身就带来了一条对实际工作有用的信息。ROCm 的结构现在仍在变化中。这意味着安装步骤和包名可能随每个大版本变化，看文档时必须对准版本号。CUDA 这一侧在这个维度上要稳定得多。

库对应表

这里要理解 AMD 特有的双重命名方式。AMD 官方文档是这样解释的：roc 前缀的库是面向 AMD GPU、用 HIP 写成的原生高性能实现，hip 前缀的库则是实现了 CUDA 对应 API 的移植用封装层。hipBLAS 把自己称为"marshalling library"，它背后既可以接 rocBLAS，也可以接 cuBLAS。

选择标准很明确。如果是从 CUDA 移植过来，就用 hip 这一侧。调用形式相同，代码改动最少。如果是以 AMD 为主要目标全新编写，就用 roc 这一侧。少一层封装，可以直接用到 AMD 专属功能。

表格最下面三行，各有一格是空的。cuDNN、NCCL、CUTLASS 没有移植用的封装层，只有另起名字的对应物。这说明它们的 API 形态不同，不改源码是换不了的；不过深度学习框架各自把这一层吸收掉了，所以对大多数用户来说不构成问题。反过来，如果你手上的代码直接调用这些库，那部分就是移植工作里核心的成本所在。

只看 LLM 推理的话，最近出现的 AITER（AI Tensor Engine for ROCm）很重要。这是 AMD 收拢 LLM 推理内核的一个仓库，vLLM 的 ROCm 安装文档把它纳入了构建流程。它大致相当于 NVIDIA 那侧 FlashInfer 或 FlashAttention 所占的位置。

移植路径 — HIPIFY 能搞定的和搞不定的

工具有两个。

实际工作中，很多时候会先用 hipify-perl 开始。因为在大型代码库里把 CUDA 头文件全部配齐太麻烦了。

# 预览单个文件(不改动原文件，只输出转换结果)
hipify-perl kernel.cu

# 原地转换整个目录，并留下备份
find . -name "*.cu" -o -name "*.cuh" | xargs hipify-perl -inplace -print-stats

# 需要准确度时用基于 clang 的工具
hipify-clang kernel.cu -- -I/usr/local/cuda/include

自动搬过去的部分

• 运行时 API 调用名称（cudaMalloc、cudaMemcpy、cudaStreamCreate 等）
• 内核限定符与启动语法（__global__、__device__、三重尖括号启动）
• 内置变量（threadIdx、blockIdx、blockDim）
• 大多数数学内置函数
• 头文件包含语句

到这里为止，按代码量算大概是 90%。

必须手动修改的部分

• 内联 PTX 汇编。PTX 是 NVIDIA 的虚拟 ISA，AMD 没有对应物。这部分要么用 HIP 内置函数重写，要么整个用 AMDGCN 内联汇编重新写。这是移植工作里最稳定吃时间的一项。
• warp 大小假设。前面整节说的全部内容。自动转换没法判断某个 32 到底是不是 warp 大小，还是别的意思，所以干脆不碰。
• warp 级原语的语义差异。__shfl_sync 系列的掩码参数、__ballot 的返回位宽、显式同步的语义都不一样。名字会被搬过去，但语义需要验证。
• 对 CUDA 专属库的调用。像 cuDNN、CUTLASS、cuBLASLt 这类没有对应封装层的库。
• 使用驱动 API 的代码。像 cuModuleLoad 这样直接调用底层驱动 API 的部分。
• 所有性能调优常量。tile 大小、block 大小、展开程度、流水线阶段数。这些会被转换，但最优值不一样。移植和优化是两件不同的事，跳过这一步的结果就是"能跑，但只有一半速度"。

移植后必须做的验证

# 1. 先做数值验证，性能放在后面。
#    重点看带有归约操作的内核。
pytest tests/ -k "reduction or attention or norm"

# 2. 用 AMD 侧的 profiler 重新测一遍真正的瓶颈。
#    不要直接照搬 NVIDIA 上调优出来的结果。
rocprofv3 --stats -- ./my_app
rocprofv3 --kernel-trace -- ./my_app

对应 Nsight Compute 的 AMD 工具是 rocprofv3 和 ROCm Compute Profiler。概念是一样的：看每个内核的耗时、内存吞吐量、缓存命中率、占用率。

现实中真正的障碍是什么

硬件明明相近，实际现场却出现差距，原因可以归纳成三条。

1. 内核生态的积累量

这是最大的一项。每当出现新的模型结构或新的量化格式，对应的内核几乎总是先在 CUDA 上出现。FlashAttention 的新变体、新的 MoE 路由内核、新的低精度 GEMM，全都是这样。AMD 的支持会晚几周到几个月才跟上。

Triton 正在实实在在地缩小这个差距。Triton 仓库的 third_party 目录里，nvidia 和 amd 两个后端并排放着，用 Triton 写的内核两边都能编译。就像之前文章里提到的，最近的自定义 attention 内核有先在 Triton 上出现的趋势，如果这个趋势持续下去，差距会持续缩小。不过，依赖手写 CUDA 内核的部分，仍然存在时间差。

2. 框架支持的滞后与支持矩阵

看一个具体例子。vLLM v0.26.0 的 ROCm 安装文档是这么写的。

• 支持 ROCm 6.3 及以上版本，预编译的 wheel 只提供 ROCm 7.0 和 ROCm 7.2.1 版本。
• 支持的 GPU 有 MI200 系列（gfx90a）、MI300（gfx942）、MI350（gfx950）、Radeon RX 7900 系列、RX 9000 系列、Ryzen AI 系列。
• MI350 要求 ROCm 7.0 及以上。

这里就能看出滞后。ROCm 最新版本是 7.14.0，而 vLLM 的预编译 wheel 只做到 7.2.1。要用最新的 ROCm，就得转向源码构建，随之而来的是要分别对齐 Triton、FlashAttention、AITER 各自验证过的分支去构建。vLLM 文档让你去 Dockerfile 里核对这些分支值，这件事本身就说明了这套组合有多脆弱。

在 NVIDIA 那一侧，同样的工作通常就是一行 pip install vllm。这个差距纯粹是软件工程层面的摩擦，会吃掉团队的时间。

实际工作中的结论很明确。在 AMD 上，用经过验证的容器镜像事实上就是默认路径。vLLM 把官方镜像 vllm/vllm-openai-rocm 发布在 Docker Hub 上，AMD 之前发布的 rocm/vllm 系列镜像，现在也并入了官方镜像。决定自己动手构建，本身就是一个会带来相当维护成本的决定。

3. 出问题时的信息量

这一项不好量化，但确实存在。搜索错误信息时出来的结果数量、Stack Overflow 上的回答数、遇到同样问题的人写的博客、已经有答案的 GitHub issue——所有这些，CUDA 这边都压倒性地多。

调试花掉的时间，是直接计入项目进度的成本。它不会出现在硬件报价单上，但会算进总拥有成本里。

什么样的负载下 AMD 是合理选择

承认上面这些障碍之后，AMD 仍然是合理选择的场景确实存在。条件很具体。

第一，用推理来服务被广泛使用的模型。用 vLLM 服务 Llama、Qwen、DeepSeek 系列的主流模型，已经是一条走得很熟的路。只要用经过验证的容器、部署标准模型，上面说的大部分摩擦，都已经有别人先踩过、解决过了。

第二，内存容量起决定作用。把大模型放进更少数量的 GPU 里，能减少张量并行的通信量，这本身就能带来性能和简洁性。AMD 在某些世代上容量领先，在那些区间里，这是实实在在的优势。

第三，需要采购与议价能力。在大规模采购中，仅仅是存在第二供应商这件事本身就有价值。哪怕实际不用，能不能把一个备选方案摆上谈判桌，是完全不同的两回事。

第四，只在软件栈上层工作的团队。如果只在 PyTorch 之上工作，手上没有自定义 CUDA 内核，那移植成本的大头，一开始就不存在。

反过来，应该避开 AMD 的情况也同样明确。

• 有代码库的性能依赖手写 CUDA 内核或内联 PTX
• 需要比别人更早用上新模型结构和新内核的研究团队
• 无法为 GPU 基础设施配备专职人员的小团队
• 核心代码直接调用 cuDNN 或 CUTLASS

补充一条实用建议：如果你现在正在用 CUDA 写新内核，先看看能不能改用 Triton 来写，这是最便宜的保险。等以后真要评估 AMD 的时候，Triton 内核只需要重新编译，CUDA 内核则会变成一个移植项目。

结语 — 差距不在硅片，而在积累下来的内核

两家公司的 GPU，在同一代产品上都达到了相近的算力和相近的内存带宽。SM 和 CU 概念上是一回事，Tensor Core 和 Matrix Core 做的事情也一样。HIP 提供了和 CUDA 几乎一样的 API，HIPIFY 能自动搬走代码里 90% 的部分。只看到这里，差距似乎并不存在。

差距在于这层之下积累的东西。过去 15 年里为 CUDA 写出并调优的内核总量，建立在这些内核之上的库，建立在这些库之上的框架，以及围绕这一切积累下来的调试经验。这是一笔无法一夜之间复制出来的资产，也是实际性能差距背后的主要原因。

而这个差距会缩小的方向，也是从这里长出来的。内核越是用 Triton 这样可移植的层来写，而不是用某个厂商专属的语言，积累的效果就越不会只堆在一边。编译器吸收厂商锁定的那一层越厚，选择就越真正自由。这也是本系列下一篇要谈的主题。

如果现在就要做判断，问题就收敛成一个：我们的性能，到底是压在自己写的 CUDA 代码上，还是压在别人写的库上。如果是后者，AMD 值得评估。如果是前者，得先老老实实把移植成本估出来。

参考资料

• ROCm 官方文档: https://rocm.docs.amd.com/
• ROCm 发布记录(确认版本与日期): https://github.com/ROCm/ROCm/releases
• HIP 移植指南(包含 warp 大小的警告): https://rocm.docs.amd.com/projects/HIP/en/latest/how-to/hip_porting_guide.html
• ROCm API 库列表(与 CUDA 的对应关系): https://rocm.docs.amd.com/en/latest/reference/api-libraries.html
• HIPIFY 仓库与文档: https://github.com/ROCm/HIPIFY
• AITER (AI Tensor Engine for ROCm): https://github.com/ROCm/aiter
• vLLM ROCm 安装文档: https://docs.vllm.ai/en/latest/getting_started/installation/gpu.html
• AMD MI300X 调优指南: https://rocm.docs.amd.com/en/latest/how-to/tuning-guides/mi300x/index.html
• AMD Instinct 产品页面(官方规格): https://www.amd.com/en/products/accelerators/instinct.html
• NVIDIA 数据中心 GPU 产品页面: https://www.nvidia.com/en-us/data-center/