Hacker News Board Hacker News, browsed like a classic message board
Live data from Hacker News

Rust SIMD on the GPU

vectorware.com

Post reply on HN
Print view View original

101–110 of 125 posts

Re: Rust SIMD on the GPU

#101

Post by amdivia »

mathisfun123 wrote: post #85
MindSpunk wrote: post #74

Earlier quoted context omitted.

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 li…

> GPUs aren't really SIMD, they're SIMT (single instruction multiple thread) 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 perform…

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]

1: https://docs.nvidia.com/cuda/cuda-programming-guide/03-advan...

Re: Rust SIMD on the GPU

#102

Post by exDM69 »

camel-cdr wrote: post #98
exDM69 wrote: post #96

Earlier quoted context omitted.

> 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.

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…

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:

    fn do_simd_stuff(x: Simd) { x.mul_add(x+x, x*x); }
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.

Re: Rust SIMD on the GPU

#104

Post by camel-cdr »

exDM69 wrote: post #102
camel-cdr wrote: post #98

Earlier quoted context omitted.

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…

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: fn do_simd_stuff (x: Simd ) { x.mul_add(x+x, x*x); } With this…

> 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:

> fn do_simd_stuff(x: Simd) { ... }

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.

Re: Rust SIMD on the GPU

#105

Post by stymaar »

LoganDark wrote: post #100
stymaar wrote: post #99

Earlier quoted context omitted.

> 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…

> 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 st…

> 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.

Re: Rust SIMD on the GPU

#106

Post by exDM69 »

camel-cdr wrote: post #104
exDM69 wrote: post #102

Earlier quoted context omitted.

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: fn do_simd_stuff (x: Simd ) { x.mul_add(x+x, x*x); } With this…

> 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: > fn do_simd_stuff (x: Simd ) { ... } My problem is that no portable_simd example code I'…

> 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.

Re: Rust SIMD on the GPU

#107

Post by penteract »

mathisfun123 wrote: post #85
MindSpunk wrote: post #74

Earlier quoted context omitted.

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 li…

> GPUs aren't really SIMD, they're SIMT (single instruction multiple thread) 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 perform…

>> GPUs aren't really SIMD, they're SIMT (single instruction multiple thread)

> 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...

Re: Rust SIMD on the GPU

#108

Post by Tuna-Fish »

dataflow wrote: post #84
corysama wrote: post #72

Earlier quoted context omitted.

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 th…

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?

> 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.

Re: Rust SIMD on the GPU

#109

Post by Jhsto »

y1n0 wrote: post #71
monocasa wrote: post #43

Earlier quoted context omitted.

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.

64 what? Bits/bytes/something bigger?

SPIR-V states its an int, float, vector n (where 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.
Post reply on HN