AMD GPU Architecture Explained: GCN, RDNA, CDNA, Infinity Fabric and HIP

AMD builds two GPU families from one idea. RDNA powers Radeon graphics cards and runs 32-wide wavefronts on SIMD32 units grouped in workgroup processors. CDNA powers Instinct accelerators, keeps 64-wide wavefronts and adds matrix cores, with MI300 and MI350 built from chiplets joined by Infinity Fabric. You program both through ROCm and HIP.

Both families descend from Graphics Core Next (GCN), presented in 2012. This guide walks from the GCN compute unit to the CDNA 4 accelerator complex die, through the ROCm stack, and ends with a CUDA kernel ported to HIP. Every specification comes from an AMD whitepaper, presentation, product page or the ROCm documentation, linked in Sources. Where AMD has not published a number, we say so.

Two caveats before we start. First, every throughput figure in this post is a peak theoretical number from AMD, calculated from clock speed and per-clock rates. None of them is a measured benchmark, and real kernels reach a fraction of them. Second, we wrote this on a Mac with no AMD GPU, so the HIP and CUDA code below follows the ROCm documentation but was not compiled or run here. Treat it as a teaching sample and test it on your own hardware.

This post is part of a series on how processors work. If you want the wider context first, read our explainers on AI chip architecture, what an AI accelerator is and how a GPU is made, from wafer to HBM package.

Key takeaways

• One lineage, two branches. GCN (2012) set the template: four SIMD16 units per compute unit running 64-wide wavefronts. RDNA moved to SIMD32 and native wave32 for graphics latency. CDNA kept wave64 for compute and added matrix cores.
• The vocabulary maps cleanly. Warp is wavefront, SM is compute unit, shared memory is the Local Data Share (LDS), thread block is workgroup. Wavefronts are 64 lanes on CDNA and 32 or 64 on RDNA.
• MI300X is a package of chiplets. Eight accelerator complex dies (XCDs) with 38 active CUs each sit on four I/O dies holding a 256 MB Infinity Cache. AMD lists 304 CUs, 192 GB of HBM3 and 5.3 TB/s.
• CDNA 4 rebalances it. MI355X: 256 CUs on N3P, two I/O dies, 288 GB of HBM3E at 8 TB/s, a 160 KB LDS, and MXFP6 and MXFP4 formats.
• Infinity Fabric is the glue, inside the package and between GPUs: 896 GB/s of peer-to-peer bandwidth per MI300X, 1,075.2 GB/s per MI355X.
• HIP is close to CUDA, not identical. HIPIFY converts API names; the bugs that survive are 32-lane assumptions in warp code.
• The GPU stack is a software supply chain. Drivers, firmware, ROCm libraries and images carry versions and CVEs and belong in your SBOM.

A dated timeline of AMD GPU compute architectures

The dates below come from AMD's own documents. Where the source we used did not state a launch date, the table says so instead of guessing.

The first CDNA generation (Instinct MI100) introduced matrix cores. The CDNA 2 whitepaper describes Matrix Core technology as "one of the fundamental innovations in the prior generation", and the ROCm hardware page says matrix units exist on "MI100 and newer". We did not fetch a first-party document with the MI100 launch date, so it is left out of the table.

CUDA terms and AMD terms

If you learned GPUs on NVIDIA hardware, most AMD terms have a direct partner. The table below maps them. The HIP documentation itself uses "warp" for what the AMD ISA calls a "wavefront", so you will see both words in ROCm material.

GCN: the compute unit that set the template

Before GCN, AMD GPUs used VLIW shader cores. The compiler had to pack independent operations into four or five slots of a very long instruction word, and code with dependent operations left slots empty. AMD's Hot Chips 24 presentation in 2012 put the change in one line: VLIW4 was "dependency limited", while the new "GCN Quad SIMD" was "occupancy limited". GCN stopped asking the compiler to find parallelism inside one thread and asked the hardware to find it across wavefronts.

Inside a GCN compute unit

The Hot Chips slides list the parts of one CU. The text-art diagram below follows the slide's block diagram.

                    GCN compute unit (Hot Chips 24, 2012)
 +--------------------------------------------------------------------+
 | Instruction fetch, arbitration, decode, branch and message unit    |
 +--------------------------------------------------------------------+
 | Scalar unit |  SIMD0   |  SIMD1   |  SIMD2   |  SIMD3   |           |
 | + SGPRs     |  16 lane |  16 lane |  16 lane |  16 lane |           |
 |             |  64 KB   |  64 KB   |  64 KB   |  64 KB   |           |
 |             |  VGPRs   |  VGPRs   |  VGPRs   |  VGPRs   |           |
 +--------------------------------------------------------------------+
 | Local Data Share, 64 KB, 32 banks | 16 KB read/write L1 vector cache|
 +--------------------------------------------------------------------+
 | 4 texture filter units, 16 load/store (texture fetch) units        |
 +--------------------------------------------------------------------+
          | 64 bytes per CU per clock to L2 (per the slides)
          v
   L2 partitions (768 KB total on the HD 7970), write-back

The pieces that matter for compute:

• Four SIMD16 vector units. Each SIMD has 16 lanes and its own 64 KB slice of vector registers, 256 KB per CU in total.
• A scalar unit. It executes once per wavefront rather than once per lane, which suits loop counters, addresses and branch conditions. The slides give conflicting sizes for the instruction and scalar caches shared by four CUs, so we do not quote them.
• A 64 KB Local Data Share with 32 banks, which is the hardware behind OpenCL local memory and HIP shared memory.
• A 16 KB read/write L1 cache per CU, write-through, backed by L2 partitions. On the Radeon HD 7970 the slides list 32 CUs, 768 KB of L2 and a 384-bit GDDR5 bus at 5.5 Gbps for 264 GB/s. That figure checks out: 384 bits times 5.5 Gbps divided by 8 bits per byte is 264 GB/s.

The four-cycle wavefront

