HIP / ROCm
Purpose
Guide agents through AMD GPU programming with HIP: the HIP runtime API, hipcc compilation, porting CUDA code with HIPIFY (hipify-perl, hipify-clang), ROCm toolchain setup, profiling with rocprof, debugging with rocgdb, HIP-vs-CUDA API mapping, and MI300X-specific optimizations.
When to Use
- Porting an existing CUDA codebase to AMD GPUs
- Setting up ROCm on Linux for MI200/MI300 hardware
- Writing native HIP kernels for AMD data center GPUs
- Profiling HIP applications with rocprof or rocprofiler-sdk
- Debugging device faults with rocgdb or compute sanitizers
- Building multi-vendor GPU code with HIP portability macros
Workflow
1. ROCm installation and verification
# Ubuntu/Debian (check ROCm docs for your distro version)
sudo apt install rocm-dev rocm-libs hip-dev
# Verify
rocminfo | head -30
hipconfig --version
hipcc --version
# List devices
rocm-smi
Set GPU target for compilation:
export AMDGPU_TARGETS=gfx942 # MI300X
export HIP_PLATFORM=amd
2. Minimal HIP kernel
// vector_add.hip
#include <hip/hip_runtime.h>
#include <stdio.h>
__global__ void vector_add(const float *a, const float *b, float *c, int n) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n)
c[i] = a[i] + b[i];
}
int main(void) {
const int n = 1 << 20;
size_t bytes = n * sizeof(float);
float *d_a, *d_b, *d_c;
hipMalloc(&d_a, bytes);
hipMalloc(&d_b, bytes);
hipMalloc(&d_c, bytes);
int threads = 256;
int blocks = (n + threads - 1) / threads;
hipLaunchKernelGGL(vector_add, dim3(blocks), dim3(threads), 0, 0,
d_a, d_b, d_c, n);
hipDeviceSynchronize();
hipFree(d_a); hipFree(d_b); hipFree(d_c);
return 0;
}
hipcc -O3 --offload-arch=gfx942 -o vector_add vector_add.hip
./vector_add
3. CUDA → HIP porting with HIPIFY
# Perl-based batch converter (quick port)
hipify-perl cuda_kernel.cu > cuda_kernel.hip
# Clang-based (more accurate, preserves structure)
hipify-clang cuda_project/ -o hip_project/ --cuda-path=/usr/local/cuda
# Convert single file in place
hipify-clang -inplace --cuda-path=/usr/local/cuda main.cu
Common API mappings:
| CUDA | HIP |
|---|---|
cudaMalloc |
hipMalloc |
cudaMemcpy |
hipMemcpy |
cudaMemcpyAsync |
hipMemcpyAsync |
cudaStream_t |
hipStream_t |
<<<grid, block>>> |
hipLaunchKernelGGL or <<<>>> (HIP supports CUDA syntax) |
__syncthreads() |
__syncthreads() (same) |
threadIdx / blockIdx |
Same builtins |
Portability header for dual compilation:
#ifdef __HIP_PLATFORM_AMD__
#include <hip/hip_runtime.h>
#else
#include <cuda_runtime.h>
#define hipMalloc cudaMalloc
#define hipMemcpy cudaMemcpy
// ... more macros
#endif
4. hipcc flags
# Target specific GPU architecture
hipcc --offload-arch=gfx942 -O3 -o app main.hip
# Multiple architectures
hipcc --offload-arch=gfx90a --offload-arch=gfx942 -o app main.hip
# Debug
hipcc -g -O0 --offload-arch=gfx942 -o app_debug main.hip
# Link with rocBLAS
hipcc -lrocblas -o app main.hip
5. rocprof profiling
# Basic kernel trace
rocprof --stats ./app
# CSV metrics output
rocprof -i input.csv -o output.csv ./app
# input.csv example:
# pmc: SQ_INSTS_VALU_ADD_F32,SQ_INSTS_VALU_MUL_F32,GRBM_COUNT
# ROCm 6.x rocprofiler-sdk (preferred for new projects)
rocprofv3 --kernel-trace -- ./app
Key metrics (AMD terminology):
- VALU utilization — compute unit activity
- LDS bank conflicts — shared memory (LDS) stalls
- Memory throughput — HBM bandwidth utilization
6. rocgdb debugging
# Build with debug symbols
hipcc -g -O0 --offload-arch=gfx942 -o app_debug main.hip
rocgdb ./app_debug
(rocgdb) break vector_add
(rocgdb) run
(rocgdb) info rocm kernels
(rocgdb) rocm thread 0 0 0
(rocgdb) print i
AMD also supports compute-sanitizer equivalents via ROCm's roc-obj-extract and memory checking tools where available.
7. MI300X optimizations
# Enable MFMA (matrix fused multiply-add) instructions
hipcc --offload-arch=gfx942 -munsafe-fp-atomics -O3 -o app main.hip
| Optimization | MI300X note |
|---|---|
| Matrix ops | Use rocBLAS/hipBLASLt for GEMM; MFMA intrinsics for custom |
| HBM bandwidth | ~5.3 TB/s peak (MI300X) — maximize memory coalescing to approach it |
| Wavefront size | 64 threads (vs CUDA warp 32) — adjust reduction patterns |
| LDS (shared mem) | 64 KB per CU; watch bank conflicts |
Wavefront-aware reduction:
__device__ float warp_reduce_sum(float val) {
// AMD wavefront = 64 lanes
for (int offset = 32; offset > 0; offset >>= 1)
val += __shfl_down(val, offset);
return val;
}
8. Library ecosystem
| NVIDIA | AMD ROCm |
|---|---|
| cuBLAS | rocBLAS / hipBLAS |
| cuDNN | MIOpen |
| NCCL | rccl |
| Thrust | hipCUB (portable) |
| cuFFT | rocFFT |
hipcc -lrocblas -o gemm_test gemm.hip
Common Problems
| Symptom | Cause | Fix |
|---|---|---|
hipErrorNoDevice |
ROCm driver not loaded | Check rocm-smi; add user to render group |
| Wrong architecture binary | Mismatched gfx* target |
rocminfo → set --offload-arch |
| hipify incomplete port | CUDA-specific APIs | Manual fix: cooperative groups, texture refs |
| Slower than CUDA reference | Wavefront 64 vs warp 32 | Tune block size to multiples of 64 |
HSA_STATUS_ERROR |
GPU busy or OOM | rocm-smi --showmeminfo; reduce allocation |
| rocprof empty output | No kernels launched | Verify hipGetLastError() after launch |
Related Skills
skills/gpu/cuda— source CUDA patterns being portedskills/gpu/cuda-profiling— Nsight concepts map to rocprofskills/gpu/gpu-memory-model— wavefront vs warp, coalescing rulesskills/gpu/triton-lang— Triton supports AMD via ROCm backendskills/compilers/llvm— HIP uses Clang/LLVM toolchain