--- title: "cuda-oxideのDisjointSliceとは｜データ競合を型で防ぐGPU並列書き込み" url: "https://www.issoh.co.jp/tech/details/12683/" published: 2026-06-23 updated: 2026-06-23 categories: ["アーキテクチャ"] publisher: "株式会社一創" --- # cuda-oxideのDisjointSliceとは｜データ競合を型で防ぐGPU並列書き込み cuda-oxideのDisjointSliceは、NVIDIAのNVlabsが2026年5月に公開したRust→PTXコンパイラ cuda-oxide v0.1.0 が備える、GPUカーネル用の安全な書き込みスライス型です。本記事では、DisjointSliceの定義、ThreadIndex証人型による型安全の仕組み、`get_mut_indexed`など3種のメソッドの使い分け、そして3層安全モデルのなかで型保証が及ばない場面までを、公式ドキュメントの一次情報に基づいて整理します。CUDA C++の生ポインタ方式との違い、cudarcや同名の旧cuda-oxideラッパーとの混同、本番採用の可否といった判断材料も扱います。 ## まとめ｜DisjointSliceが並列書き込みのデータ競合を型で封じる要点 cuda-oxideのDisjointSliceは、GPUカーネルで複数スレッドが同じメモリ番地へ書き込んでしまうデータ競合を、実行時の検査ではなく型システムのコンパイル時検査で封じる仕組みです。書き込みは必ず`ThreadIndex`という証人を経由させ、各スレッドが固有の要素にしか触れないことを保証します。これにより、CUDA C++で生の`__global__`ポインタを使うときに避けられなかった「2スレッドが同一番地へ書く」事故が、コンパイルエラーとして表に出ます。 ただし型安全が効くのは「1スレッドが1要素を書く」Tier 1の範囲に限られます。縮約や共有メモリ、`&mut [T]`パラメータといった協調・例外パターンでは`unsafe`へ落ち、保証は手作業の責任に移ります。v0.1.0はアルファで、公式自身がバグとAPI破壊を覚悟すべきと明言しています。読み進める前提として、本番の機械学習トレーニングは現状PyTorchやJAXに置き、cuda-oxideは研究・試作の対象とする線引きが妥当です。 ## DisjointSliceの定義と、raw pointer方式が招くGPUデータ競合問題 ### 可変スライスをget\_mut経由に限定するDisjointSliceの定義と境界チェック `DisjointSlice`は、可変なデバイスメモリのスライスをラップし、書き込みを`get_mut`系メソッド経由に限定する型です。戻り値は`Option<&mut T>`で、範囲外のインデックスには`None`を返します。つまり境界チェックが型の入口に組み込まれています。CUDA C++で書く`if (i < n)`のガードと同じ役割を、戻り値の`Some`と`None`が肩代わりするわけです。 カーネル引数としては`ptr: *mut T`と`len: u64`へ分解（スカラー化）されてデバイス側へ渡ります。ホスト側の型付き`#[cuda_module]`メソッドは、書き込み用スライスに`&mut DeviceBuffer`を受け取り、ポインタと長さの取り出しを自動で行います。読み取り専用の入力は`&[T]`のまま渡せます。書き込み口だけを型で締める、という非対称な設計が要点です。 ### Rust→PTXのNVlabs版と旧crates.io版ラッパーの同名プロジェクト区別 検索で「cuda-oxide」を調べると、まったく別の2つのプロジェクトが混在します。ここを取り違えると`DisjointSlice`は見つかりません。crates.ioに以前から登録されている`cuda-oxide`は、CUDAドライバを包む高レベルのRustラッパーで、公式説明も「RustをPTXにコンパイルしない」と明記しています。GPU側コードの安全性は提供せず、ホスト側のバッファやオブジェクト操作の安全性に寄ったライブラリです。 本記事が扱う`DisjointSlice`を持つのは、NVlabsが2026年5月9日にv0.1.0として公開した新しい`cuda-oxide`のほうです。こちらはRustのカーネルコードを`rustc_public`（Stable MIR）からLLVM IRを経てPTXへ直接コンパイルする、rustcのコード生成バックエンドです。中間表現にはMLIR風のRust製フレームワークPlironを採用しています。`DisjointSlice`という語が出てきたら、後者のコンパイラ側を指していると考えて差し支えありません。 ### CUDA C++の生ポインタ方式が招くデータ競合という構造的リスク CUDA C++で並列出力を書く定番は、全スレッドが添字でアクセスする生の`__global__`ポインタです。この方式は構造的に危うく、2つのスレッドが同じ番地へ書いても何も止めてくれません。書き込み順序は未定義となり、結果がラン毎に変わります。GPUのスレッドやメモリの基礎をまず押さえたい場合は、[CUDAの定義と基本的な仕組みの解説](https://www.issoh.co.jp/tech/details/4902/)を先に読むと、なぜスレッドごとの書き込み先が衝突しうるのかが具体的につかめます。 cuda-oxideは、この生ポインタ方式を`DisjointSlice`という型に置き換え、衝突の可能性そのものをコンパイル時に潰す立場を取ります。ハードウェアが各スレッドへ一意の番号を割り振る事実を、型システムの保証へ翻訳しているわけです。実行時にサニタイザで競合を探すのではなく、ビルドが通った時点で「この書き込みは競合しない」と言える状態を目指します。 ## ThreadIndexという証人型がDisjointSliceの書き込みを一意化する仕組み ### borrow checkerが多数スレッドの同時書き込みを扱えないメモリ安全課題 GPUのカーネルは、同じ関数から起動された何千ものスレッドが、同時に同じメモリを見ます。1つのSMあたり最大2048スレッドが、同一の出力バッファを指したまま走ります。RustのborrowチェッカーはCPUの所有権と借用——可変参照は1つだけ、別名なし——を前提に設計されており、この状況をそのまま扱えません。`&mut`が1つの背後ポインタを全スレッドで共有する形になり、本来なら拒否すべき別名書き込みを区別できないのです。 メモリ起因の不具合がなぜ厄介かは、[メモリリークの定義と発生原因の解説](https://www.issoh.co.jp/tech/details/3621/)でも整理されています。ただしGPUのデータ競合は「解放漏れ」とは別種です。正しく動いて見えるのに結果だけが壊れる点が質の悪さで、再現も難しい。cuda-oxideはこの溝を、型を1つ挟むことで埋めます。 ### 公開コンストラクタを持たないThreadIndex証人型とハードウェア組込変数の由来 鍵を握るのが`ThreadIndex<'kernel, IndexSpace>`です。これは`usize`を包んだ不透明な「証人」で、公開コンストラクタを一切持ちません。`ThreadIndex::new(42)`のような生成は存在せず、書いてもコンパイルが通りません。値を得る唯一の経路は、ハードウェアの組込変数（`threadIdx`・`blockIdx`・`blockDim`）から導く信頼済みの関数だけです。これらは起動時にランタイムが埋める読み取り専用レジスタで、各スレッドが固有の値を受け取ることがハードウェアで保証されます。 証人はさらに`!Send + !Sync + !Copy + !Clone`で、`'kernel`ライフタイムがマクロの挿入するスタックローカルなスコープに縛られます。あるスレッドが自分の証人を共有メモリへ置き、隣のスレッドが後から拾う——という抜け道が型レベルで塞がれています。証人はカーネル本体より長く生きられず、スレッド間で受け渡すこともできません。 ### DisjointSliceがデータ競合をコンパイル時に排除する4つの根拠 安全性は4つの事実の組み合わせから出てきます。最も効いているのは、`index_1d()`が生む値がスレッドごとに一意である点です。`blockIdx.x * blockDim.x + threadIdx.x`という線形インデックスは、グリッド全体で重複しません。次に、`get_mut()`は境界チェック付きで、範囲外のスレッドには`None`を返します。 3つ目に、`IndexSpace`パラメータが証人とレイアウトを結び付けます。`DisjointSlice>`は`ThreadIndex<'_, Index2D<256>>`を受け付けず、stride違いはコンパイルエラーになります。4つ目が、先述の`!Send + !Sync + !Copy + !Clone`と`'kernel`束縛による転送禁止です。borrowチェッカーから見えるのはスレッドあたり1つの`&mut T`、ハードウェアが保証するのは添字どうしが互いに素であること。その2つを型システムが縫い合わせます。 ## 3種のget\_mutメソッドを1D・2D・縮約で使い分ける実装基準 ### 1D既定のget\_mut\_indexedと多スライス演算で使うget\_mutの選択基準 1次元グリッドの既定は`get_mut_indexed()`です。これはスレッドごとの証人を発行し、`&mut T`へ解決するところまでを一度の呼び出しで済ませます。戻り値の`None`は、グリッド外のスレッドと範囲外インデックスの両方を覆います。典型はベクトル加算で、`if let Some((c_elem, idx)) = c.get_mut_indexed()`と書けば、本体はほぼ普通のRustコードになります。 複数のスライスに対して同じ添字で算術する場合は、明示的な2段形`let idx = thread::index_1d(); out.get_mut(idx)`を選びます。添字`idx`を入力スライス側のインデックス計算にも使い回せるためです。判断基準はシンプルで、添字を演算へ流用するかどうか。流用しないなら一度呼びの`get_mut_indexed`、流用するなら2段の`get_mut`です。 ### const stride固定の型安全と実行時strideのunsafe契約という2D分岐 2次元グリッドでは、strideをコンパイル時に固定できるかどうかで分岐します。固定できるなら`index_2d::()`を使い、証人は`Index2D`付きで返ります。`DisjointSlice>`はその`S`しか受け付けないため、stride違いはコンパイル時に弾かれます。固定できず起動時に決まるstride（ホストでしか分からない行列サイズなど）では、`unsafe`な`index_2d_runtime(s)`と`Runtime2DIndex`を使います。 | 添字関数 | 計算式 | 戻り値の型 | unsafeの要否 | | --------------------- | -------------------------------------- | -------------------------------------------- | ----------------- | | index\_1d() | blockIdx.x \* blockDim.x + threadIdx.x | ThreadIndex<‘kernel, Index1D> | 不要（無条件で一意） | | index\_2d::() | row \* S + col | Option>> | 不要（stride違いは型エラー） | | index\_2d\_runtime(s) | row \* s + col | Option> | 必要（同一sを呼び出し側が保証） | この`unsafe`が契約です。同じスライスへ証人を渡す全スレッドが、同一の`s`を使ったことを呼び出し側が約束します。型システムには証明できないためです。なお`index_2d_row()`と`index_2d_col()`は素の`usize`を返す成分アクセサで、証人にはならず添字には使えません。線形化した結果だけが証人を得ます。コンパイル時に固定できるなら`index_2d::()`を優先してください。 ### 縮約で複数スレッドが同一要素へ書くget\_unchecked\_mutというエスケープ手段 「1スレッド1要素」に収まらない代表が縮約（reduction）です。warp内で計算した部分和をレーン0だけが書き戻すような場合、書き込み先のスレッドと添字が1対1で対応しません。ここで使うのが`unsafe`な`get_unchecked_mut(usize)`です。たとえば`if warp::lane_id() == 0`のブロックで`*out.get_unchecked_mut(warp_idx)`へ書く、という形になります。 負う義務は`ThreadIndex`方式が自動で課すものと同じです。添字が範囲内であること、2つのスレッドが同じ添字を共有しないこと。それを今度は自分で証明します。`SAFETY`コメントとして証明を書けないなら、その不変条件は実は守れていない——という公式の指摘は実務的です。書けるよう構造を組み直すか、安全なAPIへ戻す判断のほうが、深夜の`compute-sanitizer`ログより安上がりです。 ## 3層安全モデルでのDisjointSliceの守備範囲と過信が招く失敗パターン ### Tier1の型保証と、unsafeへ落ちるTier2/3の境界という守備範囲 cuda-oxideは安全性を3層に整理します。`DisjointSlice`が効くのは最上層のTier 1だけです。コンパイラが何を検証できるかで層が分かれており、自分のカーネルがどこに位置するかを把握すると、`unsafe`を書くべき箇所が見えてきます。 | 層 | 内容 | unsafe | 代表例 | | ------ | --------------- | -------- | ------------------------------------------------------------------------------ | | Tier 1 | 型で誤用を防ぐ・構築段階で安全 | 不要 | DisjointSlice + ThreadIndex（vecadd、素のGEMM） | | Tier 2 | 契約付きの限定的なunsafe | 必要（範囲限定） | 共有メモリ SharedArray、warp命令 shfl\_sync / ballot\_sync、atomics、get\_unchecked\_mut | | Tier 3 | 生ハードウェア・全面的に手作業 | 必要（広範） | TMA（sm\_90+）、tcgen05（sm\_120）、WGMMA、cluster、CLC | アプリケーションのカーネルはTier 1か、Tier 1とTier 2の境目に収まるのが普通です。Tier 3は次のCUTLASSを書く性能エンジニア向けで、vecaddやGEMM、縮約を書くだけならTier 2を出ることはまずありません。`DisjointSlice`を使う限り、守られているのはこのTier 1の範囲だと割り切るのが正確な理解です。 ### &mut \[T\]やwarp収束など型安全が及ばない3つの失敗パターン 型安全を過信すると踏む穴があります。実務でまず警戒すべきは、カーネル引数の`&mut [T]`です。マクロは今のところこの型を受理しますが、ランタイムのレイアウトでは全スレッドが同じ背後ポインタを見ます。2つのスレッドが`&mut data[i]`で書けば、まさに`DisjointSlice`が防ぐはずの別名書き込みになります。公式も既知の不健全性として明記しており、現状は`&mut [T]`を取るカーネルは全行が`unsafe`だと思って扱うべきです。出力は必ず`DisjointSlice`へ書き換えてください。 次に危ういのが実行時strideで、`index_2d_runtime`の`unsafe`契約を破れば衝突は静かに起きます。そしてwarp収束は型システムの管轄外です。`shfl_sync`や`ballot_sync`でマスクに発散したスレッドを含めると、クラッシュもエラーも出ず、カーネルが永遠に終わらない——という最悪のバグになります。ここは「ケースバイケース」で逃げず、warp命令を書くなら収束をコード設計で担保すると決め切るべき箇所です。 ### v0.1.0アルファとnightly固定という制約が示す本番採用の判断基準 採用判断は、性能の魅力と成熟度を切り分けて下します。性能面では、B200（148 SM）上のGEMMが`cta_group::2`とCLC、4段パイプラインで868 TFLOPS——cublasLtのFP16 SoL比58%——に達した実測があり、アルファとしては筋の良い数字です。前提環境は重く、固定された`nightly-2026-04-03`ツールチェーンと、Hopper／Blackwell向け機能ではLLVM 21以上が要ります。ホストの`cuda-bindings`はbindgenがclang-21を要求します。 何より公式自身がv0.1.0を「バグ・未実装・API破壊を覚悟せよ」と位置づけ、`&mut [T]`の不健全性も文書化済みです。結論として、本番の機械学習トレーニングはPyTorchやJAXに残し、cuda-oxideは研究・試作・社内検証の対象に限定するのが妥当です。単一言語でホストとカーネルを書ける利点は本物ですが、ミッションクリティカルな用途は成熟したC++ CUDAに置く、という線引きを今は守るべきです。 ## よくある質問 cuda-oxideのDisjointSliceについて、実装と採用判断の場面で迷いやすい論点をまとめます。 ### DisjointSliceと生の&mut \[T\]パラメータは何が違うのですか？ `&mut [T]`はカーネル引数として受理されますが、型安全は保証されません。ランタイムでは全スレッドが同じ背後ポインタを共有するため、2スレッドが`&mut data[i]`で書けば別名書き込み——データ競合——が起きます。公式も既知の不健全性として記載しています。対して`DisjointSlice`は書き込みを`ThreadIndex`証人経由に限定し、各スレッドが固有要素にしか触れないことを型で証明します。出力には`&mut [T]`ではなく`DisjointSlice`を使ってください。 ### cuda-oxideのDisjointSliceはcudarcやrust-cudaのスライスと同じものですか？ 別物です。cudarcや、crates.ioにある旧`cuda-oxide`は、ホスト側でCUDAドライバを包むラッパーで、RustをPTXにはコンパイルしません。安全性もCPU側の操作が中心です。一方`DisjointSlice`は、NVlabsの新しい`cuda-oxide`が持つデバイス側カーネル用の型で、レイヤーが異なります。rust-cudaとは競合ではなく補完関係にあり、NVlabsは両プロジェクトの維持者と協調していると説明しています。 ### DisjointSliceの境界チェックでカーネルは遅くなりませんか？ `get_mut`が返す`Some`と`None`の判定は、CUDA C++で手書きする`if (i < n)`のガードと同じ検査で、追加の重複ではありません。`index_1d()`の1次元では添字が無条件に一意なので、範囲判定も最小限です。Hacker Newsでは、Rustの境界チェック由来のレジスタ使用が占有率を下げる懸念も挙がりました。v0.1.0はアルファのため、性能が問題になりそうな箇所は対象ハードウェアで実測して判断するのが確実です。 ### 2次元グリッドでDisjointSliceを使うにはどうすればよいですか？ strideをコンパイル時に固定できるなら、`DisjointSlice>`と宣言し、`get_mut_indexed()`または`thread::index_2d::()`で添字を取ります。strideが違えばコンパイルエラーになり、取り違えを未然に防げます。strideが起動時にしか決まらない場合は、`unsafe`な`thread::index_2d_runtime(n)`と`Runtime2DIndex`付きスライスを使い、全スレッドが同じ`n`を使うことを呼び出し側が保証します。可能ならconst strideの`index_2d::()`を優先してください。 ### cuda-oxideは本番の機械学習トレーニングに使えますか？ 現状は推奨しません。v0.1.0はアルファで、公式がバグ・未実装・API破壊を明言しています。出力に`&mut [T]`を使うカーネルの不健全性も文書化済みです。B200でGEMMが868 TFLOPS（cublasLt FP16比58%）という実測は有望ですが、本番のトレーニングはPyTorchやJAXに残すのが安全です。cuda-oxideは研究・試作や、単一言語でホストとカーネルを書きたいケースの検証用と位置づけるのが妥当です。 ## 関連記事 - [Googleトレンドデータから見るプログラミング言語人気](https://www.issoh.co.jp/tech/details/2418/)：cuda-oxideが乗るRustの採用動向を把握する手がかりになります。 --- 出典: [cuda-oxideのDisjointSliceとは｜データ競合を型で防ぐGPU並列書き込み]()（株式会社一創）