AMD vs NVIDIA

L1 — Why GPUs are shaped differently

✓ Transfers completely.

Latency-vs-throughput, oversubscription as the latency-hiding mechanism, the memory wall. This is physics and economics, not vendor strategy. Nothing in Lesson 1 is NVIDIA-specific.

L2 — Inside the SM documented

◐ Structure transfers; two facts must be retired.

The hierarchy is the same shape: a chip of many independent cores, each core partitioned into SIMD units fed by schedulers, executing a fixed-width group of threads in lockstep. Divergence still costs you, still by masking.

Retire #1: warp size is not a constant. CDNA is wave64, always. RDNA is wave32 natively and supports wave64, selected by the compiler/driver per shader. Crucially, a wave64 on RDNA is not a wider execution unit — the ISA guide states it plainly: "Wave64 waves typically issue each instruction twice: once for the low half of the wave and then again for the high half." Two sequential passes over the same SIMD32. On CDNA it's four passes over a SIMD16.

Consequences: divergence granularity is coarser in wave64 (one stray lane can idle 63 others), and occupancy arithmetic depends on which mode you compiled for. Occupancy tools that report "waves" without stating the wave size are ambiguous in a way NVIDIA's never are.

Retire #2: there is no independent thread scheduling. A wavefront has one program counter plus an EXEC mask; divergence is predication, and reconvergence is compiler-managed. This is the pre-Volta model. Two consequences that point in opposite directions: the classic divergent-spinlock deadlock (lock holder sits in the not-taken branch) is still a live hazard on AMD, while sloppy warp-synchronous code without explicit syncs still works by accident, the way it did on Pascal. Porting bugs therefore travel both ways.

disputed AMD's HIP programming-model page contains the sentence "Each thread maintains its own register state … and a private program counter" — that is programming-model abstraction language, and it contradicts the lockstep statement on AMD's own sibling hardware page. Do not cite it as evidence of independent thread scheduling. The ISA guides' EXEC-mask and s_cbranch_execz chapters are the authoritative reference.

L3 — The logical map & how it lands on hardware

✓ Transfers, with a rename.

Kernel / grid / block / thread becomes kernel / grid / workgroup / work-item. Workgroups are still independent and still scheduled onto whichever core is free — the same property that lets one binary scale across GPU sizes. The global-index formula is identical; only warpSize in any warp-relative arithmetic changes.

One wrinkle unique to RDNA: WGP mode vs CU mode. In WGP mode the 128 KB LDS is one pool across all four SIMDs; in CU mode it splits into two 64 KB halves and a workgroup is confined to one CU. That is a scheduling-granularity knob with no NVIDIA counterpart.

L4 — The physical memory hierarchy documented

◐ The ladder is the same; three structural differences matter.

Registers → scratchpad → cache → DRAM, with latency hideable and bandwidth a hard ceiling, is unchanged. Arithmetic intensity and the roofline apply verbatim — AMD's own profiler (rocprof-compute) ships an empirical hierarchical roofline. Tiling is still the answer.

1. Register files are much larger. An RDNA 3/4 WGP has 768 KiB of VGPRs (192 KiB per SIMD32 → 1536 registers per lane in wave32). A CDNA 3 CU has 512 KiB. An NVIDIA H100 SM has 256 KB. That is a real, under-discussed AMD advantage for register-hungry kernels — large GEMM tiles, flash-attention.

2. The scalar unit is the biggest conceptual gap. Every AMD CU since GCN has a dedicated scalar pipeline (SALU) with its own register file (106 usable SGPRs per wave on RDNA 3, 102 on CDNA 3), its own instruction encodings, its own memory instructions, and its own cache. It handles everything uniform across the wave: loop counters, branch conditions, base addresses, EXEC-mask manipulation. NVIDIA added a uniform datapath in Turing, but it is integer-only at introduction, has 64 uniform registers, has no dedicated memory path, and — decisively — is not exposed in CUDA or even in PTX. It is a ptxas-level optimization you cannot address. On AMD, SGPR pressure is first-class ISA state and a real, tunable occupancy input. This is the cleanest single example of "architecturally different, not just differently branded."

3. LDS and cache are physically separate. NVIDIA's L1 and shared memory are one SRAM partitioned at runtime (the carveout you set with cudaFuncSetAttribute). AMD's L0 vector cache and LDS are distinct structures — there is no trade to make. LDS sizes: 128 KB per WGP on RDNA (max 64 KB per workgroup), 64 KB per CU on CDNA 3, and 160 KB per CU on CDNA 4 — a headline change aimed squarely at attention kernels that were capacity-bound on MI300X.

