CUDA-oxide: Nvidias offizieller Rust-zu-CUDA-Compiler

 (nvlabs.github.io)
1 Punkte von GN⁺ 2 시간 전 | 1 Kommentare | Auf WhatsApp teilen
  • cuda-oxide ist ein experimenteller Compiler, mit dem sich SIMT-GPU-Kernel in sicherheitsnahen, idiomatischen Rust schreiben lassen und der Standard-Rust-Code direkt zu PTX kompiliert
  • Es wird ausschließlich Rust verwendet, ohne DSL oder Bindings zu Fremdsprachen; vorausgesetzt werden Kenntnisse zu Ownership, Traits und Generics, und für das Async-Kapitel auch Wissen über .await
  • v0.1.0 ist ein frühes Alpha-Release, daher sollte man mit Bugs, unvollständigen Funktionen und API-Breaking-Changes rechnen
  • Das Beispiel wird mit cargo oxide run vecadd ausgeführt; die Funktion #[kernel] innerhalb von #[cuda_module] führt mit thread::index_1d() eine Vektoraddition aus
  • #[cuda_module] bindet Device-Artefakte in das Host-Binary ein und erzeugt typspezifizierte Loader sowie kernelbezogene Ausführungsmethoden

Verwendung und generierter Code

  • Schnellstart

    • Nach Erfüllen der Installationsvoraussetzungen wird das Beispiel mit cargo oxide run vecadd gebaut und ausgeführt
    • Die Installationsanleitung steht unter prerequisites
    • Das Beispiel definiert die Funktion vecadd mit #[kernel] innerhalb eines #[cuda_module]-Moduls, ermittelt den Index mit thread::index_1d() und schreibt a[i] + b[i] in DisjointSlice<f32>
    • Auf der Host-Seite werden CudaContext::new(0), der Default-Stream und kernels::load(&ctx) verwendet; der Kernel wird mit DeviceBuffer::from_host, DeviceBuffer::<f32>::zeroed und LaunchConfig::for_num_elems(1024) ausgeführt
    • Das Ausführungsergebnis wird mit c.to_host_vec(&stream) geholt und result[0] == 3.0 geprüft

  • Verhalten von #[cuda_module]

    • #[cuda_module] bindet die erzeugten Device-Artefakte in das Host-Binary ein
    • Es erzeugt eine typisierte Funktion kernels::load sowie kernelbezogene Ausführungsmethoden
    • Wenn bestimmte Sidecar-Artefakte geladen oder benutzerdefinierter Ausführungscode erstellt werden müssen, können weiterhin die Low-Level-APIs load_kernel_module und cuda_launch! verwendet werden

Voraussetzungen und Ausrichtung

  • cuda-oxide zielt darauf ab, GPU-Kernel mit Rusts Typsystem und Ownership-Modell zu schreiben, wobei Sicherheit ein Ziel erster Ordnung ist
  • Da es bei GPUs subtile Besonderheiten gibt, sollte man the safety model lesen
  • Es ist keine DSL, sondern ein benutzerdefiniertes rustc-Codegen-Backend, das reines Rust zu PTX kompiliert
  • Unterstützt wird asynchrone Ausführung, bei der GPU-Arbeit als Graph verzögert ausgeführter DeviceOperations aufgebaut, in einen Stream-Pool eingeplant und per .await auf Ergebnisse gewartet wird
  • Vorausgesetzt wird Vertrautheit mit Rusts Ownership, Traits und Generics; die späteren Kapitel zu asynchroner GPU-Programmierung erfordern außerdem Kenntnisse zu async/.await und Runtimes wie tokio
  • Als Referenzmaterial werden The Rust Programming Language, Rust by Example, Async Book bereitgestellt
  • Das Release v0.1.0 befindet sich in einer frühen Alpha-Phase; mit Bugs, unvollständigen Funktionen und API-Breaking-Changes ist zu rechnen

Verwandte Beiträge

