Rust SIMD on the GPU

At VectorWare, we are building the first GPU-native software company. Today, we are excited to announce that we can successfully use Rust's portable SIMD (core::simd) on the GPU. This milestone marks a significant step towards our vision of enabling developers to write complex, high-performance applications that leverage the full power of GPU hardware using familiar Rust abstractions.

Parallelism below the thread

When we brought Rust threads to the GPU, we mapped each std::thread to a GPU warp. This let us run many concurrent threads on the GPU but did not use the parallel lanes within each thread/warp.

On the CPU, the abstraction for parallelism within a thread is SIMD. A single instruction operates on several data elements packed into a vector unit: where scalar code adds two numbers, a SIMD add takes two vectors of, say, eight f32 values and produces eight sums at once. This data parallelism is inside a single thread, below the level where the operating system schedules anything.

Rust's portable SIMD

Historically, writing SIMD in Rust meant reaching for the architecture-specific vendor intrinsics in core::arch, such as _mm256_add_ps on x86-64 or vaddq_f32 on Arm. These intrinsics are specific to a single instruction set, so a program that runs on more than one architecture needs a separate implementation for each.

Rust's portable SIMD instead adds a layer of abstraction above these intrinsics. It provides a single generic type Simd<T, N> that represents a vector of N elements of type T. A program writes its arithmetic, comparisons, reductions, and lane shuffles once against Simd and the compiler lowers them to whatever vector instructions the target CPU has.

At VectorWare, we realized the GPU is just one more piece of vector hardware for portable SIMD to target. As a bonus, portable SIMD lives in core rather than std and it does not even need the std support we brought to the GPU.

SIMT is SIMD

GPUs execute in a model NVIDIA calls SIMT, or Single Instruction, Multiple Thread. A warp issues one instruction, and each of its 32 lanes runs that instruction on its own data. One instruction operating on many data elements is exactly what SIMD means, and the per-lane addressing that SIMT adds does not change that. A warp is a wide vector unit and a portable SIMD vector maps onto that unit directly.

For example, a Simd<i16, 32> gives one i16 element to each of the warp's 32 lanes, and adding two such vectors compiles to a single warp instruction in which every lane adds its element at once.

This new mapping completes the parallelism hierarchy from our earlier work. On the CPU, a thread contains SIMD lanes, and on the GPU our std::thread is a warp whose hardware lanes play the same role. In both cases, core::simd drives those lanes.

A world first: core::simd on the GPU

As with our earlier posts, this is hard to show visually because the code is ordinary Rust. The same core::simd types that lower to x86-64 SIMD on a laptop lower to warp operations on the GPU, with no change to the source.

Here we define a small portable SIMD routine and call it from main. It exercises the core features of the model: elementwise arithmetic, a comparison that produces a lane mask, a select driven by that mask, and a horizontal reduction across lanes.

#![feature(portable_simd)]
 
use core::simd::cmp::SimdPartialOrd;
use core::simd::num::SimdFloat;
use core::simd::{Select, Simd};
 
// Portable SIMD. This exact function also compiles and runs on the CPU,
// where it lowers to x86-64, Arm, or scalar code depending on the target.
fn relu_dot(a: Simd<f32, 32>, b: Simd<f32, 32>) -> f32 {
    // Elementwise multiply: 32 products computed at once.
    let products = a * b;
 
    // Per-lane comparison produces a mask, one boolean per lane.
    let positive = products.simd_gt(Simd::splat(0.0));
 
    // Keep the positive products, replace the rest with zero.
    let clamped = positive.select(products, Simd::splat(0.0));
 
    // Horizontal add across all lanes down to a single scalar.
    clamped.reduce_sum()
}
 
fn main() {
    // Two 32-wide vectors, built with ordinary Rust.
    let a = Simd::<f32, 32>::splat(2.0);
    let b = Simd::<f32, 32>::from_array(std::array::from_fn(|i| i as f32 - 16.0));
 
    // Elementwise ops, a comparison mask, a select, and a reduction:
    // all ordinary portable SIMD, all running on the GPU.
    let result = relu_dot(a, b);
 
    // Printed from the GPU using our std support.
    println!("relu_dot = {result}");
}

The entry point is a normal fn main with no GPU-specific annotations. Our toolchain compiles it to a GPU kernel, and the result is printed from the device using our std support.

Below is a recording of the program running on the GPU, producing the exact same output as running it on the CPU.

Implementation

As previously mentioned, the mapping rests on a single observation: a warp is a vector unit whose lanes are individually addressable. Once Simd<T, N> is laid out per lane, each family of operations has a direct warp-level counterpart.

SIMD elementwise operations are the easy case. Addition, multiplication, comparison, and the other lane-wise operators come from ordinary Rust trait implementations on Simd such as Add. The GPU runs them natively.

SIMD reductions such as reduce_sum and reduce_max combine every lane into a scalar. These use the GPU's warp shuffle instructions to exchange and combine values across lanes, producing the same scalar result in every lane.

