5 ms·

Hey I was investigating the use of GPUs just out of curiosity since it's one of the few areas of computing I wasn't very familiar with. The above benchmarks a

Hey I was investigating the use of GPUs just out of curiosity since it's one of the few areas of computing I wasn't very familiar with. The above benchmarks are about matrix multiplication. Matrix multiplication seems to be a common theme in gpu api documentation. Now for someone who rarely needs to multiply matrices, are there other good applications the GPU pipeline tends to be useful for? Are there guidelines on how complex a program you can run on GPU cores and still reap benefit? What are (vague, order of magnitude) limits on program size and size of the different memories you have access to? Can you run machine learning algorithms that are not based on matrix multiplication, gradients or derivatives. Can you do monte carlo simulations, decision trees etc?

krona 5y ago
Nvidia have an entire library of books dedicated to GPGPU, which can now be (to some extent) repurposed for WebGPU in ways not possible with WebGL. The point of the matrix multiplication example is that it's easy to implement in WebGL without having to be creative; it's a largely apples-to-apples comparison between WebGL/WebGPU/CPU.
bsenftner 5y ago
The reason matrix multiplication is the focus is because linear algebra is generic, and creates a pipeline for whatever that generic computation requiring task happens to be. The obvious case being things like rendering imagery and transforming point clouds - which if you think about it are completely generic operations.
dragontamer 5y ago
> Now for someone who rarely needs to multiply matrices, are there other good applications the GPU pipeline tends to be useful for? The study of GPU algorithms is completely different and independent of regular CPU algorithms. Sorting: Parallel Bitonic Sort networks. Radix Sort. Mergepath. Hashing: Parallel Cuckoo Hash. Graphs: Strangely similar to CPUs. Depth-first search to save space, breadth-first to generate parallelism. You'll need the parallel SIMD-stack / SIMD-queue with prefix-sums to efficiently push/pull items off the stack/queue, but its not too difficult actually (but non-obvious to the CPU-programmer). Trees: BVH traversal (aka: Raytracing) is extremely common, showing that high-throughput / parallel tree traversal is possible on GPUs. -------- The problem is that GPUs are designed for 10,000 "threads" (really, SIMD-lanes). In fact, I can almost certainly say that any program with less than 1000 threads will run faster on a CPU than a GPU. (unless you're a rare, purely memory bandwidth-bound problem, like memcpy or memset, because GPU-RAM is way faster than standard DDR4) There are all sorts of techniques to discover more "threads" (or parallel streams of computation) through the usage of prefix sum. But ultimately, a "practical" GPU application must perform tens-of-thousands of calculations in parallel to beat a CPU. Indeed: something like parallel-Bitonic Sort is less efficient technically. However, because it creates so many parallel threads of compute, its a great fit as a simple GPU-sorting algorithm. (GPU Mergepath sort is within log2(processors) amount of total work as the sequential version, and is therefore considered almost as efficient as the original sequential algorithm) That's the thing: you end up generating slightly more work with any parallel program than the original sequential algorithm in most cases. Matrix Multiplication is "great" because the parallel version has exactly the same number of additions and multiplications at the original sequential version. So everyone loves it as a simple parallelism example. Its even stronger: the GPU-algorithm not only has the same number of computations, but also the same number of memory reads/writes. Its the most ideal algorithm to run in parallel. Other cases, such as sorting, graph algorithms or whatever, will end up using slightly more overall compute (so there's a tradeoff), or traversing in a different order (CPUs favor depth first search. GPUs will go far more breadth-first before their tens-of-thousands of threads fill up, so you end up traversing the graph / tree in a different order and getting slightly different results) As such, other algorithms require the programmer to make a judgement: is it worth gaining log2(processors) amount of total work, in order to use a GPU and perform the calculation in parallel? Almost certainly. But other algorithms are O(processors) or worse, at which point a purely sequential (or a low-thread count: like 32-threads) would be better than 10,000 threads. GPU lanes are much slower: they access memory at far higher latenciess, and GPUs are in-order processors (lacking branch prediction, out-of-order or superscalar tricks from the CPU world). As such, GPU threads are much much slower individually than CPU threads. You just get a ton of GPU "threads" to make up the difference.
- spekcular 5y ago
  Thanks for the very detailed and helpful reply. Do you know of any good books or references where I could learn more about these things? The books that are usually recommended seems to be CUDA-centric and out of date. I'm interested in learning the more general concepts you talk about in your answer, so that I can effectively write e.g. monte carlo simulations on arbitrary GPU hardware. (I don't have an Nvidia GPU!)
  - dragontamer 5y ago
    > The books that are usually recommended seems to be CUDA-centric and out of date. The CUDA-ones are usually best, because they're at least with a modern API. The other recommendations I got are the 80s stuff on vector computers and the CM-2. It turns out that a huge community of high-performance programmers have experimented with all of these concepts in the 1970s, 80s, 90s, and 00s, long before GPUs. All of their concepts still work today on modern GPUs. Looking up the right keywords: such as "CREW-PRAM algorithms" (Concurrent-Read Exclusive Write, Parallel RAM model) immediately gives you plenty of results for some things. (Ex: I just searched on 'CREW-PRAM DFS' and got: https://core.ac.uk/download/pdf/82490222.pdf https://core.ac.uk/download/pdf/82490222.pdf). The key is understanding what the "old word" for GPU was. That's PRAM, the Parallel-RAM model. That's how programmers from the 1970s, 1980s, and 1990s talked about algorithms written for the GPU-style called it back then. Newer articles / books talk about GPUs directly. -------------- I'd say the fundamentals are covered in the 1970s, such as "A Parallel Algorithm for the Efficient Solution of a General Class of Recurrence Equations" by Kogge, which leads to the development of Prefix sum (as it is called today). Of course, CUDA code is clearer than a theoretical CREW-PRAM discussion. So perhaps its easier if you just read GPU Compute Gems by NVidia to cover the same material. Still, I find that the 1970s writing is often times better written for a beginner (back then, fewer programmers knew how parallel programming worked, but were clearly more math-heavy than today. So I find that reading old articles like that one helps my personal brain and personal thinking pattern) --------- Ah right, the recommendation. I'd say "Data Parallel Algorithms" by Hillis / Steele (ACM Communications 1986) is an excellent introduction to general-purpose SIMD compute. It was written for CM-2, an ancient supercomputer that no longer exists, but the PRAM style applies to GPU algorithms today. Its like 15 pages, but it really opens your eyes to the possibilities. A lot of CUDA stuff is very specific to NVidia GPUs (important specific details: like bank conflicts and shared memory, which you absolutely should learn about... but such details should be studied after you learned the way of parallel-thinking / PRAM model / etc. etc.)