Plus one extra tier you don't have. Infinity Cache — AMD's internal name is MALL, Memory Attached Last Level cache — sits below L2, attached to the memory controllers. It is a memory-side cache: it caches DRAM traffic and does not participate in the GPU coherence protocol the way L2 does. 256 MB on MI300X and MI355X; 64 MB on the RX 9070 XT. The stated rationale is avoiding a wider, hotter memory bus while preserving effective bandwidth. NVIDIA's answer to the same problem was simply a very large L2.

Bank conflicts work the same way conceptually — 32 banks of 4 bytes, worst case 32-way serialization — so Lesson 5's conflict reasoning carries over directly.

L6 — __syncthreads()

✓ Transfers, same spelling.

HIP keeps the CUDA name; it lowers to s_barrier (plus s_waitcnt). The barrier is still workgroup-scoped for the same reason: workgroups are independent by design, so there is nothing to synchronize across. The two-sided tiling pattern from Lesson 6 is unchanged.

The divergent-barrier deadlock is more dangerous here, because there is no independent thread scheduling to soften it — and, as noted under L10, AMD ships nothing equivalent to synccheck to catch it.

L7 — The launch path documented

✗ The deepest divergence — and it cuts both ways.

Compilation: AMD has (had) no PTX. NVIDIA made a virtual ISA mandatory in every fatbin since CUDA 1.0, JIT-compiled by the driver — that is what gives you forward compatibility onto GPUs that didn't exist when you shipped. AMD's default has always been the opposite: compile straight to final ISA per target (gfx942, gfx1201), no JIT. AMD's docs state the contrast themselves — HIP code objects "always contain fully compiled ISA and do not require additional compilation during the load step." That is why AMD binaries enumerate explicit gfx targets and why an unlisted architecture simply fails to run.

Two things have softened this recently, and both are worth knowing:

• Generic targets (ROCm 6.4+) — gfx11-generic, gfx10-3-generic and friends: one binary across a whole family. Still native code, still no JIT; a lowest-common-denominator ISA subset, not a virtual ISA.
• SPIR-V (--offload-arch=amdgcnspirv) — the genuine PTX analogue, production as of ROCm 7.2. Portable SPIR-V is embedded in the offload bundle and JIT-compiled at first launch, cached in ~/.cache/comgr/. Caveats worth teaching: warpSize is not constexpr and architecture macros are undefined, so you migrate to runtime predicates; library coverage is still filling in.

The framing: NVIDIA made the virtual ISA the foundation and built the ABI around it. AMD is retrofitting one fifteen years later — via a Khronos standard rather than a proprietary one.

Host-side registration: identical. __hipRegisterFatBinary is a structural clone of __cudaRegisterFatBinary — same compiler-emitted constructor, same handle, same registration of kernel stubs, device code in a .hip_fatbin section. The host ABI converged; the device payload diverged. You can run the same extraction exercise Lesson 7 does:

llvm-objcopy --dump-section=.hip_fatbin=fatbin.bin ./app
clang-offload-bundler --type=o --input=fatbin.bin --unbundle \
  --targets=hip-amdgcn-amd-amdhsa--gfx1201 --output=device.o

Submission: AMD is strictly better documented — and this repairs Lesson 7's biggest caveat. Lesson 7 has to flag that NVIDIA's pushbuffer → GPFIFO → doorbell path is known largely through reverse engineering. AMD's equivalent is an open specification: AQL (Architected Queuing Language), from the HSA Foundation. Every packet is exactly 64 bytes. A ring buffer of packets lives in GPU-visible memory with explicit read_index / write_index and normative memory-ordering rules, and the doorbell is itself a specified HSA signal backed by a mapped MMIO page. The whole thing has a public C API in hsa.h.

The producer protocol is written down, not inferred: reserve a slot with hsa_queue_add_write_index_*, write the 64-byte packet with the header stored last (release), then hsa_signal_store_screlease(queue->doorbell_signal, index). Same shape as the NVIDIA mechanism Lesson 7 describes — ring buffer plus doorbell, submission in user mode after a one-time kernel-driver setup — but you can read the spec instead of trusting a reconstruction. Steady-state dispatch involves no ioctl and no syscall; the kernel driver (amdkfd, mainline Linux) only does queue creation, memory mapping and teardown.

The dispatch packet carries what you'd expect: workgroup dims, grid size, LDS and scratch bytes, a kernel_object handle, a kernarg_address pointing at the packed argument buffer (the parameter-buffer analogue, but ordinary inspectable memory), and a completion_signal. Note grid_size is in work-items, not workgroups — the opposite of <<<grid, block>>>. The runtime multiplies.

L8 — Data, completion & errors documented

◐ Transfers; AMD's completion primitives are arguably more expressive.

DMA engines, pinned memory, streams and events all have direct counterparts. hipStream_t and hipEvent_t behave like their CUDA namesakes, and the async sticky-error model is the same trap.

Underneath, AMD exposes three dependency mechanisms in the packet stream itself:

