AMD ROCm Software Stack, in depth: the layers from amdgpu and KFD to HIP and libraries, code objects, containers and triage

The stack as layers

Each layer only talks to its neighbours. Frameworks call libraries and the HIP runtime. HIP is built on ROCr, AMD's implementation of the HSA runtime, which manages devices, queues and memory. ROCr talks to the kernel through KFD, the compute interface of the amdgpu driver. Knowing this ordering matters because the layers fail differently: a missing device node fails everything, a wrong code object fails one kernel, and a missing tuning database only makes things slow.

Kernel driver: amdgpu, KFD and device nodes

At the bottom is amdgpu, the Linux kernel driver, which also handles display GPUs. Compute uses its KFD (kernel fusion driver) side. Two device files matter: /dev/kfd, a single node that every compute process opens, and one /dev/dri/renderD* node per GPU or partition. User space needs read-write access to both, which on most distributions means membership of the video and render groups.

The kernel driver owns things software above it cannot fix: firmware loading, memory reset, RAS error reporting and partitioning. MI300-class GPUs can be split into compute partitions and memory partitions, and each partition appears as its own render node and its own device to everything above. A node that reports 64 GPUs instead of 8 is usually a partition mode, not a bug; the MI300X deep dive covers the modes. When the driver is unhealthy, dmesg shows amdgpu errors and the upper layers report no devices at all.

ROCr: the HSA runtime

ROCr enumerates agents (CPUs and GPUs), creates user-mode queues in the AQL packet format, allocates memory from pools, and implements signals for completion. When HIP launches a kernel it writes a dispatch packet into a queue and rings a doorbell; the GPU's command processor reads it without a system call. That is why launch overhead is low, and also why a hung kernel does not show up as a hung syscall.

rocminfo is the ROCr-level test. It lists every agent, and for each GPU the name to look for is the ISA target, such as gfx942. If rocminfo sees no GPU agents, nothing above will either. Device selection also starts here: ROCR_VISIBLE_DEVICES filters at the ROCr level, HIP_VISIBLE_DEVICES at the HIP level, and HIP also honours CUDA_VISIBLE_DEVICES for compatibility. Setting more than one is a classic way to end up with zero visible devices, because the second filter indexes into the already-filtered list.

HIP runtime, compiler and code objects

The HIP runtime (built in the CLR repository) provides the CUDA-like API: streams, events, memory copies, module loading. The compiler is AMD's LLVM, invoked as amdclang++ or through the hipcc wrapper. Device code is compiled ahead of time into a code object for each ISA target and bundled into the host binary. On this default path there is nothing like CUDA's PTX fallback that the driver compiles for a new GPU at load time: if your binary has no code object for the GPU's target, the kernel cannot run. Recent releases add a target-agnostic SPIR-V target, amdgcnspirv, that is compiled to native code at run time; check that your release, libraries and extensions support it before relying on it.

# Build for the GPUs you actually deploy on; one code object per target.
hipcc -O3 --offload-arch=gfx90a --offload-arch=gfx942 -o saxpy saxpy.hip

# Inspect which targets a binary or shared library carries.
roc-obj-ls ./saxpy
#   expect one entry per --offload-arch, e.g. ...amdgcn-amd-amdhsa--gfx942

Target strings can carry feature flags. gfx90a:xnack+ and gfx90a:xnack- are different code objects: xnack enables retrying page faults, which unified memory needs, and a binary built for one mode will not load in the other. The rule for production is simple: list every target in your fleet explicitly, check binaries in CI (roc-obj-ls ships with hipcc, and recent LLVM builds of llvm-objdump can also list HIP offload bundles), and treat a new GPU generation as a rebuild of every extension that carries device code, including custom PyTorch ops. The common targets are gfx90a for MI200, gfx942 for MI300X, MI300A and MI325X, and gfx950 for the MI350 series.

Libraries and how frameworks reach them

Frameworks rarely call HIP for heavy work; they call libraries. GEMMs go to hipBLASLt (and rocBLAS for some paths), convolutions and normalisations to MIOpen, fused kernels to Composable Kernel, and collectives to RCCL, which keeps NCCL's API and honours most NCCL_* variables. PyTorch's ROCm build keeps the torch.cuda namespace, so torch.cuda.is_available() returns True on AMD; torch.version.hip is how you tell which build you have.

Two library behaviours surprise operators. MIOpen searches for the fastest kernel the first time it sees a shape and compiles it, keeping compiled kernels under ~/.cache/miopen and tuning results under ~/.config/miopen; in a fresh container every job pays that cost again unless both directories are on a persistent volume. PyTorch's TunableOp, enabled with PYTORCH_TUNABLEOP_ENABLED=1, benchmarks GEMM implementations per shape and writes the winners to a CSV that you can ship with the job. Both turn first-run slowness into a deployment artefact you control.

