CUDA-oxide: Nvidia'nın resmî Rust-to-CUDA derleyicisi

• cuda-oxide, güvenliğe yakın, idiomatik Rust ile SIMT GPU çekirdekleri yazmak ve standart Rust kodunu doğrudan PTX'e derlemek için deneysel bir derleyicidir
• DSL veya yabancı dil bağları olmadan yalnızca Rust kullanır; sahiplik, trait ve generics bilgisi varsayılır, async bölümleri içinse .await bilgisi de gerekir
• v0.1.0, erken alfa sürümüdür; bu nedenle hatalar, tamamlanmamış özellikler ve API'de geriye dönük uyumsuz değişiklikler beklenmelidir
• Örnek, cargo oxide run vecadd ile çalıştırılır; #[cuda_module] içindeki #[kernel] fonksiyonu thread::index_1d() ile vektör toplama işlemi yapar
• #[cuda_module], cihaz artifact'lerini host ikilisine gömer ve tür belirtilmiş loader ile çekirdek başına çalıştırma metodları üretir

---

Kullanım şekli ve üretilen kod

• Hızlı başlangıç

  • Kurulum önkoşulları sağlandıktan sonra cargo oxide run vecadd ile örnek derlenip çalıştırılır
  • Kurulum yönergeleri prerequisites sayfasında yer alır
  • Örnek, #[cuda_module] modülü içinde #[kernel] fonksiyonu vecadd'i tanımlar; thread::index_1d() ile indeksi alır ve a[i] + b[i] sonucunu DisjointSlice<f32> içine yazar
  • Host tarafında CudaContext::new(0), varsayılan stream ve kernels::load(&ctx) kullanılır; çekirdek DeviceBuffer::from_host, DeviceBuffer::<f32>::zeroed, LaunchConfig::for_num_elems(1024) ile çalıştırılır
  • Çalışma sonucu c.to_host_vec(&stream) ile alınır ve result[0] == 3.0 doğrulanır

• #[cuda_module] nasıl çalışır

  • #[cuda_module], üretilen cihaz artifact'lerini host ikilisine gömer
  • Tür belirtilmiş kernels::load fonksiyonu ile çekirdek başına çalıştırma metodları üretir
  • Belirli sidecar artifact'lerini yüklemek veya özel çalıştırma kodu yazmak gerektiğinde daha düşük seviyeli load_kernel_module ve cuda_launch! API'leri kullanılmaya devam edilebilir

Önkoşullar ve yönelim

• cuda-oxide, Rust'ın tür sistemi ve sahiplik modeliyle GPU çekirdekleri yazmayı hedefler ve güvenliği birinci sınıf bir hedef olarak ele alır
• GPU'larda incelikli noktalar bulunduğundan the safety model belgesinin okunması gerekir
• Bu bir DSL değil, saf Rust'ı PTX'e derleyen özel bir rustc kod üretim backend'idir
• GPU işlerini ertelenmiş olarak yürütülen DeviceOperation grafiği halinde kurmayı, bunları stream havuzuna zamanlamayı ve sonuçları .await ile beklemeyi sağlayan asenkron yürütmeyi destekler
• Rust'ta sahiplik, trait ve generics konularına aşinalık varsayılır; sonraki asenkron GPU programlama bölümleri ayrıca async/.await ve tokio gibi runtime'lar hakkında bilgi gerektirir
• Başvuru kaynakları olarak The Rust Programming Language, Rust by Example, Async Book sunulur
• v0.1.0 sürümü erken alfa aşamasındadır; hatalar, tamamlanmamış özellikler ve API'de geriye dönük uyumsuz değişiklikler beklenmelidir

[GN⁺]

Hacker News yorumları

