ハクソク

世界を動かす技術を、日本語で。

CUDA-oxide: Nvidiaの公式RustからCUDAへのコンパイラ

4時間前原文(nvlabs.github.io)

概要

cuda-oxide は、Rustコードを直接CUDAのPTXにコンパイルする実験的なコンパイラ。 Rust の所有権やトレイト、ジェネリクスの知識が前提。 DSLやバインディング不要 で、純粋なRustでGPUカーネルを記述可能。 v0.1.0はアルファ版 であり、バグやAPI変更の可能性あり。 Rustの型安全性 とGPUプログラミングの統合が特徴。

cuda-oxide 概要

  • cuda-oxide は、RustからCUDA用GPUカーネルを記述・コンパイルするための 実験的コンパイラ
  • 標準的なRustコード を直接PTXに変換、 DSLや外部言語バインディング不要
  • SIMT（Single Instruction, Multiple Threads） に最適化された設計
  • Rustの型システム・所有権モデル をGPUカーネルにも適用
  • 安全性 を重視しつつも、GPU特有の注意点あり（詳細はSafety Model参照）

前提知識

  • Rustプログラミング言語 の理解（所有権、トレイト、ジェネリクス）
  • 非同期プログラミング （async/.await、tokio等）の基礎知識
  • 参考資料 ：The Rust Programming Language、Rust by Example、Async Book

プロジェクトの現状

  • v0.1.0リリース は初期アルファ版
    • バグ、未実装機能、API変更の可能性
    • ユーザーからのフィードバックを歓迎

クイックスタート

  • 必要なクレートのインポート
    • cuda_device::{cuda_module, kernel, thread, DisjointSlice}
    • cuda_core::{CudaContext, DeviceBuffer, LaunchConfig}
  • カーネル定義例 （vecaddカーネル）
    • #[cuda_module]でモジュール宣言
    • #[kernel]でカーネル関数定義
    • thread::index_1d()でスレッドインデックス取得
    • DisjointSliceで出力配列の安全な参照取得
  • ホスト側の実行例
    • CudaContextでデバイス初期化
    • DeviceBufferでデータ転送
    • kernels::loadでカーネルロード
    • module.vecaddでカーネル実行
    • 結果をto_host_vecで取得・検証
  • ビルド＆実行
    • cargo oxide run vecaddでビルド・実行
  • #[cuda_module]の役割
    • デバイス用バイナリをホストバイナリに埋め込み
    • kernels::load関数と各カーネルのlaunchメソッドを自動生成
    • 低レベルAPI（load_kernel_moduleやcuda_launch!）も利用可能

cuda-oxideの特徴

  • Rust on the GPU
    • Rustの型安全性・所有権モデルによる 安全なGPUカーネル記述
    • 安全性モデルの詳細はドキュメント参照
  • A SIMT Compiler, Not a DSL
    • DSLではなく、rustcのカスタムコード生成バックエンド で純粋なRustをPTXへ
  • Async Execution
    • GPUワークを 遅延DeviceOperationグラフ として構築
    • ストリームプールでスケジューリング
    • .awaitで非同期実行結果を待機

まとめ

  • cuda-oxide は、Rustエコシステムにおける GPUプログラミングの新たな選択肢
  • 純粋なRustコード で安全かつ効率的なGPUカーネル開発を実現
  • 今後の発展とフィードバック に期待

Hackerたちの意見

すごいね！カスタムCUDAカーネルとhttps://crates.io/crates/cudarcをずっと使ってきたけど、これ、ほぼそのまま置き換えられそうだね。特にビルド時間がどうなるか気になるな。ほとんどのRustのCUDAクレートはCMakeやnvccを呼び出すから、コンパイルがめっちゃ遅くなるんだよね。偶然にも、先週ビルド時間をプロファイリングしてたら、sccacheみたいなツールがアーティファクトをキャッシュすることでリビルド時間を劇的に短縮できることがわかったんだけど、やっぱり高価なカスタムnvccの呼び出しにはお金がかかるよね（例えば、hugging faceのcandleはカーネルコンパイルでカスタムnvccコマンドを呼び出してる）：https://arpadvoros.com/posts/2026/05/05/speeding-up-rust-whi...

└

他の人もcuda-oxideがcudarcのほぼ置き換えになると思ってる？それはすごいけど、個人的には補完的ではないと思うな。cuda-oxideの特徴が何なのか気になる。nvの管理下にあることを除いて。

└

Cudarc最高！> ほとんどのRustのCUDAクレートはCMakeやnvccを呼び出すから、コンパイルがめっちゃ遅くなるんだよね。私はこの問題に遭遇したことがないけど、ビルドスクリプトを扱うために作ったcuda_setupクレートを見てみて。これは、ファイルが変更された時だけ再コンパイルするシンプルなbuild.rsで、コンパイル時間はめっちゃ短いよ（RustのCPU側コードと比べても）。

