Rendered at 22:23:38 GMT+0000 (Coordinated Universal Time) with Cloudflare Workers.
O3marchnative 1 days ago [-]
The author mentions Rust's portable SIMD library [0]. The only issue with portable SIMD is it's only available on nightly. I used it in my FFT crate, but we had to switch to the fearless_simd crate in order to get a portable SIMD solution that works on stable [1].
Pretty common for Rust to cook things in nightly for a very long time; I wouldn't consider it a bad thing, tbh.
jmalicki 5 hours ago [-]
For these sort of things, I wish they would have some semi-stable beta or prerelease tracks other than just nightly so you could use the new features on a somewhat stable branch. If something is in nightly that is too crazy for a lot of projects to really try and exercise it because so much is changing constantly. Like a monthly or quarterly stabilization would be amazing, that may still have experimental features not making it into stable, but has a period of bugfixing w/o intentionally breaking changes to settle down.
nextaccountic 8 hours ago [-]
It's a bad thing when there are no breaking changes done and they stabilize the exact same thing a few years later
zbentley 7 hours ago [-]
Eh, it's hard to prove absence, and time is a beneficial quantity here. The longer something sits on nightly, the greater the chance that bugs are identified before it reaches stable and its usage significantly increases.
A lot of bugs with SIMD libraries are in the domain of interactions, not functionality--e.g. SIMD malfunctions on rare chips, chips with previously-unseen combinations of hardware/userspace firmware/microcode behavior, compilers run in weird harnesses that lie about hardware capabilities, and so on. I assume that's the case with Rust's portable SIMD as well.
If your QA is unpredictable individual use-cases (as with most OSS projects), then there's no way to measure "testing complete" or "coverage"; letting it bake for awhile is the best approach available.
5 hours ago [-]
LoganDark 19 hours ago [-]
It's been annoying to me as an end user that so many basic things require nightly. I use nightly as my main toolchain, but enabling unstable features makes a project nightly-only, which is undesired for crates that don't already revolve around the unstable feature.
I most often encounter unstable features when I reach for a basic common-sense utility method and discover that it's not stable. Like just earlier today I would have reached for bool::toggle which not only is unstable, but is also newly added as of like a month ago! but some unstable methods have been sitting around for years.
And now that IntelliJ-Rust is proprietary, I can't even make a feature request anymore for the ability to exclude unstable features from the autocomplete. So they will taunt me forever, perfect little helpers just locked away.
stymaar 15 hours ago [-]
> It's been annoying to me as an end user that so many basic things require nightly
It used to be the case a decade ago, but now I wouldn't agree that any "basic" things require nightly (I wouldn't call portable SIMD "basic" at all for instance).
> Like just earlier today I would have reached for bool::toggle which not only is unstable, but is also newly added as of like a month ago!
This is very likely not the kind of feature that will stay on nightly for a long time, but is instead one of the many convenience feature that land on stable every release. The 6-weeks release cadence with beta in between means there's always at least 6 weeks and up to 3 months between the time a feature land on nightly and the day it reaches stable, even if the feature is as consensual as this one.
> And now that IntelliJ-Rust is proprietary, I can't even make a feature request anymore for the ability to exclude unstable features from the autocomplete.
Can't you tell it to use stable as the default target, and use nightly manually in cargo?
LoganDark 14 hours ago [-]
> This is very likely not the kind of feature that will stay on nightly for a long time, but is instead one of the many convenience feature that land on stable every release.
Easy example of a basic method that has been unstable for a really long time: [T]::as_slice [0] since 2024 [1]. Apparently, stabilization was attempted earlier this year [2] but was then rolled back [3]. While clearly it was not yet ready for stabilization, it still took over a year before the first attempt.
Another one: Option::zip_with [4] since 2020 because nobody's figured out if it's worth having over .zip(...).map(...). Option::zip was actually stabilized [5] later in 2020 but Option::zip_with has since been sitting in limbo for over five years.
Another one: <*const [T]>::as_ptr also since 2020 [6]. I can't remember if there's an alternative now but dealing with slice pointers without relying on unstable methods has historically been very difficult/annoying. I ran into a bunch of this kinda stuff while working on a crate for iterating over rows/columns of image buffer subregions, because I wanted to use and support slice pointers. (Specifically I think getting the length of the slice pointer was nearly impossible without invoking UB, because constructing a reference (which was the only safe way to access a len method) could break aliasing rules. However I think the len method on slice pointers was stabilized a while ago so that particular problem is no more.) Speaking of which, <*mut [T]>::split_at_mut has been unstable since 2022 [7]...
I'm not saying there's no reason for any of this, just that as a Rust developer it's been frustrating. There are enough of these all over the place that it feels like a real occurring problem, even if it's not reasonable to expect a volunteer open-source project to pay full attention to everything ever.
> Can't you tell it to use stable as the default target, and use nightly manually in cargo?
Are you saying it doesn't suggest unstable features when using a stable toolchain? That was not my experience before I started using nightly.
stymaar 13 hours ago [-]
> Are you saying it doesn't suggest unstable features when using a stable toolchain? That was not my experience before I started using nightly.
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.
My heard hurts - i was stupid enough to think that SIMD was a CPU only thing - I don't understand why it would be ported to GPU - huge kudos to managing to surprise me
MindSpunk 20 hours ago [-]
It's not really obvious unless you go in depth of the details on modern GPU architecture. GPUs aren't really SIMD, they're SIMT (single instruction multiple thread). The silicon looks a lot like SIMD, but the programming model is different.
If you go look at AMD's ISA docs (they're public) you'll see you don't have the equivalent of a __mm256 register like on x86. Each 'thread' just deals with single scalar values like int32 of float32. The hardware, however, groups 32 or 64 threads together which all run the same program and runs them together. Each 'thread' loosely maps to a SIMD lane. The SIMD is implicit, not explicit.
The main difference is that the 'SIMD' execution is somewhat opaque to the program. You just write plain scalar code and the hardware model dispatches it efficiently to SIMD execution units. It's not really an abstraction because to extract maximum performance you have to understand how it works. You can use this kind of programming model on a CPU too, Intel did it with [0] ISPC. It's a C-like language that has execution semantics similar to GPU shader languages but compiles to regular CPU code, and maps threads to your CPUs SIMD lanes like a GPU.
This might be more confusing than it needs to be. SIMD and SIMT are not mutually exclusive.
People commonly think of things like vector registers when they talk about SIMD, and each "thread" in a GPU warp definitely deals with local vector registers. Granted, they may be slices of superwide registers shared by the whole warp, or whatever else, but from the programmer's perspective, that's a valid way to think about it.
Put another way, it would be a mistake to think that each lane of a vec4 in a shader gets processed by a separate unit.
False. If they were threads they'd have their own PC. They do not - only the warp has a PC.
> You just write plain scalar code and the hardware model dispatches it efficiently to SIMD execution units.
Absolutely not. If you don't write coalesced loads, bank-conflict free, predication-free, cooperative code you will get worse than CPU performance.
MindSpunk 17 hours ago [-]
Yes, of course writing naive code assuming each lane in a thread group is a real thread is going to cause problems, but I didn't feel like I needed to go into that level of detail replying to someone just learning about GPU internals. I tried to cover this loosely by mentioning how you need to know how it works for maximum performance.
If you want to get more pedantic you also need to look at your target hardware and their specific micro-architectural quirks and features to get the best performance. AMD specifically benefits a lot from exploiting the scalar unit over the vector unit, you save loads of register file space if you can keep data in SGPRs over VGPRs. There's lots of traps you can fall into where you can load data from buffers into SGPRs but they get promoted to VGPRs because the scalar unit lacks an opcode for like one math operation you did to the value somewhere.
While each lane isn't truly a thread because it doesn't have its own PC the programming model definitely tries to make it seem that way. The threads can terminate at different points too. And again, the ISA isn't a vector ISA. Your register values are scalar.
pezezin 10 hours ago [-]
> While each lane isn't truly a thread because it doesn't have its own PC the programming model definitely tries to make it seem that way. The threads can terminate at different points too. And again, the ISA isn't a vector ISA. Your register values are scalar.
This is not correct. If you check AMD's documentation there are explicit mentions of vector registers (VGPR), vector ALUs, and vector instructions. The introduction to Chapter 2 describes it as a vector ISA.
> RDNA4 shader programs (kernels) are programs executed by the shader processor. Conceptually, the shader program is executed independently on every work-item, but in reality the processor groups up to 32 or 64 work-items into a wave, that executes the shader program on all 32 or 64 work-items in one pass ("wave32" or
"wave64").
A VGPR is not the same thing as a vector register like in SSE4 or AVX. Each addressed register contains a single 32-bit value. A VGPR differs from an SGPR in that each thread in a thread group can have a different value in that register. An SGPR will have a uniform value shared with all threads in a group.
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.
> False. If they were threads they'd have their own PC. They do not - only the warp has a PC.
They are using the term SIMT as it is normally used[1]. The "single instruction" part means that there is only one PC shared across multiple 'threads'.
Actually not so false anymore. (But still they don't expect you to use this knowledge while coding, and you should treat all threads in a warp as moving in lockstep)
> In GPUs of compute capability 7.0 and later, independent thread scheduling allows full concurrency between threads, regardless of warp. With independent thread scheduling, the GPU maintains execution state per thread, including a program counter and call stack... [1]
GPU "cores" are basically what a CPU would call SIMD lanes. So a GPU with 1024 'CUDA cores' might be structured as 16 relatively independent pieces that a CPU might call a core, each with a 64 wide SIMD unit.
mathisfun123 1 days ago [-]
32 wide - only AMD has a 64 wide mode
y1n0 22 hours ago [-]
64 what? Bits/bytes/something bigger?
corysama 21 hours ago [-]
When you dig through the CUDA developer docs instead of the promotional materials, you can develop a view of Nvidia GPUs as having 8-128 processing cores, each with 4 hyperthreads, running 32-lane SIMD for almost everything. Where a lane is 32 bits wide.
The promotional material likes to label the individual lanes as “cores” because it sounds more impressive. And, it’s not entirely incorrect.
Even the dev docs use the marketing terminology. The description I gave above needs a bit of piecing together.
dataflow 18 hours ago [-]
This is a fantastic explanation, thanks for writing it. It also makes me wonder something: where exactly is the biggest difference between a 32-core x86 CPU (AVX512 basically being 16 32-bit lanes) and (say) an NVIDIA GPU with ~8-16 processing cores? Like why can't the CPU compete against a GPU like that for GPU-y tasks - or can it?
jandrewrogers 17 hours ago [-]
At risk of over-simplifying, GPUs are wider with limited computational expressiveness and higher memory bandwidth while CPUs are highly expressive computationally (and better connected to I/O) but with lower memory bandwidth. GPUs are less sensitive to memory latency by necessity. Even AVX-512 is highly flexible when inter-mixed with scalar code. GPUs get their very high register width by restricting what the cores are capable of doing efficiently.
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.
Tuna-Fish 12 hours ago [-]
> Like why can't the CPU compete against a GPU like that for GPU-y tasks - or can it?
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.
monocasa 17 hours ago [-]
Not corysama, but I'll take a stab:
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.
corysama 8 hours ago [-]
On a CPU, hyperthreads are mostly replicated register banks. This allows the CPU to hold the context for 2 threads simultaneously. And, lets parts of a CPU make progress on one thread while the other thread is stalled. CPUs also has a kinda large microcode register bank that helps work around dependencies in asm instructions that reuse named registers.
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.
dahart 17 hours ago [-]
If by processing cores you mean Streaming Multiprocessors (SMs), keep in mind that each SM is 128 threads, so 16 SMs is 2k threads - or 2048 single precision math ops per cycle. The 32 core avx512-enabled CPU is 512 math ops per cycle, if you have one FMA unit per core (or 1k ops/cycle if you have 2 FMA units/core).
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.
Jhsto 11 hours ago [-]
SPIR-V states its an int, float, vector n (where n <= 4) or a matrix (2..4 cols of vector n).
It does not necessarily mean the hardware can do 4x4x64 floating point operations in a single subgroup operation, but at least the programming model supports framing it that way.
chlorion 1 days ago [-]
GPUs work on vectors and matrices very often, that's what they are good at, so it makes a lot of sense that they can operate with SIMD I think!
ismailmaj 1 days ago [-]
There is something very SIMD-coded in GPU programming which is coalesced stores/loads, if a warp (32 threads) handles contiguous memory, it will create ~4 transactions instead of 32.
hingler36 1 days ago [-]
Welcome to the lucky 10,000! SIMD is actually a pretty integral part of how GPUs are able to work efficiently, it's part of why there's such a strong focus on branchless programming in the field.
grokcodec 23 hours ago [-]
I would love to have an open source Rust SIMD library with the scope and maturity that
https://github.com/google/highway brings to C++.
raphlinus 22 hours ago [-]
This is basically the goal of fearless_simd, but of course achieving the same level of maturity will take time.
camel-cdr 1 days ago [-]
I love how ever example of portable SIMD isn't portable.
They specifies a constant SIMD width so it's non-portable. Well, not performance portable, but why are we using SIMD again?
jandrewrogers 20 hours ago [-]
The capabilities of various SIMD ISAs don't have enough intersection to be portable outside of relatively trivial cases. Many of the somewhat unique capabilities are load-bearing, so you want to use them on architectures that support them. Taken in whole, someone who cares about performance would be using different data structures and algorithms depending on the specific SIMD architecture and that is nearly impossible to abstract in a library. Too many important but complex details are idiosyncratic to the implementation.
Another way of looking at it is that our programming environments are not sufficiently powerful and expressive to create the necessary abstractions to make SIMD truly portable.
simonask 8 hours ago [-]
> The capabilities of various SIMD ISAs don't have enough intersection to be portable outside of relatively trivial cases.
I would argue that the "trivial" cases (those relating to linear algebra in 3 dimensions) are also 95% of what people want SIMD for.
If the API can achieve cross-platform and performant vector arithmetic, dot product, and matrix multiplication in the normal ways, that already covers a lot of what people actually need.
jandrewrogers 6 hours ago [-]
Most use cases for SIMD are non-arithmetic in nature and don't assume tidy arrays of homogeneous types. I also use it for some computational geometry but that is the least interesting use case.
SIMD is widely used throughout data infrastructure e.g. parsing data, complex constraint processing, parallel manipulation of heterogeneous data types, compression, etc. I even have an I/O scheduler written in AVX-512 that is many times faster than the scalar equivalent. The ability of SIMD to do complex manipulation of ordinary data structures several times faster than scalar code is under-rated.
While linear algebra is the current thing, database engines have been using SIMD heavily for over a decade and arguably represent the frontier. It is for these use cases that SIMD is non-portable and data infrastructure isn't going away.
zamadatix 24 hours ago [-]
It should really be read/advertised as "portabler SIMD". It beats hoping the compiler autovectorizes everything well forever or writing architecture specific code manually again but is going to compromise on average performance vs platform specific SIMD.
pjmlp 16 hours ago [-]
.NET and Java have three levels of SIMD support, Go's ongoing efforts, and does the upcoming C++ standard.
Autovectorization, depending on compiler's cleverness, really portable SIMD operations, and then the CPU specific SIMD ones.
So this should be perfectly doable in crate that advertises as portable, while leaving the non portable stuff to another crate.
exDM69 15 hours ago [-]
> They specifies a constant SIMD width so it's non-portable.
This is incorrect, you can use vectors wider than native SIMD width and the compiler will break them down to register size of the target cpu.
In fact it's sometimes better to used wider than native width, in some applications I see 20% better throughput with f32x16 (512 bits) on an AVX2 CPU (256 bits). It is kinda like loop unrolling it.
camel-cdr 15 hours ago [-]
Except you can't use this in actual code, because either, as is the case in this example with f32x32, you run out of registers and spill all over the place.
Or you aren't using your full vector register or could've gotten better performance by "unrolling" more often for the larger vectors.
If you use f32x16 (the avx-512 wisth), SSE now effectively has 4 registers to work with and will spill when doing anything beyond the most simple stuff.
The default should imo be relative to the native register width, so you can do 1x, 2x or sometimes 4x the native width, depensing on your register preasure.
exDM69 13 hours ago [-]
I can and I do use this is "actual code" and I've got benchmarks to prove that it's got better throughput (for the particular use case, don't extrapolate from there) and the same applies to AVX2 and AVX512: twice the native vector width has ~20% better throughput (ie. using `f32x32` on AVX-512).
I pass in the vector width as a generic parameter like this:
With this I can easily benchmark the same code for any vector width. I can also do some compile time heuristics to choose the vector width based on what's available on the compile target CPU.
> you run out of registers and spill all over the place
As usual when optimizing SIMD code, you should keep an eye on the generated disassembly and the benchmark results and watch for register pressure and the other usual things.
I'm definitely NOT saying that you always get the best perf by using 2x SIMD width, but in this particular case it was so.
This is much much easier to do with portable_simd than if you'd write the same with intrinsics, you can change the SIMD width without having to rewrite all your code (e.g. changing from SSE `_mm_add_ps` to AVX `_mm256_add_ps` etc).
It's still a partial solution, you still need to drop down to intrinsics for some special instructions every now and then (which is easy), but in my projects this accounts for much less than 1% of the lines of code. Not applicable everywhere of course.
camel-cdr 13 hours ago [-]
> twice the native vector width has ~20% better throughput
Yes, this is what I was saying, but twice the vector width of AVX-512 will perform horrible in SSE, which is why portable SIMD abstractions should make writing code relative to the native vector width simple.
> I pass in the vector width as a generic parameter like this:
My problem is that no portable_simd example code I've seen does this, which causes people to choose one specific N and run with that.
The second part of the problem is how you find the native vector length, so you can instantiate the generic function. IIRC this isn't even exposed in portable_simd and you have to use a seperate crate to get it.
exDM69 13 hours ago [-]
> The second part of the problem is how you find the native vector length, so you can instantiate the generic function. IIRC this isn't even exposed in portable_simd and you have to use a seperate crate to get it.
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.
Sure but there's no real way to use that in a portable way, at least not a way that maximises performance on every CPU you run it on. That's pretty much impossible at the moment.
krapht 23 hours ago [-]
Which is why I've never quite understood the appeal of portable SIMD libraries for performance-critical code. If I'm explicitly writing SIMD rather than relying on the auto-vectorizer, it's usually because I want access to the particular capabilities of the target ISA.
For many problems, choosing the right instruction or instruction sequence makes a large difference. Portable SIMD abstractions necessarily expose some common semantic layer, but SIMD ISAs don't actually have equivalent capabilities. Instructions like pshufb, for example, enable algorithmic tricks that don't necessarily have an equally efficient analogue on another architecture.
If maximum performance matters, I generally want intrinsics and architecture-specific implementations; if portability matters more, I'd rather move further up the abstraction stack and use something designed to target multiple architectures, such as ISPC. There are certainly cases where portable SIMD gets close enough to optimal, but I don't think there's a compiler or abstraction that can express every useful SIMD idiom and lower it equally efficiently across fundamentally different ISAs.
pjmlp 16 hours ago [-]
Because usually they achieve a very good middle ground, they are useful for when autovectorization isn't good enough, and it is possible to give a little help to the compiler.
There are many ways that performance matters without trying to win a F1 race.
Go isn't alone, .NET, Java have similar portable libraries, and C++ is in the process of getting one.
kbolino 8 hours ago [-]
Go doesn't have auto-vectorization in the first place, so its portable simd library is at least partly there to fill the gap.
MomsAVoxell 1 days ago [-]
Why should it be portable? Honest question.
SIMD seems to me, to be very platform specific. Maybe there are times one SIMD unit is not anothers' SIMD unit?
camel-cdr 1 days ago [-]
The create is called portable_simd.
There is no reason a portable_simd relu_dot implemention should need to specify the SIMD width.
But the design and documentation of portable_simd makes the fixed size syntactically easy/the default and the width agnostic code harder.
dwattttt 23 hours ago [-]
> There is no reason a portable_simd relu_dot implemention should need to specify the SIMD width.
What should it choose then? I have a Zen 3 processor, and benchmarking some simd I did recently says 32 byte or 64 byte chunks was fastest. But I'm sure I'd get a different result on a different Zen, and different again on Intel's.
How would the library decide what SIMD width I should use?
Groxx 23 hours ago [-]
It'd need some kind of compile-time hardware-feature-detection, yea? That seems probably feasible since proc macros can do essentially anything they like (worryingly).
derefr 22 hours ago [-]
Only if the end-user is the one compiling the software, on the same very system they'll be running it on. Which is true of GPU shader kernels, due to how GPU drivers work; but isn't generally true of CPU object code (unless you're on Gentoo.)
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.
dwattttt 18 hours ago [-]
If you want the library to perform that selection, you also need the "correct" / most efficient implementation to be independent of your workload. I'm not that familiar with SIMD performance characteristics, but I wouldn't be surprised if that's not always the case.
derefr 16 hours ago [-]
From how I understand it, there'd likely only be a single SIMD function impl per uarch that'd actually be fully legally executable without hitting undefined instructions. Plus increasingly-more-generic function impls compiled for lower and lower common-denominator subsets of SIMD functionality. (Ultimately grounding in a non-SIMD impl.)
If that's the case, then the selection logic would be trivial: figure out the full hierarchical ID of the uarch you're running on, then search for the longest prefix match in the table of available impls.
If things work more like you're imagining, though, then I suppose the process-boot impl-selector would narrow down the impl matrix to just the subset that are legal on the running uarch; pick one arbitrarily to be active at first; and then wrap the calls in a handler that gradually re-works the called function in a way reminiscent of a profile-guided JIT, but without the need to actually synthesize any code at runtime — instead, it'd just be a multi-armed bandit passing-through-to and re-ranking competitor impls, with decreasing sampling of the non-first-ranked impls as confidence-in-score-separation increases.
adgjlsfhk1 16 hours ago [-]
This is one of the nice things about JIT languages. you defer the compiler time decisions to runtime and this get to choose based on what the user has
exDM69 15 hours ago [-]
Counter question: why shouldn't it be portable?
It's definitely a 80% solution where you occasionally need to drop down to intrinsics (at zero runtime perf cost) for CPU specific instructions.
But just having vector types, arithmetic, swizzling, loads and stores will go a long way for basic tasks.
And with generics you can write code that is type and width agnostic. No need to rewrite your code of you want to go from SSE to AVX512, just change from f32x4 to f32x16 (or use generics) and you are done.
MomsAVoxell 3 hours ago [-]
>Counter question: why shouldn't it be portable?
Because there are platform vendors. And SIMD performance very much depends on the use-case, which is a balance of practicalities and specifications and intended deployment targets ..
I also think this is a deployment problem, not a build problem, but okay ..
nynx 1 days ago [-]
Do you have examples of complex algorithms running on the gpu with rust with competative performance? Radix sort might be a good one to start with
LegNeato 1 days ago [-]
Author here, AMA.
lbhdc 1 days ago [-]
What is vectorware's business model? Are you planning to sell support/consulting to companies using your stack? Or are you looking to sell licenses to your tool? Or something else?
LegNeato 1 days ago [-]
The tentative plan is to open source all the compiler and `std` bits with our products built on top (compilers are not good businesses). More about our products coming in the next couple of months!
lbhdc 1 days ago [-]
Looking forward to reading more about it. Good luck on the launch :)
jcranmer 1 days ago [-]
The post is kind of vague on the IR you're targeting. Can you give some examples of what the SIMD-ized IR looks like, and how it maps to the target PTX?
LegNeato 1 days ago [-]
Didn't want to go into crazy detail in the post.
Each family of operations is a trait parameterized by the operation itself:
pub trait EvaluateReduction<Operation, T>: LaneEvaluator {
/// Reduce one distributed definition to an ordinary uniform scalar.
fn evaluate_reduction(&self, value: LaneValue<Self, role::Distributed, T>) -> T;
}
Call sites name the operation:
let one = evaluator.splat::<Splat, _>(1_u32);
let two = evaluator.splat::<Splat, _>(2_u32);
let three = evaluator.binary::<Add, _>(one, two);
let total = evaluator.reduce::<Sum, u32>(three); // a uniform u32
let running = <Executor as EvaluateScan<Scan<Sum, Exclusive>, u32>>::scan(&evaluator, three);
Operations like Sum, Max, ReduceXor, Inclusive, and Exclusive are all distinct types.
As mentioned in the post, execution shape is typed too. A static shuffle takes its control as a type-level constant, and the shuffle mode constrains which controls are expressible:
// Shift down one lane, keeping our own value where the source is inactive.
let down = <Executor as EvaluateShuffle<Shuffle<Down>, DownOrSelf<1>, u32>>::shuffle(&ev, v);
// Broadcast from lane zero.
let bcast = <Executor as EvaluateShuffle<Shuffle<Broadcast>, WarpLane<0>, u32>>::shuffle(&ev, down);
// Butterfly exchange with the neighbor one bit away.
let bfly = <Executor as EvaluateShuffle<Shuffle<Xor>, Butterfly<1>, u32>>::shuffle(&ev, bcast);
For an example of errors caught, a warp-scoped executor for a device-scoped barrier is a compile error:
<ScopedWarpExecutor<'_, WarpUniform> as EvaluateBarrier<Barrier<Device>>>::barrier(evaluator)
// error[E0277]: the trait bound `Device: NvptxBarrierScope` is not satisfied
// help: the trait `NvptxBarrierScope` is implemented for `Warp`
Strip mining is typed on the amount of work and the lane capacity, and it hands back one chunk at a time along with the predicate saying which lanes live in that chunk:
// Six work items across four active lanes: two chunks, based at 0 and 4.
<Executor as EvaluateStripMine<StripMine, (WorkItems, ActiveLanes<StripMined<4>>), i32>>::
for_each_strip_mined(
&evaluator,
(WorkItems::new(6)?, ActiveLanes::new(4)?),
|index, active| {
// ...
},
);
Hopefully that gives the flavor of it.
the__alchemist 1 days ago [-]
I'm confused too. How does this fit between these approaches for paraellization:
- CUDA kernels and Tiles (e.g. Cudarc, cuda-oxide, rust-gpu etc) - SIMD on the GPU. (E.g. as in the title...)
- CPU SIMD using avx or SSE instructions (And probably thin wrappers for vectors so you can have sane syntax). Or the maybe-upcoming core simd which should abstract over architecture-specific instructions. Magic floats etc which do 4-16 computations at once, but are a bit clumsy to work with
- Rayon thread pools - arbitrary parallel computations, including SIMD, one per CPU core.
It looks like from the code samples like maybe a cleaner syntax for writing code on the GPU than CUDA kernels? E.g. without mucking with serialization, host and device by abstracting over it? And inspired by core::simd. (Good choice if so, in the interest of standardizing on syntax; I did this for my x86 SIMD vector/quaternion lib as well)
lbhdc 1 days ago [-]
This is really cool! It sounds like y'all have a compiler fork that you are using to make this work. I wanna tinker with this, is your compiler available?
LegNeato 1 days ago [-]
It is not currently available but we intend to make it available after we launch our products.
bbminner 23 hours ago [-]
If you have to express your computation using an "array programming DSL" with things like scan and gather anyways - why not opt to use torch/tensorflow/jax or anything else that targets MLIR? An example of writing a relu using an embedded array DSL is really not helping your case either - that's exactly the problem that these other solutions mentioned above are successfully solving for the past ~15y (starting with theano etc). Not sure what this brings to the table - doing that AoT instead of at runtime?
LegNeato 22 hours ago [-]
The goal of this work is to run existing unmodified CPU libraries (which may use core::simd) on the GPU. If you are manually writing ML-shaped workloads, it doesn't add any value over writing with tech like torch/tensorflow/jax which are custom built for those use-cases (except maybe familiarity if you are a CPU programmer).
peterbower 10 hours ago [-]
All well and good but where can we install it now?
Eridrus 1 days ago [-]
Given the massive demand for GPUs for LLMs, what sorts of work do you expect to economically benefit from utilizing GPUs more?
LegNeato 1 days ago [-]
Part of our thesis is that decent GPUs are in every shipping device and most software doesn't use them and should.
Eridrus 1 days ago [-]
I guess you're looking at consumer hardware then since servers have exactly what you pay for.
Can you say more about the application space you're targeting?
shay_ker 24 hours ago [-]
Hm is the intent to one day replace the CPU?
LegNeato 24 hours ago [-]
The goal is to use similar abstractions and code across both the CPU and GPU where it makes sense.
PoignardAzur 1 days ago [-]
Any thoughts about SIMD-related crates?
max-m 24 hours ago [-]
How was your day?
adityazero 24 hours ago [-]
[dead]
guess__who 1 days ago [-]
[flagged]
melodyogonna 21 hours ago [-]
Very interesting.
But GPU programming gets complicated when you start doing 3d computation on very large data, will be interesting to see how tensor abstraction is built on top of this. Another point is that this is using fixed-width SIMD vectors; unless there is a way to compute this statically based on available GPU info, performance will always be left on the table.
frollogaston 19 hours ago [-]
I've noticed a lot of articles about SIMD on the HN front page. That's cool, but just wondering, is there some reason this is more in focus lately?
vatsachak 18 hours ago [-]
SIMD is actually underrated still. Programmers should always be thinking about it. It's a free 4x in a lot of cases
pjmlp 16 hours ago [-]
The main problem with SIMD is it is a complex subject, even if not that good, autovectorization wins over what most common devs know about SIMD.
Well now you could in theory AI generate SIMD, which will be vibe coded, as those devs have no idea of its correctness.
frollogaston 18 hours ago [-]
I did see the article about that too. Don't know about "always" since there are applications like web backends where you're never going to add arrays of floats or something. Even if it's data science stuff, if that's in Python, Numpy is doing the SIMD for you.
skitsofrandom 18 hours ago [-]
I’m wondering if AI has made SIMD intrinsics much more approachable for many and so there are just more people working on abstractions for their workflow of choice right now. There’s probably a lot of code out there that could benefit from SIMD but the effort to actually use it was too high for the return.
LoganDark 19 hours ago [-]
Maybe things being on the front page reminds others? After seeing something, sometimes you can have ideas relating to it for a while.
samuell 19 hours ago [-]
Yes, this kind of thing seems to happen quite often. Popular posts spurring further posts on a theme.
dev_dan_2 23 hours ago [-]
Really exciting work and great write up, thanks a lot and all the best to your startup!
`core` instead of `std` is great too!
This will become useful in one of my sideproject where I use bitmaps to speed up pathfinding, exited to try it out!
efnx 1 days ago [-]
Congrats to the Rust-GPU folks! Nice to see the good work flowing.
I don't get what's the value of it not being enabled by default what does the toggle get us, really? Maybe I don't understand web design and it makes it harder to read for some, I am dyslexic and never had any issues.
LegNeato 23 hours ago [-]
It's just a way for us to add minutia and details that most don't care / need to know about. There are three audiences we try to make the posts accessible for: Rust people who don't know about GPUs, GPU people who don't know about Rust, and non-Rust non-GPU people. The toggle lets knowledgable readers go "wait, what about..." and hopefully the toggle answers it.
minraws 22 hours ago [-]
Does making it the default hurt anyone though? I don't think that 1 subscript adds anything to it..
LegNeato 22 hours ago [-]
Just distracting / ugly to my eyes.
24 hours ago [-]
reindeer2 8 hours ago [-]
[dead]
the__alchemist 1 days ago [-]
Hey - this is probably off-topic/meta, but what is going on with the comments here? Is it bots?
dev_l1x_be 1 days ago [-]
No idea, but it seems HN needs POW challenges.
lukan 1 days ago [-]
Could also just be trolls attracted by the Rust topic.
dev_l1x_be 1 days ago [-]
I bet some kid is bored out of his mind and wrote a bot.
1 days ago [-]
neonsunset 13 hours ago [-]
[dead]
rust-lang 1 days ago [-]
[flagged]
donald-trump 1 days ago [-]
[flagged]
donald-trump 1 days ago [-]
[flagged]
lx-user 1 days ago [-]
[flagged]
1 days ago [-]
guess__who 1 days ago [-]
[flagged]
LegNeato 1 days ago [-]
We never mention anything about superiority nor compare with other languages or programming models. This post is about making existing Rust CPU code work on the GPU.
guess__who 1 days ago [-]
[flagged]
fluffybucktsnek 24 hours ago [-]
Those aren't lines you are reading. Those are your hallucinations.
At worst, the post reads like a propaganda for VectorWare, but, overall, it reads more like their insights on the matter.
kooi 1 days ago [-]
Useless blabbering.
If you have real feedback, great, but it's useless to rip on the hard work of others without it.
fire_wheel 1 days ago [-]
I am rewriting some components in rust - I dont like such sentiment. It negatively impacts the community.
Rust is better in so many ways.
throwaway894345 1 days ago [-]
I'm not a Rust user apart from an occasional toy program here and there, but you seem really triggered about a language that other people use. What's the issue?
the__alchemist 1 days ago [-]
I care because it means I can use this in a Rust program without a FFI barrier. Regrettably, we have built computing infrastructure as a society with many barriers; programming language is one.
extrem-RAM-shor 1 days ago [-]
Rust is the best. No other language is fun enough to program.
[0] https://doc.rust-lang.org/std/simd/index.html
[1] https://github.com/linebender/fearless_simd
A lot of bugs with SIMD libraries are in the domain of interactions, not functionality--e.g. SIMD malfunctions on rare chips, chips with previously-unseen combinations of hardware/userspace firmware/microcode behavior, compilers run in weird harnesses that lie about hardware capabilities, and so on. I assume that's the case with Rust's portable SIMD as well.
If your QA is unpredictable individual use-cases (as with most OSS projects), then there's no way to measure "testing complete" or "coverage"; letting it bake for awhile is the best approach available.
I most often encounter unstable features when I reach for a basic common-sense utility method and discover that it's not stable. Like just earlier today I would have reached for bool::toggle which not only is unstable, but is also newly added as of like a month ago! but some unstable methods have been sitting around for years.
And now that IntelliJ-Rust is proprietary, I can't even make a feature request anymore for the ability to exclude unstable features from the autocomplete. So they will taunt me forever, perfect little helpers just locked away.
It used to be the case a decade ago, but now I wouldn't agree that any "basic" things require nightly (I wouldn't call portable SIMD "basic" at all for instance).
> Like just earlier today I would have reached for bool::toggle which not only is unstable, but is also newly added as of like a month ago!
This is very likely not the kind of feature that will stay on nightly for a long time, but is instead one of the many convenience feature that land on stable every release. The 6-weeks release cadence with beta in between means there's always at least 6 weeks and up to 3 months between the time a feature land on nightly and the day it reaches stable, even if the feature is as consensual as this one.
> And now that IntelliJ-Rust is proprietary, I can't even make a feature request anymore for the ability to exclude unstable features from the autocomplete.
Can't you tell it to use stable as the default target, and use nightly manually in cargo?
Easy example of a basic method that has been unstable for a really long time: [T]::as_slice [0] since 2024 [1]. Apparently, stabilization was attempted earlier this year [2] but was then rolled back [3]. While clearly it was not yet ready for stabilization, it still took over a year before the first attempt.
Another one: Option::zip_with [4] since 2020 because nobody's figured out if it's worth having over .zip(...).map(...). Option::zip was actually stabilized [5] later in 2020 but Option::zip_with has since been sitting in limbo for over five years.
Another one: <*const [T]>::as_ptr also since 2020 [6]. I can't remember if there's an alternative now but dealing with slice pointers without relying on unstable methods has historically been very difficult/annoying. I ran into a bunch of this kinda stuff while working on a crate for iterating over rows/columns of image buffer subregions, because I wanted to use and support slice pointers. (Specifically I think getting the length of the slice pointer was nearly impossible without invoking UB, because constructing a reference (which was the only safe way to access a len method) could break aliasing rules. However I think the len method on slice pointers was stabilized a while ago so that particular problem is no more.) Speaking of which, <*mut [T]>::split_at_mut has been unstable since 2022 [7]...
I'm not saying there's no reason for any of this, just that as a Rust developer it's been frustrating. There are enough of these all over the place that it feels like a real occurring problem, even if it's not reasonable to expect a volunteer open-source project to pay full attention to everything ever.
[0]: https://doc.rust-lang.org/std/primitive.slice.html#method.as...
[1]: https://github.com/rust-lang/rust/issues/130366
[2]: https://github.com/rust-lang/rust/pull/151603
[3]: https://github.com/rust-lang/rust/pull/152963
[4]: https://github.com/rust-lang/rust/issues/70086
[5]: https://github.com/rust-lang/rust/pull/72938
[6]: https://github.com/rust-lang/rust/issues/74265
[7]: https://github.com/rust-lang/rust/issues/95595
> Can't you tell it to use stable as the default target, and use nightly manually in cargo?
Are you saying it doesn't suggest unstable features when using a stable toolchain? That was not my experience before I started using nightly.
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.
If you go look at AMD's ISA docs (they're public) you'll see you don't have the equivalent of a __mm256 register like on x86. Each 'thread' just deals with single scalar values like int32 of float32. The hardware, however, groups 32 or 64 threads together which all run the same program and runs them together. Each 'thread' loosely maps to a SIMD lane. The SIMD is implicit, not explicit.
The main difference is that the 'SIMD' execution is somewhat opaque to the program. You just write plain scalar code and the hardware model dispatches it efficiently to SIMD execution units. It's not really an abstraction because to extract maximum performance you have to understand how it works. You can use this kind of programming model on a CPU too, Intel did it with [0] ISPC. It's a C-like language that has execution semantics similar to GPU shader languages but compiles to regular CPU code, and maps threads to your CPUs SIMD lanes like a GPU.
[0] https://ispc.github.io/
This might be more confusing than it needs to be. SIMD and SIMT are not mutually exclusive.
People commonly think of things like vector registers when they talk about SIMD, and each "thread" in a GPU warp definitely deals with local vector registers. Granted, they may be slices of superwide registers shared by the whole warp, or whatever else, but from the programmer's perspective, that's a valid way to think about it.
Put another way, it would be a mistake to think that each lane of a vec4 in a shader gets processed by a separate unit.
False. If they were threads they'd have their own PC. They do not - only the warp has a PC.
> You just write plain scalar code and the hardware model dispatches it efficiently to SIMD execution units.
Absolutely not. If you don't write coalesced loads, bank-conflict free, predication-free, cooperative code you will get worse than CPU performance.
If you want to get more pedantic you also need to look at your target hardware and their specific micro-architectural quirks and features to get the best performance. AMD specifically benefits a lot from exploiting the scalar unit over the vector unit, you save loads of register file space if you can keep data in SGPRs over VGPRs. There's lots of traps you can fall into where you can load data from buffers into SGPRs but they get promoted to VGPRs because the scalar unit lacks an opcode for like one math operation you did to the value somewhere.
While each lane isn't truly a thread because it doesn't have its own PC the programming model definitely tries to make it seem that way. The threads can terminate at different points too. And again, the ISA isn't a vector ISA. Your register values are scalar.
This is not correct. If you check AMD's documentation there are explicit mentions of vector registers (VGPR), vector ALUs, and vector instructions. The introduction to Chapter 2 describes it as a vector ISA.
> RDNA4 shader programs (kernels) are programs executed by the shader processor. Conceptually, the shader program is executed independently on every work-item, but in reality the processor groups up to 32 or 64 work-items into a wave, that executes the shader program on all 32 or 64 work-items in one pass ("wave32" or "wave64").
Sources:
https://gpuopen.com/amd-gpu-architecture-programming-documen...
https://docs.amd.com/v/u/en-US/rdna4-instruction-set-archite...
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.
> False. If they were threads they'd have their own PC. They do not - only the warp has a PC.
They are using the term SIMT as it is normally used[1]. The "single instruction" part means that there is only one PC shared across multiple 'threads'.
[1] https://en.wikipedia.org/wiki/Single_instruction,_multiple_t...
> In GPUs of compute capability 7.0 and later, independent thread scheduling allows full concurrency between threads, regardless of warp. With independent thread scheduling, the GPU maintains execution state per thread, including a program counter and call stack... [1]
1: https://docs.nvidia.com/cuda/cuda-programming-guide/03-advan...
The promotional material likes to label the individual lanes as “cores” because it sounds more impressive. And, it’s not entirely incorrect.
Even the dev docs use the marketing terminology. The description I gave above needs a bit of piecing together.
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.
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.
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.
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.
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.
It does not necessarily mean the hardware can do 4x4x64 floating point operations in a single subgroup operation, but at least the programming model supports framing it that way.
They specifies a constant SIMD width so it's non-portable. Well, not performance portable, but why are we using SIMD again?
Another way of looking at it is that our programming environments are not sufficiently powerful and expressive to create the necessary abstractions to make SIMD truly portable.
I would argue that the "trivial" cases (those relating to linear algebra in 3 dimensions) are also 95% of what people want SIMD for.
If the API can achieve cross-platform and performant vector arithmetic, dot product, and matrix multiplication in the normal ways, that already covers a lot of what people actually need.
SIMD is widely used throughout data infrastructure e.g. parsing data, complex constraint processing, parallel manipulation of heterogeneous data types, compression, etc. I even have an I/O scheduler written in AVX-512 that is many times faster than the scalar equivalent. The ability of SIMD to do complex manipulation of ordinary data structures several times faster than scalar code is under-rated.
While linear algebra is the current thing, database engines have been using SIMD heavily for over a decade and arguably represent the frontier. It is for these use cases that SIMD is non-portable and data infrastructure isn't going away.
Autovectorization, depending on compiler's cleverness, really portable SIMD operations, and then the CPU specific SIMD ones.
So this should be perfectly doable in crate that advertises as portable, while leaving the non portable stuff to another crate.
This is incorrect, you can use vectors wider than native SIMD width and the compiler will break them down to register size of the target cpu.
In fact it's sometimes better to used wider than native width, in some applications I see 20% better throughput with f32x16 (512 bits) on an AVX2 CPU (256 bits). It is kinda like loop unrolling it.
If you use f32x16 (the avx-512 wisth), SSE now effectively has 4 registers to work with and will spill when doing anything beyond the most simple stuff.
The default should imo be relative to the native register width, so you can do 1x, 2x or sometimes 4x the native width, depensing on your register preasure.
I pass in the vector width as a generic parameter like this:
With this I can easily benchmark the same code for any vector width. I can also do some compile time heuristics to choose the vector width based on what's available on the compile target CPU.> you run out of registers and spill all over the place
As usual when optimizing SIMD code, you should keep an eye on the generated disassembly and the benchmark results and watch for register pressure and the other usual things.
I'm definitely NOT saying that you always get the best perf by using 2x SIMD width, but in this particular case it was so.
This is much much easier to do with portable_simd than if you'd write the same with intrinsics, you can change the SIMD width without having to rewrite all your code (e.g. changing from SSE `_mm_add_ps` to AVX `_mm256_add_ps` etc).
It's still a partial solution, you still need to drop down to intrinsics for some special instructions every now and then (which is easy), but in my projects this accounts for much less than 1% of the lines of code. Not applicable everywhere of course.
Yes, this is what I was saying, but twice the vector width of AVX-512 will perform horrible in SSE, which is why portable SIMD abstractions should make writing code relative to the native vector width simple.
> I pass in the vector width as a generic parameter like this:
> fn do_simd_stuff<const N: usize>(x: Simd<f32, N>) { ... }
My problem is that no portable_simd example code I've seen does this, which causes people to choose one specific N and run with that.
The second part of the problem is how you find the native vector length, so you can instantiate the generic function. IIRC this isn't even exposed in portable_simd and you have to use a seperate crate to get it.
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.
For many problems, choosing the right instruction or instruction sequence makes a large difference. Portable SIMD abstractions necessarily expose some common semantic layer, but SIMD ISAs don't actually have equivalent capabilities. Instructions like pshufb, for example, enable algorithmic tricks that don't necessarily have an equally efficient analogue on another architecture.
If maximum performance matters, I generally want intrinsics and architecture-specific implementations; if portability matters more, I'd rather move further up the abstraction stack and use something designed to target multiple architectures, such as ISPC. There are certainly cases where portable SIMD gets close enough to optimal, but I don't think there's a compiler or abstraction that can express every useful SIMD idiom and lower it equally efficiently across fundamentally different ISAs.
There are many ways that performance matters without trying to win a F1 race.
Go isn't alone, .NET, Java have similar portable libraries, and C++ is in the process of getting one.
SIMD seems to me, to be very platform specific. Maybe there are times one SIMD unit is not anothers' SIMD unit?
There is no reason a portable_simd relu_dot implemention should need to specify the SIMD width.
But the design and documentation of portable_simd makes the fixed size syntactically easy/the default and the width agnostic code harder.
What should it choose then? I have a Zen 3 processor, and benchmarking some simd I did recently says 32 byte or 64 byte chunks was fastest. But I'm sure I'd get a different result on a different Zen, and different again on Intel's.
How would the library decide what SIMD width I should use?
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.
If that's the case, then the selection logic would be trivial: figure out the full hierarchical ID of the uarch you're running on, then search for the longest prefix match in the table of available impls.
If things work more like you're imagining, though, then I suppose the process-boot impl-selector would narrow down the impl matrix to just the subset that are legal on the running uarch; pick one arbitrarily to be active at first; and then wrap the calls in a handler that gradually re-works the called function in a way reminiscent of a profile-guided JIT, but without the need to actually synthesize any code at runtime — instead, it'd just be a multi-armed bandit passing-through-to and re-ranking competitor impls, with decreasing sampling of the non-first-ranked impls as confidence-in-score-separation increases.
It's definitely a 80% solution where you occasionally need to drop down to intrinsics (at zero runtime perf cost) for CPU specific instructions.
But just having vector types, arithmetic, swizzling, loads and stores will go a long way for basic tasks.
And with generics you can write code that is type and width agnostic. No need to rewrite your code of you want to go from SSE to AVX512, just change from f32x4 to f32x16 (or use generics) and you are done.
Because there are platform vendors. And SIMD performance very much depends on the use-case, which is a balance of practicalities and specifications and intended deployment targets ..
I also think this is a deployment problem, not a build problem, but okay ..
Each family of operations is a trait parameterized by the operation itself:
Call sites name the operation: Operations like Sum, Max, ReduceXor, Inclusive, and Exclusive are all distinct types.As mentioned in the post, execution shape is typed too. A static shuffle takes its control as a type-level constant, and the shuffle mode constrains which controls are expressible:
For an example of errors caught, a warp-scoped executor for a device-scoped barrier is a compile error: Strip mining is typed on the amount of work and the lane capacity, and it hands back one chunk at a time along with the predicate saying which lanes live in that chunk: Hopefully that gives the flavor of it.Can you say more about the application space you're targeting?
Well now you could in theory AI generate SIMD, which will be vibe coded, as those devs have no idea of its correctness.
`core` instead of `std` is great too!
This will become useful in one of my sideproject where I use bitmaps to speed up pathfinding, exited to try it out!
"Pgrust v0.2: Now faster than Postgres" (2026-06) https://news.ycombinator.com/item?id=49111925
At worst, the post reads like a propaganda for VectorWare, but, overall, it reads more like their insights on the matter.
If you have real feedback, great, but it's useless to rip on the hard work of others without it.
Rust is better in so many ways.