NVIDIA vs AMD vs Intel vs Apple GPUs: Architectures Compared

NVIDIA, AMD, Intel and Apple all build GPUs from the same idea, many simple lanes fed by a deep memory hierarchy, but they differ in the details that decide performance: the width of the lane group (32 or 64), how much matrix hardware sits next to it, which number formats that hardware accepts, how much memory bandwidth feeds it, and how mature the software stack is. This guide compares the four, one dimension at a time, using only published vendor figures.

The short version: NVIDIA's data center parts pair 32 thread warps with the deepest library and framework support; AMD's Instinct parts use 64 lane wavefronts and lead on memory capacity per package; Intel's Xe parts use wide vector and matrix engines programmed through SYCL and oneAPI; Apple's GPUs share one unified memory pool with the CPU and publish far fewer numbers. Which one fits depends on whether your workload is limited by arithmetic, by memory bandwidth, by memory capacity, or by the software you already run.

Key takeaways

• Lane groups differ. NVIDIA schedules warps of 32 threads. AMD's CDNA data center GPUs run 64 thread wavefronts; its RDNA consumer GPUs run 32 or 64. Apple's SIMD-groups are 32 wide. Intel exposes sub-groups whose sizes you query from the device. Code that hard codes 32 is a common porting bug.
• Matrix engines dominate peak numbers. On every data center part here, the matrix units deliver roughly 6x to 15x the vector FP32 rate at 16 bit precision, and more again at 8 and 4 bit.
• Format support is the real differentiator for AI. Hopper and CDNA 3 added FP8. Blackwell and CDNA 4 added 4 and 6 bit formats with block scaling (NVIDIA's NVFP4, and the OCP MX formats on AMD).
• Memory decides inference. Single stream LLM decoding is bound by bytes per second, not FLOPS. Published bandwidth runs from 153 GB/s (Apple M5) to 8 TB/s (AMD MI355X and NVIDIA Blackwell).
• Ridge points keep rising. A part needs roughly 250 to 1,900 FLOPs per byte of memory traffic, depending on vendor and precision, before its matrix engines become the bottleneck.
• Apple does not publish GPU FLOPS for its M5 family, so an honest roofline for Apple can only use the bandwidth side.
• The software stack is a supply chain. Drivers, CUDA or ROCm or oneAPI runtimes, math libraries and collective libraries ship as versioned components with their own advisories, whichever vendor you choose.

How to read this comparison

Every spec in this post comes from a vendor white paper, a vendor product page, vendor documentation or a standards body, linked in the Sources section. Three conventions matter.

Peak is not sustained. Every throughput figure below is a peak theoretical number at the vendor's stated boost clock. Real kernels reach a fraction of it. We do not quote benchmark results, and we did not run any.

Dense versus sparse. NVIDIA and AMD both publish matrix throughput with and without 2:4 structured sparsity. The sparse number is twice the dense one and only applies when half of each group of four weights is zero. We use dense figures throughout, and say so when we derive one from a sparse figure.

Derived numbers are labelled. Where a vendor publishes per clock rates and a clock but no total, we multiply them and say that we did. Every computed number in this post was checked with Python. The GPU code samples follow the vendor documentation but were not executed: this post was written on a Mac without NVIDIA, AMD or Intel data center hardware. The small Python roofline script near the end did run.

If you want the background first, our explainers on AI chip architecture and what an AI accelerator is cover the general ideas, and how a GPU is made covers the packaging and HBM side. Within this series, start with how GPUs work and the CUDA programming model.

The parts we compare

Comparing "NVIDIA versus AMD" in the abstract is meaningless, so we pick specific, documented products. The data center parts get most of the attention because that is where vendors publish the most detail.

Two warnings before the tables. NVIDIA's "Blackwell" figures differ by product, because clocks, power limits and memory differ by SKU, so we show the DGX B200 and Blackwell Ultra figures separately. And AMD's CDNA 3 white paper lists 490.3 TFLOPS for MI300X matrix TF32 while the data sheet lists 653.7; the white paper's own rate of 1,024 FLOPs per clock per CU x 304 CUs x 2.1 GHz gives 653.7, so we use the data sheet value.

Dimension 1: compute unit structure and SIMD width

Every GPU has a basic building block that owns an instruction scheduler, a register file, vector lanes, matrix hardware and a scratchpad memory. NVIDIA calls it a Streaming Multiprocessor (SM), AMD a Compute Unit (CU), Intel an Xe-core, and Apple a GPU core. The names hide real structural differences.

What each vendor publishes about its building block

The NVIDIA figures come from the Hopper architecture blog's comparison table (SMs, cores per SM, Tensor Cores, register file, shared memory, threads per warp, max warps). The AMD figures come from the CDNA 3 white paper (XCD layout, 64 KB LDS), the ROCm hardware specification table (wavefront size 64, 512 KiB VGPR file, 32 KiB L1 vector cache) and the HIP hardware implementation guide, which describes "four pools, each containing slots for up to ten warps" per CU. The Intel figures come from the oneAPI GPU optimization guide and Intel's specification page (128 Xe-cores, 1,024 XMX engines, 1,024 vector engines). The Apple figures come from Apple Newsroom (core counts), the Metal feature set tables (threadgroup limits) and Apple's WWDC22 compute session, which states that threadExecutionWidth is 32 on all Apple GPUs.

Note the asymmetry. NVIDIA and AMD publish register file sizes, warp slots and per clock throughput. Intel publishes per clock throughput and thread counts but, in the guide we used, not the Xe-HPC register file size. Apple publishes core counts, memory bandwidth and API limits, and almost nothing about the core's internals. That is a legitimate design choice for a company that sells finished computers rather than chips, but it limits how far any comparison can go.

Block diagrams, side by side

These are simplified text diagrams drawn from the vendor descriptions. They show what the vendor has published, not hidden internals.

NVIDIA H100 SM (Hopper)                     AMD MI300X CU (CDNA 3)
+--------------------------------------+    +--------------------------------------+
| Warp schedulers, dispatch            |    | Scheduler, 4 pools x 10 wave slots    |
| Register file 256 KB (64K x 32 bit)  |    | VGPR file 512 KiB (split over 4 SIMD) |
| 128 FP32 | 64 FP64 | 64 INT32 lanes  |    | 4 x SIMD16 vector units (wave64 in    |
| 4 x 4th gen Tensor Core              |    |   4 cycles per instruction)           |
| Tensor Memory Accelerator (TMA)      |    | Matrix Core unit                      |
| L1 / shared memory 256 KB            |    | LDS 64 KiB    | L1 vector cache 32 KiB|
|   (up to 228 KB as shared memory)    |    | Instruction cache 64 KB shared by 2 CUs|
+--------------------------------------+    +--------------------------------------+
  x132 per H100 SXM5, 50 MB L2 per GPU        x38 per XCD, 4 MB L2 per XCD, 8 XCDs,
                                              256 MB Infinity Cache on the IODs

Intel Xe-core (Xe-HPC, Max 1550)            Apple GPU core (M5 family)
+--------------------------------------+    +--------------------------------------+
| 8 vector engines (XVE), 512 bit each |    | Shader core (internals not published) |
|   8 hardware threads per XVE         |    | SIMD-groups of 32 threads             |
| 8 XMX matrix engines, 4096 bit each  |    | Neural Accelerator (one per core, M5) |
| L1 cache / SLM 512 KB                |    | Threadgroup memory up to 32 KB (API)  |
| Load/store 1024 B per cycle          |    | Ray tracing engine (3rd gen on M5)    |
+--------------------------------------+    +--------------------------------------+
  16 per Xe-slice, 4 slices per stack,        10 (M5) to 80 (M5 Ultra) cores,
  2 stacks, up to 408 MB last level cache     unified memory shared with the CPU

Why the 32 versus 64 lane choice matters

A warp or wavefront is the unit of instruction issue. All lanes in it execute the same instruction; when lanes take different branches, the hardware runs each path with the inactive lanes masked off. Two consequences follow from the width.

Divergence cost. With 64 lanes, a single divergent lane forces the whole wavefront through both paths, and the chance that at least one lane in a group diverges rises with group size. Branchy code, such as tree traversal or sparse indexing, tends to suit 32 wide groups better.

Latency hiding per instruction. The HIP documentation explains that CDNA's SIMD units are 16 lanes wide, so one wave64 instruction occupies a SIMD for four cycles, and with four SIMDs per CU "a new instruction can issue to each SIMD every cycle." Each issued instruction covers more work, which suits dense, regular arithmetic.

AMD's consumer RDNA parts support both: the ROCm table lists "32 or 64" as the wavefront size for RDNA 3 and RDNA 4. CDNA data center parts list 64 only.

The practical lesson is to never hard code the width. In CUDA and HIP, the built in warpSize gives the value. In SYCL you query the sub-group sizes the device supports. In Metal you read threadExecutionWidth on the pipeline state, or threads_per_simdgroup inside the shader, which is what Apple's WWDC20 Apple silicon session recommends.

// HIP: a warp-size agnostic warp reduction.
// Follows the HIP programming guide; not executed for this post.
__device__ float warp_sum(float v) {
    for (int offset = warpSize / 2; offset > 0; offset /= 2)
        v += __shfl_down(v, offset);   // 64 lanes on CDNA, 32 on NVIDIA
    return v;
}

On NVIDIA the CUDA version uses __shfl_down_sync with a 32 bit lane mask. A kernel that writes 16 as the starting offset, assuming 32 lanes, compiles fine for NVIDIA and then silently combines only half of each wavefront on CDNA if ported carelessly.

Worked occupancy example

Occupancy is the fraction of a block's thread slots that are filled with resident threads (inside an NVIDIA SM covers it in depth). More resident threads give the scheduler more work to switch to while others wait on memory. Registers are usually the limiting resource, so here is the same kernel register budget on two parts with published register files.

Assume a kernel uses 128 32 bit registers per thread.

NVIDIA H100. The SM has 65,536 32 bit registers (256 KB). At 128 registers per thread, 65,536 / 128 = 512 threads fit, which is 16 warps. The SM can hold 64 warps (2,048 threads), so occupancy is 16 / 64 = 25 percent. At 64 registers per thread, 1,024 threads fit and occupancy is 50 percent.

AMD MI300X. The CU's 512 KiB VGPR file is split across four SIMDs, 128 KiB each. A VGPR for a wave64 holds 64 lanes x 4 bytes = 256 bytes, so each SIMD holds 128 KiB / 256 B = 512 VGPRs worth of wave registers. At 128 VGPRs per wave, 512 / 128 = 4 waves fit per SIMD, 16 per CU, against 40 slots: 40 percent of slots, or 4 of 10 per SIMD. At 64 VGPRs, 8 waves fit per SIMD.

These are upper bounds from register use alone. Shared memory or LDS use, block size, and the vendor's allocation granularity can lower them; Each vendor's profiler reports the exact figure per kernel. And high occupancy is not a goal in itself. The CUDA C++ Best Practices Guide puts it directly: "Higher occupancy does not always equate to higher performance," because a lower occupancy kernel has more registers per thread and can hide latency through instruction level parallelism. The same guide notes that register allocations are rounded up to the nearest 256 registers per warp, one of the granularity effects that makes the simple division above an upper bound.

Intel and Apple do not publish register file sizes in the sources we used, so we cannot run the same arithmetic for them. Intel's guide gives 8 hardware threads per vector engine; Apple's Metal API reports the maximum threads per threadgroup for each compiled pipeline.

Dimension 2: matrix engines and their formats

Matrix engines (Tensor Cores, Matrix Cores, XMX engines, Apple's Neural Accelerators) execute small matrix multiply and accumulate operations in one instruction, usually on 16 bit or narrower inputs with wider accumulators. They are where nearly all of the advertised FLOPS live.

Peak dense matrix throughput by format

Notes on the table:

• H100. NVIDIA's product page lists sparse values with an asterisk ("With sparsity"); we halve them. The page lists FP64 Tensor Core at 67 TFLOPS and FP32 at 67 TFLOPS.
• B200. The DGX B200 page lists the whole eight GPU system: "FP4 Tensor Core: 144 PFLOPS | 72 PFLOPS" (sparse | dense) and "FP8 Tensor Core: 72 PFLOPS" with a footnote that this is sparse and "Dense performance is ½ sparse spec shown." Dividing by eight gives 9 PF dense FP4 and 4.5 PF dense FP8 per GPU. NVIDIA's Blackwell Ultra blog lists "Blackwell" at 10 PF dense NVFP4 and 5 PF dense FP8, which reflects a different SKU and power limit (it lists up to 1,200 W for Blackwell). Both are NVIDIA figures; they describe different products.
• MI300X and MI355X. These come from AMD's data sheet and CDNA 4 white paper, which list dense and sparse values separately. The CDNA 4 per clock table prints 16,834 FLOPs per clock per CU for MXFP4; 16,384 x 256 CUs x 2.4 GHz = 10.07 PF matches the published 10 PF, so the 16,834 looks like a typo in the white paper.
• Max 1550. Intel's spec page lists no FLOPS. The oneAPI optimization guide says each Xe-core "delivers 8192 int8 and 4096 FP16/BF16 operations/cycle" on its matrix engines; 4,096 x 128 Xe-cores x 1.6 GHz (the "Graphics Max Dynamic Clock" on Intel's spec page) = 838.9 TF. That is our derivation, not an Intel published total.
• Arc B580. Intel's spec page lists 233 peak INT8 TOPS.
• Apple. Apple does not publish GPU matrix throughput for the M5 family. It says M5 delivers "over 4x the peak GPU compute performance for AI compared to M4" and that M5 Ultra has "up to 4.3x the peak AI compute performance when compared to M3 Ultra." Ratios without a base value cannot go in a roofline.

Format support, generation by generation

Sources: NVIDIA's Hopper blog (FP8, FP16, BF16, TF32, FP64, INT8 Tensor Core types) and Blackwell Ultra blog (FP8, FP6, NVFP4); ROCm's precision support page (matrix core data type matrix per architecture); the CDNA 3 white paper, which notes CDNA 3 supports both FP8 variants; the Intel optimization guide; and Apple's Metal feature set tables, which list SIMD-scoped matrix multiply operations from the Apple7 GPU family and Metal 4 tensors from Apple7.

One detail that trips people up: ROCm's documentation says CDNA 3 "natively supports FP8 FNUZ (E4M3 and E5M2)," which differs from the OCP FP8 variant. FNUZ has no infinities or negative zero. CDNA 4 supports the OCP variants. A model quantized to FP8 for one encoding needs care when moved to hardware that expects the other.

Block scaled 4 bit formats, worked

At 4 bits, a single scale per tensor loses too much accuracy, so both NVIDIA and the Open Compute Project define block scaling.

OCP MX formats (used by AMD CDNA 4). The OCP Microscaling Formats v1.0 specification defines MXFP4 as blocks of k = 32 elements, each element 4 bits (E2M1), sharing one 8 bit E8M0 scale. Storage per element is (32 x 4 + 8) / 32 = 4.25 bits, or 0.53125 bytes.

NVFP4 (NVIDIA Blackwell). NVIDIA's NVFP4 article describes "a fine-grained FP8 scale to each 16-value micro-block and a second-level FP32 scalar per tensor," with the block scale in E4M3. Storage per element is (16 x 4 + 8) / 16 = 4.5 bits, or 0.5625 bytes, ignoring the single per tensor scalar.

So a 70 billion parameter model needs about 37.2 GB in MXFP4 and about 39.4 GB in NVFP4, against 70 GB in FP8 and 140 GB in BF16. The smaller block in NVFP4 costs 0.25 bits per element and buys finer scaling; the E4M3 scale can represent non power of two factors, whereas E8M0 is a pure exponent. Which trade works better depends on the model, and the vendors' own accuracy claims are claims, not results we checked.

Dimension 3: memory systems

For many real workloads, and for nearly all single stream LLM inference, memory bandwidth and capacity matter more than FLOPS. Our GPU memory and interconnects post explains HBM, GDDR, NVLink and unified memory in detail.

Three memory philosophies

Discrete HBM with a large cache (NVIDIA, AMD, Intel data center). HBM stacks sit on the same package as the GPU dies, connected through a silicon interposer or bridge. NVIDIA puts a large L2 (50 MB on H100) in front of HBM. AMD adds a further level: the CDNA 3 white paper describes a 256 MB Infinity Cache on the I/O dies, built from 128 channels of 2 MB, with a peak bandwidth of 17.2 TB/s, more than three times the HBM bandwidth behind it. Intel's Max 1550 lists up to 408 MB of last level cache. These caches matter for kernels whose working set fits: a 50 MB L2 holds a meaningful slice of a layer's activations; a 256 MB cache holds more.

Unified memory on a system on chip (Apple). Apple's M series chips put CPU, GPU and Neural Engine on one die family with one LPDDR pool. Apple's M5 release says the unified memory architecture "enables the entire chip to access a large single pool of memory." The advantage is capacity per dollar and zero copies between CPU and GPU: an M5 Ultra Mac Studio can hold 512 GB that the GPU can address directly, more than any single data center GPU in this table. The cost is bandwidth: 1.2 TB/s at most, against 8 TB/s on the newest HBM parts.

LPDDR on a discrete inference card (Intel's announced D-series). Intel's OCP 2026 page describes up to 480 GB of LPDDR5X at 1.5 TB/s with 1.3 PF of FP8 compute, positioned for "sustained, memory-intensive AI inference." That trades bandwidth for capacity and cost, much like Apple's choice but on a PCIe card. We list it as announced; we did not find independent measurements.

Worked number: how long to read all of memory once

A useful sanity figure is the time to stream a device's full memory once at peak bandwidth. It is the floor on one decode step for a model that fills the device.

The pattern is the point: as capacity grows faster than bandwidth, the "read everything" time grows too. A 512 GB unified pool can hold a very large model, but it takes over 400 ms to sweep it once, so it suits low batch, low token rate use such as local experimentation rather than high throughput serving.

Dimension 4: interconnect

Once a model or a training job spans more than one GPU, the link between GPUs becomes the next bottleneck. Each vendor has a proprietary scale up fabric, plus PCIe to the host.

A few notes. NVIDIA's Hopper blog explains that a fourth generation NVLink link uses two differential pairs per direction for 25 GB/s each way, so 18 links give 900 GB/s total. AMD's CDNA 4 paper gives 76.8 GB/s per direction per link (153.6 GB/s bidirectional); seven such links give the 1,075.2 GB/s P2P figure, and the eighth link is configured for PCIe. Both vendors count bidirectional bandwidth, so the headline numbers are comparable, though topology differs: NVIDIA's NVLink Switch builds a full crossbar domain (72 GPUs in an NVL72 rack, per NVIDIA), while AMD's eight GPU nodes connect each GPU directly to the other seven.

The GPU cost per million tokens estimate on this site shows how these memory and interconnect figures feed into serving cost.

Dimension 5: programming model and libraries

Hardware is only half the comparison. The other half is how you program it and which tuned libraries you get.

Sources: NVIDIA's cuBLAS, cuDNN, NCCL and PTX documentation; ROCm's HIP, HIPIFY, hipBLASLt, rocWMMA, MIOpen and RCCL documentation; the UXL Foundation's oneDNN, oneMath and oneCCL documentation and Khronos SYCL page; Apple's Metal Performance Shaders Graph documentation and Metal feature set tables; the PyTorch 2.14 notes for the MPS backend ("Metal Performance Shaders Graph framework and tuned kernels provided by Metal Performance Shaders") and for Intel GPUs; and the PyTorch install page, which notes that "the ROCm build of PyTorch uses the same semantics at the python API level" as CUDA.

The same kernel four ways

To show how close (and how far apart) the models are, here is SAXPY (y = a * x + y) in each. These follow the vendor programming guides and were not executed for this post.

CUDA C++

__global__ void saxpy(int n, float a, const float* x, float* y) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) y[i] = a * x[i] + y[i];
}
// Host side launch: 256 threads per block.
// saxpy<<<(n + 255) / 256, 256>>>(n, a, d_x, d_y);