dahart 5y ago
Program size limits are a thing of the past. Today you don’t need to worry about program size, and (like on the CPU) you’ll probably bump into long compile times way before running out of space. Memory access is whatever RAM is on the board, plus whatever you want to stream/transfer from CPU RAM and/or disk. Complexity doesn’t matter, your GPU programs can be arbitrarily complex. what matters is that your workload is fairly uniform - GPUs are SIMT machines that will realize performance gains any time you have lots of threads (thousands, millions, billions) that all mostly do the same thing and can execute mostly independently. Yes, you can do Monte Carlo simulations, decision trees, and ML without Neural Networks. All of these things are extremely common.
Const-me 5y ago
> are there other good applications the GPU pipeline tends to be useful for? Pretty much everything compute-bound. SpaceX presentation about rocket engines: https://www.youtube.com/watch?v=vYA0f6R5KAI https://www.youtube.com/watch?v=vYA0f6R5KAI Unreal Engine 5 presentation about rendering many millions of small triangles using compute shaders: https://www.youtube.com/watch?v=TMorJX3Nj6U https://www.youtube.com/watch?v=TMorJX3Nj6U > Are there guidelines on how complex a program you can run on GPU cores and still reap benefit? Arbitrarily complex. Many real-world GPGPU programs are split over multiple compute shaders / kernels. When implemented properly, they even run in parallel. This is critically important on Windows with DirectCompute, because 2 seconds TDR timeout. When a single compute shader takes more than 2 seconds, the OS concludes the GPU’s hang, resets the hardware and reloads the driver. > What are (vague, order of magnitude) limits on program size and size of the different memories you have access to? Technically unlimited because you can stream data from system memory, disk, or even internets. Practically, most modern GPUs come with 6-12 GB VRAM, for optimal performance you’d want your data to fit there. When doing micro-optimizations, another important number is amount of on-chip memory, the order of magnitude is 64kb per core. > Can you do monte carlo simulations, decision trees etc? Monte Carlo is a great fit for GPUs, just don’t forget to seed your RNG with SV_DispatchThreadID or an equivalent. Trees are tricky. GPUs share instruction decoder and instruction pointer over 32-64 hardware threads. A straightforward implementation of binary trees gonna be suboptimal due to divergence of these threads. Often possible to do something else instead.
- dragontamer 5y ago
  > Trees are tricky. GPUs share instruction decoder and instruction pointer over 32-64 hardware threads. A straightforward implementation of binary trees gonna be suboptimal due to divergence of these threads. Often possible to do something else instead. The GPU-straightforward implementation is to use a SIMD-stack / SIMD-queue. I don't know its proper term, but its "obvious" to anyone who programs GPUs. The following is probably wrong, but hopefully correct enough to demonstrate the idea... active-lane-workgroup(){ wavefront_activelane = prefix-sum(execution-mask); // Execution mask is 0 if not currently executing, and 1 if currently executing in your wavefront. // wavefront_activelane is available at assembly level in one clock tick in both NVidia and AMD GPUs. __shared__ int workgroup_prefix_sum[wavefronts_in_workgroup]; if(wavefront_activelane == 0){ workgroup_prefix_sum[wavefront_id] = wavefront_horizontal_max(wavefront_activelane); } __syncthreads(); return workgroup_prefix_sum[my_wavefronts_idx] + wavefront_activelane; } SIMD-push(stack, data): stack[stack.ptr + active-lane-workgroup()] = data; __syncthreads(); if(active-lane-workgroup() == 0){ stack.ptr += workgroup-prefix-max(active-lane-workgroup()); } __syncthreads(); SIMD-pop(stack): toReturn = stack[stack.ptr - active-lane-workgroup()]; __syncthreads(); if(active-lane-workgroup() == 0){ stack.ptr += workgroup-prefix-max(active-lane-workgroup()); } __syncthreads(); return toReturn; ------ Now that you have a stack, its simply a matter of pushing your DFS into the SIMD-stack, and popping it off to traverse. You're somewhat BFS, because on a 1024-wide workgroup, your threads will visit the top 1024 items of the stack each step. But the stack overall behaves in a DFS manner. The CUDA cub library (and ROCm hipCUB / ROCm rocPRIM libraries) implement these horizontal operations (glorified prefix-sum). But its not too hard to write workgroup prefix-sum or workgroup prefix-max yourself. (Indeed: I suggest beginners write their own prefix-sum to get a feel of how simple and efficient the prefix sum / scan operations can be) -------- The "pattern" is that you use horizontal operations and __syncthreads() for thread synchronization and communication. In effect, this SIMD-stack is performing load-balancing simultaneously with the DFS. That is to say, it "reaches breadth-wise" and pulls extra nodes to visit when more lanes are available, while it prefers depth to minimize memory usage.
