<story><title>Memory Bandwidth</title><url>https://fgiesen.wordpress.com/2017/04/11/memory-bandwidth/</url></story><parent_chain></parent_chain><comment><author>drewg123</author><text>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.&lt;p&gt;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&amp;#x27;re bottlenecked serving Netflix video at ~240Gb&amp;#x2F;s on AMD Rome servers (quad channel DDR4-3200, single socket).&lt;p&gt;The theoretical performance, assuming everything besides serving the video is &amp;quot;free&amp;quot; is ~260-300Gb&amp;#x2F;s (depending on numa config, etc).&lt;p&gt;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), &lt;i&gt;and&lt;/i&gt; by eliminating a CPU stall on DRAM access.&lt;p&gt;I&amp;#x27;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-&amp;gt;Gen5, 100GbE -&amp;gt; 400GbE), so we&amp;#x27;ll still be fighting the same battle.&lt;p&gt;Things like inline-kTLS (where the NIC does the TLS encryption) reduce the memory bandwidth requirements by roughly half, and is how we&amp;#x27;re able to serve up over 350Gb&amp;#x2F;s from a single socket.</text></comment>
<story><title>Memory Bandwidth</title><url>https://fgiesen.wordpress.com/2017/04/11/memory-bandwidth/</url></story><parent_chain></parent_chain><comment><author>dragontamer</author><text>This is a good start to the discussion.&lt;p&gt;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).&lt;p&gt;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&amp;#x2F;write to it at high speeds, so it serves as a good area for inter-thread communications. (&amp;quot;Thread&amp;quot; is a bit ambiguous, but... lets just simplify it and call them cuda-threads).&lt;p&gt;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).&lt;p&gt;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 &amp;quot;core&amp;quot;). I kid you not: AMD RDNA has more registers per WGP than Intel Skylake had L1 cache per core!!&lt;p&gt;---------------&lt;p&gt;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.</text></comment>