HIP C++ (AMD; also compiles for NVIDIA)

#include <hip/hip_runtime.h>
__global__ void saxpy(int n, float a, const float* x, float* y) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) y[i] = a * x[i] + y[i];
}
// hipLaunchKernelGGL(saxpy, dim3((n + 255) / 256), dim3(256), 0, 0, n, a, d_x, d_y);
// Build for MI300X:  hipcc --offload-arch=gfx942 saxpy.cpp
// Build for MI355X:  hipcc --offload-arch=gfx950 saxpy.cpp

The gfx942 and gfx950 LLVM target names come from ROCm's hardware specification table.

SYCL (Intel oneAPI DPC++; also runs on other back ends)

#include <sycl/sycl.hpp>
void saxpy(sycl::queue& q, int n, float a, const float* x, float* y) {
    q.parallel_for(sycl::range<1>(n), [=](sycl::id<1> i) {
        y[i] = a * x[i] + y[i];
    }).wait();
}
// Query sub-group sizes the device supports:
// auto sizes = q.get_device().get_info<sycl::info::device::sub_group_sizes>();

Metal Shading Language (Apple)

#include <metal_stdlib>
using namespace metal;
kernel void saxpy(device const float* x [[buffer(0)]],
                  device float*       y [[buffer(1)]],
                  constant float&     a [[buffer(2)]],
                  constant uint&      n [[buffer(3)]],
                  uint i [[thread_position_in_grid]]) {
    if (i < n) y[i] = a * x[i] + y[i];
}
// Host side (Swift or Objective-C): read pipeline.threadExecutionWidth (32 on Apple GPUs)
// and pipeline.maxTotalThreadsPerThreadgroup before choosing the threadgroup size.

