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<T, IndexSpace>は、可変なデバイスメモリのスライスをラップし、書き込みをget_mut系メソッド経由に限定する型です。戻り値はOption<&mut T>で、範囲外のインデックスにはNoneを返します。つまり境界チェックが型の入口に組み込まれています。CUDA C++で書くif (i < n)のガードと同じ役割を、戻り値のSomeとNoneが肩代わりするわけです。

カーネル引数としてはptr: *mut Tとlen: u64へ分解（スカラー化）されてデバイス側へ渡ります。ホスト側の型付き#[cuda_module]メソッドは、書き込み用スライスに&mut DeviceBuffer<T>を受け取り、ポインタと長さの取り出しを自動で行います。読み取り専用の入力は&[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の定義と基本的な仕組みの解説を先に読むと、なぜスレッドごとの書き込み先が衝突しうるのかが具体的につかめます。

cuda-oxideは、この生ポインタ方式をDisjointSliceという型に置き換え、衝突の可能性そのものをコンパイル時に潰す立場を取ります。ハードウェアが各スレッドへ一意の番号を割り振る事実を、型システムの保証へ翻訳しているわけです。実行時にサニタイザで競合を探すのではなく、ビルドが通った時点で「この書き込みは競合しない」と言える状態を目指します。

ThreadIndexという証人型がDisjointSliceの書き込みを一意化する仕組み

borrow checkerが多数スレッドの同時書き込みを扱えないメモリ安全課題

GPUのカーネルは、同じ関数から起動された何千ものスレッドが、同時に同じメモリを見ます。1つのSMあたり最大2048スレッドが、同一の出力バッファを指したまま走ります。RustのborrowチェッカーはCPUの所有権と借用——可変参照は1つだけ、別名なし——を前提に設計されており、この状況をそのまま扱えません。&mutが1つの背後ポインタを全スレッドで共有する形になり、本来なら拒否すべき別名書き込みを区別できないのです。

メモリ起因の不具合がなぜ厄介かは、メモリリークの定義と発生原因の解説でも整理されています。ただし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<T, Index2D<128>>は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::<S>()を使い、証人はIndex2D<S>付きで返ります。DisjointSlice<T, Index2D<S>>はそのSしか受け付けないため、stride違いはコンパイル時に弾かれます。固定できず起動時に決まるstride（ホストでしか分からない行列サイズなど）では、unsafeなindex_2d_runtime(s)とRuntime2DIndexを使います。

このunsafeが契約です。同じスライスへ証人を渡す全スレッドが、同一のsを使ったことを呼び出し側が約束します。型システムには証明できないためです。なおindex_2d_row()とindex_2d_col()は素のusizeを返す成分アクセサで、証人にはならず添字には使えません。線形化した結果だけが証人を得ます。コンパイル時に固定できるならindex_2d::<S>()を優先してください。

縮約で複数スレッドが同一要素へ書く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を書くべき箇所が見えてきます。

アプリケーションのカーネルは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<T, Index2D<S>>と宣言し、get_mut_indexed()またはthread::index_2d::<S>()で添字を取ります。strideが違えばコンパイルエラーになり、取り違えを未然に防げます。strideが起動時にしか決まらない場合は、unsafeなthread::index_2d_runtime(n)とRuntime2DIndex付きスライスを使い、全スレッドが同じnを使うことを呼び出し側が保証します。可能ならconst strideのindex_2d::<S>()を優先してください。

cuda-oxideは本番の機械学習トレーニングに使えますか？

現状は推奨しません。v0.1.0はアルファで、公式がバグ・未実装・API破壊を明言しています。出力に&mut [T]を使うカーネルの不健全性も文書化済みです。B200でGEMMが868 TFLOPS（cublasLt FP16比58%）という実測は有望ですが、本番のトレーニングはPyTorchやJAXに残すのが安全です。cuda-oxideは研究・試作や、単一言語でホストとカーネルを書きたいケースの検証用と位置づけるのが妥当です。

関連記事

• Googleトレンドデータから見るプログラミング言語人気：cuda-oxideが乗るRustの採用動向を把握する手がかりになります。