RTX 50番台を持っていて、RustでGPUカーネルを書いてみたい人向けです。
NVIDIAがRustをそのままPTXへ落とすコンパイラを公開したので、同梱の例177本をRTX 5070 Ti（sm_120）で1本ずつ実行しました。
実行して結果まで検証できたのは153本、実際に落ちたのは6本です。

データセンター向けのB100やB200、あるいはHopperしか触らないなら、境界の話は当てはまりません。
ただしPTXの読み方は世代に依存しないので、後半だけ拾って読み飛ばしてください。

公式ドキュメントが「消費者向けBlackwellでは動かない」と書いている境界は、実機ではそこに無かった。

買えるBlackwellで、公式の例は何本動くのか

cuda-oxideはNVIDIA Labsが2026-05-07に公開したRust→PTXコンパイラで、DSLでもFFIでもなく、普通のRustがrustcのコード生成バックエンドを通ってPTXになります。

まず動かすところで4回止まりました。
READMEはsudo apt install clang-21を前提にしていますが、このPCでsudoは使えません。

cargo oxide doctorが✗を出したのはclangの1行だけで、apt.llvm.orgのdebをdpkg-deb -xで展開すれば片付きます。
2回目はバックエンドのリンクでunable to find dynamic system library 'ffi'、3回目はvecaddのリンクで同じ形の'cuda'が出ます。
文面は同じでも原因は別で、libffi.soは-devパッケージが張る側なので.so.8しか入っておらず、libcuda.soは/usr/lib/wsl/libに実在するのにリンカがそこを探索パスに入れていません。
どちらも1つのディレクトリにsymlinkを集めてRUSTFLAGS="-L <dir>"で渡せば通ります。

4回目がいちばん意外で、bindgenがcuda.h:56:10: fatal error: 'stdlib.h' file not foundを返します。
/usr/includeを覗くと、stdlib.hどころかlibcの開発ヘッダが1つも入っていません。
そういうマシンでRustを書き続けていたと、ここで初めて気づきました…。

4つともsudo無しで回避できたので、手順を置いておきます。

uv pip install nvidia-cuda-nvcc nvidia-cuda-cccl        # nvcc・libNVVM・nvJitLink・libdevice・ヘッダ一式
rustup component add llvm-tools --toolchain nightly-2026-04-03   # llc はここに入っている（LLVM 22.1.2）
# clang 21 は apt.llvm.org の deb を展開するだけ（dpkg-deb -x は root を要求しない）
for p in clang-21 libclang-common-21-dev libclang1-21 libclang-cpp21 libllvm21; do
  curl -sSLO "$BASE/${p}_${VER}_amd64.deb" && dpkg-deb -x "${p}_${VER}_amd64.deb" .
done
# libc・libstdc++ のヘッダはコンテナから抜き出す
docker run --rm -v "$PWD:/out" ubuntu:24.04 bash -c \
  'apt-get update -qq && apt-get install -y -qq libc6-dev g++ && tar cf - /usr/include | (cd /out && tar xf -)'
# libffi は .so が未作成、libcuda は探索パスの外。symlink を 1 か所に集めて -L で渡す
ln -sf /usr/lib/x86_64-linux-gnu/libffi.so.8 extralib/libffi.so
ln -sf /usr/lib/wsl/lib/libcuda.so.1         extralib/libcuda.so

READMEに書かれていない最後の2行だけは、エラーが探索パスの羅列なので気づきにくく、unable to find dynamic system library 'X'が出たらまずこれを疑ってください。

ここまで通るとcargo oxide doctorが全項目に✓を返し、cargo oxide run vecaddがGPUで正しい答えを出します。
このGPUは研究用に置いてあるもので、今回の動機も好奇心です。
1本動いたところで止める理由が無いので、同梱の177本を順に実行するループに移ります。

161本成功、と一度は書いた

最初は終了コードだけを数えて、161本が0を返したので成功率91%と書きかけたところで手が止まった。
wgmmaのログにこう出ていたからです。

ℹ️  PTX load failed (expected on non-Hopper): DriverError(218, "a PTX JIT compilation failed")

PTXのロードに失敗したのを想定内として飲み込み、終了コード0で終わっています。
数えていたのは「例のプログラムがエラーを返さなかったこと」であって、「カーネルがGPUで走ったこと」ではありません。

177本のログを1本ずつ読み直して、4つに分けました。

分類	本数	中身
実行して結果を検証できた	153	カーネルが走り、例が値を突き合わせて成功と表示した
成功が返るがカーネルは動いていない	8	PTXの生成だけ、あるいはロード失敗を想定内として飲み込んだ
落ちて正しい	10	`error_*`。コンパイラが弾くことを見せる例で、非ゼロ終了が正解
実際に落ちた	6	下で個別に見る