Four observations:

1. CUDA and HIP are nearly identical by design. AMD's HIP documentation describes "an interface closely aligned with NVIDIA CUDA" so developers can "efficiently migrate existing CUDA code," and the HIPIFY tools convert CUDA source to HIP. The differences show up in warp width, inline PTX, and library coverage.
2. SYCL is single source standard C++. Kernels are lambdas, and the same code targets Intel GPUs, CPUs and, through other back ends, other vendors' GPUs. The cost is that peak performance on any one vendor usually needs that vendor's extensions.
3. Metal separates host and device code. Kernels are compiled from the Metal Shading Language, and the host API is Swift or Objective-C. There is no single source C++ path comparable to CUDA's.
4. Higher level languages sit above all of them. Triton describes itself as "a language and compiler for parallel programming" that provides a Python based environment for writing custom DNN kernels. Tools like this reduce, but do not remove, the cost of moving between vendors; check each tool's documentation for the back ends it supports today.

Matrix engines from code

A portable way to reach the matrix units is to call the vendor BLAS or DNN library. Writing to them directly differs per vendor:

// CUDA WMMA (warp-level matrix multiply-accumulate), 16x16x16 FP16 -> FP32.
// From the CUDA C++ Programming Guide's WMMA API; not executed for this post.
#include <mma.h>
using namespace nvcuda;
__global__ void tile_mma(const half* A, const half* B, float* C) {
    wmma::fragment<wmma::matrix_a, 16, 16, 16, half, wmma::row_major> a;
    wmma::fragment<wmma::matrix_b, 16, 16, 16, half, wmma::col_major> b;
    wmma::fragment<wmma::accumulator, 16, 16, 16, float> c;
    wmma::fill_fragment(c, 0.0f);
    wmma::load_matrix_sync(a, A, 16);
    wmma::load_matrix_sync(b, B, 16);
    wmma::mma_sync(c, a, b, c);
    wmma::store_matrix_sync(C, c, 16, wmma::mem_row_major);
}

