NVIDIA AI hat gerade cuda-oxide veröffentlicht: Ein experimentelles Rust-to-CUDA-Compiler-Backend, das SIMT-GPU-Kernel direkt in PTX kompiliert

Richten Sie die Rust Nightly Toolchain ein

Cuda-Oxid erfordert Rost Nacht- mit zwei zusätzlichen Komponenten: rust-src Und rustc-dev. Die Toolchain ist angeheftet an nightly-2026-04-03 über rust-toolchain.toml im Repository – es wird automatisch installiert, wenn Sie zum ersten Mal einen Construct im Repository ausführen.

Wenn Sie es manuell installieren müssen:

# Set up the pinned nightly toolchain
rustup toolchain set up nightly-2026-04-03

# Add required elements
rustup element add rust-src rustc-dev 
  --toolchain nightly-2026-04-03

# Affirm the toolchain is energetic
rustup present

ⓘ Warum diese Komponenten?
rustc-dev Macht die internen Compiler-APIs verfügbar, in die sich das benutzerdefinierte Codegen-Backend einklinkt. rust-src wird benötigt, damit der Compiler seine eigenen Standardbibliotheksquellen für das Geräteziel finden und kompilieren kann.

Installieren Sie LLVM 21 mit dem NVPTX-Backend

Die Cuda-Oxide-Pipeline sendet textuelle LLVM-IR (.ll Dateien) und übergibt sie an die Externe llc binär, um PTX zu erzeugen. Sie benötigen LLVM 21 oder höher mit aktiviertem NVPTX-Backend.

# Ubuntu/Debian
sudo apt set up llvm-21

# Confirm the NVPTX backend is current
llc-21 --version | grep nvptx

Die Pipeline erkennt automatisch llc-22 Und llc-21 auf deinem PATH in dieser Reihenfolge. Um eine bestimmte Binärdatei anzuheften, legen Sie die Umgebungsvariable fest:

# Pin to a selected llc binary
export CUDA_OXIDE_LLC=/usr/bin/llc-21

⚠ Häufiger Fehler
Wenn NVPTX nicht in der Ausgabe von erscheint llc-21 --versionIhr LLVM-Construct wurde ohne das NVPTX-Ziel kompiliert. Installieren Sie vom offiziellen LLVM apt-Repository und nicht von den Standardpaketen Ihrer Distribution, die möglicherweise GPU-Backends weglassen.

Installieren Sie Clang 21 für die Cuda-Bindings-Kiste

Der cuda-bindings Kiste verwendet bindgen um FFI-Bindungen zu generieren cuda.h zur Bauzeit. bindgen Bedürfnisse libclang – und insbesondere benötigt es Clangs eigenes Ressourcenverzeichnis (einschließlich stddef.h). Eine nackte libclang1-* Laufzeitpaket ist nicht genug.

# Set up the complete clang-21 package deal (contains useful resource headers)
sudo apt set up clang-21

# Alternatively, the -dev header package deal additionally works
sudo apt set up libclang-common-21-dev

⚠ Symptom eines fehlenden Klangs
Wenn Sie nur die Laufzeit, aber nicht die Header installieren, schlägt der Host-Construct mit einem kryptischen Fehler fehl 'stddef.h' file not discovered Fehler beim Binden. Laufen cargo oxide physician im nächsten Schritt, um dies zu erkennen, bevor Sie einen Construct versuchen.

Klonen Sie das Repo und installieren Sie Cargo-Oxide

cargo-oxide ist ein Cargo-Unterbefehl, der die gesamte Construct-Pipeline steuert – ausgeführt cargo oxide construct, cargo oxide run, cargo oxide debugUnd cargo oxide pipeline.

Im Repo (zum Ausprobieren von Beispielen):

git clone https://github.com/NVlabs/cuda-oxide.git
cd cuda-oxide

# cargo oxide works out of the field through a workspace alias
cargo oxide run vecadd

Außerhalb des Repos (für Ihre eigenen Projekte):

# Set up globally from the git supply
cargo set up 
  --git https://github.com/NVlabs/cuda-oxide.git 
  cargo-oxide

# On first run, cargo-oxide fetches and builds the codegen backend

Überprüfen Sie anschließend mit der integrierten Gesundheitsprüfung, ob alle Voraussetzungen erfüllt sind:

cargo oxide physician

ⓘ Was der Arzt überprüft
Es validiert Ihre Rust-Toolchain (nightly, rust-src, rustc-dev), Ihr CUDA-Toolkit, Ihre LLVM-Model und NVPTX-Unterstützung, Clang/libclang-Header und die Codegen-Backend-Binärdatei. Korrigieren Sie alle roten Elemente, bevor Sie fortfahren.

Erstellen Sie das vecadd-Beispiel und führen Sie es aus

Das kanonische erste Beispiel ist vecadd – ein Vektoradditionskernel, der zwei Arrays von 1.024 hinzufügt f32 Werte auf der GPU und überprüft das Ergebnis auf dem Host.

# Construct and run end-to-end
cargo oxide run vecadd

Wenn alles richtig konfiguriert ist, sehen Sie:

✓ SUCCESS: All 1024 parts appropriate!

Um die vollständige Kompilierungspipeline anzuzeigen – von Rust MIR über jeden Pliron-Dialekt bis hin zu PTX – führen Sie Folgendes aus:

# Print the complete Rust MIR — dialect-mir — mem2reg — dialect-llvm — LLVM IR — PTX hint
cargo oxide pipeline vecadd

Zum Debuggen cuda-gdb:

cargo oxide debug vecadd --tui