Tools for each question

# Kernel trace plus a per-kernel summary for one training step.
rocprofv3 --kernel-trace --stats -d ./prof -- python train.py --steps 20

Releases and version pinning

ROCm ships as a versioned set of components that are tested together. As of this writing, AMD's documentation lists ROCm Core SDK 10.1.0 (released 2026-10-05), and says that since ROCm 7.14 the stack is built and released through TheRock, with a transition guide for the older release stream. Version numbers move quickly, so check the compatibility matrix for your GPU and distribution rather than trusting any article, including this one.

The practical rule is that the user-space stack (ROCr, HIP, libraries) lives in your container image and moves with your application, while the kernel driver lives on the host and moves with your fleet. AMD documents which driver and user-space combinations are supported; pin both, record them in job metadata, and upgrade the host driver before the containers that need it.

Containers and Kubernetes

A container needs the device nodes, the groups and enough shared memory for data loaders and RCCL:

docker run -it --rm \
  --device=/dev/kfd --device=/dev/dri \
  --group-add video --group-add "$(getent group render | cut -d: -f3)" \
  --ipc=host --shm-size=16g \
  -v $HOME/.cache/miopen:/root/.cache/miopen \
  -v $HOME/.config/miopen:/root/.config/miopen \
  rocm/pytorch:latest \
  python -c "import torch; print(torch.version.hip, torch.cuda.device_count(),
             torch.cuda.get_device_properties(0).gcnArchName)"

Passing all of /dev/dri exposes every GPU; to give a container one GPU, pass /dev/kfd plus that GPU's render node only. Profilers may additionally need a relaxed seccomp profile. The render group is added by numeric ID because many images have no group of that name, and docker refuses an unknown group name. In Kubernetes, the AMD GPU device plugin advertises amd.com/gpu as a resource and mounts the right nodes, and the AMD GPU Operator manages the driver, plugin and metrics exporter together. Pods then request amd.com/gpu: 8 exactly as they would request NVIDIA GPUs. See GPUs in containers for the general pattern.

Worked example: bring up a node layer by layer

Here is a bring-up check that tests each layer in order and stops at the first failure. Run it on every new node image and inside every new container image:

#!/usr/bin/env bash
set -euo pipefail
echo "1. kernel driver";  lsmod | grep -q '^amdgpu' && ls -l /dev/kfd /dev/dri/renderD*
echo "2. permissions";    test -r /dev/kfd -a -w /dev/kfd
echo "3. ROCr agents";    rocminfo | grep -E '^\s+Name:\s+gfx' | sort | uniq -c
echo "4. health";         amd-smi list
echo "5. HIP + torch";    python - <<'EOF'
import torch
assert torch.version.hip, "not a ROCm build of PyTorch"
n = torch.cuda.device_count(); assert n > 0, "HIP sees no devices"
a = torch.randn(4096, 4096, device="cuda", dtype=torch.bfloat16)
torch.cuda.synchronize(); print(n, "GPUs", torch.cuda.get_device_properties(0).gcnArchName,
      float((a @ a).float().abs().mean()))
EOF
echo "6. collectives";    echo "run rccl-tests all_reduce_perf across all local GPUs next"

On an eight-GPU MI300X node, step 3 should print a count of 8 next to gfx942 (with the default partition mode), and step 5 should report 8 GPUs and a finite mean. If step 3 passes and step 5 fails, the problem is in the container's user-space stack or environment variables, not the host. For multi-node jobs, finish with collective benchmarks using rccl-tests before scheduling real work.

Failure modes by layer

Trade-offs

The ROCm stack is open from the driver up, which means you can read, patch and rebuild every layer, and its device count, partitioning and memory capacity are strong for large models. The cost is that it has fewer forgiving defaults than CUDA: native code objects per GPU target by default, more first-run tuning, and faster-moving versions. Teams that do well on it treat the stack as part of their build: pinned images, explicit targets, persisted tuning caches, and a layer-by-layer health check in every node and container pipeline. For per-GPU details of the newest parts, see the MI325X deep dive.

What to do next

1. Run the bring-up script on one node and one container image; record the gfx target, partition mode, driver and ROCm versions.
2. Add roc-obj-ls checks to CI for every binary and Python extension with device code, failing the build if a fleet target is missing.
3. Pick one device-selection variable for your scheduler and remove the others from job templates.
4. Mount a persistent MIOpen cache and, for GEMM-heavy jobs, try TunableOp and ship the results file.
5. Profile one training step with rocprofv3 --kernel-trace --stats and list the top five kernels by time.
6. Run rccl-tests across nodes before the first multi-node job, and keep the numbers as a baseline.