7 ms·

I write a lot of compute kernels in CUDA, and my litmus test is prefix sum, for two reasons. First, you can implement it in pure CUDA C++, and max out the memo

I write a lot of compute kernels in CUDA, and my litmus test is prefix sum, for two reasons. First, you can implement it in pure CUDA C++, and max out the memory bandwidth of any nvidia or AMD GPU. The CUB library provides a state of the art implementation (using decoupled-lookback) that one can compare against new programming languages. Second, it is one of the most basic parallel algorithm building blocks. Many parallel algorithms use it, and many parallel algorithms are "prefix-sum-like". If I am not able to write prefix sum from scratch efficiently in your programming language / library, I can't use it. Every time someone shows a new programming language for compute, the examples provided are super basic (e.g. `map(...).fold(...)`), but I have yet to see a new programming language that can be used to implement the 2-3 most fundamental parallel algorithms from any parallel algorithms graduate course. For example, Futhark provides a prefix sum intrinsic, that just calls CUB - if you want to implement a prefix-sum-like algorithm, you are out of luck. In WGPU, it appears that prefix-sum will be an intrinsic of WHSL, which sounds like you would be out-of-luck too. You mentioned WGPU and Vulkan. Do you know how to implement prefix-sum from scratch on these? If so, do you know how the performance compare against CUB?

raphlinus 7y ago
That sounds like a great project, one I heartily recommend to people wanting to get their feet wet. First, the basic principles of how to write an efficient prefix sum are the same as before. Definitely read [1]. Also see [2] for an exploration into how much subgroups help over threadgroup shared memory. In fact, [2] has much of the answer you seek, as it's already written in Vulkan and GLSL (compiling to SPIR-V), though it does miss a direct performance comparison to CUB. Second, doing this depends on WebGPU exposing threadgroup shared memory and subgroups operations. I'm not sure where the standard and implementations are on this; there are significant challenges. For example, subgroup shuffle is not available in DX12, though subgroup reduce operations such as WavePrefixSum are. So in general an implementation will need to do runtime detection of capabilities and choose a most-optimum kernel based on that. In theory, it should be possible to achieve comparable performance. But it's likely that WebGPU implementations haven't really been tuned for performance yet, while CUB has seen a ton of work. So, as always, the thing to do is implement and measure. [1]: https://developer.nvidia.com/gpugems/gpugems3/part-vi-gpu-computing/chapter-39-parallel-prefix-sum-scan-cuda https://developer.nvidia.com/gpugems/gpugems3/part-vi-gpu-co... [2]: https://cachemiss.xyz/blog/parallel-reduce-and-scan-on-the-GPU https://cachemiss.xyz/blog/parallel-reduce-and-scan-on-the-G...
- fluffything 7y ago
  > First, the basic principles of how to write an efficient prefix sum are the same as before. Definitely read [1]. That's slightly outdated. The state-of-the art is: https://research.nvidia.com/publication/single-pass-parallel-prefix-scan-decoupled-look-back https://research.nvidia.com/publication/single-pass-parallel... (EDIT: well the state of the art is CUB, but that paper is a bit more approachable than the CUB docs, which are still very good). > though it does miss a direct performance comparison to CUB. For perspective, CUB prefix-sum is as fast as a memcpy on the GPU. It achieves the peak memory bandwidth of the device.
  - raphlinus 7y ago
    Great paper. I believe it's possible to do this efficiently in Vulkan + SPIR-V, and also adapt it to WebGPU given the caveats mentioned above, but it's hard to say for sure without doing it. Clearly that paper uses very sophisticated techniques and would not be trivial to implement - it goes considerably past the [2] I linked. I agree 100% that prefix sum is an excellent litmus test for this platform. I'll give one example to fill in more specifics. The fence-free descriptor updates described in section 4.4 of your linked paper depend on 64 bit atomics. These are available using the VK_KHR_shader_atomic_int64 extension, and appear (using vulkan.gpuinfo.org) to be available on NV and AMD hardware but not Intel (and basically not at all on mobile). On Vulkan, you query at runtime for the presence of this extension, and swap in the appropriate kernel based on the result of the query; optimizing for hardware without 64 bit atomics might give you a different approach. Not having to deal with all this is the reason NVidia owns the space :)
  - raphlinus 7y ago
    Followup: I did a very rough prototype on my current Vulkan runner; all the good stuff is in the prefix.comp kernel code. It's not doing the fancy stuff, in particular it doesn't parallelize the lookback code, so that ends up taking a bunch of time. But it does give an example of how to use the Vulkan memory model, for which I've seen very little in the way of published material. If anyone is interested in tuning this and doing a performance comparison, please get in touch. It'd make a good blog post, I think, but already the time I've spent on it is a distraction from the renderer. https://github.com/linebender/piet-gpu/tree/prefix https://github.com/linebender/piet-gpu/tree/prefix
