GPU Memory and Interconnects: HBM, GDDR, NVLink, PCIe, CXL

A modern GPU is usually limited by how fast it moves bytes, not by how fast it multiplies. High Bandwidth Memory (HBM) stacks DRAM beside the GPU on a 1,024 or 2,048 bit bus; GDDR runs a narrow bus very fast on the board. PCIe, NVLink, NVSwitch, Infinity Fabric and CXL link GPUs to CPUs and to each other.

This post explains each layer from the DRAM die to the rack, with the bandwidth arithmetic worked out and checked in Python, and then uses the roofline model to show when a kernel is limited by memory and when it is limited by compute. Every specification below comes from a vendor document, a standards body announcement or a peer reviewed paper, linked in the Sources section. Where a vendor has not published a figure, we say so rather than guess. The code samples follow vendor documentation but were not executed for this article: we wrote it on a Mac with no NVIDIA or AMD data center GPU attached.

Key takeaways

• Bandwidth is bus width times data rate. One HBM3 stack is 1,024 bits at up to 6.4 Gb/s per pin, which is 819.2 GB/s. One HBM4 stack doubles the width to 2,048 bits; at 8 Gb/s that is 2,048 GB/s, which JEDEC rounds to 2 TB/s.
• HBM wins on width, GDDR wins on cost and simplicity. HBM needs a silicon interposer or bridge and advanced packaging such as TSMC's CoWoS. GDDR sits on an ordinary printed circuit board and gets its speed from signalling tricks: PAM4 in GDDR6X and PAM3 in GDDR7.
• PCIe is the slow lane by design. A PCIe 5.0 x16 link carries about 63 GB/s each way after encoding. An H100 SXM5 module reads its own HBM3 at 3.35 TB/s, which is over 50 times faster.
• NVLink and NVSwitch exist because PCIe is too slow for GPU to GPU traffic. NVIDIA lists 900 GB/s per GPU for fourth generation NVLink (Hopper), 1,800 GB/s for the fifth (Blackwell) and a preliminary 3,000 GB/s for the sixth (Vera Rubin). A GB200 NVL72 rack connects 72 GPUs with 130 TB/s of total NVLink bandwidth.
• AMD uses Infinity Fabric links. An MI300X has seven x16 Infinity Fabric links to its peers in an 8 GPU board, 896 GB/s in total by AMD's figures, plus one PCIe 5.0 x16 link to the host.
• CXL is coherent memory over PCIe wires. It matters more for CPU memory expansion and pooling than for feeding GPU tensor cores today.
• Unified memory is a programming model, not free bandwidth. On software coherent systems a page fault migrates data across PCIe; on hardware coherent systems such as Grace Hopper, the GPU can read CPU memory over a 900 GB/s NVLink-C2C link without migrating it.
• The roofline model tells you which wall you hit. On an H100 SXM5, a kernel needs about 295 FP16 operations per byte of HBM traffic before tensor cores, not memory, become the limit. A large matrix multiply clears that easily; a matrix vector product used in LLM token generation does not.

Why memory and interconnect now decide GPU performance

Peak arithmetic has grown much faster than the ability to feed it. Take NVIDIA's own figures. The Volta V100 whitepaper lists 900 GB/s of HBM2 bandwidth. The H100 SXM5 product page lists 3.35 TB/s of HBM3 and 1,979 teraFLOPS of FP16 tensor throughput with sparsity, which NVIDIA states is double the dense figure, so about 989.5 dense teraFLOPS. Bandwidth grew about 3.7 times from V100 to H100; dense tensor throughput grew far more. That gap is the reason the memory system now gets as much engineering attention as the shader cores.

The same pattern repeats one level up. A single GPU no longer holds the largest models, so GPUs must exchange activations, gradients and key value caches with each other many times per step. If that exchange crosses a PCIe link, it runs at tens of gigabytes per second. If it crosses NVLink or Infinity Fabric, it runs at hundreds. That difference often decides whether adding a GPU speeds a job up at all.

So when people say a GPU is "memory bound", they usually mean one of three different things:

1. HBM or GDDR bound: the kernel streams data from device memory faster than it can use it, or rather it cannot get data from device memory fast enough to keep the arithmetic units busy.
2. Interconnect bound: the job waits on data from another GPU (NVLink, Infinity Fabric, network) or from the host (PCIe, NVLink-C2C).
3. Capacity bound: the model, the batch or the cache does not fit, so the job must shard, offload or recompute.

The rest of this post takes each of these in turn.

The memory hierarchy from register to rack

Before the details, here is the whole path a byte can take, with the figures we could source. Sizes and bandwidths are vendor peak figures, not sustained measurements.

                       FASTEST, SMALLEST
  +--------------------------------------------------------------+
  | Registers (per SM)          private to each thread           |
  | Shared memory / L1 (per SM) shared by a thread block         |
  +--------------------------------------------------------------+
  | L2 cache (whole GPU)        H100: 50 MB   A100: 40 MB        |
  | AMD Infinity Cache          MI300X: 256 MB, 17.2 TB/s peak   |
  +--------------------------------------------------------------+
  | Device memory (HBM / GDDR)  H100 SXM5: 80 GB HBM3, 3.35 TB/s |
  |                             MI300X:  192 GB HBM3, 5.3 TB/s   |
  |                             RTX 4090: 24 GB GDDR6X, 1008 GB/s|
  +--------------------------------------------------------------+
  | Peer GPU memory             NVLink 4: 900 GB/s per GPU       |
  |  (same node)                MI300X Infinity Fabric: 896 GB/s |
  +--------------------------------------------------------------+
  | Host CPU memory             PCIe 5.0 x16: 128 GB/s (both ways)|
  |                             NVLink-C2C (GH200): 900 GB/s     |
  +--------------------------------------------------------------+
  | Other nodes                 network fabric (out of scope)    |
  +--------------------------------------------------------------+
                       SLOWEST, LARGEST

The L2 sizes come from NVIDIA's Ampere and Hopper documents; the Infinity Cache figures from AMD's CDNA 3 whitepaper; device memory from the H100, MI300X and Ada documents; link figures from the NVLink, CDNA 3 and Grace Hopper pages. On-chip registers and shared memory are covered in the CUDA programming model post in this series; here we start at the L2 and work outward.

HBM explained: stacks, TSVs, base dies, interposers and CoWoS

How an HBM stack is built

HBM is ordinary DRAM technology rearranged in three dimensions. Several DRAM dies are thinned, stacked on top of a logic base die and connected vertically by through silicon vias (TSVs), which are copper filled holes etched through each die. JEDEC's announcement for the JESD235A update (January 2016) describes HBM as using "Wide I/O and TSV technologies". The stack talks to the GPU through a very wide interface: 1,024 data bits per stack for HBM2, HBM3 and HBM3E, and 2,048 bits for HBM4. Micron's HBM3E page lists 1,024 IO pins per placement.

A bus that wide cannot run across a normal circuit board. Thousands of wires at micrometre pitch need a substrate with silicon level wiring density. So the GPU die and its HBM stacks sit side by side on an interposer, a slab of silicon (or organic material with silicon bridges) that carries the short, dense wires between them.

   Side view of one GPU package with HBM (not to scale)

     HBM stack            GPU die                 HBM stack
   +-----------+   +-----------------------+   +-----------+
   | DRAM die  |   |                       |   | DRAM die  |
   | DRAM die  |   |   SMs / CUs, L2,      |   | DRAM die  |
   | DRAM die  |   |   memory controllers  |   | DRAM die  |
   | DRAM die  |   |                       |   | DRAM die  |
   | base die  |   |                       |   | base die  |
   +--||||||---+   +---|||||||||||||||||---+   +--||||||---+
   ===============  interposer (dense wiring)  ===============
   ------------------- package substrate ---------------------
   ...................... board (BGA balls) ..................

   |||| = microbumps     inside each stack, TSVs run vertically