// Metal: SIMD-group 8x8 matrix multiply (Apple7 family and later).
// Types per the Metal Shading Language spec; not executed for this post.
#include <metal_stdlib>
using namespace metal;
kernel void tile8(device const half* A [[buffer(0)]],
                  device const half* B [[buffer(1)]],
                  device float*      C [[buffer(2)]]) {
    simdgroup_half8x8  a, b;
    simdgroup_float8x8 c = make_filled_simdgroup_matrix<float, 8, 8>(0.0f);
    simdgroup_load(a, A, 8);
    simdgroup_load(b, B, 8);
    simdgroup_multiply_accumulate(c, a, b, c);
    simdgroup_store(c, C, 8);
}

The tile shapes alone show the difference in granularity: a CUDA WMMA fragment covers 16x16x16 per warp, Apple's SIMD-group matrix covers 8x8 per SIMD-group, and on Hopper the newer wgmma PTX instructions operate across a warpgroup of four warps. PTX defines a warpgroup as "a set of four contiguous warps." On AMD, the rocWMMA library offers a WMMA style API over the matrix cores. On Intel, oneDNN is the documented route to XMX in the sources we used.

Dimension 6: ecosystem maturity

"Ecosystem maturity" is easy to argue about and hard to measure. We use evidence that a reader can check: which frameworks document official support, which libraries exist, and how much of the stack is vendor maintained.

