CUDA-oxide: kompiler Rust-to-CUDA resmi dari Nvidia

 (nvlabs.github.io)
1 poin oleh GN⁺ 2 jam lalu | 1 komentar | Bagikan ke WhatsApp
  • cuda-oxide adalah kompiler eksperimental untuk menulis kernel GPU SIMT dengan Rust idiomatis yang mendekati aman dan mengompilasi kode Rust standar langsung ke PTX
  • Hanya menggunakan Rust tanpa DSL atau binding bahasa asing, dengan asumsi pemahaman tentang ownership, trait, dan generic, dan bab async juga memerlukan pengetahuan .await
  • v0.1.0 adalah rilis alpha awal, jadi bug, fitur yang belum selesai, dan perubahan API yang merusak kompatibilitas harus diantisipasi
  • Contoh dijalankan dengan cargo oxide run vecadd, dan fungsi #[kernel] di dalam #[cuda_module] melakukan penjumlahan vektor menggunakan thread::index_1d()
  • #[cuda_module] menyertakan artefak device ke dalam biner host, serta menghasilkan loader bertipe dan metode eksekusi per kernel

Cara penggunaan dan kode yang dihasilkan

  • Mulai cepat

    • Setelah memenuhi prasyarat instalasi, bangun dan jalankan contoh dengan cargo oxide run vecadd
    • Panduan instalasi ada di prerequisites
    • Contoh mendefinisikan fungsi #[kernel] bernama vecadd di dalam modul #[cuda_module], mengambil indeks dengan thread::index_1d(), lalu menulis a[i] + b[i] ke DisjointSlice<f32>
    • Di sisi host, kernel dijalankan menggunakan CudaContext::new(0), stream default, kernels::load(&ctx), DeviceBuffer::from_host, DeviceBuffer::<f32>::zeroed, dan LaunchConfig::for_num_elems(1024)
    • Hasil eksekusi diambil dengan c.to_host_vec(&stream) dan memeriksa bahwa result[0] == 3.0

  • Cara kerja #[cuda_module]

    • #[cuda_module] menyertakan artefak device yang dihasilkan ke dalam biner host
    • Menghasilkan fungsi kernels::load yang bertipe dan metode eksekusi per kernel
    • Saat perlu memuat artefak sidecar tertentu atau membuat kode eksekusi kustom, API level rendah load_kernel_module dan cuda_launch! tetap bisa digunakan

Prasyarat dan arah pengembangan

  • cuda-oxide bertujuan menulis kernel GPU dengan sistem tipe dan model ownership Rust, dengan keamanan sebagai tujuan utama
  • Karena GPU memiliki banyak detail yang halus, perlu membaca the safety model
  • Ini adalah backend codegen rustc kustom yang mengompilasi Rust murni ke PTX, bukan DSL
  • Mendukung eksekusi asinkron dengan menyusun pekerjaan GPU sebagai graf DeviceOperation yang dieksekusi tertunda, menjadwalkannya ke pool stream, dan menunggu hasil dengan .await
  • Diasumsikan pembaca sudah familier dengan ownership, trait, dan generic Rust; bab pemrograman GPU async berikutnya juga memerlukan pengetahuan tentang async/.await dan runtime seperti tokio
  • Materi referensi yang disediakan adalah The Rust Programming Language, Rust by Example, dan Async Book
  • Rilis v0.1.0 masih berada pada tahap alpha awal, jadi bug, fitur yang belum selesai, dan perubahan API yang merusak kompatibilitas harus diantisipasi

Bacaan terkait

