2006年、NVIDIAがCUDA(Compute Unified Device Architecture)をリリースしたとき、GPUはグラフィックス描画専用の部品から汎用並列計算機へと変貌を遂げた。それ以来、GPU上で動くカーネル(計算の中核を担う関数)を書く手段は、事実上C/C++の拡張であるCUDAか、AMD陣営のHIPに限られてきた。どちらもメモリ安全を保証しない。ポインタの指す先が有効かどうか、スレッド間でデータ競合が起きていないかどうかは、すべてプログラマの責任だ。

この状況は、GPUがAI学習や科学技術計算の中心インフラになった2020年代に入って、いっそう重くのしかかっている。カーネルのバグはセグメンテーションフォルトや未定義動作を引き起こし、大規模クラスタでは数日の計算を無駄にする。NVIDIAのソフトウェアエコシステムは250以上のCUDAライブラリを抱え、「CUDAの堀(moat)」と呼ばれてきた。競合ハードウェアが登場しても、ソフトウェアの成熟度が追いつかなければ乗り換えは起きない。

AD

Rustの所有権モデルがGPUに届かなかった理由

Rustは2015年の1.0リリース以降、コンパイル時の所有権チェックによってメモリ安全とデータ競合の排除を保証する言語として、システムプログラミングの分野で存在感を高めてきた。Lawrence Livermore国立研究所やUniversity of Torontoなど、HPC(高性能計算)コミュニティでも採用が進んでいる。

ところが、GPUカーネルに限れば事情が違った。Rustの所有権システムは単一のCPUスレッドが順序立ててメモリにアクセスすることを前提に設計されており、数千のスレッドが同時に動くGPUの並列実行モデルとは構造的に噛み合わない。既存の試みは三つの方向に分かれていた。

アプローチ 代表例 制約
SPIR-Vベースのコンパイル rust-gpu 汎用ポインタ非対応。グラフィックスから計算への転向中で成熟度不足
ベンダーAPIの薄いバインディング rust-cuda 全カーネルでunsafe必須。メモリ安全の保証なし
特定ベンダーへの密結合 cuda-oxide NVIDIA専用。ポータビリティなし

「安全で、ポータブルで、十分速い」Rust GPUインターフェースは存在しなかった。

LLVM Offloadという共通基盤に乗せる

この空白を埋めるのが、Manuel S. Drehwaldら5名のチームによる論文「GPU Offload in Rust: Portable, Safe, and Fast」(arXiv: 2608.13759)だ。DrehwaldはUniversity of Torontoの博士課程に在籍しつつ、Lawrence Livermore国立研究所(LLNL)でフルタイムのGPU offload開発に従事している。共著者のJohannes DoerfertはLLVM Offloadプロジェクトの主要開発者であり、Alán Aspuru-GuzikはToronto大学で計算化学とAIを率いる。

彼らの手法の核は、Rustコンパイラ(rustc)自体にGPUオフロード機能を組み込み、LLVMのOffload基盤を経由してNVIDIAとAMD両方のGPU向けネイティブコードを生成する点にある。LLVM Offloadは元々OpenMPのGPU実行支援として開発されたが、後にOpenMPから分離され、任意の言語フロントエンドが利用できる汎用基盤として再設計された。C++/Fortranでの実績があるこの基盤にRustを接続することで、ベンダー固有のツールチェーンに依存しない道を開いた。

AD

所有権情報がコンパイラ最適化の「材料」になる

技術的な核心は、Rustの型システムと所有権モデルがLLVMの中間表現(IR)に提供する情報の豊かさにある。

安全なRustの参照(reference)は、デフォルトで$noalias$メタデータ付きでLLVM IRにloweredされる。$noalias$とは「このポインタが指すメモリ領域に、他のポインタからアクセスしない」というコンパイラへの宣言だ。C/C++ではプログラマがrestrictキーワードを手動で付与しなければ得られないこの情報が、Rustでは安全なコードを書くだけで自動的に付く。LLVMバックエンドはこの情報を使って、メモリアクセスの並べ替えや不要なロードの削除といった最適化をより積極的に行える。

データ転送の自動生成も所有権モデルに支えられている。コンパイラはRustのMIR(Mid-level Intermediate Representation)段階でカーネル引数の型、レイアウト、可変性(mutability)を走査し、ホストからデバイスへどのデータを何バイト転送するかを自動的に決定する。不変参照(&T)は読み取り専用としてデバイスに送り、可変参照(&mut T)は書き戻し対象として扱う。プログラマがOpenMPの#pragma omp target mapやCUDAのcudaMemcpyを明示的に書く必要がない。

GPUカーネル内部の安全な並列アクセスには「Region」という抽象化を導入した。数千スレッドが同じスライス(slice)の異なる要素にアクセスする典型的なパターンを、所有権のルールを破らずに表現する仕組みだ。フロントエンド内部では生ポインタ(raw pointer)を使うが、ユーザーが書くコードは安全なRustのままになる。

RAJAPerfでCUDA/HIPと正面比較

評価には、LLNLが開発したRAJAPerfベンチマークスイートを使用した。RAJAPerfはHPCアプリケーションから抽出したループベースの計算カーネルを集めたもので、OpenMP、CUDA、HIPの各実装との性能比較が標準化されている。Drehwaldらはこのスイートの一部を純粋なRustに移植し、AMD MI250X、NVIDIA H100、NVIDIA RTX A2000の3環境で測定した。コンパイラはLLVM 23.1.0-rc1ベースの拡張rustcを使っている。

