diff --git a/projects/hip/README.md b/projects/hip/README.md index 145f950fef..2bffd12162 100644 --- a/projects/hip/README.md +++ b/projects/hip/README.md @@ -6,7 +6,7 @@ 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. +* The [HIPIFY](hipify-clang/README.md) tools automatically convert 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. @@ -36,7 +36,7 @@ HIP releases are typically of two types. The tag naming convention is different - [HIP Profiling ](docs/markdown/hip_profiling.md) - [HIP Debugging](docs/markdown/hip_debugging.md) - [HIP Terminology](docs/markdown/hip_terms.md) (including Rosetta Stone of GPU computing terms across CUDA/HIP/HC/AMP/OpenCL) -- [hipify-clang](hipify-clang/README.md) +- [HIPIFY](hipify-clang/README.md) - Supported CUDA APIs: * [Runtime API](docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md) * [Driver API](docs/markdown/CUDA_Driver_API_functions_supported_by_HIP.md) @@ -71,7 +71,7 @@ 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); +hipMemcpy(C_h, C_d, Nbytes, hipMemcpyDeviceToHost); ``` @@ -88,7 +88,7 @@ __global__ void vector_square(T *C_d, const T *A_d, size_t N) { size_t offset = (hipBlockIdx_x * hipBlockDim_x + hipThreadIdx_x); - size_t stride = hipBlockDim_x * hipGridDim_x ; + size_t stride = hipBlockDim_x * hipGridDim_x; for (size_t i=offset; i +- [hipify-clang](#clang) + * [Usage](#hipify-clang-usage) +- [hipify-perl](#perl) + * [Usage](#hipify-perl-usage) - [Supported CUDA APIs](#cuda-apis) - [Dependencies](#dependencies) - [Build and install](#build-and-install) @@ -13,12 +15,92 @@ * [Testing](#testing) * [Linux](#linux) * [Windows](#windows) -- [Running and using hipify-clang](#running-and-using-hipify-clang) - * [hipify-perl](#perl) - [Disclaimer](#disclaimer) +## hipify-clang + +`hipify-clang` is a clang-based tool for translation CUDA sources into HIP sources. +It translates CUDA source into an abstract syntax tree, which is being traversed by transformation matchers. +After applying all the matchers, the output HIP source is produced. + +**Advantages:** + +1. It is a translator; thus, any even very complicated constructs will be parsed successfully, or an error will be reported. +2. It supports clang options like -I, -D, --cuda-path, etc. +3. Seamless support of new CUDA versions as it is clang's responsibility. +4. Ease in support. + +**Disadvantages:** + +1. The main advantage is also the main disadvantage: the input CUDA code should be correct; incorrect code wouldn't be translated to HIP. +2. CUDA should be installed and provided in case of multiple installations by --cuda-path option. +3. All the includes and defines should be provided to transform code successfully. + +### hipify-clang usage + +To process a file, `hipify-clang` needs access to the same headers that would be required to compile it with clang. + +For example: + +```shell +./hipify-clang square.cu --cuda-path=/usr/local/cuda-10.1 -I /usr/local/cuda-10.1/samples/common/inc +``` + +`hipify-clang` arguments are given first, followed by a separator '--', and then the arguments you'd pass to `clang` if you +were compiling the input file. For example: + +```shell +./hipify-clang cpp17.cu --cuda-path=/usr/local/cuda-10.1 -- -std=c++17 +``` + +The [Clang manual for compiling CUDA](https://llvm.org/docs/CompileCudaWithLLVM.html#compiling-cuda-code) may be useful. + +For a list of `hipify-clang` options, run `hipify-clang --help`. + +## hipify-perl + +`hipify-perl` is autogenerated perl-based script which heavily uses regular expressions. + +**Advantages:** + +1. Ease in use. + +2. It doesn't check the input source CUDA code for correctness. + +3. It doesn't have dependencies on 3rd party tools, including CUDA. + +**Disadvantages:** + +1. Current disability (and difficulty in implementing) of transforming the following constructs: + + * macros expansion; + + * namespaces: + + - redefines of CUDA entities in user namespaces; + + - using directive; + + * templates (some cases); + + * device/host function calls distinguishing; + + * header files correct injection; + + * complicated argument lists parsing. + +2. Difficulties in supporting. + +### hipify-perl usage + +To generate `hipify-perl`, run `hipify-clang --perl`. + +```shell +perl hipify-perl square.cu > square.cu.hip +``` + ## Supported CUDA APIs - [Runtime API](../docs/markdown/CUDA_Runtime_API_functions_supported_by_HIP.md) @@ -391,47 +473,6 @@ cmake -- Generating done -- Build files have been written to: f:/HIP/hipify-clang/build ``` - -## Running and using hipify-clang - -To process a file, `hipify-clang` needs access to the same headers that would be needed to compile it with clang. - -For example: - -```shell -./hipify-clang square.cu --cuda-path=/usr/local/cuda-10.1 -I /usr/local/cuda-10.1/samples/common/inc -``` - -`hipify-clang` arguments are given first, followed by a separator, and then the arguments you'd pass to `clang` if you -were compiling the input file. The [Clang manual for compiling CUDA](https://llvm.org/docs/CompileCudaWithLLVM.html#compiling-cuda-code) -may be useful. - -For a list of `hipify-clang` options, run `hipify-clang --help`. - -### hipify-perl - -To produce a Perl-based script `hipify-perl`, run `hipify-clang --perl`. - -The `hipify-perl` script, unlike the `hipify-clang`, being based on regular expressions, and not on an abstract syntax tree, has several gaps: - -1. macros expansion; - -2. namespaces: - - - redefines of CUDA entities in user namespaces; - - - using directive; - -3. templates (some cases); - -4. device/host function calls distinguishing; - -5. header files correct injection; - -6. complicated argument lists parsing. - -Nonetheless, `hipify-perl` is easy in use and doesn't check the input source CUDA code for correctness. - ## Disclaimer The information contained herein is for informational purposes only, and is subject to change without notice. While every precaution has been taken in the preparation of this document, it may contain technical inaccuracies, omissions and typographical errors, and AMD is under no obligation to update or otherwise correct this information. Advanced Micro Devices, Inc. makes no representations or warranties with respect to the accuracy or completeness of the contents of this document, and assumes no liability of any kind, including the implied warranties of noninfringement, merchantability or fitness for particular purposes, with respect to the operation or use of AMD hardware, software or other products described herein. No license, including implied or arising by estoppel, to any intellectual property rights is granted by this document. Terms and limitations applicable to the purchase or use of AMD's products are as set forth in a signed agreement between the parties or in AMD's Standard Terms and Conditions of Sale.