AMD ROCm/HIP + Apple MLX — awareness-level

Scope: you are not going to be an AMD or Apple kernel specialist by end of month 13. You are going to be able to have credible technical conversations about non-NVIDIA stacks, port a Triton kernel to ROCm, and understand why Meta serves Llama 3.1 405B on MI300X. That’s this document’s target.

NVIDIA is not the only game in town, and pretending otherwise is a career risk. Skip nothing here, but time-box: 1 week total for both stacks.

---

AMD ROCm / HIP

The one-paragraph summary

AMD’s data-center GPUs are CDNA architecture (MI300X = CDNA3, MI325X/MI355X = CDNA4 family). Consumer GPUs are RDNA (RX 9070 = RDNA4). The software stack is ROCm, an open-source CUDA-alike. HIP is the C++ language layer — nearly 1-to-1 with CUDA, with a hipify tool to convert .cu → .hip. rocBLAS is the BLAS. Composable Kernel (CK) is the CUTLASS analogue. AITER is AMD’s inference kernel library.

Why it matters (concrete)

• MI300X: 192 GB HBM3 @ 5.3 TB/s, 750W. Highest single-GPU memory in the industry. This is why AMD wins on memory-bound decode of large models. Meta reportedly runs 100% of Llama 3.1 405B live traffic on MI300X (per vLLM blog).

