GPUs have far more bandwidth, but CPUs beat them in latency. Being able to AVX512 your L1 cached data for a memcpy will always be superior to passing data to the GPU.
With Ice Lake's 1MB L2 cache, pretty much all tasks smaller than 1MB will be superior in AVX512 rather than sending it to a GPU. Sorting 250,000 Float32 elements? Better to SIMD Bitonic sort / SIMD Mergepath (https://web.cs.ucdavis.edu/~amenta/f15/GPUmp.pdf) on your AVX512 rather than spend a 5us PCIe 4.0 traversal to the GPU.
It is better to keep the data hot in your L2 / L3 cache, rather than pipe it to a remote computer (even if the 16x PCIe 4.0 pipe is 32GB/s and the HBM2 RAM is high bandwidth once it gets there).
--------
But similarly: CPU SIMD can never compete against GPGPUs at what they do. GPUs have access to 8GBs @500GB/s VRAM on the low-end and 40GBs @1000GB/s on the high end (NVidia's A100). EDIT: Some responses have reminded me about the 80GB @ 2000GB/s models NVidia recently released.
CPUs barely scratch 200GB/s on the high end, since DDR4 is just slower than GPU-RAM. For any problem where data-bandwidth and parallelism is the bottleneck, that fits inside of GPU-VRAM (such as many-many sequences of large scale matrix multiplications), it will pretty much always be better to compute that sort of thing on a GPU.
-------
It should be noted that most CPUs do use DDR4 for good reason: they have an expectation to use more than the capacity of HBM2. GPUs (and a supercomputer CPU, like the A64FX) can rely upon the fact that their workloads are designed for distributed-compute and/or otherwise fit inside of 32GBs or 40GBs.
CPUs on the other hand, could very well be reading/writing to a 2TB in-memory database these days, even on a single-socket single-node system like AMD EPYC. That flexibility to have as much (or as little) RAM as the customer demands is a major advantage of traditional CPUs like EPYC or Xeon, which the A64FX cannot partake in.
If you need more than 32GBs of RAM, A64Fx is a no-go. That HBM2 is soldered directly onto the package and cannot expand.
Although an A100 still has twice the bandwidth and over 3x the GFLOPS without tensor cores and more than 6x with.
Yeah. And at 3200 Mbit/sec, that comes out to 200GB/s. (3200 MHz x 8-bytes (aka 64-bit) == 25GB/s. x8 channels == 200GB/s).
> where IIRC the biggest NVIDIA part is still at 6.
That's 6x *1024-bit* HBM2 channels. Total bandwidth is 2000GBps, or over 10x the speed of the "8x channel EPYC". Yeah, HBM2 is fat, extremely fat.
----------
*ONE* HBM2 channel offers over 300GBps bandwidth. And the A100 has *SIX* of them. Literally ONE HBM2 channel beats the speed of all 8x DDR4 EPYC memory channels working in parallel.
25600 is the channel rate in (EDIT) MB/sec of the stick of RAM. That's 25GB/s for a 3200 MHz DDR4 stick. x8 (for 8-channels working in parallel) is 200GB/s.
-----------
This has been measured in practice by Netflix: https://2019.eurobsdcon.org/slides/NUMA%20Optimizations%20in...
As you can see, Netflix's FreeBSD optimizations have allowed EPYC to reach 194GB/s measured performance (or just under the 200GB/s theoretical). And only with VERY careful NUMA-tuning and extreme optimizations were they able to get there.
BVH-tree traversals are done on the GPU now for a reason. GPUs are better at latency hiding and taking advantage of larger sets of bandwidth than CPUs. Yes, even on things like pointer-chasing through a BVH-tree for AABB bounds checking.
GPUs have pushed latency down and latency-hiding up to unimaginable figures. In terms of absolute latency, you're right, GPUs are still higher latency than CPUs. But in terms of "practical" effects (once accounting for latency hiding tricks on the GPU, such as 8x way occupancy (similar to hyperthreading), as well as some dedicated datastructures / programming tricks (largely taking advantage of the millions of rays processed in parallel per frame), it turns out that you can convert many latency-bound problems into bandwidth-constrained problems.
-----------
That's the funny thing about computer science. It turns out that with enough RAM and enough parallelism, you can convert ANY latency-bound problem into a bandwidth-bound problem. You just need enough cache to hold the results in the meantime, while you process other stuff in parallel.
Raytracing is an excellent example of this form of latency hiding. Bouncing a ray off of your global data-structure of objects involved traversing pointers down the BVH tree. A ton of linked-list like current_node = current_node->next like operations (depending on which current_node->child the ray hit).
From the perspective of any ray, it looks like its latency-bound. But from the perspective of processing 2.073 million rays across a 1920 x 1080 video game scene with realtime-raytracing enabled, its bandwidth bound.
> (which is what I think ajross keeps alluding to)
It seems like ajross is accusing me of underestimating CPU-bandwidth. At least, that's my interpretation of the discussion so far. As you've pointed out however, I'm overestimating it.
EDIT: But I'm overestimating it on both sides. A100 2000 TB/s is the "channel bandwidth" as well, as the CAS and RAS commands still need to go through the channel and get interpreted.
Whereas Nvidia's A100 has over 2000 GB/s of memory bandwidth. That's 10-fold better.
The two last apps I worked on have been GPU-only. The CPU process starts running and launches GPU work, and that's it, the GPU does all the work until the process exits.
There is no need to "pass data to the GPU" because data is never on CPU memory, so there is nothing to pass from there. All network and file I/O goes directly to the GPU.
Once all your software runs on the GPU, passing data to the CPU for some small task doesn't make much sense either.
However, the famous "Moana" scene for Disney-level productions is a 93GB (!!!!) scene statically, with another 131GBs (!!!) of animation data (trees blowing in the winds, waves moving on the shore, etc. etc.).
That's simply never going to fit on a 8GB, 40GB, or even 80GB high-end GPU. The only way to work with that kind of data is to think about how to split it up, and have the CPU store lots of the data, while the GPU processes pieces of the data in parallel.
https://www.render-blog.com/2020/10/03/gpu-motunui/
Which has been done before, mind you. But it should be noted that the discussion point for GPU-scale compute runs into practical RAM-capacity constraints today, even on movie-scale problems from 5 years ago (Moana was released in 2016, and had to be rendered on hardware years older than 2016).
Moana scene is here if you're curious: https://www.disneyanimation.com/resources/moana-island-scene...
----------
But yes, if your data fits within the 8GBs GPU (or you can afford a 40GB or 80GB VRAM GPU and your data fits in that), doing everything on the GPU is absolutely an option.
If you’re Disney, you can afford boxes with 10+ A100s with NVLink sharing the memory in a single 400+ GB pool. Unknown if that ends up being more economical than the equivalent CPU version, but it’s important to understand in order to evaluate the future of GPUs.
There is always a problem size that does not fit into memory.
Whether that memory is the GPU memory, or the CPU memory, doesn't really matter.
We have been solving this problem for 60 years already. It isn't rocket science.
---
The CPU doesn't have to do anything.
The GPU can map a file stored on hard disk to VRAM memory, do random access into it, process chunks of it, write the results into network sockets and send them over the network, etc.
The only thing the CPU has to do is launch a kernel:
int main(args...) {
main_kernel<<<...>>>(args...);
synchronize();
return 0;
}
---
> However, the famous "Moana" scene for Disney-level productions is a 93GB (!!!!) scene statically, with another 131GBs [...] That's simply never going to fit on a high-end GPU.
LOL.
V100 with 32Gbs and 8x per rack gave you 256 Gb of VRAM addressable from any GPU in the rack.
A100 with 80GB and 16x per rack give you 1.3 TB of VRAM addressable from any GPU in the rack.
You can fit Moana in GPU VRAM in a now old DGX-2.
If you are willing to bet cash on Moana never fitting on a single GPU, I'd take you on that bet. Sounds like free money to me.
This person rendered the Moana scene on just 8GBs of GPU VRAM. It does this by rendering 6.7GB chunks at a time on the GPU, with the CPU keeping the RAM-heavy "big picture" in mind. (EDITED paragraph. First wording of this paragraph was poor).
------
Its not that these problems "cannot be solved", its that these problems "become grossly more complicated" when under RAM / VRAM constraints. They're still solvable, but now you have to do strange techniques.
------
With regards to a Ray-tracer, tracing the ray-of-light that's bouncing around could theoretically touch ANY of the 93GBs of static object data (which could have been shifted by any of the 131 GBs of animation data). That is to say: a ray that bounces off of a any leaf on any tree could bounce in any direction, hitting potentially any other geometry in the scene.
That pretty much forces you to keep the geometry in high-speed RAM, and not do an I/O cycle between each ray-bounce.
As a rough reminder of the target performance: Raytracers aim at ~30 million to 30-billion ray-bounces per second, depending on movie-grade vs video-game optimized. Either way, that level of performance is really only ever going to be solved by keeping all of the geometry data in RAM.
> A100 with 80GB and 16x per rack give you 1.3 TB of VRAM addressable from any GPU in the rack.
That doesn't mean it makes sense to traverse a BVH-tree across a relatively high-latency NVLink connection off-chip. I know GPUs have decent latency hiding but... that's a lot of latency to hide.
Again: your CPU-renderers can hit 10s of millions of rays per second. I'm not sure if you're gonna get something pragmatic by just dropping the entire geometry into distributed NVSwitch'd memory and hoping for the best.
Honestly, that's where the 8GB CPU+GPU team becomes interesting to me. A methodology for clearly separating the geometry and splitting up which local compute-devices are responsible for handling which rays is going to scale better than a naive dump reliant on remote-connections pretending to be RAM.
Video games hit Billions of rays/second. The promise of GPU-compute is on that order, and I just doubt that remote RAM accesses over NVLink will get you there.
> If you are willing to bet cash on Moana never fitting on a single GPU, I'd take you on that bet. Sounds like free money to me.
The issue is not Moana (or other movies from 2016), the issue are the movies that will be made in 2022 and into the future. Especially if they're near photorealistic like Marvel-movies or Star Wars.
----------
The other problem is: what's cheaper? A DGX-system could very well be faster than one CPU system. But would it be faster than a cluster of Ice-Lake Xeons with AVX512 each with the precise amount of RAM needed for the problem? (Ex: 512GBs in some hypothetical future movie?)
A team probably would be better: CPUs have expandable RAM, that's their biggest advantage. GPUs have fixed RAM. Slicing the problem up so that pieces of Raytracing fits on GPUs, while the other, "bulkier" bits fit on CPU DDR4 (or DDR5), would probably be the most cost-efficient way at solving the raytracing problem.
The GPU-Moana experiment showed that "collecting rays that bounce outside of RAM" is an efficient methodology. Slice the scene into 8GB chunks, process the rays that are within that chunk, and the collate the rays together to find where the rays go.
This is very interesting - do you have a link that explains how it works / is implemented?
https://www.nvidia.com/en-us/geforce/news/rtx-io-gpu-acceler...
https://www.amd.com/en/products/professional-graphics/radeon...
The GPU itself can send PCIe 4.0 messages out. So why not have the GPU make I/O requests on behalf of itself? Its a bit obscure, but this feature has been around for a number of years now. The idea is to remove the CPU and DDR4 from the loop entirely, because those just bottleneck / slowdown the GPU.
--------
From an absolute performance perspective, it seems good. But CPUs are really good and standardized at accessing I/O in very efficient ways. I'm personally of the opinion that blocking and/or event driven I/O from the CPU (with the full benefit of threads / OS-level concepts) would be easier to think about than high-performance GPU-code.
But still, its a neat concept, and it seems like there's a big demand for it (see PS5 / XBox Series X).
I think the Radeon + Premier Pro documentation makes it clear how it works: https://www.amd.com/system/files/documents/radeon-pro-ssg-pr...
As you can see, the GPU is attached to the x16 slot, and the 4x NVMe SSDs are attached to the GPU. When the CPU wants to store data on the SSD, it communicates first to the GPU, which then pass-throughs the data to the four SSDs.
That's the simpler example.
--------------
In NVidia's case, they're building on top of GPUDirect Storage (https://developer.nvidia.com/blog/gpudirect-storage/), which seems to be based on enterprise technology where PCIe switches were used.
NVidia's GPUs would command the PCIe switch to grab data, without having the PCIe switch send data to the CPU (which would most likely be dropped in DDR4, or maybe L3 in an optimized situation).
That's a pretty big "if" and often not a workable.
You just write normal C++, and it works.
Another aspect that seems like a hidden assumption in CPU-GPU discussions is that you have the time-energy-expertise budget to (re)build your application to fit GPUs.
40TBs+ -- Storage-only solutions. "External Tape Merge sort algorithm", "Sequential Table Scan", etc. etc. (SSDs or even Hard drives if you go big enough)
4TB to 40TBs -- Multi-socket DDR4 RAM is king (8-way Ice Lake Xeon Scalable Platinum will probably reach 40TBs). Single-node distributed memory with NUMA / UPI to scale.
1TB to 4TB -- Single Socket DDR4 RAM (EPYC, even if at 4x NUMA. Or Single-node Ice Lake).
80GB to 1TB -- DGX / NVlink distributed memory A100 ganging up HBM2 together. GPU-distributed RAM is king.
256MBs to 80GBs -- HBM2 / GDDR6 Graphics RAM is king (80GB A100 2TB/s).
1.5MBs to 256MBs -- L3 cache is king (8x32MBs EPYC L3 cache, or POWER9 110MB+ L3 cache unified)
128kB to 1.5MBs -- L2 cache is king (1.25MB Ice Lake Xeons L2, this article)
1kB to 128kB -- L1 cache is king. (128kB L1 cache on Apple M1). Note: "GPU __Shared__" is a close analog to L1 and competes against it, but is shared between 32 to 256 GPU threads, so its not an apples-to-apples comparison.
1kB and below -- The realm of register-space solutions. (See 64-bit chess engine bitboards and the like). Almost fully CPU-constrained / GPU-constrained programming. 256x 32-bit GPU registers per GPU-thread / SIMD thread. CPUs have fewer nominal registers, but many "out of order" buffers or "reorder buffers" that practically count as register storage in a practical / pragmatic sense. CPUs just use their "real registers" as a mechanism to automatically discover parallelism in otherwise single-thread written code.
------------
As you can see: GPUs win in some categories, but CPUs win in others. And these numbers change every few months as a new CPU and/or GPU comes out. And at the lowest levels: CPUs and GPUs cannot be compared due to fundamental differences in architecture.
For example: GPU __shared__ memory has gather/scatter capabilities (the NVidia PTX instructions / AMD GCN instructions permute vs bpermute), while CPUs traditionally only accelerate gather capabilities (pshufb), and leave vgather/vscatter instructions to the L1 cache instead. GPUs have 32x ports to __shared__, so every one of the 32-threads in a wave-front can read/write every single clock-tick (as long as all 32 they are on different ports/alignment, or you have a special one-to-all broadcast). CPUs only have 2 or 4 ports, so vscatter and vgather operate slowly, as if a single thread were reading/writing each of the memory locations.
But CPU L1 cache has store-forwarding, MESI + cache coherence, and other acceleration features that GPUs don't have.
GPUs are therefore more efficient at sharing data within workgroups of ~256 threads, but CPUs are more efficient at sharing data between cores, or even among out-of-die NUMA solutions, thanks to robust MESI messaging.
https://www.phoronix.com/scan.php?page=news_item&px=Linus-To...
The problem is that for the workloads that need specialization, you sometimes really need it. You could also delete the vectorized AES units in your Intel machines too and the general purpose performance wouldn't be affected much, but cryptographic performance specifically would tank, and it turns out, that matters a lot in aggregate for many people.
Ultimately there are literally dozens of specialized inactive units on any CPU at any given time that could be "better spent on general purpose units" (which also isn't necessarily true if other architectural choices prevent those units from being utilized effectively). People just like complaining about AVX-512 because it's easily digestible water cooler chat they read about on a blog.
It impedes latency-sensitive applications due to forced down-clocking taking place for AVX-intensive instruction streams. The other side effect is that it quickly produces extreme heat that not only forces further down-clocking of the core but also taints neighbouring cores with dissipated heat and prevents them from going into turbo.
Not everything is about performance in laptops.
Only interesting if you care about price (dollars spent per inference)
For raw speed (no matter the price) the GPU wins
SIMD throughput?