• Gerçekten etkileyici. Uzun zamandır özel CUDA kernel’leri ve https://crates.io/crates/cudarc kullanıyorum; bu neredeyse drop-in bir alternatif olabilir gibi görünüyor
  Özellikle derleme sürelerinin nasıl kıyaslanacağını merak ediyorum. Rust CUDA crate’lerinin çoğu CMake ya da nvcc çağrılarına dayanıyor, bu da derlemeyi can sıkıcı derecede yavaşlatabiliyor
  Geçen hafta tam da derleme sürelerini profillerken sccache gibi araçların çıktı önbelleklemesiyle yeniden derleme süresini ciddi biçimde azaltabildiğini gördüm, ama özel nvcc çağrılarının maliyeti yine kalıyor. Örneğin Hugging Face’in candle’ı da kernel derleme sırasında özel nvcc komutları çağırıyor: https://arpadvoros.com/posts/2026/05/05/speeding-up-rust-whi...
  • Cudarc gerçekten çok iyi
    Rust CUDA crate’lerinin çoğunun CMake ya da nvcc çağırdığı için derlemenin yavaş olduğu kısmını kişisel olarak pek yaşamadım. Derleme betiklerini ele almak için yaptığım cuda_setup crate’ine bakarsanız, basit bir build.rs; dosyalar değiştiğinde yeniden derleniyor ve Rust’ın CPU tarafındaki koda kıyasla derleme süresi oldukça küçük kalıyor
  • Başkaları da cuda-oxide’ın cudarc için neredeyse drop-in bir alternatif gibi görünüp görünmediğini merak ediyor mu
    Öyle olsa harika olurdu, ama bana kalırsa muhtemelen tamamlayıcı olmaya daha yakın. cuda-oxide’ı farklılaştıran şeyin ne olduğunu da merak ediyorum; NVIDIA’nın tamamen kontrol ediyor olması dışında başka ne var, onu da merak ediyorum
• “Doğrudan PTX’e” gitmek tuhaf geliyor. Yakın dönemde NVIDIA MLIR de epey iyi ve hızlı. Ya da CuTile’ın kullandığı daha kolay ve daha güncel Tile IR [1] hedeflenebilirdi
  Tile IR biraz daha üst seviye, bu yüzden hedeflemesi çok daha kolay; sadece epilog füzyonu gibi yerlerde kayıp yaşarsınız
  [1] https://docs.nvidia.com/cuda/tile-ir/
  [2] https://developer.nvidia.com/cuda/tile