alfalfasprout 5y ago
tl;dr problems that are compute bound but trivially parallelizable tend to be good choices for GPU computation. Eg; when you're running a series of small computations on a lot of data. Or a lot of small computations on a moderate amount of data. GPUs tend to be pretty ineffective for computations where there are a lot of data dependencies since the individual shader units are slow compared to traditional CPU cores... so most of the GPU will be idle. Or for ultra low latency applications (since moving data between to and from the GPU is costly even with direct memory access).

krona 5y ago

Nvidia have an entire library of books dedicated to GPGPU, which can now be (to some extent) repurposed for WebGPU in ways not possible with WebGL. The point of the matrix multiplication example is that it's easy to implement in WebGL without having to be creative; it's a largely apples-to-apples comparison between WebGL/WebGPU/CPU.

bsenftner 5y ago

The reason matrix multiplication is the focus is because linear algebra is generic, and creates a pipeline for whatever that generic computation requiring task happens to be. The obvious case being things like rendering imagery and transforming point clouds - which if you think about it are completely generic operations.

dragontamer 5y ago

> Now for someone who rarely needs to multiply matrices, are there other good applications the GPU pipeline tends to be useful for? The study of GPU algorithms is completely different and independent of regular CPU algorithms. Sorting: Parallel Bitonic Sort networks. Radix Sort. Mergepath. Hashing: Parallel Cuckoo Hash. Graphs: Strangely similar to CPUs. Depth-first search to save space, breadth-first to generate parallelism. You'll need the parallel SIMD-stack / SIMD-queue with prefix-sums to efficiently push/pull items off the stack/queue, but its not too difficult actually (but non-obvious to the CPU-programmer). Trees: BVH traversal (aka: Raytracing) is extremely common, showing that high-throughput / parallel tree traversal is possible on GPUs. -------- The problem is that GPUs are designed for 10,000 "threads" (really, SIMD-lanes). In fact, I can almost certainly say that any program with less than 1000 threads will run faster on a CPU than a GPU. (unless you're a rare, purely memory bandwidth-bound problem, like memcpy or memset, because GPU-RAM is way faster than standard DDR4) There are all sorts of techniques to discover more "threads" (or parallel streams of computation) through the usage of prefix sum. But ultimately, a "practical" GPU application must perform tens-of-thousands of calculations in parallel to beat a CPU. Indeed: something like parallel-Bitonic Sort is less efficient technically. However, because it creates so many parallel threads of compute, its a great fit as a simple GPU-sorting algorithm. (GPU Mergepath sort is within log2(processors) amount of total work as the sequential version, and is therefore considered almost as efficient as the original sequential algorithm) That's the thing: you end up generating slightly more work with any parallel program than the original sequential algorithm in most cases. Matrix Multiplication is "great" because the parallel version has exactly the same number of additions and multiplications at the original sequential version. So everyone loves it as a simple parallelism example. Its even stronger: the GPU-algorithm not only has the same number of computations, but also the same number of memory reads/writes. Its the most ideal algorithm to run in parallel. Other cases, such as sorting, graph algorithms or whatever, will end up using slightly more overall compute (so there's a tradeoff), or traversing in a different order (CPUs favor depth first search. GPUs will go far more breadth-first before their tens-of-thousands of threads fill up, so you end up traversing the graph / tree in a different order and getting slightly different results) As such, other algorithms require the programmer to make a judgement: is it worth gaining log2(processors) amount of total work, in order to use a GPU and perform the calculation in parallel? Almost certainly. But other algorithms are O(processors) or worse, at which point a purely sequential (or a low-thread count: like 32-threads) would be better than 10,000 threads. GPU lanes are much slower: they access memory at far higher latenciess, and GPUs are in-order processors (lacking branch prediction, out-of-order or superscalar tricks from the CPU world). As such, GPU threads are much much slower individually than CPU threads. You just get a ton of GPU "threads" to make up the difference.