1 komentar

 
GN⁺ 2 jam lalu
Komentar Hacker News
  • Benar-benar luar biasa. Sudah lama saya memakai kernel CUDA kustom dan https://crates.io/crates/cudarc, dan ini tampaknya bisa menjadi hampir pengganti drop-in
    Saya terutama penasaran bagaimana perbandingan waktu build-nya. Kebanyakan crate Rust CUDA bergantung pada pemanggilan CMake atau nvcc sehingga kompilasinya bisa terasa sangat lambat
    Kebetulan minggu lalu saya sedang memprofilkan waktu build dan melihat bahwa alat seperti sccache bisa sangat mengurangi waktu rebuild lewat caching artefak, tetapi biaya pemanggilan nvcc kustom tetap ada. Misalnya candle dari Hugging Face juga memanggil perintah nvcc kustom saat kompilasi kernel: https://arpadvoros.com/posts/2026/05/05/speeding-up-rust-whi...
    • Cudarc memang sangat bagus
      Soal kebanyakan crate Rust CUDA memanggil CMake atau nvcc sehingga kompilasinya lambat, saya pribadi tidak terlalu mengalaminya. Kalau melihat crate cuda_setup yang saya buat untuk menangani build script, itu cuma build.rs sederhana jadi hanya dikompilasi ulang saat file berubah, dan waktu kompilasinya sangat kecil dibanding kode Rust di sisi CPU
    • Saya penasaran apakah orang lain juga merasa cuda-oxide tampak seperti hampir pengganti drop-in untuk cudarc
      Kalau begitu akan sangat bagus, tetapi saya pribadi merasa ini mungkin lebih dekat ke pelengkap. Saya juga penasaran apa pembeda cuda-oxide, selain fakta bahwa NVIDIA mengontrolnya sepenuhnya
  • Agak aneh bahwa ini memilih jalur “langsung ke PTX”. NVIDIA MLIR belakangan juga cukup bagus dan cepat. Atau bisa saja menargetkan Tile IR [1] yang lebih mudah dan sedang lebih tren, seperti yang dipakai CuTile
    Tile IR sedikit lebih tingkat tinggi jadi jauh lebih mudah untuk ditargetkan, dan hanya rugi di hal-hal seperti fusi epilog
    [1] https://docs.nvidia.com/cuda/tile-ir/
    [2] https://developer.nvidia.com/cuda/tile
  • Saya cukup penasaran bagaimana model memori Rust disesuaikan dengan semantik CUDA. Saya juga ingin tahu apa yang berbeda dibanding CUDA C++, dan apakah sistem tipe Rust benar-benar bisa memberi keamanan tambahan di CUDA
    Menurut saya menulis kernel GPU pada dasarnya tidak aman. Sangat sulit membuat bahasa yang aman karena cara kerja perangkat kerasnya dan karena kita selalu harus melakukan optimasi ekstrem
    • Ada empat perbedaan besar yang terlihat. Pertama, ini menangani use-after-free dan semantik drop, tidak seperti pendekatan pemanggilan cudaFree secara manual
      Kedua, argumen void* di C++ hanyalah array pointer dan hanya jumlahnya yang divalidasi, sedangkan di sini argumen kernel dipaksakan lewat cuda_launch!
      Ketiga, ada masalah aliasing pada penulisan yang dapat berubah. Di C++, kode yang membuat dua atau lebih thread menulis ke out[i] dengan i yang sama tetap bisa dikompilasi, tetapi DisjointSlice dan ThreadIndex tidak punya konstruktor publik, dan hanya API https://github.com/NVlabs/cuda-oxide/blob/2a03dfd9d5f3ecba52... index_1d, index_2d, index_2d_runtime yang bisa dipakai
      Keempat, di C++ Anda bisa cuda memcpy std::string atau praktis POD apa pun lalu merusak state, tetapi di sini hanya menerima DisjointSlice, skalar, dan closure https://nvlabs.github.io/cuda-oxide/gpu-programming/memory-a...
      Rinciannya ada di https://nvlabs.github.io/cuda-oxide/gpu-safety/the-safety-mo... dan https://nvlabs.github.io/cuda-oxide/gpu-programming/memory-a.... Tentu tidak menangkap semuanya, tetapi tampaknya memberi jauh lebih banyak guardrail terhadap perilaku tak terdefinisi dibanding file .cu mentah
    • Sebagai referensi, model memori Rust memang sengaja hampir sepenuhnya sama dengan C++. Operasi atomik-nya juga sama, dan ada juga konsep seperti provenance
      Apakah ini bahasa yang nyaman untuk pemrograman GPU masih harus dilihat, tetapi saya tidak akan kaget kalau ternyata bisa dibuat API mirip DSL yang cukup baik untuk menulis kode aman sambil tetap memanfaatkan semua keanehan khas GPU. Bukankah CUDA pada akhirnya juga seperti itu
    • Dokumentasinya membahas ini dengan cukup detail. Ada lapisan aman, lapisan yang sebagian besar aman, dan lapisan tidak aman
      Untuk pekerjaan yang aman tetapi paralel dan sulit dimasukkan dengan rapi ke model Send/Sync Rust, memang diperlukan sedikit kekasaran
    • Mungkin tergantung tujuannya. Dari sudut pandang orang yang menulis aplikasi dalam Rust dan sesekali ingin memakai komputasi GPU di dalamnya, sejujurnya saya tidak terlalu peduli
      Akan bagus kalau model memori atau model kepemilikan bisa dimanfaatkan dengan gesekan rendah. Tetapi kalau itu justru membuat pengalaman pemakaian jauh lebih tidak nyaman, saya tidak menginginkan pendekatan seperti itu
      Menurut saya baseline-nya adalah cara Cudarc bekerja sekarang. Manajemen memorinya tidak terlalu ikut campur, hanya sintaks imperatif yang membungkus FFI dan beberapa baris build script yang memanggil nvcc saat kernel berubah
  • Saya penasaran apa artinya ini bagi Slang[0]. Intinya tampaknya orang ingin melakukan pemrograman GPU dengan bahasa yang lebih modern, dan sekarang seolah-olah cukup pakai Rust
    Sebagai catatan, saya cukup suka Slang
    [0]: https://shader-slang.org/
    • Menulis shader, setidaknya untuk saat ini, secara praktis berbeda dari menulis kernel CUDA. Shader sekaligus lebih tinggi tingkatnya dan lebih rendah tingkatnya, dan punya banyak keanehan karena dirancang untuk himpunan fitur driver/GPU yang spesifik dan terbatas
      Misalnya descriptor set, register resource, batas dispatch, dan hal-hal semacam itu
    • Targetnya berbeda. Pihak Slang lebih tertarik pada pemrograman grafis daripada algoritme AI
      Bahasa shading juga lebih ramah pengguna dari sisi fitur. Lagipula NVIDIA sudah memakai Slang di produksi, dan orang-orang itu tidak akan menulis ulang pipeline shader mereka ke Rust
  • Terkait pembahasan Rust dan bahasa pemrograman yang “aman”, saya penasaran apakah ada yang tahu lebih banyak tentang bagaimana NVIDIA memakai Spark/Ada
    Satu-satunya yang bisa saya temukan hanya ini
    https://www.adacore.com/case-studies/nvidia-adoption-of-spar...
  • Dari kalimat “tidak ada DSL, tidak ada binding bahasa asing, cuma Rust”, rasanya seperti port CUDA resmi ini bahkan tidak terlalu dipikirkan baik-baik di paragraf pembuka
    Saya tetap mencoba mengabaikannya dan membaca dokumentasinya, tetapi ketika mulai menarik karena ada IR kustom, lalu muncul kalimat seperti “implementasi MLIR berarti C++ dengan tambahan TableGen, sistem build yang mengharuskan kompilasi seluruh LLVM, dan sesi debugging yang membuat Anda mempertanyakan pilihan karier”, saya jadi makin sulit menganggap serius industri ini
    • Seluruh codebase-nya tampaknya sebagian besar ditulis AI
    • Kalau halaman webnya tidak memakai AI, reaksinya mungkin justru jadi “kenapa NVIDIA tidak menulis situs web dan dokumentasinya sendiri dengan AI? Apa mereka sendiri tidak percaya pada cerita mereka sebagai pengelola pabrik AI dan ribuan agen?”
      Ini tampak seperti bentuk dogfooding yang tepat dari perusahaan yang sangat membesar-besarkan AI
    • Namanya juga CUDA-oxide, yang menunjukkan mereka bahkan tidak tahu bahwa asal nama bahasa Rust bukan dari oksidasi melainkan jamur
    • Saya tidak benar-benar paham apa masalahnya. Apakah karena seseorang mengamati bahwa MLIR sangat kompleks dan bergantung pada LLVM?
  • Hal-hal seperti TileLang https://github.com/tile-ai/tilelang dan Tile Kernels https://github.com/deepseek-ai/TileKernels suatu saat akan membuat CUDA usang
    • CUDA sudah hampir 20 tahun, dan tidak akan hilang dalam beberapa tahun ke depan
    • Klaimnya cukup besar untuk landasan yang sangat minim
  • Kalau melihat dokumen https://nvlabs.github.io/cuda-oxide/gpu-safety/the-safety-mo..., kernel GPU berjalan di ribuan thread yang melihat memori yang sama secara bersamaan, dan di CPU Rust mencegah data race dengan ownership dan borrowing, tetapi di GPU ada 2048 thread per SM yang mulai dari fungsi yang sama dan menunjuk ke buffer keluaran yang sama, sehingga borrow checker memang tidak dirancang untuk itu
    cuda-oxide membuat kasus umum “satu thread menulis satu elemen” aman secara struktural, lalu untuk kasus yang lebih jarang seperti shared memory, warp shuffle, dan hardware intrinsic, ia meminta unsafe dengan kontrak yang terdokumentasi, sementara fitur garis depan seperti TMA, tensor core, dan komunikasi tingkat cluster dibiarkan sepenuhnya manual agar sesuai dengan kompleksitas perangkat kerasnya
    Namun ini terasa tidak terlalu khas Rust. Di Rust, kalau abstraksi yang ada tidak cocok dengan masalahnya, biasanya dibuat abstraksi aman yang baru. Rust for Linux adalah contohnya
    Kalau tidak aman, jadi apa alasan memakai Rust? Menyediakan API unsafe untuk orang yang ingin memeras performa terakhir itu masuk akal, tetapi itu seharusnya bukan default
    Saya jadi membandingkannya dengan pustaka user-space untuk API seperti io_uring atau Vulkan. Mendesain API yang aman untuk hal-hal seperti itu memang cukup sulit, dan memang ada upaya yang ternyata tidak sound
  • Saya penasaran apakah ini akan memungkinkan berbagi struct antara host dan device. Itu selama ini bagian besar yang hilang dari workflow Rust/CUDA yang ada sekarang
    Demikian juga penghalang serialisasi/byte di antaranya
  • Hal yang membuat saya waspada saat memakai Rust di CUDA adalah Rust menambahkan sedikit overhead yang biasanya bisa diabaikan, tetapi di sini mungkin penting
    Misalnya saya penasaran apakah pemeriksaan batas array bisa memicu penggunaan register tambahan sehingga menurunkan konkruensi kernel