CDNA & MFMA — A Distil

01 The roofline model

Two hardware ceilings — compute throughput and memory bandwidth — place an upper bound on any kernel. Arithmetic intensity (FLOPs per byte) decides which one you hit.

• AI = FLOPs ÷ bytes transferred. Left of the ridge = memory-bound → optimize data reuse. Right of the ridge = compute-bound → optimize ALU use.
• Algorithm AI scaling: SAXPY O(1) · FFT O(log N) · SGEMM O(N).
• The roofline is the only diagnostic that tells you which optimization to try first. Profile, plot, then tune.

02 Profiling — which metrics, and why

rocprofv3 collects hundreds of counters. Three families diagnose a kernel: timing, compute, memory.

rocprofv3 --kernel-trace \
  --pmc SQ_WAVES VALUBusy MemUnitBusy FetchSize \
  -- ./my_kernel
# writes  <pid>_kernel_trace.csv          (timing)
#         <pid>_counter_collection.csv    (counters)

03 ROCm — the software stack underneath

ROCm = Radeon Open Compute — AMD's open-source software stack for GPU computing.

04 ROCm releases & compatibility

Which product runs which architecture, and how that maps to an LLVM target.

05 HIP — one source, two targets

Compiler that targets both AMD and NVIDIA GPUs from a single source file.

                       my_kernel.hip
                            │
                          hipcc
                       ┌────┴────┐
                       ▼         ▼
                   amdclang     nvcc
                       │         │
                       ▼         ▼
                   AMD ISA    NVIDIA PTX
        

• Same source: same .hip compiles on either platform — no #ifdefs, no per-vendor branches.
• CUDA-like syntax: __global__, <<<grid,block>>>, threadIdx, blockIdx.
• Most CUDA code ports: hipify-perl renames CUDA → HIP automatically.

06 HIP vs CUDA — side by side

Minimal vector-add kernel in both. The shape is identical — only the API prefix differs.

#include <hip/hip_runtime.h>
__global__ void add(float* a, float* b, float* c) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  c[i] = a[i] + b[i];
}
int main() {
  float *a, *b, *c;
  hipMalloc(&a, N*sizeof(float));
  hipMalloc(&b, N*sizeof(float));
  add<<<blocks, threads>>>(a,b,c);
  hipDeviceSynchronize();
}

#include <cuda_runtime.h>
__global__ void add(float* a, float* b, float* c) {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  c[i] = a[i] + b[i];
}
int main() {
  float *a, *b, *c;
  cudaMalloc(&a, N*sizeof(float));
  cudaMalloc(&b, N*sizeof(float));
  add<<<blocks, threads>>>(a,b,c);
  cudaDeviceSynchronize();
}

Kernel body is identical — only the API prefix differs.

07 ROCm tool tour

Six tools cover the full cycle: monitor → query → compile → profile → trace → compute.

08 How your code maps to the chip

The kitchen analogy — how software entities land on the hardware:

09 SIMT — the GPU execution model

Single Instruction, Multiple Threads — write code for ONE thread; the GPU runs thousands of copies, grouped into 64-thread wavefronts on CDNA.

 kernel pc → ┐
             ▼
 T0  T1  T2  T3  T4  ... T60 T61 T62 T63   ← 64 lanes, all advancing one
 ▏   ▏   ▏   ▏   ▏        ▏   ▏   ▏   ▏      instruction per cycle
 ───────────── 1 wavefront ─────────────
        

• Write code for one thread. The kernel describes work for a single thread; the hardware spawns thousands of copies — one per work-item.
• 64 threads = 1 wavefront. The scheduler issues 64 threads together. They share a program counter and advance one instruction per cycle on a SIMD.
• MI300X scale: 304 CUs × 64 ALUs — the GPU can keep ~20,000 threads in flight at once.

10 Wavefronts & latency hiding

A wavefront is the GPU's unit of scheduling: 64 threads in lockstep. Keeping many wavefronts in flight hides memory latency — on memory stall, the CU scheduler switches to another wavefront with zero context-switch cost.

11 Three bottleneck types

Every kernel sits in one of three regimes. The optimization recipe is different for each:

