CUDA-oxide: คอมไพเลอร์ Rust-to-CUDA อย่างเป็นทางการของ Nvidia

 (nvlabs.github.io)
1 คะแนน โดย GN⁺ 2 시간 전 | 1 ความคิดเห็น | แชร์ทาง WhatsApp
  • cuda-oxide เป็นคอมไพเลอร์เชิงทดลองสำหรับเขียนเคอร์เนล SIMT GPU ด้วย Rust แบบ idiomatic ที่ใกล้เคียงความปลอดภัย และคอมไพล์โค้ด Rust มาตรฐานเป็น PTX ได้โดยตรง
  • ใช้เพียง Rust โดยไม่มี DSL หรือ foreign-language binding และตั้งอยู่บนสมมติฐานว่าผู้ใช้เข้าใจ ownership, trait และ generic ส่วนบท async ก็ต้องมีความรู้เรื่อง .await ด้วย
  • v0.1.0 เป็นรีลีส อัลฟาระยะแรก จึงควรคาดว่าจะมีบั๊ก ฟีเจอร์ที่ยังไม่เสร็จสมบูรณ์ และการเปลี่ยนแปลง API แบบทำลายความเข้ากันได้
  • ตัวอย่างรันด้วย cargo oxide run vecadd และฟังก์ชัน #[kernel] ภายใน #[cuda_module] จะทำการบวกเวกเตอร์ด้วย thread::index_1d()
  • #[cuda_module] จะฝัง device artifact ไว้ในโฮสต์ไบนารี และสร้างตัวโหลดแบบระบุชนิดพร้อมเมธอดรันแยกตามเคอร์เนล

วิธีใช้งานและโค้ดที่สร้างขึ้น

  • เริ่มต้นอย่างรวดเร็ว

    • หลังเตรียมข้อกำหนดเบื้องต้นแล้ว ให้บิลด์และรันตัวอย่างด้วย cargo oxide run vecadd
    • คำแนะนำการติดตั้งอยู่ที่ prerequisites
    • ตัวอย่างจะกำหนดฟังก์ชัน vecadd ที่เป็น #[kernel] ภายในโมดูล #[cuda_module] แล้วใช้ thread::index_1d() เพื่อรับอินเด็กซ์ จากนั้นเขียนค่า a[i] + b[i] ลงใน DisjointSlice<f32>
    • ฝั่งโฮสต์จะใช้ CudaContext::new(0), default stream และ kernels::load(&ctx) จากนั้นรันเคอร์เนลด้วย DeviceBuffer::from_host, DeviceBuffer::<f32>::zeroed, LaunchConfig::for_num_elems(1024)
    • ดึงผลลัพธ์กลับมาด้วย c.to_host_vec(&stream) และตรวจสอบว่า result[0] == 3.0

  • การทำงานของ #[cuda_module]

    • #[cuda_module] จะฝัง device artifact ที่สร้างขึ้นไว้ในโฮสต์ไบนารี
    • สร้างฟังก์ชัน kernels::load แบบระบุชนิด และเมธอดรันแยกตามเคอร์เนล
    • หากต้องการโหลด sidecar artifact เฉพาะ หรือสร้างโค้ดรันแบบคัสตอม ก็ยังสามารถใช้ API ระดับล่างอย่าง load_kernel_module และ cuda_launch! ได้ต่อไป

