For decades, GPU programming has meant one of two things: writing CUDA kernels in C++ or wrestling with OpenCL. Both require you to think in a fundamentally different paradigm than CPU programming. But VectorWare just changed that by making Rust's portable SIMD — core::simd
— work natively on the GPU.
If that sounds like a niche technical achievement, it isn't. It's the first crack in a wall that has separated CPU and GPU programming for twenty years.
Modern processors offer two levels of parallelism:
Thread-level parallelism is what most developers know. You spawn threads, they run concurrently, the OS schedules them. This works the same on CPU and GPU — VectorWare already demonstrated Rust threads running on GPUs earlier this year.
SIMD (Single Instruction, Multiple Data) is the harder level. A single instruction operates on multiple data elements simultaneously — a vector of 8 floats added to another vector of 8 floats in one clock cycle. On CPUs, this is how you get peak performance from numerical code. On GPUs, the equivalent is called SIMT (Single Instruction, Multiple Thread), where a "warp" of 32 lanes executes one instruction, each on its own data.
The problem: writing SIMD code has historically meant choosing a specific CPU architecture. x86 has AVX (via _mm256_add_ps
). ARM has NEON (via vaddq_f32
). You write different code for each.
Rust's portable SIMD (core::simd
) solves this on the CPU side — you write Simd<T, N>
once and the compiler lowers it to whatever vector instructions the target has. But until now, it didn't work on GPUs.
VectorWare realized something elegant: a GPU warp is a wide vector unit. A Simd<i16, 32>
maps perfectly onto a 32-lane warp. Adding two such vectors compiles to a single warp instruction where every lane adds its element simultaneously.
This means the same Rust SIMD code that runs on an x86 CPU with AVX now also runs on an NVIDIA GPU — without changes. The abstraction layers correctly.
let a: Simd<i16, 32> = Simd::from_slice(&data_a);
let b: Simd<i16, 32> = Simd::from_slice(&data_b);
let result = a + b; // One warp instruction on GPU, one AVX instruction on x86
No CUDA kernels. No OpenCL boilerplate. No separate codebase for GPU and CPU. The same Rust code compiles and runs on both.
For developers: GPU programming just got more accessible. You don't need to learn CUDA or OpenCL. You write Rust SIMD code you already know, and it runs on the GPU. The learning curve drops from "learn a new paradigm" to "learn one new type."
For portability: Code written against core::simd
now runs across x86, ARM, and GPU. Three targets, one codebase. This is unprecedented — until now, GPU code was always a separate, platform-specific artifact.
For Rust: This validates Rust's approach to portable abstractions. The same language that gives you memory safety also gives you portable SIMD that spans CPU and GPU. No other language offers this.
For performance: GPUs have massive parallelism that most applications can't tap because the programming model is too different. Portable SIMD on GPU removes that barrier. Any Rust program using SIMD for numerical workloads can now benefit from GPU acceleration without rewriting.
The key realization is that SIMT (GPU's model) is actually SIMD in disguise. NVIDIA calls it SIMT because each lane can diverge — branch independently — which pure SIMD doesn't allow. But when lanes don't diverge, a warp is exactly a SIMD vector. And Simd<T, N>
in Rust gives you explicit control over lane operations, so you can write code that stays in the fast non-divergent path.
VectorWare's implementation maps each Simd<T, N>
to a warp, with N matching the warp width (32 on NVIDIA, 64 on AMD). Operations on these vectors compile directly to warp instructions. Lane shuffles, reductions, and comparisons all map cleanly.
The elegance is that this isn't an emulation layer — it's a direct mapping. The Rust compiler already knows how to lower Simd
operations to vector instructions. VectorWare taught it that a GPU warp is just another vector target.
The immediate applications are in numerical computing: linear algebra, image processing, signal processing, simulations. Any Rust code already using portable SIMD for CPU acceleration can now run on GPU with minimal changes.
The longer-term implications are bigger. If core::simd
works on GPU, then any Rust crate built on top of it — numerical libraries, ML frameworks, game engines — gains GPU support for free. The ecosystem leverage is enormous.
And it points toward a future where the CPU/GPU divide is less of a wall and more of a gradient. You write code once, and the compiler decides where it runs best. Not through some magic auto-parallelization, but through a clean abstraction that works across both.
This is early. VectorWare is a startup building "the first GPU-native software company," so they have an interest in promoting this capability. The implementation requires their GPU runtime — it's not something you can use with stock rustc
today. Performance data is limited. And warp-level programming still requires understanding GPU memory hierarchies and execution models for best results.
But the direction is right. The fact that it works at all — that Rust's portable SIMD can target a GPU warp as easily as an x86 AVX unit — is a proof of concept that the programming language community has been waiting for.
For twenty years, GPU programming has required leaving your comfortable CPU programming model behind and learning an entirely new paradigm. VectorWare just showed that it doesn't have to be that way. Rust's portable SIMD works on GPUs because a GPU warp is, at its core, a wide vector unit — and Simd<T, N>
is the right abstraction for it.
This won't replace CUDA for maximum-performance GPU code any time soon. But for the 90% of applications that need GPU acceleration without dedicating a team to GPU programming, portable SIMD in Rust just became the easiest path forward. And that's a bigger deal than it sounds.
Based on VectorWare's announcement of Rust SIMD on GPU support. Read the original article.