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 Topaz.
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 HIPhipify-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:
#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.cu
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:
#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 hello.hip.cu -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 Topaz:
$ salloc -N 1 -n 1 --gres=gpu:1 -p gpuq-dev -t 0:01:00 $ srun -n 1 --gres=gpu:1 --export=all ./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
- For ROCm/HIP tutorials, see the external ROCm learning centre
- For links to the HIP Programming Guide and other documentation around HIP, see the HIP homepage
- For documentation around the ROCm platform, see the external ROCm homepage
- AMD shares ROCm/HIP source codes on GitHub: