CUDA-oxide: Nvidia का आधिकारिक Rust-to-CUDA compiler

 (nvlabs.github.io)
1 पॉइंट द्वारा GN⁺ 2 시간 전 | 1 टिप्पणियां | WhatsApp पर शेयर करें
  • cuda-oxide एक experimental compiler है जो लगभग सुरक्षित, idiomatic Rust में SIMT GPU kernel लिखने और standard Rust code को सीधे PTX में compile करने देता है
  • यह DSL या foreign language bindings के बिना केवल Rust का उपयोग करता है, और ownership, traits, generics की समझ मानकर चलता है; async अध्यायों के लिए .await का ज्ञान भी चाहिए
  • v0.1.0 एक early alpha release है, इसलिए bugs, अधूरे features और API में breaking changes की अपेक्षा रखनी चाहिए
  • उदाहरण को cargo oxide run vecadd से चलाया जाता है, और #[cuda_module] के भीतर #[kernel] function thread::index_1d() के साथ vector addition करता है
  • #[cuda_module] device artifacts को host binary में शामिल करता है, और typed loader तथा kernel-विशिष्ट execution methods generate करता है

उपयोग का तरीका और generated code

  • त्वरित शुरुआत

    • installation prerequisites पूरी करने के बाद cargo oxide run vecadd से उदाहरण को build और run किया जाता है
    • installation guide prerequisites में है
    • उदाहरण #[cuda_module] module के अंदर #[kernel] function vecadd को define करता है, और thread::index_1d() से index लेकर a[i] + b[i] को DisjointSlice<f32> में लिखता है
    • host side पर CudaContext::new(0), default stream, और kernels::load(&ctx) का उपयोग होता है, तथा DeviceBuffer::from_host, DeviceBuffer::<f32>::zeroed, LaunchConfig::for_num_elems(1024) से kernel चलाया जाता है
    • execution result को c.to_host_vec(&stream) से लाकर result[0] == 3.0 की पुष्टि की जाती है

  • #[cuda_module] कैसे काम करता है

    • #[cuda_module] generated device artifacts को host binary में शामिल करता है
    • यह typed kernels::load function और हर kernel के लिए execution methods generate करता है
    • यदि किसी खास sidecar artifact को load करना हो या custom execution code बनाना हो, तो lower-level load_kernel_module और cuda_launch! API का उपयोग अभी भी किया जा सकता है

पूर्वधारणाएँ और दिशा

  • cuda-oxide का लक्ष्य Rust के type system और ownership model के साथ GPU kernel लिखना है, और safety को first-class goal मानता है
  • GPU में कई सूक्ष्म पहलू होते हैं, इसलिए the safety model पढ़ना आवश्यक है
  • यह DSL नहीं, बल्कि pure Rust को PTX में compile करने वाला custom rustc codegen backend है
  • यह GPU work को deferred-execution DeviceOperation graph के रूप में compose करता है, stream pool पर schedule करता है, और .await से result का इंतज़ार करने वाली asynchronous execution को support करता है
  • यह मानकर चलता है कि पाठक Rust ownership, traits, और generics से परिचित हैं; आगे के async GPU programming अध्यायों के लिए async/.await और tokio जैसे runtime की जानकारी भी चाहिए
  • संदर्भ सामग्री के रूप में The Rust Programming Language, Rust by Example, Async Book दी गई हैं
  • v0.1.0 release शुरुआती alpha चरण में है, इसलिए bugs, अधूरे features, और API में breaking changes की अपेक्षा रखनी चाहिए

संबंधित पढ़ाई

