What is this repository for?
HIP allows developers to convert CUDA code to portable C++. The same source code can be compiled to run on NVIDIA or AMD GPUs. Key features include:
- HIP is very thin and has little or no performance impact over coding directly in CUDA or hcc “HC” mode.
- HIP allows coding in a single-source C++ programming language including features such as templates, C++11 lambdas, classes, namespaces, and more.
- HIP allows developers to use the “best” development environment and tools on each target platform.
- The “hipify” tool automatically converts source from CUDA to HIP.
- Developers can specialize for the platform (CUDA or hcc) to tune for performance or handle tricky cases
New projects can be developed directly in the portable HIP C++ language and can run on either NVIDIA or AMD platforms. Additionally, HIP provides porting tools which make it easy to port existing CUDA codes to the HIP layer, with no loss of performance as compared to the original CUDA application. HIP is not intended to be a drop-in replacement for CUDA, and developers should expect to do some manual coding and performance tuning work to complete the port.
The HIP repository maintains several branches. The branches that are of importance are:
- master branch: This is the stable branch. All stable releases are based on this branch.
- developer-preview branch: This is the branch were the new features still under development are visible. While this maybe of interest to many, it should be noted that this branch and the features under development might not be stable.
HIP releases are typically of two types. The tag naming convention is different for both types of releases to help differentiate them.
- release_x.yy.zzzz: These are the stable releases based on the master branch. This type of release is typically made once a month.
- preview_x.yy.zzzz: These denote pre-release code and are based on the developer-preview branch. This type of release is typically made once a week.
How do I get set up?
See the Installation notes.
The HIP API includes functions such as hipMalloc, hipMemcpy, and hipFree.
Programmers familiar with CUDA will also be able to quickly learn and start coding with the HIP API. Compute kernels are launched with the “hipLaunchKernel” macro call. Here is simple example showing a snippet of HIP API code:
hipMalloc(&C_d, Nbytes)); hipMemcpy(A_d, A_h, Nbytes, hipMemcpyHostToDevice); const unsigned blocks = 512;
const unsigned threadsPerBlock = 256;
hipLaunchKernel(vector_square, /* compute kernel*/ dim3(blocks), dim3(threadsPerBlock), 0/*dynamic shared*/, 0/*stream*/, /* launch config*/ C_d, A_d, N); /* arguments to the compute kernel */ hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost);