WebGPU computations performance in comparison to WebGL
pixelscommander.com
pixelscommander.com
https://www.khronos.org/registry/webgl/specs/latest/2.0-comp...
https://www.khronos.org/assets/uploads/developers/presentati...
https://bugs.chromium.org/p/chromium/issues/detail?id=113199...
> macOS' OpenGL implementation never supported compute shaders
which is a pretty valid reason to drop it. Nothing to do with Chrome. Further, OpenGL is dead, every graphics developer knows this. All the browser vendors are working on WebGPU. No reason to waste time on implementing outdating things. There are finite resources. No other browser vendor named any plans to support webgl-compute. So stop with the ridiculous spin
Nothing on the standard requires OpenGL backend, only that the semantics are preserved, GL ES, Metal and DirectX 11 all have compute shaders support.
OpenGL is not dead until Khronos comes up with an API that is actually usable without an EE degree on GPU design and shader compilers.
WebGPU is years away to become usable, even if it happens to be released during 2022, it will be a 1.0 MVP, years behind of what Metal, Vulkan, DX 12, NVN and LibGNM are capable of in 2021.
> In order to reclaim code space in Chromium's installer that is needed by WebGPU, the webgl2-compute context must be removed.
I bet the Chrome team wouldn't bother with code space installer for other Google critical features that are part of Project Fugu.
WebGL on Linux requires the OpenGL backend until eventually Vulkan is supported well enough.
Apple's WebGL2-on-Metal approach involves the same translator engine I work on, and was very much a collaborative effort between Google, Apple, and external contributors. Chrome will adopt this in the future after integration work is finished.
I can confirm that browser GPU programmers are definitely spread pretty thin just ensuring compatibility and reliability across all platforms
No wonder the focus on making streaming work for 3D rendering with current APIs on modern hardware.
It would be even further away if all the devs working on it were instead spending years on WebGL-Compute. It's not like you snap your fingers and it works. It takes years of cross platform work and testing. That time is better spent moving forward.
As further proof, the only names on the compute spec are Intel and one Google name. Compare to the WebGPU spec which clearly has buy in from all the browsers.
So if Chrome had shipped it you'd have the gallery of comments bitching about Chrome going too fast, implementing too many things no other browser plans to implement.
BTW: WebGL2 in Safari does not run on Metal (yet)
As a counterpoint, I've been using WebGPU (through wgpu-rs) for the past 1.5 years. It's been a pleasure to use. For instance, here's the CPU-side code for a glow post-process shader using 4 render passes https://github.com/JMS55/sandbox/blob/master/src/glow_post_p....
A plugin framework for 3D APIs is probably the easiest part to implement on a game engine, specially when API like OpenGL and Vulkan already require implementing one due to extension spaghetti.
Why would this ever happen? It seems like there is nothing in the works, nothing on the horizon, and no demand for a higher-level, less-performant API to become a new standard. Even OpenGL itself has been getting lower level and more detailed ever since version 1. People can build/use a wrapper API or game engine or something else if they want easy. It seems weird to say this right after defending Apple’s use of Metal to implement WebGL. Apple’s moves to get rid of OpenGL from the Mac ecosystem are one of the strongest forces pushing OpenGL out.
FWIW as someone who exclusively uses OpenGL for 3D graphics, this actually makes me push Apple out :-P
Regarding Metal, indeed those that leave OpenGL are more likely to move into middleware than forcing Vulkan upon themselves.
Hence why Khronos started ANARI, as most visualisation products and CAD/CAM people couldn't care less that Vulkan on its present state exists.
In many other benchmarks, it loses. That being said, it still manages decent performance, which is extremely impressive for a one man project using only the Vulkan interface. It wouldn't surprise me if it eventually becomes the default OpenGL driver in the open source driver stack (for HW capable enough to support a Vulkan driver, obviously).
As for using middleware, GPU's are vastly more capable and complicated today than 30y ago when OpenGL 1 appeared. In most cases it makes sense to use a higher level interface, specialized for the particular type of application you're writing, be it ANARI, a game engine, some scenegraph library, or whatever.
That isn't going to happen because everyone is, in fact, moving away from the idea of a "graphics API" altogether and simply allowing the compute systems to calculate everything.
See: the Nanite renderer from EPIC: https://www.youtube.com/watch?v=eviSykqSUUw
To a first and second order approximation, no one cares about graphics that aren't related to games.
OP is suggesting that Chrome should have spearheaded a spec that no one but Intel and Google ever showed any interest in. How would that not have been literally "they like a standard (or propose one), they implement it, use it on their projects like google doc, and let the competition deal with the mess"[1]?
Instead they're implementing a spec with broad involvement and support across browsers, it's just taking longer. Seems like exactly how we'd want things to go.
Yes. Exactly what they've been doing with plenty of other "standards"
Also, none of specific reasons people opposed Microsoft’s policies with IE, like it’s attempts to lock people into windows-only APIs (ActiveX etc) apply today.
In this case, google can pull off something like OP says because it's dominating the market.
> Also, none of specific reasons people opposed Microsoft’s policies with IE, like it’s attempts to lock people into windows-only APIs (ActiveX etc) apply today.
Chrome implements plenty of chrome only API. They like a standard (or propose one), they implement it, use it on their projects like google doc, and let the competition deal with the mess.
This is nothing like the situation was with IE6. Back then basic functionality was half-broken and it was hard to get anything done without a pile of hacks.
Causes are not 1 to 1, but consequences are the same: monopoly leads to abuse, abuse lead to some site not working with Firefox, consommers needs being ignored by google and them abusing their dominant position to pass what they want as standard, or just destroy API they don't like (see the latest adblock scandal).
edit: GLES3.2 => GLES 3.1
WebGPU is still on the horizon during the next couple of years.
Now, years later, Apple actually implemented WebGL on Metal, so today we could think about implementing WebGL compute on Mac. However WebGPU is now in origin trials. It's very unlikely that Apple would put any effort into a compute shader implementation for WebGL now. And Chrome is not going to go implement a major WebGL feature that has no prospect of ever being supported in Safari.
But for WebGL it matters what Apple does?
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?
Pretty much everything compute-bound. SpaceX presentation about rocket engines: 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
> 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.
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.
I agree. Haven’t personally did that exact thing because wave intrinsics require D3D 12.2 (or new enough CUDA like in your example), but I did similar stuff with group shared memory which is more compatible, only requires feature level 11.0 hardware.
However, that GPU-straightforward implementation is not that straightforward for programmers with CPU background starting to use GPUs. Even simpler things like reduction (dot product, or matrix*vector in linear algebra) are rather tricky to implement efficiently on GPUs, very different than even manually-vectorized SIMD CPU code.
But if you read older PRAM stuff from the 1980s or 1990s, they talk about "execution-masks" and manipulating it directly like this. So its helpful to know how this older technique relates to modern APIs like cooperative groups.
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.
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.
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.
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!)
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).
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.)
They are all closely related to the PRAM-model however. So algorithms study isn't really about learning CUDA-details or whatever, but learning the generic (and cross-language) concepts of parallel programming.
------
So it really doesn't matter if you're reading CUDA, or C-star (1980s code for the old CM-2 supercomputer). They're both PRAM in concept and therefore somewhat compatible in your brain.
It helps to know GPU-quirks for highest levels of optimization (wavefront programming, uniform branches, __shared__ memory), but you can learn these details after learning the generic PRAM stuff known for the past decades.
Regarding learning GPU Architectures and Programming, usually in the Graphics Gems books there is an introductory section devoted to compute. But you are on own regarding streaming, tracing, tuning, multi gpu. All still very much dark arts ;)
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).
Exploitable vulnerabilities will probably come in faster than the performance...
You can't easily turn it off in Chrome/Blink/Edge
Same as WebRTC.
I know WebGPU is an subset, but I always wanted to how many features are missing compared to modern Direct X 12 and Metal ( As well as possibly GNMX, Vulkan etc )
https://twitter.com/BenHouston3D/status/1466054449898659842
There is a separate project from Google for just WebGPU in Node.js here but it isn't production ready:
Also Apple keeps threatening to kill OpenGL so any day now all the people developing on Mac might lose the ability to dev this headless component locally. (And this was always the situation with Vulkan.)
That hit me a few times, but ads and malware ruined that for everyone. (and there is no simple permission dialog possible, where you ask the user to allow it)
https://developer.mozilla.org/en-US/docs/Web/API/WebGL_API/W...
Sounds good (along with the performance graph).
But I would also just like a asynchronous getPixelData for the canvas now.
Does anyone know if there is a chance of this coming anytime soon? From my user perspective, it does not seem too complicated to implement, but the potential benefits seem great.
Because WebGL is working quite stable now and WebGPU not yet.
Well, not the fastest path for reading back the rendertarget, so I can see why not doing it would be more efficient