4 ms·

This takes me back to some blissful days spent optimizing an integer compression scheme in AVX2. It's really a difficult instruction set to work with. The autho

This takes me back to some blissful days spent optimizing an integer compression scheme in AVX2. It's really a difficult instruction set to work with. The authors comment about a set of oddly shaped Legos is very apt. It often doesn't have instructions that cross lanes and often lacks the thing you need, so you find yourself hacking around it with bit operations and lookup tables. AVX512 is a huge improvement, but it still hasn't caught on in a big way over seven years later.

dragontamer 4y ago
Play with Intel ISPC or CUDA or OpenCL. I think what programmers are missing is the "mental model" that scales well to large problems. Once you have such a model, working your way back down to SIMD-assembly is way easier. --------- The general mental model is "perform work independently. Then perform prefix/scan to synchronize and coordinate between threads", combined with thread-barriers. There are other models of course, but this paradigm works for a surprisingly huge number of problems. EDIT: Compact, the solution from this blogpost, is itself a prefix-scan operation. See http://www.cse.chalmers.se/~uffe/streamcompaction.pdf http://www.cse.chalmers.se/~uffe/streamcompaction.pdf . In case people weren't aware how prefix-sum is underlying so many of today's parallel algorithms in practice. Intel decided to make a specific assembly instruction for stream-compaction, but its just the prefix-sum steps we've known and loved for ages.
- paulmd 4y ago
  100% the mental model I've always used. Few notes: - instead of using the "assign object X to bin Y by placing a pointer to X in binY.objectVector", that operation is equivalent to "sort+prefix-scan". Build an array of [0..N] (thrust counting_iterator is golden) and a parallel "values" array of [binX, ..., binZ], sort the index array using val array as the keyval. - Once you've sorted the bins, the prefix scan finds the "start" of each sub-array and the length is start[x+1] - start[x]. If you have 0 size, that sub-array is empty. - if the data is sparse, you can filter before you launch - so you'd build an array containing the (sorted) indexes of any Y that has data. - all of these concepts become much easier working with "offset from an array" rather than a pointer directly - structure-of-arrays rather than array-of-structs - to keep memory access aligned and maximized within a warp - if you are looking for "each object outputs between 0 and N items as a result", that's basically prefix-scan within your warp/grid. Have the 0th thread do an atomic-CAS/atomic-increment to increment a global counter that offsets the whole warp/grid with the appropriate number of items, then every thread with an item writes to arr[GLOBAL_COUNT + myPrefixSum]. Again, in many cases you will probably write TWO parallel arrays at this time (index of the item, and the actual output value), or more. - Allocate your arrays at program startup and allocate them as big as you think they'll need to be. GPUs don't do memory fragmentation, you'll want big contiguous allocations. - Thrust Framework makes all this super easy with zip-iterators, counting-iterators, sort, prefix-scan, etc. If needed you can get pointers down to the actual arrays and intermingle all this with conventional low-level CUDA __kernel__ or __device__ functions. Check out their quickstart example at this link. https://thrust.github.io/ https://thrust.github.io/ (radix/postal) sorting and prefix-sum is, by far, the most successful way I've seen to port "conventional" logic flows to GPUs, those instructions run VERY fast due to being parallel and aligned. Your code is really mostly just "glue" to build those values for either warp-wide or device-wide sorts/prefix sums - the GPU is performing the heavy lifting of "re-aligning" the data and it does it very efficiently.
  - dragontamer 4y ago
    > - Thrust Framework makes all this super easy with zip-iterators, counting-iterators, sort, prefix-scan, etc. If needed you can get pointers down to the actual arrays and intermingle all this with conventional low-level CUDA __kernel__ or __device__ functions. I find that CUDA's cub library is better if you're doing prefix-sums within a kernel. "Thrust" is more of a quick-and-dirty prototype kind of code, which has weaker performance than people expect. You gotta get down and dirty with the __kernel__ functions. And cub is the library for that (not Thrust). Thrust is great for GPU-prototypes / general scan/reduce prototyping IMO. Probably good enough for a lot of problems, but its a bit slow in practice. Thrust has the mental model down pat, but it just doesn't have enough performance. https://docs.nvidia.com/cuda/cub/index.html https://docs.nvidia.com/cuda/cub/index.html
    - paulmd 4y ago
      > I find that CUDA's cub library is better if you're doing prefix-sums within a kernel. Thrust doesn't have a __device__ prefix-sum iirc, just the global call /laugh > "Thrust" is more of a quick-and-dirty prototype kind of code, which has weaker performance than people expect. Yes, absolutely, they're complimentary and in many cases CUB does things slightly better, or does things that Thrust doesn't support. But Thrust is fantastic for "I want to allocate some GPU arrays, set up some data, and run sort+prefix sum, then hand it off to something else to run the actual algorithm. It's glue that helps you get started (eg see those quickstarts - those are very short even by CUDA standards let alone OpenCL) and figure out if your idea is going to work. And there's very little penalty to keeping the "global steps" inside thrust, eg if you're just doing "fill this index-array with 0..N and then sort(arr1,arr2)" that is not much slower than doing everything raw, or writing one big function that tries to do everything without intermediate computations. It's also easy to get Thrust containers to give you a real pointer and at that point you can call CUB or real kernels or do whatever else you want. As far as performance... eh, CUB is a little faster but not like incredibly much so, maybe 10% or so from what I remember, it wasn't huge. Thrust algorithms are usually not in-place so CUB can provide slightly higher problem size in most situations (since you don't have to allocate a scratch buffer). I actually found the CUB in-place sort was slower than Thrust non-inplace though (understandable, that's a common penalty, and CUB non-inplace might be even faster). More fundamentally, Thrust really works at the level of iterators and not kernels/grids, so you can't really do warp-level operations at all using global "sort this shit" type commands. Thrust doesn't expose the grid information to you and doesn't make guarantees about what grid topology will be executed (there is an OpenMP backend!). But if there is some general "per-item" function in your algorithm, you can call it using the map-iterator (can't remember what it's called but like, pass this object to this function) and either pass the object to work on, or have the value passed be an index of a work-item and your function loads it (store a pointer to the array start in the map-iterator). And in that case you inherit some of the occupancy auto-tuning that Thrust does, which is nice just as a basic thing to get off the ground - it'll try to use as wide a grid as is feasible given the occupancy/utilization. I seem to remember that I did find a way to kinda work around it somehow, like what I was iterating was grid launches instead of work-items, and obviously those can use warp-collective calls etc, but yeah at some point you'll have to make the hop to a proper kernel launch, Thrust just lets you push it off a bit. I was just seeing if I could do it to leverage Thrust's occupancy auto-tuning. Maybe it was that I'd stride the object space (eg launch an iterator for every 32 items) and do a kernel launch on each chunk, or something like that.

