How Does Rust Code Run on a GPU? Tracing cuda-oxide and PTX

A developer traced how a Rust vector-addition kernel compiled with NVIDIA's cuda-oxide toolchain becomes PTX, examining how thread indexing and array accesses are lowered for the GPU. The kernel uses thread::index_1d() and DisjointSlice to assign one output element per thread with bounds checks enabled, built against cuda-rust commit a0cc6cc0 on an RTX 5060 Ti targeting sm_120a. CUDA Rust is described as an early-stage research project with APIs still under development.

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. 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 https://github.com/NVIDIA/cuda-rust . 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 https://nvidia.github.io/cuda-rust/cuda-oxide/ . 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 https://github.com/NVIDIA/cuda-rust . As someone who writes Rust, I find this a very intriguing project. 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 https://github.com/NVIDIA/cuda-rust . 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. 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 https://docs.nvidia.com/cuda/parallel-thread-execution/index.html introduction . 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. 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 https://nvidia.github.io/cuda-rust/cuda-oxide/compiler/architecture-overview.html . 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. 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. 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 https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/writing-cuda-kernels.html 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. 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. 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. 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 https://docs.nvidia.com/cuda/parallel-thread-execution/index.html miscellaneous-instructions-trap . 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. 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 https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html just-in-time-compilation . The overall flow is: php 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.