spekcular 5y ago
Thanks for the very detailed and helpful reply. Do you know of any good books or references where I could learn more about these things? The books that are usually recommended seems to be CUDA-centric and out of date. I'm interested in learning the more general concepts you talk about in your answer, so that I can effectively write e.g. monte carlo simulations on arbitrary GPU hardware. (I don't have an Nvidia GPU!)
- dragontamer 5y ago
  > The books that are usually recommended seems to be CUDA-centric and out of date. The CUDA-ones are usually best, because they're at least with a modern API. The other recommendations I got are the 80s stuff on vector computers and the CM-2. It turns out that a huge community of high-performance programmers have experimented with all of these concepts in the 1970s, 80s, 90s, and 00s, long before GPUs. All of their concepts still work today on modern GPUs. Looking up the right keywords: such as "CREW-PRAM algorithms" (Concurrent-Read Exclusive Write, Parallel RAM model) immediately gives you plenty of results for some things. (Ex: I just searched on 'CREW-PRAM DFS' and got: https://core.ac.uk/download/pdf/82490222.pdf https://core.ac.uk/download/pdf/82490222.pdf). The key is understanding what the "old word" for GPU was. That's PRAM, the Parallel-RAM model. That's how programmers from the 1970s, 1980s, and 1990s talked about algorithms written for the GPU-style called it back then. Newer articles / books talk about GPUs directly. -------------- I'd say the fundamentals are covered in the 1970s, such as "A Parallel Algorithm for the Efficient Solution of a General Class of Recurrence Equations" by Kogge, which leads to the development of Prefix sum (as it is called today). Of course, CUDA code is clearer than a theoretical CREW-PRAM discussion. So perhaps its easier if you just read GPU Compute Gems by NVidia to cover the same material. Still, I find that the 1970s writing is often times better written for a beginner (back then, fewer programmers knew how parallel programming worked, but were clearly more math-heavy than today. So I find that reading old articles like that one helps my personal brain and personal thinking pattern) --------- Ah right, the recommendation. I'd say "Data Parallel Algorithms" by Hillis / Steele (ACM Communications 1986) is an excellent introduction to general-purpose SIMD compute. It was written for CM-2, an ancient supercomputer that no longer exists, but the PRAM style applies to GPU algorithms today. Its like 15 pages, but it really opens your eyes to the possibilities. A lot of CUDA stuff is very specific to NVidia GPUs (important specific details: like bank conflicts and shared memory, which you absolutely should learn about... but such details should be studied after you learned the way of parallel-thinking / PRAM model / etc. etc.)