12 MI300 — package, XCD, compute unit

From the full board (package + HBM + I/O) down to one CU's internals. Three zoom levels:

 [ MI300X package ]
 ┌──────────────────────────────────────────────────────┐
 │  HBM3        XCD       XCD       XCD       XCD  HBM3 │
 │  HBM3        XCD       XCD       XCD       XCD  HBM3 │
 │  HBM3            PCIe 5  +  Infinity Fabric I/O  HBM3│
 │  HBM3                                             HBM3│
 └──────────────────────────────────────────────────────┘
        8 XCDs · 8 HBM3 stacks · 192 GB total HBM3

 [ one XCD ]                          [ one CU ]
 ┌────────────────────┐               ┌─────────────────┐
 │ ▢ ▢ ▢ ▢ ▢ ▢ ▢ ▢   │               │ SIMD0 SIMD1     │
 │ ▢ ▢ ▢ ▢ ▢ ▢ ▢ ▢   │   ────→       │ SIMD2 SIMD3     │
 │ ▢ ▢ ▢ ▢ ▢ ▢ ▢ ▢   │               │ + MFMA  + scalar│
 │ ▢ ▢ ▢ ▢ ▢ ▢       │               │ LDS 64 KB · L1  │
 └────────────────────┘               └─────────────────┘
   ~38 CUs / XCD                       4 SIMDs × 16 ALUs = 64 lanes
   304 active CUs total                + matrix engine + LDS + L1
        

13 MFMA — one instruction, one full matrix tile

The matrix engine computes a 16×16 tile multiply-accumulate in a single instruction — the building block of every GEMM.

• What it is: fused matrix op — one instruction does a full tile MMA.
• How it works: each CU has its own MFMA unit, pipelined alongside SIMD ALUs.
• Why it matters: 70–90% of AI compute is GEMM — rocBLAS / MIOpen use MFMA automatically.
• Variants: 16×16×16 fp16, 32×32×8 bf16, FP8 / FP6 / FP4 on CDNA 3+.

14 CDNA vs RDNA — two architectures

Same parent, different goals. The differences matter because porting CUDA code or tuning across families hits real pitfalls.

Same HIP source, different machine code: hipcc reads the ISA family.

15 Why CDNA has no graphics pipeline

Every transistor on a CDNA die goes to compute. The result is a radically different chip from a gaming GPU — even though both are “GPUs.”

The result:

16 Memory hierarchy — CDNA edition

The 10–50× bandwidth gap from registers to HBM drives every optimization strategy. Keep hot data as close to the CU as possible.

Every step down the ladder is at least an order of magnitude slower. Keep hot data as close to the CU as possible. If a kernel touches the same byte twice, it should come from LDS or L1 the second time — not HBM.

17 Memory optimization recipes

Workflow: profile first (rocprofv3) → identify the bottleneck → apply the matching technique.

18 Summary — what to remember

• CDNA wave width is 64. Every __shfl_*, occupancy calculation, and MFMA builtin assumes 64 lanes. The number flips to 32 on RDNA.
• MFMA is the CDNA matrix unit. One instruction does a full 16×16 (or 32×32) tile MAC. 70–90% of AI compute is GEMM, so MFMA is the load-bearing piece.
• HBM3 is the only DRAM tier on CDNA. 192 GB on MI300X, ~5 TB/s aggregate. No GDDR, no Infinity Cache on most CDNA parts — latency hiding via occupancy is mandatory.
• The roofline tells you which optimization to try first. Memory-bound → coalesce / fuse / reuse. Compute-bound → MFMA / lower precision / ILP. Overhead-bound → fewer fatter kernels / HIP graphs.
• rocprofv3 gives you the three counters that matter: VALUBusy, MemUnitBusy, L2CacheHit. Three numbers and you can place a kernel on the roofline.
• ROCm is open source, top to bottom. HIP is CUDA-portable, hipify-perl auto-renames most CUDA code, and the entire stack lives on GitHub.
• The same HIP source compiles for either family — but a CDNA kernel ported to RDNA needs new matrix intrinsics and a wavefront-width audit. That's the whole reason the companion deck exists.