Frequently asked questions

This topic provides answers to frequently asked questions from new HIP users and users familiar with NVIDIA CUDA.

HIP support

What hardware does HIP support?

HIP supports AMD GPUs. See prerequisites of the install guide for detailed information.

What operating systems does HIP support?

Linux as well as Windows are supported by ROCm. The exact versions are listed in the system requirements for Supported operating systems and rocm-install-on-windows:supported-skus-win.

Note

Not all HIP runtime API functions are yet supported on Windows. A note is added to those functions’ documentation in the HIP runtime API reference.

What libraries does HIP provide?

HIP provides key math and AI libraries. See ROCm libraries for the full list.

What CUDA features can be ported to HIP?

The NVIDIA CUDA runtime API supported by HIP and NVIDIA CUDA driver API supported by HIP pages describe which CUDA APIs can be automatically converted to HIP equivalents. The HIP runtime API reference describes each HIP API and its limitations, if any, compared with the equivalent CUDA API.

The kernel language features are documented in the HIP C++ language extensions page.

Relation to other GPGPU frameworks

How easy is it to port CUDA code to HIP?

The HIPIFY tools can automatically convert almost all CUDA runtime code to HIP. Most device code needs no additional conversion because HIP and CUDA have the same signatures for math and built-in functions except for the name. Once ported to HIP, code can be optimized for AMD GPU architectures.

Additional porting might be required to deal with architecture feature queries or CUDA capabilities that HIP doesn’t support.

To better understand the syntax differences, see CUDA to HIP API Function Comparison or the HIP porting guide.

Compiler related questions

How to use HIP-Clang to build HIP programs?

hipcc is a compiler driver. This means it is not a compiler but calls the appropriate compiler and sets some options.

The underlying compiler is amdclang++, and can be called directly.

What is HIP-Clang?

HIP-Clang is a Clang/LLVM-based compiler used to compile HIP programs for AMD platforms. The executable is named amdclang++ on Linux and clang++ on Windows.

Can I link HIP device code with host code compiled with another compiler such as gcc, icc, or clang?

Yes. HIP generates object code that conforms to the GCC ABI, and links with libstdc++. This means you can compile host code with the compiler of your choice and link the generated host object code with device code.

Can HIP applications be compiled with a C compiler?

HIP is a C/C++ API that can be used with C compilers. However, this applies only to the API itself. Device code and the syntax for calling kernels must be compiled with a supported compiler like hipcc. The code objects that are generated with hipcc can, however, be used with a C compiler, as shown in the code examples below.

The following is the HIP device code, assumed to be saved in device.hip:

#include <hip/hip_runtime.h>

__global__ void kernel(double* array, size_t size){
    const int x = threadIdx.x + blockIdx.x * blockDim.x;
    if(x < size){array[x] = x;}
};

extern "C"{
    hipError_t callKernel(int blocks, int threadsPerBlock, double* array, size_t size){
        kernel<<<blocks, threadsPerBlock, 0, hipStreamDefault>>>(array, size);
        return hipGetLastError();
    }
}

The following is the host code, written in C, saved in host.c:

#include <hip/hip_runtime_api.h>
#include <stdio.h>
#include <stdlib.h>

#define HIP_CHECK(c) {                                \
   if (c != hipSuccess){                              \
      printf("HIP Error : %s", hipGetErrorString(c)); \
      printf(" %s %d\n", __FILE__, __LINE__);         \
      exit(c);                                        \
   }                                                  \
}

// Forward declaration - the implementation needs to be compiled with
// a device compiler like hipcc or amdclang++
hipError_t callKernel(int blocks, int threadsPerBlock, double* array, size_t size);

int main(int argc, char** argv) {
    int blocks = 1024;
    int threadsPerBlock = 256;
    size_t arraySize = blocks * threadsPerBlock;
    double* d_array;
    double* h_array;
    h_array = (double*)malloc(arraySize * sizeof(double));

    HIP_CHECK(hipMalloc((void**)&d_array, arraySize * sizeof(double)));
    HIP_CHECK(callKernel(blocks, threadsPerBlock, d_array, arraySize));
    HIP_CHECK(hipMemcpy(h_array, d_array, arraySize * sizeof(double), hipMemcpyDeviceToHost));
    HIP_CHECK(hipFree(d_array));

    free(h_array);
    return 0;
}

These files are then compiled and linked using

hipcc -c device.hip
gcc host.c device.o $(hipconfig --cpp_config) -L/opt/rocm/lib -lamdhip64

assuming the default installation of ROCm in /opt/rocm.

How to guard code specific to the host or the GPU?

The compiler defines the __HIP_DEVICE_COMPILE__ macro only when compiling device code.

Refer to the Porting CUDA code to HIP for more information.