How does this compare to nvidia's CUDA-oxide? The latter is similar in syntax to CUDARC on on host side, but replace's the normal-cuda-kernel (in c++ish) on device side with rust.
I ask, because I use CUDA in rust (kernels via cudarc; ML with burn and candle, and cuFFT with FFI), so I am trying to figure out how I would fit this into a workflow.
You got it right: This project exposes our tile programming model, whereas cuda-oxide exposes our lower level CUDA-like programming model. Our tile programming model is higher-level: It compiles to what looks like the CUDA-like programming model.
If you're using burn and candle, and you're writing custom kernels, you can probably write most of your kernels in cutile-rs and let the Tile IR compiler optimize your kernel.
That said, if you're used to writing CUDA, then there is a bit of a learning curve. We have tutorials available that walk you through how it works here: https://nvlabs.github.io/cutile-rs/0.2.0/index.html
Familiarity with numpy helps substantially (it's supposed to have a numpy-like feel), but if you're coming from CUDA and want to leverage the safety features this project provides, then you should jump straight to the "useful mental models" page, which touches on how this compares to CUDA: https://nvlabs.github.io/cutile-rs/main/guide/useful-mental-...
I'm excited to see what cuTile-rs unlocks. Like the direction of HuggingFace's grout https://github.com/huggingface/grout project for local LLM inference:
- state of the art performance
- codebase that fits in a context window (including kernel definitions!)
- single binary deployment
Similar to antirez's ds4.c, but in Rust and with cuTile making kernels both easier to author and higher performance.
Hey! Eric here, one of the folks behind Grout (from HF). The small codebase was a deliberate goal, as the whole engine including kernels is meant to be minimal and readable end to end, which is only practical because cuTile lets us write the kernels in Rust instead of a separate CUDA file. I think this makes things super promising for Rust + CUDA development and rapid iteration!
Yes! That's exactly the spirit. The readable, single-binary, kernels-included codebase is a big part of what makes Grout fun, and the antirez parallel is accurate.
There's a parallel to ThunderKittens too, on the kernel-authoring side: tile-based abstractions for writing fast kernels. The twist with cuTile Rust is the safety layer on top, carrying Rust's ownership model into the kernels: it's a safe, high-performance programming model, not just a perf DSL. The safe surface API is fairly domain-specific today (dense tensor/tile ops), and the Tile IR compiler is still maturing, but it's showing real promise for sparse and multi-GPU. Excited to see where those go. :)
Burn's great, and we cite it in the paper's related work. Different layer, though: Burn is the high-level DL framework (tensor ops, autodiff, backends), whereas cuTile Rust is the kernel-authoring layer underneath. The closer neighbor is actually CubeCL, Tracel's compute layer that Burn sits on, and there's an open issue to integrate cuTile Rust into CubeCL, so we're already looking at how these fit together. See https://github.com/tracel-ai/cubecl/issues/1282
Yeah, I linked to Burn since it was the parent framework, but yes CubeCL. Still blows my mind that they basically implemented a subset of Rust inside of a Rust macro so you can write your compute kernels in Rust.
I built cuTile Rust and just posted the paper preprint. Happy to answer questions.
TL;DR: Rust gives you fearless concurrency on the CPU, but GPU kernel programming still requires unsafe code. cuTile Rust carries Rust's ownership model across the launch boundary and maintains it via a safe tile-based programming model that compiles to Tile IR. The host-side GPU work you write composes into synchronous launches, async pipelines, or a CUDA graph you capture once and replay.
It works by letting you partition mutable output tensors into disjoint pieces on the host: Each tile program gets an exclusive `&mut` view of its piece and the inputs as shared `&` reads. Kernels are written with single-threaded semantics, and the compiler maps that to thread blocks and manages shared memory. Because the pieces are provably disjoint and ordering is threaded through mutable references, you get compile-time data-race freedom when using the safe surface API. As with any other Rust tool, safety remains extensible: Any functionality not yet exposed safely can be made available by writing your own safe abstractions over unsafe code where you supply the invariants yourself. The lower-level Tile IR operation surface (as of CUDA 13.3) is exposed through unsafe intrinsics.
On the B200, an optimized safe GEMM kernel is competitive with cuBLAS: It's within 0.3% of a hand-written low-level (Tile IR) variant at ~92% of the GPU's dense f16 peak. Element-wise is ~7 TB/s. So on these kernels the safety is effectively free. We also worked with Hugging Face to evaluate Grout, a Qwen3 inference engine built on cuTile Rust. In batch-1 Qwen3 decode, Grout reaches 171 tokens/s for Qwen3-4B on an NVIDIA GeForce RTX 5090 and 82 tokens/s for Qwen3-32B on a B200, showing strong performance on memory-bound inference tasks, consistent with our HBM roofline analysis. Benchmarks (harnesses + CSVs), hardware, and clock settings can be found in the repo. The full methodology is available in the paper: https://arxiv.org/abs/2606.15991
This is an early-stage research release: the tensor API is young (a few patterns still need raw pointers), GEMM trails cuBLAS at some sizes, and the tile model gives up SIMT-level control (explicit warp primitives, manual shared memory) in exchange for the semantics that make safety checkable. The recent 0.2.0 release added low-precision support (FP4 packing and block-scaled MMA on CUDA 13.3). These only just landed, so we haven't benchmarked them yet, but we expect them to perform well. Also, while Tile IR is not portable across GPU vendors, it is portable across NVIDIA GPU architectures: What you write in cutile-rs will work on sm_80+ (Ampere and up) with CUDA 13.3, but hardware-specific features such as native FP4 require architectures that support them.
The latest release is on crates.io. After setting up CUDA 13, you can `cargo add cutile` to pull it into your own project, or clone the repo and `cargo run -p cutile-examples --example hello_world` to try it out. Most of what's here has co-evolved with community feedback and contributions since we made the repo public. We read everything that comes in, so anything you raise will shape the direction of this project. If you have a cool feature idea, open an issue or a PR and let's discuss.
cuTile Rust lowers through CUDA Tile IR, which is NVIDIA-specific, so any of those targets (ROCm, Vulkan, Metal, NPUs, OpenCL) would be a new compiler backend. The programming model itself isn't NVIDIA-bound, but the lowering is.
Any thoughts on layering on-GPU work stealing or cudf on top?
For gfql (graph query language mapping down to cudf calls), we're trying to jettison the hot loop of python->cpu->gpu, so been loosely watching cuTile evolve!
I think the answer's actually yes. I'll give the short rationale here, but it's a deep topic and I'd be happy to take it further.
The hot loop you're describing is the whole reason you'd want to reach for a systems language over Python: once the orchestration lives in the interpreter, you pay for the round-trip every iteration. There are now two tile-based ways to do that in a compiled language: CUDA C++ Tile and cuTile Rust. Both get you out of Python and into the same tile model. What cuTile Rust adds is the safety layer: Rust's ownership model gives you compile-time memory safety and data-race freedom (we prove it in the paper), which the C++ path doesn't. So if you're moving off Python anyway, the real question is whether you want those guarantees.
How much of gfql maps cleanly depends on your specific ops and how the pipeline's structured. Happy to dig into it if you open an issue.
Looks interesting. For someone already comfortable with Triton, what's the strongest argument for choosing cuTile Rust instead?
Is the main differentiator the ownership-based safety model, tighter integration with Rust ecosystems, or idk are there workloads where the Tile IR pipeline has advantages beyond language ergonomics?
Yes, the ownership-based safety model and tight Rust integration, plus a pretty fleshed-out host-side tensor API for different workloads (sync, async, replayable CUDA graphs).
Though Triton's a Python DSL, so there's no official Triton Rust to do an apples-to-apples comparison. Being pointer-based, it's unclear how one would reason about memory safety and data-race freedom. As a point-of-comparison, cutile-rs supports both pointers and tensors. We're only able to statically reason about safety when folks operate on tensors directly. The pointer path remains unsafe.
How does this compare to nvidia's CUDA-oxide? The latter is similar in syntax to CUDARC on on host side, but replace's the normal-cuda-kernel (in c++ish) on device side with rust.
I ask, because I use CUDA in rust (kernels via cudarc; ML with burn and candle, and cuFFT with FFI), so I am trying to figure out how I would fit this into a workflow.
You got it right: This project exposes our tile programming model, whereas cuda-oxide exposes our lower level CUDA-like programming model. Our tile programming model is higher-level: It compiles to what looks like the CUDA-like programming model.
If you're using burn and candle, and you're writing custom kernels, you can probably write most of your kernels in cutile-rs and let the Tile IR compiler optimize your kernel.
That said, if you're used to writing CUDA, then there is a bit of a learning curve. We have tutorials available that walk you through how it works here: https://nvlabs.github.io/cutile-rs/0.2.0/index.html
Familiarity with numpy helps substantially (it's supposed to have a numpy-like feel), but if you're coming from CUDA and want to leverage the safety features this project provides, then you should jump straight to the "useful mental models" page, which touches on how this compares to CUDA: https://nvlabs.github.io/cutile-rs/main/guide/useful-mental-...
I'm excited to see what cuTile-rs unlocks. Like the direction of HuggingFace's grout https://github.com/huggingface/grout project for local LLM inference:
- state of the art performance
- codebase that fits in a context window (including kernel definitions!)
- single binary deployment
Similar to antirez's ds4.c, but in Rust and with cuTile making kernels both easier to author and higher performance.
Hey! Eric here, one of the folks behind Grout (from HF). The small codebase was a deliberate goal, as the whole engine including kernels is meant to be minimal and readable end to end, which is only practical because cuTile lets us write the kernels in Rust instead of a separate CUDA file. I think this makes things super promising for Rust + CUDA development and rapid iteration!
Yes! That's exactly the spirit. The readable, single-binary, kernels-included codebase is a big part of what makes Grout fun, and the antirez parallel is accurate.
There's a parallel to ThunderKittens too, on the kernel-authoring side: tile-based abstractions for writing fast kernels. The twist with cuTile Rust is the safety layer on top, carrying Rust's ownership model into the kernels: it's a safe, high-performance programming model, not just a perf DSL. The safe surface API is fairly domain-specific today (dense tensor/tile ops), and the Tile IR compiler is still maturing, but it's showing real promise for sparse and multi-GPU. Excited to see where those go. :)
So Rust is getting cuTile as well.
Wondering when it will make its way into Nsight tooling as well.
See also https://github.com/tracel-ai/burn
Burn's great, and we cite it in the paper's related work. Different layer, though: Burn is the high-level DL framework (tensor ops, autodiff, backends), whereas cuTile Rust is the kernel-authoring layer underneath. The closer neighbor is actually CubeCL, Tracel's compute layer that Burn sits on, and there's an open issue to integrate cuTile Rust into CubeCL, so we're already looking at how these fit together. See https://github.com/tracel-ai/cubecl/issues/1282
Yeah, I linked to Burn since it was the parent framework, but yes CubeCL. Still blows my mind that they basically implemented a subset of Rust inside of a Rust macro so you can write your compute kernels in Rust.
Nathaniel Simard: "Rust for AI & Accelerated Computing" | RustConf 2025 https://www.youtube.com/watch?v=RaSxyRQ7egU
What's going on in this thread? Why so many replies are flagged/dead!!
Its corpo PR
Hello,
I built cuTile Rust and just posted the paper preprint. Happy to answer questions.
TL;DR: Rust gives you fearless concurrency on the CPU, but GPU kernel programming still requires unsafe code. cuTile Rust carries Rust's ownership model across the launch boundary and maintains it via a safe tile-based programming model that compiles to Tile IR. The host-side GPU work you write composes into synchronous launches, async pipelines, or a CUDA graph you capture once and replay.
It works by letting you partition mutable output tensors into disjoint pieces on the host: Each tile program gets an exclusive `&mut` view of its piece and the inputs as shared `&` reads. Kernels are written with single-threaded semantics, and the compiler maps that to thread blocks and manages shared memory. Because the pieces are provably disjoint and ordering is threaded through mutable references, you get compile-time data-race freedom when using the safe surface API. As with any other Rust tool, safety remains extensible: Any functionality not yet exposed safely can be made available by writing your own safe abstractions over unsafe code where you supply the invariants yourself. The lower-level Tile IR operation surface (as of CUDA 13.3) is exposed through unsafe intrinsics.
On the B200, an optimized safe GEMM kernel is competitive with cuBLAS: It's within 0.3% of a hand-written low-level (Tile IR) variant at ~92% of the GPU's dense f16 peak. Element-wise is ~7 TB/s. So on these kernels the safety is effectively free. We also worked with Hugging Face to evaluate Grout, a Qwen3 inference engine built on cuTile Rust. In batch-1 Qwen3 decode, Grout reaches 171 tokens/s for Qwen3-4B on an NVIDIA GeForce RTX 5090 and 82 tokens/s for Qwen3-32B on a B200, showing strong performance on memory-bound inference tasks, consistent with our HBM roofline analysis. Benchmarks (harnesses + CSVs), hardware, and clock settings can be found in the repo. The full methodology is available in the paper: https://arxiv.org/abs/2606.15991
This is an early-stage research release: the tensor API is young (a few patterns still need raw pointers), GEMM trails cuBLAS at some sizes, and the tile model gives up SIMT-level control (explicit warp primitives, manual shared memory) in exchange for the semantics that make safety checkable. The recent 0.2.0 release added low-precision support (FP4 packing and block-scaled MMA on CUDA 13.3). These only just landed, so we haven't benchmarked them yet, but we expect them to perform well. Also, while Tile IR is not portable across GPU vendors, it is portable across NVIDIA GPU architectures: What you write in cutile-rs will work on sm_80+ (Ampere and up) with CUDA 13.3, but hardware-specific features such as native FP4 require architectures that support them.
The latest release is on crates.io. After setting up CUDA 13, you can `cargo add cutile` to pull it into your own project, or clone the repo and `cargo run -p cutile-examples --example hello_world` to try it out. Most of what's here has co-evolved with community feedback and contributions since we made the repo public. We read everything that comes in, so anything you raise will shape the direction of this project. If you have a cool feature idea, open an issue or a PR and let's discuss.
Has there been any thought/exploration in non-nvidia platforms? ROCm, vulkan, metal, npus, opencl, etc?
cuTile Rust lowers through CUDA Tile IR, which is NVIDIA-specific, so any of those targets (ROCm, Vulkan, Metal, NPUs, OpenCL) would be a new compiler backend. The programming model itself isn't NVIDIA-bound, but the lowering is.
Any thoughts on layering on-GPU work stealing or cudf on top?
For gfql (graph query language mapping down to cudf calls), we're trying to jettison the hot loop of python->cpu->gpu, so been loosely watching cuTile evolve!
I think the answer's actually yes. I'll give the short rationale here, but it's a deep topic and I'd be happy to take it further.
The hot loop you're describing is the whole reason you'd want to reach for a systems language over Python: once the orchestration lives in the interpreter, you pay for the round-trip every iteration. There are now two tile-based ways to do that in a compiled language: CUDA C++ Tile and cuTile Rust. Both get you out of Python and into the same tile model. What cuTile Rust adds is the safety layer: Rust's ownership model gives you compile-time memory safety and data-race freedom (we prove it in the paper), which the C++ path doesn't. So if you're moving off Python anyway, the real question is whether you want those guarantees.
How much of gfql maps cleanly depends on your specific ops and how the pipeline's structured. Happy to dig into it if you open an issue.
Looks interesting. For someone already comfortable with Triton, what's the strongest argument for choosing cuTile Rust instead?
Is the main differentiator the ownership-based safety model, tighter integration with Rust ecosystems, or idk are there workloads where the Tile IR pipeline has advantages beyond language ergonomics?
Yes, the ownership-based safety model and tight Rust integration, plus a pretty fleshed-out host-side tensor API for different workloads (sync, async, replayable CUDA graphs).
Though Triton's a Python DSL, so there's no official Triton Rust to do an apples-to-apples comparison. Being pointer-based, it's unclear how one would reason about memory safety and data-race freedom. As a point-of-comparison, cutile-rs supports both pointers and tensors. We're only able to statically reason about safety when folks operate on tensors directly. The pointer path remains unsafe.