Introduction to HIP Programming

Marcus DuPont, author of SIMBI
Spitzer & FFPS Fellow, Princeton University

Overview

What is HIP?

Basic Idea

Heterogeneous-compute Interface for Portability (HIP) is a C++ runtime API and kernel language that allows developers to create portable applications for AMD and NVIDIA GPUs from a single codebase.

It is part of a larger initiative to ease the transition to GPU computing via a large software distro called the Radeon Open Compute Platform, or ROCm.

HIP:

• is open-source
• is syntactically similar to CUDA (can even find-replace cuda functions with hip, and most simple programs will still work.)
• supports a good chunk of the CUDA runtime out of the box (e.g., cudaMalloc, cudaMemset, cudaMemcpy)

Comparison Table

CUDA vs HIP Runtime Calls

Warps on HIP (AMD) Devices

On AMD devices, the concept similar to CUDA's "warp" is called a "wavefront." A wavefront consists of 64 threads, whereas a warp in CUDA consists of 32 threads. This difference in the number of threads per execution unit is one of the key distinctions between AMD and NVIDIA architectures.

Here are some key differences between warps in CUDA and wavefronts in HIP:

• Number of Threads: A warp in CUDA consists of 32 threads, while a wavefront in HIP consists of 64 threads.
• Execution Model: Both warps and wavefronts execute in a SIMD (Single Instruction, Multiple Data) fashion, meaning all threads within a warp or wavefront execute the same instruction simultaneously.
• Divergence Handling: In both CUDA and HIP, if threads within a warp or wavefront diverge (i.e., follow different execution paths), the execution is serialized, which can lead to performance penalties. However, the larger size of wavefronts in HIP can make divergence more costly compared to CUDA.
• Synchronization: Synchronization within a warp or wavefront is implicit, meaning that all threads within a warp or wavefront are always synchronized. Explicit synchronization is required only when synchronizing threads across different warps or wavefronts.

Understanding these differences is crucial for optimizing performance on AMD and NVIDIA GPUs, as the size of warps and wavefronts can impact how you structure your parallel algorithms and manage thread divergence.

Installation Guide

To install HIP, you need to install the ROCm software stack, which includes the HIP runtime and compiler. The ROCm software stack is available for Linux operating systems and supports a wide range of AMD GPUs.

To install ROCm, follow these steps:

1. Download the ROCm software stack from the official ROCm website.
2. Follow the installation instructions provided on the ROCm website for your specific Linux distribution.
3. Once ROCm is installed, you can start using HIP to develop GPU-accelerated applications.

Once HIP is installed, check that it works with the following command

hipcc --version

If that works, then compiling goes like:

hipcc -o hello_world hello_world.cpp

Checking the platform

To check the platform, you can use the following code snippet:

hipconfig --platform

This will output the platform information, including the platform name, version, and other details. For example, it might output:

Platform: ROCm
Version: 4.0.0
Device: AMD Radeon RX 5700 XT

If the platform is autodetected incorrectly, it can be fixed by setting the correct environment variables. For example, if needing an Nvidia environment:

export HIP_PLATFORM=nvidia; export HIP_RUNTIME=cuda; export HIP_COMPILER=nvcc

Portability

HIP allows developers to write code that can run on both AMD and NVIDIA GPUs. This is achieved by providing a set of tools that allow developers to write code in a single-source manner, and then compile it for either platform.

For example, the following code snippet shows how to write a simple vector addition kernel in HIP:

						__global__ void vector_add(int *a, int *b, int *c, int n) {
							int i = blockIdx.x * blockDim.x + threadIdx.x;
							if (i < n) {
								c[i] = a[i] + b[i];
							}
						}
					

It's the same as CUDA (!)

HIP vs CUDA

HIP is designed to be very similar to CUDA, making it easier for developers to port CUDA applications to run on AMD hardware. Here are some key comparisons:

• HIP provides a similar API to CUDA, with many functions having the same names and parameters.
• HIP supports both AMD and NVIDIA GPUs, while CUDA is specific to NVIDIA GPUs.
• HIP code can be compiled with both HIPCC (for AMD GPUs) and NVCC (for NVIDIA GPUs).

HIP Runtime API Examples

Vector Addition Example

Exploiting the C++ toolset

Since HIP code is C++ at its core, we can utilize key C++ constructs to make code a bit safer than C. For example, we can make custom vector classes utilize C++'s context switching to handle the memory management for us. The C++ context switching is called Resource Acquisition Is Initialization (RAII).

>

Bonus: Towards heterogeneous programming

To reduce redundancy on writing CPU/GPU code, one can utilize macros and C++ lambda functions to achieve a single-implementation style program.

Bonus^2: Code Generation

One could think even more modular and use code generation libraries. For example, the python-based cogapp does this. To generate cpu or gpu code based on the platform, one could use the following code snippet within some, say, main.cpp.cog file:

The "some_cpu_code.cog" file might look like

and so forth...

Then generating the code with:

GPU_PLATFORM=amd cog -d -o main.cpp main.cpp.cog

Conclusion

HIP and ROCm provide a powerful platform for developing GPU-accelerated applications that can run on both AMD and NVIDIA hardware. By leveraging the similarities between HIP and CUDA, developers can more easily port their existing CUDA applications to run on AMD GPUs. Furthermore, lambda functions make generic programming on cpu/gpu architectures very streamlined. This is the bread and butter of third-party libraries like Kokkos.