Read the table with the trade-offs in mind.

• NVIDIA has the longest track record of the four in general purpose GPU computing, and its library set (cuBLAS, cuDNN, NCCL and more) is the broadest. The cost is lock in: inline PTX and CUDA only libraries do not move.
• AMD has made the HIP path close to CUDA, and ROCm's libraries are documented publicly per GPU target. The practical risk is that a given model, kernel or library version may land on CUDA first, so teams should check that the exact versions they need are supported on the exact gfx target they buy.
• Intel bets on open standards (SYCL, oneAPI under the UXL Foundation); see our Intel GPU architecture post for the Xe-core detail. That makes code portable, but Intel's data center GPU roadmap has shifted: the Max 1550 launched in Q1 2023, and Intel's newest announced data center GPU, the D-series code-named Crescent Island, is an inference part with LPDDR5X rather than HBM. Plan around products that are shipping, not roadmaps.
• Apple is a single vendor, single platform stack. Metal, MPS and MLX are well integrated and capable for local inference and development, and the unified memory pool is genuinely unusual. But there is no multi GPU fabric, no published FLOPS, and no path to run Metal code on non Apple hardware.

Worked rooflines from published peaks

A roofline model says that attainable throughput is the lower of two limits: peak compute, and memory bandwidth times arithmetic intensity (FLOPs per byte of memory traffic). The crossover is the ridge point: peak FLOPS divided by peak bandwidth. Kernels with intensity below the ridge are memory bound; above it, compute bound.

