Memory bandwidth and latency is a real problem. As the referenced article points out, both have gotten worse as compared to CPU speeds as time goes on.
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.
> Things like inline-kTLS (where the NIC does the TLS encryption) reduce the memory bandwidth requirements by roughly half
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.
I've been looking at this. When using inline hardware kTLS offload, this is an amazing solution on paper.
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.
We've spoken to at least one small NIC vendor that puts ~48GB of RAM on their FPGA based smart nic. That let's them handle TCP rexmits without paying the price of re-DMA'ing data from the host to re-encrypt it. I wish the FPGA didn't blow our power budget, because their solution would be ideal for P2P DMA.
I imagine for some reason their encryption or some other processing step is not performed by their hardware (you can get NICs with crypto offload) or they've found it cheaper or more flexible to go with doing it on the CPU.
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.
> 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?
Dumb question: since memory bandwidth is clearly a bottleneck in many server workloads, why don't chip manufacturers make CPUs with 20 memory channels instead of just 6-10, as in current CPUs?
> Dumb question: since memory bandwidth is clearly a bottleneck in many server workloads, why don't chip manufacturers make CPUs with 20 memory channels instead of just 6-10, as in current CPUs?
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.
"There are two hard things in computer science: cache invalidation, naming things, and off-by-one errors" - Phil Karlton
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.
Fascinating! Have you looked at other ways to get more TLS performance out of the system as a whole? How does that performance even compare to the Netflix services side of things (aka in AWS)? How do various TLS offload strategies stack up?
how do you eliminate cache misses in this workload? seems like long runs of sequential reads/writes will result in eviction by the time the “byte” is ready for the next stage of the process.
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.
> AMD has shown us the future: 64MB SRAM being grafted onto these chips to maybe add +64MB per die.
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.
Isn't this a bit dated? I certainly agree that the shared memory semantics are a critical distinction, but in addition to the 256Kb per SM shared memory, NVIDIA Volta GPUs have 128kB L1$ per SM, and a unified 6MB L2$. I don't think these caches are entirely /not worth discussing/.
> I don't think these caches are entirely /not worth discussing/
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).
A fun exercise is to see how bad you can get effective bandwidth to be in the minimal amount of code.
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.
When profiling a certain compiler I found a O(n^2) algorithm which literally did almost exactly what you said. 1 cacheline accessed for every 3 iterations due to some enormous structs, whose members index other arrays etc. Etc.
It turned out 87% of all straight-to-DRAM accesses happened on one line.
The golden rule for CPU is Latency and Bandwidth for GPU. Which is why having higher L2 and L3 cache is much more important than say raw DRAM bandwidth. The extra bandwidth on SoC / APU are mostly for GPU purposes.
And DDR5 / LPDDR5 seems to have enough head room for us for another 5 years. So we seem to be doing fine ( For now )
> sounds good, but each SM drives 128 “CUDA cores”, each corresponding to a thread in the SIMT programming model. Per thread, we get about 0.09 bytes of memory bandwidth per cycle – or perhaps less awkward at this scale, one byte every 11 instructions.
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.
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.