HIP

Owned by Cristian Di Pietrantonio
Last updated: May 01, 2024 by Emily Kahl
3 min read

HIP is a C++ Runtime API and Kernel Language that allows developers to create portable applications for both AMD and NVIDIA GPUs from a single source code.

Prerequisites

This page builds on GPU programming concepts introduced in CUDA.

An introduction to HIP

Heterogeneous-computing Interface for Portability (HIP) is an open-source C++ Runtime API and Kernel Language by AMD that allows developers to create portable applications for both AMD and NVIDIA GPUs from a single source code. HIP can use different backends depending on whether the code needs to be run on AMD or NVIDIA GPUs, namely the AMD open-source framework for GPU computing, Radeon Open Compute (ROCm), and the NVIDIA proprietary CUDA Toolkit, respectively.

HIP is available as a module on GPU-enabled Pawsey systems. It uses the ROCm backend on Setonix and the CUDA backend on Garrawarla.

Hipification of CUDA code

HIP provides two tools to help converting existing CUDA code to make use of HIP:

  • hipify-perl : a Perl script that makes use of regular expressions to convert CUDA syntax to HIP
  • hipify-clang : a clang-based translator that interprets the code and will identify errors as they occur.

As an example, let's use the "hello world" code from the CUDA page:

Listing 1. "Hello world" program in CUDA
#include <stdio.h>


__global__ void helloGPU() {

  int tid = threadIdx.x;
  int bid = blockIdx.x;
  printf("Hello from GPU thread %d in block %d\n",tid,bid);

}


int main(int argc, char *argv[]) {

  int no_blocks = 4;
  int no_threads = 5;

  helloGPU<<<no_blocks,no_threads>>>();

  cudaDeviceSynchronize();

}

The following command converts the CUDA source file, hello.cu, to a new HIP source file, hello.hip.cu:

$ hipify-perl hello.cu > hello.hip.cpp

When using a CUDA backend for compiling, note how the .cu extension has to be kept, to allow for the subsequent compilation with a CUDA backend. On the other hand, when using a ROCm backend for compiling, HIP source files typically have the .cpp extension.

The contents of hello.hip.cu are as follows:

Listing 2. "Hello world" program in HIP
#include "hip/hip_runtime.h"
#include <stdio.h>


__global__ void helloGPU() {

  int tid = threadIdx.x;
  int bid = blockIdx.x;
  printf("Hello from GPU thread %d in block %d\n",tid,bid);

}


int main(int argc, char *argv[]) {

  int no_blocks = 4;
  int no_threads = 5;

  hipLaunchKernelGGL(helloGPU, dim3(no_blocks), dim3(no_threads), 0, 0);

  hipDeviceSynchronize();

}

Note that a new HIP include statement has appeared, and that the CUDA kernel calls and API function calls have been replaced by HIP equivalents.

Compile and run HIP code

HIP code can be compiled by using the hipcc compiler provided by AMD ROCm:

$ hipcc --offload-arch=<architecture> hello.hip.cpp -o helloHIP

The code can be executed in an interactive SLURM session or within a batch job. An explicit request to use one or more GPUs is required. Terminal 1 shows an interactive session on Setonix:

Terminal 1. Run HIP code
$ salloc -N 1 --gres=gpu:1 -p gpu-dev -A <project>-gpu -t 0:10:00

$ module load rocm craype-accel-amd-gfx90a

$ hipify-perl hello.cu > hello.hip.cpp

$ hipcc --offload-arch=gfx90a hello.hip.cpp -o helloHIP

$ export OMP_NUM_THREADS=1

$ srun -N 1 -n 1 -c 8 --gres=gpu:1 --gpus-per-task=1 --gpu-bind=closest ./helloHIP
Hello from GPU thread 0 in block 0
Hello from GPU thread 1 in block 0
Hello from GPU thread 2 in block 0
Hello from GPU thread 3 in block 0
Hello from GPU thread 4 in block 0
Hello from GPU thread 0 in block 1
Hello from GPU thread 1 in block 1
Hello from GPU thread 2 in block 1
Hello from GPU thread 3 in block 1
Hello from GPU thread 4 in block 1
Hello from GPU thread 0 in block 3
Hello from GPU thread 1 in block 3
Hello from GPU thread 2 in block 3
Hello from GPU thread 3 in block 3
Hello from GPU thread 4 in block 3
Hello from GPU thread 0 in block 2
Hello from GPU thread 1 in block 2
Hello from GPU thread 2 in block 2
Hello from GPU thread 3 in block 2
Hello from GPU thread 4 in block 2

Related pages

  • For detailed information on how to compile HIP software on Pawsey systems, see Compiling.

External links