ⓘ Ausgabeartefakte
Ein erfolgreicher Construct erzeugt zwei Dateien: goal/debug/vecadd (die Host-Binärdatei) und goal/debug/vecadd.ptx (der Gerätecode). Die Host-Binärdatei lädt die PTX-Datei zur Laufzeit über den CUDA-Treiber.

Schreiben Sie Ihre eigene #(Kernel)-Funktion

Eine Kernelfunktion wird mit annotiert #(kernel). Verwenden DisjointSlice<T> für veränderliche Ausgänge und &(T) für schreibgeschützte Eingaben. Greifen Sie mit auf den eindeutigen {Hardware}-Index des Threads zu thread::index_1d().

use cuda_device::{kernel, thread, DisjointSlice};

// Tier 1 security: race-free by building, no `unsafe` wanted.
// DisjointSlice::get_mut() solely accepts a ThreadIndex —
// a hardware-derived opaque kind guaranteeing distinctive writes per thread.
#(kernel)
pub fn scale(enter: &(f32), issue: f32, mut out: DisjointSlice<f32>) {
    let idx = thread::index_1d();
    if let Some(elem) = out.get_mut(idx) {
        *elem = enter(idx.get()) * issue;
    }
}

ⓘ Tier-1-Sicherheit – wie es funktioniert
ThreadIndex ist ein undurchsichtiger neuer Typ usize das nur aus in die {Hardware} integrierten Registern erstellt werden kann (threadIdx, blockIdx, blockDim). Da jeder Thread einen eindeutigen Wert erhält, und DisjointSlice::get_mut() akzeptiert nur a ThreadIndexschreibt sind konstruktionsbedingt rassenfrei – nein unsafe irgendwo im Kernel.

Starten des Kernels über den Hostcode

Host- und Gerätecode leben im selben .rs Datei. Die Hostseite verwendet CudaContext, DeviceBufferund die cuda_launch! Makro zur Verwaltung des GPU-Speichers und der Zuteilung.

use cuda_core::{CudaContext, DeviceBuffer, LaunchConfig};
use cuda_host::{cuda_launch, load_kernel_module};

fn essential() {
    // Initialize GPU context on gadget 0
    let ctx    = CudaContext::new(0).unwrap();
    let stream = ctx.default_stream();
    let module = load_kernel_module(&ctx, "scale_example").unwrap();

    // Add enter knowledge to GPU reminiscence
    let knowledge: Vec<f32> = (0..1024).map(|i| i as f32).acquire();
    let enter  = DeviceBuffer::from_host(&stream, &knowledge).unwrap();
    let mut output = DeviceBuffer::<f32>::zeroed(&stream, 1024).unwrap();

    // Dispatch the kernel — LaunchConfig auto-sizes blocks/grids
    cuda_launch! {
        kernel: scale,
        stream: stream,
        module: module,
        config: LaunchConfig::for_num_elems(1024),
        args: (slice(enter), 2.5f32, slice_mut(output))
    }.unwrap();

    // Obtain end result again to host
    let end result = output.to_host_vec(&stream).unwrap();
    assert!((end result(1) - 2.5).abs() < 1e-5);
    println!("✓ Kernel ran efficiently!");
}

ⓘ Was für ein cuda_launch! tut
Es skaliert die Argumentliste – reduziert Slices, Skalare und erfasste Abschlüsse – in PTX-Kernelparameter und sendet den Kernel an den angegebenen Stream. Es ist kein manuelles Argument-Marshalling erforderlich.

Was Sie als Nächstes erkunden sollten

Sie haben ein funktionierendes Cuda-Oxid-Setup. Hier sind die hochwertigen Zukunftspfade, sortiert nach Komplexität:

• Generische Kernel mit Monomorphisierung – Probieren Sie es aus generic Beispiel (cargo oxide run generic), um zu sehen, wie fn scale<T: Copy> kompiliert, um PTX-Kernel professional Typ zu trennen.
• Schließungen mit Erfassungen – Die host_closure Beispiel zeigt, wie a transfer |x: f32| x * issue Der Abschluss wird skalarisiert und automatisch als PTX-Kernelparameter übergeben.
• Asynchrone GPU-Ausführung — cuda_launch_async! gibt ein faul zurück DeviceOperation das ausgeführt wird .sync() oder .await. Siehe die async_mlp Und async_vecadd Beispiele.
• Gemeinsamer Speicher und Warp-Intrinsics – diese erfordern einen Gültigkeitsbereich unsafe Blöcke mit dokumentierten Sicherheitsverträgen. Siehe Tier 2 in der Dokumentation zum Sicherheitsmodell.
• GEMM mit Lichtgeschwindigkeit – Die gemm_sol Beispiel erreicht 868 TFLOPS auf B200 (58 % von cuBLAS SoL). cta_group::2CLC und eine 4-stufige Pipeline.
• Blackwell-Tensorkerne – Die tcgen05 Beispiel zielt auf sm_100a mit TMEM, MMA und ab cta_group::2. Erfordert LLVM 21+.

ⓘ Bekannte Einschränkung in Model 0.1.0
index_2d(stride) wird als derzeit fehlerhaft dokumentiert – wenn Threads im selben Kernel unterschiedliche Schrittwerte verwenden, können zwei Threads erhalten &mut T zum gleichen Component mit Nr unsafe in Sichtweite. Bis der Repair landet (Schritt in einen Typparameter heben), binden Sie Schritt an einen einzelnen let binden und an jedem Anrufstandort wiederverwenden.

Vollständige Dokumentation: nvlabs.github.io/cuda-oxide · Quelle: github.com/NVlabs/cuda-oxide