DPC++/SYCL might be more interesting in this area (GPU code looking like evolved C++).
ISPC is an interesting take on 'building a vector DSL which can easily be integrated in a C++ build-chain'.
Though calling GPUs SIMD nowadays is kind of reductive, since you also have to contain with a very constrained memory hierarchy, and you don't write or think your code as SIMD much but more. Also they don't give access to most bit-tricks, and the rigmarole of shuffles one would expect from a SIMD processor.
> Though calling GPUs SIMD nowadays is kind of reductive, since you also have to contain with a very constrained memory hierarchy, and you don't write or think your code as SIMD much but more. Also they don't give access to most bit-tricks,
Oh rly?
* Popcnt is defined in ROCm and CUDA.
* Shifts, XORs, AND, OR, NOT are all defined in GPUs.
* You can actually do AES and a lot of other bit-level twiddling in GPU space very efficiently. (aka: see any mining software for highly-efficient bit-twiddling).
Yeah, not as many tricks (ex: pext / pdep) as a CPU. But popcnt is the "big" one. GPUs also have bit-reverse instructions (which reduces the need for LS1B instructions like CPUs because you can do LS1B tricks on "both ends" by just bit-reversing).
But honestly, GPUs are so parallel and Register-space is so plentiful that you can probably just build whatever you want out of AND/OR/NOT/XOR instructions alone.
And since GPUs have single-cycle "popcnt", any symmetric bit-twiddling function can be built off of popcnt within a few cycles.
> and the rigmarole of shuffles one would expect from a SIMD processor.
Abuse of __shared__ memory and the full crossbar of a GPU-core means that you can arbitrarily shuffle data through __shared__ memory in like 5 clock ticks (well... in certain ways... if you do it wrong it'd take many clock ticks).
GPUs are actually far superior to AVX512 in this front. NVidia's shfl.bfly instruction can implement FFTs, even-odd sorting, bitonic sorts, and more as high-speed communications.
Trust me on this case: GPU data-movement is far, far, far superior to CPU data-movement.
Intel does NOT have the bpermute or permute instructions like a GPU does. Meanwhile, both AMD and NVidia GPUs have bpermute / permute for both gather-and-scatter-like operation / distribution of data across lanes.
AMD/NVidia also have guarantees upon the speed of broadcasts across __shared__ memory. And butterfly-shuffles can theoretically implement any arbitrary shuffle you want in log2(SIMD-Width) steps, so we have incredible amounts of tools in GPU space that CPU-programmers have no idea about.
Thanks for the enlightenment here, some of those I didn't know. Being used to explicit vectors, going from avx512 proficiency, pdep/pext, gfni rabbit holes, I kind of forgot all of that when going to cuda, trying to avoid the trap of 'doing C in Rust (or Ada, hopefully the meaning is clear)', jumped to new idioms and I must say most of those you cited never appeared in most high performance code I've read, and some I only saw perusing ptx and lower level compiled code, which... I was never sure nvidia would maintain over time. Seems for us it's cub, barriers, atomics and ballots.
It seems I have lots of reading to do and lots of ways to improve my sorting networks / counting sort implementations.
> I was never sure nvidia would maintain over time.
PTX is maintained over time. Its a high-level assembly so to speak, the full details of the machine remain abstracted so that code can be more portable.
SASS is not. SASS changes from architecture-to-architecture. SASS is the actual machine code of NVidia cards. There's an overall understanding of SASS in the GPU world but its not really documented and you "shouldn't" want to learn about it.
--------
I should note that Intel's "pshufb" instruction is very similar to the permute instruction in NVidia/AMD. So yeah, there's a high-speed generic shuffle that's key to Intel/AMD AVX512 code.
But having the backwards-direction (bpermute) available too, as well as __shared__ memory for all other cases is great.
I saw the proof in some Binary-decision-diagram book that I've forgotten the name of.
The gist is that if you have a symmetric function (ie: order doesn't matter) of say, 8-bits, then that means f(10101010) == f(11110000) == f(00001111), and all such combinations. Because this is the very definition of "order doesn't matter".
This necessarily implies that f(bits) can be rewritten into the form of f(bits) == g(popcnt(bits)).
Something to do with counting up all the possibilities of "order doesn't matter" and then pigeon-holing them into popcnt combinations or something really mathematical like that. I'm sorry I don't remember the proof, but hopefully this is enough to give you the gist of the idea.
------------------
For example: XOR is the simplest symmetric function. We can see that XOR can be rewritten as XOR(bits) == g(popcnt(bits)). Where G is "take the bottom-bit from popcnt".
A lot of the "power" of BDDs is their ability to uncannily decompose functions into symmetric and non-symmetric parts. Not consistently mind you, but if a BDD is well-behaved (low-memory space, high-speeds, etc. etc.), its likely because a large part of the calculations happened to be symmetric.
This can be beneficial with say addFourDWORDs(a, b, c, d), where a+b+c+d has a "partial" level of symmetry. (a_0 XOR b_0 XOR c_0 XOR d_0 determines output_0 bit... while the 1st bit is related to XOR (a1,b1,c1,d1), etc. etc.). So there's all kinds of 'hidden popcounts" that could pop up in practice for random functions (IE: "partially symmetric"), though its a puzzle on how to exactly decompose arbitrary bit-functions into this form.
And even if you do, its no guarantee that the popcnt form was actually faster either. But it does give you some "trick" to try when trying to optimize an arbitrary bitwise function into a high-speed routine.
-----------------
Other symmetric functions (assuming 8-bit numbers)
Or ROCm (basically CUDA but for AMD).
I always was a fan of Microsoft's C++AMP though. I thought that was easiest to get into. Too bad it never stuck though.