中身が無い8本のうち7本はテンソルコア関連（tcgen05系が4本、WGMMA系が2本、Blackwellのsparse MMAが1本）で、残る1本はMathDx SDKのダウンロード待ちです。
テンソルコア系がまとめて外れて、それ以外はほぼ全部動いた形になります。

実際に落ちた6本のうち3本は同じ根に行き着きました。
cpp_consumes_rust_deviceとsmall_type_ffi_testと、それに引きずられた経路です。

small_type_ffi_test.ll: error: IR version 1.0 incompatible with current version 2.0
Compiling with options: -arch=compute_120 -gen-lto -fma=1

--archをsm_90、sm_100、sm_120と振っても同じ行が出るので、これはGPUの世代の話ではありません。
cuda-oxideが出力するNVVM IRのバージョンが1.0で、手元のCUDA 13.3に入っているlibNVVMが2.0を要求しています。
NVVM IR経路を使う例だけが、ツールチェーンの版のずれで落ちています。

残る3本はabi_hmmとhost_closureとcutile_inter_kernelです。
cutile_inter_kernelはcuTileのコンパイラtileirasがPyPIのCUDAには入っていないので環境不足、abi_hmmはホストメモリをGPUから直接読むHMMの例で、WSL2ではDriverError(700)になります。
host_closureだけは中身が違い、テスト1から8まで通ったあとテスト9の「moveしない参照キャプチャ」でDriverError(700)が出ます。
これだけは環境の都合で説明がつかないので、追試するなら再現手順付きのissueにする価値がある1本です。

