- Published on
AMD と NVIDIA、何が違うのか — ハードウェアよりスタックが問題である理由
- Authors
- Name
- Youngju Kim
- @fjvbn20031
- はじめに — HIPIFY が 90 パーセントを移植した後に残る 10 パーセント
- ハードウェア — SM と CU、Tensor Core と Matrix Core
- ワープ 32 対 ウェーブフロント 64 — もっとも高くつく一行
- ソフトウェアスタック — CUDA と ROCm、HIP
- 移植の経路 — HIPIFY で済むものと済まないもの
- 現実に何が障害になるのか
- どんなワークロードなら AMD が合理的なのか
- おわりに — 差はシリコンではなく、蓄積されたカーネルにある
- 参考資料
はじめに — HIPIFY が 90 パーセントを移植した後に残る 10 パーセント
CUDA コードベースを AMD に移植する作業は、大体こんな流れで進みます。hipify-perl を実行すると、ほとんどのファイルが数秒で変換されます。cudaMalloc は hipMalloc になり、__global__ はそのままで、カーネル起動の文法も同じです。ビルドが通り、小さなテストが通り、ここまでで半日です。
そこから 2 週間かかります。結果が微妙に間違っているカーネルを一つ追いかけていると、__shfl_xor の後ろに 32 がハードコードされているのを見つけ、性能が半分しか出ないカーネルにつまずいていると、共有メモリのタイルサイズがワープ 32 を前提に決められていたことに気づきます。インライン PTX が含まれるファイルが一つあれば、それは丸ごと書き直すしかありません。
この記事では、その 10 パーセントが正確には何なのか、そしてそれよりずっと重要なエコシステムの差が何なのかを扱います。応援も非難もせず、構造だけを見ます。確認の基準は ROCm 7.14.0(2026 年 7 月 16 日リリース)、vLLM v0.26.0 のドキュメント、AMD ROCm の公式ドキュメントです。
ハードウェア — SM と CU、Tensor Core と Matrix Core
用語のところから対応させます。概念のレイヤーはほぼそのまま重なります。
| NVIDIA | AMD | 役割 |
|---|---|---|
| SM(Streaming Multiprocessor) | CU(Compute Unit) | 独立してスケジュールされる実行単位 |
| ワープ(32 スレッド) | ウェーブフロント(CDNA 64、RDNA 32) | 一緒に発行されるスレッドの束 |
| スレッドブロック | ワークグループ | 共有メモリを共有する単位 |
| 共有メモリ | LDS(Local Data Share) | SM/CU 内部のプログラマ管理 SRAM |
| Tensor Core | Matrix Core | 行列積和演算専用ユニット |
| NVLink | Infinity Fabric | GPU 間の高速リンク |
| コンピュート能力(sm_90) | gfx コード(gfx942) | 命令セット世代の識別子 |
ビルド時にターゲットアーキテクチャを明示しなければならない点も同じです。NVIDIA で -arch=sm_90 を指定するように、ROCm では gfx942 のような値を指定します。自分のデバイスの値は rocminfo で確認します。
rocminfo | grep gfx # 例: gfx942 (MI300 系列), gfx950 (MI350 系列)
rocm-smi # nvidia-smi に相当。使用率・温度・電力
今の市場で出会う世代を整理すると、次のようになります。以下のハードウェア数値はベンダーの発表と二次報道を総合したもので、筆者自身が実測したものではありません。購入判断の際は、必ずベンダーの公式スペックシートを確認してください。
| 部品 | 世代 | メモリ | 状態 |
|---|---|---|---|
| AMD MI300X | CDNA 3 | HBM3 192GB | 広く普及 |
| AMD MI355X | CDNA 4 | HBM3E 288GB、約 8TB/s | 展開中 |
| AMD MI455X(MI400 系列) | 次世代 | HBM4 | 2026 年 7 月発表、下半期出荷予定 |
| NVIDIA B200 | Blackwell | HBM3E 192GB | 広く普及 |
| NVIDIA B300 | Blackwell Ultra | HBM3E 288GB | 展開中 |
| NVIDIA VR200(Rubin) | Rubin | HBM4 | 2026 年下半期に量産予定 |
ここで読み取るべきは個別の数字ではありません。両社のハードウェアが、同じ世代でおおむね似たメモリ容量と帯域幅のクラスに達しているという事実です。AMD が容量で先行していた世代もあれば、NVIDIA が追いついた世代もあります。この軸で決定的な差はありません。
したがって、実務上の差は別のところから生まれます。その別の場所が、この記事の残りの部分です。
ワープ 32 対 ウェーブフロント 64 — もっとも高くつく一行
ハードウェアの違いの中で、コードに直接傷を残すのはこれ一つです。NVIDIA のワープは 32 スレッドです。AMD の CDNA 系データセンター GPU のウェーブフロントは 64 です。RDNA 系は 32 です。
HIP のドキュメントはこの点を明示的に警告しています。
Code should not assume a warp size of 32 or 64, as AMD GPU architectures have different warp sizes. The
warpSizebuilt-in should be used in device code.
この一行が実際に生むバグは 3 種類あります。
第一に、シャッフルリダクションの反復回数です。以前の記事で書いたワープリダクションをもう一度見てみます。
// NVIDIA を前提に書いたコード。AMD では半分しかリダクションされない。
__inline__ __device__ float warpSum(float v) {
for (int off = 16; off > 0; off >>= 1) // 32 を前提にした 16
v += __shfl_xor_sync(0xffffffffu, v, off);
return v;
}
ウェーブフロントが 64 なら、off は 32 から始めなければなりません。16 から始めると、前半 32 レーンと後半 32 レーンがそれぞれ別々に合算され、最終結果は半分しか合いません。コンパイルも通り、クラッシュもせず、ただ間違っている。これが移植作業でもっとも長く足止めされる種類のバグです。
移植可能な形にすると、こうなります。
// 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;
}
第二に、レーンマスクの幅です。NVIDIA ではアクティブレーンマスクが 32 ビットなので 0xffffffff が自然です。ウェーブフロント 64 では 64 ビットが必要になります。HIP のドキュメントは、ここにもう一つ具体的な落とし穴を書き残しています。64 幅のウェーブフロントで 32 ビット整数を 31 より大きい値だけシフトすると、レジスタが 0 にクリアされるというものです。解決策は、レーンマスクに uint64_t を使うことです。
第三に、タイルサイズとブロックサイズに埋め込まれた暗黙の前提です。ブロックサイズ 256 を「8 ワープ」と考えて、共有メモリの部分和配列を 8 マスで確保していたとすると、AMD ではウェーブフロントが 4 つしかないため 4 マスしか使われず、ロジックが狂います。逆に、32 の倍数にだけ合わせたブロックサイズは、AMD では 64 の倍数にならず、最後のウェーブフロントの半分が遊んでしまうことがあります。
原則は一つです。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 は自らを「マーシャリングライブラリ」と呼んでおり、背後に rocBLAS を置くことも cuBLAS を置くこともできます。
| NVIDIA | AMD ネイティブ(roc) | AMD 移植ラッパー(hip) |
|---|---|---|
| cuBLAS | rocBLAS | hipBLAS |
| cuFFT | rocFFT | hipFFT |
| cuRAND | rocRAND | hipRAND |
| cuSOLVER | rocSOLVER | hipSOLVER |
| cuSPARSE | rocSPARSE | hipSPARSE |
| CUB / Thrust | rocPRIM | hipCUB |
| cuDNN | MIOpen | なし |
| NCCL | RCCL | なし |
| CUTLASS | Composable Kernel | なし |
選ぶ基準は明確です。CUDA から移植している最中なら hip 側を使います。呼び出しの形が同じなので、コード変更は最小限で済みます。AMD を主ターゲットとして新規に書くなら roc 側を使います。ラッパーの層が一つなく、AMD 専用機能に直接アクセスできます。
表の下 3 行は対応が一マスずつ空いています。cuDNN、NCCL、CUTLASS には移植用ラッパーがなく、別名の対応物しかありません。これは API の形が違うためソース修正なしには置き換えられないという意味であり、ディープラーニングフレームワークがこの層をそれぞれ吸収してくれるため、ほとんどのユーザーには問題になりません。逆に、これらのライブラリを直接呼び出すコードを抱えているなら、その部分が移植作業の中心的なコストになります。
LLM 推論に限定すると、最近登場した AITER(AI Tensor Engine for ROCm)が重要になります。AMD が LLM 推論用のカーネルを集めたリポジトリで、vLLM の ROCm インストールドキュメントがビルド手順に含めています。NVIDIA 側の FlashInfer や FlashAttention が占める位置に相当します。
移植の経路 — HIPIFY で済むものと済まないもの
ツールは 2 つあります。
| ツール | 方式 | 必要なもの | 性格 |
|---|---|---|---|
hipify-perl | パターン置換 | なし | 速くて粗い。文法が壊れたコードも処理できる |
hipify-clang | Clang AST を解析して再生成 | CUDA のインストールとヘッダー | 正確。ビルド可能なコードである必要がある |
実務では 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 のインラインアセンブリとして新しく書く必要があります。移植作業でもっとも確実に時間を食う項目です。
- ワープサイズの前提。前節のすべてです。自動変換は、32 という数字がワープサイズを指すのか別の意味なのか判別できないため、手を付けません。
- ワープレベルプリミティブの意味の違い。
__shfl_sync系のマスク引数、__ballotの戻り値の幅、明示的な同期のセマンティクスが異なります。名前は移植されますが、意味は検証が必要です。 - CUDA 専用ライブラリの呼び出し。cuDNN、CUTLASS、cuBLASLt のように対応ラッパーがないもの。
- ドライバー API を使うコード。
cuModuleLoadのように低レベルのドライバー API を直接使う部分。 - 性能チューニング用の定数すべて。タイルサイズ、ブロックサイズ、アンロール度、パイプラインステージ数。変換はされますが、最適値は異なります。移植と最適化は別の作業であり、ここを飛ばすと「動くけれど速度は半分」で終わります。
移植後に必ず行うべき検証
# 1. まず数値検証から。性能はその後。
# 特にリダクションを含むカーネルを重点的に見る。
pytest tests/ -k "reduction or attention or norm"
# 2. AMD 側のプロファイラで実際のボトルネックを測り直す。
# NVIDIA でのチューニング結果をそのまま信じてはいけない。
rocprofv3 --stats -- ./my_app
rocprofv3 --kernel-trace -- ./my_app
Nsight Compute に相当する AMD のツールは rocprofv3 と ROCm Compute Profiler です。考え方は同じです。カーネルごとの時間、メモリスループット、キャッシュヒット率、占有率を見ます。
現実に何が障害になるのか
ハードウェアが似ているのに現場で差が出る理由を、3 つに整理します。
1. カーネルエコシステムの蓄積量
これがもっとも大きな項目です。新しいモデル構造や新しい量子化フォーマットが登場すると、そのためのカーネルはほぼ必ず CUDA から先に出てきます。FlashAttention の新しい派生形、新しい MoE ルーティングカーネル、新しい低精度 GEMM、すべてそうです。AMD 対応は数週間から数か月遅れて追いついてきます。
Triton がこの差を実質的に縮めています。Triton リポジトリの third_party には nvidia と amd のバックエンドが並んで入っており、Triton で書いたカーネルは両方でコンパイルできます。以前の記事で扱ったとおり、最近のカスタムアテンションカーネルは Triton で先に出てくる傾向があるため、この傾向が続けば差は縮まり続けます。ただし、手書きの CUDA カーネルに依存する部分には、依然として時差があります。
2. フレームワーク対応の時差とサポートマトリクス
具体的な例を一つ見てみます。vLLM v0.26.0 の ROCm インストールドキュメントには、こう書かれています。
- ROCm 6.3 以上をサポートしており、ビルド済みホイールは 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 のビルド済みホイールは 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 が合理的な選択になる場合は確かにあります。条件は具体的です。
第一に、広く使われているモデルを推論でサービングする場合です。Llama、Qwen、DeepSeek 系の主流モデルを vLLM でサービングする作業は、すでに十分に踏み固められた経路です。検証済みのコンテナを使い、標準的なモデルを載せている限り、上で述べた摩擦の大部分はすでに他の誰かが経験して解決済みです。
第二に、メモリ容量が決定的な場合です。大きなモデルをより少ない数の 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 移植ガイド(ワープサイズの警告を含む): 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/