rocm

ROCm GPU Development

Version: rocm-7.2.4 Language: C/C++ with HIP (.hip or .cpp files) License: MIT (examples and most libraries); some components Apache-2.0 Docs: https://rocm.docs.amd.com

Overview

ROCm is AMD's open-source GPU computing stack. HIP (Heterogeneous-compute Interface for Portability) is the primary programming API — it is syntactically close to CUDA and compiles on both AMD and NVIDIA GPUs. ROCm includes optimized compute libraries (BLAS, FFT, RNG, solvers, sparse, tensor contractions, structured sparsity), HIP Graphs for low-overhead kernel replay, and profiling/debugging tools. As of ROCm 7.2, hipcc is deprecated in favor of amdclang++ (same flags, direct invocation recommended).

Quick Start

// minimal.hip
#include <hip/hip_runtime.h>
#include <stdio.h>

__global__ void hello(int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) printf("thread %d\n", i);
}

int main() {
    hello<<<dim3(4), dim3(8)>>>(32);    // 4 blocks × 8 threads
    hipDeviceSynchronize();
}

# Compile and run (hipcc still works but is deprecated in 7.2+)
hipcc -O2 -o hello minimal.hip && ./hello
# Preferred: amdclang++ -O2 -x hip -o hello minimal.hip && ./hello

# CMake (preferred for larger projects)
# cmake -S . -B build && cmake --build build

Error-check macro (use throughout your code):

#define HIP_CHECK(expr) do { \
    hipError_t err = (expr); \
    if (err != hipSuccess) { \
        fprintf(stderr, "HIP error: %s at %s:%d\n", \
                hipGetErrorString(err), __FILE__, __LINE__); \
        exit(EXIT_FAILURE); \
    } \
} while(0)

// Usage:
HIP_CHECK(hipMalloc(&d_ptr, bytes));
HIP_CHECK(hipMemcpy(d_ptr, h_ptr, bytes, hipMemcpyHostToDevice));
my_kernel<<<grid, block>>>(args);
HIP_CHECK(hipGetLastError());       // catch launch errors
HIP_CHECK(hipDeviceSynchronize());  // catch runtime errors

Core HIP Concepts

Device Management

int n_devices;
hipGetDeviceCount(&n_devices);
hipSetDevice(0);                          // select GPU

hipDeviceProp_t props;
hipGetDeviceProperties(&props, 0);
// props.name, .totalGlobalMem, .warpSize (64 on AMD!),
// .maxThreadsPerBlock, .gcnArchName