今回いちばん引っかかったのはここで、測り方が甘いと91%という数字が出ます。
同じことを試すなら、終了コードではなくログに検証行（✓ All ... correctのような突き合わせの結果）があるかを見てください。
grep -L "✓" logs/*.logで中身の無い成功が拾えます。

「sm_120では動かない」と書いてある例が動いた

分類しながら、通ってはいけないはずの例が通っていることに気づきました。
tma_multicastです。

この例のREADMEには、消費者向けBlackwellでの実行結果としてこう書かれています。

### On Consumer Blackwell (sm_120) or Hopper (sm_90):
GPU Compute Capability: sm_120
✗ PTX load failed: DriverError(CUDA_ERROR_INVALID_PTX, ...)
  Consumer Blackwell (sm_120) does NOT support multicast.

同じ文字列がソース側にもあり、ハードウェア要件の節には「NOT supported: Consumer Blackwell (sm_120)」と書かれています。
手元のRTX 5070 Tiはまさにそのsm_120なので、この例は落ちるはずでした。

実際の出力はこうです。

=== TMA Multicast Example (sm_100a) ===
GPU Compute Capability: sm_120
✓ PTX loaded successfully
2. Launching tma_multicast_test (cluster=(4,1,1), block=256)...
   ✓ All 4 CTAs have identical tile data (4096 values each)!
🎉 TMA multicast successful — one load, 4 CTAs served!

一度はフォールバックを疑いました。
multicastを使わない普通のTMAに落として、4つのCTAがそれぞれ読んだ結果がたまたま一致しただけ、という筋です。
生成されたPTXを見て、この筋は消えた。

$ grep -oE "cp\.async\.bulk\.tensor[a-z0-9_.:]*" tma_multicast.ptx | sort | uniq -c
      1 cp.async.bulk.tensor.2d.shared::cluster.global.tile.mbarrier::complete_tx::bytes.multicast::cluster

末尾のmulticast::cluster修飾子がそのまま出ていて、.targetもsm_120aなので、sm_100a向けの成果物が混ざったわけではなく、2回続けて実行しても同じ結果です。

なぜREADMEと食い違うのか。
select_targetの実装を読むと、機能から出力ターゲットを決める階段が見えます。

左の機能に対して右の.targetが決まる階段で、右端のsm_120がそのどこにも無いことだけ見てください。

tcgen05やFP8 multimemを使うカーネルはsm_100aへ、wgmmaはsm_90aへ、というように最低要件で決まります。
この階段にsm_120の段はありません。
消費者向けBlackwellは、nvidia-smiが返す12.0を助言として渡したときだけ選ばれる作りで、しかもその助言は「そのGPUがカーネルを実行できるときだけ」採用されます。

TMA multicastの要件は、コード上ではmajor >= 9。
コメントにも「PTX ISAが要求するのはsm_90以上で、接尾辞付きのターゲットは性能のための推奨」とあり、実装は最初からsm_120で動く前提になっています。
取り残されていたのはREADME側の記述だけです。

これは粗探しではなく、順番として妥当な取り残しに見えます。
tcgen05のようにTensor Memoryを使う命令は本当にsm_120に無く、そちらの線引きは実装もドキュメントも正確です。
multicastのほうは「datacenter推奨」と「datacenter必須」の区別が文章から抜けただけで、間違えたら被害が大きい種類の誤りではありません。

ただし読者としての結論は変わりません。
cuda-oxideを消費者向けBlackwellで試すなら、READMEの「NOT supported」は自分のGPUで確かめ直したほうがいい。
cargo oxide run <例名>を1回実行して、✓ PTX loaded successfullyが出るかどうかを見れば済みます。

ドキュメントが保守的なのは意図的で、未記載の挙動に寄りかかるのは危ない、という反論はあるはずです。
今日動いたのはドライバ595.97でのことなので、次のドライバでCUDA_ERROR_INVALID_PTXに戻る可能性も残ったままでしょう。
それでも、alphaのプロジェクトで「動かないと書いてあるから動かない」を採用すると、動く機能を捨てることになります。

Rustの境界チェックはPTXに残り、trapになる

ここからは世代に依存しない話です。
cargo oxide newが作る雛形のカーネルを、そのままPTXまで落として読みました。

use cuda_core::{CudaContext, DeviceBuffer, LaunchConfig1D};                // ホスト側の文脈・バッファ・起動設定use cuda_device::{kernel, launch_bounds, launch_contract, thread, DisjointSlice};
use cuda_host::cuda_module;                                               // デバイスコードを実行ファイルに埋め込む

#[cuda_module]                                                            // このモジュールをまとめてPTXへ落とすmod kernels {
    use super::*;

    #[kernel]                                                             // このfnをカーネルとして切り出す
    #[launch_bounds(256)]                                                 // 1ブロック256スレッドを上限として宣言
    #[launch_contract(domain = 1, block = (256, 1, 1))]                   // 1次元・ブロック形状を型で固定する
    pub fn vecadd(a: &[f32], b: &[f32], mut c: DisjointSlice<f32>) {
        let idx = thread::index_1d();                                     // 通し番号を型付きで受け取る
        let idx_raw = idx.get();                                          // 生のusizeへ落とす
        if let Some(c_elem) = c.get_mut(idx) {                            // 出力側は境界内のときだけ書き込み口が得られる
            *c_elem = a[idx_raw] + b[idx_raw];                            // 入力側は普通のスライス添字
        }
    }
}

fn main() -> Result<(), Box<dyn std::error::Error>> {
    let ctx = CudaContext::new(0)?;                                       // GPU 0の文脈
    let stream = ctx.default_stream();                                    // 既定ストリーム
    const N: usize = 1024;                                                // 要素数
    let a_host: Vec<f32> = (0..N).map(|i| i as f32).collect();            // 入力a
    let b_host: Vec<f32> = (0..N).map(|i| (i * 2) as f32).collect();      // 入力b
    let a_dev = DeviceBuffer::from_host(&stream, &a_host)?;               // aをGPUへ
    let b_dev = DeviceBuffer::from_host(&stream, &b_host)?;               // bをGPUへ
    let mut c_dev = DeviceBuffer::<f32>::zeroed(&stream, N)?;             // 出力cを0で確保

    // SAFETY: 上のkernelsモジュールから埋め込まれたデバイスコードは、このパッケージが所有している
    let module = unsafe { kernels::load(&ctx)? };                         // 埋め込み済みモジュールを読み込む
    let cfg = LaunchConfig1D::new((N as u32).div_ceil(256), 256, 0);      // ブロック数・スレッド数・共有メモリ
    let prepared = module.prepare_vecadd(cfg)?;                           // 起動設定を型で検証しておく
    module.vecadd(&stream, &prepared, &a_dev, &b_dev, &mut c_dev)?;       // 生成された型付きラッパで起動

    let c_host = c_dev.to_host_vec(&stream)?;                             // 結果を戻す
    let errors = (0..N).filter(|&i| (c_host[i] - (a_host[i] + b_host[i])).abs() > 1e-5).count();
    println!("errors = {errors}");                                        // 0なら全要素一致
    Ok(())
}

出力側のDisjointSliceは他のスレッドと重ならない書き込み口であることを型で示すもので、入力側は&[f32]のままなので、見た目には「書き込みだけ守られている」ように読めます。

左のスライス3本がPTXでは(ポインタ, 要素数)の6引数に開かれ、右の青い3行が入力側に付いた検査です。

実際にはa[idx_raw]とb[idx_raw]にも境界チェックが付いていて、違うのは外れたときの行き先です。
出力側のget_mutがNoneを返す経路はretへ飛んで何も書かずに終わりますが、入力側の2本はtrapへ飛びます。

trapが何になるかは、入力aだけを半分の長さにして確かめました。

Error: DriverError(719, "unspecified launch failure")

CPUならindex out of bounds: the len is 512 but the index is 512とファイル名と行番号が出るところです。
GPUではメッセージも添字も行番号も残りません。
bだけを短くしても文字列は1文字も変わらないので、719だけを見てもどちらのスライスが外れたかは分かりません。

診断の道具は両方とも空振りで、compute-sanitizerはPyPIから入るものの、WSL2ではTarget application terminated before first instrumented API callで起動せず、cuda-gdbにはそもそもwheelがありません。

そこで残るのは生成PTXを読む方法で、cargo oxide inspectを実行するとtrapのブロックが2つ出ます。
それぞれを守っているsetpの右辺がどの引数かを見れば、どのスライスの検査かが分かります（図の%rd5がaの要素数、%rd6がbの要素数）。
719が出たらinspectを1回実行してこの対応を作るのが、今のところいちばん早い道です。

inspectには1つ引っかかる点があって、runはGPUを見てsm_120a（PTX ISA 8.7）を出すのに対し、inspectはカーネルの最低要件だけを見てsm_80（ISA 7.0）を出します。
inspectの出力だけで「sm_80向けだ」と判断すると外れるので、走る側は--arch sm_120aを明示してください。
中身の差は.versionと.target、それにtrap; exit;がtrap;になる2箇所だけで、命令列は同一です。

安全のぶんは命令12本、時間の差は3点でしか言えない

同じ計算をCUDA C++で書いてnvccでPTXにし、並べました。

点が比の中央値で、縦の棒は9往復の最小から最大まで、つまりその差を判定できる精度があるかを表します。

命令は36対24、レジスタは32対24、分岐は4本対1本で、CUDA C++の通例では出力側だけをif (i < clen)で見るので、増えている3本が入力側2本とポインタのnull検査に対応します。

実行時間はcudarcで両方のPTXをロードし、要素数2^10から2^26まで2段ずつの9点で、同じバッファに交互に9往復ずつ投げました。
比は隣り合う往復どうしを組にして取っています。

最初は絶対時間を両対数で並べたのですが、6.7µsから1,327µsまでの中に数%の差を乗せても線が重なるだけでした。
比だけを線形の軸に取り直すと、探していたものとは別のものが出てきます。
小さい側では比が0.81から1.23まで、幅にして0.42も振れます。
9往復のうち速い回と遅い回で20%以上違うので、数%の差など最初から判定できていません。

幅が0.06を切るのは2^22からで、そこから右では比が1.00から1.03に収まります。
つまり「境界チェックのぶんは実行時間に出ない」と言えるのは右端の3点だけで、左側の6点は差が無いのではなく、測れていないだけです。

測っているのはカーネルの起動から完了までで、計算が支配的な形や内側のループで境界チェックが繰り返される形は測っていません。

nvcc側でも1つ引っかかりました。
PyPIから入るnvccは13.3で、何も指定しないとPTXのISAが9.3になり、ドライバ595.97（CUDA 13.2相当）ではCUDA_ERROR_UNSUPPORTED_PTX_VERSIONでロードできません。
今回は.versionの行だけを8.7に書き換えています（命令列が同一であることはdiffで確認しました）。
cuda-oxide側はLLVMのNVPTXバックエンドを通るぶんISAが保守的で、そのまま通ります。

測れていないこと

177本はそれぞれ1回ずつで、再実行によるばらつきを見ていません（時間の比だけは9往復して中央値と最小〜最大を出しています）。

GPUは1枚だけなので、sm_120で動いたtma_multicastが他のRTX 50番台でも動くのか、ドライバを上げても動き続けるのかは確かめていません。

NVVM IR経路の失敗は、CUDA 13.3のlibNVVMとの組み合わせで見た結果なので、12.x系なら通る可能性が残ります。
PyPIで13系しか入れていないため、そちらは確かめていません。

所感

NVIDIA自身がRustでカーネルを書けるようにしてきたことについては、ついにRustにも脚光が、というのが最初の反応です。

落ちた6本の原因は、どれもRustの外にあります。
これが今回いちばん腑に落ちた点です。
内訳はNVVM IRの版のずれ、cuTileのコンパイラの不在、WSL2のHMM、そして参照キャプチャ1つで、Rustからカーネルを書くこと自体は153本ぶんの例が普通に動く程度には成立しています。

検証環境：WSL2（Ubuntu 24.04.4、kernel 6.6.87.2）/ NVIDIA GeForce RTX 5070 Ti（compute capability 12.0、VRAM 16303MiB）/ ドライバ595.97 / cuda-oxide 0810ba9（2026-07-31）/ rustc 1.96.0-nightly（55e86c996、nightly-2026-04-03）/ llc LLVM 22.1.2（rustupのllvm-tools）/ clang 21.1.8 / nvcc 13.3.73（PyPI）/ cudarc 0.19.8