Rust SIMD on the GPU

At VectorWare, engineers have successfully implemented Rust's portable SIMD (core::simd) within GPU hardware. This accomplishment signals a substantial stride toward their objective of empowering developers to craft sophisticated, high-performance applications utilizing the full potential of GPU capabilities through familiar Rust abstractions.

Previously, VectorWare had integrated Rust threads into the GPU by mapping each std::thread to a GPU warp, thereby facilitating numerous concurrent threads on the GPU. However, this approach did not capitalize on the parallel lanes within each thread/warp. In contrast, SIMD offers a crucial abstraction for parallelism within a single thread, enabling a single instruction to operate on multiple data elements simultaneously.

Traditionally, implementing SIMD in Rust required utilizing architecture-specific vendor intrinsics like _mm256_add_ps on x86-64 or vaddq_f32 on Arm, leading to the need for separate implementations for multiple architectures. Rust's portable SIMD introduced an abstraction layer above these intrinsics, providing a uniform type Simd T, N that represents a vector of N elements of type T. By writing arithmetic, comparisons, reductions, and lane shuffles once using Simd, the compiler translates them into the corresponding vector instructions of the target CPU.

Recognizing that the GPU is merely another piece of vector hardware for portable SIMD, VectorWare found a way to leverage this technology on the GPU. The SIMD model used by NVIDIA, named SIMT (Single Instruction, Multiple Thread), issues a single instruction with each of its 32 lanes executing it on individual data. This concept is fundamentally the same as SIMD, where a single instruction operates on multiple data elements.

Mapping Rust's Simd onto the GPU's warp unit is straightforward. For instance, a Simd i16, 32 gives each of the warp's 32 lanes an i16 element, and adding two such vectors compiles into a single warp instruction where every lane performs the addition concurrently. This development completes the parallelism hierarchy established in VectorWare's earlier work.

On the CPU, a thread contains SIMD lanes, and on the GPU, a std::thread warp acts as a warp whose hardware lanes serve the same purpose. Core::simd is responsible for driving these lanes in both scenarios. The simplicity of the Rust code, which employs core::simd types, is identical when compiled to a GPU kernel, with no modifications to the source code.

A compact SIMD routine is defined and invoked from main, demonstrating the core features: elementwise arithmetic, comparisons yielding lane masks, selections driven by those masks, and horizontal reductions across lanes. The program's output is printed to the device using standard support. This GPU implementation is validated by a recording showing the exact same output as when executed on the CPU.

The mapping process hinges on the observation that a warp functions as a vector unit with individually addressable lanes. Once Simd T, N is organized per lane, each family of operations corresponds to a direct warp-level counterpart. SIMD elementwise operations, such as addition, multiplication, comparisons, and other lane-wise operators, are natively executed by the GPU.

SIMD reductions, like reduce_sum and reduce_max, aggregate values from all lanes into a scalar, utilizing warp shuffle instructions to exchange and combine values across lanes, resulting in the same scalar output in each lane. SIMD cross-lane shuffles, such as simd_swizzle! and rotates, move elements between lanes, aligning with the GPU's warp shuffle primitives for efficient data exchange.

Similarly, SIMD masks, which provide a predicate to each SIMD lane, map seamlessly onto GPU vote and ballot instructions, enabling operations like any and all. Rust's types differentiate between scalar and varying values. For instance, a plain f32 is uniform across all lanes, while Simd f32, 32 is varying, with each lane holding a distinct value.

The abstraction and hardware do not perfectly align in terms of lane count. While Rust's Simd T, N allows any N value from 1 to 64 on the CPU, the GPU hardware has a fixed lane width: 32 on NVIDIA and 32 or 64 on AMD. This mapping is one-to-one only when N matches the warp width. When N is smaller than the warp width, some lanes remain idle; if N is larger, certain lanes process multiple elements.

In cases where the warp is underutilized, a mechanism is needed to assign work to lanes following specific rules. At VectorWare, an intermediate representation (IR) is devised for this purpose, encoded within Rust's type system using types, generics, const generics, and trait bounds. The program is composed of typed operations, including ballots, shuffles, reductions, and scans, which are executed within the constraints of the warp's architecture.