borune 7y ago
I am currently working on implementing the decoupled-lookback in Futhark, as my master thesis. The development is under the branch `onepassscan`. It is still missing some features, like it cannot do map'o'scan, meaning include the map function in the scan optimization. It is likely that the implementation will only be used for Nvidia, since the cache guarentees only are prommised to work for Nvidia. The current version of Futhark is using the reduce-then-scan strategy.
- raphlinus 7y ago
  Awesome! I'll note that these "cache guarantees" in the Vulkan world are the Vulkan 1.2 memory model, and are supported by the latest drivers for AMD, Intel, and NVidia [1]. This recent change is one big reason I'm saying Vulkan + SPIR-V is getting ready for real compute workloads, and this wasn't the case even months ago. [1]: https://vulkan.gpuinfo.org/listdevicescoverage.php?extension=VK_KHR_vulkan_memory_model&platform=windows https://vulkan.gpuinfo.org/listdevicescoverage.php?extension...
- fluffything 7y ago
  Sounds like a great project.
Athas 7y ago
> For example, Futhark provides a prefix sum intrinsic, that just calls CUB - if you want to implement a prefix-sum-like algorithm, you are out of luck. Futhark doesn't call CUB (it also targets OpenCL), and here is an asymptotically work-efficient prefix sum in pure Futhark: https://gist.github.com/athas/611e6b9a76b382ec079f979ec7fb85f1 https://gist.github.com/athas/611e6b9a76b382ec079f979ec7fb85... It's not as fast as a hand-written scan in CUDA, nor as fast as the built-in scan in Futhark, but you can do it. Futhark isn't a low-level or GPU-specific language, so you can't take direct advantage of hardware-specific features. Generally, Futhark doesn't compete performance-wise primitive-for-primitive, but can do application-level transformations that are infeasible for libraries, and impractical to do by hand. For example, you can 'map' a Futhark 'scan' over a two-dimensional array and it will just work pretty well, while there's no way to turn the CUB scan into a regular segmented scan, unless the original author thought of it. Similarly, if you are 'scan'ing the same array twice in a Futhark program, but with different operators or pre-scan transforms, the compiler will fuse the two operations into one, to reduce the number of accesses.
- fluffything 7y ago
  > Futhark doesn't call CUB (it also targets OpenCL), If the builtin scan doesn't call CUB, how fast is a prefix sum in Futhark compared to Cuda using CUB ? > For example, you can 'map' a Futhark 'scan' over a two-dimensional array and it will just work pretty well, What does "pretty well" mean? 95% of peak device throughput ? 60% ? When a compute device costs 8.000$, an efficiency of 90% means I'm throwing 800$ out of the window.
  - Athas 7y ago
    > If the builtin scan doesn't call CUB, how fast is a prefix sum in Futhark compared to Cuda using CUB ? Last I checked, it's about 50% the performance of CUB, but a sibling post describes a current effort to implement the decoupled lookback algorithm in Futhark's CUDA backend. > What does "pretty well" mean? 95% of peak device throughput ? 60% ? It's as fast as Futhark's single-dimensional scan. In exotic cases (when a single segment or "scan instance" fits in a thread block) it can be even faster. > When a compute device costs 8.000$, an efficiency of 90% means I'm throwing 800$ out of the window. Certainly! But people run much less than 90% efficient code on $8000 CPUs all the time, so I think there is definitely room for a high-level programming language that runs on GPUs. It's the usual performance-versus-productivity tradeoff, although Futhark's ambition (and to a large extent, reality) puts more emphasis on the performance side of things than when people usually use that argument. While Futhark can in principle not beat hand-written code, I will note that empirically it is not unusual for us to see performance increases when we take published GPU benchmark code and rewrite it straightforwardly in Futhark. Much GPU code out there is not perfectly written at the level of CUB, and while the Futhark compiler doesn't exactly generate highly inspired code, it gets the basics consistently right.
