Memory Bandwidth
fgiesen.wordpress.com
fgiesen.wordpress.com
The biggest bottleneck in my favorite workload (Netflix CDN server) is memory bandwidth. When doing software kTLS, we do 4 access per byte served to customers. One access to DMA the contents of a video file to DRAM from storage. A second access to read the byte for encryption. A third to write the encrypted data to DRAM. And a fourth to DMA the encrypted data to the NIC. That means we're bottlenecked serving Netflix video at ~240Gb/s on AMD Rome servers (quad channel DDR4-3200, single socket).
The theoretical performance, assuming everything besides serving the video is "free" is ~260-300Gb/s (depending on numa config, etc).
Every cache miss that I can eliminate leads to a small performance increase by both decreasing the memory system load (and leaving more headroom for the workload), and by eliminating a CPU stall on DRAM access.
I'm thrilled that next generation servers will have DDR5. That should help a lot with bandwidth, but the PCIe and network speeds will go up as well (Gen4->Gen5, 100GbE -> 400GbE), so we'll still be fighting the same battle.
Things like inline-kTLS (where the NIC does the TLS encryption) reduce the memory bandwidth requirements by roughly half, and is how we're able to serve up over 350Gb/s from a single socket.
I'd argue it's the only problem more times than not, although I don't work at a particularly high level of abstraction so YMMV
> That means we're bottlenecked serving Netflix video at ~240Gb/s on AMD Rome servers (quad channel DDR4-3200, single socket).
AMD Rome supports octa-channel DDR4 even in single-socket servers. Did you mean octa-channel DDR4, or is there some other limitation in the system that limits you to quad-channel configurations?
If storage is accessed via PCI-e, couldn't one just pipe it straight from the storage to the NIC, using p2pdma[1][2] or similar? Seems you could skip all of the main DRAM round-trip this way.
Or aren't things that simple?
[1]: https://www.kernel.org/doc/html/latest/driver-api/pci/p2pdma...
But for other, this is what is done. Which is probably why their use case here isn't optimized all that well. There have been ongoing efforts on and off to make DMA directly to/from CPU cache work well. And it can be made to work with some hardware, but I don't know if it's performed to potential.
In theory his workload requires no DRAM accesses for the data if the working set could fit in cache. DMA from disk to dirty cache line, CPU loads from there encrypts, stores it in place or to new dirty line, NIC loads from that dirty data, and the dirty cache lines are invalidated or overwritten in a ring buffer with new data to avoid the mandatory writeback to DRAM.
Some of the new CXL protocols coming down the pipeline might improve this situation, but they also might not.
However, NVME speaks in 4KB blocks, but NICs speak in ~1400 byte MSS chunks. According to a colleague who knows the internal architecture of NVME drives well, making 1400b requests is a non-starter.
So to use P2P DMA, you need some place to cache things. There turn out to be 3 places to cache things w/o going to the host:
1) On the NVME drive itself. This is known as controller memory buffer (CMB). Unfortunately, in order to get a CMB that's large enough to be useful, you need to buy gold-plated enterprise drives, which would be fine for a science experiment, but does not scale.
2) On the NIC. We use Mellanox CX6-DX NICs, and they have essentially no storage.
3) On something on the PCIe bus that has memory, like a GPU or custom "box of ram". The GPU would blow our thermal budget, and I can't think of a low cost, low-power "box of RAM" that can be attached at full speed to PCIe.
Ah, I had imagined the NIC TLS stuff to be more stream based, clearly out of my comfort zone this stuff.
> I can't think of a low cost, low-power "box of RAM" that can be attached at full speed to PCIe.
Yeah, I mean some FPGA cards could probably do the trick, but would not be a low cost option and I guess not low power if you want full speed PCIe.
Close but no cigar for p2pdma then...
It's the same problem when solving technical difficulties in life: I fill my pockets and backpack with adaptors, because the cost of a cache miss (going back to my room to get it) means it's no longer useful (for plugging in the laptop to the projector so these people can see the slides).
If only we could predict the future, then the cache could pre-load! And usually things do happen as planned, but sometimes free will gets involved and we branch another way, and need to fetch something else from memory, and that takes a while.
It gets difficult to fit it all and make it work. An 8-channel DDR4 memory controller already has 512 pins just for data bits, and that's before you deal with all of the necessary grounding pins and all of the other signal lines required to interface with memory chips.
There are rumors that AMD's next generation server CPUs will support 12-channel DDR5. The expanded memory channels and other features require a socket with 6096 pins. That's a lot of interfaces to squeeze through a socket.
A hypothetical 20-channel CPU would also require you to install RAM in multiples of 20 sticks to populate every channel.
You can theoretically get 16 channels of memory bandwidth in a single, dual-socket system if each socket has an 8-channel memory controller. Dual-socket AMD systems have a theoretical upper limit memory bandwidth around 300GB (big B) per second.
In reality, not as many applications are constrained my memory bandwidth as you might expect. CPU vendors are adding additional cache to work around the memory issue somewhat. As long as the application isn't constantly streaming bytes into and out of memory (as in the CDN example), it will eventually have to stop to work on the data, allowing the CPU a chance to cache some additional data from main memory and hide the memory bandwidth bottleneck somewhat.
Rome has 8 memory channels. https://en.wikichip.org/wiki/amd/cores/rome
Let me add my $0.02. CPUs seem to hide memory issues with more-and-more cache. The most recent AMD EPYC chips have 256MB of L3 cache. Intel has less L3 cache, but has something like 1MB of L2 cache per core now (!!!). AMD has shown us the future: 64MB SRAM being grafted onto these chips to maybe add +64MB per die (or maybe 768 MB of L3 cache in a hypothetical EPYC in the future).
GPUs have a different design. GPUs have almost no caches worth discussing. Instead, GPUs have manually managed scratch RAM, also known as __shared__ memory (which roughly fits in the L1 cache speeds). __shared__ is pretty small though, maybe only 64kB or so. The real important thing about __shared__ is that any SIMD lane can read/write to it at high speeds, so it serves as a good area for inter-thread communications. ("Thread" is a bit ambiguous, but... lets just simplify it and call them cuda-threads).
The other bit GPUs have are a ridiculous number of registers. AMD RDNA has 1024-registers, of which up to 256 can be allocated into a kernel. (The device driver of modern GPUs will further parallelize loads by requesting only some of the registers per invocation. A well written memcpy kernel may only need 64 vector-registers for example... the device driver may load 16 parallel instances of this hypothetical kernel to run in parallel on the GPU core so all 1024 registers are used).
These 1024-registers are 32-bits and 32-SIMD-lanes wide on AMD RDNA, or 128kB of registers per WGP (workgroup processor, an AMD RDNA "core"). I kid you not: AMD RDNA has more registers per WGP than Intel Skylake had L1 cache per core!!
---------------
As such, we can immediately see how modern GPUs and CPUs get around the memory bandwidth issue. CPUs try to hide it with increasingly large caches... while GPU register sizes are so large that it borders on insanity.
Hmmm... NVidia's is clearly aiming at just punching through the memory-bandwidth problem with GDDR6x (2-bits per clock tick since its got 4-level encoding).
That's the thing, NVidia isn't really pushing memory bandwidth or size limits on the L2 or even L1 cache IMO. Even the 128kB L1$ per SM is only roughly the size of the SM's register space. Their most interesting move really is GDDR6x, which is the brute-force way to solve that problem.
--------
AMD's "Infinity Cache" on RDNA2 is 128MB of L3$, but AMD is using only standard GDDR6 (1-bit per clock tick transferred). AMD's RDNA2 is very strange: L0, L1, L2, and L3 caches, when AMD GCN was just L1 and L2 layers of cache.
That "infinity cache" is worth talking about I guess... its large enough to be relevant in a number of gaming situations.
------
I guess AMD and NVidia are both using HBM at the high end for 1TBps to 2TBps bandwidths. But those chips aren't in the consumer realm anymore. The ultimate brute force solution: spend more money.
You're right in that the L1 and L2 caches (and L0 and L3 caches of AMD) probably do affect performance in real ways.
AMD does not. L1 and __shared__ are different pools on GCN, CDNA, and RDNA architectures. I believe shared is actually higher bandwidth than L1 on AMD systems, especially with atomics.
> Registers and cache are used together, though
But not for the same purpose, or the same way as CPUs. The cache is non-cohesive, large amounts of the cache are "K$", constant space that's non-cohesive. Etc. etc.
Its a bit different. Some level of caching will improve effective memory bandwidth, but it seems like the GPU's primary strategy is to "float" register loads. GPUs are still an in-order processor but... the load-register assembly instructions clearly execute in an async-like manner.
That load-register assembly instruction could be from L1 cache, it could be from GDDR6x, it could be from another GPU over NVLink / NVSwitch, or it could be even from PCIe (!!), being stored on the CPU's DDR4 RAM all the way across the motherboard.
Doesn't matter: the load-register instruction will appropriately start loading the data, and will indicate to the core when the register is ready (and the GPU core will task-switch to other kernels while waiting for that request to be completed).
That's probably the likely future, but an alternate future is 0 (L3) cache on die, plus the 3d cache on top. Save space on the main die which uses the most expensive fab. Depends on how costly adding cache on top is, I guess.
Read a byte, use it to feed FNV-1a, use that to look up another byte in a fixed-size array. The number of bytes per second you end up reading ends up comedically low.
It turned out 87% of all straight-to-DRAM accesses happened on one line.
And DDR5 / LPDDR5 seems to have enough head room for us for another 5 years. So we seem to be doing fine ( For now )
Edit: Read Dragontamer's post above / below.
This is a weird/incorrect way of looking at memory bandwidth on a GPU. You can saturate the memory bandwidth without using even close to all of the threads available. I'm not sure why the author chose to use this metric rather than simply running a benchmark.