ข้อกำหนดเบื้องต้นและทิศทาง

  • cuda-oxide มีเป้าหมายเพื่อให้เขียน GPU kernel ด้วยระบบชนิดและโมเดล ownership ของ Rust โดยยกระดับความปลอดภัยเป็นเป้าหมายชั้นหนึ่ง
  • เนื่องจาก GPU มีรายละเอียดที่ละเอียดอ่อน จึงควรอ่าน the safety model
  • มันไม่ใช่ DSL แต่เป็น custom rustc codegen backend ที่คอมไพล์ Rust ล้วนเป็น PTX
  • รองรับ การรันแบบอะซิงก์ โดยประกอบงาน GPU เป็นกราฟ DeviceOperation ที่ถูกประเมินผลแบบหน่วงเวลา จัดตารางลงบน stream pool และรอผลด้วย .await
  • ตั้งอยู่บนสมมติฐานว่าผู้ใช้คุ้นเคยกับ ownership, trait และ generic ของ Rust และในบทถัดไปเกี่ยวกับ async GPU programming ก็ต้องมีความรู้เรื่อง async/.await และรันไทม์อย่าง tokio ด้วย
  • มีเอกสารอ้างอิงคือ The Rust Programming Language, Rust by Example, Async Book
  • รีลีส v0.1.0 ยังอยู่ในช่วงอัลฟาระยะแรก และควรคาดว่าจะมี บั๊ก ฟีเจอร์ที่ยังไม่สมบูรณ์ และการเปลี่ยนแปลง API แบบทำลายความเข้ากันได้

บทความที่เกี่ยวข้อง

