discussion: what about alterntives?

[Redhawk18]

Have you heard about cubecl/burn or rust-gpu? Cubecl is extremely similar to this project and has gotten to the point where the maintainers have a collection of kernels called cubek.

[dollspace-gay]

https://github.com/dollspace-gay/ferrotorch

[nihalpasham]

Looks interesting

[dollspace-gay]

Theres probably a lot of prior art in there you could benefit from borrowing, Its MIT/Apache dual licenced so feel free.

[nihalpasham]

Hi @Redhawk18 — thanks for raising this. We try to address it directly in the book's Rust + GPU Ecosystem appendix. I just realized CubeCL isn't called out by name there yet — I'll fix that.

To answer your question, there's some surface-level similarity in that both projects let you write GPU kernels in Rust. But up close, as we read the two projects, the design centers point in different directions, and we don't see them as alternatives so much as two projects working on different slices of the same Rust + GPU problem.

CubeCL's design center is portability across hardware. It's an embedded DSL: #[cube] is a proc-macro that parses the function body via syn and emits a second Rust function whose job is to build CubeCL's own IR. That IR is optimized and then lowered by per-backend codegens — to CUDA C++, HIP C++, and Metal C++ source (NVIDIA goes through NVRTC at JIT time to PTX), to SPIR-V, to WGSL, and to CPU. One kernel, six targets — and their own motivation page frames it the same way. The trade-off is that the body of a #[cube] function works with CubeCL types (Array<F>, Vector<F, N>, ABSOLUTE_POS, the Float trait) rather than ordinary Rust ones — it has to be expressible by every backend, including the most constrained.

cuda-oxide's design center is the CUDA programming model in safe Rust on NVIDIA. We're a rustc codegen backend — regular Rust compiles through the standard rustc pipeline to stable MIR, which we import into pliron dialects and lower to PTX. No nvcc, no NVRTC — no source-to-source step. A __global__ written in cuda-oxide reads as ordinary Rust that happens to run on a GPU, with first-class access to NVIDIA-specific intrinsics — warp shuffles, TMA, cluster launch, async pipelines. The trade-off is the obvious one: we target NVIDIA only.

The way we think about it: if your target list is "as many GPUs as possible" and the portable subset of features works for you, CubeCL is the right tool — and cubek on top of it already gives you tuned matmul and friends. If you want the full CUDA programming model expressed in safe Rust on NVIDIA — PTX-level intrinsics, cooperative groups, async pipelines, clusters — that's what we're building.

Rust-GPU, which you also asked about, is covered in the ecosystem page's Other neighbors section — different again: rustc → SPIR-V, design center is graphics shaders, with compute supported but not the primary focus.

Both directions are valuable, and we expect the projects to remain complementary — different parts of the same Rust + GPU stack rather than an alternative. I'll update the appendix to call out CubeCL specifically.

[Redhawk18]

Maybe I misunderstand the example in the readme, but it seems to are still using a rust DSL with a proc macro with your kernel macro. Where instead you use rustc's llvm IR directly to PTX?

[nihalpasham]

To clarify, the #[kernel] macro reads like a DSL but it isn't one. Here's roughly what it expands to:

#[kernel]
fn vecadd(a: &[f32], b: &[f32], mut c: DisjointSlice<f32>) { /* body */ }

becomes:

#[unsafe(no_mangle)]
fn cuda_oxide_kernel_<hash>_vecadd(a: &[f32], b: &[f32], mut c: DisjointSlice<f32>) {
    /* body unchanged */
}
// + a host-side `impl CudaKernel for vecadd` carrying the PTX entry name,
//   used by `cuda_launch!` at the call site.

Three small things: we rename the kernel, so the codegen backend can spot entry points, pin the symbol with no_mangle, and emit a host-side trait impl carrying the PTX entry name. The function body is plain Rust, untouched (the one exception is a single let __scope = ... injected at the top of the block if you use thread::index_*, purely for lifetime hygiene).

The contrast with CubeCL's #[cube] is structural: #[cube] walks the body with syn and rewrites it into a second function whose job is to build CubeCL IR at runtime. Types like Array<F> and ABSOLUTE_POS aren't real values, they're IR-building handles. The body is metaprogramming. In cuda-oxide the body is normal Rust — for/if/closures/generics/?/&[T]/Option — compiled by rustc through its normal pipeline.