dahart 5y ago

Program size limits are a thing of the past. Today you don’t need to worry about program size, and (like on the CPU) you’ll probably bump into long compile times way before running out of space. Memory access is whatever RAM is on the board, plus whatever you want to stream/transfer from CPU RAM and/or disk. Complexity doesn’t matter, your GPU programs can be arbitrarily complex. what matters is that your workload is fairly uniform - GPUs are SIMT machines that will realize performance gains any time you have lots of threads (thousands, millions, billions) that all mostly do the same thing and can execute mostly independently. Yes, you can do Monte Carlo simulations, decision trees, and ML without Neural Networks. All of these things are extremely common.

Const-me 5y ago

> are there other good applications the GPU pipeline tends to be useful for? Pretty much everything compute-bound. SpaceX presentation about rocket engines: https://www.youtube.com/watch?v=vYA0f6R5KAI https://www.youtube.com/watch?v=vYA0f6R5KAI Unreal Engine 5 presentation about rendering many millions of small triangles using compute shaders: https://www.youtube.com/watch?v=TMorJX3Nj6U https://www.youtube.com/watch?v=TMorJX3Nj6U > Are there guidelines on how complex a program you can run on GPU cores and still reap benefit? Arbitrarily complex. Many real-world GPGPU programs are split over multiple compute shaders / kernels. When implemented properly, they even run in parallel. This is critically important on Windows with DirectCompute, because 2 seconds TDR timeout. When a single compute shader takes more than 2 seconds, the OS concludes the GPU’s hang, resets the hardware and reloads the driver. > What are (vague, order of magnitude) limits on program size and size of the different memories you have access to? Technically unlimited because you can stream data from system memory, disk, or even internets. Practically, most modern GPUs come with 6-12 GB VRAM, for optimal performance you’d want your data to fit there. When doing micro-optimizations, another important number is amount of on-chip memory, the order of magnitude is 64kb per core. > Can you do monte carlo simulations, decision trees etc? Monte Carlo is a great fit for GPUs, just don’t forget to seed your RNG with SV_DispatchThreadID or an equivalent. Trees are tricky. GPUs share instruction decoder and instruction pointer over 32-64 hardware threads. A straightforward implementation of binary trees gonna be suboptimal due to divergence of these threads. Often possible to do something else instead.

