1 điểm bởi GN⁺ 2 giờ trước | 1 bình luận | Chia sẻ qua WhatsApp
  • cuda-oxide là một trình biên dịch thử nghiệm cho phép viết kernel GPU SIMT bằng Rust thành ngữ, gần với mức an toàn, và biên dịch trực tiếp mã Rust tiêu chuẩn sang PTX
  • Chỉ dùng Rust mà không cần DSL hay binding ngôn ngữ ngoài, nhưng giả định người dùng hiểu về ownership, trait, generic, và các phần về async cũng cần kiến thức về .await
  • v0.1.0 là bản phát hành alpha ban đầu, nên cần dự kiến sẽ có lỗi, tính năng chưa hoàn thiện và các thay đổi API gây phá vỡ tương thích
  • Ví dụ được chạy bằng cargo oxide run vecadd, trong đó hàm #[kernel] bên trong #[cuda_module] thực hiện phép cộng vector bằng thread::index_1d()
  • #[cuda_module] sẽ nhúng artifact thiết bị vào binary host, đồng thời tạo loader có định kiểu và các phương thức chạy theo từng kernel

Cách dùng và mã được sinh ra

  • Bắt đầu nhanh

    • Sau khi đáp ứng các điều kiện cài đặt, có thể build và chạy ví dụ bằng cargo oxide run vecadd
    • Hướng dẫn cài đặt có tại prerequisites
    • Ví dụ định nghĩa hàm #[kernel]vecadd bên trong module #[cuda_module], lấy chỉ số bằng thread::index_1d() rồi ghi a[i] + b[i] vào DisjointSlice<f32>
    • Ở phía host, dùng CudaContext::new(0), stream mặc định, kernels::load(&ctx), rồi chạy kernel với DeviceBuffer::from_host, DeviceBuffer::<f32>::zeroed, LaunchConfig::for_num_elems(1024)
    • Kết quả thực thi được lấy bằng c.to_host_vec(&stream) và kiểm tra result[0] == 3.0
  • Cách #[cuda_module] hoạt động

    • #[cuda_module] nhúng artifact thiết bị đã được tạo vào binary host
    • Nó tạo ra hàm kernels::load có định kiểu và các phương thức thực thi theo từng kernel
    • Khi cần tải artifact sidecar cụ thể hoặc tự tạo mã thực thi tùy biến, vẫn có thể tiếp tục dùng API mức thấp load_kernel_modulecuda_launch!

Điều kiện và định hướng

  • cuda-oxide hướng tới việc viết kernel GPU bằng hệ thống kiểu và mô hình ownership của Rust, đồng thời đặt an toàn làm mục tiêu hàng đầu
  • Vì GPU có nhiều chi tiết tinh vi, cần đọc the safety model
  • Đây không phải DSL mà là một backend sinh mã rustc tùy biến, biên dịch Rust thuần sang PTX
  • Hỗ trợ thực thi bất đồng bộ bằng cách cấu thành công việc GPU thành đồ thị DeviceOperation được thực thi trì hoãn, lên lịch trên pool stream, rồi chờ kết quả bằng .await
  • Tài liệu giả định người đọc đã quen với ownership, trait, generic của Rust; các chương về lập trình GPU async phía sau còn cần kiến thức về async/.await và runtime như tokio
  • Tài liệu tham khảo gồm The Rust Programming Language, Rust by Example, Async Book
  • Bản phát hành v0.1.0 đang ở giai đoạn alpha ban đầu, nên cần dự kiến có lỗi, tính năng chưa hoàn thiện và thay đổi API gây phá vỡ tương thích