9q9 7y ago
2-3 most fundamental parallel algorithms Would you be able to list them? I'm asking because I do develop programming languages for parallel hardware. It would be very useful for me to look at your examples and apply them to my PL design ideas. An alternative way I might be rendering my question could be: what primitives do you recommend a language should provide that's currently missing in the languages you evaluate?
- fluffything 7y ago
  > Would you be able to list them? scan (e.g. prefix sum), merge, partition, reduce (e.g. minimum), tree contraction, sort > what primitives do you recommend a language should provide that's currently missing in the languages you evaluate? Hardware primitives and powerful capabilities for defining safe abstractions over those. Chances are that, e.g., I need to implement my own "primitive-like" operation (e.g. my own sort or scan). If your language provides a "partition" primitive, but not the tools to implement my own, I can't use your language. When people create languages for GPUs, for some reason they add the C++ STL or similar as primitives, instead of providing users the tools to write such libraries themselves. The consequence is that those languages end up not being practical to use, and nobody uses them.
  - 9q9 7y ago
    Thanks. How to abstract "hardware primitives" in a way that can be instantiated to many GPU architectures without performance penalty and be useful for higher-level programming is not so clear. How would you, to take an example from the CPU world, fruitfully abstract the CPU's memory model? As far as I'm aware that's not a solved problem in April 2020, and write papers on this subject are still appearing in top conferences .
    - deleted 7y ago
      [deleted]
    - fluffything 7y ago
      There exists a language that can optimally target AMD's ROCm, NVIDIA's PTX, and multicore x86-64 CPUs: CUDA C++. This language has 3 production quality compilers for these targets (nvcc, hip, and pgi). Unfortunately, this language is from 2007, and the state-of-the art has not really been improved since then (except for the evolution of CUDA proper). It would be cool for people to work on improving the state of the art, providing languages that are simpler, perform better, or are more high-level than CUDA, without sacrificing anything (i.e. true improvements). There have been some notable experiments in this regard, e.g., Sequoia for roadrunner scratchpads had very interesting abstraction capabilities over the cache hierarchy, that could have led to a net total improvement over CUDA __shared__ memory. Most newer languages have the right goal of trying to simplify CUDA, but they end up picking trade-offs that only allow them to do so by sacrificing a lot of performance. That's not an interesting proposition for most CUDA developers - the reason they pick up CUDA is performance, and sacrificing ~5% might be acceptable if the productivity gains are there, but a 20-30% perf loss isn't acceptable - too much money involved. One can simplify CUDA while retaining performance by restricting a language to GPUs and a particular domain, e.g., computer graphics / gfx shaders or even stencil codes or map-reduce, etc. However, a lot of the widely advertised languages try to (1) target both GPUs and CPUs, compromising on some minimum common denominator of features, (2) support general-purpose compute kernels, and (3) significantly simplify CUDA, often doing so by just removing the explicit memory transfers, which compromises performance. These languages are quite neat, but they aren't really practical, because they aren't really true improvements over CUDA. Most companies I work with that use CUDA today are using the state of the art (C++17 or 20 extensions), experimental libraries, drivers, etc. So it isn't hard to sell them into a new technology, _if it is better than what they are already using_.

raphlinus 7y ago

That sounds like a great project, one I heartily recommend to people wanting to get their feet wet. First, the basic principles of how to write an efficient prefix sum are the same as before. Definitely read [1]. Also see [2] for an exploration into how much subgroups help over threadgroup shared memory. In fact, [2] has much of the answer you seek, as it's already written in Vulkan and GLSL (compiling to SPIR-V), though it does miss a direct performance comparison to CUB. Second, doing this depends on WebGPU exposing threadgroup shared memory and subgroups operations. I'm not sure where the standard and implementations are on this; there are significant challenges. For example, subgroup shuffle is not available in DX12, though subgroup reduce operations such as WavePrefixSum are. So in general an implementation will need to do runtime detection of capabilities and choose a most-optimum kernel based on that. In theory, it should be possible to achieve comparable performance. But it's likely that WebGPU implementations haven't really been tuned for performance yet, while CUB has seen a ton of work. So, as always, the thing to do is implement and measure. [1]: https://developer.nvidia.com/gpugems/gpugems3/part-vi-gpu-computing/chapter-39-parallel-prefix-sum-scan-cuda https://developer.nvidia.com/gpugems/gpugems3/part-vi-gpu-co... [2]: https://cachemiss.xyz/blog/parallel-reduce-and-scan-on-the-GPU https://cachemiss.xyz/blog/parallel-reduce-and-scan-on-the-G...

