Ranked #2 on Hacker News with 97 points and 48 comments.
GPU code can now use Rust's portable SIMD. We share the implementation approach and what this unlocks for GPU programming.
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.
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.
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.