GoKawiilAt 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.
CPU thread SIMD op 0 1 2 N ⋯ SIMD lanes CPU thread
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
... continue reading