• Rust’ın bellek modelini CUDA semantiğine nasıl uyarladıklarını epey merak ediyorum. CUDA C++ ile kıyaslandığında neyin farklı olduğunu ve Rust tür sisteminin gerçekten CUDA’ya daha fazla güvenlik sağlayıp sağlayamayacağını da bilmek isterim
  GPU kernel yazmanın doğası gereği güvensiz olduğunu düşünüyorum. Donanımın çalışma biçimi ve sürekli en uç optimizasyonların gerekmesi nedeniyle güvenli bir dil yapmak fazlasıyla zor
  • Büyük görünen farklar dört tane. Birincisi, cudaFree’yi elle çağırma yaklaşımının aksine use-after-free ve drop semantiğini ele alıyor
    İkincisi, C++’taki void* argümanları yalnızca bir işaretçi dizisi ve sayı doğrulaması iken burada cuda_launch! ile kernel argümanları zorunlu kılınıyor
    Üçüncüsü, değiştirilebilir yazımın aliasing sorunu. C++’ta iki ya da daha fazla thread aynı i ile out[i] üzerine yazan kodu yine de derler, ama DisjointSlice ve ThreadIndex için açık yapıcı yok ve https://github.com/NVlabs/cuda-oxide/blob/2a03dfd9d5f3ecba52... yalnızca index_1d, index_2d, index_2d_runtime API’lerinin kullanılmasına izin veriyor
    Dördüncüsü, C++’ta std::string ya da pratikte herhangi bir POD’u cuda memcpy ile kopyalayıp durumu bozabilirsiniz, ama burada yalnızca DisjointSlice, skalerler ve closure’lar kabul ediliyor https://nvlabs.github.io/cuda-oxide/gpu-programming/memory-a...
    Ayrıntılar https://nvlabs.github.io/cuda-oxide/gpu-safety/the-safety-mo... ve https://nvlabs.github.io/cuda-oxide/gpu-programming/memory-a... içinde var. Elbette her şeyi yakalamıyor, ama çıplak .cu dosyalarına kıyasla tanımsız davranışı önlemek için çok daha fazla guardrail sağlıyor gibi görünüyor
  • Bu arada Rust’ın bellek modeli kasıtlı olarak C++ ile neredeyse tamamen aynı. Atomik işlemler de aynı ve provenance gibi kavramlar da var
    GPU programlama için uygun bir dil olup olmadığını zaman gösterecek, ama GPU’ya özgü tüm garipliklerden yararlanırken güvenli kod yazmaya izin veren, makul bir DSL benzeri API yapılabilirse şaşırmam. Sonuçta CUDA da biraz böyle değil mi?
  • Belgelerde oldukça ayrıntılı anlatılıyor. Güvenli katman, çoğunlukla güvenli katman ve güvensiz katman var
    Rust’ın Send/Sync modeline kolayca oturtulamayan güvenli ama paralel işler için biraz hantallık gerekiyor
  • Sanırım hedefe göre değişir. Rust ile uygulama yazıp arada bir içinde GPU hesaplaması kullanmak isteyen biri olarak dürüst olmak gerekirse bunu çok umursamıyorum
    Bellek modeli ya da sahiplik modelinden düşük sürtünmeyle yararlanabilmek güzel olur. Ama bunun karşılığında kullanım deneyimi çok rahatsız edici hale gelirse, o yaklaşımı istemem
    Mevcut taban çizgisinin Cudarc’ın yaptığı şey olduğunu düşünüyorum. Bellek yönetimine fazla karışmıyor; FFI’yi saran emirsel bir sözdizimi ve kernel değiştiğinde nvcc çağıran birkaç satırlık derleme betiği var
• Bunun Slang[0] için ne anlama geldiğini merak ediyorum. İnsanların GPU programlamayı daha modern bir dille yapmak istemesi asıl mesele gibi görünüyor; şimdi de sanki Rust kullanmaları yeterliymiş gibi duruyor
  Bu arada Slang’i epey seviyorum
  [0]: https://shader-slang.org/
  • Shader yazmak, en azından şu an için, CUDA kernel yazmaktan pratikte farklı. Shader’lar hem daha üst seviye hem de daha alt seviye ve belirli, sınırlı bir sürücü/GPU özellik kümesi için tasarlandıklarından birçok tuhaf yanları var
    Örneğin descriptor set’ler, kaynak yazmaçları, dispatch sınırları gibi şeyler
  • Hedef farklı. Slang tarafı AI algoritmalarından çok grafik programlama ile ilgileniyor
    Shader dilleri özellikler açısından da daha kullanıcı dostu. Ayrıca NVIDIA zaten Slang’i prodüksiyonda kullanıyor; o ekiplerin shader pipeline’larını Rust ile baştan yazmaları pek olası değil
• Rust ve “güvenli” programlama dilleri konusu açılmışken, NVIDIA’nın Spark/Ada’yı nasıl kullandığı hakkında daha fazla ayrıntı bilen var mı merak ediyorum
  Bulabildiğim tek şey aşağıdaki
  https://www.adacore.com/case-studies/nvidia-adoption-of-spar...
  • Geçen DEF CON’da bununla ilgili ayrıntılı bir sunum yaptılar: https://www.youtube.com/watch?v=KhWtkZmOPn4
  • AdaCore konferansında NVIDIA’nın 2020’de yaptığı “Securing the Future of Safety and Security of Embedded Software” sunumunun kaydı var
    https://www.youtube.com/watch?v=2YoPoNx3L5E
