【10月5日公開】ROCm 10.1のGPU直結I/Oを分解。16MiBの中継バッファが残る条件と、今すぐ試す3手順

Posted at 2026-10-11

AMD GPUでモデルやチェックポイントを読み込んでいるなら、そのデータがホストRAMを通っていないか、一度確かめてほしい。

ROCm 10.1では、ストレージからGPUメモリへ直接データを送る経路が非同期I/Oにも広がった。AMDが10月5日に出した技術解説の主役は、演算器にデータを届ける仕組みだ。

HIP streamに読み書きを投入し、ホスト側の中継コピーを省ける。
直結できない要求には、ホストメモリを経由する別経路もある。
非同期の互換経路は、streamごとに既定で16MiBの中継バッファを再利用する。

それが、ROCm 10.1に含まれる「hipFile 0.5.0」だ。16MiBは環境変数の仕様に載った実際の既定値である。

たとえば64GiBのファイルをホストRAM経由でGPUへ転送すると仮定すると、ホスト側の論理的な読み書きは合計128GiBになる。これは経路を理解するための例示計算であり、実測したDRAM転送量ではない。中継を省く意味は、この往復を作らないことにある。

「Asyncを呼んだ」「読み込みに成功した」だけでは、データが直結経路を通ったとは言えない。

結論から言うと

1. まず、実際のI/Oをfastpathとfallbackに分けて観測すること。
対応環境の確認にはais-check、実行中の経路別転送量にはais-statsを使う。

2. 16MiBはファイルサイズの上限ではなく、非同期fallbackの中継バッファの既定容量だ。
大きな転送を分割して運ぶため、ファイル全体がホストRAMに収まる必要はない。

3. ROCm 10.1の非同期fastpathは、失敗後にfallbackへ自動再試行しない。
互換経路が存在することと、直結転送の失敗を救済できることは別の仕様だ。

生成イラスト。琥珀色の中継槽を通る経路と、そこを省く経路を表した比喩であり、実機の配線図ではない。

資料確認日は2026年10月11日。新規性は10月5日のAMD解説、版と制約はROCm Core SDK 10.1.0のリリースノートに基づく。細部は確認時点で0.5.0と表示されたhipFile公式資料を参照した。対象は同リリースで示されたLinux・Instinctの構成である。実機ベンチマークは未実施、以下のHIPコードとコマンドも未実行だ。64GiBの例は独自の単純計算で、性能値ではない。

数字で見るhipFileのデータ経路

公式のバックエンド仕様、環境変数、非同期API、公式サンプルを同日に照合して整理した。

項目	値・仕様	読み違えると困る点
今回の対象版	hipFile 0.5.0	ROCm全体の版番号とは異なる
主な転送経路	2種類	fastpathとfallback
fastpathの適格スコア	100	帯域のスコアではない
fallbackの適格スコア	1	処理できないバックエンドは−1
非同期fallbackの中継バッファ	streamごとに既定16MiB	`0`を指定しても既定容量を使う
`ALLOW_COMPAT_MODE`の既定	`true`	経路選択時にfallbackを許す
`FORCE_COMPAT_MODE`の既定	`false`	直結が可能なら候補になる
非同期APIの1回の最大I/O量	2³¹−システムページサイズ byte	ページサイズ4KiBなら2,147,479,552 byte
公式の往復サンプル	4 stream × 1MiB = 4MiB	大容量ストレージの性能試験ではない
ページ整列を申告するstreamフラグ	4KiB整列	あらゆるファイルシステムの必須整列値とは限らない

ここまでは公開仕様の話。面白いのはその先だ。

何を省くのか。CPUではなく、ペイロードの中継を省く

AMD Infinity Storageの説明では、fastpathはストレージとGPUメモリの間をDMAで転送する。CPUは要求の投入や完了処理に関わるが、ペイロードをホスト側の中継領域へコピーしない。

fallback: ストレージ → ホスト中継バッファ → GPUメモリ
fastpath: ストレージ ─────────────────→ GPUメモリ
          CPUは要求の投入・完了処理を担当

「CPUを一切使わない」わけではない。ストレージ上のデータをGPUへ渡す際に、ホストRAMへ一度置く工程をなくす設計だ。

先ほどの64GiBの例で数えたのは、中継領域への64GiBの書き込みと、そこからの64GiBの読み出しである。ページキャッシュやキャッシュ階層の挙動まで再現したモデルではない。この128GiBを省ける可能性と、処理時間が半分になることは同義ではない。 ストレージ、PCIe、演算のどこが律速するかで効果が変わる。

100対1で経路を選ぶ。ただし「適格」と「完了」は違う

要求を受けると、各バックエンドは処理可能性をスコアで返す。直結できるfastpathが100、互換経路が1、拒否が−1。高いスコアの経路が選ばれる。

