hip-rocm

name	hip-rocm
description	HIP and ROCm skill for AMD GPU programming. Use when writing HIP kernels with hipcc, porting CUDA code via HIPIFY, profiling with rocprof, debugging with rocgdb, or optimizing for MI300X. Activates on queries about HIP, ROCm, hipify, hipcc, rocprof, or CUDA to AMD porting.

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 ported
• skills/gpu/cuda-profiling — Nsight concepts map to rocprof
• skills/gpu/gpu-memory-model — wavefront vs warp, coalescing rules
• skills/gpu/triton-lang — Triton supports AMD via ROCm backend
• skills/compilers/llvm — HIP uses Clang/LLVM toolchain