6 ms·

Memory Bandwidth

114 points by luu 5y ago

dragontamer 5y ago
This is a good start to the discussion. 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.
- sephamorr 5y ago
  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/.
  - dragontamer 5y ago
    > 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.
    - shaklee3 5y ago
      Registers and cache are used together, though. And shared memory is the same as the L1 cache, but manually controlled.
      - dragontamer 5y ago
        NVidia __shared__ takes from L1 cache. 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).
- toast0 5y ago
  > 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.
  - ksec 5y ago
    Doesn't work with Heat Spot. Although we can increase L2 / L3 Cache with Die Stacking from top to bottom.
ksec 5y ago
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 ) Edit: Read Dragontamer's post above / below.
corysama 5y ago
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.
- mhh__ 5y ago
  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.
drewg123 5y ago
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.
- mhh__ 5y ago
  > real problem 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
- PragmaticPulp 5y ago
  That's a fascinating project to work on. > 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?
- magicalhippo 5y ago
  > 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. Or aren't things that simple? [1]: https://www.kernel.org/doc/html/latest/driver-api/pci/p2pdma.html https://www.kernel.org/doc/html/latest/driver-api/pci/p2pdma... [2]: https://youtu.be/LDOlqgUZtHE?t=612 https://youtu.be/LDOlqgUZtHE?t=612
nyanpasu64 5y ago
(2017)?
sydthrowaway 5y ago
One of the best blogs for systems programming. Anyone know others?
shaklee3 5y ago
> 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.
zozbot234 5y ago
(2017), it would be nice to see updated data on this. But of course the basic point stands.

114 points by luu 5y ago

This is a good start to the discussion. 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.

sephamorr 5y ago
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/.
- dragontamer 5y ago
  > 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.
  - shaklee3 5y ago
    Registers and cache are used together, though. And shared memory is the same as the L1 cache, but manually controlled.
    - dragontamer 5y ago
      NVidia __shared__ takes from L1 cache. 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).
toast0 5y ago
> 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.
- ksec 5y ago
  Doesn't work with Heat Spot. Although we can increase L2 / L3 Cache with Die Stacking from top to bottom.

ksec 5y ago

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 ) Edit: Read Dragontamer's post above / below.

corysama 5y ago

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.

mhh__ 5y ago
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.

drewg123 5y ago

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.

mhh__ 5y ago
> real problem 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
PragmaticPulp 5y ago
That's a fascinating project to work on. > 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?
magicalhippo 5y ago
> 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. Or aren't things that simple? [1]: https://www.kernel.org/doc/html/latest/driver-api/pci/p2pdma.html https://www.kernel.org/doc/html/latest/driver-api/pci/p2pdma... [2]: https://youtu.be/LDOlqgUZtHE?t=612 https://youtu.be/LDOlqgUZtHE?t=612

nyanpasu64 5y ago

(2017)?

sydthrowaway 5y ago

One of the best blogs for systems programming. Anyone know others?

shaklee3 5y ago

> 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.

zozbot234 5y ago

(2017), it would be nice to see updated data on this. But of course the basic point stands.