How Does Rust Code Run on a GPU? Tracing cuda-oxide and PTX
I used cuda-oxide to run a vector addition kernel written in Rust on a GPU.
Pass in two arrays and get back their element-wise sums. The calculation is simple, but I wanted to understand what form that Rust function takes before it reaches the GPU.
Among the files generated during the build was some code called PTX. This time, I will read it to see how Rust's thread indexing and array accesses are represented.
What is cuda-oxide?
cuda-oxide is part of NVIDIA's CUDA Rust project. It provides a compiler and related tools for writing kernels for NVIDIA GPUs in Rust. See the official CUDA Rust repository.
The official book introduces it this way:
> cuda-oxide is a Rust-to-CUDA compiler that lets you write (SIMT) GPU kernels in safe(ish), idiomatic Rust.
Source: The cuda-oxide Book.
In other words, you can write the actual function that runs on the GPU in idiomatic Rust and compile it. SIMT stands for Single Instruction, Multiple Threads. It is the programming model used in this example, where many threads execute the same kernel.
As of October 2026, CUDA Rust describes itself as an early-stage research project, with its features and APIs still under development. Here, I will focus on a small vector addition example and look at what the Rust code becomes. See the project status.
As someone who writes Rust, I find this a very intriguing project.
Adding arrays: who computes each element?
Here is the kernel we will inspect. A kernel is a function that serves as an entry point for execution on the GPU.
I verified the code and output on November 1, 2026, under the following conditions. I built the vecadd example from cuda-rust commit a0cc6cc0bc18722ec5cdfbea136c118a7c051977 in a source tree separate from my ongoing changes.
| Item | Verified environment or setting |
| --- | --- |
| Host | Linux, x86-64 |
| GPU | NVIDIA GeForce RTX 5060 Ti (Compute Capability 12.0) |
| NVIDIA driver | 595.99.02 |
| Rust toolchain | nightly-2026-08-28, rustc 1.100.0-nightly (e457a7b0d) |
| LLVM | Version 23.1.0 bundled with the toolchain |
| Build | Release, target sm_120a, ordinary PTX output with bounds checks enabled |
use cuda_device::{DisjointSlice, cuda_module, kernel, thread};
#[cuda_module]
mod kernels {
use super::*;
#[kernel]
pub fn vecadd(a: &[f32], b: &[f32], mut c: DisjointSlice<f32>) {
let idx = thread::index_1d();
let idx_raw = idx.get();
if let Some(c_elem) = c.get_mut(idx) {
*c_elem = a[idx_raw] + b[idx_raw];
}
}
}
This is based on the vecadd example included in NVIDIA's cuda-rust repository. The original code is published under the Apache-2.0 license. See the cuda-rust source.
The calculation is c[i] = a[i] + b[i]. However, there is no for loop inside the function to traverse the whole array.
On a CPU, we could compute the elements sequentially like this:
for i in 0..n {
c[i] = a[i] + b[i];
}
In this GPU kernel, each thread works out the index of the element it is responsible for and computes that one element. Many threads execute the same kernel.
Thread responsible for element 0: c[0] = a[0] + b[0]
Thread responsible for element 1: c[1] = a[1] + b[1]
Thread responsible for element 2: c[2] = a[2] + b[2]
...
thread::index_1d() produces this index. Its return value, idx, is not just an integer: it has cuda-oxide's ThreadIndex type. idx.get() extracts the integer used to index the input arrays.
On the output side, DisjointSlice accepts this thread index and provides a mutable reference to the element assigned to the thread. If c.get_mut(idx) returns Some, the thread writes its result to that element. I will leave the details of these types aside for now and focus on which elements each thread reads and where it writes.
On the host, meaning the CPU side, the program transfers the input data to GPU buffers, specifies the thread configuration, launches the kernel, and retrieves the result. We need to distinguish the whole program, including that preparation, from the kernel running on the GPU.
PTX is a representation before GPU machine code
So what is the PTX generated from this kernel?
PTX stands for Parallel Thread Execution. It is a virtual instruction set for NVIDIA GPUs. It is not specific to Rust: compilers for languages such as CUDA C++ generate PTX too. See the NVIDIA PTX specification.
PTX contains instruction names and register-like names, but another conversion step remains before we have the machine code the actual GPU executes. You will also encounter the name SASS for the instructions of the physical GPU.
Here, I will use PTX to see how the Rust expression is broken down into operations. I will not treat it as output that already fixes the final machine-instruction count or physical-register allocation.
The vecadd.ptx file starts with these directives:
.version 8.7
.target sm_120a
.address_size 64
They specify the PTX version, target GPU architecture, and address width, respectively. PTX has version and target requirements, so the same file is not necessarily usable with every NVIDIA GPU or driver.
After these directives come the kernel entry point, .visible .entry vecadd(, its parameters, and register-variable declarations. Let's look at what happens in the body.
Where does cuda-oxide connect to rustc?
Before reading the PTX, let's establish the path from Rust source to this point.
cuda-oxide provides a compiler, libraries, and build tools for writing CUDA kernels in Rust. At its center is rustc-codegen-cuda, a custom code generation backend for rustc.
It uses rustc for tasks such as parsing Rust, type checking, and borrow checking. GPU code generation takes the resulting MIR and proceeds by transforming it into other representations. See the cuda-oxide architecture.
The ordinary PTX generation path we are following can be simplified like this:
Rust kernel
| rustc performs parsing, type checking, borrow checking, and other work
v
MIR
| cuda-oxide transforms it using Pliron IR
v
LLVM IR
| LLVM's llc generates code with the NVPTX backend
v
PTX
Pliron is a framework for building and transforming intermediate representations. cuda-oxide uses it to handle representations of Rust MIR and lower them to LLVM IR.
LLVM produces the PTX at the end of this path. cuda-oxide writes out LLVM IR and passes it to LLVM's code generation tool, llc. Here too, LLVM provides the infrastructure for generating code for a different target.
We can break down the roles as follows:
| Name | Role in this path |
| --- | --- |
| rustc | Rust analysis, checking, and MIR processing |
| rustc-codegen-cuda and related crates | Extract GPU-side work and transform MIR into LLVM IR |
| LLVM's llc and NVPTX backend | Generate PTX from LLVM IR |
| cargo-oxide | Provide commands for building and related operations |
Host-side code in the same file, such as main, goes to the ordinary LLVM backend and becomes CPU code. In this example, the host executable also contains the GPU's PTX.
Calling the body of a #[kernel] function as an ordinary Rust function on the CPU does not move execution onto the GPU. The host uses a launch API to invoke it as a GPU kernel.
Finding one line of Rust in the PTX
For this verification, I built and ran the example from the cuda-oxide workspace, specifying the example and target. --verbose also displays information such as the tools used to generate LLVM IR and PTX.
cargo oxide run vecadd --arch sm_120a --verbose
The final output was as follows. The program compared the results with values calculated on the CPU, and all 1024 elements matched.
✓ SUCCESS: All 1024 elements correct!
The same build generated crates/rustc-codegen-cuda/examples/vecadd/vecadd.ptx. The following excerpts come from that file, with whitespace adjusted for readability.
Working out the element index
First, let's find the instructions related to thread::index_1d(). Omitting an intervening parameter-load instruction gives us this sequence:
mov.u32 %r1, %ctaid.x;
mov.u32 %r2, %ntid.x;
mov.u32 %r3, %tid.x;
mul.wide.u32 %rd11, %r1, %r2;
cvt.u64.u32 %rd12, %r3;
add.s64 %rd13, %rd11, %rd12;
CUDA groups threads into blocks and arranges those blocks in a grid. In one dimension, the global element index is the block index multiplied by the number of threads per block, plus the thread's index within its block. See CUDA's thread hierarchy.
The first three mov instructions read the values used in that calculation.
| Special register | Value read |
| --- | --- |
| %ctaid.x | Block index in the x direction |
| %ntid.x | Number of threads per block in the x direction |
| %tid.x | Thread index within the block in the x direction |
For example, with 256 threads per block, block index 2, and thread index 7 within the block, the element index is 2 * 256 + 7 = 519.
In the PTX, mul.wide.u32 multiplies two 32-bit values and places the 64-bit result in %rd11. The next cvt converts the thread index within the block to 64 bits, and add.s64 adds the values. %rd13 now holds the element index.
What was a single function call in Rust has become reads of the block and thread coordinates, a multiplication, and an addition.
Names such as %r1 and %rd13 refer to virtual register variables in PTX. We do not need to interpret their numbers as physical GPU register numbers.
From an element index to a memory location
Next come the array accesses.
The index we just calculated is in %rd13, but the later instructions use %rd14. This instruction sits between them:
selp.b64 %rd14, %rd13, -1, %p7;
selp selects between two values according to a condition. If %p7 is true, it places the value of %rd13 in %rd14; otherwise, it selects -1 as an invalid value. In this output, %p7 represents the one-dimensional launch condition: the y and z dimensions of both the grid and the block must all be 1.
When that condition holds, %rd14 contains the element index calculated earlier. One instruction that uses it is:
shl.b64 %rd15, %rd14, 2;
shl is a left shift. Here, it shifts the element index left by two bits, multiplying it by four.
Each element of these arrays is an f32, which occupies four bytes. For element 519, the distance from the start of the array is 519 * 4 = 2076 bytes.
The PTX also adds that distance to the starting addresses. The relevant instructions are:
add.s64 %rd3, %rd8, %rd15;
add.s64 %rd18, %rd1, %rd15;
add.s64 %rd19, %rd2, %rd15;
In this output, %rd1 points to the start of input a, %rd2 to the start of input b, and %rd8 to the start of output c. Adding the byte offset produces the addresses to read from and write to.
The compact expression a[idx_raw] hides the element size and the address calculation.
Read two values, add them, and write back
Now that we have the addresses, let's look at the calculation. In the actual output, input-address calculations alternate with loads.
add.s64 %rd18, %rd1, %rd15;
ld.global.b32 %r8, [%rd18];
add.s64 %rd19, %rd2, %rd15;
ld.global.b32 %r9, [%rd19];
add.f32 %r10, %r8, %r9;
st.global.b32 [%rd3], %r10;
ld.global.b32 loads 32 bits from global memory. Square brackets, as in [%rd18], indicate an access to the value at that address. The two inputs are loaded into %r8 and %r9.
Next, add.f32 performs a 32-bit floating-point addition, placing the sum in %r10. The .b32 loads handle 32-bit bit patterns, which the addition interprets as f32 values. There is no integer-to-floating-point conversion in between.
Finally, st.global.b32 writes the result to the output address in %rd3. Unlike a load, the bracketed destination appears on the left.
We can now see the work corresponding to this line of Rust:
*c_elem = a[idx_raw] + b[idx_raw];
The addition itself is one instruction, but the surrounding code computes the element index, calculates memory addresses, and reads and writes the values.
Bounds checks remain as instructions too
Around that calculation are instructions that test conditions and branch.
For example, before accessing input a, we find:
ld.param.b64 %rd4, [vecadd_param_1];
setp.ge.u64 %p13, %rd14, %rd4;
@%p13 bra $L__BB0_5;
In this output, vecadd_param_1 is the length of input a. It is passed as a parameter separate from the starting address.
setp.ge.u64 tests whether the element index in %rd14 is greater than or equal to the length in %rd4 and places the result in the predicate variable %p13. The next instruction, @%p13 bra, branches to the named label if that condition is true.
What is at the destination?
$L__BB0_5:
trap;
trap aborts execution and notifies the host of an error. In this output, an out-of-bounds access to input a takes this path. Input b has a similar check. The bounds checks associated with Rust's indexing are visible as PTX comparisons and branches. See the PTX trap instruction.
The output is handled differently: the Rust code writes only when c.get_mut(idx) returns Some. Skipping the calculation when the output index is out of bounds and trapping when an input index is out of bounds are two different paths.
The number of GPU threads can be slightly larger than the number of elements. For example, launching four blocks of 256 threads for 1000 elements gives 1024 threads. The extra 24 threads must not write past the end of the array. These index checks have a practical role.
What remains after PTX is generated?
We have read the PTX, but the journey to GPU execution is not finished.
On the path used here, where PTX is passed to the CUDA driver for loading, the driver's JIT compiler translates it into machine code for the target GPU. It can also cache the compiled result for reuse. See NVIDIA's explanation of JIT compilation.
The overall flow is:
Build time
Rust kernel -> MIR -> Pliron IR -> LLVM IR -> PTX
Rust host code -> CPU code
|
v
Embedded in the host executable
Run time
The host program starts
| Pass PTX to the CUDA driver and load the module
v
Machine code for the target GPU (using JIT compilation or its cache)
| Launch the kernel from the host
v
GPU threads compute and write results to memory
|
v
The host waits for completion and retrieves the results
There are also paths that build a GPU binary from PTX ahead of time. Here, I am following the path that loads PTX.
Running Rust code on a GPU involves more than sending the Rust source to it. There is checking in rustc, transformation in cuda-oxide, PTX generation in LLVM, and then GPU-specific compilation and launch work.
What we found in the PTX was thread-index acquisition, element-address calculation, two loads, one addition, and a store. That short array expression is becoming readable as the work each GPU thread performs.
Mentions
CUDA Rust: NVIDIA's CUDA platform for Rust (3.9k stars) What it does All text below is quoted directly from the project README. >NVIDIA's CUDA platform for Rust. Host runtime crates plus Tile (cutile-rs) and SIMT (cuda-oxide) kernel programming models in idiomatic Rust. CUDA Rust: - cuda-oxide: A custom rustc codegen backend, cargo helper utility, and crates to help users target the traditional CUDA SIMT (Single Instruction, Multiple Threads) kernel programming model. Here users have access to every detail of device-side CUDA programming. - cutile-rs: A proc-macro plugin and crates to help users write CUDA kernels using the CUDA Tile programming model. - Runtime: cuda-bindings, cuda-core (with cuda-core-derive), and cuda-async provide a shared runtime: driver bindings, contexts, streams, device buffers, events, and lazy, composable device operations. - Both kernel programming models extend Rust's ownership rules across the GPU launch boundary: mutable buffers are partitioned into disjoint pieces before launch, immutable buffers are shared, and launchers keep those borrows alive while GPU work is in flight. - Kernels from both models interoperate: a cutile-rs Tile kernel and a cuda-oxide SIMT kernel can run on the same stream over shared device buffers. >CUDA Rust is an early-stage research project. Expect bugs, incomplete features, and API breakage. Repository NVIDIA/cuda-rust Repository stats |Metric|Value| |:-|:-| |Stars|3,901| |Forks|350| |Last release|v0.2.1 (2026-06-10)| |Languages|Rust 93.1%, Python 3%, LLVM 2.1%| |Contributors|82| |AI contributor accounts (Claude / Codex / Cursor / Copilot)|@claude (15 commits)| |AI co-authored commits|Claude: 17| Main contributor |Field|Value| |:-|:-| |GitHub|@nihalpasham| |Commits to this repo|790 (top contributor)| |Bio|Rust @NVIDIA cuda-oxide author Proof-oriented programming| |Company / Location|Nvidia / Bangalore, India| |Followers|365| ────────────────────────────── 📌 If you want to browse more repos like this, we've organized them by category and rankings here: bestgithubrepositories.com 🗂️ We're also archiving repos where possible, so useful projects don't disappear when they're removed from GitHub.
How Does Rust Code Run on a GPU? Tracing cuda-oxide and PTX I used cuda-oxide to run a vector addition kernel written in Rust on a GPU. Pass in two arrays and get back their element-wise sums. The calculation is simple, but I wanted to understand what form that Rust function takes before it reaches the GPU. Among the files generated during the build was some code called PTX. This time, I will read it to see how Rust's thread indexing and array accesses are represented. What is cuda-oxide? cuda-oxide is part of NVIDIA's CUDA Rust project. It provides a compiler and related tools for writing kernels for NVIDIA GPUs in Rust. See the official CUDA Rust repository. The official book introduces it this way: > cuda-oxide is a Rust-to-CUDA compiler that lets you write (SIMT) GPU kernels in safe(ish), idiomatic Rust. Source: The cuda-oxide Book. In other words, you can write the actual function that runs on the GPU in idiomatic Rust and compile it. SIMT stands for Single Instruction, Multiple Threads. It is the programming model used in this example, where many threads execute the same kernel. As of October 2026, CUDA Rust describes itself as an early-stage research project, with its features and APIs still under development. Here, I will focus on a small vector addition example and look at what the Rust code becomes. See the project status. As someone who writes Rust, I find this a very intriguing project. Adding arrays: who computes each element? Here is the kernel we will inspect. A kernel is a function that serves as an entry point for execution on the GPU. I verified the code and output on November 1, 2026, under the following conditions. I built the vecadd example from cuda-rust commit a0cc6cc0bc18722ec5cdfbea136c118a7c051977 in a source tree separate from my ongoing changes. | Item | Verified environment or setting | | --- | --- | | Host | Linux, x86-64 | | GPU | NVIDIA GeForce RTX 5060 Ti (Compute Capability 12.0) | | NVIDIA driver | 595.99.02 | | Rust toolchain | nightly-2026-08-28, rustc 1.100.0-nightly (e457a7b0d) | | LLVM | Version 23.1.0 bundled with the toolchain | | Build | Release, target sm_120a, ordinary PTX output with bounds checks enabled | use cuda_device::{DisjointSlice, cuda_module, kernel, thread}; #[cuda_module] mod kernels { use super::*; #[kernel] pub fn vecadd(a: &[f32], b: &[f32], mut c: DisjointSlice<f32>) { let idx = thread::index_1d(); let idx_raw = idx.get(); if let Some(c_elem) = c.get_mut(idx) { *c_elem = a[idx_raw] + b[idx_raw]; } } } This is based on the vecadd example included in NVIDIA's cuda-rust repository. The original code is published under the Apache-2.0 license. See the cuda-rust source. The calculation is c[i] = a[i] + b[i]. However, there is no for loop inside the function to traverse the whole array. On a CPU, we could compute the elements sequentially like this: for i in 0..n { c[i] = a[i] + b[i]; } In this GPU kernel, each thread works out the index of the element it is responsible for and computes that one element. Many threads execute the same kernel. Thread responsible for element 0: c[0] = a[0] + b[0] Thread responsible for element 1: c[1] = a[1] + b[1] Thread responsible for element 2: c[2] = a[2] + b[2] ... thread::index_1d() produces this index. Its return value, idx, is not just an integer: it has cuda-oxide's ThreadIndex type. idx.get() extracts the integer used to index the input arrays. On the output side, DisjointSlice accepts this thread index and provides a mutable reference to the element assigned to the thread. If c.get_mut(idx) returns Some, the thread writes its result to that element. I will leave the details of these types aside for now and focus on which elements each thread reads and where it writes. On the host, meaning the CPU side, the program transfers the input data to GPU buffers, specifies the thread configuration, launches the kernel, and retrieves the result. We need to distinguish the whole program, including that preparation, from the kernel running on the GPU. PTX is a representation before GPU machine code So what is the PTX generated from this kernel? PTX stands for Parallel Thread Execution. It is a virtual instruction set for NVIDIA GPUs. It is not specific to Rust: compilers for languages such as CUDA C++ generate PTX too. See the NVIDIA PTX specification. PTX contains instruction names and register-like names, but another conversion step remains before we have the machine code the actual GPU executes. You will also encounter the name SASS for the instructions of the physical GPU. Here, I will use PTX to see how the Rust expression is broken down into operations. I will not treat it as output that already fixes the final machine-instruction count or physical-register allocation. The vecadd.ptx file starts with these directives: .version 8.7 .target sm_120a .address_size 64 They specify the PTX version, target GPU architecture, and address width, respectively. PTX has version and target requirements, so the same file is not necessarily usable with every NVIDIA GPU or driver. After these directives come the kernel entry point, .visible .entry vecadd(, its parameters, and register-variable declarations. Let's look at what happens in the body. Where does cuda-oxide connect to rustc? Before reading the PTX, let's establish the path from Rust source to this point. cuda-oxide provides a compiler, libraries, and build tools for writing CUDA kernels in Rust. At its center is rustc-codegen-cuda, a custom code generation backend for rustc. It uses rustc for tasks such as parsing Rust, type checking, and borrow checking. GPU code generation takes the resulting MIR and proceeds by transforming it into other representations. See the cuda-oxide architecture. The ordinary PTX generation path we are following can be simplified like this: Rust kernel | rustc performs parsing, type checking, borrow checking, and other work v MIR | cuda-oxide transforms it using Pliron IR v LLVM IR | LLVM's llc generates code with the NVPTX backend v PTX Pliron is a framework for building and transforming intermediate representations. cuda-oxide uses it to handle representations of Rust MIR and lower them to LLVM IR. LLVM produces the PTX at the end of this path. cuda-oxide writes out LLVM IR and passes it to LLVM's code generation tool, llc. Here too, LLVM provides the infrastructure for generating code for a different target. We can break down the roles as follows: | Name | Role in this path | | --- | --- | | rustc | Rust analysis, checking, and MIR processing | | rustc-codegen-cuda and related crates | Extract GPU-side work and transform MIR into LLVM IR | | LLVM's llc and NVPTX backend | Generate PTX from LLVM IR | | cargo-oxide | Provide commands for building and related operations | Host-side code in the same file, such as main, goes to the ordinary LLVM backend and becomes CPU code. In this example, the host executable also contains the GPU's PTX. Calling the body of a #[kernel] function as an ordinary Rust function on the CPU does not move execution onto the GPU. The host uses a launch API to invoke it as a GPU kernel. Finding one line of Rust in the PTX For this verification, I built and ran the example from the cuda-oxide workspace, specifying the example and target. --verbose also displays information such as the tools used to generate LLVM IR and PTX. cargo oxide run vecadd --arch sm_120a --verbose The final output was as follows. The program compared the results with values calculated on the CPU, and all 1024 elements matched. ✓ SUCCESS: All 1024 elements correct! The same build generated crates/rustc-codegen-cuda/examples/vecadd/vecadd.ptx. The following excerpts come from that file, with whitespace adjusted for readability. Working out the element index First, let's find the instructions related to thread::index_1d(). Omitting an intervening parameter-load instruction gives us this sequence: mov.u32 %r1, %ctaid.x; mov.u32 %r2, %ntid.x; mov.u32 %r3, %tid.x; mul.wide.u32 %rd11, %r1, %r2; cvt.u64.u32 %rd12, %r3; add.s64 %rd13, %rd11, %rd12; CUDA groups threads into blocks and arranges those blocks in a grid. In one dimension, the global element index is the block index multiplied by the number of threads per block, plus the thread's index within its block. See CUDA's thread hierarchy. The first three mov instructions read the values used in that calculation. | Special register | Value read | | --- | --- | | %ctaid.x | Block index in the x direction | | %ntid.x | Number of threads per block in the x direction | | %tid.x | Thread index within the block in the x direction | For example, with 256 threads per block, block index 2, and thread index 7 within the block, the element index is 2 * 256 + 7 = 519. In the PTX, mul.wide.u32 multiplies two 32-bit values and places the 64-bit result in %rd11. The next cvt converts the thread index within the block to 64 bits, and add.s64 adds the values. %rd13 now holds the element index. What was a single function call in Rust has become reads of the block and thread coordinates, a multiplication, and an addition. Names such as %r1 and %rd13 refer to virtual register variables in PTX. We do not need to interpret their numbers as physical GPU register numbers. From an element index to a memory location Next come the array accesses. The index we just calculated is in %rd13, but the later instructions use %rd14. This instruction sits between them: selp.b64 %rd14, %rd13, -1, %p7; selp selects between two values according to a condition. If %p7 is true, it places the value of %rd13 in %rd14; otherwise, it selects -1 as an invalid value. In this output, %p7 represents the one-dimensional launch condition: the y and z dimensions of both the grid and the block must all be 1. When that condition holds, %rd14 contains the element index calculated earlier. One instruction that uses it is: shl.b64 %rd15, %rd14, 2; shl is a left shift. Here, it shifts the element index left by two bits, multiplying it by four. Each element of these arrays is an f32, which occupies four bytes. For element 519, the distance from the start of the array is 519 * 4 = 2076 bytes. The PTX also adds that distance to the starting addresses. The relevant instructions are: add.s64 %rd3, %rd8, %rd15; add.s64 %rd18, %rd1, %rd15; add.s64 %rd19, %rd2, %rd15; In this output, %rd1 points to the start of input a, %rd2 to the start of input b, and %rd8 to the start of output c. Adding the byte offset produces the addresses to read from and write to. The compact expression a[idx_raw] hides the element size and the address calculation. Read two values, add them, and write back Now that we have the addresses, let's look at the calculation. In the actual output, input-address calculations alternate with loads. add.s64 %rd18, %rd1, %rd15; ld.global.b32 %r8, [%rd18]; add.s64 %rd19, %rd2, %rd15; ld.global.b32 %r9, [%rd19]; add.f32 %r10, %r8, %r9; st.global.b32 [%rd3], %r10; ld.global.b32 loads 32 bits from global memory. Square brackets, as in [%rd18], indicate an access to the value at that address. The two inputs are loaded into %r8 and %r9. Next, add.f32 performs a 32-bit floating-point addition, placing the sum in %r10. The .b32 loads handle 32-bit bit patterns, which the addition interprets as f32 values. There is no integer-to-floating-point conversion in between. Finally, st.global.b32 writes the result to the output address in %rd3. Unlike a load, the bracketed destination appears on the left. We can now see the work corresponding to this line of Rust: *c_elem = a[idx_raw] + b[idx_raw]; The addition itself is one instruction, but the surrounding code computes the element index, calculates memory addresses, and reads and writes the values. Bounds checks remain as instructions too Around that calculation are instructions that test conditions and branch. For example, before accessing input a, we find: ld.param.b64 %rd4, [vecadd_param_1]; setp.ge.u64 %p13, %rd14, %rd4; @%p13 bra $L__BB0_5; In this output, vecadd_param_1 is the length of input a. It is passed as a parameter separate from the starting address. setp.ge.u64 tests whether the element index in %rd14 is greater than or equal to the length in %rd4 and places the result in the predicate variable %p13. The next instruction, @%p13 bra, branches to the named label if that condition is true. What is at the destination? $L__BB0_5: trap; trap aborts execution and notifies the host of an error. In this output, an out-of-bounds access to input a takes this path. Input b has a similar check. The bounds checks associated with Rust's indexing are visible as PTX comparisons and branches. See the PTX trap instruction. The output is handled differently: the Rust code writes only when c.get_mut(idx) returns Some. Skipping the calculation when the output index is out of bounds and trapping when an input index is out of bounds are two different paths. The number of GPU threads can be slightly larger than the number of elements. For example, launching four blocks of 256 threads for 1000 elements gives 1024 threads. The extra 24 threads must not write past the end of the array. These index checks have a practical role. What remains after PTX is generated? We have read the PTX, but the journey to GPU execution is not finished. On the path used here, where PTX is passed to the CUDA driver for loading, the driver's JIT compiler translates it into machine code for the target GPU. It can also cache the compiled result for reuse. See NVIDIA's explanation of JIT compilation. The overall flow is: Build time Rust kernel -> MIR -> Pliron IR -> LLVM IR -> PTX Rust host code -> CPU code | v Embedded in the host executable Run time The host program starts | Pass PTX to the CUDA driver and load the module v Machine code for the target GPU (using JIT compilation or its cache) | Launch the kernel from the host v GPU threads compute and write results to memory | v The host waits for completion and retrieves the results There are also paths that build a GPU binary from PTX ahead of time. Here, I am following the path that loads PTX. Running Rust code on a GPU involves more than sending the Rust source to it. There is checking in rustc, transformation in cuda-oxide, PTX generation in LLVM, and then GPU-specific compilation and launch work. What we found in the PTX was thread-index acquisition, element-address calculation, two loads, one addition, and a store. That short array expression is becoming readable as the work each GPU thread performs.
// AI Infrastructure 7/10 NVIDIA open-sourced CUDA for Rust with both Tile and SIMT kernel models plus host runtime crates. Rust folks have wanted this for years, curious how idiomatic it actually feels. NVIDIA/cuda-rust
Rust 写 GPU 内核，还得绑 C++？ NVIDIA 的 cuda-rust 是 Rust 版 CUDA 平台：SIMT 和 Tile 两种内核写法，所有权跟到 GPU。 适合 Rust 写 GPU 的人，CUDA 老手走 SIMT。 选 SIMT 还是 Tile？ #Rust #CUDA @grok NVIDIA/cuda-rust
Let's see if this gonna Merge or Kicked ! cuda-rust#1627 t.co
I added Rust support to GPUFlight's CUDA Performance Lab Hi all, I previously share GPUFlight here, a CUDA profiler I'm building. I've now added Rust support to its Performance Lab and Workbench. After seeing NVIDIA’s CUDA Rust announcement, I wanted to try writing CUDA kernels in Rust using cuda-oxide. I was also curious whether GPUFlight's gpufl trace command could profile them. It runs your CUDA application and collects profiling data that you can inspect in the GPUFlight UI. It worked after I fixed a few issues, so I added Rust support to Performance Lab. You can now solve the lab problems in Rust and profile your kernels. GPUFlight’s Workbench also supports Rust if you want to experiment with your own code. I've marked it as Beta because I'm still working on benchmark calibration, but it's ready to try. Try Performance Lab If you'd like to learn more about "gpufl trace", check out this guide to profiling a CUDA application: Profile a CUDA app with gpufl trace i.redd.it
CUDA Rust: NVIDIA Compila Rust Directo a PTX — Dos Tracks para Kernels GPU NVIDIA acaba de anunciar CUDA Rust: escribir kernels GPU directamente en Rust, compilado nativo a PTX. Dos proyectos open source, Apache 2.0, cubriendo los dos modelos de programación de CUDA. #CUDA #Rust #GPU #NVIDIA #CUDA-Rust #PTX #B200 #HPC #OpenSource ⏱️ CAPÍTULOS 00:00 CUDA en Rust — Rust real compilado a PTX 00:17 Dos tracks: SIMT y Tile 00:34 CUDA-Oxide — Intercepta la compilación MIR→PTX 00:52 Memoria segura en GPU con DisjointSlice 01:11 Cutile-rs — Rust estable, tile-based 01:30 Ownership antes que unsafe 01:43 2 PFLOP/s en B200 — Benchmarks contundentes 02:04 Rust seguro vs unsafe: menos del 0.3% 02:21 Qwen3 en Rust con Grout — HuggingFace usa cutile 02:41 Requisitos: Linux, GPU CC 8.0+, CUDA 03:03 Rust en infraestructura NVIDIA 03:26 Roadmap hasta 2027 y más allá 03:43 Fearless Concurrency on the GPU 03:59 Los dos proyectos: CUDA-Oxide y Cutile-rs 📊 DATOS CLAVE • CUDA-Oxide: MIR → Pliron IR → LLVM IR → PTX nativo, early alpha • Cutile-rs: Tile-based, JIT a CUDA Tile IR, Rust estable 1.89+, ya en crates.io • Benchmarks B200: 7 TB/s element-wise, 2.07 PFLOP/s GEMM (96.4% cuBLAS) • Rust seguro vs unsafe: diferencia menor al 0.3%, sin overhead medible • HuggingFace usa cutile-rs en Grout: Qwen3-4B a 171 tok/s en RTX 5090 • Apache 2.0, open source, repositorios en NVlabs en GitHub 🔗 FUENTES 📄 NVIDIA Developer Blog: Introducing CUDA Rust — developer.nvidia.com 📄 CUDA-Oxide en GitHub — NVlabs/cuda-oxide 📄 Cutile-rs en GitHub — nvlabs/cutile-rs 📄 Cutile-rs Docs — nvlabs.github.io 📄 CUDA-Oxide Book — nvlabs.github.io 👍 Si te interesa el futuro de los kernels GPU en Rust, dale like, suscríbete y activa la campanita para no perderte nada. 🌐 nramos.dev · 💼 linkedin.com/in/nramosdev · 🐙 github.com/nicolasramos #CUDA #Rust #NVIDIA #GPU #PTX #MachineLearning #HPC #OpenSource #CUDA-Rust #B200 #Quantum #AI
정말 인연이 없을것 같은 곳에 피알이 들어가고 있다. cuda-rust#1346
NVIDIA/cuda-rust 从 nvlabs里出来了，合并了之前的 cuda oxide和 cutile rs