1 ความคิดเห็น

 
GN⁺ 2 시간 전
ความคิดเห็นจาก Hacker News
  • สุดยอดมาก ใช้ custom CUDA kernel กับ https://crates.io/crates/cudarc มานานแล้ว และนี่ก็ดูเหมือนจะเป็น ตัวแทนแบบ drop-in ได้เกือบเลย
    โดยเฉพาะอยากรู้ว่า เวลา build เทียบกันแล้วเป็นอย่างไร เพราะ Rust CUDA crate ส่วนใหญ่พึ่ง CMake หรือการเรียก nvcc ทำให้การคอมไพล์ช้าจนน่าทรมานได้
    เมื่อสัปดาห์ก่อนเพิ่งโปรไฟล์เวลา build แล้วเห็นว่าเครื่องมืออย่าง sccache ช่วยลดเวลา rebuild ได้มากด้วยการแคชผลลัพธ์ แต่ต้นทุนของการเรียก nvcc แบบ custom ก็ยังอยู่ ตัวอย่างเช่น candle ของ Hugging Face ก็เรียกคำสั่ง nvcc แบบ custom ตอนคอมไพล์ kernel: https://arpadvoros.com/posts/2026/05/05/speeding-up-rust-whi...
    • Cudarc ดีมากจริงๆ
      ส่วนที่บอกว่า Rust CUDA crate ส่วนใหญ่มักเรียก CMake หรือ nvcc จนคอมไพล์ช้านั้น ส่วนตัวไม่ค่อยเจอหนักขนาดนั้น ถ้าดู cuda_setup crate ที่ทำมาเพื่อจัดการ build script มันก็เป็นแค่ build.rs ธรรมดา จึงคอมไพล์ใหม่เฉพาะตอนที่ไฟล์เปลี่ยน และเมื่อเทียบกับโค้ดฝั่ง CPU ของ Rust แล้ว เวลาคอมไพล์ก็น้อยมาก
    • สงสัยเหมือนกันว่าคนอื่นๆ มองว่า cuda-oxide ดูเป็นตัวแทนแบบ drop-in ของ cudarc ได้เกือบเลยหรือเปล่า
      ถ้าเป็นแบบนั้นก็คงดีมาก แต่ส่วนตัวคิดว่าน่าจะเป็น ตัวเสริม มากกว่า และก็อยากรู้เหมือนกันว่าจุดที่ทำให้ cuda-oxide ต่างออกไปคืออะไร นอกจากเรื่องที่ NVIDIA ควบคุมได้ทั้งหมด
  • การไป “ตรงสู่ PTX” ดูแปลกดี ช่วงนี้ NVIDIA MLIR ก็ดีและเร็วพอสมควร หรือจะไปลงที่ Tile IR [1] ซึ่งใหม่กว่าและใช้ง่ายกว่าที่ CuTile ใช้อยู่ก็ยังได้
    Tile IR มีระดับนามธรรมสูงกว่าเล็กน้อย จึง target ได้ง่ายกว่ามาก และจะเสียเปรียบก็แค่เรื่องอย่าง epilogue fusion
    [1] https://docs.nvidia.com/cuda/tile-ir/
    [2] https://developer.nvidia.com/cuda/tile
  • ค่อนข้างอยากรู้ว่าเขาปรับ memory model ของ Rust ให้เข้ากับ semantics ของ CUDA อย่างไร เมื่อเทียบกับ CUDA C++ แล้วต่างกันอย่างไร และ type system ของ Rust จะให้ความปลอดภัยกับ CUDA ได้มากขึ้นจริงไหม
    ผมคิดว่าการเขียน GPU kernel โดยธรรมชาติแล้วไม่ปลอดภัย เพราะทั้งวิธีที่ฮาร์ดแวร์ทำงานและความจำเป็นที่ต้อง optimize จนสุด ทำให้สร้างภาษาที่ปลอดภัยได้ยากมาก
    • ความต่างใหญ่ๆ ที่เห็นมี 4 อย่าง อย่างแรกคือมันจัดการ use-after-free และ semantics ของ drop แทนที่จะต้องเรียก cudaFree เองด้วยมือ
      อย่างที่สอง ใน C++ อาร์กิวเมนต์ void* เป็นแค่อาร์เรย์ของพอยน์เตอร์และตรวจแค่จำนวน แต่ที่นี่บังคับอาร์กิวเมนต์ kernel ผ่าน cuda_launch!
      อย่างที่สามคือปัญหา aliasing ของการเขียนแบบ mutable ใน C++ โค้ดที่สองเธรดขึ้นไปเขียน out[i] ด้วยค่า i เดียวกันก็ยังคอมไพล์ได้ แต่ DisjointSlice และ ThreadIndex ไม่มี constructor แบบ public และมีให้ใช้แค่ API index_1d, index_2d, index_2d_runtime https://github.com/NVlabs/cuda-oxide/blob/2a03dfd9d5f3ecba52...
      อย่างที่สี่ ใน C++ คุณสามารถ cuda memcpy std::string หรือแทบจะ POD อะไรก็ได้จนสถานะพัง แต่ที่นี่รับได้แค่ DisjointSlice, scalar, และ closure 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... แน่นอนว่ามันจับได้ไม่ทุกอย่าง แต่ดูเหมือนให้ guardrail ป้องกัน undefined behavior ได้มากกว่าไฟล์ .cu แบบ raw มาก
    • เผื่อเป็นข้อมูลอ้างอิง memory model ของ Rust ตั้งใจให้เกือบเหมือน C++ แทบทั้งหมดอยู่แล้ว atomic operation ก็เหมือนกัน และก็มีแนวคิดอย่าง provenance ด้วย
      จะเหมาะเป็นภาษาสำหรับ GPU programming แค่ไหนคงต้องดูกันต่อไป แต่ก็ไม่น่าแปลกใจถ้าจะสร้าง API คล้าย DSL ที่ใช้เขียนโค้ดอย่างปลอดภัยได้พอสมควร ขณะเดียวกันก็ยังเข้าถึงความแปลกเฉพาะทางของ GPU ได้ครบ CUDA เองก็ดูเหมือนจะเป็นแบบนั้นไม่ใช่หรือ
    • ในเอกสารอธิบายไว้ค่อนข้างละเอียด มีทั้ง safe layer, ชั้นที่ปลอดภัยเป็นส่วนใหญ่, และชั้นที่ไม่ปลอดภัย
      งานที่ปลอดภัยแต่ขนานกัน และใส่ให้เข้ากับโมเดล Send/Sync ของ Rust ได้ไม่ง่ายนัก จำเป็นต้องยอมรับความขรุขระอยู่บ้าง
    • น่าจะขึ้นอยู่กับเป้าหมาย ถ้ามองจากคนที่เขียนแอปด้วย Rust แล้วอยากใช้การคำนวณบน GPU เป็นครั้งคราวในนั้น ตรงๆ คือไม่ได้สนใจมากขนาดนั้น
      ถ้าใช้ memory model หรือ ownership model ได้แบบ friction ต่ำก็ดี แต่ถ้าต้องแลกกับประสบการณ์ใช้งานที่แย่ลงมากก็ไม่เอา
      baseline สำหรับผมคือแนวทางปัจจุบันของ Cudarc ที่ไม่ได้แทรกแซงการจัดการหน่วยความจำมากนัก เป็นแค่ไวยากรณ์เชิง imperative ที่ห่อ FFI ไว้ กับ build script ไม่กี่บรรทัดที่เรียก nvcc ตอน kernel เปลี่ยน
  • สงสัยว่าสิ่งนี้จะมีความหมายอย่างไรกับ Slang[0] แก่นของเรื่องคงเป็นการที่คนอยากเขียนโปรแกรม GPU ด้วยภาษาที่ทันสมัยกว่า และตอนนี้ก็ดูเหมือนว่าใช้ Rust แทนได้แล้ว
    เอาจริงผมก็ชอบ Slang พอสมควร
    [0]: https://shader-slang.org/
    • อย่างน้อยตอนนี้ การเขียน shader ก็แตกต่างจากการเขียน CUDA kernel อย่างมีนัยสำคัญอยู่ดี shader ทั้งเป็นระดับสูงกว่าและต่ำกว่าพร้อมกัน และมีความแปลกหลายอย่างจากการที่มันถูกออกแบบมาสำหรับชุดความสามารถของ driver/GPU ที่เฉพาะและจำกัด
      เช่น descriptor set, resource register, และข้อจำกัดของ dispatch
    • เป้าหมายคนละแบบ ฝั่ง Slang สนใจ graphics programming มากกว่าอัลกอริทึม AI
      ภาษาสำหรับ shader ก็เป็นมิตรกับผู้ใช้มากกว่าในแง่ฟีเจอร์ด้วย แถม NVIDIA ก็ใช้ Slang ใน production อยู่แล้ว และคนกลุ่มนั้นคงไม่เขียน pipeline ของ shader ใหม่ด้วย Rust หรอก
  • พูดถึง Rust และภาษาโปรแกรมแบบ “ปลอดภัย” ก็สงสัยว่ามีใครรู้รายละเอียดมากกว่านี้ไหมว่า NVIDIA ใช้ Spark/Ada อย่างไร
    ที่หาเจอมีแค่นี้
    https://www.adacore.com/case-studies/nvidia-adoption-of-spar...
  • แค่คำโปรยว่า “ไม่มี DSL, ไม่มี foreign language binding, แค่ Rust” ก็ดูเหมือนเป็น CUDA port อย่างเป็นทางการที่ไม่ได้ใส่ใจกับย่อหน้าเปิดตัวเท่าไรแล้ว
    ถึงอย่างนั้นก็ยังพยายามมองข้ามและไปอ่านเอกสารต่อ แต่พอเริ่มน่าสนใจจากการมี custom IR แล้วกลับมาเจอประโยคทำนองว่า “การทำ MLIR นั้นเป็น C++ ที่โรย TableGen พร้อม build system ที่ต้องคอมไพล์ LLVM ทั้งก้อน และ debugging session ที่ทำให้คุณตั้งคำถามกับเส้นทางอาชีพของตัวเอง” ก็ทำให้ยากจะมองวงการนี้อย่างจริงจังต่อไป
    • โค้ดเบสทั้งหมดดูเหมือนจะ เขียนด้วย AI เป็นส่วนใหญ่
    • ถ้าเว็บเพจนั้นไม่ใช้ AI ก็คงมีคนพูดว่า “ทำไม NVIDIA ไม่ใช้ AI เขียนเว็บไซต์และเอกสารของตัวเองล่ะ? หรือพวกเขาไม่เชื่อเรื่องที่ตัวเองพูดเกี่ยวกับโรงงาน AI และพนักงานที่ดูแลเอเจนต์นับพัน?”
      นี่ดูเหมือนเป็นตัวอย่าง กินอาหารหม้อเดียวกับที่ขาย ได้ตรงตามสไตล์บริษัทที่โหมกระแส AI พอดี
    • ยังตั้งชื่อว่า CUDA-oxide อีก เหมือนไม่รู้ว่าที่มาของชื่อภาษา Rust ไม่ได้มาจากออกไซด์ แต่เป็น เชื้อรา
    • ไม่แน่ใจว่าคุณไม่พอใจอะไรกันแน่ เพราะมีคนสังเกตว่า 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... เขาบอกว่า GPU kernel รันอยู่บนเธรดหลายพันตัวที่มองหน่วยความจำเดียวกัน ขณะที่ Rust บน CPU ใช้ ownership และ borrowing เพื่อป้องกัน data race แต่บน GPU มีได้ถึง 2048 เธรดต่อ SM ที่เริ่มจากฟังก์ชันเดียวกันและชี้ไปยัง output buffer เดียวกัน ดังนั้น borrow checker จึงไม่ได้ถูกออกแบบมาเพื่อสิ่งนี้
    cuda-oxide ทำให้กรณีที่พบบ่อยอย่าง “หนึ่งเธรดเขียนหนึ่งองค์ประกอบ” ปลอดภัยเชิงโครงสร้าง ส่วนกรณีที่พบไม่บ่อยอย่าง shared memory, warp shuffle, และ hardware intrinsic จะต้องใช้ unsafe พร้อมสัญญาการใช้งานที่มีเอกสารกำกับ และฟีเจอร์แนวหน้าพวก TMA, tensor core, และการสื่อสารระดับคลัสเตอร์ ก็ปล่อยให้ทำแบบ manual ทั้งหมดตามความซับซ้อนของฮาร์ดแวร์
    แต่ทั้งหมดนี้ดู ไม่ค่อยเป็น Rust เท่าไร ใน Rust ถ้า abstraction เดิมไม่เข้ากับปัญหา ก็มักจะสร้าง abstraction ใหม่ที่ปลอดภัยขึ้นมา Rust for Linux ก็เป็นตัวอย่าง
    ถ้ามันไม่ปลอดภัย แล้วจะใช้ Rust ไปทำไมกัน ผมโอเคกับการมี unsafe API ให้คนที่ต้องรีด performance ขั้นสุดท้ายใช้ แต่ไม่ควรให้สิ่งนั้นกลายเป็นค่าเริ่มต้น
    มันทำให้นึกถึงไลบรารี user-space สำหรับ API อย่าง io_uring หรือ Vulkan การออกแบบ safe API สำหรับของพวกนี้ยากมาก และก็มีความพยายามหลายครั้งที่สุดท้ายไม่ sound จริง
  • มีใครรู้ไหมว่าสิ่งนี้จะช่วยให้ แชร์ struct ระหว่าง host กับ device ได้หรือเปล่า นี่เป็นส่วนใหญ่ที่หายไปจาก workflow Rust/CUDA แบบเดิมมาตลอด
    เช่นเดียวกับกำแพงเรื่อง serialization/byte ระหว่างสองฝั่งนั้น
  • สิ่งที่ผมระวังเสมอเวลาใช้ Rust กับ CUDA คือ Rust อาจเพิ่ม overhead เล็กๆ น้อยๆ ที่ปกติไม่สำคัญจนแทบมองข้ามได้ แต่ในที่นี้อาจมีผล
    ตัวอย่างเช่น สงสัยว่าการตรวจขอบเขตของอาร์เรย์อาจทำให้ต้องใช้รีจิสเตอร์เพิ่มขึ้น จนลด concurrency ของ kernel ได้ไหม