CDNA & MFMA

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; a CDNA SIMD is 16 lanes wide, so one wave64 instruction occupies it for 4 cycles.
• MI300X scale: 304 CUs × 4 SIMDs × 16 lanes = 19,456 ALU lanes executing at once, with many more threads resident per SIMD so a stalled wave can be swapped out.

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: GEMMs dominate transformer prefill and training, and rocBLAS / hipBLASLt route them to MFMA.
• Variants: 16×16×16 fp16, 32×32×8 bf16, FP8 on CDNA 3, FP6 / FP4 on CDNA 4.

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

Every step from registers to HBM is a big drop in bandwidth and a big rise in latency, and that gap drives every optimization strategy: keep hot data as close to the CU as possible.

Capacities from AMD's ROCm GPU specification tables; HBM size and peak bandwidth from AMD's Instinct product pages.

If a kernel touches the same byte twice, the second touch should come from registers, LDS or L1, not HBM.

17 Memory optimization recipes

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

18 Where MFMA sits in hipfire 0.4.1

hipfire's tuned kernel families target RDNA3/3.5 and RDNA4 through WMMA. On CDNA the engine runs correctly through a portable wave64 path. It knows gfx908 and gfx940–gfx942 (plus Vega gfx906) as wave64-native targets, and reaches MFMA in two ways:

• rocBLAS on CDNA3, by default. On gfx940–gfx942, when the rocBLAS library is available and outside graph capture, prefill GEMMs of at least 4 rows (HIPFIRE_ROCBLAS_MIN_BATCH) dequantize the 4-bit weights into a cached FP16 shadow and call rocBLAS, which issues MFMA. Below 4 rows, launch overhead outweighs the compute win and the hand-written wave64 GEMV stays in charge.
• Hand-written MFMA kernels, opt-in. A set of gfx942 MFMA prefill kernels for HFQ4 weights reads the quantized weights directly, skipping the FP16 shadow (HIPFIRE_GFX942_MFMA_PREFILL=1…4, at least 16 rows). They are not a default.

// crates/rdna-compute/src/dispatch.rs (0.4.1)
fn rocblas_arch_eligible(&self) -> bool { /* … */ self.arch_caps.is_cdna3() }
fn rocblas_min_batch(&self) -> usize { /* HIPFIRE_ROCBLAS_OFF → usize::MAX */ self.flags.rocblas_min_batch.unwrap_or(4) }

19 Summary — what to remember

• CDNA wave width is 64. Every __shfl_*, occupancy calculation and MFMA builtin assumes 64 lanes; a SIMD is 16 lanes wide, so a wave64 instruction takes 4 cycles. The number flips to 32 on RDNA.
• MFMA is the CDNA matrix unit. One instruction does a full 16×16 (or 32×32) tile multiply-accumulate; rocBLAS and hipBLASLt route GEMMs to it.
• HBM is the DRAM tier: 192 GB at 5.3 TB/s on MI300X. MI300X and MI355X add a 256 MiB Infinity Cache; earlier CDNA parts have none.
• 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 counters that matter: VALUBusy, MemUnitBusy, L2CacheHit.
• The same HIP source compiles for either family, but a CDNA kernel ported to RDNA needs new matrix intrinsics and a wavefront-width audit.
• In hipfire 0.4.1, CDNA runs the portable wave64 path, with MFMA through rocBLAS for CDNA3 prefill.

20 Sources

• AMD — ROCm GPU hardware specifications: CU counts, XCDs, LDS, caches, VGPR file.
• AMD — Instinct MI300X and Instinct MI355X: HBM capacity and bandwidth, peak matrix throughput.
• hipfire v0.4.1 — crates/rdna-compute/src/dispatch.rs, gemm.rs, arch_caps.rs; the GPU support table in the README.
• hipfire v0.4.1 — kernels/src/*.gfx942.hip: the opt-in MFMA kernels.