判定に使うのは、ファイルの開き方、GPUバッファの種類、ファイルシステム、整列条件などだ。通常のローカルNVMe構成では、O_DIRECT、device memory、対応するext4/XFSなどが問題になる。

一方、選ばれた後にもデバイス探索やP2P DMAが失敗する余地がある。ここで10.1.0のリリースノートは、次のように明記している。

async fastpath operations do not currently retry on the fallback backend if the fastpath fails in this release.

直結が無理なら最初からfallbackを選ぶ場合と、直結を始めてから失敗する場合を分けて考えること。

見つけた証拠	言えること	言えないこと
fastpathの適格スコア100	その要求を受け付ける条件がそろった	転送が必ず完了する
`ais-check`の終了コード0	対応するボリュームが少なくとも1つある	対象ファイルも、そのGPUも直結できる
`ais-stats`の経路別byte数	どの経路でどれだけ転送したか	その経路が常に最速である
非同期処理後のbyte数と内容の一致	要求したデータの転送結果を確認できた	ホスト側の中継がなかった

Asyncの役目は順序をつなぐこと

非同期APIは、HIP streamへI/Oを並べる。同じstream上では、読み込みの後に投入した処理がその読み込みより後に動く。別々のnon-blocking stream間には、その順序関係は自動では生まれない。

以下は非同期APIの引数を使った順序説明用の擬似コードである。未実行。ファイル・GPUバッファ・streamの登録処理とエラー処理の実装は省略した。

// 擬似コード。サイズ・offset・完了byte数は同期完了まで生存させる。
size_t bytes = 1 << 20;
hoff_t file_offset = 0, buffer_offset = 0;
ssize_t completed = 0;

auto status = hipFileReadAsync(
    file_handle, device_buffer, &bytes,
    &file_offset, &buffer_offset, &completed, stream);
check_submission(status);                 // 独自の擬似関数
check_stream(hipStreamSynchronize(stream)); // 独自の擬似関数
require(completed == bytes);              // 独自の擬似関数
// 完了とbyte数を確認してから、このデータを利用する。

このコードから分かるのは、要求の投入とデータの利用可能化を分ける必要があることだ。経路がfastpathだったかどうかは、別途統計で確認する。

資料から説明できるのは経路とAPIの契約までで、手元の構成で得られる帯域や短縮時間は実機測定が必要だ。

今すぐ試すべき3手順

前提は、対応するLinux・Instinct環境にhipFileとツールが導入済みであること。導入先は公式インストール手順に従って選ぶ。以下の/path/to/...やアプリ名は、自分の検証環境の値へ置き換えること。

1. 対象ファイルの置き場所を調べる

公式の互換性確認手順のコマンドで、ボリュームごとのHIPFILE列を読む。

# 未実行。正確な診断にはroot権限が必要。
sudo ais-check -v

終了コードだけを見るのは早い。モデルファイルを置くマウントポイントの行でyesになっているか、kernel・HIP runtime・amdgpuの条件もそろうかを確認する。

2. 4MiBをGPUへ送り、ファイルへ戻す

公式の非同期サンプルは4つのstreamを使う。各streamで1MiBを読み、そのstreamで書き戻す。次のコマンドは、対応版のソース一式から公式のCMake手順でサンプルをビルドし、実行ファイルのあるディレクトリに移動してから使う。

このサンプルは入力ファイルも生成し、既存内容を置き換える。実際のモデルファイルを渡さないこと。 検証したいボリューム上の専用ディレクトリを使う。

# 未実行。両方とも検証用の新しいファイルパスへ置き換えること。
./roundtrip-async-multi-stream-registered \
  /path/to/readfile /path/to/writefile 0

サンプルのbyte数・ハッシュ検証が通ることと、直結で動くことは別だ。次の手順を重ねる。

3. 同じ処理のfastpathとfallbackを数える

公式の統計ツールはアプリを起動して、そのI/Oを経路別に集計できる。./your_hipfile_appはhipFileを使う検証アプリへ置き換える。

# 未実行。検証アプリの必要な引数も末尾へ付けること。
HIPFILE_STATS_LEVEL=1 ais-stats ./your_hipfile_app

終了後のレポートで、fastpathとfallbackそれぞれの読み込みbyte数を比較する。fallbackが増えていたら、その時点で「GPU直結の実験」と呼ぶのをやめ、置き場所や要求条件を調べる。

統計収集には負荷があるため、性能を比較する両条件では収集設定もそろえる。転送サイズだけでなく、エラー数と非整列I/Oの件数も合わせて読む。

あなたの環境では、ファイルの置き場所、整列、GPUの仮想化のどれが最初の壁になりそうだろうか。あなたの予想もコメントで教えてほしい。