• completion_signal in the dispatch packet — hardware decrements it on retirement; the host waits with hsa_signal_wait_*.
• The barrier bit in the packet header — this packet waits for all preceding packets in the queue. Clear it and packets in one queue may overlap. This is what gives an AQL queue its in-order semantics.
• Barrier-AND / barrier-OR packets — wait on up to five arbitrary signals, resolved in hardware, across queues and agents with no host involvement.

That last one is more expressive than cudaStreamWaitEvent at the hardware interface; CUDA reaches comparable power only through the graph API. AQL packet headers also carry per-dispatch acquire/release fence scopes (system / agent / workgroup), making the memory model explicit at launch granularity — CUDA has no public equivalent.

One caution the docs are candid about: HIP streams are not 1:1 with hardware queues. Many streams multiplex onto a bounded number of hardware queues, which is where the CUDA stream mental model can mislead. inferred the exact multiplexing policy is not clearly documented.

L9 — Sharing the GPU documented

◐ Same problems, different solutions.

Queue limits. AMD documents a hardware limit of 24 user-mode compute queues on current Instinct parts — a close parallel to Hyper-Q's 32. Beyond it, the hardware scheduler round-robins, and AMD explicitly warns that evicting an inactive queue "can be in the order of milliseconds, during which the GPU is effectively blocked." That is a documented latency cliff with no NVIDIA equivalent in the public docs.

Partitioning follows physical seams. MIG partitions a monolithic die for isolation and QoS. AMD's CPX mode simply exposes each physical XCD as its own device — MI300X is eight chiplets, so the boundary already exists in silicon. Conceptually cleaner; but the granularity is fixed by die count rather than a profile catalogue, and mode changes require the GPU to be idle. NPS memory partitioning is the orthogonal axis, controlling how many NUMA domains the HBM stacks expose.

Why it exists at all is the better teaching point: each XCD is physically nearest only a couple of the eight HBM stacks, so an MI300X in single-partition mode presents as uniform but isn't. Partitioning lets you pin a tenant to its local memory. MIG solves isolation; AMD's partitioning solves isolation and genuine non-uniformity.

disputed The intermediate mode names vary across AMD documents — SPX (whole GPU) and CPX (one partition per XCD) are the reliable anchors; sources disagree on whether MI300X offers DPX only or DPX and QPX, and TPX appears in MI300A material. Check the partitioning docs for the specific part rather than repeating a mode list.

Virtualization. AMD does fractional GPU allocation through standard PCIe SR-IOV (MxGPU), one virtual function per partition. NVIDIA needed MIG as a separate proprietary construct layered over vGPU.

Graphs. hipGraph mirrors the CUDA API — stream capture, node/edge construction, instantiate, launch — and even adds hipStreamBeginCaptureToGraph, which CUDA lacks. AMD's docs state plainly that "not all features supported by CUDA graphs are yet supported in HIP." Treat it as high coverage, not parity.

L10 — Seeing inside documented

◐ Tool-for-tool parity, with one asymmetry in each direction.

Lesson 10's central discipline — match the symptom to the data source — transfers unchanged. The zoom ladder (driver telemetry → system trace → kernel counters) exists on both sides, and the counter-replay mechanism that makes Nsight Compute slow works the same way in rocprof-compute.

NVIDIA	AMD / ROCm	Notes
nvcc	hipcc / amdclang++	--offload-arch=gfx942
cuobjdump	llvm-objdump --offloading	roc-obj* deprecated in 6.4
nvdisasm	llvm-objdump -d --mcpu=gfx942	ISA is publicly documented
nvidia-smi / NVML	amd-smi / amdsmi	rocm-smi superseded, no EoL date
Nsight Systems (nsys)	rocprof-systems	was Omnitrace; renamed ROCm 6.3
Nsight Compute (ncu)	rocprof-compute	was Omniperf; renamed ROCm 6.3
CUPTI	rocprofiler-sdk	rocprof/v2/roctracer deprecated
NVTX	rocTX	no domains, colors, or payloads
cuda-gdb	rocgdb	first-class info lanes
compute-sanitizer memcheck	GPU AddressSanitizer	requires the xnack+ target suffix
racecheck · initcheck · synccheck	nothing	the largest correctness gap
DCGM	RDC + Device Metrics Exporter
nvtop	nvtop (supports AMD), amdgpu_top	read the kernel driver, not ROCm
CUDA_LAUNCH_BLOCKING	HIP_LAUNCH_BLOCKING, AMD_SERIALIZE_KERNEL=3	plus AMD_LOG_LEVEL

The gap in AMD's favour of NVIDIA: no data-race detector, no uninitialized-memory detector, no barrier-divergence checker. GPU ASan covers memory addressing only, and explicitly not LDS or private memory. This compounds badly with the wave64 porting hazards below — the exact bugs a port introduces are the ones no AMD tool will find for you.