• vLLM on MI300X vs TGI on same HW: 1.5× throughput, 1.7× TTFT for Llama 3.1 405B (vLLM benchmark, https://vllm.ai/blog/2024-10-23-vllm-serving-amd).

• The software gap is closing but real: SemiAnalysis’s MI300X vs H100/H200 training benchmark showed the HW advantage but exposed ROCm’s training-side maturity gap. On inference the story is much better.

Current state (July 2026)

• ROCm 7.2 (March 2026): first release with RDNA 4 (RX 9070) support, single Windows+Linux installer, Ollama/LM Studio/llama.cpp/vLLM parity. https://localaimaster.com/blog/amd-rocm-local-llm-setup

• ROCm 7.0: MI300X/Llama3/DeepSeek benchmarks on llama.cpp landed. AMD publishing their own perf blog series: https://rocm.blogs.amd.com/ecosystems-and-partners/llama-cpp-oct2025/README.html

• Triton on AMD: Triton 3.7 added warp specialization on AMD, RDNA4 code paths, TDM. Most Triton kernels port with zero changes.

• HipKittens (Nov 9 2025): Stanford Hazy Research released AMD-optimized tile primitives — https://hazyresearch.stanford.edu/blog/2025-11-09-amd-brr — significant because it’s the first serious cross-vendor kernel library from that team. Uses 64-thread waves (vs NVIDIA 32), 8-wave / 4-wave patterns instead of NVIDIA wave-specialization, chiplet-aware cache reuse.

What to actually learn (5-day plan)

Day 1 — HIP porting reflex. Take one of your CUDA kernels from PMPP (transpose, reduction). Run hipify-perl kernel.cu > kernel.hip. Compile with hipcc. On a rented MI300X (Modal/Runpod/Hot Aisle), run it. Note what breaks (warp size 64 vs 32 is the big one).

Day 2 — Wavefront model. On AMD:

• Wavefront = 64 threads (CDNA) or 32/64 (RDNA4 configurable). Not 32.

• LDS = SMEM equivalent.

• Matrix cores = tensor cores equivalent, mfma_* instructions.

• → read the AMD Matrix Instruction Calculator: https://github.com/ROCm/amd_matrix_instruction_calculator to know which mfma variants exist and their throughputs.

Day 3 — Triton on AMD. Take one of your Triton kernels (fused softmax or RMSNorm). Set TRITON_PRINT_AUTOTUNING=1 and run on MI300X. Observe the autotune sweeps differ. Read: https://github.com/ROCm/triton — the ROCm-forked Triton docs.

Day 4 — vLLM on AMD. Deploy Llama 3.1 8B via vLLM on an MI300X instance (Modal has them). Compare tokens/sec vs H100 rental. Read the vLLM AMD blog end to end: https://vllm.ai/blog/2024-10-23-vllm-serving-amd.

Day 5 — HipKittens read. Read the HazyResearch HipKittens blog fully. Understand why 8-wave patterns replace NVIDIA-style wave-specialization on CDNA. This is the current AMD kernel-authoring frontier.

Vocabulary crash course

CUDA	HIP / ROCm
SM	CU (Compute Unit)
warp (32)	wavefront (64 on CDNA)
SMEM	LDS (Local Data Share)
Tensor Core	Matrix Core
mma.sync / wgmma	mfma_* (Matrix FMA)
cuBLAS	rocBLAS
cuDNN	MIOpen
CUTLASS	Composable Kernel (CK)
Nsight Compute	rocprof / omniperf
CUDA graphs	HIP graphs

Profiling on AMD

• rocprof — command-line profiler. Analog of ncu but coarser.

• omniperf (aka rocprofiler-compute) — higher-level UI, roofline analysis, closer to ncu experience.

• omnitrace — timeline analog of nsys.

AMD references

• ROCm docs: https://rocm.docs.amd.com

• HIP porting guide: https://rocm.docs.amd.com/projects/HIP/en/latest/how-to/hip_porting_guide.html

• vLLM x AMD blog: https://vllm.ai/blog/2024-10-23-vllm-serving-amd

• AMD ROCm blogs (llama.cpp benchmarks): https://rocm.blogs.amd.com

• HipKittens: https://hazyresearch.stanford.edu/blog/2025-11-09-amd-brr

• ROCm 7 setup guide: https://localaimaster.com/blog/amd-rocm-local-llm-setup

• State of PyTorch HW acceleration 2025 (CUDA vs ROCm vs XLA vs MPS): https://tunguz.github.io/PyTorch_Hardware_2025

---

Apple MLX

The one-paragraph summary

Apple Silicon uses a unified memory architecture — CPU and GPU share the same RAM at the same bandwidth. MLX is Apple’s array library, NumPy-like Python + C++ API, with a lazy execution graph and JIT-compiled Metal kernels. Metal is Apple’s low-level GPU API (equivalent to CUDA C++). Apple GPUs use SIMD groups of 32 threads (like NVIDIA warps).

Why it matters (concrete)

• On-device inference on iPhone/Mac is a real product surface. Nobody serves production models on Mac Studios, but many people prototype, dev, and ship consumer apps on Apple hardware.

• Unified memory changes what’s cheap: no hostToDevice copy, but bandwidth is shared with CPU workloads.

• Ollama, llama.cpp, MLX-LM all target Apple Silicon. If a job or study involves consumer AI apps, this matters.

Current state (July 2026)

• MLX active development: https://github.com/ml-explore/mlx — Apple + open-source contributors, regular releases.

• MLX-LM: https://github.com/ml-explore/mlx-examples — reference LLM implementations (Llama, Mistral, Qwen, Phi).

• Apple Silicon lineup: M4 Max (128 GB unified), M3 Ultra (192 GB, 819 GB/s BW).

• Metal 3 + MPS Graph: the underlying stack, unified in torch.mps for PyTorch users.

What to actually learn (2-day plan)

Day 1 — MLX Hello World. On a Mac (M-series), pip install mlx mlx-lm. Load Llama 3.2 3B, generate a few tokens. Time it. Understand mx.eval() (lazy execution barrier).

Day 2 — The mental model. Read the MLX docs’ “How MLX Works”: lazy graphs, unified memory (no .to(device)), function transformations (grad, vmap, compile). Contrast with PyTorch.

Optional: write a Metal Performance Shaders (MPS) kernel via mlx.core.metal. Non-critical.

Vocabulary

CUDA	Metal / MLX
Warp (32)	SIMD group (32)
Block	Threadgroup
SMEM	Threadgroup memory
Grid	Grid (same)
torch.tensor on device	mx.array (unified memory)

Apple references

• MLX repo: https://github.com/ml-explore/mlx

• MLX docs: https://ml-explore.github.io/mlx/build/html/index.html

• MLX examples (LLMs): https://github.com/ml-explore/mlx-examples

• Metal shading language: https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf

• PyTorch MPS backend: https://pytorch.org/docs/stable/notes/mps.html

---

Time-box discipline

• AMD: 5 days. Do the plan. Don’t get sucked in.

• Apple: 2 days. Enough to know how MLX feels; move on.

• Do not attempt to become an AMD kernel author in month 3. The path is: NVIDIA fluency first (Triton + CUTLASS), then port. HipKittens exists specifically because “port fluent NVIDIA thinking to AMD” is the winning move for that team.