A GCN wavefront has 64 work-items, but each SIMD is only 16 lanes wide. So one vector instruction for one wavefront takes four clock cycles: lanes 0 to 15 in the first cycle, 16 to 31 in the second, and so on. The CU's instruction issue rotates across the four SIMDs, one per cycle, so every SIMD receives a new instruction every fourth cycle. AMD's RDNA deck redraws this exactly:

 GCN issue, one SIMD16 per cycle in rotation (from AMD's RDNA deck)

 Cycle    0      1      2      3      4      5      6      7
 Issue    SIMD0  SIMD1  SIMD2  SIMD3  SIMD0  SIMD1  SIMD2  SIMD3
 SIMD0    0-15   16-31  32-47  48-63  0-15   16-31  32-47  48-63
 SIMD1    0-15   16-31  32-47  48-63  0-15   16-31  32-47  48-63
 SIMD2    0-15   16-31  32-47  48-63  0-15   16-31  32-47  48-63
 SIMD3    0-15   16-31  32-47  48-63  0-15   16-31  32-47  48-63

Two consequences follow. A dependent ALU instruction in the same wavefront rarely stalls, because four cycles pass before the SIMD returns to it. And the machine needs many threads to stay busy. The RDNA deck states that each wave is assigned to one SIMD16 with up to 10 waves per SIMD, and that keeping two GCN CUs fully busy needs 2 times 4 times 64, or 512 threads, before you count the extra waves needed to hide memory latency.

GCN's memory model

The Hot Chips slides describe a relaxed memory model: work-items in one workgroup share an L1, different workgroups may not, and the L2 is the point of coherence where global atomics run. HIP inherits this, which is why data written by one workgroup and read by another needs a fence or a kernel boundary, not just __syncthreads().

RDNA: the workgroup processor, wave32 and Infinity Cache

RDNA was designed for graphics, where many dispatches are small and latency matters more than raw throughput. AMD's own summary slide puts it bluntly: "GCN was compute/throughput focused, RDNA is graphics/latency focused."

The workgroup processor

The basic RDNA block is the workgroup processor (WGP), sometimes called a dual compute unit. One WGP holds two CUs, and each CU holds two SIMD32 units, each with its own scalar ALU. The two CUs in a WGP share an instruction cache, a scalar cache and a 128 KB LDS.

             RDNA workgroup processor (from AMD's RDNA deck)
 +-------------------------------------------------------------------+
 | Instruction cache 32 KB      |  Scalar cache 16 KB                 |
 +-------------------------------------------------------------------+
 | CU 0                         |  CU 1                               |
 |  SALU + SIMD32  SALU + SIMD32|   SALU + SIMD32  SALU + SIMD32      |
 |  L0 vector cache 16 KB       |   L0 vector cache 16 KB             |
 |  texture addresser and data  |   texture addresser and data        |
 +-------------------------------------------------------------------+
 |          Local Data Share, 128 KB per WGP, up to 64 KB per        |
 |          workgroup, 32 banks (RDNA 1)                             |
 +-------------------------------------------------------------------+
                 |
                 v
     L1 cache, 128 KB per shader array  ->  L2, 4 MB (RX 5700 XT)

The RX 5700 XT ("Navi") in that deck has 20 WGPs, which AMD equates to 40 CUs.

Wave32, single-cycle issue and why it lowers latency

Each RDNA SIMD32 issues one instruction per cycle. A wave32 instruction runs across 32 lanes in a single pass, so its vector throughput is one instruction per cycle per SIMD, where a GCN SIMD16 needed four cycles for a wave64 instruction. AMD lists up to 20 waves per SIMD32, and five cycles of exposed latency that the hardware checks automatically. Other waves, or independent instructions in the same wave, fill those cycles.

This makes instruction-level parallelism matter again. AMD's deck shows the same six instructions in two orders. When each v_fmac_f32 depends on the v_fma_f32 just before it, the sequence takes about 18 cycles; reorder them so three independent FMAs run first and the same work fits in about 12.

; dependent order, each pair stalls on the previous result
v_fma_f32  v4, v0, s0, s3
v_fmac_f32 v4, v1, s1
v_fma_f32  v5, v0, s4, s7
v_fmac_f32 v5, v1, s5
v_fma_f32  v6, v0, s8, s11
v_fmac_f32 v6, v1, s9

; interleaved order, three independent chains hide the latency
v_fma_f32  v4, v0, s0, s3
v_fma_f32  v5, v0, s4, s7
v_fma_f32  v6, v0, s8, s11
v_fmac_f32 v4, v1, s1
v_fmac_f32 v5, v1, s5
v_fmac_f32 v6, v1, s9

RDNA still supports wave64. The hardware runs a wave64 vector instruction as two wave32 halves, and if one half of the execution mask is zero, that half is skipped. AMD's guidance is that the compiler picks the mode: compute and vertex shaders usually compile as wave32, pixel shaders usually as wave64. Wave32 gives lower latency and quicker ramp-up after barriers. Wave64 gives more threads in flight per lane and is more efficient for some memory patterns.

For small dispatches this matters: two GCN CUs need 512 threads to fill every SIMD, one RDNA WGP needs 4 times 32, or 128. AMD notes that 256 threads per WGP often reached over 90% ALU utilization in one game trace it studied, an observation rather than a benchmark.

Occupancy as threads per SIMD lane

RDNA's register file is where people get occupancy wrong. Each SIMD32 has 1,024 physical vector registers, each 32 lanes wide, divided among waves at up to 256 per wave. A wave64 "counts double". AMD's examples:

The register file size per SIMD32 works out to 1,024 registers times 32 lanes times 4 bytes, which is 128 KB. AMD's advice is to think about occupancy as "number of threads per SIMD lane", because that number means the same thing on GCN and RDNA. GCN occupancy 4 equals 16 threads per lane, and the RDNA equivalent is 16 wave32 or 8 wave64.

Worked: a wave32 kernel using V VGPRs fits at most 1,024 divided by V waves per SIMD32 (capped at the wave slots), so 8 waves at 128 VGPRs and 4 at 256. Allocation granules can lower this slightly.

RDNA's cache hierarchy, with numbers

AMD's deck compares the RX Vega 64 (GCN) and RX 5700 XT (RDNA):

RDNA also doubled LDS throughput to 32 dwords per cycle, kept 32 banks (AMD's slide says "Mind the bank conflicts!"), and moved loads to 32 addresses per clock against GCN's 4 to 16 coalesced addresses. AMD's practical advice: prefer coalesced stores of at least 256 bytes per wave, and keep workgroup sizes a multiple of 64 so the same code runs well in either wave mode.

RDNA 2: Infinity Cache

RDNA 2 added a large memory-side cache that AMD calls Infinity Cache. The RX 6900 XT product page lists 80 CUs, 128 MB of Infinity Cache, 16 GB of GDDR6 on a 256-bit bus at up to 16 Gbps, and up to 512 GB/s of memory bandwidth. Check the bandwidth: 256 bits times 16 Gbps divided by 8 is 512 GB/s. The idea is to hit in a large on-die cache often enough that a narrow, cheaper memory bus behaves like a wider one. The ROCm hardware page notes that RDNA 2, 3 and 4 LDS designs use 64 banks of 4 bytes, 256 bytes per cycle, up from RDNA 1's 32 banks.

The peak FP32 figure on the product page also checks out from first principles. RDNA 2 does 128 FP32 FLOPs per CU per cycle (64 lanes, two FLOPs per fused multiply-add). At a 2,250 MHz boost clock, 80 CUs times 2.25 GHz times 128 is 23.04 TFLOPS, matching AMD's "23.04 TFLOPs".

RDNA 3: chiplets and dual issue

AMD announced RDNA 3 on 3 November 2022 with the Radeon RX 7900 XTX and XT, calling them "the world's first gaming graphics cards to feature an advanced AMD chiplet design". The press release describes:

• a 5 nm Graphics Compute Die (GCD) with up to 96 compute units;
• six 6 nm Memory Cache Dies (MCDs), each with up to 16 MB of second-generation Infinity Cache, for 96 MB in total;
• Infinity Links and fanout packaging between the chiplets at up to 5.3 TB/s;
• up to 24 GB of GDDR6 on a 384-bit interface.

          RDNA 3 package (RX 7900 XTX), per AMD's press release
   +------+  +------+  +------+
   | MCD  |  | MCD  |  | MCD  |   each MCD: 16 MB Infinity Cache
   | 6 nm |  | 6 nm |  | 6 nm |   + 64-bit GDDR6 interface
   +--+---+  +--+---+  +--+---+
      |         |         |
   +--+---------+---------+------------------------------+
   |              GCD, 5 nm, up to 96 CUs                |
   |   shader engines, WGPs, L0, L1, L2, front end       |
   +--+---------+---------+------------------------------+
      |         |         |      chiplet links up to 5.3 TB/s
   +--+---+  +--+---+  +--+---+
   | MCD  |  | MCD  |  | MCD  |
   +------+  +------+  +------+

AMD describes the 5 nm and 6 nm dies as "each optimized for specific jobs": compute logic on the newer node, cache and memory interfaces on the mature one. The press release gives one number for the 384-bit bus, and six MCDs, which means each MCD carries 64 bits of the interface.

RDNA 3 also introduced dual-issue vector instructions. The ROCm Compute Profiler documentation for RDNA 3.5 (gfx115x) explains the encoding, called VOPD: two independent VALU operations packed in one instruction, issued in the same cycle in wave32 mode. It gives the peak as 128 FP32 FLOPs per CU per cycle single-issue and 256 with VOPD. That doubling is visible in the RX 7900 XTX product page: 96 CUs times 2.5 GHz boost times 256 FLOPs per cycle is 61.44 TFLOPS, which matches AMD's "61.4 TFLOPs" FP32 vector figure. In other words, the headline number assumes the compiler finds pairs of independent operations every cycle. Kernels that cannot be paired see half of it.

The same page lists up to 960 GB/s of memory bandwidth (384 bits times 20 Gbps divided by 8) and a separate "effective memory bandwidth" of up to 3,500 GB/s that counts Infinity Cache hits. Plan rooflines on the first; the second depends on hit rate.

RDNA 4 on a product page

For RDNA 4 we have only the product page, not a whitepaper. AMD lists the RX 9070 XT with 64 CUs, a 2,970 MHz boost clock, 64 MB of Infinity Cache, 16 GB of GDDR6 at 20 Gbps on a 256-bit bus (640 GB/s), 48.7 TFLOPS FP32 vector, 195 TFLOPS FP16 matrix, and 389 TFLOPS of FP8 matrix in both E5M2 and E4M3, doubling with structured sparsity. The FP32 figure again implies dual issue: 64 times 2.97 GHz times 256 is 48.66 TFLOPS. AMD has not published per-CU matrix rates for RDNA 4 in a document we could find, so we do not derive them.

CDNA: AMD's compute-only branch

CDNA drops the fixed-function graphics hardware (rasterizers, render backends, display engines) and spends the area on compute, matrix units, cache and HBM. Its CU is still recognisably GCN: four SIMDs, 64-wide wavefronts, a scalar unit, an LDS, and a vector L1.

CDNA 2: two dies in one package

The CDNA 2 whitepaper describes the Instinct MI200 Graphics Compute Die (GCD). Each GCD has 112 physical CUs in four arrays; products enable 104 (MI250 and MI210) or 110 (MI250X). The MI250 and MI250X put two GCDs in one OCP Accelerator Module (OAM) package, which is why an MI250X shows up as two GPUs to software.

Three CDNA 2 details stand out. Double precision runs at the single-precision rate, 64 FMAs per clock per CU, with 64-wide wavefronts executed over four cycles on four 16-wide pipelines, the GCN pattern. Packed FP32 doubles FP32 to 128 FMAs per cycle per CU, but only when packed instructions are used. And matrix cores, introduced in the first CDNA, gained FP64 support for scientific codes.

Memory and links per the whitepaper: 8 MB of L2 per GCD in 32 slices of 16-way set-associative cache, 64 GB of HBM2e per GCD (128 GB per MI250X), 3.2 TB/s aggregate across both GCDs. The in-package Infinity Fabric between the two GCDs runs at 25 Gbps for up to 400 GB/s bidirectional, and eight external Infinity Fabric links give up to 800 GB/s. AMD's endnote spells out the memory bandwidth formula: an 8,192-bit bus (4,096 bits times two dies) at 3.2 Gbps, divided by 8, is 3,276.8 GB/s.

CDNA 3: XCDs, IODs and the MI300 family

The CDNA 3 whitepaper describes a heterogeneous package: up to eight accelerator complex dies (XCDs) on TSMC 5 nm, stacked on four I/O dies (IODs) on TSMC 6 nm, wired together with Infinity Fabric and connected to eight stacks of HBM. The MI300X uses eight XCDs. The MI300A APU replaces two XCDs with three CPU complex dies (CCDs) holding 24 "Zen 4" cores.

              AMD Instinct MI300X package (after the CDNA 3 whitepaper)

   HBM3  HBM3            HBM3  HBM3            HBM3  HBM3            HBM3  HBM3
     \    /                \    /                \    /                \    /
 +----------------+  +----------------+  +----------------+  +----------------+
 | XCD      XCD   |  | XCD      XCD   |  | XCD      XCD   |  | XCD      XCD   |
 | 38 CU    38 CU |  | 38 CU    38 CU |  | 38 CU    38 CU |  | 38 CU    38 CU |
 | 4MB L2   4MB L2|  | 4MB L2   4MB L2|  | 4MB L2   4MB L2|  | 4MB L2   4MB L2|
 |================|==|================|==|================|==|================|
 | IOD 0, 6 nm    |  | IOD 1          |  | IOD 2          |  | IOD 3          |
 | 64 MB Infinity |  | 64 MB          |  | 64 MB          |  | 64 MB          |
 | Cache, 2 HBM   |  | Infinity Cache |  | Infinity Cache |  | Infinity Cache |
 | controllers    |  |                |  |                |  |                |
 +----------------+  +----------------+  +----------------+  +----------------+
     Infinity Fabric across all four IODs; 7 external x16 links + 1 x16 PCIe Gen 5

The 64 MB per IOD in the sketch is our division of AMD's 256 MB total across four IODs; the whitepaper describes the cache as 16 channels per HBM stack, 128 channels in all, totalling 256 MB.

The XCD. Each XCD has 40 physical CUs with 2 disabled for yield, leaving 38 active, plus four Asynchronous Compute Engines (ACEs), a hardware scheduler and a shared 4 MB L2. Eight XCDs give 304 CUs. AMD lists 19,456 stream processors (304 times 64) and 1,216 matrix cores (four per CU).

The CU. The instruction cache is now shared between two CUs and doubled to 64 KB. The vector L1 cache doubled its line size to 128 bytes and its capacity to 32 KB. The LDS stays at 64 KB. Most of the new area went to matrix cores. The whitepaper's per-CU, per-clock rates:

New data types. CDNA 3 matrix cores support TF32 (19 bits: sign, 8-bit exponent, 10-bit mantissa) and both OCP FP8 variants, E5M2 for training and E4M3 for inference. They also exploit 4:2 structured sparsity: when at least two of every four input values are zero, the non-zero values are stored compactly with metadata and throughput doubles, up to "8K operations per clock for a CU".

The L2. The 4 MB L2 in each XCD is 16-way set-associative with 16 channels of 256 KB. Each channel reads a 128-byte line per clock, so one XCD's L2 can deliver 2,048 bytes per clock. It is write-back and write-allocate, coherent within its XCD, and connects to the IOD over 16 channels of 64 bytes each.

Infinity Cache. The 256 MB Infinity Cache on the IODs is a memory-side cache. It does not take part in coherence, so it never absorbs snoop traffic, but it holds a snoop filter covering the XCD L2 caches so that most cross-XCD coherence requests resolve there. AMD gives aggregate peak bandwidths in a figure: 17.2 TB/s from Infinity Cache to the XCDs, and 5.3 TB/s from HBM.

HBM. Each IOD drives two HBM3 stacks. The MI300X has 192 GB and the MI300A 128 GB, both at 5.3 TB/s peak. AMD's endnote again gives the formula: an 8,192-bit bus (1,024 bits times eight stacks) at 5.2 Gbps, divided by 8, is 5,324.8 GB/s. The MI325X keeps the CDNA 3 compute die but moves to 256 GB of HBM3E at 6.0 GT/s for 6.0 TB/s.

The MI300A APU and unified memory

The MI300A swaps two XCDs for three Zen 4 CCDs, which leaves 6 XCDs and 228 CUs alongside 24 CPU cores. The whitepaper is explicit that on the MI300A "the HBM3 memory is unified and shared between the GPU and" CPU. There is no separate host DRAM pool to copy into device memory, and both kinds of core sit behind the same Infinity Cache. For HPC codes that bounce data between CPU and GPU phases, removing the explicit copy and the PCIe hop is the point of the design. In a four-APU node, each MI300A connects to its peers with two Infinity Fabric links, with AMD quoting 256 GB/s between processors and a peer-to-peer ring aggregate of 384 GB/s.

Partitioning

Both MI300 parts can be partitioned along XCD lines: up to eight partitions on the MI300X, one per XCD, and up to three on the MI300A, two XCDs each. Each partition appears as a separate GPU, and a kernel on one XCD sees that XCD's 4 MB L2, not all 32 MB.

CDNA 4: MI350X and MI355X

The CDNA 4 whitepaper (2025) keeps the chiplet recipe but rebalances it:

• XCDs on TSMC N3P. Each XCD has 36 CUs in four arrays of nine, with 32 active. Eight XCDs give 256 CUs, fewer than MI300X's 304, but with more matrix throughput per CU.
• Two IODs instead of four, still on N6. AMD says the simpler direct connection between IODs is roughly 14% faster than CDNA 3's, a figure AMD attributes to its own testing.
• Matrix rates doubled for 16-bit and 8-bit types, to 4,096 FLOPs per clock per CU for FP16 and BF16 and 8,192 for FP8. Two new microscaling formats from the OCP MX specification, MXFP6 (E3M2 and E2M3) and MXFP4 (E2M1), run at twice the FP8 rate.
• TF32 left the hardware. AMD writes that TF32 "has been moved out of hardware and is supported through software emulation utilizing the BF16 datatype".
• FP64 matrix halved. Matrix FP64 drops from 256 to 128 FLOPs per clock per CU, so MI355X FP64 matrix peak (78.6 TF) is about half of MI300X's (163.4 TF). For double-precision HPC this is a real trade-off, and AMD's table shows it plainly.
• A 160 KB LDS. The LDS grows from 64 KB to 160 KB, read bandwidth doubles to 256 bytes per clock, and the LDS can load directly from the L1 data cache without passing through vector registers. The ROCm hardware page lists 64 banks for CDNA 4.
• HBM3E at 8 Gbps. Eight 12-high stacks of 36 GB give 288 GB, and the bandwidth is 8,192 bits times 8 Gbps divided by 8, or 8,192 GB/s, which AMD rounds to 8.0 TB/s.

The MI350X (1,000 W, air cooled) and MI355X (1,400 W, direct liquid cooled) share the silicon and differ in clock: 2,200 MHz against 2,400 MHz.

Deriving the peak numbers yourself

AMD's endnotes describe the method: peak FLOPS equals engine clock times CUs times FLOPs per clock per CU. Every figure below was recomputed in Python from AMD's per-clock rates and clocks.

Doing this arithmetic catches errors in source documents. Two examples from AMD's own whitepapers:

• The CDNA 3 whitepaper's Table 1 lists MI300X TF32 as 490.3 TFLOPS, but its Table 2 and the endnote give 653.7. 490.3 is exactly the MI300A figure (228 CUs times 2.1 GHz times 1,024), so Table 1 appears to carry the MI300A value. Use 653.7 for MI300X.
• The CDNA 4 whitepaper lists MXFP6 and MXFP4 at "16834" FLOPs per clock per CU. Twice the FP8 rate of 8,192 is 16,384, and only 16,384 reproduces AMD's 10 PF, so "16834" is a transposition.

None of these is sustained throughput; power limits, memory stalls and instruction mix pull real kernels below them.

Infinity Fabric: inside the package and across the node

Infinity Fabric is AMD's name for the family of interconnects it uses between dies in a package, between packages, and to host CPUs. In GPU contexts you will also see the term xGMI for the GPU-to-GPU links. It is not one fixed wire: rate, width and protocol differ by generation and by role.

Checking the link arithmetic

The CDNA 4 whitepaper says each link is 16 bits wide, fully bidirectional, at 38.4 Gbps, "for a total link bandwidth of 76.8GB/s in each direction". That is 16 times 38.4 divided by 8. Seven such links to the other seven GPUs give 7 times 76.8 times 2 directions, which is 1,075.2 GB/s, matching AMD's P2P aggregate. Add one x16 PCIe Gen 5 link at 64 GB/s each way (128 GB/s total) and you get AMD's "total peak aggregate I/O bandwidth" of 1,203.2 GB/s.

CDNA 3 works the same way at 32 Gbps: 16 times 32 divided by 8 is 64 GB/s each way, times 7 links times 2 directions is 896 GB/s, plus 128 GB/s of PCIe gives AMD's 1,024 GB/s.

The 8-GPU node

Seven links per GPU is exactly enough to connect each of eight GPUs to every other with no switch: 8 times 7 divided by 2 gives 28 pairs, each with one dedicated link, on an OCP Universal Base Board. AMD says this makes allreduce and allgather more efficient than its previous generation. The trade-off of a switchless mesh is that one pair can use only one link, 76.8 GB/s each way on MI355X, and the scale-up domain stops at eight GPUs. Larger jobs go over the scale-out network, which RCCL handles.

Worked numbers: rooflines on MI300X and MI355X

The roofline model says a kernel's attainable throughput is the smaller of the compute peak and the memory bandwidth times its arithmetic intensity (FLOPs per byte moved from memory). The ridge point, where the two limits meet, is peak FLOPS divided by bandwidth.

Example 1: SAXPY. y[i] = a * x[i] + y[i] in FP32 reads 8 bytes, writes 4, and does one multiply and one add. Its arithmetic intensity is 2 FLOPs divided by 12 bytes, about 0.167. On MI300X that gives an attainable ceiling of 0.167 times 5.3248 TB/s, about 0.89 TFLOPS: roughly half of one percent of the FP32 vector peak. No amount of compute tuning helps; only bandwidth does. The 256 MB Infinity Cache helps only if the vectors fit and are reused.

Example 2: a large BF16 GEMM. Multiplying two 8,192 by 8,192 matrices takes 2 times 8,192 cubed FLOPs, about 1.10 TFLOP. If each matrix crossed HBM once at 2 bytes per element, traffic would be about 403 MB, an ideal intensity of about 2,731 FLOPs per byte, far right of the 245.5 ridge. At peak the math takes about 0.84 ms; moving 403 MB at 5.3 TB/s takes about 0.076 ms. Real GEMMs re-read tiles, lowering intensity, but at this size they stay compute bound.

Example 3: why tiling matters. A naive tiled kernel that loads two 16 by 16 FP32 tiles from global memory per 16-step inner product does 2 times 16 cubed FLOPs per 2 times 16 squared times 4 bytes, an intensity of 4 FLOPs per byte, well below MI300X's FP32 ridge of 30.7. Doubling the tile to 32 by 32 doubles intensity to 8. This is why production GEMM libraries use large tiles, register blocking and matrix instructions, and why CDNA 4's bigger LDS is aimed at "matrix multiply routines because of the extensive data reuse", in AMD's words.

These rooflines use HBM bandwidth. A working set that fits the 256 MB Infinity Cache sees a ceiling nearer the 17.2 TB/s AMD gives for that cache; prefer measured ceilings when you have them.

Occupancy worked on CDNA 3

Occupancy on CDNA is limited by four resources the workgroup manager tracks, according to the ROCm hardware page: wave slots, VGPRs, SGPRs and LDS. The page says each SIMD pool has slots for up to ten waves (eight on MI200), and that MFMA units use both standard VGPRs and accumulation VGPRs (AGPRs), with "up to 512 KiB of combined register storage per CU".

That 512 KiB spread over four SIMDs is 128 KiB per SIMD, which at 64 lanes of 4 bytes is 512 registers per lane. If a wave uses R combined VGPRs and AGPRs, a SIMD can hold at most 512 divided by R such waves, rounded down, before slot limits apply:

The table ignores allocation granularity, which rounds each wave's request up to a block size defined in the ISA, so real numbers can be a little lower. The jump from 96 to 128 registers costs one wave per SIMD; the jump from 128 to 256 halves occupancy.

LDS is the other common limit. With 64 KB per CU on CDNA 3, a workgroup that allocates 16 KB of shared memory allows at most 4 resident workgroups per CU. On CDNA 4, the same workgroup fits 10 times into 160 KB, so a kernel limited by LDS can more than double its resident workgroups without changes, or keep the same occupancy with larger tiles.

A matrix core instruction, up close

The ROCm hardware page shows an example CDNA instruction:

v_mfma_f32_16x16x4f16 v[0:15], v[16:31], v[32:47], v[0:15]

It computes D = A x B + C on a 16 by 16 output tile with an inner dimension of 4, reading FP16 inputs and accumulating in FP32. Count the work: 16 times 16 outputs, each needing 4 multiply-adds, is 1,024 multiply-adds or 2,048 FLOPs for the whole wavefront in one instruction. The 64 lanes each hold a slice of the fragments in their registers, which is why the operand ranges span 16 registers. Programmers reach these instructions through libraries (rocBLAS, hipBLASLt, Composable Kernel), through rocWMMA, through compiler intrinsics, or through Triton, rather than by writing assembly.

RDNA 3 and RDNA 4 use a different family of WMMA instructions for their "AI accelerators". The shapes and rates differ from CDNA's MFMA, so code tuned for one family rarely runs at full speed on the other without retuning.

The ROCm software stack

ROCm is AMD's open software stack for GPU computing. At the time of writing, the ROCm documentation describes ROCm 10.1.0, built through TheRock, AMD's open build and release system that replaced the previous monolithic release process. The layers, from bottom to top:

The hip* libraries are thin portability layers that can sit on either the roc* implementation or the CUDA library, so the same source can build for both vendors. The roc* libraries are the AMD implementations. When you port, calling hipBLAS keeps an NVIDIA build path open; calling rocBLAS directly gives access to AMD-specific options.

Porting a CUDA kernel to HIP, worked

We will port a block-wide sum reduction. It is short, it uses shared memory, and it uses warp shuffles, which is where most CUDA-to-HIP bugs hide.

The CUDA original

// reduce.cu  (CUDA, illustrative; not compiled or run for this post)
#include <cuda_runtime.h>
#include <cstdio>

__inline__ __device__ float warpReduceSum(float v) {
    // Assumes 32 threads per warp and a 32-bit full mask.
    for (int offset = 16; offset > 0; offset /= 2)
        v += __shfl_down_sync(0xffffffff, v, offset);
    return v;
}

__global__ void blockSum(const float* in, float* out, int n) {
    __shared__ float partial[32];               // one slot per warp (1024/32)
    int tid  = blockIdx.x * blockDim.x + threadIdx.x;
    float v  = (tid < n) ? in[tid] : 0.0f;

    v = warpReduceSum(v);
    int lane = threadIdx.x % 32;
    int warp = threadIdx.x / 32;
    if (lane == 0) partial[warp] = v;
    __syncthreads();

    if (warp == 0) {
        v = (threadIdx.x < blockDim.x / 32) ? partial[lane] : 0.0f;
        v = warpReduceSum(v);
        if (lane == 0) atomicAdd(out, v);
    }
}

int main() {
    const int n = 1 << 24;
    float *d_in, *d_out;
    cudaMalloc(&d_in, n * sizeof(float));
    cudaMalloc(&d_out, sizeof(float));
    cudaMemset(d_out, 0, sizeof(float));
    // ... fill d_in ...
    blockSum<<<(n + 255) / 256, 256>>>(d_in, d_out, n);
    cudaDeviceSynchronize();
    cudaFree(d_in); cudaFree(d_out);
    return 0;
}

On a 32-wide warp this is correct. Count the hidden assumptions: the loop starts at 16, the mask is the 32-bit 0xffffffff, lane and warp divide by 32, and partial has 32 slots.

Step 1: run HIPIFY

HIPIFY ships two tools. hipify-clang parses the code with Clang, so it understands the program and reports what it cannot convert; it needs a CUDA installation to parse headers. hipify-perl uses regular expressions generated from hipify-clang's tables and needs nothing else, but the HIPIFY documentation warns that it "offers less error detection".

# Text-based translation, quick and dependency free
hipify-perl reduce.cu > reduce.hip

# Clang-based translation, needs CUDA headers to parse the source
hipify-clang reduce.cu --cuda-path=/usr/local/cuda -o reduce.hip

HIPIFY rewrites API names, such as cudaMalloc to hipMalloc and the runtime header to hip/hip_runtime.h. The triple-chevron launch, blockSum<<<grid, block>>>(...), is valid HIP as written; the HIP porting guide says HIP-Clang supports it alongside hipLaunchKernelGGL. The HIPIFY documentation is clear about the limit: it "does not automatically convert all CUDA code into HIP code seamlessly", and you should review, test and tune afterwards.

Step 2: fix what HIPIFY cannot see

HIPIFY translates names. It cannot know that your algorithm assumes 32 lanes. The HIP documentation lists the problems directly:

• "Code should not assume a warp size of 32 or 64." Use the warpSize built-in in device code, and query it from device properties on the host.
• HIP's __ballot and __activemask return 64-bit values, "unlike the 32-bit value returned by the CUDA built-ins".
• The _sync warp functions are "available in ROCm 7.0 (and enabled by default, unlike in previous versions)" and "require a 64-bit unsigned integer mask argument". The implementation includes a static assert that checks the mask type, so passing the 32-bit 0xffffffff fails to compile rather than silently misbehaving.
• Lane masks built with 1 << lane are 32-bit; on a 64-wide wavefront, lanes 32 to 63 shift the bit out. The porting guide recommends a lane_mask_t that is 64-bit on GFX8 and GFX9 targets (GCN and CDNA) and 32-bit otherwise.

On a wave64 MI300X, a reduction whose shuffle starts at offset 16 never joins lanes 32 to 63 with lanes 0 to 31, so each wave yields two partial sums and the per-warp indexing is wrong. If it compiled, the answer would be silently incorrect.

Step 3: the portable HIP version

// reduce.hip  (HIP, written to the ROCm docs; not compiled or run for this post)
#include <hip/hip_runtime.h>
#include <cstdint>

__inline__ __device__ float warpReduceSum(float v) {
    // Works for 32- or 64-wide waves. HIP masks are 64-bit.
    const uint64_t full = ~0ull;
    for (int offset = warpSize / 2; offset > 0; offset /= 2)
        v += __shfl_down_sync(full, v, offset);
    return v;
}

// 1024 threads / 32 lanes = 32 waves is the worst case on any target.
constexpr int kMaxWavesPerBlock = 32;

__global__ void __launch_bounds__(256)
blockSum(const float* __restrict__ in, float* __restrict__ out, int n) {
    __shared__ float partial[kMaxWavesPerBlock];
    int tid  = blockIdx.x * blockDim.x + threadIdx.x;
    float v  = (tid < n) ? in[tid] : 0.0f;

    v = warpReduceSum(v);
    int lane   = threadIdx.x % warpSize;
    int wave   = threadIdx.x / warpSize;
    int nWaves = (blockDim.x + warpSize - 1) / warpSize;
    if (lane == 0) partial[wave] = v;
    __syncthreads();

    if (wave == 0) {
        v = (lane < nWaves) ? partial[lane] : 0.0f;
        v = warpReduceSum(v);
        if (lane == 0) atomicAdd(out, v);
    }
}

int main() {
    const int n = 1 << 24;
    float *d_in = nullptr, *d_out = nullptr;
    (void)hipMalloc(&d_in, n * sizeof(float));
    (void)hipMalloc(&d_out, sizeof(float));
    (void)hipMemset(d_out, 0, sizeof(float));
    // ... fill d_in ...
    blockSum<<<(n + 255) / 256, 256>>>(d_in, d_out, n);
    (void)hipDeviceSynchronize();
    (void)hipFree(d_in); (void)hipFree(d_out);
    return 0;
}

On ROCm releases before 7.0, where the _sync variants were not enabled by default, the non-sync __shfl_down(v, offset) is the equivalent call; the HIP documentation notes that AMD GPUs execute all lanes of a wavefront in lockstep. Production code should also check every HIP return value instead of casting it away as this sample does for brevity.

Step 4: build for specific targets

AMD GPUs are identified by gfx target names, and a binary without code for your GPU's target will not run on it. The HIP programming model page names gfx942 for the MI300X. For any other part, take the target from the ROCm compatibility matrix or from rocminfo on the machine.

# Build for an MI300X; add one --offload-arch per extra target in your fleet
hipcc -O3 --offload-arch=gfx942 reduce.hip -o reduce

# See which gfx target the installed GPU reports
rocminfo | grep -m1 -o 'gfx[0-9a-f]*'

hipcc is a compiler driver that invokes amdclang++ with the right include paths and the __HIP_PLATFORM_AMD__ macro, per the porting guide. You can call amdclang++ directly if your build system prefers.

Step 5: translate launch bounds and register limits

CUDA's __launch_bounds__(MAX_THREADS, MIN_BLOCKS_PER_SM) second argument means blocks per SM. HIP's second argument means waves per execution unit. The porting guide gives the conversion:

• In CU mode, typical of CDNA: MIN_WARPS_PER_EU = (MIN_BLOCKS * MAX_THREADS) / (warpSize * 2).
• In WGP mode, an RDNA feature: MIN_WARPS_PER_EU = (MIN_BLOCKS * MAX_THREADS) / (warpSize * 4).

Worked example: a CUDA kernel declared __launch_bounds__(256, 4) on CDNA becomes 4 times 256 divided by 64 times 2, which is 8, so __launch_bounds__(256, 8). On RDNA in WGP mode with wave32, the same CUDA bounds give 4 times 256 divided by 32 times 4, which is also 8. The porting guide also notes that amdclang++ does not support nvcc's --maxregcount flag and points you to __launch_bounds__ instead.

Step 6: map the libraries

Use the ROCm stack table above to swap libraries: cuBLAS for hipBLAS or rocBLAS, cuDNN for MIOpen, NCCL for RCCL (which keeps NCCL's API), CUB and Thrust for hipCUB and rocThrust, CUTLASS-style kernels for Composable Kernel. Framework users get this for free: PyTorch's ROCm builds use these libraries under the same Python API, and torch.version.hip reports the HIP version a build targets.

Choosing between RDNA and CDNA for compute

This is a trade-off, not a ranking. The right part depends on the precision, memory size and interconnect your workload needs.

If you develop on a Radeon card and deploy to Instinct, remember that the wave width changes from 32 to 64 between the two. Code written with warpSize survives that; code with a hard-coded 32 does not.

The software supply chain around an AMD GPU

A GPU job depends on far more than the kernel you wrote. Between your Python script and the silicon sit a kernel driver, firmware, the ROCm runtime and compilers, a stack of math and communication libraries, a framework build and a container image. Each has a version, each receives security fixes, and they must match each other within the ROCm compatibility matrix. For a broader view of where these costs and parts sit in an AI deployment, see our self-hosting cost estimate.

What to record for every GPU host and image: the amdgpu kernel driver, GPU firmware, the ROCm release, the HIP runtime and compiler, the math and communication libraries (rocBLAS, hipBLASLt, MIOpen, RCCL), the framework build, and the container base image digest. On Linux, commands like these read them; they were not run for this post:

modinfo amdgpu | grep -i '^version'        # out-of-tree driver; in-tree follows uname -r
uname -r                                   # kernel, which carries the in-tree driver
dpkg -l | grep -Ei 'firmware|rocm|hip|miopen|rccl'   # Debian or Ubuntu packages
amd-smi version                            # AMD SMI, driver and ROCm versions
hipconfig --version                        # HIP version
python -c "import torch; print(torch.__version__, torch.version.hip)"
docker inspect --format '{{.RepoDigests}}' rocm/pytorch:YOUR_TAG

Three habits help. Pin images by digest, not by a moving tag, so you know exactly which ROCm and library versions a job ran with. Generate an SBOM for each image you build, in SPDX or CycloneDX, so a new advisory against one library can be matched to every image that carries it. And keep the host side (kernel driver and firmware) on the same patch cadence as the images, because a container cannot fix a vulnerable host driver.

Common mistakes

1. Hard-coding 32. Any 32, 16 start offset, 0xffffffff mask or 1 << lane in warp-level code breaks on wave64 CDNA. Use warpSize and 64-bit masks.
2. Assuming __ballot is 32-bit. In HIP it returns 64 bits; storing it in an unsigned int drops half the lanes.
3. Trusting HIPIFY output without tests. HIPIFY changes names, not algorithms. Review, run correctness tests on each target, and only then tune.
4. Reading peak FLOPS as delivered FLOPS. Every AMD figure quoted here is peak theoretical. Profile your own kernel and compare against a measured roofline.
5. Ignoring dual issue in RDNA 3 and 4 peaks. The headline FP32 numbers assume VOPD pairing. Code that cannot pair runs at half.
6. Treating an MI250X as one GPU. It is two GCDs and appears as two devices. Size work and collectives accordingly.
7. Forgetting partition modes. A partitioned MI300X exposes up to eight smaller GPUs, each with one XCD's L2. Benchmarks taken in one mode do not transfer to another.
8. Building for the wrong gfx target. A binary without code for your GPU's target will not run there. Build fat binaries or match targets to your fleet.
9. Letting the GPU stack drift. Mixing a new ROCm user space with an old kernel driver, or rebuilding an image on a moving tag, produces failures that look like hardware faults and leaves unpatched components in place.

FAQ

What is the difference between RDNA and CDNA?

RDNA is AMD's graphics architecture for Radeon cards: SIMD32 units, native wave32, workgroup processors of two CUs, GDDR memory and Infinity Cache. CDNA is its compute architecture for Instinct accelerators: no graphics hardware, 64-wide wavefronts, matrix cores and HBM. Both run HIP code through ROCm.

What is a wavefront on an AMD GPU?

The group of work-items that execute one instruction together, AMD's equivalent of a CUDA warp. It is 64 work-items on GCN and CDNA, and 32 or 64 on RDNA, where the compiler chooses. HIP exposes the width as warpSize.

Is an AMD compute unit the same as an NVIDIA SM?

It plays the same role, running workgroups and holding registers, shared memory and an L1 cache. Inside, a CDNA CU has four SIMDs, a scalar unit and matrix cores, and an RDNA WGP pairs two CUs that share an LDS. Counts are not comparable across vendors.

What is Infinity Cache?

A large memory-side last-level cache: 128 MB on the RX 6900 XT, 96 MB across the memory cache dies of the RX 7900 XTX, and 256 MB on the MI300X I/O dies with 17.2 TB/s of aggregate peak bandwidth. It raises effective bandwidth when the working set fits.

What is Infinity Fabric?

AMD's interconnect family. It links chiplets inside a package and GPUs to each other. On MI355X, seven x16 links at 38.4 Gbps per GPU form a fully connected eight-GPU node with 1,075.2 GB/s of peak peer-to-peer bandwidth per GPU.

How many compute units does the MI300X have?

304: eight XCDs with 38 active CUs each, out of 40 physical. The MI300A has 228 CUs on six XCDs plus 24 Zen 4 cores. The MI355X has 256 CUs, 32 active per XCD.

Can I run CUDA code on AMD GPUs?

Not directly. You port it to HIP, usually with HIPIFY, which converts most runtime API calls and keeps the triple-chevron launch. Expect manual work on warp-size assumptions, inline PTX and NVIDIA-only libraries.

Does the MI300A really share memory between CPU and GPU?

Yes. AMD's CDNA 3 whitepaper says the MI300A's 128 GB of HBM3 is "unified and shared" between CPU and GPU, behind a shared Infinity Cache, so code can skip explicit host-to-device copies. Data placement still affects performance.

Where Safeguard fits

Safeguard does not design, program, tune or speed up GPUs, and nothing in this post depends on Safeguard. The fit is with the software around the GPU. ROCm libraries, framework builds and the container images that bundle them are open source components with versions and CVEs. Safeguard generates and analyses SBOMs in SPDX and CycloneDX, scans repositories and container images (for example safeguard scan --image against a ROCm-based training image), matches components against vulnerability databases, shows which projects carry a given library version, and ranks findings with reachability analysis. Autonomous remediation can open fix pull requests when a fixed version exists, and it ships switched off. What Safeguard does not cover here: the amdgpu driver and firmware on a host outside any image you scan, and the GPU hardware itself. Those updates come from your Linux distribution and from AMD.

A GPU cluster is rebuilt and re-pulled constantly, and every rebuild fetches a fresh set of ROCm libraries, Python wheels and base images that may carry new advisories. The steps below apply to those images and dependencies, not to the silicon or the host driver.

More in this series

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

• CPU vs GPU vs QPU Explained: How Each Works and Why They Coexist
• CUDA Programming Model Explained: Grids, Blocks, Warps and Memory
• GPU Memory and Interconnects: HBM, GDDR, NVLink, PCIe, CXL
• 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. Asked to fix a lodash CVE, it reads the finding and the advisory, asks before any command that changes something, edits package.json, runs npm install and npm test, and pushes nothing. The same flow applies to the requirements file or Dockerfile of a ROCm training service.

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

The second demo sets up the local runner on your own machine, approves it with a device code, and then scans from the terminal. That suits teams who build GPU images on internal hosts and want to scan them where they are built.

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 chooses which tools Claude may use. While porting a CUDA service, you can ask in the same conversation which images still carry a library version with an open advisory.

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

To try the CLI yourself, download Safeguard and follow the CLI documentation.

Quick check

Sources

• AMD, Introducing AMD CDNA 3 Architecture (white paper)
• AMD, Introducing AMD CDNA 4 Architecture (white paper)
• AMD, AMD CDNA 2 Architecture (white paper)
• AMD GPUOpen, RDNA Architecture (presentation, 2019)
• AMD GPUOpen, RDNA 3: Beyond the current gen (presentation, 2023)
• Mike Mantor, AMD Radeon HD 7970 with Graphics Core Next (GCN) Architecture, Hot Chips 24 (2012)
• AMD press release, AMD Unveils World's Most Advanced Gaming Graphics Cards, Built on Groundbreaking AMD RDNA 3 Architecture with Chiplet Design (3 November 2022)
• AMD, Radeon RX 7900 XTX product specifications
• AMD, Radeon RX 6900 XT product specifications
• AMD, Radeon RX 9070 XT product specifications
• ROCm documentation, What is ROCm?
• ROCm HIP documentation, Hardware implementation
• ROCm HIP documentation, Programming model
• ROCm HIP documentation, HIP porting guide
• ROCm HIP documentation, C++ language extensions (warp functions)
• ROCm HIPIFY documentation
• ROCm Compute Profiler documentation, RDNA workgroup processor metrics

Last reviewed: 25 October 2026. Specifications are AMD's published peak theoretical figures; check AMD's current documentation before relying on them.