1 Kommentare

 
GN⁺ 2 시간 전
Hacker-News-Kommentare
  • Wirklich beeindruckend. Ich nutze schon lange benutzerdefinierte CUDA-Kernel und https://crates.io/crates/cudarc, und das hier sieht fast so aus, als könnte es ein Drop-in-Ersatz werden
    Ich bin vor allem gespannt, wie sich die Build-Zeiten vergleichen. Die meisten Rust-CUDA-Crates verlassen sich auf CMake- oder nvcc-Aufrufe, wodurch das Kompilieren schmerzhaft langsam werden kann
    Ich habe erst letzte Woche beim Profiling von Build-Zeiten gesehen, dass Tools wie sccache durch Artefakt-Caching die Rebuild-Zeiten stark verkürzen können, aber die Kosten benutzerdefinierter nvcc-Aufrufe bleiben trotzdem. Zum Beispiel ruft auch candle von Hugging Face bei der Kernel-Kompilierung benutzerdefinierte nvcc-Befehle auf: https://arpadvoros.com/posts/2026/05/05/speeding-up-rust-whi...
    • Cudarc ist wirklich gut
      Dass die meisten Rust-CUDA-Crates CMake oder nvcc aufrufen und deshalb langsam kompilieren, habe ich persönlich nicht so stark erlebt. Wenn man sich das cuda_setup-Crate ansieht, das ich für die Verarbeitung von Build-Skripten erstellt habe: Das ist nur ein einfaches build.rs, es wird also nur neu kompiliert, wenn sich Dateien ändern, und verglichen mit dem Rust-CPU-Code ist die Kompilierzeit sehr gering
    • Ich frage mich, ob andere auch finden, dass cuda-oxide fast wie ein Drop-in-Ersatz für cudarc aussieht
      Das wäre wirklich schön, aber ich vermute eher, dass es ein komplementäres Werkzeug ist. Ich frage mich auch, wodurch sich cuda-oxide eigentlich abhebt, abgesehen davon, dass NVIDIA die vollständige Kontrolle darüber hat
  • Direkt auf PTX zu gehen, wirkt seltsam. Das aktuelle NVIDIA MLIR ist ebenfalls ziemlich gut und schnell. Oder man hätte auf das einfachere und derzeit modernere Tile IR [1] zielen können, das CuTile verwendet
    Tile IR ist etwas höherstufig und daher viel einfacher als Ziel zu verwenden; man verliert eigentlich nur bei Dingen wie Epilog-Fusion
    [1] https://docs.nvidia.com/cuda/tile-ir/
    [2] https://developer.nvidia.com/cuda/tile
  • Mich interessiert ziemlich, wie das Speichermodell von Rust an die CUDA-Semantik angepasst wurde. Ich würde auch gern wissen, was im Vergleich zu CUDA C++ anders ist und ob das Rust-Typsystem CUDA tatsächlich mehr Sicherheit geben kann
    Ich halte das Schreiben von GPU-Kernels für inhärent unsicher. Wegen der Funktionsweise der Hardware und weil man ständig bis ans Limit optimieren muss, ist es einfach zu schwer, dafür eine sichere Sprache zu bauen
    • Vier große Unterschiede springen ins Auge. Erstens behandelt es use-after-free und Drop-Semantik, statt wie bei manuellen cudaFree-Aufrufen zu arbeiten
      Zweitens: Während C++-void*-Argumente nur ein Pointer-Array sind und lediglich die Anzahl validiert wird, erzwingt man hier die Kernel-Argumente mit cuda_launch!
      Drittens geht es um Alias-Probleme bei veränderbaren Schreibzugriffen. In C++ wird auch Code kompiliert, in dem zwei oder mehr Threads mit demselben i nach out[i] schreiben, aber DisjointSlice und ThreadIndex haben keine öffentlichen Konstruktoren, und man soll nur die APIs https://github.com/NVlabs/cuda-oxide/blob/2a03dfd9d5f3ecba52... index_1d, index_2d, index_2d_runtime verwenden
      Viertens kann man in C++ std::string oder praktisch jeden POD per cuda memcpy kopieren und damit den Zustand beschädigen, während hier nur DisjointSlice, Skalare und Closures akzeptiert werden https://nvlabs.github.io/cuda-oxide/gpu-programming/memory-a...
      Mehr Details stehen unter https://nvlabs.github.io/cuda-oxide/gpu-safety/the-safety-mo... und https://nvlabs.github.io/cuda-oxide/gpu-programming/memory-a.... Es fängt natürlich nicht alles ab, aber es scheint viel mehr Leitplanken gegen undefiniertes Verhalten zu geben als bei rohen .cu-Dateien
    • Zur Einordnung: Das Speichermodell von Rust ist absichtlich fast vollständig identisch mit dem von C++. Atomare Operationen sind ebenfalls gleich, und Konzepte wie Provenance gibt es auch
      Ob es eine gute Sprache für GPU-Programmierung ist, muss sich noch zeigen, aber es würde mich nicht überraschen, wenn man eine halbwegs vernünftige DSL-artige API bauen kann, mit der man sicheren Code schreibt und gleichzeitig all die seltsamen GPU-spezifischen Dinge nutzt. Ist CUDA am Ende nicht genau so etwas?
    • In der Dokumentation wird das ziemlich ausführlich erklärt. Es gibt eine sichere Schicht, eine überwiegend sichere Schicht und eine unsichere Schicht
      Für sichere, aber parallele Aufgaben, die sich nicht leicht in Rusts Send/Sync-Modell einfügen lassen, braucht es etwas unbeholfene Konstruktionen
    • Das hängt wohl von den Zielen ab. Wenn ich Anwendungen in Rust schreibe und darin gelegentlich GPU-Berechnungen nutzen will, ist mir das ehrlich gesagt nicht besonders wichtig
      Es wäre gut, wenn sich Speicher- oder Ownership-Modell mit wenig Reibung nutzen ließen. Aber wenn dadurch die Nutzungserfahrung deutlich schlechter wird, will ich das auch nicht
      Die Messlatte ist für mich das, was Cudarc aktuell macht: relativ wenig Eingriff in die Speicherverwaltung, eine imperative Syntax als FFI-Wrapper und ein paar Zeilen Build-Skript, die nvcc aufrufen, wenn sich der Kernel ändert
  • Ich frage mich, was das für Slang[0] bedeutet. Der Kern der Sache scheint doch zu sein, dass Leute GPU-Programmierung in moderneren Sprachen machen wollen, und jetzt sieht es so aus, als könnte man einfach Rust verwenden
    Zur Einordnung: Ich mag Slang ziemlich gern
    [0]: https://shader-slang.org/
    • Shader zu schreiben unterscheidet sich zumindest derzeit praktisch von CUDA-Kernels. Shader sind gleichzeitig höher- und niedrigerstufig, und weil sie für einen spezifischen und begrenzten Satz an Treiber-/GPU-Funktionen ausgelegt sind, haben sie viele Eigenheiten
      Zum Beispiel Descriptor Sets, Ressourcenregister oder Dispatch-Beschränkungen
    • Das Ziel ist ein anderes. Bei Slang geht es eher um Grafikprogrammierung als um KI-Algorithmen
      Shader-Sprachen sind funktional außerdem benutzerfreundlicher. Und NVIDIA setzt Slang bereits produktiv ein; diese Leute werden ihre Shader-Pipeline kaum in Rust neu schreiben
  • Im Zusammenhang mit Rust und „sicheren“ Programmiersprachen frage ich mich, ob jemand mehr darüber weiß, wie NVIDIA Spark/Ada einsetzt
    Alles, was ich dazu finden kann, ist Folgendes
    https://www.adacore.com/case-studies/nvidia-adoption-of-spar...
  • Schon die Formulierung „keine DSL, keine Fremdsprachen-Bindings, einfach nur Rust“ wirkt für einen offiziellen CUDA-Port so, als hätte man sich nicht einmal beim Einführungstext Mühe gegeben
    Ich wollte das ignorieren und trotzdem die Doku lesen, aber gerade als es mit dem benutzerdefinierten IR interessant zu werden begann, kam so ein Satz wie „eine MLIR-Implementierung in C++ mit etwas TableGen dazu, einem Build-System, das das gesamte LLVM kompiliert, und Debugging-Sessions, die einen die eigene Berufswahl hinterfragen lassen“, und danach fiel es mir schwer, diese Branche noch ernst zu nehmen
    • Die gesamte Codebasis wirkt größtenteils KI-generiert
    • Wenn sie auf der Website keine KI eingesetzt hätten, wäre die Reaktion gewesen: „Warum nutzt NVIDIA für die eigene Website und Dokumentation keine KI? Glauben sie ihrer eigenen Geschichte über KI-Fabriken und Mitarbeitende, die Tausende Agenten verwalten, etwa selbst nicht?“
      Das sieht nach genau dem erwartbaren dogfooding eines KI-Hype-Unternehmens aus
    • Schon der Name CUDA-oxide zeigt, dass man offenbar nicht weiß, dass sich der Name der Sprache Rust nicht von Oxidation, sondern von einem Pilz ableitet
    • Ich verstehe nicht genau, was das Problem ist. Ist es einfach die Feststellung, dass MLIR sehr komplex und von LLVM abhängig ist?
  • Dinge wie TileLang https://github.com/tile-ai/tilelang und Tile Kernels https://github.com/deepseek-ai/TileKernels werden CUDA irgendwann alt aussehen lassen
    • CUDA ist fast 20 Jahre alt und wird auch in den nächsten Jahren nicht verschwinden
    • Das ist eine ziemlich große Behauptung für so wenig Belege
  • Wenn man das Dokument https://nvlabs.github.io/cuda-oxide/gpu-safety/the-safety-mo... liest, steht dort, dass GPU-Kernels auf Tausenden Threads laufen, die gleichzeitig denselben Speicher sehen; auf der CPU verhindert Rust mit Ownership und Borrowing Datenrennen, aber auf der GPU starten 2048 Threads pro SM in derselben Funktion und zeigen auf denselben Ausgabepuffer, wofür der Borrow Checker nicht ausgelegt ist
    cuda-oxide mache den häufigen Fall „ein Thread schreibt ein Element“ strukturell sicher, verlange für seltenere Fälle wie Shared Memory, Warp Shuffles oder Hardware-Intrinsics unsafe mit dokumentierten Verträgen und lasse ganz vorn liegende Features wie TMA, Tensor Cores oder Kommunikation auf Cluster-Ebene vollständig manuell, entsprechend der Hardware-Komplexität
    Aber das wirkt nicht besonders Rust-typisch. In Rust baut man neue sichere Abstraktionen, wenn bestehende für ein Problem nicht gut passen. Rust for Linux ist dafür ein Beispiel
    Wenn es nicht sicher ist, fragt man sich, wozu man dann Rust verwendet. Es ist in Ordnung, Leuten, die das letzte Quäntchen Leistung herausholen wollen, eine unsichere API zu geben, aber das sollte nicht der Standard sein
    Das erinnert mich an User-Space-Bibliotheken für APIs wie io_uring oder Vulkan. Dafür sichere APIs zu entwerfen, ist ziemlich schwierig, und es gab tatsächlich auch unsounde Versuche
  • Weiß jemand, ob das gemeinsam genutzte Structs zwischen Host und Device ermöglichen wird? Das war bisher ein großer fehlender Baustein in Rust/CUDA-Workflows
    Gleiches gilt für die Serialisierungs-/Byte-Barriere dazwischen
  • Meine Sorge bei Rust in CUDA war, dass Rust etwas Overhead hinzufügt, der sonst meist ignorierbar ist, hier aber wichtig werden könnte
    Zum Beispiel frage ich mich, ob Bounds Checks bei Arrays zusätzlichen Registerverbrauch verursachen und dadurch die Kernel-Konkurrenzfähigkeit verringern könnten