Each stack is split into independent channels so that many requests can be in flight at once. JEDEC's HBM3 announcement (January 2022) doubled independent channels from 8 in HBM2 to 16, each with two pseudo channels, for 32 in total. The HBM4 announcement (April 2025) doubles again to 32 channels per stack, each with two pseudo channels. HBM3 also added on-die, symbol based ECC and real time error reporting.

The interposer and CoWoS

The interposer is where much of the cost and supply constraint sits. TSMC's name for its family is CoWoS (Chip on Wafer on Substrate). Per TSMC's 3DFabric material, CoWoS-S uses a silicon interposer that can reach about 3.3 times the reticle size (about 2,700 square millimetres). CoWoS-R uses an organic redistribution layer interposer, and CoWoS-L combines redistribution layers with small local silicon interconnect bridges where dense wiring is needed, which TSMC recommends above 3.3 times reticle size. TSMC states it completed certification of interposers at 5.5 times reticle size in 2025, with volume production planned for 2026. We could not fetch TSMC's CoWoS page directly (it blocks scripted clients), so these figures come from TSMC's text as indexed by search; treat them as TSMC's statements, not independent measurements.

The reason interposer area matters is simple geometry. Each extra HBM stack needs space next to the GPU die. NVIDIA's Blackwell architecture page states that every Blackwell product uses two reticle limited dies joined by a 10 TB/s chip to chip link, presented to software as one GPU. Two large dies plus eight HBM stacks need an interposer several reticles in size. The companion post how a GPU is made, from wafer to HBM package walks through that packaging flow step by step.

HBM generations at a glance

Two notes on that table. First, "HBM2E" and "HBM3E" are names the memory makers use for faster parts within or beyond a JEDEC generation; JEDEC's press releases we found title the documents HBM, HBM2 (as a JESD235 update), HBM3 and HBM4. Second, the HBM4 stack height and die density detail appears in only some coverage of JEDEC's release, so treat it as reported rather than confirmed until you read JESD270-4 itself, which JEDEC offers as a download after registration.

Worked example: bandwidth per stack

Bandwidth per stack is the interface width in bits times the per pin data rate, divided by eight to get bytes:

# Peak bandwidth per HBM stack, GB/s = width_bits * Gbps_per_pin / 8
stacks = {
    "HBM2 (JESD235A)":   (1024, 2.0),
    "HBM3 (JESD238)":    (1024, 6.4),
    "HBM3E at 9.2 Gb/s": (1024, 9.2),
    "HBM4 (JESD270-4)":  (2048, 8.0),
}
for name, (width, gbps) in stacks.items():
    print(f"{name:20s} {width * gbps / 8:8.1f} GB/s")

Output, checked with python3:

HBM2 (JESD235A)         256.0 GB/s
HBM3 (JESD238)          819.2 GB/s
HBM3E at 9.2 Gb/s      1177.6 GB/s
HBM4 (JESD270-4)       2048.0 GB/s

The HBM3 result matches JEDEC's 819 GB/s, and HBM4 matches its "up to 2 TB/s". The HBM3E line shows something useful: 9.2 Gb/s on 1,024 pins gives 1,177.6 GB/s, which is just under Micron's "more than 1.2 TB/s". Working backwards, 1.2 TB/s needs about 9.4 Gb/s per pin, which is consistent with Micron's wording of "greater than 9.2 Gbps". Doing this arithmetic is a good habit; it catches marketing figures that round in your favour.

Worked example: what shipping GPUs actually run

Vendors publish total bandwidth and bus width, so you can derive the per pin rate each product runs its memory at. The width and bandwidth columns below are the vendors' figures; the last column is our arithmetic.

AMD's MI300X page does the arithmetic for you: an 8,192 bit bus at 5.2 Gb/s gives 5.325 TB/s. Our python3 check gives 5,324.8 GB/s. Notice that both NVIDIA's A100 and H100 dies have six HBM sites, but the products enable five; NVIDIA's Ampere whitepaper shows six stacks and twelve 512 bit controllers on the full GA100 and five stacks and ten controllers on the A100 product. Disabling a stack is a yield and cost decision, and it is why the bus is 5,120 bits rather than 6,144.

Capacity follows the same pattern. The H100 SXM5 has 80 GB across five stacks, so 16 GB per stack. The MI300X has 192 GB across eight stacks, so 24 GB per stack. HBM3's ceiling of 64 GB per stack, from JEDEC's announcement, is 16 dies of 32 Gb each: 16 times 32 Gb is 512 Gb, which is 64 GB.

Newer parts push both numbers. NVIDIA's DGX B200 page lists 8 Blackwell GPUs with 1,440 GB of HBM3e and 64 TB/s in total, which is 180 GB and 8 TB/s per GPU. The GB200 NVL72 specification table lists 13.4 TB of HBM3E and 576 TB/s across 72 GPUs, which is about 186 GB and exactly 8 TB/s per GPU. The two pages give slightly different capacities per GPU; we report each as published rather than pick one.

GDDR6, GDDR6X and GDDR7: speed on an ordinary circuit board

Why GDDR exists

GDDR (graphics double data rate) memory takes the opposite approach to HBM. Each GDDR device is a separate package soldered onto the graphics card around the GPU, with a narrow 32 bit interface. No interposer is needed. To get high bandwidth, GDDR runs each pin very fast and the GPU uses many devices in parallel. The trade-off is power and signal integrity: pushing tens of gigabits per second down a few centimetres of board trace is hard, and the energy per bit is higher than for HBM's short interposer wires.

Micron's GDDR6X page gives a clear example of how the totals add up. Its comparison table lists 12 placements (devices), 19 to 24 Gb/s per pin, 76 to 96 GB/s per placement and 912 to 1,152 GB/s per system on a 384 bit configuration. The arithmetic checks: 32 bits times 24 Gb/s is 96 GB/s per device, and 12 devices make 1,152 GB/s.

NRZ, PAM4 and PAM3 signalling

The big change in recent GDDR generations is how many bits each symbol carries.

• NRZ (non return to zero) uses two voltage levels, so one bit per symbol. GDDR6 uses NRZ.
• PAM4 (four level pulse amplitude modulation) uses four levels, so two bits per symbol. Micron introduced it in GDDR6X; its GDDR6X page describes "PAM4 signal techniques".
• PAM3 uses three levels. JEDEC's GDDR7 announcement (JESD239, 5 March 2024) calls GDDR7 the first JEDEC standard DRAM to use pulse amplitude modulation, and says PAM3 improves the signal to noise ratio for high frequency operation. Coverage of the release describes the encoding as 3 bits over 2 cycles.

   Voltage levels per symbol (idealised eye diagrams)

   NRZ (GDDR6)        PAM3 (GDDR7)        PAM4 (GDDR6X, PCIe 6/7)
   ----- 1            ----- +1            ----- 11
                      ----- 0             ----- 10
   ----- 0            ----- -1            ----- 01
                                          ----- 00
   1 bit/symbol       1.5 bits/symbol     2 bits/symbol
   widest eye         wider than PAM4     narrowest eye

