Rust SIMD on the GPU

Oh really? I've never used any Jetbrain product so I don't know but if it's indeed the case even when you don't even use a nightly toolchain that sounds like a very bad design.

The major one is that there's one layer of indirection that exists on GPUs that doesn't really on CPUs. There's one giant vector register file per for each of these processing cores (that'll be something like 2048 rows x 32 lanes x 32 bits). An individual shader invocation might only need say, 16 rows. While there's hardware for issuing 4 hyperthreads at any given time, there can be a variable number of thread states in the register file. So for the case of each invocation only needing 16 rows, you might be able to fit 128 hyperthread states into processing core. Those four hyperthreads then hardware schedule those 128 states and will execute any that are ready, as well as allocate more from other scheduling hardware as gaps in the register file appear from shaders completing.

Because of this massive amount thread state, you don't depend nearly as much on a cache hierarchy to deal with DRAM latency. There's ostensibly some other thread state sitting around that can be serviced while others wait for the hundreds of cycles of latency to access DRAM.

So the whole model of how you account for the discrepancy between ALU cycle time, and DRAM latency changes versus a CPU. Where a modern CPU spends a lot of area on complex cache hierarchies, speculation, etc, to hide the latency to memory, a GPU focuses on having a lot of thread state around and a lot of ALUs, but balanced ideally, so there's always ALU work to do while other thread states are waiting on memory.

Now, over time, GPUs have gotten more complex hardware, and more complex cache hierarchies to cover the cases that aren't handled well by extremely long access times. But those tend to be very explicit. Additionally, CPU vector files have gotten more similar to GPU cores as architectural features like lane masking/predicates have been added to have the equivalent of CUDA threads in the same warp that take different paths through control flow blocks. That's a lot of what people mean when they say that AVX-512 adds a lot more than just 512-bit registers. The K mask registers let you do a lot of GPU shader tricks to have effective partial residency, and not have to use all the lanes if the data doesn't line up with that.

Current CPU cores do two AVX-512 operations per cycle. If you can saturate this you’ll often run out of memory bandwidth on CPUs because of lower bandwidth compared to GPUs. In principle, if you bought a 192-core processor you’d have 6,144 GPU-ish cores of 32-bit operations, and they would run at a significantly higher clock rate than a GPU. It would not be competitive with a GPU for the kinds of things GPUs are good at it but it wouldn’t be as far off as you might assume. For some types of code, AVX-512 is unambiguously better.

Horses for courses. GPUs and CPUs were optimized for different things but their capabilities have slowly been converging over time. They all work from the same transistor budgets, the differences are where the tradeoffs are made.

There is a pithy silicon architecture tradeoff trilemma to be made regarding CPUs, GPUs, and barrel processors.

Note that modern NVIDA GPUs like the 5090 are actually more like 170 SMs on a chip, or 21,760 flop/cycle, or ~20-40x more ops/cycle than your example CPU.

Put GDDR6 vs DDR5 memory on top of that, and it’s easy to see why the GPU can churn through math so fast … as long as it’s GPU-y. For stuff the GPU does well, the CPU typically can’t compete, the GPU is often more than 10x faster. But GPU-y tasks are a subset, and there are CPU-y things the GPU can’t compete on, despite (or even because of) the thread count discrepancy.

Others have taken a stab at the actual differences, but there is a deeper fundamental reason.

A CPU is optimized for low latency of operations. They are designed to complete a given piece of code as fast as possible. There are some affordances for throughput, such as SIMD, but even those are designed to only be as good as they can without compromising the low-latency design of the core.

And the reason this cannot compete with GPUs in throughput loads is that after a point, completing a single task 2x as fast costs a lot more than 2x the transistors and power. CPUs chase that curve as high as practical, GPUs stop once it no longer makes sense for throughput. This is not just clock speed (though it is also clock speed, modern GPUs hang around in the 2.5GHz area while CPUs are about twice that), but especially their ability to hide memory latency, and ILP. CPUs spend big on being able to issue, execute and retire multiple instructions from the same stream, with complex reordering and more than half a dozen execution units per thread, while GPUs are either scalar within a thread, or maybe dual issue. A CPU has a cache hierarchy optimized for bringing average memory latency down, while GPUs just juggle more threads and use them to get something to execute when waiting for memory.

On the GPU however, the hyperthreads are just a round-robin execution queue to take advantage of instruction pipelining. The register bank of a single GPU core is huge and can be flexibly divided across a variable number of thread contexts when a kernel is launched. Many thread contexts can be held in registers simultaneously in a single GPU core. That makes stalling on memory latency much less of a problem. The hardware can focus on delivering raw bandwidth with high latency and get great overall performance. This throughput-instead-of-latency trade-off extends to many other aspects of GPU design.

An add instruction on an AMD GPU adds two scalar values. If they're in a VGPR then each thread will add two values unique to that thread. A SIMD ISA as is common on a CPU is different because an add instruction explicitly adds a vector of values. xmm1 stores 128-bits of data. VGPR[1] stores 32-bits of data vectored over 32-64 threads in a thread group.

Without special instructions a thread can't access the VGPR values stored in other threads.

What you'd actually want is a matrix of variant implementations burned into the binary, with runtime (or process-boot-time) hardware detection that swaps symbols out to point to the correct variant.

https://github.com/seclorum/SIMDSynth

(uses SIMD for optimizing voice allocation and filter calculations on both ARM and x86, same codebase ..)

This is trivial (but not pretty!) to do with something like `#[cfg(target_feature = "avx2")] const SIMD_WIDTH: usize = 8`. You need a few lines of ugly cfg logic to configure this.

A somewhat orthogonal and much more difficult problem is how to select it at runtime. You would either need to have different binaries built with different compiler options, link object files built with different compiler options to same binary, or dynamically link the correct code at runtime.

This is actually one of the (IMO only) cases where intrinsics are more practical: you can use `_mm256_add_ps` from AVX2 intrinsics regardless of whether you've configured your compiler to support AVX2 or not. As long as you check at runtime before calling the code so you don't get illegal instruction exceptions.