3 ms·

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 b

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.

dragontamer 3y ago
> 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. AMD has similar instructions: https://gpuopen.com/learn/amd-gcn-assembly-cross-lane-operations/ https://gpuopen.com/learn/amd-gcn-assembly-cross-lane-operat... 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.
- Aardwolf 3y ago
  > And since GPUs have single-cycle "popcnt", any symmetric bit-twiddling function can be built off of popcnt within a few cycles This sounds interesting, do you have a reference to this?
  - dragontamer 3y ago
    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) AND(bits) == (popcnt(bits) == 8). OR(bits) == (popcnt(bits) > 0) Etc. etc. EDIT: Looks like Wikipedia has a link to this concept: https://en.wikipedia.org/wiki/Symmetric_Boolean_function https://en.wikipedia.org/wiki/Symmetric_Boolean_function
- touisteur 3y ago

dragontamer 3y ago

> 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. AMD has similar instructions: https://gpuopen.com/learn/amd-gcn-assembly-cross-lane-operations/ https://gpuopen.com/learn/amd-gcn-assembly-cross-lane-operat... 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.

Aardwolf 3y ago
> And since GPUs have single-cycle "popcnt", any symmetric bit-twiddling function can be built off of popcnt within a few cycles This sounds interesting, do you have a reference to this?
- dragontamer 3y ago
  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) AND(bits) == (popcnt(bits) == 8). OR(bits) == (popcnt(bits) > 0) Etc. etc. EDIT: Looks like Wikipedia has a link to this concept: https://en.wikipedia.org/wiki/Symmetric_Boolean_function https://en.wikipedia.org/wiki/Symmetric_Boolean_function
touisteur 3y ago