dragontamer 4y ago

Play with Intel ISPC or CUDA or OpenCL. I think what programmers are missing is the "mental model" that scales well to large problems. Once you have such a model, working your way back down to SIMD-assembly is way easier. --------- The general mental model is "perform work independently. Then perform prefix/scan to synchronize and coordinate between threads", combined with thread-barriers. There are other models of course, but this paradigm works for a surprisingly huge number of problems. EDIT: Compact, the solution from this blogpost, is itself a prefix-scan operation. See http://www.cse.chalmers.se/~uffe/streamcompaction.pdf http://www.cse.chalmers.se/~uffe/streamcompaction.pdf . In case people weren't aware how prefix-sum is underlying so many of today's parallel algorithms in practice. Intel decided to make a specific assembly instruction for stream-compaction, but its just the prefix-sum steps we've known and loved for ages.

paulmd 4y ago
100% the mental model I've always used. Few notes: - instead of using the "assign object X to bin Y by placing a pointer to X in binY.objectVector", that operation is equivalent to "sort+prefix-scan". Build an array of [0..N] (thrust counting_iterator is golden) and a parallel "values" array of [binX, ..., binZ], sort the index array using val array as the keyval. - Once you've sorted the bins, the prefix scan finds the "start" of each sub-array and the length is start[x+1] - start[x]. If you have 0 size, that sub-array is empty. - if the data is sparse, you can filter before you launch - so you'd build an array containing the (sorted) indexes of any Y that has data. - all of these concepts become much easier working with "offset from an array" rather than a pointer directly - structure-of-arrays rather than array-of-structs - to keep memory access aligned and maximized within a warp - if you are looking for "each object outputs between 0 and N items as a result", that's basically prefix-scan within your warp/grid. Have the 0th thread do an atomic-CAS/atomic-increment to increment a global counter that offsets the whole warp/grid with the appropriate number of items, then every thread with an item writes to arr[GLOBAL_COUNT + myPrefixSum]. Again, in many cases you will probably write TWO parallel arrays at this time (index of the item, and the actual output value), or more. - Allocate your arrays at program startup and allocate them as big as you think they'll need to be. GPUs don't do memory fragmentation, you'll want big contiguous allocations. - Thrust Framework makes all this super easy with zip-iterators, counting-iterators, sort, prefix-scan, etc. If needed you can get pointers down to the actual arrays and intermingle all this with conventional low-level CUDA __kernel__ or __device__ functions. Check out their quickstart example at this link. https://thrust.github.io/ https://thrust.github.io/ (radix/postal) sorting and prefix-sum is, by far, the most successful way I've seen to port "conventional" logic flows to GPUs, those instructions run VERY fast due to being parallel and aligned. Your code is really mostly just "glue" to build those values for either warp-wide or device-wide sorts/prefix sums - the GPU is performing the heavy lifting of "re-aligning" the data and it does it very efficiently.
- dragontamer 4y ago
  > - Thrust Framework makes all this super easy with zip-iterators, counting-iterators, sort, prefix-scan, etc. If needed you can get pointers down to the actual arrays and intermingle all this with conventional low-level CUDA __kernel__ or __device__ functions. I find that CUDA's cub library is better if you're doing prefix-sums within a kernel. "Thrust" is more of a quick-and-dirty prototype kind of code, which has weaker performance than people expect. You gotta get down and dirty with the __kernel__ functions. And cub is the library for that (not Thrust). Thrust is great for GPU-prototypes / general scan/reduce prototyping IMO. Probably good enough for a lot of problems, but its a bit slow in practice. Thrust has the mental model down pat, but it just doesn't have enough performance. https://docs.nvidia.com/cuda/cub/index.html https://docs.nvidia.com/cuda/cub/index.html
  - paulmd 4y ago
    > I find that CUDA's cub library is better if you're doing prefix-sums within a kernel. Thrust doesn't have a __device__ prefix-sum iirc, just the global call /laugh > "Thrust" is more of a quick-and-dirty prototype kind of code, which has weaker performance than people expect. Yes, absolutely, they're complimentary and in many cases CUB does things slightly better, or does things that Thrust doesn't support. But Thrust is fantastic for "I want to allocate some GPU arrays, set up some data, and run sort+prefix sum, then hand it off to something else to run the actual algorithm. It's glue that helps you get started (eg see those quickstarts - those are very short even by CUDA standards let alone OpenCL) and figure out if your idea is going to work. And there's very little penalty to keeping the "global steps" inside thrust, eg if you're just doing "fill this index-array with 0..N and then sort(arr1,arr2)" that is not much slower than doing everything raw, or writing one big function that tries to do everything without intermediate computations. It's also easy to get Thrust containers to give you a real pointer and at that point you can call CUB or real kernels or do whatever else you want. As far as performance... eh, CUB is a little faster but not like incredibly much so, maybe 10% or so from what I remember, it wasn't huge. Thrust algorithms are usually not in-place so CUB can provide slightly higher problem size in most situations (since you don't have to allocate a scratch buffer). I actually found the CUB in-place sort was slower than Thrust non-inplace though (understandable, that's a common penalty, and CUB non-inplace might be even faster). More fundamentally, Thrust really works at the level of iterators and not kernels/grids, so you can't really do warp-level operations at all using global "sort this shit" type commands. Thrust doesn't expose the grid information to you and doesn't make guarantees about what grid topology will be executed (there is an OpenMP backend!). But if there is some general "per-item" function in your algorithm, you can call it using the map-iterator (can't remember what it's called but like, pass this object to this function) and either pass the object to work on, or have the value passed be an index of a work-item and your function loads it (store a pointer to the array start in the map-iterator). And in that case you inherit some of the occupancy auto-tuning that Thrust does, which is nice just as a basic thing to get off the ground - it'll try to use as wide a grid as is feasible given the occupancy/utilization. I seem to remember that I did find a way to kinda work around it somehow, like what I was iterating was grid launches instead of work-items, and obviously those can use warp-collective calls etc, but yeah at some point you'll have to make the hop to a proper kernel launch, Thrust just lets you push it off a bit. I was just seeing if I could do it to leverage Thrust's occupancy auto-tuning. Maybe it was that I'd stride the object space (eg launch an iterator for every 32 items) and do a kernel launch on each chunk, or something like that.