1 bình luận

 
Ý kiến trên Hacker News
  • Thật sự quá ấn tượng. Tôi đã dùng các kernel CUDA tùy chỉnh và https://crates.io/crates/cudarc từ lâu rồi, và cái này có vẻ gần như có thể trở thành một phương án thay thế drop-in
    Điều tôi đặc biệt tò mò là thời gian build sẽ so sánh ra sao. Hầu hết các crate Rust CUDA đều phụ thuộc vào CMake hoặc gọi nvcc, nên việc biên dịch có thể chậm đến mức rất khó chịu
    Tình cờ tuần trước khi profiling thời gian build, tôi thấy các công cụ như sccache có thể giảm đáng kể thời gian build lại nhờ cache đầu ra, nhưng chi phí của các lệnh gọi nvcc tùy chỉnh vẫn còn đó. Ví dụ, candle của Hugging Face cũng gọi lệnh nvcc tùy chỉnh trong lúc biên dịch kernel: https://arpadvoros.com/posts/2026/05/05/speeding-up-rust-whi...
    • Cudarc thực sự rất tốt
      Còn chuyện đa số crate Rust CUDA gọi CMake hay nvcc nên compile chậm thì cá nhân tôi chưa gặp quá nhiều. Nếu nhìn vào crate cuda_setup được tạo để xử lý script build, nó chỉ là một build.rs đơn giản nên chỉ biên dịch lại khi file thay đổi, và thời gian compile của nó rất nhỏ nếu so với phần mã CPU của Rust
    • Tôi cũng tò mò không biết người khác có nghĩ cuda-oxide trông gần như là một phương án thay thế drop-in cho cudarc hay không
      Nếu đúng vậy thì quá tuyệt, nhưng cá nhân tôi đoán nó có lẽ gần với một công cụ bổ sung hơn. Tôi cũng muốn biết điểm khác biệt của cuda-oxide là gì, ngoài việc nó được NVIDIA kiểm soát hoàn toàn
  • Việc đi “thẳng tới PTX” nghe khá lạ. NVIDIA MLIR gần đây cũng khá tốt và nhanh. Hoặc cũng có thể nhắm tới Tile IR [1], thứ dễ dùng hơn và đang thịnh hành hơn mà CuTile sử dụng
    Tile IR ở mức trừu tượng cao hơn một chút nên nhắm mục tiêu sẽ dễ hơn nhiều, chỉ phải đánh đổi ở những chỗ như fusion phần epilogue
    [1] https://docs.nvidia.com/cuda/tile-ir/
    [2] https://developer.nvidia.com/cuda/tile
  • Tôi khá tò mò họ đã ánh xạ mô hình bộ nhớ của Rust sang ngữ nghĩa CUDA như thế nào. Cũng muốn biết nó khác gì so với CUDA C++, và liệu hệ thống kiểu của Rust có thực sự mang lại thêm độ an toàn cho CUDA hay không
    Tôi nghĩ việc viết GPU kernel về bản chất là không an toàn. Do cách phần cứng hoạt động và vì luôn phải tối ưu đến cực hạn nên rất khó tạo ra một ngôn ngữ an toàn cho việc này
    • Có 4 khác biệt lớn dễ thấy. Thứ nhất, nó xử lý use-after-free và ngữ nghĩa drop, thay vì gọi cudaFree thủ công
      Thứ hai, thay vì đối số void* của C++ chỉ là một mảng con trỏ và chỉ kiểm tra số lượng, ở đây cuda_launch! ép buộc kiểu đối số kernel
      Thứ ba là vấn đề aliasing khi ghi có thể thay đổi được. Trong C++, mã mà từ hai hay nhiều thread cùng ghi vào out[i] với cùng i vẫn compile được, nhưng DisjointSliceThreadIndex không có constructor công khai, và chỉ cho phép dùng các API index_1d, index_2d, index_2d_runtime https://github.com/NVlabs/cuda-oxide/blob/2a03dfd9d5f3ecba52...
      Thứ tư, trong C++ bạn có thể cuda memcpy std::string hay thực ra gần như bất kỳ POD nào để làm hỏng trạng thái, nhưng ở đây chỉ nhận DisjointSlice, scalar và closure https://nvlabs.github.io/cuda-oxide/gpu-programming/memory-a...
      Chi tiết có trong https://nvlabs.github.io/cuda-oxide/gpu-safety/the-safety-mo...https://nvlabs.github.io/cuda-oxide/gpu-programming/memory-a.... Dĩ nhiên nó không bắt được mọi thứ, nhưng có vẻ cung cấp nhiều rào chắn an toàn hơn hẳn so với các file .cu raw thuần túy
    • Tham khảo thêm thì mô hình bộ nhớ của Rust được cố ý làm gần như giống hệt C++. Các phép toán nguyên tử cũng giống nhau, và cũng có những khái niệm như provenance
      Liệu đây có phải là ngôn ngữ thuận tiện cho lập trình GPU hay không thì còn phải xem, nhưng sẽ không ngạc nhiên nếu có thể tạo ra một API kiểu DSL đủ ổn để viết mã an toàn mà vẫn khai thác được mọi điểm kỳ quặc riêng của GPU. CUDA rốt cuộc có lẽ cũng là như vậy
    • Trong tài liệu có giải thích khá kỹ. Có lớp an toàn, lớp phần lớn an toàn, và lớp không an toàn
      Với những tác vụ an toàn nhưng song song mà khó nhét gọn vào mô hình Send/Sync của Rust, sẽ cần một chút sự thô ráp
    • Có lẽ còn tùy mục tiêu. Nếu bạn viết ứng dụng bằng Rust và thỉnh thoảng chỉ muốn dùng GPU compute trong đó, thì thật lòng tôi không quá bận tâm
      Sẽ tốt nếu có thể tận dụng mô hình bộ nhớ hay mô hình ownership với mức ma sát thấp. Nhưng nếu vì thế mà trải nghiệm sử dụng trở nên quá bất tiện, thì tôi không muốn kiểu đó
      Đường cơ sở ở đây theo tôi là cách Cudarc làm hiện nay. Quản lý bộ nhớ không can thiệp quá nhiều, chỉ là cú pháp mệnh lệnh bọc quanh FFI và vài dòng script build gọi nvcc khi kernel thay đổi
  • Tôi tự hỏi điều này có ý nghĩa gì với Slang[0]. Trọng tâm có vẻ là mọi người muốn lập trình GPU bằng các ngôn ngữ hiện đại hơn, và giờ thì có vẻ chỉ cần dùng Rust là được
    Nói thêm là tôi khá thích Slang
    [0]: https://shader-slang.org/
    • Việc viết shader, ít nhất là ở hiện tại, về thực chất khác với viết CUDA kernel. Shader vừa ở mức cao hơn vừa ở mức thấp hơn, và có nhiều điểm kỳ lạ do được thiết kế cho một tập tính năng driver/GPU cụ thể và bị giới hạn
      Ví dụ như descriptor set, thanh ghi tài nguyên, hay các giới hạn dispatch
    • Mục tiêu khác nhau. Phía Slang quan tâm đến lập trình đồ họa hơn là các thuật toán AI
      Các ngôn ngữ shader cũng thân thiện hơn với người dùng về mặt tính năng. Hơn nữa, NVIDIA đã dùng Slang trong production rồi, và những người đó sẽ không viết lại pipeline shader của họ bằng Rust đâu
  • Nhân nói về Rust và các ngôn ngữ lập trình “an toàn”, tôi muốn biết liệu có ai hiểu rõ hơn cách NVIDIA dùng Spark/Ada hay không
    Thứ duy nhất tôi tìm được là nội dung dưới đây
    https://www.adacore.com/case-studies/nvidia-adoption-of-spar...
  • Ngay từ câu “không DSL, không binding ngoại ngữ, chỉ là Rust” đã cho cảm giác đây là một cổng CUDA chính thức nhưng đoạn giới thiệu còn không được chăm chút tử tế
    Dù vậy tôi vẫn cố bỏ qua và đọc tài liệu, rồi đúng lúc bắt đầu thấy hứng thú vì có IR tùy chỉnh thì lại gặp kiểu câu như “triển khai MLIR là C++ kèm TableGen, build system bắt bạn phải compile cả LLVM và những phiên debug khiến bạn nghi ngờ lựa chọn sự nghiệp của mình”, từ đó tôi khó mà tiếp tục xem ngành này một cách nghiêm túc
    • Toàn bộ codebase nhìn chung có vẻ như được AI viết
    • Nếu trang web không dùng AI thì đã lại có phản ứng kiểu “tại sao NVIDIA không dùng AI để viết website và tài liệu của chính họ? Họ không tin câu chuyện tự kể về nhân viên quản lý nhà máy AI và hàng nghìn agent à?”
      Điều này trông đúng kiểu dogfooding mà ta có thể kỳ vọng ở một công ty thổi phồng AI
    • Họ còn đặt tên là CUDA-oxide, như thể không biết tên ngôn ngữ Rust bắt nguồn từ nấm chứ không phải quá trình oxy hóa
    • Tôi không rõ chính xác bạn đang khó chịu điều gì. Là vì ai đó nhận xét rằng MLIR rất phức tạp và phụ thuộc vào LLVM sao?
  • Những thứ như TileLang https://github.com/tile-ai/tilelang và Tile Kernels https://github.com/deepseek-ai/TileKernels rồi sẽ khiến CUDA trở nên lỗi thời
    • CUDA đã gần 20 năm tuổi rồi, và nó sẽ không biến mất trong vài năm tới đâu
    • Đây là một tuyên bố khá lớn nhưng cơ sở lại quá ít
  • Khi đọc tài liệu https://nvlabs.github.io/cuda-oxide/gpu-safety/the-safety-mo..., họ nói rằng GPU kernel chạy trên hàng nghìn thread cùng nhìn vào một vùng nhớ, trong khi trên CPU, Rust dùng ownership và borrowing để ngăn data race, nhưng trên GPU thì có tới 2048 thread mỗi SM bắt đầu từ cùng một hàm và cùng trỏ tới một buffer đầu ra, nên borrow checker không được thiết kế cho chuyện đó
    cuda-oxide làm cho trường hợp phổ biến là “mỗi thread ghi một phần tử” trở nên an toàn về mặt cấu trúc, còn các trường hợp hiếm hơn như shared memory, warp shuffle hay hardware intrinsic thì yêu cầu unsafe với các hợp đồng được ghi rõ trong tài liệu, và các tính năng tuyến đầu như TMA, tensor core hay giao tiếp cấp cluster thì để hoàn toàn thủ công cho tương xứng với độ phức tạp của phần cứng
    Nhưng điều này lại không quá đậm chất Rust. Trong Rust, nếu các trừu tượng hiện có không hợp với bài toán thì người ta sẽ tạo các trừu tượng an toàn mới. Rust for Linux là ví dụ như vậy
    Nếu không an toàn thì tôi tự hỏi lý do dùng Rust là gì. Việc cung cấp API unsafe cho những người cần vắt kiệt hiệu năng cuối cùng thì ổn, nhưng nó không nên là mặc định
    Nó khiến tôi liên tưởng đến các thư viện không gian người dùng cho các API như io_uring hay Vulkan. Thiết kế API an toàn cho những thứ đó khá khó, và thực tế đã có những nỗ lực không sound
  • Tôi muốn biết liệu cái này có cho phép chia sẻ struct giữa host và device hay không. Đó chính là phần lớn còn thiếu từ trước đến nay trong quy trình làm việc Rust/CUDA
    Rào cản serialize/byte giữa hai bên cũng vậy
  • Điều khiến tôi dè dặt khi dùng Rust trong CUDA là Rust có thể thêm một chút overhead mà bình thường có thể bỏ qua, nhưng ở đây lại có thể quan trọng
    Ví dụ, tôi muốn biết liệu kiểm tra biên mảng có thể làm phát sinh thêm việc dùng thanh ghi, từ đó làm giảm mức độ đồng thời của kernel hay không