SIMD cross-lane shuffles, such as simd_swizzle! and rotates, move elements between lanes. Because a SIMD lane is a GPU warp lane, these map onto the same warp shuffle primitives that make GPU lanes so good at exchanging data.

SIMD masks map just as cleanly. A Mask<T, N> gives one predicate to each SIMD lane. Mask::select performs a selection in every warp lane. Horizontal mask queries such as any and all use GPU vote and ballot instructions.

Scalar values in the surrounding code, such as a loop counter or a constant, are computed identically by every lane and so are simply replicated across the warp just like in ordinary CUDA. This is the same uniform-versus-varying distinction that data-parallel languages like ISPC make explicit, except here it falls out of Rust's own types: a plain f32 is uniform, a Simd<f32, 32> is varying.

Working with lanes

The one place the abstraction and the hardware do not line up is lane count. On the CPU a Simd<T, N> allows any N from 1 through 64, but GPU hardware has a fixed width: 32 lanes on NVIDIA and 32 or 64 on AMD. The mapping is one to one only when N matches that width. A smaller N leaves some lanes idle while a larger N gives some or all lanes more than one element to process.

When there is more work than the warp is wide, we need a way to say which lanes do what. It helps to think of the warp as a small "machine" of its own: a fixed set of primitives for moving and combining data across lanes, plus invariants about which lanes are active and how much data each one holds. "Programming" it means placing work onto lanes within those rules.

At VectorWare, we give that machine an IR. Rather than a standalone data structure, we encode it in Rust's type system using types, generics, const generics, and trait bounds. A program is composed of typed operations: ballots, shuffles, reductions, scans, gathers, scatters, atomics, and strip mining for vectors wider than the warp. Operands, execution shape, and capacity are typed too. Because the operations carry their shape in the types, many invalid programs cannot be constructed at all.

The IR needs no interpreter on the GPU. Each operation lowers straight to the corresponding instructions with zero cost over hand-written PTX. The same types let us run it on the CPU too. We built a reference interpreter that executes the IR deterministically, a kind of Miri for warp-lane programming. We use it to simulate GPU code and for differential testing.

Our work targets NVIDIA today, but nothing here is CUDA specific. AMD wavefronts and Vulkan subgroups expose similar primitives and semantics. The IR itself is architecture-agnostic Rust.

Benefits

The same source runs on the CPU and the GPU. Code and libraries that already use portable SIMD become candidates for GPU execution without a rewrite.

Unmodified CPU code can use GPU lane-level parallelism. GPU-aware code can still go further by using core::arch intrinsics that map directly to PTX.

A Simd<T, N> is an ordinary owned value. The borrow checker, lifetimes, and type checking apply to it exactly as they do on the CPU. We are not adding a GPU-specific vector type or a new set of annotations. We are mapping Rust's existing portable SIMD onto the GPU's native execution model. At VectorWare, we are making GPUs behave like a normal Rust platform.

Downsides

Portable SIMD is still unstable in Rust. It requires the nightly #![feature(portable_simd)], and its surface may change before it stabilizes.

Vectors narrower than the warp leave lanes idle, and vectors wider than the warp turn each operation into more instructions. The abstraction is only zero cost when the vector width matches the number of warp lanes.

Not every cross-lane operation maps to an efficient warp instruction. Shuffles that match the hardware's supported patterns are cheap, but arbitrary permutations may need several instructions or a trip through shared memory. Horizontal operations like reductions and all/any also act as synchronization points within the warp, which constrains how freely the scheduler can overlap work.

We had to change the compiler to make the abstraction sound when interacting with other Rust features. As this is uncharted territory, we are not yet confident we have covered every case.

Future work

With SIMD, threads, and async all mapped onto the GPU, the natural next step is composing them: threads spreading work across warps, core::simd spreading data across the lanes within each warp, and async structuring the concurrency between them.

We are also interested in lowering matrix-shaped SIMD onto the GPU's tensor cores, and in auto-vectorizing ordinary scalar Rust loops into Simd operations so that code gets warp-level parallelism without being written against core::simd at all. As members of the Rust compiler team, we are keen to explore how much of this can happen in the compiler itself.

A vector representation shared across the CPU and the GPU is valuable, though it is not clear that today's portable SIMD types are the right basis for one. For one thing, they largely sit in a world of their own within the core and std APIs. More exploration is necessary.

Is VectorWare only focused on Rust?

The speed at which we are able to make progress on the GPU is a testament to the power of Rust's abstractions and ecosystem.

As a company, we understand that not everyone uses Rust. Our future products will support multiple programming languages and runtimes. However, we believe Rust is uniquely well suited to building high-performance, reliable GPU-native applications and that is what we are most excited about.

Follow along

Follow us on X, Bluesky, LinkedIn, or subscribe to our blog to stay updated on our progress. We will be sharing more about our work in the coming months. You can also reach us at hello@vectorware.com.