Rustのメモリモデルをどう扱ったのか、すごく興味あるな。CUDAのセマンティクスにうまくマッピングできないかもしれないし。CUDA C++と比べてどんな違いがあるのか、Rustの型システムがCUDAにもっと安全性をもたらせるのか気になる。GPUカーネルを書くのは本質的に危険だと思うし、ハードウェアの動きのせいで安全な言語を作るのは難しいから、常にハイパー最適化してるのもあるしね。

└

目的によると思う。パターンマッチング脳の私が言うには、これに対処することに興味がある人はいると思う。Rustでアプリを書く立場から言うと、GPUコンピュートを使いたい時もあるけど、あんまり気にしないかな。メモリモデルやオーナーシップモデルを低フリクションで活用できるならそれでいいし、高フリクションな体験になるなら、そういうやり方は避けたい。基本的には、今のCudarcのやり方が基準だと思う。メモリ管理はあんまり関わってないし、FFIをラップする命令型の構文と、カーネルが変わった時にnvccを呼び出すためのビルドスクリプトの数行だけだと思う。

└

これはドキュメントで詳しく説明されてるよ。安全なレイヤー、ほぼ安全なレイヤー、そしてunsafeなレイヤーがある。一部のもどかしさは、安全だけど並列処理を行うために必要で、RustのSend/Syncモデルに簡単には収まらなかったんだ。

└

見える主要な4つは：1. use-after-free、ドロップセマンティクス vs 手動のcudaFree 2. cuda_launch!を使ったカーネル引数の強制、一方でCPPのvoid引数はポインタの配列で、カウントだけを検証する 3. エイリアスの可変書き込み。例えば、CPPでは同じiで複数のスレッドがout[i]に書き込むことができて、これがコンパイルされる。でも、DisjointSliceとThreadIndexには公開されたコンストラクタがない（参照： https://github.com/NVlabs/cuda-oxide/blob/2a03dfd9d5f3ecba52...）し、index_1d、index_2d、index_2d_runtimeのAPIしか使えない 4. std::stringや他のPODをcuda memcpyして、状態を「壊して」使えなくすることができると思う。ここでは、DisjointSlice、スカラー、クロージャしか受け付けない（https://nvlabs.github.io/cuda-oxide/gpu-programming/memory-a...）。でも、細かいところはこのセクションにあるよ。 https://nvlabs.github.io/cuda-oxide/gpu-safety/the-safety-mo... * https://nvlabs.github.io/cuda-oxide/gpu-programming/memory-a... 編集：とはいえ、これがすべてをカバーするわけじゃないけど、.cuファイルのUBに対するガードレールをかなり増やしてくれる感じ。

めっちゃいいね！Cudarc（カーネル）とFFI（cuFFT）でやってたよ。バイト配列とRustのデータ構造の間で手動で[デ]シリアライズしてる。これでフリクションが減るといいな！

Re: Rust（と「安全な」プログラミング言語）。NVIDIAのSpark/Adaの利用について詳しい情報を持ってる人いる？私が見つけたのは以下に載ってることだけだよ：https://www.adacore.com/case-studies/nvidia-adoption-of-spar...

└

2020年のNVIDIAによるプレゼンテーション「埋め込みソフトウェアの安全性とセキュリティの未来を守る」の録画があるよ。AdaCoreカンファレンスでのものだね。: https://www.youtube.com/watch?v=2YoPoNx3L5E

└

彼らは昨年のDEF CONで詳しいトークをしてたよ。: https://www.youtube.com/watch?v=KhWtkZmOPn4

誰か、これがホストとデバイス間で構造体を共有できるか知ってる？今のところ、既存のRust/CUDAのワークフローで一番足りない部分だよね。（あと、シリアル化やバイトの壁もあるし）

CUDA用のRustで気になってるのは、Rustが追加するちょっとしたオーバーヘッドなんだ。通常は無視できるけど、ここでは影響があるかも。例えば、配列の境界チェックとか。これが原因で追加のレジスタが使われて、カーネルの並行性が下がることはないかな？

そうだね、ステンレスがあるから、Linuxのコードは錆びないってことだよね。次は、リンを作ってくれる人が必要だな。そうすれば、錆びたコードを古い鉄に変えられる。GPLファンはRustボックスやステンレスマシン、未来に対応した鉄の作業馬を使えるようになる。ソフトウェアは3つのエディションで提供できるよ。錆びたことのないステンレスドライバー、既存のコードでRustを使った酸化ドライバー、そして新しいリンのツールを使ってRustをCに戻したIronエディション。多様性が私たちの強みになるかも。錆びたC/C++コードを作るのは、酸洗いって呼べるかもね。

└

誰か必要だよね > それでGPLファンがチェックアウトするってことか。

Slang[0]にとって何を意味するんだろう。おそらく、みんながもっとモダンな言語でGPUプログラミングをしたいってことだよね。でも、今はRustを使えばいいんじゃないかな…（免責事項：私はSlangが大好きです。）[0]: https://shader-slang.org/

すごい！

TileLang https://github.com/tile-ai/tilelang とか、Tile Kernels https://github.com/deepseek-ai/TileKernelsみたいなものが、いつかCUDAを時代遅れにするだろうね。