Ridge points

All values are dense peaks from the sources above; the Max 1550 values are our derivation from Intel's per cycle figures and clock.

Two things jump out. First, the vector FP32 ridge points (16 to 31 FLOP per byte) are an order of magnitude below the matrix ridge points, which is why ordinary FP32 code is so often memory bound while matrix multiplies are not. Second, every step down in precision doubles the compute peak without changing bandwidth, so the ridge doubles too. Low precision formats make kernels faster only if the kernel had enough reuse to be compute bound in the first place; otherwise their main benefit is that the data itself is smaller.

A roofline calculator you can run

This script runs on any machine with Python 3; we ran it on a Mac to produce the tables in this section.

# roofline.py: attainable throughput from published peaks (dense).
PARTS = {
    # name: (peak FLOP/s, peak bytes/s)
    "H100 SXM BF16":   (989.5e12, 3.35e12),
    "H100 SXM FP8":    (1979e12,  3.35e12),
    "MI300X BF16":     (1307.4e12, 5.3e12),
    "MI300X FP8":      (2614.9e12, 5.3e12),
    "MI355X BF16":     (2.5e15,   8.0e12),
    "MI355X FP8":      (5.0e15,   8.0e12),
    "Max 1550 BF16":   (838.9e12, 3.2768e12),   # derived from per-cycle figures
}

def attainable(peak, bw, intensity):
    return min(peak, bw * intensity)

for name, (peak, bw) in PARTS.items():
    ridge = peak / bw
    row = [attainable(peak, bw, ai) / 1e12 for ai in (1, 16, 128, 512)]
    print(f"{name:16s} ridge {ridge:7.1f}  "
          + "  ".join(f"AI={ai}:{v:8.1f} TF" for ai, v in zip((1, 16, 128, 512), row)))

At an intensity of 1 FLOP per byte (typical of batch 1 LLM decoding in BF16), every part delivers only its bandwidth in TFLOPS: 3.35 TF on H100, 5.3 TF on MI300X, 8 TF on MI355X. That is under 1 percent of their BF16 matrix peaks.

Arithmetic intensity of a square matrix multiply

For C = A x B with all three matrices n x n in BF16, the multiply does 2n³ FLOPs and must at minimum read A and B and write C: 3n² elements of 2 bytes. Intensity is 2n³ / 6n² = n / 3 FLOP per byte.

This is a lower bound on traffic, assuming perfect on chip reuse; real kernels approach it through tiling in shared memory or LDS and registers. The table shows why large GEMMs reach a high fraction of peak on every vendor's part, and why small ones do not; our CUDA matrix multiply walkthrough shows the tiling that approaches this bound.

LLM decoding: a memory bound worked example

When a transformer generates one token for one user, every weight is read once and used for about two FLOPs (a multiply and an add). With BF16 weights that is about 1 FLOP per byte; with FP8 weights, about 2. Both are far below every ridge point above, so the token rate is capped by bandwidth:

upper bound on tokens per second = memory bandwidth / bytes of weights read per token

For a dense 70 billion parameter model this ignores the KV cache, activations and all overheads, so real numbers are lower. It is a ceiling, not a benchmark.

