a very gentle introduction to AMD GPU programming and transition from CUDA kernels https://x.com/i/article/2118910057626873856
Intro to AMD GPUs (Part 1) - From CUDA to CDNA
Most of the GPU posts I've written have been about the NVIDIA stack. I have spent a fair amount of time writing kernels on CUDA, DSLs like CUTEDSL, Triton and a tons of profiling work. But I've been meaning to explore the AMD side of the stack and this is the first post in that series.
For people who already know CUDA, this blog should be extremely easy to follow. For the people new to GPU programming, do follow along as I will still be breaking down the kernels and the architecture in follow up posts.
Glossary
The first distinct term you've probably already come across is CDNA and RDNA. CDNA is simply the hardware architecture family used by AMD's Instinct compute accelerators. This means they are designed for HPC and AI workloads. RDNA is the architecture family behind Radeon graphics GPUs and mainly aimed towards gamers. For this series, we are interested in the AI workloads. If you want to have a quick comparison on what they mean on NVIDIA terms, they are similar to the generations like Ampere, Hopper and Blackwell.
ROCm is the software platform basically the compilers, runtimes, libraries, and tools that let us use that hardware. HIP is a C++ programming interface within that platform, with a kernel language and runtime APIs that will look very familiar to a CUDA programmer. AMD's ROCm overview explains how these pieces fit together.
So if someone says “this GPU uses ROCm,” that tells us about the software environment. We still need to know the GPU architecture before deciding how to organize a kernel.
Let's get hands on
Starting with the good ol addition where each thread reads one element from A and B, adds them, and writes the result to C.
Oh wow, this looks almost exactly like CUDA, right? We have a grid, blocks, and threads. HIP keeps those names in the code. The AMD architecture documentation also calls a block a workgroup and a thread a work-item. The HIP language reference documents these built-ins and the launch syntax.
On NVIDIA, a block is assigned to an SM, or Streaming Multiprocessor. For our ordinary HIP launch on MI300X, the corresponding place to think about is a CU, or Compute Unit. A workgroup's threads execute on the same CU, which lets them cooperate using its local resources.
This is a useful correspondence, but the resource counts and execution pipelines inside a CU differ from those inside an SM. We cannot compare the number of CUs with the number of SMs and use that ratio to predict performance.
Our 256 threads become eight warps. Or four wavefronts.
On NVIDIA, the block is divided into groups of 32 threads called warps. On CDNA 3, the workgroup is divided into groups of 64 threads called wavefronts, often shortened to waves.
Let's take a look at that with our 256-thread block.
We still launched 256 threads on both GPUs. Having half as many waves does not mean we did half the work, or that each AMD wave takes twice as long.
A warp's width only describes its thread grouping while the instruction latency and throughput depend on the execution pipeline and the particular instruction.
Also, 64 is our CDNA assumption. RDNA supports different wave modes, so “AMD always means 64” is a bad rule to bake into portable code. Use warpSize in HIP device code and query the device properties on the host, as described in the HIP porting guide.
Our shared memory tiles now live in LDS
In the tiled matmul post, we loaded tiles of A and B into shared memory so that threads could reuse them. On AMD, the on chip memory for that is called LDS, or Local Data Share. HIP still uses __shared__ to declare it and __syncthreads() for a block barrier.
Let's do a small calculation. Suppose a matmul block computes a 128 x 128 output tile, stepping through K in chunks of 32, with FP16 inputs. Its A tile is 128 x 32 and its B tile is 32 x 128.
A tile = 128 x 32 x 2 bytes = 8 KiB
B tile = 32 x 128 x 2 bytes = 8 KiB
One stage = 16 KiB
Two buffered stages = 32 KiB
On MI300X, those two stages already use half the CU's LDS capacity, before padding or other shared storage. LDS capacity alone would allow at most two such workgroups to reside on a CU. Registers and other limits may reduce that further. A tile shape that worked well on H100 might not work here and you might need to think of a different strategy.
What replaces Tensor Cores?
For matrix multiplication, CDNA has dedicated matrix cores. We access them through MFMA, or Matrix Fused Multiply-Add, instructions that compute small matrix products and add them to accumulators. This serves the same broad purpose as the Tensor Core operations we used on NVIDIA.
The key detail is the distribution of matrix elements across lanes. An MFMA is a cooperative wave operation. Each lane supplies part of the operands and owns part of the result. AMD's CDNA matrix-core guide works through those layouts for specific instructions.
This is pretty similar to the WMMA post where one thread did not independently own the whole matrix tile. We had to get the fragments into the form expected by the instruction.
So an optimized port from CUDA needs to choose the AMD matrix instruction first, then work out how to load and distribute its operands. Copying a CUDA fragment layout and changing the instruction name is not enough. We will work through an actual MFMA tile later in the series.
Nvidia vs AMD existing stack
If we call a matrix multiplication from PyTorch, the framework can dispatch to a library implementation. If we write our own HIP kernel, we choose the grid and block configuration ourselves. Either way, the compiled work eventually runs on the CUs we have been discussing.
Can we convert our existing CUDA code?
We wrote the vector-add kernel in HIP earlier, but suppose we already have a working CUDA program. Do we have to go through the entire file changing every allocation, copy, and stream call ourselves?
AMD provides HIPIFY to automate supported CUDA-to-HIP source conversions. It handles familiar changes such as replacing the runtime header and translating CUDA API names to their HIP counterparts. Here are a few examples from the API mapping reference:
For example, this host-side copy changes as follows. Assume the source and destination buffers are already allocated:
Our kernel's threadIdx.x, blockIdx.x, and __global__ syntax are already familiar to HIP. Simple code can therefore need relatively few changes. HIPIFY also knows supported library mappings, but the coverage depends on the API and tool version.
hipify-perl and hipify-clang
There are two tools worth knowing about. hipify-perl uses pattern matching and is easy to try on a small source file. It does not require a CUDA installation to perform the translation, but complex macros and C++ constructs can need manual attention.
hipify-clang parses the CUDA source into an abstract syntax tree, which gives it more context for the conversion. It needs a compatible CUDA installation and the headers and compiler definitions required to parse the project. For a larger codebase, supplying the project's compilation database can help provide that context.
Given a complete CUDA program in vector_add.cu, with HIPIFY and the ROCm development tools installed, a small MI300X port could start like this:
Notice that translation and compilation are separate steps. HIPIFY produces HIP source. Then hipcc invokes the toolchain to compile the host and device code for our selected target. We are recompiling the program for AMD hardware, an existing NVIDIA executable does not become an AMD executable through API renaming. The HIP compiler guide describes this compilation path.
What next
I want to explore the MI300X architecture in more detail, so the next post will be a walkthrough of the hardware design in detail. Until then, ciao!