カーネル実行時間: ほぼ同等、一部で優位

カーネル単体の実行時間では、Rust OffloadはRAJA(C++実装)とほぼ同等の性能を示した。FIRとLTIMESという2つのベンチマークではRustがCUDA比で44%、46%遅いが、これらは数回の乗算/加算しか含まない微小ループであり、コンパイラのアンロール判断の違いが結果に直結する。同じ2ベンチマークで、RustはHIP(AMD)比では15%、32%速い。つまり遅いケースはコード生成の微調整で詰める余地があるとチームは分析している。

全体実行時間: 同期オーバーヘッドが課題

カーネル起動から同期までの全体実行時間では差が広がる。MI250X上ではCUDA/RAJA比で32%高速から43%低速の範囲、H100上では11%高速から46%低速の範囲に分布した。チームはホスト-デバイス間の同期効率に改善の余地があると認めている。

データ転送量と転送時間

H100上の全ベンチマーク合計で、Rustはホストからデバイスへ53回・計423 MBを転送した。RAJAの55回・468 MBより少ない。にもかかわらず転送にかかった時間はRustが46 ms、RAJAが16 msと逆転している。チームはメモリ種別(memory kind)の違いと非同期転送の未実装を原因と推定している。

指標 (H100, RAJAPerf全体) Rust Offload RAJA (CUDA C++)
ホスト→デバイス転送回数 53回 55回
ホスト→デバイス転送量 423 MB 468 MB
ホスト→デバイス転送時間 46 ms 16 ms
デバイス→ホスト転送量 69 MB 99 MB
カーネル実行時間 (中央値付近) ほぼ同等 ベースライン

もう一つ重要な数値がある。データ転送をカーネル起動ごとに繰り返す素朴な実装(インターフェースAの単純適用)は、最適化した実装と比べてMI250X上で400倍以上遅い(比較元・比較先の絶対実行時間は論文の図中にのみ示されており、テキスト中には記載がない)。この差を埋めるために、チームはLLVMのOpenMP-optパスを拡張し、繰り返される自動転送をまとめて排除する最適化のプロトタイプを構築した。

AD

安全なコードが速いコードになりうる構造的理由

この結果が意味するのは、「安全だから遅い」という通念の逆転だ。Rustの所有権システムはプログラマにとっての制約であると同時に、コンパイラにとっての情報源でもある。$noalias$保証が自動的に付与されることで、C/C++では手動アノテーションやプロファイリングに基づくチューニングでしか得られなかった最適化機会が、安全なコードを書くだけで手に入る。

もちろん、これは「Rustなら常にCUDAより速い」という主張ではない。論文が示したのは、LLVMバックエンドが生成するIRの品質がCUDA/HIPコンパイラと競合する水準に到達したという事実であり、差分が残るベンチマークも存在する。チーム自身、FIRとLTIMESの遅延はコンパイラのアンロール判断の違いに起因すると述べ、コード生成の改善で縮められる見込みを示している。

rustcへの統合は進行中、安定版までは道半ば

この研究は学術的な発表だけで完結していない。Drehwaldはrustc本体へのGPU offload機能の統合を並行して進めており、2024年10月にトラッキングIssue(rust-lang/rust#131513)が開設された。2025年半ばにはホスト側のコード生成PRがマージされ、カーネル起動、コンパイル手順の簡素化(最終的にcargoコマンド2回とclang-linker-wrapper 1回まで削減)、CIでのテスト有効化が段階的に進んでいる。2025年後半のRust Project Goalsにも「std::offloadモジュールの完成」が採択された。

ただし、現時点ではnightlyチャンネルの実験的機能であり、安定版に載る時期は明示されていない。Intel GPU対応はLLVM側のバックエンド成熟を待って追加される予定で、Apple Siliconへの言及もある。

残された課題と、この研究が問うていること

未解決の点は少なくない。第一に、ホスト-デバイス間の同期オーバーヘッド。カーネル実行自体は同等でも、全体の実行時間で差がつく主因はここにある。非同期転送とカーネル起動の統合が次の目標だ。第二に、Rustの標準ライブラリ(std)をGPU上で動かす仕組みがまだない。チームはLLVMのlibc-for-gpuプロジェクトをRust向けに移植する方針を示しているが、実装はこれからだ。第三に、複数GPUを跨ぐマルチデバイス環境への対応も計画段階にとどまる。

再現性の面では、この論文はarXiv上のプレプリントであり、ピアレビューを経ていない。RAJAPerfのサブセットを用いた評価であり、全カーネルカテゴリを網羅したわけではない。

それでも、この研究が提起する問いはGPUコンピューティングの根幹に触れている。2006年以来、GPU上で動くコードを書くということは、メモリ安全の保証を放棄することを意味してきた。CUDAが築いた20年のエコシステムは、その放棄の上に成り立っている。もし安全な言語で書いたカーネルが性能面で遜色ないとすれば、GPUソフトウェアの設計思想そのものが変わる可能性がある。次の焦点は、この結果がより多様なアプリケーションとハードウェア構成で再現されるかどうか、そしてrustcの安定版に載るまでに同期オーバーヘッドをどこまで削れるかにかかっている。