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.

GitHub
安装命令
npx skhub add mohitmishra786/hip-rocm
Markdown
SKILL.md

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:

CUDAHIP
cudaMallochipMalloc
cudaMemcpyhipMemcpy
cudaMemcpyAsynchipMemcpyAsync
cudaStream_thipStream_t
<<<grid, block>>>hipLaunchKernelGGL or <<<>>> (HIP supports CUDA syntax)
__syncthreads()__syncthreads() (same)
threadIdx / blockIdxSame 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
OptimizationMI300X note
Matrix opsUse rocBLAS/hipBLASLt for GEMM; MFMA intrinsics for custom
HBM bandwidth~5.3 TB/s peak (MI300X) — maximize memory coalescing to approach it
Wavefront size64 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

NVIDIAAMD ROCm
cuBLASrocBLAS / hipBLAS
cuDNNMIOpen
NCCLrccl
ThrusthipCUB (portable)
cuFFTrocFFT
hipcc -lrocblas -o gemm_test gemm.hip

Common Problems

SymptomCauseFix
hipErrorNoDeviceROCm driver not loadedCheck rocm-smi; add user to render group
Wrong architecture binaryMismatched gfx* targetrocminfo → set --offload-arch
hipify incomplete portCUDA-specific APIsManual fix: cooperative groups, texture refs
Slower than CUDA referenceWavefront 64 vs warp 32Tune block size to multiples of 64
HSA_STATUS_ERRORGPU busy or OOMrocm-smi --showmeminfo; reduce allocation
rocprof empty outputNo kernels launchedVerify 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
发现
标签

此技能尚未发布标签。

版本
最新版本元数据

版本

v2026.09.24

发布时间

2026年9月24日

分类

未分类

许可证

MIT

源路径

skills/gpu/hip-rocm

默认分支

main

最新提交

bdc5847

Tree SHA

1178323