ROCm and the AMD GPU Training Ecosystem: From Porting to Tuning

Note: This is a review article (no first-hand experimental data), based on public docs and migration practice.

Experiment Objective

AMD GPUs hold a dominant share of the Top500 supercomputers (Frontier topped the list for years), and the training ecosystem (ROCm) is maturing fast. But "can CUDA code run directly on AMD cards?" is still the most asked question. Goal: dissect the ROCm software stack, the real cost of HIP porting (which code auto-converts, which needs hand work), the right way to use PyTorch ROCm builds, and the CUDA-equivalent tuning map. After reading you can answer: to move this CUDA project to AMD, how much must change, what performance gap to expect, and what to tune first.

Code Walkthrough

Part 1: HIP porting — "semi-automatic translation" of CUDA code

// Original CUDA (example)
__global__ void saxpy(float *y, float a, float *x, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) y[i] = a * x[i] + y[i];
}

// HIP version: almost line-for-line identical (hipify auto-converts)
__global__ void saxpy(float *y, float a, float *x, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) y[i] = a * x[i] + y[i];
}

// Auto conversion: hipify-perl / hipify-clang
hipify-perl saxpy.cu > saxpy_hip.cpp

Key points:

1. Core CUDA syntax maps 1:1 in HIP: __global__/blockIdx/threadIdx/__syncthreads are identically named — HIP is designed as a "superset of the CUDA dialect," so simple kernels migrate at near-zero cost.
2. APIs have a mapping table: cudaMalloc→hipMalloc, cudaMemcpy→hipMemcpy, cudaStreamCreate→hipStreamCreate; hipify replaces them by table — pure CUDA-API + kernel code migrates in hours.
3. The real hand work is library dependencies: cuBLAS→rocBLAS, cuDNN→MIOpen, NCCL→RCCL, TensorRT→the ROCm inference stack — in deep-learning projects these library calls are the migration bulk, not kernel syntax.

Part 2: The right way to use PyTorch ROCm

# Option 1 (recommended): official prebuilt wheel
pip install torch --index-url https://download.pytorch.org/whl/rocm6.2

# Option 2: build from source (must match driver/ROCm version exactly)
export ROCM_PATH=/opt/rocm
export HCC_AMDGPU_TARGET=gfx90a        # MI250X; gfx942 for MI300X

Key points:

• Prebuilt wheel first: PyTorch publishes official ROCm wheels (rocm5.x/6.x), used exactly like the CUDA build (torch.cuda maps to the ROCm backend) — for pure PyTorch projects (no custom CUDA extensions), moving to AMD means changing the install source.
• HCC_AMDGPU_TARGET decides the kernel target: gfx90a (MI200 family) / gfx942 (MI300 family) cannot be mixed; the wrong target yields "installs but won't run" binaries.
• Custom CUDA extensions are the watershed: .cu extensions must be hipify'd and rebuilt, and CUDA libraries (cuDNN/cuBLAS) replaced with ROCm equivalents — this step decides whether migration takes hours or weeks.

Principles

ROCm software stack layers

PyTorch/TensorFlow (ROCm build)
  → HIP runtime (HIP is ROCm's "CUDA")
  → libraries: rocBLAS(matmul) / MIOpen(conv) / RCCL(collectives) / rocSOLVER
  → compiler: HIP-Clang (LLVM-based)
  → driver: amdgpu + ROCm kernel modules
  → AMD GPU (gfx architectures)

Layer-for-layer CUDA correspondence: CUDA runtime→HIP runtime, cuBLAS→rocBLAS, cuDNN→MIOpen, NCCL→RCCL, nvcc→hipcc (HIP-Clang). Once you understand this map, "which library to substitute" is obvious during migration.

Where the performance gap really is (honest CUDA comparison)

Area	Status (estimated)	Notes
standard ops (GEMM/conv)	90%-105%	rocBLAS/MIOpen aligned; some ops lag
fused/custom kernels	70%-95%	community kernel count/quality still catching up
collectives (large scale)	85%-100%	RCCL maturing fast; large clusters still need validation
ecosystem tools (profiler/debug)	60%-90%	rocprof/Omnitrace usable but thinner than Nsight

Key insight: AMD's per-op performance is already good; the gap is ecosystem depth (op coverage, tuning tools, library release cadence), not hardware peak — so "how fast it runs" depends more on how deeply the project uses CUDA-specific features.

Tuning correspondence

CUDA's tuning trio all have ROCm counterparts: rocprof (≈ nsys/ncu: kernel timing and hardware counters), environment variables like HSA_XNACK (≈ CUDA env tuning), and MIOpen's tuner (auto-selecting conv implementations). Teams migrating from CUDA should first run rocprof to confirm op coverage and bottlenecks, then apply ch08's Nsight methodology — the methodology is universal, the tool names change.

Results

(Estimated, not measured; typical migration cases)

Project type	Migration effort	End-to-end perf (vs same-tier CUDA)
pure PyTorch (official ROCm wheel)	hours	90%-105%
PyTorch + a few custom kernels	days	85%-100%
deep CUDA-library-dependent project	weeks	70%-95%

Engineering Takeaways

1. Assess "library-dependency depth" before promising a schedule: pure API/kernel code auto-converts; library dependencies (cuDNN/NCCL etc.) are the workload — list dependencies before migrating.
2. Official wheels first, source builds second: wheels match driver/ROCm versions, avoiding compile pits; only build from source for custom kernels/optimization.
3. Verify performance per-op: beyond end-to-end comparison, use rocprof for op-level comparison (echoing ch08's profiling method) — knowing whether the gap is "library coverage" or "kernel quality" determines the fix.
4. Coexist with heterogeneous scheduling: ch07 covers mixed deployment; on AMD cards run ROCm jobs, on NVIDIA cards run CUDA jobs — the scheduler partitions by architecture (SLURM partitions/GRES).

Next article preview: with the training-side AMD ecosystem covered, back to inference operations — "Inference Capacity Planning: From QPS to GPU Instance Count."