Serving many users changes the picture. With a batch of B requests, each weight read serves B tokens, so intensity rises to roughly B FLOP per byte in BF16 (2B with FP8 weights). On H100 the BF16 ridge of 295.4 means a batch around 295 before the matrix engines, rather than memory, become the limit; on MI300X, about 247. In practice KV cache traffic grows with batch and context length, so the crossover comes later, and memory capacity (how much KV cache fits) often becomes the binding constraint first. That is a big reason the large capacity parts (MI300X, MI355X, Blackwell Ultra, and at the low end Apple's 512 GB M5 Ultra) are attractive for inference.

Which to choose for what

Here is how the published figures map to common jobs, with the main trade-off for each.

Large scale training

What matters: dense BF16 and FP8 matrix throughput, scale up bandwidth, collective library maturity, and the cluster software around them.

• NVIDIA's Hopper and Blackwell parts offer FP8 (and on Blackwell, FP4 and FP6), NVLink domains up to 576 GPUs, and NCCL. The trade-off is cost and dependence on one vendor's closed libraries.
• AMD's MI300X and MI355X match or exceed the NVIDIA parts on several published peaks and on memory capacity, with RCCL and open source ROCm. The trade-off is that you must confirm your training stack's exact versions on the exact gfx target, and eight GPU direct connect topologies differ from switched NVLink domains.
• Intel's Max 1550 offers BF16 XMX and Xe Link, but Intel's newest announced data center GPU is aimed at inference. Evaluate support timelines before committing training infrastructure.
• Apple is not a training platform at cluster scale; there is no GPU to GPU fabric.

Inference serving

What matters: memory bandwidth (token rate), memory capacity (model size and KV cache), low precision format support, and serving software.

• Large capacity HBM parts (MI300X at 192 GB, MI355X and Blackwell Ultra at 288 GB) let a single GPU hold larger models and longer contexts.
• FP8 and 4 bit formats halve or quarter the bytes per weight and therefore raise the bandwidth bound, but they require matching hardware support and a quantization scheme the model tolerates. Remember the FNUZ versus OCP FP8 encoding difference on AMD.
• Intel's announced D-series targets this segment with high capacity LPDDR5X at lower bandwidth. It is an announcement until it ships.

Graphics and gaming

What matters: rasterization, ray tracing, drivers, game support. Data center parts are a poor fit; H100's architecture blog notes that only two TPCs on the H100 SXM5 and PCIe parts are graphics capable. Consumer parts (AMD RDNA 4, Intel Arc Xe2, NVIDIA's GeForce line) and Apple's integrated GPUs with hardware ray tracing are built for this. The trade-off between them is mostly drivers, price and feature support, which this post does not rank.

Mobile and laptop

What matters: performance per watt and memory shared with the CPU. Apple's M series and other integrated GPUs win here by design. Apple's unified memory is especially useful for local model work: a laptop with 128 GB of unified memory can load models that need two data center GPUs. The trade-off is bandwidth (614 GB/s at most on M5 Max) and Metal only software.

HPC and scientific computing

What matters: FP64. Here the picture is the reverse of AI. MI300X lists 163.4 TF of FP64 matrix and 81.7 TF FP64 vector; H100 SXM lists 67 TF FP64 Tensor Core and 34 TF FP64. On CDNA 4, AMD halved the FP64 matrix rate per CU (78.6 TF on MI355X) while doubling the low precision rates, and the Blackwell Ultra chip comparison does not lead with FP64 at all. Teams with FP64 heavy codes should read each vendor's FP64 line specifically, not the AI headline.

A compact decision table

Common mistakes

1. Comparing sparse figures to dense ones. NVIDIA's H100 page leads with sparse numbers marked by an asterisk. AMD's data sheet lists both. Halve the sparse figure before comparing.
2. Treating peak as achievable. Peaks assume the boost clock and perfect instruction mix. Sustained figures under power limits are lower, and vendors do not publish them as specs.
3. Hard coding 32 lanes. CDNA wavefronts are 64. Use warpSize, sub-group queries or threadExecutionWidth.
4. Ignoring memory capacity. A part with higher FLOPS but too little memory for the model and KV cache will split across GPUs and pay interconnect costs.
5. Assuming FP8 means one format. E4M3 and E5M2 differ, and AMD CDNA 3 uses the FNUZ variants while CDNA 4 and NVIDIA use OCP style encodings.
6. Reading Apple ratios as FLOPS. "4x the peak GPU compute for AI compared to M4" is a ratio; without Apple publishing a base figure you cannot place it on a roofline.
7. Comparing system figures to chip figures. The DGX B200 page lists eight GPU totals. Divide before comparing to a single MI355X or H100.

FAQ

Is NVIDIA faster than AMD for AI?

It depends on the part, the precision and the workload. On published dense peaks, MI300X lists 1,307.4 TF BF16 against H100 SXM's 989.5 TF, and MI355X lists 2.5 PF BF16. NVIDIA's Blackwell parts lead on FP4 peaks. Real results depend on software maturity, which this post does not benchmark.

What is the difference between a warp and a wavefront?

Both are the group of threads that execute one instruction together. NVIDIA's warp is 32 threads. AMD's CDNA wavefront is 64 threads; RDNA supports 32 or 64. Apple's SIMD-group is 32.

Do Apple GPUs have tensor cores?

Apple says each GPU core in the M5 family has a Neural Accelerator, which Apple positions for GPU based AI work. Earlier Apple GPUs (Apple7 family onward) support SIMD-group 8x8 matrix multiply in Metal. Apple does not publish the throughput of either.

Can CUDA code run on AMD GPUs?

Not directly. HIP is a CUDA like C++ API, and AMD's HIPIFY tools convert CUDA source to HIP, which then compiles for AMD GPUs. Inline PTX and CUDA only libraries need replacements.

What is SYCL and why does Intel use it?

SYCL is a Khronos standard for single source C++ heterogeneous programming. Intel's oneAPI DPC++ compiler implements it for Intel GPUs and CPUs. Libraries such as oneDNN and oneCCL (now hosted by the UXL Foundation) sit on top of it.