Why would GDDR7 choose three levels when four carry more bits? Because the eye (the gap between adjacent levels) gets smaller as you add levels, and a smaller eye is harder to read without errors at high speed. Three levels split the swing into two gaps instead of three, so each gap is larger. Two PAM3 symbols have nine combinations, enough to encode three bits (eight combinations) with one to spare: 1.5 bits per symbol. The theoretical maximum for three levels is log base 2 of 3, about 1.585 bits; we checked that with python3. GDDR7 trades a little of PAM4's density for a cleaner signal it can then run faster.

GDDR generations at a glance

Worked example: two GeForce cards

NVIDIA's Ada architecture whitepaper lists the RTX 4090 with 24 GB of GDDR6X at 21 Gbps on a 384 bit interface and 1,008 GB/s. The check: 384 times 21 divided by 8 is 1,008 GB/s.

For the RTX 5090, NVIDIA's product page lists 32 GB of GDDR7 on a 512 bit interface but, when we fetched it, did not list the data rate or bandwidth. NVIDIA's launch figure as reported by several outlets is 28 Gbps and 1,792 GB/s, and the arithmetic agrees: 512 times 28 divided by 8 is 1,792 GB/s. Because we could not find that figure on an NVIDIA page, treat it as reported rather than vendor confirmed.

The GDDR7 per device figure works the same way. A 32 bit device at 192 GB/s implies 48 Gb/s per pin: 32 times 48 divided by 8 is 192. That is the ceiling JEDEC wrote into the standard, not what any particular card ships at today.

HBM or GDDR: the trade-off

Neither is "better". HBM gives more bandwidth and capacity per square millimetre of package at a higher cost and with a constrained packaging supply. GDDR gives respectable bandwidth with simpler, cheaper manufacturing. That is why one vendor can ship both: HBM on the data center parts and GDDR on the graphics cards.

A reusable bandwidth calculator

Every link and memory in this article follows the same two formulas:

• Raw bandwidth = lanes (or pins) times transfer rate per lane, divided by 8.
• Usable bandwidth = raw bandwidth times encoding efficiency, minus protocol overhead.

Here is a small Python script that reproduces every derived figure in this article. It only does arithmetic; it does not measure anything on your machine.

#!/usr/bin/env python3
"""Peak bandwidth arithmetic for GPU memory and links.
All inputs are vendor or standards body figures cited in the article.
Results are theoretical peaks, never measurements."""

def mem_gbs(width_bits, gbps_per_pin):
    return width_bits * gbps_per_pin / 8

def pcie_gbs_per_dir(gt_per_s, lanes=16, encoding=128/130):
    # PCIe 3.0 to 5.0 use 128b/130b encoding. PCIe 6.0 and 7.0 use PAM4
    # with FLIT mode; we report raw rate for those (encoding=1.0).
    return gt_per_s * lanes * encoding / 8

print("RTX 4090 GDDR6X:", mem_gbs(384, 21), "GB/s")
print("MI300X HBM3:    ", mem_gbs(8192, 5.2), "GB/s")
for gen, rate in [(3, 8), (4, 16), (5, 32)]:
    print(f"PCIe {gen}.0 x16:", round(pcie_gbs_per_dir(rate), 2), "GB/s each way")
for gen, rate in [(6, 64), (7, 128)]:
    print(f"PCIe {gen}.0 x16 raw:", pcie_gbs_per_dir(rate, encoding=1.0), "GB/s each way")

Output, checked with python3:

RTX 4090 GDDR6X: 1008.0 GB/s
MI300X HBM3:     5324.8 GB/s
PCIe 3.0 x16: 15.75 GB/s each way
PCIe 4.0 x16: 31.51 GB/s each way
PCIe 5.0 x16: 63.02 GB/s each way
PCIe 6.0 x16 raw: 128.0 GB/s each way
PCIe 7.0 x16 raw: 256.0 GB/s each way

The PCIe 4.0 line agrees with NVIDIA's Ampere whitepaper, which quotes 31.5 GB/s for a Gen 4 x16 connection against 15.75 GB/s for Gen 3.

PCIe: the host link every GPU has

What PCIe does for a GPU

Every discrete GPU, including the ones that also have NVLink or Infinity Fabric, connects to its host CPU over PCI Express. That link carries the driver's commands, the copies your program makes with cudaMemcpy or hipMemcpy, and page migrations under unified memory. On many servers it also carries traffic to network cards and NVMe drives that the GPU reads directly.

PCIe generations

PCI-SIG's headline x16 numbers count both directions together. That is why its PCIe 6.0 release says "up to 256 GB/s" while our calculator says 128 GB/s each way. Always check which convention a figure uses before you compare it with a memory bandwidth, which is usually quoted in one direction (a read stream).

PCIe 6.0 is the first generation to use PAM4, the same four level signalling as GDDR6X. Per PCI-SIG's release, the higher error rate that PAM4 brings is handled with lightweight forward error correction (FEC) and a cyclic redundancy check, and the link moves to FLIT mode, where data travels in fixed size flow control units rather than variable length packets. PCIe 7.0 keeps PAM4 and FLIT and doubles the rate again. PCI-SIG's site blocks scripted clients, so we confirmed these releases through syndicated copies of PCI-SIG's own press releases rather than its web pages.

Worked example: how long does PCIe take?

Say you need to load 80 GB of model weights from host memory into an H100's HBM. At PCIe 5.0 x16's 64 GB/s per direction (NVIDIA's H100 page lists 128 GB/s total for PCIe Gen5), the transfer takes at least 80 divided by 64, which is 1.25 seconds. The same data read once from HBM at 3.35 TB/s takes about 24 milliseconds. That ratio, about 50 to 1, is why frameworks try hard to load data once and keep it on the GPU, and why pinned (page locked) host memory matters: the GPU's copy engines can only reach full PCIe speed from memory the operating system will not page out.

Checking what your GPU actually negotiated

A GPU in the wrong slot, behind a slow riser or a power saving state can negotiate a lower link generation or width. These read-only commands show what the link trained to. They follow vendor documentation; we could not run them here.

# NVIDIA: current and maximum PCIe generation and width
nvidia-smi --query-gpu=name,pcie.link.gen.current,pcie.link.gen.max,pcie.link.width.current,pcie.link.width.max --format=csv

# Any vendor on Linux: find the GPU, then read the link capability and status
lspci | grep -i -E "vga|3d|display"
sudo lspci -vv -s <bus:device.function> | grep -E "LnkCap|LnkSta"

# AMD ROCm: bus information per GPU
rocm-smi --showbus

Note that many GPUs drop to a lower PCIe generation when idle to save power, so read the current values while a workload is running.

NVLink: GPU to GPU without the host

Why NVLink exists

PCIe is a general purpose bus designed around a host. When two GPUs on the same board need to swap gigabytes many times per second, routing that traffic through PCIe switches and the CPU's root complex costs bandwidth and latency. NVLink is NVIDIA's point to point link for GPU to GPU (and on some systems GPU to CPU) traffic.

NVLink generations

The totals are simple products, and our python3 check reproduces each one: 4 times 40, 6 times 50, 12 times 50 and 18 times 50 give 160, 300, 600 and 900 GB/s. NVIDIA's NVLink page lists the fifth and sixth generations directly; we did not find per link figures for those on the pages we fetched, so we have not derived them. The sixth generation figures are labelled "preliminary specifications; may be subject to change" on NVIDIA's own page.

All of these totals count both directions. The Hopper in-depth blog spells this out for NVLink 4: each link carries 25 GB/s in each direction over two differential pairs per direction, and 18 links make 900 GB/s. So the one way rate between two H100s connected directly through all their links is 450 GB/s.

Comparing NVLink with PCIe fairly

NVIDIA's NVLink page compares the sixth generation with PCIe Gen6 and calls it "12x the bandwidth". With both sides counted the same way (3,000 GB/s against PCI-SIG's 256 GB/s for a Gen6 x16 link), the ratio is 11.7, which rounds to 12. For the fifth generation, 1,800 against PCIe Gen5's 128 GB/s is about 14 times. These comparisons are fair only because both figures use the same both directions convention.