fluffything 7y ago
> First, the basic principles of how to write an efficient prefix sum are the same as before. Definitely read [1]. That's slightly outdated. The state-of-the art is: https://research.nvidia.com/publication/single-pass-parallel-prefix-scan-decoupled-look-back https://research.nvidia.com/publication/single-pass-parallel... (EDIT: well the state of the art is CUB, but that paper is a bit more approachable than the CUB docs, which are still very good). > though it does miss a direct performance comparison to CUB. For perspective, CUB prefix-sum is as fast as a memcpy on the GPU. It achieves the peak memory bandwidth of the device.
- raphlinus 7y ago
  Great paper. I believe it's possible to do this efficiently in Vulkan + SPIR-V, and also adapt it to WebGPU given the caveats mentioned above, but it's hard to say for sure without doing it. Clearly that paper uses very sophisticated techniques and would not be trivial to implement - it goes considerably past the [2] I linked. I agree 100% that prefix sum is an excellent litmus test for this platform. I'll give one example to fill in more specifics. The fence-free descriptor updates described in section 4.4 of your linked paper depend on 64 bit atomics. These are available using the VK_KHR_shader_atomic_int64 extension, and appear (using vulkan.gpuinfo.org) to be available on NV and AMD hardware but not Intel (and basically not at all on mobile). On Vulkan, you query at runtime for the presence of this extension, and swap in the appropriate kernel based on the result of the query; optimizing for hardware without 64 bit atomics might give you a different approach. Not having to deal with all this is the reason NVidia owns the space :)
- raphlinus 7y ago
  Followup: I did a very rough prototype on my current Vulkan runner; all the good stuff is in the prefix.comp kernel code. It's not doing the fancy stuff, in particular it doesn't parallelize the lookback code, so that ends up taking a bunch of time. But it does give an example of how to use the Vulkan memory model, for which I've seen very little in the way of published material. If anyone is interested in tuning this and doing a performance comparison, please get in touch. It'd make a good blog post, I think, but already the time I've spent on it is a distraction from the renderer. https://github.com/linebender/piet-gpu/tree/prefix https://github.com/linebender/piet-gpu/tree/prefix

borune 7y ago

I am currently working on implementing the decoupled-lookback in Futhark, as my master thesis. The development is under the branch `onepassscan`. It is still missing some features, like it cannot do map'o'scan, meaning include the map function in the scan optimization. It is likely that the implementation will only be used for Nvidia, since the cache guarentees only are prommised to work for Nvidia. The current version of Futhark is using the reduce-then-scan strategy.

raphlinus 7y ago
Awesome! I'll note that these "cache guarantees" in the Vulkan world are the Vulkan 1.2 memory model, and are supported by the latest drivers for AMD, Intel, and NVidia [1]. This recent change is one big reason I'm saying Vulkan + SPIR-V is getting ready for real compute workloads, and this wasn't the case even months ago. [1]: https://vulkan.gpuinfo.org/listdevicescoverage.php?extension=VK_KHR_vulkan_memory_model&platform=windows https://vulkan.gpuinfo.org/listdevicescoverage.php?extension...
fluffything 7y ago
Sounds like a great project.

Athas 7y ago

> For example, Futhark provides a prefix sum intrinsic, that just calls CUB - if you want to implement a prefix-sum-like algorithm, you are out of luck. Futhark doesn't call CUB (it also targets OpenCL), and here is an asymptotically work-efficient prefix sum in pure Futhark: https://gist.github.com/athas/611e6b9a76b382ec079f979ec7fb85f1 https://gist.github.com/athas/611e6b9a76b382ec079f979ec7fb85... It's not as fast as a hand-written scan in CUDA, nor as fast as the built-in scan in Futhark, but you can do it. Futhark isn't a low-level or GPU-specific language, so you can't take direct advantage of hardware-specific features. Generally, Futhark doesn't compete performance-wise primitive-for-primitive, but can do application-level transformations that are infeasible for libraries, and impractical to do by hand. For example, you can 'map' a Futhark 'scan' over a two-dimensional array and it will just work pretty well, while there's no way to turn the CUB scan into a regular segmented scan, unless the original author thought of it. Similarly, if you are 'scan'ing the same array twice in a Futhark program, but with different operators or pre-scan transforms, the compiler will fuse the two operations into one, to reduce the number of accesses.