知らないと損する4個の罠

罠1：XFSなら、下に何があっても直結できると思う

制約一覧は、LVM/device-mapper、multipath、dm-crypt、MD RAID、loopbackといった中間ブロック層を挙げている。ファイルシステムの種類だけをそろえても、その下の経路が条件を満たすとは限らない。

対策：対象マウントの背後にあるブロック層まで確認すること。暗号化や冗長化を外す判断を、速度だけで決めないこと。

罠2：4KiB単位なら何でもよいと思う

HIPFILE_STREAM_PAGE_ALIGNED_INPUTSは4KiB整列を申告するフラグだ。一方、バックエンドが検査する条件では、statx()が返すoffset側とメモリアドレス側の整列値を区別する。両者が同じとは限らない。

サイズだけ整えても、file offsetやbuffer addressが外れていれば意図した経路を使えない。

対策：サイズ・file offset・buffer offset／addressをまとめて検査し、フラグで事実と異なる申告をしないこと。

罠3：互換モードを許可したら、非同期の失敗も救われると思う

HIPFILE_ALLOW_COMPAT_MODE=trueでも、0.5.0の既知の問題には非同期fastpathの自動再試行がないと明記されている。デバイス探索やP2P DMAが失敗すれば、その非同期操作は失敗する。

対策：投入時の戻り値、stream完了、転送byte数を確認すること。互換経路を検証したい場合は、後述の強制設定で最初から選ぶこと。

罠4：GPUが見え、診断が0終了なら十分だと思う

同じ既知の問題では、SR-IOVのVFでfastpathは無効になる。PFだけが対応対象だ。ais-checkのVF警告は終了コードを変えない。

つまり「GPUを認識できた」と「GPUとストレージが直結できた」は別の確認になる。

対策：amd-smiでVFかどうかを確認し、診断の警告と経路別byte数を読むこと。

比較：自動選択・直結必須・互換強制の3つを使い分ける

別の製品へ移る前に、同じバイナリの経路だけを変えて比較できる。以下は10月11日に確認した公式環境変数から組み立てた表だ。列名のALLOWとFORCEは、それぞれHIPFILE_ALLOW_COMPAT_MODEとHIPFILE_FORCE_COMPAT_MODEを指す。

比較条件	ALLOW	FORCE	経路選択	向く用途
通常の自動選択	`true`	`false`	fastpathを優先し、不適格なら受付可能なfallback	通常設定での経路観測
直結必須	`false`	`false`	fallbackを許さず、直結不可なら失敗	中継コピーを隠さない検証
互換強制	`true`	`true`	最初からfallback	ホスト経由との比較対照

両経路とも要求を受け付けられなければ失敗する。自動選択の行も、非同期fastpathの途中失敗を救済する設定ではない。ALLOW=falseとFORCE=trueの組み合わせは、fallbackを強制しながら禁止するためI/Oが失敗する。

表から読み取れること：

直結の検証では、成功率より経路を固定することが先だ。 自動fallbackが混ざった比較では、何を測ったかが変わる。
互換強制は、消すべき機能とは限らない。 環境制約のある構成や比較対照として役割がある。
16MiBを大きくしてもfastpathにはならない。 中継容量の調整と経路の選択を混同しないこと。

検証用アプリを直結必須で起動する設定は、次の形になる。未実行。このコマンドは公式環境変数とais-statsの起動方法を組み合わせたものだ。

HIPFILE_ALLOW_COMPAT_MODE=false \
HIPFILE_FORCE_COMPAT_MODE=false \
HIPFILE_STATS_LEVEL=1 \
ais-stats ./your_hipfile_app

比較時は上の2変数だけを表どおり変更し、ファイル、GPU、I/Oサイズ、stream数、測定区間を固定する。最初に経路別byte数を確かめ、その後で帯域と待ち時間を比べる。

教訓：省けたコピーを、成功ログから推測しない

✗ "AsyncならGPU直結" → 非同期fallbackにも中継バッファがある
✗ "16MiBより大きなファイルは読めない" → 中継容量と転送全体の大きさは別
✗ "診断の終了コード0なら完成" → 対応ボリュームの存在と実際の要求は別
✗ "失敗すれば自動でCPU経由に戻る" → 10.1の非同期fastpathは再試行しない
✗ "コピーが減れば必ず2倍速い" → 残った律速と実測時間を確認する必要がある

hipFileの面白さは、アプリが欲しい「GPUメモリ上のデータ」を、ホストRAMへの中継なしで届けられる点にある。その道が開くと、演算器の最適化だけでは届かなかった待ち時間に手を入れられる。

ただし、API名も終了コードも、データが走った道の証拠にはならない。今日、検証用の読み込みを1つ選び、fastpathとfallbackの転送byte数を並べてみてほしい。