• “DSL yok, yabancı dil binding’leri yok, sadece Rust” ifadesinden başlayınca, resmi bir CUDA portu diye sunulmasına rağmen giriş bölümüne bile özen gösterilmemiş hissi veriyor
  Yine de bunu görmezden gelip belgeleri okumaya çalıştım, özel IR ortaya çıkınca ilginçleşecek gibi oldu ama sonra “MLIR implementasyonu C++ içinde biraz TableGen ile geliyor, LLVM’in tamamını derlemeniz gereken bir build sistemi var ve kariyer seçimlerinizi sorgulatan debug seansları yaşatıyor” tarzı ifadeleri görünce bu sektörü ciddiye almakta daha da zorlandım
  • Tüm kod tabanı büyük ölçüde AI ile yazılmış gibi görünüyor
  • Web sayfasında AI kullanılmasaydı “NVIDIA neden kendi web sitesini ve belgelerini AI ile yazmıyor? AI fabrikaları ve binlerce agent yöneten çalışanlar hakkındaki kendi anlatılarına inanmıyorlar mı?” tepkisi gelirdi
    Bu, AI abartısı yapan bir şirketten beklenecek kadar yerinde bir kendi ürününü kullanma örneği gibi duruyor
  • Adını da CUDA-oxide koymuşlar; sanki Rust dilinin adının kökeninin oksitlenme değil mantar olduğunu bilmiyorlar gibi
  • Tam olarak neye kızdığını anlamadım. Birinin MLIR’ın çok karmaşık ve LLVM’e bağımlı olduğunu söylemesine mi?
• TileLang https://github.com/tile-ai/tilelang ve Tile Kernels https://github.com/deepseek-ai/TileKernels gibi şeyler bir gün CUDA’yı demode hale getirecek
  • CUDA neredeyse 20 yıllık ve daha yıllarca ortadan kaybolmayacaktır
  • Oldukça büyük bir iddia ama destekleyen kanıt çok az
• https://nvlabs.github.io/cuda-oxide/gpu-safety/the-safety-mo... belgesine bakınca, GPU kernel’lerinin aynı belleği aynı anda gören binlerce thread üzerinde çalıştığı, CPU tarafında Rust’ın sahiplik ve ödünç alma ile veri yarışlarını önlediği ama GPU’da her SM başına 2048 thread aynı fonksiyonda başlayıp aynı çıktı tamponuna işaret ettiğinden borrow checker’ın bunun için tasarlanmadığı söyleniyor
  cuda-oxide, yaygın olan “bir thread bir eleman yazar” durumunu yapısal olarak güvenli hale getiriyor; paylaşımlı bellek, warp shuffle ve donanım iç gömülü işlevler gibi daha nadir durumlarda belgelenmiş sözleşmeleri olan unsafe gerektiriyor; TMA, tensor core ve küme düzeyinde iletişim gibi en ileri özellikleri ise donanımın karmaşıklığına uygun şekilde tamamen manuel bırakıyor
  Ama bu pek de Rust usulü gelmiyor. Rust’ta mevcut soyutlamalar probleme uymadığında yeni güvenli soyutlamalar yapılır. Rust for Linux bunun bir örneği
  Güvenli değilse Rust kullanmanın anlamı ne, diye düşündürüyor. Son damla performansı sıkmak isteyenlere unsafe API vermek tamam, ama bunun varsayılan olmaması gerekir
  İnsan ister istemez io_uring ya da Vulkan gibi API’ler için yazılmış kullanıcı alanı kütüphaneleriyle kıyaslıyor. Bunlar için güvenli API tasarlamak oldukça zor ve gerçekten sound olmayan denemeler de oldu
• Bunun host ve device arasında struct paylaşımına izin verip vermeyeceğini bilen var mı merak ediyorum. Mevcut Rust/CUDA iş akışlarında şimdiye kadar eksik olan büyük parçalardan biri buydu
  Aradaki serileştirme/bayt bariyeri de buna dahil
• CUDA’da Rust kullanırken çekindiğim noktalardan biri, Rust’ın normalde göz ardı edilebilecek ama burada önemli olabilecek ufak bir ek yük getirmesi
  Örneğin dizi sınır kontrolünün ek yazmaç kullanımına yol açıp kernel eşzamanlılığını düşürüp düşüremeyeceğini merak ediyorum