⏱️ Lectura: 10 min
VectorWare, a startup dedicated to building native software for GPUs, got SIMD on GPU to run with exactly the same Rust code already used on CPU. The company announced this on its official blog, showing a function that compiles to a vector instruction on x86-64 on the laptop and to a warp instruction on the graphics card, without changing a single line.
📑 En este artículo
The achievement matters because until now, programming SIMD on GPU meant learning a different shader language or writing separate CUDA kernels. With this announcement, Rust’s Simd<T, N> abstraction is no longer exclusive to the CPU.
TL;DR
- VectorWare managed to compile Rust’s portable SIMD (core::simd) so it also runs on the GPU.
- The same function, unchanged, compiles to a CPU instruction and to a GPU warp instruction.
- A 32-element i16 Simd fills all 32 lanes of a warp: vpaddw on CPU, add.s16 on GPU (PTX).
- The announcement builds on VectorWare’s earlier work, which already mapped std::thread to a GPU warp.
- core::simd lives in core, not std, and doesn’t depend on the GPU support for std the company built earlier.
- NVIDIA calls its model SIMT, but VectorWare argues a warp is, in essence, a SIMD unit.
- The API is still gated behind the nightly portable_simd feature, with no public stabilization date in Rust.
What happened
VectorWare published a technical post demonstrating that Rust’s generic type Simd<T, N>, originally designed to vectorize CPU code, also compiles to warp instructions on the GPU. The example they show defines a function that multiplies two 32-element vectors, applies a mask via a comparison, and reduces the result to a scalar. That same function, without touching a line, runs on the author’s CPU and on an NVIDIA GPU.
The company had already solved a similar problem before: mapping Rust’s std::thread to a GPU warp, so that each logical thread in the program occupies a physical warp. That earlier work, according to VectorWare’s blog, is the foundation on which they built GPU SIMD support.
Context and history
For years, writing SIMD in Rust meant using architecture-specific intrinsics in core::arch: functions like _mm256_add_ps on x86-64 or vaddq_f32 on ARM. Each architecture had its own function name, so a cross-platform program needed a separate implementation for each one.
Rust’s portable-simd project solved that problem by adding an abstraction layer: a generic type Simd<T, N> representing a vector of N elements of type T. The programmer writes arithmetic, comparisons, and reductions once against Simd, and the compiler decides which vector instructions to use based on the compilation target.
VectorWare discovered that the GPU is, quite simply, another piece of vector hardware that core::simd can target. As an added benefit, core::simd lives in core rather than std, so it doesn’t even depend on the GPU support for std the company built in its earlier work.
Technical details: how SIMD on GPU works
NVIDIA calls its execution model SIMT, or Single Instruction, Multiple Thread. A warp issues a single instruction and each of its 32 lanes executes it on its own data. That’s essentially SIMD: one instruction operating on many elements at once. The per-lane addressing SIMT adds doesn’t change that underlying nature. A warp is a wide vector unit, and a portable Simd vector maps directly onto that unit.
A Simd<i16, 32> gives one i16 element to each of the warp’s 32 lanes. Adding two vectors of that type compiles, on CPU, to a single AVX-512 vpaddw instruction. On GPU, the same addition translates to add.s16 in PTX assembly. The Rust source code is identical in both cases.
#![feature(portable_simd)]
use core::simd::Simd;
fn sumar_vectores(a: Simd<f32, 8>, b: Simd<f32, 8>) -> Simd<f32, 8> {
a + b
}
fn main() {
let a = Simd::<f32, 8>::splat(3.0);
let b = Simd::from_array([1.0, 2.0, 3.0, 4.0, 5.0, 6.0, 7.0, 8.0]);
println!("{:?}", sumar_vectores(a, b));
}
This function runs today on any x86-64 or ARM CPU with Rust nightly, and compiles to a single vector instruction. It’s the kind of primitive VectorWare now also compiles for GPU.
VectorWare’s real example combines several operations: element-wise multiplication, a comparison that generates a per-lane boolean mask, a select that filters based on that mask, and a horizontal reduction. It’s the pattern behind a ReLU-style activation in a neural network layer:
use core::simd::cmp::SimdPartialOrd;
use core::simd::num::SimdFloat;
use core::simd::{Select, Simd};
fn producto_relu(pesos: Simd<f32, 32>, entradas: Simd<f32, 32>) -> f32 {
let productos = pesos * entradas;
let son_positivos = productos.simd_gt(Simd::splat(0.0));
let filtrados = son_positivos.select(productos, Simd::splat(0.0));
filtrados.reduce_sum()
}
producto_relu multiplies 32 weights by 32 inputs in parallel, discards negative products, and sums what’s left into a single scalar. According to VectorWare, this function compiles and runs the same on CPU and GPU, without any special annotation.
💭 Key point: the entry point is still a normal fn main, with no kernel attributes or GPU annotations. VectorWare’s toolchain decides behind the scenes which parts of the program to compile as a kernel.
Here’s the full parallelism hierarchy, with core::simd handling the lowest level in both worlds:
flowchart TD
subgraph CPU["CPU"]
T["thread"] --> L0["lane 0"]
T --> L1["lane 1"]
T --> L2["lane 2"]
T --> LN["lane N"]
end
subgraph GPU["GPU"]
W["warp"] --> G0["lane 0"]
W --> G1["lane 1"]
W --> G2["lane 2"]
W --> GN["lane N"]
end
L0 -.-> G0
L1 -.-> G1
L2 -.-> G2
LN -.-> GN
How to start testing it
The frontend of this work, core::simd, is part of official Rust, and anyone can try it today on CPU without waiting for VectorWare’s toolchain. The steps are the same on Windows, macOS, and Linux since they depend on rustup, not the operating system:
# Windows, macOS, and Linux (same sequence with rustup)
rustup toolchain install nightly
rustup override set nightly
cargo new simd-demo
cd simd-demo
Then, add #![feature(portable_simd)] as the first line of src/main.rs, paste the sumar_vectores example from above, and run cargo run.
To confirm the compiler generated a vector instruction and not a scalar loop, install cargo-show-asm and inspect the generated assembly:
cargo install cargo-show-asm
cargo asm --release simd_demo::sumar_vectores
If a single instruction like vaddps (or addps without AVX) shows up in the output, vectorization worked. That same method, checking the generated assembly, is what VectorWare used to show that its compiler produces add.s16 in PTX on the GPU side.
⚠️ Heads up: the compiler that ports this code to the GPU is VectorWare’s own and isn’t part of official Rust. What you can reproduce today with rustup is only half the announcement: the CPU execution.
Impact and analysis
Rust already has several paths for GPU programming: rust-gpu compiles to SPIR-V for use with Vulkan, wgpu runs compute shaders written in WGSL, and crates like cudarc expose direct bindings to CUDA. What’s different about VectorWare’s approach is that there’s no separate shader language: the same Rust function with core::simd targets both backends.
| Option | When to use it | Advantage | Limitation |
|---|---|---|---|
| core::simd + VectorWare | Code that already uses portable SIMD and wants to also run on GPU | A single source for CPU and GPU | Proprietary toolchain, not integrated into official rustc |
| rust-gpu | Graphics or compute shaders inside a Vulkan pipeline | Generates standard SPIR-V | Requires thinking in terms of shaders, not portable SIMD |
| wgpu (WGSL) | Cross-platform apps already using wgpu for graphics | Runs on web, desktop, and mobile | Compute is written in a separate language, WGSL |
| cudarc | Fine-grained control over specific NVIDIA hardware | Direct access to the CUDA API | Tied to NVIDIA, no portability to other vendors |
A real limitation, not mentioned in the announcement but inherent to the model: the 32-lane warp size is specific to NVIDIA. A Simd<T, 32> designed for that number doesn’t fully use a 64-lane wavefront on another vendor’s hardware, so the portable vector’s size still ties the code to a specific architecture, even if the syntax is portable.
What’s next
core::simd is still gated behind a nightly feature, with no public stabilization date in Rust’s official repository. That means any project adopting it today, for CPU or GPU, accepts the risk that the API could change before reaching stable Rust. VectorWare didn’t specify in its post whether it plans to open-source its GPU compiler or keep it as a proprietary product, nor which architectures besides NVIDIA it intends to support.
📖 Summary on Telegram: See summary
Try it yourself: install Rust nightly, add #![feature(portable_simd)] to a new crate, and use cargo asm to compare the assembly core::simd generates on your own CPU.
Frequently Asked Questions
What is Rust’s portable SIMD (core::simd)?
It’s a generic type, Simd<T, N>, representing a vector of N elements of type T. It replaces architecture-specific intrinsics, like _mm256_add_ps on x86-64, with a single API that the compiler translates based on the target hardware.
What’s the difference between SIMD and SIMT?
SIMD executes one instruction on several pieces of data within a single thread. SIMT, the model NVIDIA uses in its GPUs, adds per-lane addressing within a warp, but according to VectorWare’s analysis it’s still SIMD underneath: a 32-lane warp is a wide vector unit.
Do I need to rewrite my code to run it on the GPU?
According to VectorWare’s announcement, no. The same function that uses Simd<T, N> compiles the same way for CPU or GPU; the only thing that changes is the compiler backend that processes it.
Is core::simd available in stable Rust?
No. It requires the nightly feature #![feature(portable_simd)] and has no public stabilization date in Rust’s official repository.
What alternatives exist today for GPU programming in Rust?
Among the most used are rust-gpu (compiles to SPIR-V), wgpu with WGSL shaders, and cudarc with direct CUDA bindings.
Is VectorWare’s GPU compiler open source?
VectorWare didn’t publish the code for its GPU toolchain alongside this announcement. Only core::simd, the frontend that does run on CPU today, is part of Rust’s official portable-simd project.
References
- VectorWare, “Rust SIMD on the GPU”: the original announcement with the code example and lane-to-lane mapping diagrams.
- rust-lang/portable-simd on GitHub: the official repository for the
core::simdimplementation. - NVIDIA, “Parallel Thread Execution ISA”: official documentation on the SIMT model and PTX assembly.
- Wikipedia, “Single instruction, multiple data”: historical context on the SIMD model.
📱 Like this content? Follow @programacion on Telegram for daily tech content in Spanish: quick summaries, fresh content every day. @programacion
Imagen destacada: Foto de Anne Nygård en Unsplash
0 Comments