One correction on "rustc's LLVM IR directly to PTX": cuda-oxide is a rustc codegen backend. We plug in at stable MIR (not LLVM IR), import it into pliron dialects, and lower to PTX. So the front of the compilation pipeline is the same rustc that compiles any Rust crate; we replace LLVM at the codegen tail.

[NiekAukes]

Hi all,
it's very nice to see this project! For the past two years, we've been prototyping a similar compiler, and it's great to see this idea as a mainline NVidia project. It seems like we've converged on a similar idea, but your project seems to be quite further along. Perhaps, we could learn a thing or two from one another? Below is a repo with our code (still, at a prototype stage) and some papers related to it, with the topmost giving the best overview, and the second going a bit more in-depth on the internals.

https://github.com/NiekAukes/rust-gpu-hybrid-compiler (modified compiler)

https://benevol2025.github.io/pre/paper04.pdf (summary paper + small case study)
https://purl.utwente.nl/essays/100981 (initial iteration)
https://purl.utwente.nl/essays/107371 (working prototype)
https://purl.utwente.nl/essays/107371 (benchmarks)

[nihalpasham]

Thanks for sharing, I skimmed the repo and the paper, and it does seem like we've made a lot of the same design choices: #[kernel] as the entry marker, rustc's front-end shared with the host, and the split happening at MIR.

A couple of curious questions on the architecture:

Compiler integration: Looking at rustc_kernel/rustc_kernel_codegen and the additions to rustc_middle, it looks like you have deeper hooks into rustc than a -Z codegen-backend=... plugin can give. Any reason why?

Device IR: I see ModuleNVVM / FunctionNVVM / TyNVVM / ValueNVVM in rustc_kernel_codegen building textual NVVM IR for libnvvm to JIT at runtime. Did you consider a multi-dialect IR framework (like pliron), or was the single textual-NVVM builder a deliberate choice to keep scope simple?

[NiekAukes]

It is incredibly similar. Most parts are nearly identical between both our projects. I haven't been able to fully dig through your work because I'm a bit busy, but I'll look at it sometime later :).

A few answers to your questions:
The deep hooks into queries and mir lowering were/are necessary in our compiler version. I started with a dev version of 1.80 in early 2024, we've had a contributor try to bump it to 1.89 a few months ago, but he quit the effort. In essence, the codegen backend choice was crate-wide, and would not allow a subsection of the code be built for a different target/architecture than another part of the code. rust-cuda and rust-gpu solve this by putting the GPU and shared code in different crates. However a single-source programming model was one of our primary requirements.

The initial effort focused on allowing the compiler to build for multiple targets at once. In my initial design I was targeting SPIR-V and PTX via the existing backends, but that proved very hard due to version differences and how the compiler and the backends worked internally. I therefore decided to make my own backend, I specifically wanted to avoid any FFI tools since it was very difficult to integrate those into the compiler bootstrapping. Also having 2 different versions of LLVM was very unpractical, as there are FFI conflicts. The choice for NVVM in-text was mostly a practical choice, bytecode LLVM is not very well documented (people just use the LLVM framework), textual NVVM IR was the most understandable, and quickly developed in-time for our submission. The backend we developed is very experimental, and it was the idea that rust-cuda and rust-gpu could become compatible with our compiler (for example for us to specify multiple backends in the compiler)

I hope this answers your questions, feel free to contact me anytime :).

[juntyr]

Great project! Please also mention https://github.com/juntyr/rust-cuda as prior work (different from the rust-cuda you already list). My rust-cuda was, to the best of my knowledge, the first to support generic kernels (though with more hacks required) and safe-by-design kernel launches with slice wrappers (similar to yours) but also support for allocated memory that is moved to or shared with the GPU (something you could get inspired by and copy). My rust-cuda was built for a specific project and I haven’t been able to actively maintain it, so I’m happy that you’re carrying on many of the ideas in a much cleaner form, and it would be awesome if you’d also adopt some of the other ones (e.g. memory transfers). You can find an example of a kernel using all functionality and cross-crate code in https://github.com/juntyr/necsim-rust/blob/c8c3023a114aeacb97e130a60eff779f9c7cb539/rustcoalescence/algorithms/cuda/gpu-kernel/src/lib.rs#L47