Two things AMD has that NVIDIA doesn't: (1) fully public ISA reference manuals for every architecture generation — you can read disassembly against an authoritative document, which is simply impossible with SASS; and (2) ATT / thread trace, near-cycle-accurate per-wave instruction tracing with exact execution path and stall attribution, and an open-source decoder. Nothing public from NVIDIA compares.

Lane masks are 64 bits

The most common silent-corruption bug. __ballot() returns a 64-bit mask on AMD. CUDA code doing 1 << laneId with a 32-bit int is undefined behaviour for lanes 32–63. AMD's documented fix is 1ull, and abstracting a lane_mask_t over both wave sizes.

warpSize is not a compile-time constant

AMD's guidance is explicit: "Code should not assume a warp size of 32 or 64." Any __shared__ float smem[BLOCK/32] idiom breaks, and under SPIR-V warpSize is not even constexpr.

HIPIFY silently drops __shfl_sync masks inferred

HIP has no __syncwarp() with CUDA's semantics and no mask arguments on the *_sync intrinsics, so __shfl_sync(mask, …) maps to __shfl(…). On post-Volta CUDA that mask is load-bearing; on AMD it is meaningless. The code happens to work on AMD — and has lost the annotation that made it correct on NVIDIA. It will not survive a round trip.

__shfl's default width follows warpSize inferred

Code that omitted width on CUDA (implicitly 32) silently changes semantics on CDNA.

Don't translate __CUDA_ARCH__ numerically

AMD's guidance is to replace architecture checks with feature macros (__HIP_ARCH_HAS_DOUBLES__) or runtime hipGetDeviceProperties(). Under SPIR-V, architecture macros are undefined entirely.

Inline PTX has no translation path

The number-one hard blocker in real codebases. It must be rewritten as AMDGCN inline assembly or intrinsics — no tool does this for you.

Hardcoded shared-memory sizes

48 KB / 96 KB / 164 KB CUDA assumptions break against 64 KB (CDNA 3), 160 KB (CDNA 4) and RDNA's per-WGP allocation.

"UDNA" / the unification story disputed

The only on-record AMD statement is an executive interview at IFA in September 2024 saying the architectures would unify under the name UDNA. AMD's November 2025 Financial Analyst Day roadmap did not name the next consumer architecture at all. "RDNA 5," "CDNA 5," "UDNA 6" as concrete generations are analyst inference layered on roadmap slides. A GFX13 target existing in compiler code is verifiable; everything attached to it is not. As of 2026, RDNA 4 and CDNA 4 are still two architectures.

Is there an MPS equivalent? disputed

ROCm markets no MPS analogue, and a peer-reviewed source states multi-process workloads time-slice rather than run concurrently. But KFD runlists can hold queues from different processes simultaneously, which would in principle permit spatial sharing without any daemon. No AMD documentation resolves it cleanly. Treat as genuinely open.

RDNA 4's L1 cache disputed

Two AMD-sourced signals conflict: the RDNA 4 ISA guide still describes an L2 → L1 → L0 path, while AMD's Hot Chips 2025 deck lists only an aggregate figure and third-party analysis argues the mid-level L1 was replaced by a coalescing buffer. Do not state an RDNA 4 L1 size.

An AMD-documented fact worth sitting with

CDNA 4 cut FP64 by roughly half. AMD's own whitepaper table shows matrix FP64 going from 163.4 TFLOPS on MI300X to 78.6 TFLOPS on MI355X, labelled "~0.5×". Matrix FP64 no longer runs at twice the vector rate; the die area for wide FP64 multipliers went to FP4 and FP6 support. NVIDIA made the same call with GB200. Two vendors independently trading scientific-computing capability for AI throughput is the industry telling you something about where the volume is — and it means MI300X remains the better FP64 part.

RDNA's "AI Accelerators" are not Tensor Cores

GPUOpen states it directly: RDNA 3 "doesn't include dedicated execution units for AI acceleration like matrix cores … Wave MMA instructions effectively reuse the vector ALUs." They share schedulers, registers and issue ports with the regular shader ALUs — closer to an AVX-style ISA extension than to a Tensor Core. So "192 AI Accelerators vs 128 Tensor Cores" reads as a favourable comparison and is not one. CDNA's Matrix Cores (MFMA instructions, with a separate AGPR accumulation register file) are dedicated hardware.

Worth knowing for the optimization arc: MFMA has no asynchronous-from-LDS path and no dedicated tensor scratchpad, so operands stage through VGPRs — where NVIDIA went Volta HMMA → Hopper WGMMA (async, reads shared memory directly) → Blackwell tcgen05 with dedicated tensor memory. AMD compensates with a much larger register file and more CUs than NVIDIA has SMs. Different solutions to the same operand-bandwidth problem, and the main reason optimized AMD kernels look structurally different from optimized CUDA ones.