NVLink-C2C: the CPU joins the club

NVLink-C2C is a chip to chip variant that connects NVIDIA's Grace CPU to a Hopper or Blackwell GPU in the same module. NVIDIA's Grace Hopper page describes "900 gigabytes per second (GB/s) of coherent interface" and says that is "7X faster than PCIe Gen5"; 900 divided by 128 is 7.03, so the claim checks against PCIe's both directions figure. The word coherent is the important one, and we come back to it under unified memory.

NVSwitch, NVLink domains and NVL72 racks

From point to point to a switch

With point to point links, each GPU can only spend its links on a few neighbours. NVSwitch is a switch chip that lets every GPU in a domain reach every other GPU at full NVLink bandwidth. The Hopper in-depth blog describes the third generation NVSwitch used in H100 systems: 64 fourth generation NVLink ports per switch and 13.6 Tbit/s of total switch throughput. It also adds two features that matter for distributed training:

• Hardware multicast, so one GPU can write the same data to many peers in one operation.
• SHARP in-network reductions, which perform parts of collective operations such as all-reduce inside the switch rather than on the GPUs' SMs. NVIDIA's blog reports up to 2 times throughput for small block collectives compared with NCCL on A100. That is NVIDIA's figure, not one we can verify.

Domain sizes

NVIDIA's NVLink page lists the NVLink Switch "GPU domain" sizes and totals:

These totals are per GPU bandwidth times GPU count: 8 times 0.9 TB/s is 7.2 TB/s; 72 times 1.8 TB/s is 129.6 TB/s, which NVIDIA rounds to 130; 72 times 3 TB/s is 216 TB/s. Our python3 check reproduces each.

What a GB200 NVL72 rack contains