dragontamer 5y ago
> Trees are tricky. GPUs share instruction decoder and instruction pointer over 32-64 hardware threads. A straightforward implementation of binary trees gonna be suboptimal due to divergence of these threads. Often possible to do something else instead. The GPU-straightforward implementation is to use a SIMD-stack / SIMD-queue. I don't know its proper term, but its "obvious" to anyone who programs GPUs. The following is probably wrong, but hopefully correct enough to demonstrate the idea... active-lane-workgroup(){ wavefront_activelane = prefix-sum(execution-mask); // Execution mask is 0 if not currently executing, and 1 if currently executing in your wavefront. // wavefront_activelane is available at assembly level in one clock tick in both NVidia and AMD GPUs. __shared__ int workgroup_prefix_sum[wavefronts_in_workgroup]; if(wavefront_activelane == 0){ workgroup_prefix_sum[wavefront_id] = wavefront_horizontal_max(wavefront_activelane); } __syncthreads(); return workgroup_prefix_sum[my_wavefronts_idx] + wavefront_activelane; } SIMD-push(stack, data): stack[stack.ptr + active-lane-workgroup()] = data; __syncthreads(); if(active-lane-workgroup() == 0){ stack.ptr += workgroup-prefix-max(active-lane-workgroup()); } __syncthreads(); SIMD-pop(stack): toReturn = stack[stack.ptr - active-lane-workgroup()]; __syncthreads(); if(active-lane-workgroup() == 0){ stack.ptr += workgroup-prefix-max(active-lane-workgroup()); } __syncthreads(); return toReturn; ------ Now that you have a stack, its simply a matter of pushing your DFS into the SIMD-stack, and popping it off to traverse. You're somewhat BFS, because on a 1024-wide workgroup, your threads will visit the top 1024 items of the stack each step. But the stack overall behaves in a DFS manner. The CUDA cub library (and ROCm hipCUB / ROCm rocPRIM libraries) implement these horizontal operations (glorified prefix-sum). But its not too hard to write workgroup prefix-sum or workgroup prefix-max yourself. (Indeed: I suggest beginners write their own prefix-sum to get a feel of how simple and efficient the prefix sum / scan operations can be) -------- The "pattern" is that you use horizontal operations and __syncthreads() for thread synchronization and communication. In effect, this SIMD-stack is performing load-balancing simultaneously with the DFS. That is to say, it "reaches breadth-wise" and pulls extra nodes to visit when more lanes are available, while it prefers depth to minimize memory usage.

alfalfasprout 5y ago

tl;dr problems that are compute bound but trivially parallelizable tend to be good choices for GPU computation. Eg; when you're running a series of small computations on a lot of data. Or a lot of small computations on a moderate amount of data. GPUs tend to be pretty ineffective for computations where there are a lot of data dependencies since the individual shader units are slow compared to traditional CPU cores... so most of the GPU will be idle. Or for ultra low latency applications (since moving data between to and from the GPU is costly even with direct memory access).