--- title: "cuda-oxideのDisjointSliceとは｜RustでGPU並列書き込みを型安全にする仕組み" url: "https://www.issoh.co.jp/tech/details/12684/" published: 2026-06-23 updated: 2026-06-23 categories: ["アーキテクチャ"] publisher: "株式会社一創" --- # cuda-oxideのDisjointSliceとは｜RustでGPU並列書き込みを型安全にする仕組み cuda-oxideのDisjointSliceは、GPUカーネルで複数のスレッドが同じ出力先へ書き込むデータ競合を、Rustの型システムだけでコンパイル時に防ぐスライス型です。本記事では、2026年5月7日にv0.1.0が公開されたNVlabs版cuda-oxideにおけるDisjointSliceの役割、ThreadIndexとIndexSpaceによる安全機構、`get_mut_indexed`や`get_mut`といったAPIの使い分け、そして型が守らない`&mut [T]`引数の落とし穴までを一次情報に基づいて整理します。生ポインタを使うCUDA C++との違いや、alpha段階での採用判断の基準も扱います。 ## まとめ：DisjointSliceがGPUのデータ競合を型で防ぐ核心と注意点 結論を先に示します。DisjointSliceは証人型ThreadIndexと組で使い、「1スレッドが1要素だけに書く」という並列出力を、呼び出し側にunsafeを書かずに型で保証します。これがcuda-oxideの安全モデルの中核（Tier1）であり、生ポインタを全スレッドが添字でアクセスするCUDA C++のデータ競合を、コンパイル時に消し去ります。 ただし万能ではありません。`&mut [T]`をカーネル引数に取ると型の保護は効かず、共有メモリ・ワープ命令・アトミックはTier2のunsafeが必要です。cuda-oxide本体は現在もalphaで、Rust nightlyとLLVM 21+を要します。研究・試作で単一言語のRustでGPUを書く用途なら早期参入の好機ですが、本番のML学習や基幹HPCは当面PyTorch/JAXや成熟したC++ CUDAに置くのが堅実です。 ## cuda-oxideでのDisjointSliceの役割と生ポインタ方式との根本的な違い DisjointSliceを理解するには、まずそれが乗っているcuda-oxide本体と、置き換える対象である生ポインタ方式の危うさを押さえる必要があります。 ### RustをPTXへ直接コンパイルするcuda-oxide本体とv0.1.0の位置づけ cuda-oxideは、NVIDIA Labs（NVlabs）が2026年5月7日にv0.1.0を公開した、Rustのソースコードを直接PTX（NVIDIAのGPU向け中間表現）へコンパイルする実験的なコンパイラです。DSLや外部言語バインディングを介さず、`#[kernel]`を付けた通常のRust関数をそのままGPUカーネルにします。内部のパイプラインはRustソース→rustc\_public（Stable MIR）→dialect-mir→mem2reg→dialect-llvm→LLVM IR→PTXで、中間表現にはMLIR風のRust製フレームワークPlironを使います。DisjointSliceはこのcuda-oxideが提供する型の一つで、デバイス側クレート`cuda_device`に含まれます。公開時のv0.1.0はalphaで、以後v0.2系（2026年6月時点で最新はv0.2.1）へ更新されましたが、いずれもバグ・未実装・API破壊が前提のalpha段階です。 ### NVlabs版とcrates.io旧ラッパーで混同しやすい同名cuda-oxideの区別 検索時に注意したいのが、cuda-oxideという名前が二つの別物に使われている点です。本記事が扱うのはNVlabsのcuda-oxide、すなわちRustを直接PTXへ変換するコンパイラです。一方、crates.ioに登録された旧来の「cuda-oxide」は、CUDAライブラリをRustから扱う高レベルラッパーで、こちらはRustをPTXへコンパイルしません。DisjointSliceが登場するのは前者だけで、後者には存在しない型です。NVlabs版は`cargo +nightly-2026-04-03 install --git https://github.com/NVlabs/cuda-oxide.git cargo-oxide`のようにGit経由で導入し、crates.ioの同名パッケージとは切り離して扱います。 ### 生\_\_global\_\_ポインタが二重書き込みを許すCUDA C++の構造的危険 CUDA C++で並列出力を書く定番は、生の`__global__`ポインタを全スレッドが添字でアクセスする方法です。この方式自体には、二つのスレッドが同じ位置へ書き込むのを止める仕組みがありません。添字計算を一つ誤れば、複数スレッドが同一要素を上書きするデータ競合になります。生ポインタの誤用は、[ダングリングポインタやNULL参照によるセグメンテーションフォルト](https://www.issoh.co.jp/tech/details/7982/)と同種のメモリ不正を、GPU上で大量のスレッドにわたって再現させます。ポインタそのものの基礎は[ポインタの基本概念](https://www.issoh.co.jp/tech/details/4290/)を参照してください。 ### 可変スライスを包みTier1の安全を担うDisjointSliceの役割 DisjointSliceは、この生ポインタ方式の安全な代替です。可変スライスを包み、自分と同じIndexSpace（索引空間）を持つThreadIndexを通じてしか書き込みを許しません。各スレッドが必ず一意の要素にアクセスすることを型で保証するため、呼び出し側にunsafeを書かずにデータ競合のないカーネルが書けます。cuda-oxideはカーネルの安全性をTier1〜3に分け、DisjointSliceとThreadIndexの組み合わせは「型システムが誤用を防ぐTier1（unsafe不要）」に位置づけられます。vecaddやGEMMの素朴な実装など、多くのアプリケーションカーネルはこのTier1に収まります。 ## ThreadIndexの証人型とIndexSpaceがデータ競合を封じる型の仕掛け 「型で安全」と言える根拠は、証人型ThreadIndexとIndexSpaceという二つの型にあります。仕掛けを分解します。 ### 公開コンストラクタを持たないThreadIndex証人型の唯一の生成経路 DisjointSliceの安全性を支えるのが`ThreadIndex<'kernel, IndexSpace>`です。これはusizeを包む不透明な「証人（witness）」型で、公開されたコンストラクタを持ちません。手に入れる唯一の方法は、ハードウェア組み込み変数（`threadIdx`・`blockIdx`・`blockDim`）から値を導く信頼された関数（`index_1d`や`index_2d::`）を呼ぶことだけです。これらの組み込み変数とスレッド・ブロック・グリッドの階層は[CUDAの基本的な仕組み](https://www.issoh.co.jp/tech/details/4902/)で整理しています。1Dグリッドなら`blockIdx.x * blockDim.x + threadIdx.x`が全スレッドで一意になるため、証人は無条件にスレッドごとに異なる値を表します。 ### !Send等のマーカーと’kernelライフタイムが封じる証人の使い回し 証人型をスレッド間で使い回せないことも、安全性の柱です。ThreadIndexは`!Send + !Sync + !Copy + !Clone`であり、`'kernel`ライフタイムはマクロが挿入するスタックローカルなスコープに縛られます。このため、あるスレッドが自分の索引を共有メモリに置いて隣のスレッドに拾わせる「証人のロンダリング」ができません。証人はカーネル本体より長く生存することもできません。ハードウェアが各スレッドに一意な索引を渡し、型システムがその索引を持ち出し不能にする。この二段構えで、別スレッドの索引を使って他人の領域へ書き込む経路を塞いでいます。 ### Index2Dのストライド不一致をコンパイルエラーにする型一致 2次元カーネルでは、IndexSpaceの型がストライド（行幅）の取り違えを防ぎます。`index_2d::()`は行幅Sについてconstジェネリックで、戻り値は`ThreadIndex<'kernel, Index2D>`です。`DisjointSlice>`は`Index2D<128>`の証人しか受け取らないため、`Index2D<256>`の証人を渡すコードはコンパイルエラーになります。異なるストライドの証人を同じスライスに混在させるというデータ競合の温床が、型エラーとして弾かれるわけです。GEMMやステンシルのように行幅をコンパイル時に固定できる場合は、この定数ストライド形が安全な既定になります。 ### 一意な索引と境界検査が競合をコンパイル時に消す4つの根拠 DisjointSliceの安全が成り立つ根拠は、四つの事実の組み合わせとして整理できます。第一に`index_1d()`はハードウェア保証によりスレッドごとに一意の値を返します。第二に`get_mut()`は境界検査つきで、範囲外のスレッドには`None`を返します。第三にIndexSpaceの型が各証人をレイアウトに結びつけ、ストライド不一致を型エラーにします。第四に証人は持ち出し不能で、共有メモリ経由のロンダリングを許しません。借用チェッカはスレッドごとに単一の`&mut T`を見て、ハードウェアは索引が互いに素であることを保証する。この二つを型システムが結びつけているわけです。 ## get\_mut\_indexedとget\_mutの使い分けと1D・2D索引の書き方 実装で迷うのは、どのメソッドと索引関数を選ぶかです。次元と索引の再利用有無で機械的に決められます。 ### 既定で使うget\_mut\_indexedの一発取得とNoneが弾く2つの条件 1Dカーネルの既定は`get_mut_indexed()`です。証人の生成と`&mut T`への解決を一度に行う一発取得形で、`if let Some((out, idx)) = out.get_mut_indexed()`のように書きます。`None`が返るのは二つの条件、すなわちグリッド外のスレッド（2Dで`col`が行幅以上になる場合など）と、スライス範囲外の索引のときです。どちらも手動チェックなしに弾けるため、CUDA C++で書いていた`if (col < n)`のようなガードを型側に肩代わりさせられます。索引そのものが要らない単純な変換処理では、この一発形が最短で安全です。 ### 複数スライスの添字演算で必要になるget\_mutの2段階の書き方 索引の値を別の計算に使いたいときは、2段階形を選びます。`let idx = thread::index_1d();`で証人を取り、`out.get_mut(idx)`で書き込み先を得る形です。複数の入力スライスに同じ索引でアクセスする場合や、`idx.get()`で得た添字を使って算術を行う場合に向きます。vecaddで`let i = idx.get(); *c_elem = a[i] + b[i];`と書く実装が典型例です。一発形のget\_mut\_indexedと2段階形のget\_mutは、索引を後段で再利用するかどうかで選び分けます。 ### index\_1dとindex\_2d::で書き分ける1Dと定数ストライド2D 次元ごとに索引関数を使い分けます。1Dは`index_1d()`で`Index1D`の証人を、定数ストライドの2Dは`index_2d::()`で`Index2D`の証人を得ます。`index_2d_row()`や`index_2d_col()`は素のusizeを返すだけで証人にはならず、これ単体ではDisjointSliceの添字に使えません。行と列を算術に使いつつ、書き込みは証人経由にするという分担です。スライス側も`DisjointSlice>`のようにストライドを型に載せ、`get_mut_indexed()`で範囲内が保証された要素だけを書き換えます。 ### 実行時ストライド向けのindex\_2d\_runtimeとunsafeの扱い 行幅が実行時にしか決まらない場合は、`unsafe`な`thread::index_2d_runtime(n)`と`Runtime2DIndex`を使います。unsafeが付くのは、「同じカーネル内の全スレッドが同じ`n`を渡す」という、型では証明できない約束を呼び出し側が引き受けるからです。コンパイル時に行幅を固定できるなら`index_2d::()`が安全で優先されます。縮約のようにワープ内の1スレッドだけが結果を書くパターンでは、`get_unchecked_mut`（unsafe）が抜け道になります。索引が一意であることを自分でSAFETYコメントとして証明する点が、型に任せる通常形との違いです。 ## DisjointSliceが保証しない範囲とmut参照スライス引数の落とし穴 DisjointSliceの守備範囲には明確な限界があります。ここを曖昧にせず、頼ってよい場面と頼るべきでない場面を切り分けます。 ### 型では守られない&mut \[T\]引数というcuda-oxide最大の落とし穴 cuda-oxideで最も注意すべき落とし穴は、`&mut [T]`をカーネル引数に取る場合です。マクロは現状この型を受け付けますが、実行時のレイアウトでは全スレッドが同じ背後ポインタを見るため、二つのスレッドが`&mut data[i]`経由で書き込むと、DisjointSliceが防ぐはずのエイリアシングがそのまま起こります。型システムはこれを止めません。公式の安全性モデルもこれを次に塞ぐべき未解決のギャップと位置づけ、「`&mut [T]`を取るカーネルは全行がunsafeだと思って扱え」と明言しています。可変出力は必ずDisjointSliceで受けるのが鉄則です。 ### 共有メモリ・ワープ命令・アトミックでunsafeが要るTier2の境界 1スレッド1要素から外れる処理は、Tier2の限定的なunsafeに入ります。共有メモリ（`static mut`の`SharedArray`）はブロック内の全スレッドから見えるため借用チェッカが扱えず、`thread::sync_threads()`による同期を前提にunsafeで書きます。ワープシャッフル（`shfl_sync`など）はスレッドの収束を型で検査できず、発散したスレッドにマスクを渡すと、クラッシュせず静かにハングする最悪のバグになります。アトミックは生ポインタからの構築時だけがunsafeで、操作自体は安全に呼べます。Tier2のunsafeが局所的で監査可能な点が、全体が手動になる生のC++とは異なります。 ### 縮約や協調パターンでDisjointSlice単独に頼るべきでない条件 立場を明確にします。DisjointSlice単独に頼ってよいのは「各スレッドが自分の索引にだけ書く」パターンに限られます。縮約・スキャン・生産者消費者のように、複数スレッドが同期をはさんで重なり合う領域へ意図的にアクセスする協調パターンでは、DisjointSliceは答えになりません。ここで無理に型へ寄せず、同期を明示したunsafe（`get_unchecked_mut`やTier2 API）に切り替えるのが正解です。ワープ収束が絡む箇所も型では守られないため、収束の前提が崩れうるなら、DisjointSliceの安全保証はそこには及ばないと割り切ります。 ## 生ポインタや他のRust GPU手段と比較したcuda-oxide採用の判断基準 採用の可否は、何と比べて何を得るのかで決まります。生ポインタ方式、他のRust GPU手段、そしてalpha段階という現実の三点で判断します。 ### 手動境界検査が要るCUDA C++と型で守るcuda-oxideの並列出力比較 並列出力の書き方をCUDA C++の生ポインタ方式とcuda-oxideのDisjointSlice方式で比べると、安全検査の置き場所が決定的に違います。 | 観点 | CUDA C++（生ポインタ） | cuda-oxide（DisjointSlice） | | ------------ | ------------------------ | ---------------------------------------------- | | 並列出力の表現 | 生\_\_global\_\_ポインタ＋手動添字 | DisjointSlice＋ThreadIndex | | 境界外アクセス | 手動チェック（漏れると未定義動作） | get\_mut系がNoneを返す | | データ競合の検知 | 実行時にcompute-sanitizer等 | コンパイル時に型で排除（1D・定数2D） | | ホスト/デバイスのコード | .cuに分離＋nvcc | 単一.rsファイル＋cargo | | ビルド依存 | nvcc・CMake・FFI | cargo中心（nvcc・CMake不要／CUDA Toolkit・clang等は別途必要） | 生成されるPTXは手書きと同等で、安全網はコード生成時に消えます。型による検査はコンパイル時のみで、実行時コストは加わりません。 ### rust-cudaやwgpuなど他のRust GPU手段との用途の住み分け Rust×GPUには複数の選択肢があり、cuda-oxideはそれらと競合ではなく住み分けます。rust-cudaは既存のRust製CUDA環境で、NVlabsはこれを補完的なプロジェクトとして連携しています。rust-gpu（Embark Studios）はVulkanやMetal向けのグラフィックスシェーダ、CubeCLは制御されたDSLによるベンダー横断のGPU計算、wgpuはブラウザと各ネイティブ環境（Vulkan・Metal・DX12）を横断するクロスプラットフォームのWebGPU実装という棲み分けです。cuda-oxideの軸足は、ホストとデバイスのコードを単一のRustリポジトリに置き、バリアやワープ命令といったCUDAのプログラミングモデルをRustでそのまま表現する点にあります。NVIDIAのGPUに絞ってC++ツールチェーンを避けたいなら、cuda-oxideが噛み合います。 ### v0.1.0 alphaを試作で採用してよい条件と本番で見送る判断基準 導入の可否は段階で判断します。cuda-oxideは現在もalphaで、バグ・未実装・API破壊が前提です。ビルドにはRust nightly（`nightly-2026-04-03`にピン留め）とLLVM 21+が要り、TMAやtcgen05などのテンソルコア機能はLLVM 20以前のllcでは扱えません。研究・試作で、単一言語のRustでGPUカーネルを書きたい用途なら、今が早期参入の好機です。一方、本番のML学習や基幹のHPCワークロードは、当面PyTorch/JAXや成熟したC++ CUDAに置くのが堅実でしょう。性能面ではB200でGEMMが868 TFLOPS（cublasLtのFP16 SoL比58%）に達しており、alphaとしては十分競争力のある水準を示しています。 ## cuda-oxideのDisjointSliceに関するよくある質問 導入や実装でつまずきやすい点を、実際の検索質問に沿って5つ整理します。 ### DisjointSliceはどのクレートに含まれ、どうインポートしますか？ DisjointSliceはデバイス側クレート`cuda_device`に含まれます。カーネルファイルの先頭で`use cuda_device::{cuda_module, kernel, thread, DisjointSlice};`のようにインポートし、ホスト側は`cuda_core`（同期実行）や`cuda_async`（非同期実行）から`CudaContext`や`LaunchConfig`などを取り込みます。crates.ioの同名パッケージには存在しない型なので、依存はNVlabsのリポジトリからGit経由で取得します。 ### DisjointSliceを使えばカーネルからunsafeを完全になくせますか？ いいえ。1スレッドが1要素に書くTier1のパターンならunsafe不要ですが、共有メモリ・ワープ命令・アトミック・縮約などはTier2以上でunsafeが必要です。特に`&mut [T]`引数は型で守られないため、DisjointSliceで受けない可変出力にはエイリアシングの危険が残ります。unsafeをゼロにできるかどうかは、カーネルがTier1に収まるかで決まります。 ### get\_mut\_indexedとget\_mutはどちらを使うべきですか？ 索引の値を後段の算術に使わないなら、一発取得の`get_mut_indexed()`が既定で最も簡潔です。複数の入力スライスに同じ索引でアクセスしたり、`idx.get()`で得た添字を計算に使う場合は、`let idx = thread::index_1d();`で証人を取ってから`get_mut(idx)`を呼ぶ2段階形を選びます。どちらも`None`で境界外を弾く点は同じです。 ### crates.ioのcuda-oxideとNVlabsのcuda-oxideは同じものですか？ 別物です。crates.ioの旧「cuda-oxide」はCUDAライブラリの高レベルRustラッパーで、RustをPTXへコンパイルしません。本記事のNVlabs版cuda-oxideは、Rustを直接PTXへ変換する実験的コンパイラで、DisjointSliceはこちらにしか存在しません。名前が同じでも提供する機能が異なるため、導入時はNVlabsのリポジトリを参照しているか確認してください。 ### cuda-oxideは本番のML学習やHPCで使えますか？ 現時点では推奨されません。最新版も依然alphaで、既知の不健全性やAPI破壊が前提です。本番のML学習はPyTorchやJAX、基幹のHPCは成熟したC++ CUDAに置くのが堅実です。cuda-oxideは研究や試作、単一言語のRustでGPUカーネルを書く検証用途に向いており、エコシステムが安定した先のbeta・本番化を見据えた早期評価に適しています。 ## 関連記事 - [Rustにおけるasync/await構文の基本](https://www.issoh.co.jp/tech/details/8681/)：cuda-asyncがRustのasync/await（Future）を基盤とし、Tokio等のランタイムと組み合わせて使うため、非同期GPU処理の前提知識として役立ちます - [NVIDIA CUDA Toolkitに4件の深刻な脆弱性発覚](https://www.issoh.co.jp/tech/details/10712/)：GPU開発でもメモリ安全が問われる文脈の具体例です --- 出典: [cuda-oxideのDisjointSliceとは｜RustでGPU並列書き込みを型安全にする仕組み]()（株式会社一創）