Rust SIMD on the GPU

Dispatches

At VectorWare, we are building the first GPU-native package company. Today, we are excited to announce that we tin successfully usage Rust's portable SIMD (core::simd) connected the GPU. This milestone marks a important measurement towards our imagination of enabling developers to write complex, high-performance applications that leverage the afloat powerfulness of GPU hardware using acquainted Rust abstractions.

Parallelism beneath the thread

When we brought Rust threads to the GPU, we mapped each std::thread to a GPU warp. This fto america tally galore concurrent threads connected the GPU but did not usage the parallel lanes within each thread/warp.

On the CPU, the abstraction for parallelism wrong a thread is SIMD. A azygous instruction operates connected respective information elements packed into a vector unit: wherever scalar codification adds two numbers, a SIMD adhd takes 2 vectors of, say, 8 f32 values and produces eight sums astatine once. This information parallelism is inside a azygous thread, beneath the level wherever the operating strategy schedules anything.

CPU threadSIMD op012N⋯SIMD lanesCPU thread

Rust's portable SIMD

Historically, penning SIMD successful Rust meant reaching for the architecture-specific vendor intrinsics successful core::arch, specified as _mm256_add_ps on x86-64 aliases vaddq_f32 on Arm. These intrinsics are circumstantial to a azygous instruction set, truthful a program that runs connected much than 1 architecture needs a abstracted implementation for each.

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

At VectorWare, we realized the GPU is conscionable 1 much portion of vector hardware for portable SIMD to target. As a bonus, portable SIMD lives successful halfway alternatively than std and it does not moreover request the std support we brought to the GPU.

SIMT is SIMD

GPUs execute successful a exemplary NVIDIA calls SIMT, aliases Single Instruction, Multiple Thread. A warp issues 1 instruction, and each of its 32 lanes runs that instruction connected its ain data. One instruction operating connected galore information elements is exactly what SIMD means, and the per-lane addressing that SIMT adds does not change that. A warp is simply a wide vector portion and a portable SIMD vector maps onto that portion directly.

CPU thread012N⋯SIMD lanes≈GPU warp012N⋯warp lanes

For example, a Simd<i16, 32> gives one i16 constituent to each of the warp's 32 lanes, and adding 2 specified vectors compiles to a azygous warp instruction successful which each lane adds its constituent astatine once.

This caller mapping completes the parallelism level from our earlier work. On the CPU, a thread contains SIMD lanes, and connected the GPU our std::thread is a warp whose hardware lanes play the aforesaid role. In some cases, core::simd drives those lanes.

CPU⋯thread 0012N⋯thread 1012N⋯thread N012N⋯SIMD lanes≈GPU⋯warp 0012N⋯warp 1012N⋯warp N012N⋯warp lanes

A world first: core::simd connected the GPU

As pinch our earlier posts, this is difficult to show visually because the codification is ordinary Rust. The aforesaid core::simd types that little to x86-64 SIMD connected a laptop little to warp operations connected the GPU, pinch nary alteration to the source.

Here we specify a mini portable SIMD regular and telephone it from main. It exercises the halfway features of the model: elementwise arithmetic, a comparison that produces a lane mask, a prime driven by that mask, and a horizontal simplification crossed lanes.

The introduction constituent is simply a normal fn main pinch nary GPU-specific annotations. Our toolchain compiles it to a GPU kernel, and the consequence is printed from the instrumentality utilizing our std support.

Below is simply a signaling of the programme moving connected the GPU, producing the nonstop aforesaid output as running it connected the CPU.

Implementation

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

SIMD elementwise operations are the easy case. Addition, multiplication, comparison, and the different lane-wise operators travel from mean Rust trait implementations connected Simd specified as Add. The GPU runs them natively.

SIMD reductions specified as reduce_sum and reduce_max combine each lane into a scalar. These usage the GPU's warp shuffle instructions to exchange and harvester values crossed lanes, producing the aforesaid scalar consequence successful each lane.

SIMD cross-lane shuffles, specified as simd_swizzle! and rotates, move elements betwixt lanes. Because a SIMD lane is simply a GPU warp lane, these representation onto the same warp shuffle primitives that make GPU lanes truthful bully astatine exchanging data.

SIMD masks representation conscionable arsenic cleanly. A Mask<T, N> gives 1 predicate to each SIMD lane. Mask::select performs a action successful each warp lane. Horizontal disguise queries specified as any and all usage GPU vote and ballot instructions.