NVIDIA's GB200 NVL72 specification table lists 36 Grace CPUs and 72 Blackwell GPUs in one rack scale system, with 13.4 TB of HBM3E at 576 TB/s, 130 TB/s of NVLink bandwidth and 17 TB of LPDDR5X CPU memory at 14 TB/s. Each GB200 superchip pairs one Grace CPU with two Blackwell GPUs over NVLink-C2C.

   GB200 NVL72, logical view (per NVIDIA's specification table)

   +---------------------+   +---------------------+         +---------------------+
   | GB200 superchip 1   |   | GB200 superchip 2   |  ...    | GB200 superchip 36  |
   | Grace CPU           |   | Grace CPU           |         | Grace CPU           |
   |   | NVLink-C2C      |   |   | NVLink-C2C      |         |   | NVLink-C2C      |
   | GPU 1     GPU 2     |   | GPU 3     GPU 4     |         | GPU 71    GPU 72    |
   +--|---------|--------+   +--|---------|--------+         +--|---------|--------+
      |         |               |         |                     |         |
   ===================== NVLink Switch fabric (fifth generation) ===================
     every GPU reaches every other GPU at 1.8 TB/s; 72 GPUs, 130 TB/s in total

   Memory: 72 x HBM3E (13.4 TB, 576 TB/s)  +  36 x LPDDR5X (17 TB, 14 TB/s)

Two things stand out. First, the HBM bandwidth across the rack (576 TB/s) is more than four times the NVLink total (130 TB/s), so even inside an NVL72 a GPU reads its own memory much faster than a neighbour's. Locality still matters. Second, the CPU memory (17 TB) is larger than the HBM (13.4 TB) but much slower (14 TB/s against 576 TB/s). That is the capacity versus bandwidth trade-off again, now at rack scale.

Software you need for NVSwitch systems

On HGX and DGX class servers with NVSwitch, NVIDIA's Fabric Manager service configures the switch fabric before GPUs can use it, according to NVIDIA's Fabric Manager user guide. Its version rules have changed over time: the guide states that for driver releases earlier than 620 the Fabric Manager packages must be the same version as the driver package, and that starting with driver release 620 Fabric Manager is versioned standalone. Either way it is one more component to track in the server's software inventory, which we return to in the Safeguard section.

AMD Infinity Fabric and xGMI

AMD's equivalent of NVLink for its Instinct accelerators is Infinity Fabric; AMD's ROCm documentation and tools call the GPU to GPU links xGMI. AMD's CDNA 3 whitepaper describes the MI300X and MI325X using seven high bandwidth, low latency Infinity Fabric links to form a fully connected 8 GPU system, with each GPU also connected to the host CPU by a PCIe 5.0 x16 link. The whitepaper's table gives 7 x16 Infinity Fabric links and 896 GB/s for the 8 GPU configuration. AMD's MI300X product page lists a peak of 128 GB/s per Infinity Fabric link; 7 times 128 is 896, which matches. AMD's ROCm blog on RCCL and xGMI (2 March 2025) gives the same link as 64 GB/s in each direction, theoretical, and reports 45 to 48 GB/s realised per link per direction in AMD's own tests, after protocol overhead and CRC. That gap between a vendor peak and a vendor measurement, roughly 25 to 30 percent, is a useful reminder of how far peak figures sit from what software sees.

   MI300X 8-GPU board: fully connected mesh (per AMD CDNA 3 whitepaper)

          G0 ----- G1
         /| \     / |\
        / |  \   /  | \        each GPU: 7 Infinity Fabric links,
      G7  |   \ /   |  G2      one to every other GPU, 128 GB/s each
       |\ |    X    | /|
       | \|   / \   |/ |       plus 1 x PCIe 5.0 x16 to the host CPU
      G6  |  /   \  |  G3
        \ | /     \ | /        (not every link is drawn)
         \|/       \|/
          G5 ----- G4

The design differs from an NVSwitch domain. In a full mesh with no switch, each pair of GPUs has one dedicated link. A transfer between two specific GPUs uses that one link, 128 GB/s, unless software spreads it over multi-hop paths; an all-to-all pattern, where every GPU talks to every other at once, uses all seven links at the same time. In an NVSwitch domain, any single pair can use a GPU's whole NVLink bandwidth. Which design suits a workload depends on its communication pattern, which is why collective libraries (NCCL on NVIDIA, RCCL on AMD) pick algorithms per topology.

To see the topology on a running system, these read-only commands follow AMD's and NVIDIA's documentation (not run here):

# AMD: link type (XGMI or PCIE) and hop count between GPUs
rocm-smi --showtopo        # deprecated in favour of amd-smi
amd-smi topology

# NVIDIA: matrix of NVLink (NV#), PCIe switch and CPU paths between GPUs
nvidia-smi topo -m

AMD's MI300X system optimisation guide explains the hop values: 1 means a direct xGMI connection, 2 means the GPUs share a CPU socket and 3 means they sit on different sockets.

Inside the package, AMD uses Infinity Fabric too. The CDNA 3 whitepaper describes eight HBM stacks across four I/O dies, with a 256 MB Infinity Cache in front of them that AMD rates at a peak of 17.2 TB/s. On the MI300A APU, AMD states that HBM3 is unified and shared between the CPU and GPU cores in the same package, which is AMD's route to the same goal as NVLink-C2C: one memory both kinds of processor can use without copies.

CXL: coherent memory over PCIe wires

What CXL is

Compute Express Link (CXL) is an open standard from the CXL Consortium that builds on PCIe and adds cache coherent protocols on top. The consortium's CXL 3.0 white paper describes CXL as multiplexing three kinds of traffic: I/O (CXL.io, based on PCIe), caching (CXL.cache) and memory (CXL.memory), and says CXL keeps a unified, coherent memory space between the CPU and memory on the attached device. Devices are commonly grouped into three types: Type 1 devices cache host memory but have none of their own, Type 2 devices are accelerators with their own memory (the white paper calls these "commonly known as CXL Type2 devices"), and Type 3 devices are memory expanders.

CXL versions

The CXL Consortium's own pages state that CXL 4.0 "doubles the bandwidth to 128GTs with zero added latency" and keeps full backward compatibility with CXL 3.x, 2.0, 1.1 and 1.0. We did not find a consortium primary source for the CXL 3.1 date, so it is not in the table.

Where CXL fits for GPUs, honestly

CXL is often mentioned in the same breath as NVLink, but they solve different problems today. NVLink and Infinity Fabric move data between accelerators at hundreds of gigabytes per second per GPU. A CXL link is at most a PCIe link's bandwidth: 64 GT/s or 128 GT/s per lane. Its strength is coherence and capacity: a server can add terabytes of DRAM through Type 3 memory expanders, or pool memory across hosts with CXL 2.0 and later switching. For GPU work, that makes CXL most relevant to the CPU side of the system (bigger host memory for staging data, KV cache offload or preprocessing) rather than as a replacement for HBM. Whether a specific GPU can act as a CXL Type 2 device is a product question; none of the GPU product pages we fetched for this article list CXL support, so check the vendor's documentation for the part you are considering.

Collectives: the arithmetic of GPUs talking to each other

Most multi GPU traffic in training is collective operations: all-reduce for gradients, all-gather and reduce-scatter for sharded parameters. NVIDIA's nccl-tests documentation explains how to reason about all-reduce. With n ranks and S bytes, an all-reduce needs 2(n minus 1) data transfers per element, and the time on n links of bandwidth B is:

t = (S * 2 * (n - 1)) / (n * B)

So each GPU sends and receives about 2(n minus 1)/n times S bytes. The nccl-tests documentation calls the rate normalised by that factor "bus bandwidth", so results can be compared with link speed.

Worked example: an all-reduce over NVLink versus PCIe

Take 16 GB of gradients (for example 8 billion parameters in BF16) across 8 GPUs. Each GPU moves 2 times 7 divided by 8 times 16 GB, which is 28 GB.

These are lower bounds from peak link rates, checked with python3, not benchmark results. Real all-reduces add latency per step, protocol overhead and contention, and NCCL picks tree or ring algorithms (and, on NVSwitch systems, in-network reduction) depending on message size. The point of the exercise is the ratio: about seven times, which is the gap between the two link types. If your training step spends more time in all-reduce than in compute, the interconnect, not the GPU, sets your speed.

Unified memory and managed memory

The idea

Unified memory gives the CPU and GPU one address space, so a pointer means the same thing on both. It removes explicit cudaMemcpy calls from your code. It does not remove the physics: if the data sits in host DRAM and the kernel runs on the GPU, the bytes still have to cross a link.

CUDA's unified memory paradigms

The current CUDA Programming Guide groups systems into paradigms:

• Full unified memory support with hardware coherence. The guide's example is NVIDIA Grace Hopper, which has a logically combined page table for CPU and GPU. The GPU can access CPU memory directly over NVLink-C2C without a page fault, and system allocated memory (plain malloc) works in kernels.
• Full unified memory support with software coherence. The guide's example is modern Linux with Heterogeneous Memory Management (HMM). The guide states the requirements: Linux kernel 6.1.24 or later (or 6.2.11 or later, or 6.3 or later), a GPU of compute capability 7.5 or higher, and a CUDA driver 535 or later installed with the open kernel modules. Here CPU and GPU have separate page tables, and a GPU access to a CPU resident page faults and migrates it.
• CUDA managed memory only. Memory from cudaMallocManaged migrates on demand between host and device.
• Limited unified memory. The guide lists devices below compute capability 6.0, Windows (including WSL) and Tegra, where concurrentManagedAccess is 0, with extra restrictions on concurrent CPU and GPU access.

Your program can query the relevant device attributes (pageableMemoryAccess, concurrentManagedAccess, directManagedMemAccessFromHost) at run time and pick a strategy.

Example: managed memory with prefetch

The following follows the patterns in the CUDA Programming Guide's unified memory chapter. We did not compile or run it.

#include <cuda_runtime.h>
#include <cstdio>

__global__ void scale(float* x, float a, size_t n) {
    size_t i = blockIdx.x * (size_t)blockDim.x + threadIdx.x;
    if (i < n) x[i] *= a;
}

int main() {
    const size_t n = 1ull << 28;                 // 268,435,456 floats = 1 GiB
    float* x = nullptr;
    cudaMallocManaged(&x, n * sizeof(float));    // one pointer for CPU and GPU

    for (size_t i = 0; i < n; ++i) x[i] = 1.0f;  // pages first touched on the CPU

    int dev = 0;
    cudaStream_t s;
    cudaStreamCreate(&s);

    // Move the pages to the GPU before the kernel runs, instead of
    // paying a page fault and migration on first touch.
    cudaMemLocation gpu = {};
    gpu.type = cudaMemLocationTypeDevice;
    gpu.id = dev;
    cudaMemPrefetchAsync(x, n * sizeof(float), gpu, 0, s);

    const unsigned threads = 256;
    const unsigned blocks = (unsigned)((n + threads - 1) / threads);
    scale<<<blocks, threads, 0, s>>>(x, 2.0f, n);

    // Bring the result back to the CPU before reading it there.
    cudaMemLocation host = {};
    host.type = cudaMemLocationTypeHost;
    cudaMemPrefetchAsync(x, n * sizeof(float), host, 0, s);
    cudaStreamSynchronize(s);

    printf("x[0] = %f\n", x[0]);
    cudaFree(x);
    cudaStreamDestroy(s);
    return 0;
}

A version note: the current CUDA Programming Guide documents cudaMemPrefetchAsync taking a cudaMemLocation and a flags argument. Older CUDA 12 code passes a device number instead (for example cudaMemPrefetchAsync(x, bytes, dev, s)). Check the signature in the toolkit version you build with; this is one of many small reasons to pin and record the CUDA toolkit version in your build.

What the numbers say about prefetching

That 1 GiB buffer crosses PCIe 5.0 x16 in at least about 17 milliseconds at 64 GB/s each way. The kernel itself reads and writes 1 GiB from HBM, about 2.1 GB of traffic, which at 3.35 TB/s is about 0.6 milliseconds. Without the prefetch, the first kernel launch pays for the migration in small page sized pieces, each with fault handling overhead, so it is usually slower than one bulk prefetch. With hardware coherence on Grace Hopper, the GPU can instead read the host copy directly over NVLink-C2C; the CUDA guide notes that for data both processors use heavily, explicit copies or cudaMemAdvise hints such as preferred location can still beat leaving it to the defaults.

Unified memory outside NVIDIA

• AMD MI300A: per the CDNA 3 whitepaper, CPU and GPU share the same HBM3 in one package, so there is no host to device copy at all.
• Apple silicon: Apple's M4 family announcement (30 October 2024) describes "unified memory" shared by CPU and GPU, lists 120 GB/s of memory bandwidth for M4 and "over half a terabyte per second" for M4 Max, with up to 128 GB of unified memory on M4 Max. That is far less bandwidth than an HBM accelerator, but far more capacity per GPU than most graphics cards, which is why such machines can hold large models that would not fit in 24 or 32 GB of GDDR.

The roofline model, worked

The model in one formula

The roofline model, from Williams, Waterman and Patterson's paper in Communications of the ACM (volume 52, issue 4, April 2009), bounds the performance of a kernel by two numbers:

attainable FLOP/s = min( peak FLOP/s,  arithmetic intensity x peak bandwidth )

arithmetic intensity (AI) = FLOPs performed / bytes moved to and from memory

Plot attainable performance against arithmetic intensity on log scales and you get a sloped line (memory bound) that meets a flat line (compute bound). The meeting point is the ridge point:

ridge point = peak FLOP/s / peak bandwidth      (FLOPs per byte)

A kernel to the left of the ridge is memory bound; to the right, compute bound.

The ridge points of real GPUs

Using vendor peak figures, checked with python3:

These ridge points are high. A kernel must do roughly 250 to 600 floating point operations for every byte it reads from HBM before the tensor cores become the bottleneck. Very few kernels outside dense matrix multiply and convolution get there.

   Roofline for H100 SXM5, FP16 tensor (log-log, schematic)

   TFLOP/s
   1000 |                         ___________________________ 989.5 peak
        |                       /   * GEMM 8192^3 (AI ~2731)
        |                     /
    100 |                   /
        |                 /
     10 |               /
        |    * GEMV   /   slope = 3.35 TB/s
      1 |  (AI ~1)  /
        |         /
    0.1 |  * SAXPY (AI ~0.17, FP32 CUDA cores in practice)
        +-------+--------+--------+--------+--------+------> FLOP/byte
              0.1       1        10      100    295   1000
                                             ridge

Worked example 1: a large matrix multiply (compute bound)

Multiply two 8,192 by 8,192 FP16 matrices into an 8,192 by 8,192 result. Matrix multiply does 2 times M times N times K floating point operations. The minimum memory traffic is reading A and B once and writing C once, at 2 bytes per element.

M = N = K = 8192
flops = 2 * M * N * K                     # 1,099,511,627,776
bytes_moved = 2 * (M*K + K*N + M*N)       # 402,653,184
ai = flops / bytes_moved                  # about 2,730.7 FLOP/byte
t_compute_ms = flops / 989.5e12 * 1e3     # about 1.11 ms
t_memory_ms  = bytes_moved / 3.35e12 * 1e3  # about 0.12 ms

The arithmetic intensity of about 2,731 FLOP/byte is nine times the H100's ridge point, so this multiply is compute bound: at peak, the arithmetic takes about 1.11 ms and the minimum memory traffic only about 0.12 ms. In practice the kernel does not achieve the minimum traffic, because the matrices are far larger than on-chip memory and tiles of A and B are re-read. Good kernels choose tile sizes so that the real intensity stays well above the ridge; that is what the shared memory tiling and register blocking stages of a matmul optimisation are for. We do not quote an achieved TFLOP/s figure here, because we did not run it and will not invent one.

Worked example 2: a matrix vector product (memory bound)

Now multiply the same 8,192 by 8,192 FP16 matrix by a single vector. That is the core operation of generating one token in a large language model at batch size 1.

n = 8192
flops = 2 * n * n                         # 134,217,728
bytes_moved = 2 * n * n + 2 * n * 2       # matrix + vectors, about 134.25 MB
ai = flops / bytes_moved                  # about 1.0 FLOP/byte
t_us = bytes_moved / 3.35e12 * 1e6        # about 40 microseconds
attainable_tflops = 3.35 * ai             # about 3.35 TFLOP/s

The intensity is about 1 FLOP/byte, roughly 295 times below the ridge. The best the H100 can do is about 3.35 TFLOP/s on this operation, about 0.34 percent of its dense FP16 peak. No amount of tensor core cleverness helps; only more bandwidth, fewer bytes (quantised weights) or more work per byte (batching several requests so each weight read serves several tokens) moves the number.

Worked example 3: SAXPY (deeply memory bound)

y = a * x + y in FP32 does 2 FLOPs per element and moves 12 bytes (read x, read y, write y, 4 bytes each). That is an intensity of about 0.17 FLOP/byte, so the attainable rate on an H100 is at most 0.17 times 3.35 TB/s, about 0.56 TFLOP/s. The useful metric for a kernel like this is achieved bandwidth as a fraction of peak, not FLOP/s.

Worked example 4: an upper bound on LLM decode speed

Generating each token at batch size 1 requires reading every weight once. That gives a hard upper bound on tokens per second from memory bandwidth alone:

These are ceilings from arithmetic, checked with python3, not benchmarks. Real decode adds key value cache reads, attention, kernel launch gaps and framework overhead, so measured numbers are lower. Note also that a 140 GB model does not fit in an 80 GB H100 at all, which is the capacity wall rather than the bandwidth wall. Our post on estimating GPU cost per million tokens when self hosting turns this kind of ceiling into a cost estimate.

A roofline script you can adapt

#!/usr/bin/env python3
"""Roofline bound for a kernel. Inputs are vendor peak figures.
Prints a bound, not a prediction or a measurement."""

def roofline(flops, bytes_moved, peak_tflops, peak_tbs):
    ai = flops / bytes_moved
    ridge = peak_tflops / peak_tbs
    bound = min(peak_tflops, ai * peak_tbs)
    kind = "compute bound" if ai >= ridge else "memory bound"
    return ai, ridge, bound, kind

H100_FP16 = (989.5, 3.35)   # dense TFLOP/s, TB/s (NVIDIA H100 page)

for name, f, b in [
    ("GEMM 8192^3 FP16", 2 * 8192**3, 2 * 3 * 8192**2),
    ("GEMV 8192 FP16",   2 * 8192**2, 2 * 8192**2 + 4 * 8192),
]:
    ai, ridge, bound, kind = roofline(f, b, *H100_FP16)
    print(f"{name:18s} AI={ai:8.1f}  ridge={ridge:6.1f}  bound={bound:7.1f} TFLOP/s  {kind}")

Output, checked with python3:

GEMM 8192^3 FP16   AI=  2730.7  ridge= 295.4  bound=  989.5 TFLOP/s  compute bound
GEMV 8192 FP16     AI=     1.0  ridge= 295.4  bound=    3.3 TFLOP/s  memory bound

To place a real kernel on the chart, measure its FLOPs and DRAM bytes with a profiler. NVIDIA Nsight Compute includes a roofline chart in its reports, and AMD's ROCm tools offer similar counters; use the measured bytes, not the theoretical minimum, because caches and re-reads change the real intensity.

Hierarchical rooflines

The same reasoning applies to each level of the hierarchy. A kernel can be compute bound relative to HBM but bound by L2 or shared memory bandwidth, or bound by NVLink when it reads from a peer. Draw one sloped roof per level (L2, HBM, NVLink, PCIe) using that level's bandwidth, and the lowest roof at your kernel's intensity for that level is the one that binds. On the same H100, the PCIe roof (64 GB/s each way) is about 52 times lower than the HBM roof, which is the roofline view of why host to device transfers in the inner loop are so costly.

Putting it together: a decision table

Common mistakes

• Mixing directions. PCI-SIG and NVIDIA usually quote link bandwidth with both directions added; memory bandwidth is usually one stream. Compare like with like, or you overstate a link by two times.
• Using sparse figures as dense. NVIDIA's H100 and GB200 tables mark sparsity figures with a footnote; dense is half. Ridge points computed from sparse peaks are double what dense kernels can reach.
• Treating peak as sustained. Every figure in this post is a vendor peak. Sustained bandwidth on real kernels is lower; NVIDIA's Volta whitepaper, for example, claims over 95 percent memory bandwidth efficiency on many workloads for V100, which is itself a vendor claim about one product.
• Forgetting the PCIe generation that actually trained. A data center GPU in a slot that negotiated Gen 3 x8 has a quarter of the Gen 4 x16 bandwidth. Check with nvidia-smi or lspci under load.
• Assuming unified memory is free. On software coherent systems, first touch on the wrong processor causes faults and migrations over PCIe. Prefetch, or use hints.
• Assuming every GPU pair in a box is equally connected. An NVSwitch domain and an Infinity Fabric mesh behave differently for point to point and all-to-all traffic; PCIe only boxes are different again. Read the topology (nvidia-smi topo -m on NVIDIA systems) before placing ranks.
• Treating CXL as an HBM substitute. CXL attached memory runs at PCIe class bandwidth. It adds capacity, not tensor core feeding speed.
• Ignoring the software that runs the fabric. NVSwitch systems need Fabric Manager matched to the driver; multi GPU jobs depend on NCCL or RCCL versions. Unpinned versions cause both failures and unpatched vulnerabilities.

Where Safeguard fits

Safeguard does not design, program or speed up GPUs, memory or interconnects, and nothing in its product makes a kernel faster. The honest connection is the software that runs all of this hardware. A GPU server's image or container typically carries a GPU driver user space, a CUDA or ROCm toolkit, cuDNN or MIOpen, NCCL or RCCL, Fabric Manager on NVSwitch systems, a container toolkit, and a framework and model serving stack on top. Each is a versioned component, and each gets security advisories.

Safeguard's generally available tooling applies to that software layer: SBOM generation and analysis in SPDX and CycloneDX, scanning of repositories, container images and SBOM files, vulnerability matching enriched with EPSS and CISA KEV data, reachability analysis to rank the findings your code actually calls, and autonomous remediation that opens tested fix pull requests, switched off by default. In practice that means you can scan a CUDA or ROCm base image, see which images carry a given NCCL or container toolkit version, and get a pull request that moves a pinned version forward when a fix exists. It does not scan GPU firmware, NVSwitch firmware or the host kernel driver installed outside your images; those updates still come from the hardware vendor and your operating system distribution.

FAQ

What is the difference between HBM and GDDR memory?

HBM stacks several DRAM dies on a base die and connects them to the GPU through a 1,024 bit (HBM3, HBM3E) or 2,048 bit (HBM4) interface on a silicon interposer. GDDR uses separate 32 bit chips on the circuit board and runs each pin much faster. HBM gives more bandwidth and capacity per package at higher cost; GDDR is cheaper and simpler to build.

How do you calculate GPU memory bandwidth?

Multiply the memory bus width in bits by the data rate per pin in Gb/s and divide by 8. An RTX 4090's 384 bit bus at 21 Gb/s gives 1,008 GB/s, and an MI300X's 8,192 bit bus at 5.2 Gb/s gives about 5.3 TB/s. The result is a theoretical peak; sustained bandwidth is lower.

Why does GDDR7 use PAM3 instead of PAM4?

Three voltage levels leave larger gaps between levels than four, which improves the signal to noise ratio at high frequencies. JEDEC says PAM3 improves signal to noise for high frequency operation. GDDR7 carries 3 bits over 2 cycles, 1.5 bits per symbol, less dense than PAM4's 2 bits but easier to run fast.

How much faster is NVLink than PCIe?

Using both directions for both: fourth generation NVLink gives 900 GB/s per H100 against 128 GB/s for PCIe 5.0 x16, about 7 times; fifth generation gives 1,800 GB/s, about 14 times PCIe 5.0. NVIDIA describes sixth generation NVLink at 3 TB/s as 12 times PCIe Gen6.

What is NVSwitch and why does NVL72 need it?

NVSwitch is a switch chip that connects many GPUs' NVLink ports so any GPU can reach any other at full NVLink bandwidth. In a GB200 NVL72, NVLink switches connect 72 GPUs into one domain with 1.8 TB/s per GPU and 130 TB/s in total, per NVIDIA's specification table.

Is AMD Infinity Fabric the same as NVLink?

It plays the same role. On an MI300X board, each GPU has seven Infinity Fabric links, one to each of the other seven GPUs, for 896 GB/s in total by AMD's figures, in a full mesh with no switch. NVLink systems with NVSwitch route through switches instead. The topology affects which communication patterns run fastest.

Does CXL replace HBM or NVLink for GPUs?

No. CXL runs on PCIe physical layers, so its bandwidth per link is in the PCIe class (64 GT/s per lane for CXL 3.0, 128 GT/s for CXL 4.0). Its strength is coherent memory expansion and pooling, mostly on the CPU side. It adds capacity, not the hundreds of gigabytes per second per GPU that HBM and NVLink provide.

What is unified memory in CUDA?

It is a single address space shared by CPU and GPU. With cudaMallocManaged, pages migrate on demand. On hardware coherent systems such as Grace Hopper, the GPU can access CPU memory directly over NVLink-C2C. On Linux with HMM, plain malloc memory also works in kernels, given the kernel, GPU and driver versions the CUDA Programming Guide lists.

What is the roofline model for GPUs?

It bounds a kernel's speed by the lower of peak compute and arithmetic intensity times memory bandwidth. On an H100 SXM5 at dense FP16, the ridge point is about 295 FLOP per byte, so a kernel below that intensity is memory bound. Large matrix multiplies are compute bound; matrix vector products are memory bound.

Why is LLM token generation memory bound?

At batch size 1, generating each token reads every weight once but does only about one multiply and add per weight, about 1 FLOP per byte in FP16. That is far below any modern GPU's ridge point, so memory bandwidth, not compute, sets the speed. Batching more requests raises the intensity.

None of the bandwidth arithmetic above changes the fact that every GPU node runs a large stack of versioned open source and vendor software, from collective libraries to container toolkits and serving frameworks, and those components get security fixes on their own schedules. The practices below apply to the parts of that stack you build into your own images and repositories; GPU firmware and host drivers installed outside them still come from the vendor.

More in this series

Other posts in our GPUs, CPUs, QPUs and lithography series:

• AMD GPU Architecture Explained: GCN, RDNA, CDNA, Infinity Fabric and HIP
• CPU vs GPU vs QPU Explained: How Each Works and Why They Coexist
• CUDA Programming Model Explained: Grids, Blocks, Warps and Memory
• CUDA vs HIP vs SYCL vs OpenCL vs Metal vs Vulkan vs Triton Compared
• GPU Software Supply Chain: Drivers, CUDA, cuDNN, NCCL and Your SBOM
• How a QPU Works: Qubit Chips, Fridges, Control Pulses and Readout
• How CPUs Work: Pipelines, Out-of-Order, Caches, Branches, SIMD
• How EUV Lithography Machines Work: Tin Plasma, Mirrors and High-NA
• How GPUs Work: Graphics Pipeline, Shader Cores, Ray Tracing and Compute
• Inside an NVIDIA SM: Warp Schedulers, Registers, Tensor Cores
• Intel GPU Architecture: Xe Cores, XMX, Ponte Vecchio and oneAPI SYCL
• NVIDIA GPU Architecture History: From Tesla and Fermi to Blackwell
• NVIDIA vs AMD vs Intel vs Apple GPUs: Architectures Compared
• Optimizing a CUDA Matrix Multiply Step by Step, From Naive to Tensor Cores
• Photolithography Explained: Contact Printing to DUV Immersion and Overlay
• Tile-Based GPUs Explained: Apple, Mali, Adreno, PowerVR and TBDR

Fix it without waiting for upstream

Whatever the vulnerability, most of the risk sits in two gaps. The first is the time between a fix being published and that fix actually running in your production builds. The second, and worse, is the time before any fix exists at all, while a maintainer works out a patch and a release schedule. Waiting through both gaps is a habit, not a requirement. Safeguard gives teams three ways to close them.

1. Autonomous remediation: pull requests, not tickets

When a fixed version exists, Safeguard's Griffin models open a tested, policy gated pull request that moves the affected dependency to the safe version. You can run it on one repository or in bulk across every project that carries the vulnerable version. Remediation covers eight categories, including dependencies, licenses, security posture and findings. Reachability analysis ranks the vulnerable code paths your application actually calls first, so the first pull requests you merge are the ones that matter most. Autonomous remediation ships switched off, with every category set to Manual, so your team decides how much runs on its own and what still needs a human approval.

2. Gold Open Source: the safe version in seconds, free

gold.safeguard.sh is free and needs no login. Search any CVE or package to see affected and fixed versions, CVSS, EPSS, CISA KEV status and a recommended safe version, or use the "Am I affected?" checker. The same data is available through a free read API, README status badges and RSS threat feeds. The safeguard-sh/gold-scan-action GitHub Action can fail a build on a known exploited vulnerability, on any known vulnerability, or never, so a version you have just removed cannot quietly come back in next week's merge.

3. Gold Registry: start from components that are already fixed

registry.safeguard.sh provides hardened, drop in replacements for upstream packages and container images. Every Gold artifact is held continuously to zero exploitable critical or high CVEs, zero malware and a completed license check, with SLSA L3 or higher provenance, a Sigstore signature and a public evidence page you can hand to an auditor. An artifact loses its Gold badge if it regresses. Teams can point their own package directory at the registry so hardened equivalents resolve without changing how developers install anything.

When upstream has no fix yet, or will not ship one, Safeguard first works to get the fix accepted upstream. If the issue is severe or upstream will not move, Safeguard builds a fork, fixes the vulnerability there and publishes a hardened Gold build, so you are not stuck on a maintainer's release schedule. You can also request a hardened, compatibility tested build of a specific package. This fork and patch flow is in early access today, so ask which packages it covers before you plan around it.

One honest limit: when the vulnerable code lives in a vendor appliance, a hosted service or an operating system runtime you do not build yourself, the fix still has to come from that vendor or distribution. The steps above apply to the open source dependencies and container images you build and ship, which is where most teams carry the largest and least tracked share of their exposure.

See how Safeguard would help

The first demo is a real Safeguard CLI session. You install the CLI, open the interactive terminal and ask it to fix a lodash CVE in a repository; it reads the finding and the advisory, asks before any command that changes something, edits package.json, runs the install and tests with your approval and pushes nothing. The same flow applies to the Python and container dependencies of a GPU training or serving repository, and safeguard scan --image takes a CUDA or ROCm based image.

Click through it at your own pace, or open it full screen.

The second demo sets up the Safeguard runner on your own machine, approves it with a device code and then scans from the terminal. GPU teams often keep images and model code on their own build hosts; the runner lets those scans happen on hardware you control.

Click through it at your own pace, or open it full screen.

The third demo connects Claude to the Safeguard MCP server with OAuth and shows how you choose which Safeguard tools Claude can use. Once connected, you can ask in plain language which of your projects or images carry a given NCCL, CUDA toolkit or container toolkit version.

Click through it at your own pace, or open it full screen.

To try it on your own images, download the Safeguard CLI and follow the CLI documentation. For background on accelerators in general, see AI chip architecture and what an AI accelerator is and how it works.

Quick check

Sources

• JEDEC: JESD238 HBM3 announcement, 27 January 2022 (via Business Wire)
• JEDEC: HBM3 update to the HBM standard (republished release)
• JEDEC: JESD270-4 HBM4 standard, 16 April 2025 (via Business Wire)
• JEDEC: GDDR7 graphics memory standard JESD239, 5 March 2024
• JEDEC: JESD235A HBM update, 12 January 2016 (republished release)
• Micron: HBM3E product page
• Micron: GDDR6X product page
• TSMC 3DFabric: CoWoS
• PCI-SIG: PCIe 5.0 specification, May 2019 (republished release)
• PCI-SIG: PCIe 6.0 specification, 11 January 2022 (via insideHPC)
• PCI-SIG: PCIe 7.0 specification, 11 June 2025 (via HPCwire)
• NVIDIA: Pascal architecture whitepaper (Tesla P100)
• NVIDIA: Volta architecture whitepaper (Tesla V100)
• NVIDIA: Ampere architecture whitepaper (A100)
• NVIDIA: A100 product page
• NVIDIA Technical Blog: NVIDIA Hopper Architecture In-Depth, 22 March 2022
• NVIDIA: H100 product page
• NVIDIA: Ada GPU architecture whitepaper
• NVIDIA: GeForce RTX 5090 product page
• NVIDIA: Blackwell architecture
• NVIDIA: DGX B200
• NVIDIA: NVLink and NVLink Switch
• NVIDIA: GB200 NVL72
• NVIDIA: Grace Hopper Superchip
• NVIDIA: Fabric Manager user guide
• NVIDIA: CUDA Programming Guide, unified memory
• NVIDIA: nccl-tests performance notes
• AMD: CDNA 3 architecture whitepaper
• AMD: Instinct MI300X product page
• CXL Consortium: about CXL
• CXL Consortium: CXL 3.0 specification release
• CXL Consortium: CXL 4.0 white paper
• CXL Consortium: CXL 3.0 white paper
• AMD ROCm Blogs: Understanding RCCL bandwidth and xGMI performance on MI300X, 2 March 2025
• AMD: MI300X system optimization guide
• Apple Newsroom: MacBook Pro with the M4 family, 30 October 2024
• Williams, Waterman and Patterson: Roofline, Communications of the ACM 52(4), April 2009

Last reviewed: 25 October 2026. Vendor specifications change between product revisions and some figures here are marked preliminary by their vendors; check the linked documents before you size a system on them.