Why does memory bandwidth matter more than FLOPS for LLM inference?

Generating one token for one user reads every weight once and does about two FLOPs per weight, so the arithmetic intensity is about 1 to 2 FLOPs per byte, far below every GPU's ridge point. The token rate is capped at bandwidth divided by model bytes.

What are MXFP4 and NVFP4?

Both are 4 bit floating point formats with block scaling. OCP's MXFP4 shares an 8 bit E8M0 scale across 32 elements; NVIDIA's NVFP4 shares an E4M3 FP8 scale across 16 elements plus a per tensor FP32 scale. AMD CDNA 4 supports MX formats; NVIDIA Blackwell supports NVFP4.

Which GPU is best for a local LLM on a laptop?

For model size, unified memory capacity matters most, which favours Apple's M series with large memory configurations. For token rate, bandwidth matters most. A discrete laptop GPU has more bandwidth but usually much less memory.

Where Safeguard fits

Safeguard does not design, program or speed up GPUs, and nothing in this comparison depends on it. The honest connection is the software that every one of these GPUs needs. A CUDA, ROCm, oneAPI or Metal based workload ships a kernel driver, a user space runtime, math libraries such as cuBLAS or hipBLASLt, DNN libraries such as cuDNN, MIOpen or oneDNN, collective libraries such as NCCL or RCCL, a framework build such as PyTorch, and usually a container base image. Each is a versioned component, and each has had published advisories.

Safeguard generates and analyses SBOMs in SPDX and CycloneDX, scans repositories, container images and SBOM files (for example safeguard scan --image <name> for a serving image), matches components against vulnerability databases, shows which projects contain a given component and version, applies reachability analysis, and enforces policies and gates. That covers the open source frameworks, Python packages and container layers in a GPU workload. It does not cover firmware on the GPU board, the host kernel driver installed outside your images, or closed vendor binaries you cannot inventory; those updates still come from the GPU vendor and your OS distribution, so track their bulletins separately.

The GPU software stack changes quickly: a new CUDA or ROCm release pulls new library versions into every image that rebuilds on it, and a model serving framework upgrade can drag in dozens of Python dependencies. Keeping that stack inventoried and patched is ordinary software supply chain work, and it is where the steps below apply.

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
• 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
• 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. It opens the interactive terminal, reads a dependency finding and its advisory, asks before any command that changes something, edits the manifest to a fixed version and runs the install and tests with your approval. The same flow applies to a Python requirements file in a model serving repository that pins a vulnerable framework or library version.

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 scans from the terminal. That suits a GPU workstation or build host where you want to scan a CUDA or ROCm based container image locally before it is pushed.

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 lets you choose which tools Claude can use. It is a quick way to ask questions about the components in your GPU serving images without leaving a chat.

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

To try it on your own stack, download Safeguard or read the CLI documentation.

Quick check

Sources

• NVIDIA Hopper Architecture In-Depth (NVIDIA Technical Blog, 22 March 2022)
• NVIDIA H100 Tensor Core GPU product page and specifications
• NVIDIA DGX B200 specifications
• NVIDIA Blackwell architecture page
• Inside NVIDIA Blackwell Ultra (NVIDIA Technical Blog, 22 August 2025)
• Introducing NVFP4 (NVIDIA Technical Blog)
• CUDA C++ Programming Guide
• CUDA C++ Best Practices Guide
• PTX ISA
• cuBLAS documentation
• cuDNN documentation
• NCCL documentation
• AMD CDNA 3 architecture white paper
• AMD Instinct MI300X data sheet
• AMD CDNA 4 architecture white paper
• ROCm GPU hardware specifications
• ROCm data type and precision support
• HIP hardware implementation
• HIP documentation
• hipBLASLt documentation
• MIOpen documentation
• RCCL documentation
• rocWMMA documentation
• HIPIFY documentation
• Intel oneAPI GPU Optimization Guide: Intel Xe GPU architecture
• Intel Data Center GPU Max 1550 specifications
• Intel Arc B580 Graphics specifications
• Intel at OCP 2026 (Intel Newsroom)
• oneDNN documentation (UXL Foundation)
• oneMath documentation (UXL Foundation)
• oneCCL documentation (UXL Foundation)
• SYCL (Khronos Group)
• Apple introduces M4 Pro and M4 Max (Apple Newsroom, 30 October 2024)
• Apple unleashes M5 (Apple Newsroom, 15 October 2025)
• Apple introduces new Mac Studio with M5 Max and M5 Ultra (Apple Newsroom, August 2026)
• Metal feature set tables (Apple)
• Scale compute workloads across Apple GPUs (WWDC22)
• Bring your Metal app to Apple silicon Macs (WWDC20)
• Metal Shading Language Specification
• Metal Performance Shaders Graph documentation
• MLX documentation
• PyTorch: MPS backend notes
• PyTorch: Getting started on Intel GPU
• PyTorch: Get started locally
• Triton language documentation
• OCP Microscaling Formats (MX) v1.0 specification

Last reviewed: 25 October 2026. Vendor specifications change between SKUs and revisions; check the linked vendor pages before relying on a specific figure.