Scalar values successful the surrounding code, specified arsenic a loop antagonistic aliases a constant, are computed identically by each lane and truthful are simply replicated crossed the warp conscionable for illustration successful ordinary CUDA. This is the aforesaid uniform-versus-varying favoritism that data-parallel languages like ISPC make explicit, isolated from present it falls retired of Rust's ain types: a plain f32 is uniform, a Simd<f32, 32> is varying.

Working pinch lanes

The 1 spot the abstraction and the hardware do not statement up is lane count. On the CPU a Simd<T, N> allows immoderate N from 1 done 64, but GPU hardware has a fixed width: 32 lanes connected NVIDIA and 32 aliases 64 connected AMD. The mapping is 1 to 1 only erstwhile N matches that width. A smaller N leaves immoderate lanes idle while a larger N gives immoderate or all lanes much than 1 constituent to process.

When location is much activity than the warp is wide, we request a measurement to opportunity which lanes do what. It helps to deliberation of the warp arsenic a mini "machine" of its own: a fixed group of primitives for moving and combining information crossed lanes, positive invariants about which lanes are progressive and really overmuch information each 1 holds. "Programming" it intends placing work onto lanes wrong those rules.

At VectorWare, we springiness that instrumentality an IR. Rather than a standalone information structure, we encode it successful Rust's type strategy utilizing types, generics, const generics, and trait bounds. A programme 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 transportation their style successful the types, galore invalid programs cannot beryllium constructed astatine all.

The IR needs nary expert connected the GPU. Each cognition lowers consecutive to the corresponding instructions pinch zero costs complete hand-written PTX. The aforesaid types fto america tally it connected the CPU too. We built a reference expert that executes the IR deterministically, a benignant of Miri for warp-lane programming. We usage it to simulate GPU codification and for differential testing.

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

Benefits

The aforesaid root runs connected the CPU and the GPU. Code and libraries that already usage portable SIMD go candidates for GPU execution without a rewrite.

Unmodified CPU codification tin usage GPU lane-level parallelism. GPU-aware codification tin still spell further by using core::arch intrinsics that representation straight to PTX.

A Simd<T, N> is an mean owned value. The get checker, lifetimes, and type checking use to it precisely arsenic they do connected the CPU. We are not adding a GPU-specific vector type aliases a caller group of annotations. We are mapping Rust's existing portable SIMD onto the GPU's autochthonal execution model. At VectorWare, we are making GPUs behave for illustration a normal Rust platform.

Downsides

Portable SIMD is still unstable successful Rust. It requires the nightly #![feature(portable_simd)], and its aboveground whitethorn alteration earlier it stabilizes.

Vectors narrower than the warp time off lanes idle, and vectors wider than the warp turn each cognition into much instructions. The abstraction is only zero costs erstwhile the vector width matches the number of warp lanes.

Not each cross-lane cognition maps to an businesslike warp instruction. Shuffles that match the hardware's supported patterns are cheap, but arbitrary permutations whitethorn request several instructions aliases a travel done shared memory. Horizontal operations for illustration reductions and all/any besides enactment arsenic synchronization points wrong the warp, which constrains how freely the scheduler tin overlap work.

We had to alteration the compiler to make the abstraction sound erstwhile interacting pinch other Rust features. As this is uncharted territory, we are not yet assured we person covered every case.

Future work

With SIMD, threads, and async all mapped onto the GPU, the earthy adjacent measurement is composing them: threads spreading activity crossed warps, core::simd spreading information crossed the lanes within each warp, and async structuring the concurrency betwixt them.

We are besides willing successful lowering matrix-shaped SIMD onto the GPU's tensor cores, and successful auto-vectorizing ordinary scalar Rust loops into Simd operations truthful that codification gets warp-level parallelism without being written against core::simd astatine all. As members of the Rust compiler team, we are keen to research really overmuch of this tin hap successful the compiler itself.

A vector practice shared crossed the CPU and the GPU is valuable, though it is not clear that today's portable SIMD types are the correct ground for one. For 1 thing, they mostly beryllium successful a world of their ain wrong the halfway and std APIs. More exploration is necessary.

Is VectorWare only focused connected Rust?

The velocity astatine which we are capable to make advancement connected the GPU is simply a testament to the powerfulness of Rust's abstractions and ecosystem.

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

Follow along

Follow america connected X, Bluesky, LinkedIn, aliases subscribe to our blog to enactment updated connected our progress. We will beryllium sharing much astir our activity in the coming months. You tin besides scope america astatine [email protected].