fluffything 7y ago
> Futhark doesn't call CUB (it also targets OpenCL), If the builtin scan doesn't call CUB, how fast is a prefix sum in Futhark compared to Cuda using CUB ? > For example, you can 'map' a Futhark 'scan' over a two-dimensional array and it will just work pretty well, What does "pretty well" mean? 95% of peak device throughput ? 60% ? When a compute device costs 8.000$, an efficiency of 90% means I'm throwing 800$ out of the window.
- Athas 7y ago
  > If the builtin scan doesn't call CUB, how fast is a prefix sum in Futhark compared to Cuda using CUB ? Last I checked, it's about 50% the performance of CUB, but a sibling post describes a current effort to implement the decoupled lookback algorithm in Futhark's CUDA backend. > What does "pretty well" mean? 95% of peak device throughput ? 60% ? It's as fast as Futhark's single-dimensional scan. In exotic cases (when a single segment or "scan instance" fits in a thread block) it can be even faster. > When a compute device costs 8.000$, an efficiency of 90% means I'm throwing 800$ out of the window. Certainly! But people run much less than 90% efficient code on $8000 CPUs all the time, so I think there is definitely room for a high-level programming language that runs on GPUs. It's the usual performance-versus-productivity tradeoff, although Futhark's ambition (and to a large extent, reality) puts more emphasis on the performance side of things than when people usually use that argument. While Futhark can in principle not beat hand-written code, I will note that empirically it is not unusual for us to see performance increases when we take published GPU benchmark code and rewrite it straightforwardly in Futhark. Much GPU code out there is not perfectly written at the level of CUB, and while the Futhark compiler doesn't exactly generate highly inspired code, it gets the basics consistently right.

9q9 7y ago

2-3 most fundamental parallel algorithms Would you be able to list them? I'm asking because I do develop programming languages for parallel hardware. It would be very useful for me to look at your examples and apply them to my PL design ideas. An alternative way I might be rendering my question could be: what primitives do you recommend a language should provide that's currently missing in the languages you evaluate?

fluffything 7y ago
> Would you be able to list them? scan (e.g. prefix sum), merge, partition, reduce (e.g. minimum), tree contraction, sort > what primitives do you recommend a language should provide that's currently missing in the languages you evaluate? Hardware primitives and powerful capabilities for defining safe abstractions over those. Chances are that, e.g., I need to implement my own "primitive-like" operation (e.g. my own sort or scan). If your language provides a "partition" primitive, but not the tools to implement my own, I can't use your language. When people create languages for GPUs, for some reason they add the C++ STL or similar as primitives, instead of providing users the tools to write such libraries themselves. The consequence is that those languages end up not being practical to use, and nobody uses them.
- 9q9 7y ago
  Thanks. How to abstract "hardware primitives" in a way that can be instantiated to many GPU architectures without performance penalty and be useful for higher-level programming is not so clear. How would you, to take an example from the CPU world, fruitfully abstract the CPU's memory model? As far as I'm aware that's not a solved problem in April 2020, and write papers on this subject are still appearing in top conferences .
  - deleted 7y ago
    [deleted]
  - fluffything 7y ago
    There exists a language that can optimally target AMD's ROCm, NVIDIA's PTX, and multicore x86-64 CPUs: CUDA C++. This language has 3 production quality compilers for these targets (nvcc, hip, and pgi). Unfortunately, this language is from 2007, and the state-of-the art has not really been improved since then (except for the evolution of CUDA proper). It would be cool for people to work on improving the state of the art, providing languages that are simpler, perform better, or are more high-level than CUDA, without sacrificing anything (i.e. true improvements). There have been some notable experiments in this regard, e.g., Sequoia for roadrunner scratchpads had very interesting abstraction capabilities over the cache hierarchy, that could have led to a net total improvement over CUDA __shared__ memory. Most newer languages have the right goal of trying to simplify CUDA, but they end up picking trade-offs that only allow them to do so by sacrificing a lot of performance. That's not an interesting proposition for most CUDA developers - the reason they pick up CUDA is performance, and sacrificing ~5% might be acceptable if the productivity gains are there, but a 20-30% perf loss isn't acceptable - too much money involved. One can simplify CUDA while retaining performance by restricting a language to GPUs and a particular domain, e.g., computer graphics / gfx shaders or even stencil codes or map-reduce, etc. However, a lot of the widely advertised languages try to (1) target both GPUs and CPUs, compromising on some minimum common denominator of features, (2) support general-purpose compute kernels, and (3) significantly simplify CUDA, often doing so by just removing the explicit memory transfers, which compromises performance. These languages are quite neat, but they aren't really practical, because they aren't really true improvements over CUDA. Most companies I work with that use CUDA today are using the state of the art (C++17 or 20 extensions), experimental libraries, drivers, etc. So it isn't hard to sell them into a new technology, _if it is better than what they are already using_.