1 टिप्पणियां

 
GN⁺ 2 시간 전
Hacker News टिप्पणियाँ
  • वाकई कमाल है। मैं लंबे समय से custom CUDA kernels और https://crates.io/crates/cudarc इस्तेमाल करता आ रहा हूँ, और यह लगभग drop-in replacement बन सकता है
    खास तौर पर यह देखना दिलचस्प होगा कि build time की तुलना कैसी रहती है। ज़्यादातर Rust CUDA crates CMake या nvcc invocation पर निर्भर करती हैं, इसलिए compilation काफ़ी दर्दनाक रूप से धीमा हो सकता है
    अभी पिछले हफ्ते build time profile करते हुए देखा कि sccache जैसे tools artifact caching के ज़रिए rebuild time को काफ़ी घटा सकते हैं, लेकिन custom nvcc invocation की लागत फिर भी रहती है। उदाहरण के लिए Hugging Face का candle भी kernel compilation में custom nvcc command चलाता है: https://arpadvoros.com/posts/2026/05/05/speeding-up-rust-whi...
    • Cudarc वाकई अच्छा है
      Rust CUDA crates में ज़्यादातर compilation CMake या nvcc call की वजह से धीमा हो जाता है—यह मैंने व्यक्तिगत रूप से उतना ज़्यादा महसूस नहीं किया। अगर cuda_setup crate को देखें, जो build scripts संभालने के लिए बना है, तो वह बस एक साधारण build.rs है, इसलिए केवल file बदलने पर ही दोबारा compile करता है, और Rust के CPU-side code की तुलना में उसका compile time बहुत छोटा है
    • मैं जानना चाहूँगा कि क्या बाकी लोगों को भी cuda-oxide, cudarc का लगभग drop-in replacement लग रहा है
      ऐसा हो तो बहुत अच्छा होगा, लेकिन मुझे निजी तौर पर लगता है कि यह शायद replacement से ज़्यादा complement होगा। यह भी जानना दिलचस्प होगा कि cuda-oxide की अलग पहचान क्या है, NVIDIA के पूर्ण नियंत्रण से आगे इसमें और क्या खास है
  • सीधे “PTX को target” करना थोड़ा अजीब लगता है। हाल का NVIDIA MLIR भी काफ़ी अच्छा और तेज़ है। या फिर CuTile जिस ज़्यादा आसान और हाल में लोकप्रिय Tile IR [1] का उपयोग करता है, उसे target किया जा सकता था
    Tile IR थोड़ा higher-level है, इसलिए उसे target करना कहीं आसान है, और नुकसान बस epilogue fusion जैसी चीज़ों में होता है
    [1] https://docs.nvidia.com/cuda/tile-ir/
    [2] https://developer.nvidia.com/cuda/tile
  • मुझे काफ़ी जिज्ञासा है कि Rust के memory model को CUDA semantics के साथ कैसे fit किया गया है। CUDA C++ की तुलना में क्या अलग है, और क्या Rust type system वास्तव में CUDA में ज़्यादा safety दे सकता है
    GPU kernel लिखना अपने आप में unsafe लगता है। hardware के काम करने के तरीके और हर समय extreme optimization की ज़रूरत की वजह से safe language बनाना बहुत मुश्किल है
    • ऊपर से देखने पर चार बड़े फ़र्क दिखते हैं। पहला, मैन्युअली cudaFree call करने के बजाय यह use-after-free और drop semantics को संभालता है
      दूसरा, C++ में void* arguments सिर्फ pointer array होते हैं और केवल count verify होता है, जबकि यहाँ cuda_launch! kernel arguments को enforce करता है
      तीसरा, mutable writes में aliasing की समस्या है। C++ में ऐसा code भी compile हो जाता है जहाँ दो या अधिक threads एक ही i के साथ out[i] में लिखते हैं, लेकिन DisjointSlice और ThreadIndex के public constructors नहीं हैं, और https://github.com/NVlabs/cuda-oxide/blob/2a03dfd9d5f3ecba52... केवल index_1d, index_2d, index_2d_runtime API के उपयोग की अनुमति देता है
      चौथा, C++ में std::string या लगभग कोई भी POD cuda memcpy करके उसकी state ख़राब की जा सकती है, लेकिन यहाँ सिर्फ DisjointSlice, scalar, और closures स्वीकार किए जाते हैं https://nvlabs.github.io/cuda-oxide/gpu-programming/memory-a...
      विवरण https://nvlabs.github.io/cuda-oxide/gpu-safety/the-safety-mo... और https://nvlabs.github.io/cuda-oxide/gpu-programming/memory-a... में है। बेशक यह सब कुछ नहीं पकड़ता, लेकिन raw .cu files की तुलना में undefined behavior रोकने के लिए काफ़ी ज़्यादा guardrails देता है
    • वैसे Rust का memory model जानबूझकर C++ के लगभग पूरी तरह समान है। Atomic operations भी वही हैं, और provenance जैसी अवधारणाएँ भी मौजूद हैं
      यह GPU programming के लिए सुविधाजनक language बनेगा या नहीं, यह अभी देखना बाकी है, लेकिन मुझे हैरानी नहीं होगी अगर GPU की सारी अजीब विशेषताओं का लाभ लेते हुए safe code लिखने के लिए कोई ठीक-ठाक DSL-जैसा API बनाया जा सके। शायद CUDA भी आख़िरकार कुछ ऐसा ही है
    • यह docs में काफ़ी विस्तार से समझाया गया है। एक safe layer, एक mostly-safe layer, और एक unsafe layer है
      जो काम सुरक्षित तो हैं लेकिन parallel भी हैं, और जिन्हें Rust के Send/Sync model में आसानी से fit नहीं किया जा सकता, उनके लिए थोड़ी भद्दी बनावट की ज़रूरत पड़ती है
    • यह शायद आपके लक्ष्य पर निर्भर करता है। अगर आप Rust में application लिखना चाहते हैं और बीच-बीच में उसमें GPU compute इस्तेमाल करना चाहते हैं, तो सच कहूँ तो मुझे यह बात ज़्यादा परेशान नहीं करती
      अगर memory model या ownership model को कम friction के साथ इस्तेमाल किया जा सके तो अच्छा है। लेकिन अगर उसकी वजह से user experience बहुत असुविधाजनक हो जाए, तो मैं वैसा नहीं चाहूँगा
      मेरे हिसाब से baseline वह है जो Cudarc अभी करता है। memory management में बहुत ज़्यादा दखल नहीं, बस FFI को wrap करने वाला imperative syntax और कुछ lines की build script जो kernel बदलने पर nvcc call करती है
  • मैं सोच रहा हूँ कि इसका Slang[0] के लिए क्या मतलब है। लगता है असली बात यह है कि लोग GPU programming एक ज़्यादा modern language में करना चाहते हैं, और अब ऐसा लग रहा है कि बस Rust इस्तेमाल किया जा सकता है
    वैसे, मुझे Slang काफ़ी पसंद है
    [0]: https://shader-slang.org/
    • shader लिखना, कम-से-कम अभी के लिए, CUDA kernels लिखने से व्यवहारिक रूप से अलग है। shaders एक साथ higher-level भी हैं और lower-level भी, और किसी खास, सीमित driver/GPU feature set के लिए डिज़ाइन होने के कारण उनमें कई अजीब बातें हैं
      जैसे descriptor sets, resource registers, dispatch limits वगैरह
    • target अलग है। Slang वाली दुनिया AI algorithms से ज़्यादा graphics programming में रुचि रखती है
      shading languages feature के लिहाज़ से भी ज़्यादा user-friendly हैं। और NVIDIA पहले से production में Slang इस्तेमाल कर रहा है, इसलिए वे लोग अपने shader pipeline को Rust में दोबारा नहीं लिखने वाले
  • Rust और “safe” programming languages की बात के संदर्भ में, क्या किसी को यह ज़्यादा पता है कि NVIDIA Spark/Ada का इस्तेमाल कैसे करता है
    मुझे सिर्फ़ नीचे दिया गया यह संदर्भ मिला
    https://www.adacore.com/case-studies/nvidia-adoption-of-spar...
  • “कोई DSL नहीं, कोई foreign-language bindings नहीं, बस Rust” जैसी पंक्ति से ही लगा कि आधिकारिक CUDA port होने के बावजूद introduction paragraph पर ठीक से ध्यान नहीं दिया गया
    फिर भी मैंने इसे नज़रअंदाज़ कर docs पढ़ने की कोशिश की, लेकिन जैसे ही custom IR आया और बात दिलचस्प होने लगी, तुरंत “MLIR implementation C++ में TableGen के साथ है, और इसके लिए पूरे LLVM को compile करना पड़ता है, साथ में ऐसे debugging sessions भी हैं जो आपको अपने career choices पर शक करा दें” जैसी पंक्ति दिखी, और उसके बाद इस industry को गंभीरता से लेना मुश्किल हो गया
    • पूरा codebase काफ़ी हद तक AI-generated लगता है
    • अगर web page पर AI इस्तेमाल न किया गया होता, तो लोग शायद कहते, “NVIDIA अपनी website और docs AI से क्यों नहीं लिखवाता? क्या उसे खुद अपनी उस कहानी पर भरोसा नहीं कि उसके कर्मचारी AI factories और हज़ारों agents मैनेज करते हैं?”
      यह तो AI hype company से अपेक्षित बिल्कुल सटीक dogfooding जैसा लगता है
    • नाम भी CUDA-oxide रखा है, मानो उन्हें यह पता ही न हो कि Rust language के नाम की उत्पत्ति oxidation से नहीं बल्कि fungus से है
    • मुझे ठीक-ठीक समझ नहीं आ रहा कि शिकायत किस बात की है। क्या सिर्फ़ इस observation से कि MLIR बहुत जटिल है और LLVM पर निर्भर है?
  • TileLang https://github.com/tile-ai/tilelang और Tile Kernels https://github.com/deepseek-ai/TileKernels जैसी चीज़ें कभी न कभी CUDA को पुराना बना देंगी
    • CUDA लगभग 20 साल पुराना है, और कई सालों तक कहीं जाने वाला नहीं है
    • बहुत कम सबूत के साथ यह काफ़ी बड़ा दावा है
  • https://nvlabs.github.io/cuda-oxide/gpu-safety/the-safety-mo... docs के अनुसार, GPU kernels उन हज़ारों threads पर चलते हैं जो एक ही memory को एक साथ देखते हैं। CPU पर Rust ownership और borrowing से data races रोकता है, लेकिन GPU में प्रति SM 2048 threads एक ही function से शुरू होकर उसी output buffer की ओर इशारा करते हैं, इसलिए borrow checker इसके लिए डिज़ाइन नहीं किया गया था
    cuda-oxide आम मामले—“एक thread, एक element लिखता है”—को संरचनात्मक रूप से safe बनाता है, और shared memory, warp shuffles, hardware intrinsics जैसे कम आम मामलों के लिए documented contracts के साथ unsafe माँगता है; जबकि TMA, tensor cores, और cluster-level communication जैसे frontier features को hardware complexity के अनुरूप पूरी तरह manual छोड़ देता है
    लेकिन यह ज़्यादा Rust-जैसा नहीं लगता। Rust में अगर मौजूदा abstractions समस्या पर ठीक से fit नहीं बैठते, तो नई safe abstractions बनाई जाती हैं। Rust for Linux इसका उदाहरण है
    अगर यह safe नहीं है, तो फिर Rust इस्तेमाल करने की वजह क्या है—यह सवाल उठता है। उन लोगों के लिए unsafe APIs देना ठीक है जिन्हें आख़िरी performance निचोड़नी है, लेकिन वह default नहीं होना चाहिए
    इससे io_uring या Vulkan जैसे APIs की user-space libraries की याद आती है। उनके लिए safe APIs डिज़ाइन करना काफ़ी मुश्किल है, और कुछ कोशिशें वास्तव में sound भी नहीं रही हैं
  • क्या किसी को पता है कि क्या यह host और device के बीच structs share करने देगा। मौजूदा Rust/CUDA workflow में अब तक यही एक बड़ी कमी रही है
    उनके बीच की serialization/byte barrier भी उतनी ही बड़ी समस्या है
  • CUDA में Rust इस्तेमाल करते समय मेरी चिंता यह रही है कि Rust कुछ हल्का-सा overhead जोड़ देता है, जो आम तौर पर नज़रअंदाज़ किया जा सकता है, लेकिन यहाँ महत्वपूर्ण हो सकता है
    उदाहरण के लिए, क्या array bounds checks अतिरिक्त register usage पैदा कर सकते हैं, जिससे kernel concurrency कम हो जाए—यह जानने की जिज्ञासा है