diff --git a/projects/hip/bin/hipify-perl b/projects/hip/bin/hipify-perl
index 25e02352ad..2e391ab8d2 100755
--- a/projects/hip/bin/hipify-perl
+++ b/projects/hip/bin/hipify-perl
@@ -769,6 +769,7 @@ sub simpleSubstitutions {
$ft{'include'} += s/\btexture_fetch_functions.h\b//g;
$ft{'include'} += s/\bvector_types.h\b/hip\/hip_vector_types.h/g;
$ft{'include_cuda_main_header'} += s/\bcuComplex.h\b/hip\/hip_complex.h/g;
+ $ft{'include_cuda_main_header'} += s/\bcub\/cub.cuh\b/hipcub\/hipcub.hpp/g;
$ft{'include_cuda_main_header'} += s/\bcublas.h\b/hipblas.h/g;
$ft{'include_cuda_main_header'} += s/\bcublas_v2.h\b/hipblas.h/g;
$ft{'include_cuda_main_header'} += s/\bcuda.h\b/hip\/hip_runtime.h/g;
diff --git a/projects/hip/hipify-clang/CMakeLists.txt b/projects/hip/hipify-clang/CMakeLists.txt
index 5eb71b3a3e..fd172a9c8d 100644
--- a/projects/hip/hipify-clang/CMakeLists.txt
+++ b/projects/hip/hipify-clang/CMakeLists.txt
@@ -1,8 +1,4 @@
-if (CUDA_VERSION VERSION_GREATER "9.2")
- cmake_minimum_required(VERSION 3.12.3)
-else()
- cmake_minimum_required(VERSION 3.7.2)
-endif()
+cmake_minimum_required(VERSION 3.5.2)
project(hipify-clang)
@@ -58,7 +54,7 @@ target_link_libraries(hipify-clang PRIVATE
LLVMOption
LLVMCore)
-if (LLVM_PACKAGE_VERSION VERSION_GREATER_EQUAL "7")
+if (LLVM_PACKAGE_VERSION VERSION_GREATER "6.0.1")
target_link_libraries(hipify-clang PRIVATE clangToolingInclusions)
endif()
@@ -72,7 +68,7 @@ else()
set(StdCpp "-std=c++")
endif()
-if (LLVM_PACKAGE_VERSION VERSION_GREATER_EQUAL "10.0")
+if (LLVM_PACKAGE_VERSION VERSION_GREATER "9.0")
string(APPEND StdCpp "14")
# MSVC starting from 1900 (VS 2015) supports only the following c++ std values: c++14|c++17|c++latest
elseif (MSVC)
diff --git a/projects/hip/hipify-clang/README.md b/projects/hip/hipify-clang/README.md
index ab52c288c8..bb9654d8b4 100644
--- a/projects/hip/hipify-clang/README.md
+++ b/projects/hip/hipify-clang/README.md
@@ -63,7 +63,7 @@
In most cases, you can get a suitable version of LLVM+CLANG with your package manager.
Failing that or having multiple versions of LLVM, you can [download a release archive](http://releases.llvm.org/), build or install it, and set
-[CMAKE_PREFIX_PATH](https://cmake.org/cmake/help/v3.12/variable/CMAKE_PREFIX_PATH.html) so `cmake` can find it; for instance: `-DCMAKE_PREFIX_PATH=f:\LLVM\9.0.0\dist`
+[CMAKE_PREFIX_PATH](https://cmake.org/cmake/help/v3.5/variable/CMAKE_PREFIX_PATH.html) so `cmake` can find it; for instance: `-DCMAKE_PREFIX_PATH=f:\LLVM\9.0.0\dist`
## Build and install
@@ -119,7 +119,8 @@ To run it:
- **Windows**:
```shell
cmake \
- -G "Visual Studio 16 2019 Win64" \
+ -G "Visual Studio 16 2019" \
+ -A x64 \
-DCMAKE_INSTALL_PREFIX=../dist \
-DLLVM_SOURCE_DIR=../llvm \
-DLLVM_TARGETS_TO_BUILD="X86;NVPTX" \
@@ -149,6 +150,14 @@ To run it:
- Windows: `-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-10.1-windows10-x64-v7.6.4.38`
+5. Ensure [`CUB`](https://github.com/NVlabs/cub) of the version corresponding to CUDA's version is installed.
+
+ * Path to CUB should be specified by the `CUDA_CUB_ROOT_DIR` option:
+
+ - Linux: `-DCUDA_CUB_ROOT_DIR=/srv/CUB`
+
+ - Windows: `-DCUDA_CUB_ROOT_DIR=f:/GIT/cub`
+
5. Ensure [`python`](https://www.python.org/downloads) of minimum required version 2.7 is installed.
6. Ensure `lit` and `FileCheck` are installed - these are distributed with LLVM.
@@ -187,7 +196,7 @@ Ubuntu 16-18: LLVM 8.0.0 - 9.0.0, CUDA 8.0 - 10.1, cudnn-5.1.10 - cudnn-7.6.4.38
Build system for the above configurations:
-Python 2.7 (min), cmake 3.12.3 (min), GNU C/C++ 5.4.0 (min).
+Python 2.7 (min), cmake 3.5.2 (min), GNU C/C++ 5.4.0 (min).
Here is an example of building `hipify-clang` with testing support on `Ubuntu 16.04`:
@@ -199,6 +208,7 @@ cmake
-DCMAKE_PREFIX_PATH=/srv/git/LLVM/9.0.0/dist \
-DCUDA_TOOLKIT_ROOT_DIR=/usr/local/cuda-10.1 \
-DCUDA_DNN_ROOT_DIR=/srv/CUDNN/cudnn-10.1-v7.6.4.38 \
+ -DCUDA_CUB_ROOT_DIR=/srv/CUB \
-DLLVM_EXTERNAL_LIT=/srv/git/LLVM/9.0.0/build/bin/llvm-lit \
..
```
@@ -254,72 +264,73 @@ Linux 5.2.0 - Platform OS
64 - hipify-clang binary bitness
64 - python 2.7.12 binary bitness
========================================
--- Testing: 63 tests, 12 threads --
-PASS: hipify :: unit_tests/casts/reinterpret_cast.cu (1 of 63)
-PASS: hipify :: unit_tests/headers/headers_test_01.cu (2 of 63)
-PASS: hipify :: unit_tests/headers/headers_test_03.cu (3 of 63)
-PASS: hipify :: unit_tests/headers/headers_test_02.cu (4 of 63)
-PASS: hipify :: unit_tests/headers/headers_test_05.cu (5 of 63)
-PASS: hipify :: unit_tests/device/math_functions.cu (6 of 63)
-PASS: hipify :: unit_tests/device/atomics.cu (7 of 63)
-PASS: hipify :: unit_tests/headers/headers_test_07.cu (8 of 63)
-PASS: hipify :: unit_tests/headers/headers_test_06.cu (9 of 63)
-PASS: hipify :: unit_tests/headers/headers_test_04.cu (10 of 63)
-PASS: hipify :: unit_tests/device/device_symbols.cu (11 of 63)
-PASS: hipify :: unit_tests/headers/headers_test_10.cu (12 of 63)
-PASS: hipify :: unit_tests/headers/headers_test_11.cu (13 of 63)
-PASS: hipify :: unit_tests/headers/headers_test_08.cu (14 of 63)
-PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_02.cu (15 of 63)
-PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_0_based_indexing.cu (16 of 63)
-PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_1_based_indexing.cu (17 of 63)
-PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_sgemm_matrix_multiplication.cu (18 of 63)
-PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_0_based_indexing_rocblas.cu (19 of 63)
-PASS: hipify :: unit_tests/libraries/cuComplex/cuComplex_Julia.cu (20 of 63)
-PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_01.cu (21 of 63)
-PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_softmax.cu (22 of 63)
-PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_1_based_indexing_rocblas.cu (23 of 63)
-PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_sgemm_matrix_multiplication_rocblas.cu (24 of 63)
-PASS: hipify :: unit_tests/libraries/cuFFT/simple_cufft.cu (25 of 63)
-PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu (26 of 63)
-PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_01.cu (27 of 63)
-PASS: hipify :: unit_tests/headers/headers_test_09.cu (28 of 63)
-PASS: hipify :: unit_tests/libraries/cuRAND/poisson_api_example.cu (29 of 63)
-PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu (30 of 63)
-PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_02.cu (31 of 63)
-PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_05.cu (32 of 63)
-PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_04.cu (33 of 63)
-PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_06.cu (34 of 63)
-PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_07.cu (35 of 63)
-PASS: hipify :: unit_tests/namespace/ns_kernel_launch.cu (36 of 63)
-PASS: hipify :: unit_tests/pp/pp_if_else_conditionals_01.cu (37 of 63)
-PASS: hipify :: unit_tests/pp/pp_if_else_conditionals.cu (38 of 63)
-PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_08.cu (39 of 63)
-PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_11.cu (40 of 63)
-PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp (41 of 63)
-PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu (42 of 63)
-PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_09.cu (43 of 63)
-PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp (44 of 63)
-PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_12.cu (45 of 63)
-PASS: hipify :: unit_tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp (46 of 63)
-PASS: hipify :: unit_tests/samples/allocators.cu (47 of 63)
-PASS: hipify :: unit_tests/samples/MallocManaged.cpp (48 of 63)
-PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp (49 of 63)
-PASS: hipify :: unit_tests/samples/coalescing.cu (50 of 63)
-PASS: hipify :: unit_tests/samples/2_Cookbook/1_hipEvent/hipEvent.cpp (51 of 63)
-PASS: hipify :: unit_tests/samples/2_Cookbook/13_occupancy/occupancy.cpp (52 of 63)
-PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp (53 of 63)
-PASS: hipify :: unit_tests/samples/2_Cookbook/2_Profiler/Profiler.cpp (54 of 63)
-PASS: hipify :: unit_tests/samples/2_Cookbook/7_streams/stream.cpp (55 of 63)
-PASS: hipify :: unit_tests/samples/2_Cookbook/8_peer2peer/peer2peer.cpp (56 of 63)
-PASS: hipify :: unit_tests/samples/intro.cu (57 of 63)
-PASS: hipify :: unit_tests/samples/dynamic_shared_memory.cu (58 of 63)
-PASS: hipify :: unit_tests/samples/axpy.cu (59 of 63)
-PASS: hipify :: unit_tests/samples/square.cu (60 of 63)
-PASS: hipify :: unit_tests/samples/vec_add.cu (61 of 63)
-PASS: hipify :: unit_tests/samples/static_shared_memory.cu (62 of 63)
-PASS: hipify :: unit_tests/samples/cudaRegister.cu (63 of 63)
-Testing Time: 2.91s
- Expected Passes : 63
+-- Testing: 64 tests, 12 threads --
+PASS: hipify :: unit_tests/casts/reinterpret_cast.cu (1 of 64)
+PASS: hipify :: unit_tests/device/math_functions.cu (2 of 64)
+PASS: hipify :: unit_tests/device/atomics.cu (3 of 64)
+PASS: hipify :: unit_tests/device/device_symbols.cu (4 of 64)
+PASS: hipify :: unit_tests/headers/headers_test_02.cu (5 of 64)
+PASS: hipify :: unit_tests/headers/headers_test_03.cu (6 of 64)
+PASS: hipify :: unit_tests/headers/headers_test_01.cu (7 of 64)
+PASS: hipify :: unit_tests/headers/headers_test_04.cu (8 of 64)
+PASS: hipify :: unit_tests/headers/headers_test_05.cu (9 of 64)
+PASS: hipify :: unit_tests/headers/headers_test_07.cu (10 of 64)
+PASS: hipify :: unit_tests/headers/headers_test_06.cu (11 of 64)
+PASS: hipify :: unit_tests/headers/headers_test_11.cu (12 of 64)
+PASS: hipify :: unit_tests/headers/headers_test_08.cu (13 of 64)
+PASS: hipify :: unit_tests/headers/headers_test_10.cu (14 of 64)
+PASS: hipify :: unit_tests/headers/headers_test_09.cu (15 of 64)
+PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_02.cu (16 of 64)
+PASS: hipify :: unit_tests/libraries/CAFFE2/caffe2_01.cu (17 of 64)
+PASS: hipify :: unit_tests/libraries/CUB/cub_01.cu (18 of 64)
+PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_1_based_indexing.cu (19 of 64)
+PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_0_based_indexing.cu (20 of 64)
+PASS: hipify :: unit_tests/libraries/cuBLAS/cublas_sgemm_matrix_multiplication.cu (21 of 64)
+PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_0_based_indexing_rocblas.cu (22 of 64)
+PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_sgemm_matrix_multiplication_rocblas.cu (23 of 64)
+PASS: hipify :: unit_tests/libraries/cuBLAS/rocBLAS/cublas_1_based_indexing_rocblas.cu (24 of 64)
+PASS: hipify :: unit_tests/libraries/cuComplex/cuComplex_Julia.cu (25 of 64)
+PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_softmax.cu (26 of 64)
+PASS: hipify :: unit_tests/libraries/cuDNN/cudnn_convolution_forward.cu (27 of 64)
+PASS: hipify :: unit_tests/libraries/cuFFT/simple_cufft.cu (28 of 64)
+PASS: hipify :: unit_tests/libraries/cuRAND/poisson_api_example.cu (29 of 64)
+PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_01.cu (30 of 64)
+PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_generate.cpp (31 of 64)
+PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_02.cu (32 of 64)
+PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_03.cu (33 of 64)
+PASS: hipify :: unit_tests/libraries/cuRAND/benchmark_curand_kernel.cpp (34 of 64)
+PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_04.cu (35 of 64)
+PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_05.cu (36 of 64)
+PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_06.cu (37 of 64)
+PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_07.cu (38 of 64)
+PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_09.cu (39 of 64)
+PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_10.cu (40 of 64)
+PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_08.cu (41 of 64)
+PASS: hipify :: unit_tests/namespace/ns_kernel_launch.cu (42 of 64)
+PASS: hipify :: unit_tests/pp/pp_if_else_conditionals.cu (43 of 64)
+PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_11.cu (44 of 64)
+PASS: hipify :: unit_tests/libraries/cuSPARSE/cuSPARSE_12.cu (45 of 64)
+PASS: hipify :: unit_tests/pp/pp_if_else_conditionals_01.cu (46 of 64)
+PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/tex2dKernel.cpp (47 of 64)
+PASS: hipify :: unit_tests/samples/2_Cookbook/0_MatrixTranspose/MatrixTranspose.cpp (48 of 64)
+PASS: hipify :: unit_tests/samples/2_Cookbook/11_texture_driver/texture2dDrv.cpp (49 of 64)
+PASS: hipify :: unit_tests/samples/2_Cookbook/13_occupancy/occupancy.cpp (50 of 64)
+PASS: hipify :: unit_tests/samples/2_Cookbook/1_hipEvent/hipEvent.cpp (51 of 64)
+PASS: hipify :: unit_tests/samples/2_Cookbook/2_Profiler/Profiler.cpp (52 of 64)
+PASS: hipify :: unit_tests/samples/MallocManaged.cpp (53 of 64)
+PASS: hipify :: unit_tests/samples/2_Cookbook/7_streams/stream.cpp (54 of 64)
+PASS: hipify :: unit_tests/samples/2_Cookbook/8_peer2peer/peer2peer.cpp (55 of 64)
+PASS: hipify :: unit_tests/samples/allocators.cu (56 of 64)
+PASS: hipify :: unit_tests/samples/coalescing.cu (57 of 64)
+PASS: hipify :: unit_tests/samples/dynamic_shared_memory.cu (58 of 64)
+PASS: hipify :: unit_tests/samples/axpy.cu (59 of 64)
+PASS: hipify :: unit_tests/samples/cudaRegister.cu (60 of 64)
+PASS: hipify :: unit_tests/samples/intro.cu (61 of 64)
+PASS: hipify :: unit_tests/samples/square.cu (62 of 64)
+PASS: hipify :: unit_tests/samples/static_shared_memory.cu (63 of 64)
+PASS: hipify :: unit_tests/samples/vec_add.cu (64 of 64)
+Testing Time: 2.98s
+ Expected Passes : 64
[100%] Built target test-hipify
```
### Windows
@@ -334,13 +345,14 @@ LLVM 7.0.0 - 9.0.0, CUDA 7.5 - 10.1, cudnn-7.0.5.15 - cudnn-7.6.4.38
Build system for the above configurations:
-Python 3.6 - 3.7.4, cmake 3.12.3 - 3.15.4, Visual Studio 2017 (15.5.2) - 2019 (16.3.2).
+Python 3.6 - 3.7.4, cmake 3.5.2 - 3.15.5, Visual Studio 2017 (15.5.2) - 2019 (16.3.4).
Here is an example of building `hipify-clang` with testing support on `Windows 10` by `Visual Studio 16 2019`:
```shell
cmake
- -G "Visual Studio 16 2019 Win64" \
+ -G "Visual Studio 16 2019" \
+ -A x64 \
-DHIPIFY_CLANG_TESTS=1 \
-DCMAKE_BUILD_TYPE=Release \
-DCMAKE_INSTALL_PREFIX=../dist \
@@ -348,6 +360,7 @@ cmake
-DCUDA_TOOLKIT_ROOT_DIR="c:/Program Files/NVIDIA GPU Computing Toolkit/CUDA/v10.1" \
-DCUDA_SDK_ROOT_DIR="c:/ProgramData/NVIDIA Corporation/CUDA Samples/v10.1" \
-DCUDA_DNN_ROOT_DIR=f:/CUDNN/cudnn-10.1-windows10-x64-v7.6.4.38 \
+ -DCUDA_CUB_ROOT_DIR=f:/GIT/cub \
-DLLVM_EXTERNAL_LIT=f:/LLVM/9.0.0/build/Release/bin/llvm-lit.py \
-Thost=x64
..
diff --git a/projects/hip/hipify-clang/src/CUDA2HIP.cpp b/projects/hip/hipify-clang/src/CUDA2HIP.cpp
index d2f11df458..50f4682af2 100644
--- a/projects/hip/hipify-clang/src/CUDA2HIP.cpp
+++ b/projects/hip/hipify-clang/src/CUDA2HIP.cpp
@@ -67,6 +67,8 @@ const std::map CUDA_INCLUDE_MAP{
// cuSPARSE includes
{"cusparse.h", {"hipsparse.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_SPARSE}},
{"cusparse_v2.h", {"hipsparse.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_SPARSE}},
+ // CUB includes
+ {"cub/cub.cuh", {"hipcub/hipcub.hpp", "", CONV_INCLUDE_CUDA_MAIN_H, API_CUB}},
// CAFFE2 includes
{"caffe2/core/common_gpu.h", {"caffe2/core/hip/common_gpu.h", "", CONV_INCLUDE, API_CAFFE2, UNSUPPORTED}},
{"caffe2/core/context_gpu.h", {"caffe2/core/hip/context_gpu.h", "", CONV_INCLUDE, API_CAFFE2, UNSUPPORTED}},
diff --git a/projects/hip/hipify-clang/src/CUDA2HIP.h b/projects/hip/hipify-clang/src/CUDA2HIP.h
index acddd23a0d..b02e7f1f3e 100644
--- a/projects/hip/hipify-clang/src/CUDA2HIP.h
+++ b/projects/hip/hipify-clang/src/CUDA2HIP.h
@@ -67,6 +67,8 @@ extern const std::map CUDA_CAFFE2_TYPE_NAME_MAP;
extern const std::map CUDA_CAFFE2_FUNCTION_MAP;
// Maps the names of CUDA Device functions to the corresponding HIP functions
extern const std::map CUDA_DEVICE_FUNC_MAP;
+// Maps the names of CUDA CUB API types to the corresponding HIP types
+extern const std::map CUDA_CUB_TYPE_NAME_MAP;
/**
* The union of all the above maps, except includes.
diff --git a/projects/hip/hipify-clang/src/CUDA2HIP_CUB_API_types.cpp b/projects/hip/hipify-clang/src/CUDA2HIP_CUB_API_types.cpp
new file mode 100644
index 0000000000..0ef1912b54
--- /dev/null
+++ b/projects/hip/hipify-clang/src/CUDA2HIP_CUB_API_types.cpp
@@ -0,0 +1,28 @@
+/*
+Copyright (c) 2015 - present Advanced Micro Devices, Inc. All rights reserved.
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in
+all copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN
+THE SOFTWARE.
+*/
+
+#include "CUDA2HIP.h"
+
+// Maps the names of CUDA CUB API types to the corresponding HIP types
+const std::map CUDA_CUB_TYPE_NAME_MAP{
+ {"cub", {"hipcub", "", CONV_TYPE, API_CUB}},
+};
diff --git a/projects/hip/hipify-clang/src/HipifyAction.cpp b/projects/hip/hipify-clang/src/HipifyAction.cpp
index 5c5663985a..930f3ec8c5 100644
--- a/projects/hip/hipify-clang/src/HipifyAction.cpp
+++ b/projects/hip/hipify-clang/src/HipifyAction.cpp
@@ -33,16 +33,20 @@ THE SOFTWARE.
#include "StringUtils.h"
#include "ArgParse.h"
-namespace ct = clang::tooling;
-namespace mat = clang::ast_matchers;
-
+const std::string sHIP = "HIP";
+const std::string sROC = "ROC";
+const std::string sCub = "cub";
const std::string sHIP_DYNAMIC_SHARED = "HIP_DYNAMIC_SHARED";
+const std::string sHIP_KERNEL_NAME = "HIP_KERNEL_NAME";
std::string sHIP_SYMBOL = "HIP_SYMBOL";
-std::string sHIP_KERNEL_NAME = "HIP_KERNEL_NAME";
std::string s_reinterpret_cast = "reinterpret_cast";
-const std::string sHipLaunchKernelGGL = "hipLaunchKernelGGL(";
+const std::string sHipLaunchKernelGGL = "hipLaunchKernelGGL";
const std::string sDim3 = "dim3(";
-
+const std::string s_hiprand_kernel_h = "hiprand_kernel.h";
+const std::string s_hiprand_h = "hiprand.h";
+const std::string sOnce = "once";
+const std::string s_string_literal = "[string literal]";
+// CUDA identifiers, used in matchers
const std::string sCudaMemcpyToSymbol = "cudaMemcpyToSymbol";
const std::string sCudaMemcpyToSymbolAsync = "cudaMemcpyToSymbolAsync";
const std::string sCudaGetSymbolSize = "cudaGetSymbolSize";
@@ -51,6 +55,12 @@ const std::string sCudaMemcpyFromSymbol = "cudaMemcpyFromSymbol";
const std::string sCudaMemcpyFromSymbolAsync = "cudaMemcpyFromSymbolAsync";
const std::string sCudaFuncSetCacheConfig = "cudaFuncSetCacheConfig";
const std::string sCudaFuncGetAttributes = "cudaFuncGetAttributes";
+// Matchers' names
+const StringRef sCudaSharedIncompleteArrayVar = "cudaSharedIncompleteArrayVar";
+const StringRef sCudaLaunchKernel = "cudaLaunchKernel";
+const StringRef sCudaHostFuncCall = "cudaHostFuncCall";
+const StringRef sCudaDeviceFuncCall = "cudaDeviceFuncCall";
+const StringRef sCubNamespacePrefix = "cubNamespacePrefix";
std::set DeviceSymbolFunctions0 {
{sCudaMemcpyToSymbol},
@@ -78,7 +88,7 @@ std::set ReinterpretFunctions1{
};
void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) {
- clang::SourceManager& SM = getCompilerInstance().getSourceManager();
+ auto &SM = getCompilerInstance().getSourceManager();
size_t begin = 0;
while ((begin = s.find("cu", begin)) != StringRef::npos) {
const size_t end = s.find_first_of(" ", begin + 4);
@@ -86,7 +96,7 @@ void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) {
const auto found = CUDA_RENAMES_MAP().find(name);
if (found != CUDA_RENAMES_MAP().end()) {
StringRef repName = Statistics::isToRoc(found->second) ? found->second.rocName : found->second.hipName;
- hipCounter counter = {"[string literal]", "", ConvTypes::CONV_LITERAL, ApiTypes::API_RUNTIME, found->second.supportDegree};
+ hipCounter counter = {s_string_literal, "", ConvTypes::CONV_LITERAL, ApiTypes::API_RUNTIME, found->second.supportDegree};
Statistics::current().incrementCounter(counter, name.str());
if (!Statistics::isUnsupported(counter)) {
clang::SourceLocation sl = start.getLocWithOffset(begin + 1);
@@ -95,9 +105,7 @@ void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) {
insertReplacement(Rep, fullSL);
}
}
- if (end == StringRef::npos) {
- break;
- }
+ if (end == StringRef::npos) break;
begin = end + 1;
}
}
@@ -109,7 +117,7 @@ void HipifyAction::RewriteString(StringRef s, clang::SourceLocation start) {
* If it's an unsupported CUDA identifier, a warning is emitted.
* Otherwise, the source file is updated with the corresponding hipification.
*/
-void HipifyAction::RewriteToken(const clang::Token& t) {
+void HipifyAction::RewriteToken(const clang::Token &t) {
// String literals containing CUDA references need fixing.
if (t.is(clang::tok::string_literal)) {
StringRef s(t.getLiteralData(), t.getLength());
@@ -124,20 +132,21 @@ void HipifyAction::RewriteToken(const clang::Token& t) {
FindAndReplace(name, sl, CUDA_RENAMES_MAP());
}
-void HipifyAction::FindAndReplace(llvm::StringRef name,
+void HipifyAction::FindAndReplace(StringRef name,
clang::SourceLocation sl,
- const std::map& repMap, bool bReplace) {
+ const std::map &repMap,
+ bool bReplace) {
const auto found = repMap.find(name);
if (found == repMap.end()) {
// So it's an identifier, but not CUDA? Boring.
return;
}
Statistics::current().incrementCounter(found->second, name.str());
- clang::DiagnosticsEngine& DE = getCompilerInstance().getDiagnostics();
+ clang::DiagnosticsEngine &DE = getCompilerInstance().getDiagnostics();
// Warn the user about unsupported identifier.
if (Statistics::isUnsupported(found->second)) {
std::string sWarn;
- Statistics::isToRoc(found->second) ? sWarn = "ROC" : sWarn = "HIP";
+ Statistics::isToRoc(found->second) ? sWarn = sROC : sWarn = sHIP;
sWarn = "" + sWarn;
const auto ID = DE.getCustomDiagID(clang::DiagnosticsEngine::Warning, "CUDA identifier is unsupported in %0.");
DE.Report(sl, ID) << sWarn;
@@ -147,7 +156,7 @@ void HipifyAction::FindAndReplace(llvm::StringRef name,
return;
}
StringRef repName = Statistics::isToRoc(found->second) ? found->second.rocName : found->second.hipName;
- clang::SourceManager& SM = getCompilerInstance().getSourceManager();
+ auto &SM = getCompilerInstance().getSourceManager();
ct::Replacement Rep(SM, sl, name.size(), repName.str());
clang::FullSourceLoc fullSL(sl, SM);
insertReplacement(Rep, fullSL);
@@ -155,7 +164,7 @@ void HipifyAction::FindAndReplace(llvm::StringRef name,
namespace {
-clang::SourceRange getReadRange(clang::SourceManager& SM, const clang::SourceRange& exprRange) {
+clang::SourceRange getReadRange(clang::SourceManager &SM, const clang::SourceRange &exprRange) {
clang::SourceLocation begin = exprRange.getBegin();
clang::SourceLocation end = exprRange.getEnd();
bool beginSafe = !SM.isMacroBodyExpansion(begin) || clang::Lexer::isAtStartOfMacroExpansion(begin, SM, clang::LangOptions{});
@@ -167,7 +176,7 @@ clang::SourceRange getReadRange(clang::SourceManager& SM, const clang::SourceRan
}
}
-clang::SourceRange getWriteRange(clang::SourceManager& SM, const clang::SourceRange& exprRange) {
+clang::SourceRange getWriteRange(clang::SourceManager &SM, const clang::SourceRange &exprRange) {
clang::SourceLocation begin = exprRange.getBegin();
clang::SourceLocation end = exprRange.getEnd();
// If the range is contained within a macro, update the macro definition.
@@ -178,7 +187,7 @@ clang::SourceRange getWriteRange(clang::SourceManager& SM, const clang::SourceRa
return {SM.getSpellingLoc(begin), SM.getSpellingLoc(end)};
}
-StringRef readSourceText(clang::SourceManager& SM, const clang::SourceRange& exprRange) {
+StringRef readSourceText(clang::SourceManager &SM, const clang::SourceRange &exprRange) {
return clang::Lexer::getSourceText(clang::CharSourceRange::getTokenRange(getReadRange(SM, exprRange)), SM, clang::LangOptions(), nullptr);
}
@@ -186,53 +195,50 @@ StringRef readSourceText(clang::SourceManager& SM, const clang::SourceRange& exp
* Get a string representation of the expression `arg`, unless it's a defaulting function
* call argument, in which case get a 0. Used for building argument lists to kernel calls.
*/
-std::string stringifyZeroDefaultedArg(clang::SourceManager& SM, const clang::Expr* arg) {
- if (clang::isa(arg)) {
- return "0";
- } else {
- return readSourceText(SM, arg->getSourceRange());
- }
+std::string stringifyZeroDefaultedArg(clang::SourceManager &SM, const clang::Expr *arg) {
+ if (clang::isa(arg)) return "0";
+ else return readSourceText(SM, arg->getSourceRange());
}
} // anonymous namespace
-bool HipifyAction::Exclude(const hipCounter & hipToken) {
+bool HipifyAction::Exclude(const hipCounter &hipToken) {
switch (hipToken.type) {
case CONV_INCLUDE_CUDA_MAIN_H:
switch (hipToken.apiType) {
case API_DRIVER:
case API_RUNTIME:
- if (insertedRuntimeHeader) { return true; }
+ if (insertedRuntimeHeader) return true;
insertedRuntimeHeader = true;
return false;
case API_BLAS:
- if (insertedBLASHeader) { return true; }
+ if (insertedBLASHeader) return true;
insertedBLASHeader = true;
return false;
case API_RAND:
- if (hipToken.hipName == "hiprand_kernel.h") {
- if (insertedRAND_kernelHeader) { return true; }
+ if (hipToken.hipName == s_hiprand_kernel_h) {
+ if (insertedRAND_kernelHeader) return true;
insertedRAND_kernelHeader = true;
return false;
- } else if (hipToken.hipName == "hiprand.h") {
- if (insertedRANDHeader) { return true; }
+ } else if (hipToken.hipName == s_hiprand_h) {
+ if (insertedRANDHeader) return true;
insertedRANDHeader = true;
return false;
}
case API_DNN:
- if (insertedDNNHeader) { return true; }
+ if (insertedDNNHeader) return true;
insertedDNNHeader = true;
return false;
case API_FFT:
- if (insertedFFTHeader) { return true; }
+ if (insertedFFTHeader) return true;
insertedFFTHeader = true;
return false;
case API_COMPLEX:
- if (insertedComplexHeader) { return true; }
+ if (insertedComplexHeader) return true;
insertedComplexHeader = true;
return false;
case API_SPARSE:
- if (insertedSPARSEHeader) { return true; }
+ if (insertedSPARSEHeader) return true;
insertedSPARSEHeader = true;
return false;
default:
@@ -240,13 +246,11 @@ bool HipifyAction::Exclude(const hipCounter & hipToken) {
}
return false;
case CONV_INCLUDE:
- if (hipToken.hipName.empty()) {
- return true;
- }
+ if (hipToken.hipName.empty()) return true;
switch (hipToken.apiType) {
case API_RAND:
- if (hipToken.hipName == "hiprand_kernel.h") {
- if (insertedRAND_kernelHeader) { return true; }
+ if (hipToken.hipName == s_hiprand_kernel_h) {
+ if (insertedRAND_kernelHeader) return true;
insertedRAND_kernelHeader = true;
}
return false;
@@ -267,24 +271,19 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc,
clang::CharSourceRange filename_range,
const clang::FileEntry*, StringRef,
StringRef, const clang::Module*) {
- clang::SourceManager& SM = getCompilerInstance().getSourceManager();
- if (!SM.isWrittenInMainFile(hash_loc)) {
- return;
- }
+ auto &SM = getCompilerInstance().getSourceManager();
+ if (!SM.isWrittenInMainFile(hash_loc)) return;
if (!firstHeader) {
firstHeader = true;
firstHeaderLoc = hash_loc;
}
const auto found = CUDA_INCLUDE_MAP.find(file_name);
- if (found == CUDA_INCLUDE_MAP.end()) {
- return;
- }
+ if (found == CUDA_INCLUDE_MAP.end()) return;
bool exclude = Exclude(found->second);
Statistics::current().incrementCounter(found->second, file_name.str());
-
clang::SourceLocation sl = filename_range.getBegin();
if (Statistics::isUnsupported(found->second)) {
- clang::DiagnosticsEngine& DE = getCompilerInstance().getDiagnostics();
+ clang::DiagnosticsEngine &DE = getCompilerInstance().getDiagnostics();
DE.Report(sl, DE.getCustomDiagID(clang::DiagnosticsEngine::Warning, "Unsupported CUDA header."));
return;
}
@@ -293,11 +292,8 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc,
if (!exclude) {
clang::SmallString<128> includeBuffer;
llvm::StringRef name = Statistics::isToRoc(found->second) ? found->second.rocName : found->second.hipName;
- if (is_angled) {
- newInclude = llvm::Twine("<" + name+ ">").toStringRef(includeBuffer);
- } else {
- newInclude = llvm::Twine("\"" + name + "\"").toStringRef(includeBuffer);
- }
+ if (is_angled) newInclude = llvm::Twine("<" + name+ ">").toStringRef(includeBuffer);
+ else newInclude = llvm::Twine("\"" + name + "\"").toStringRef(includeBuffer);
} else {
// hashLoc is location of the '#', thus replacing the whole include directive by empty newInclude starting with '#'.
sl = hash_loc;
@@ -309,46 +305,33 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc,
}
void HipifyAction::PragmaDirective(clang::SourceLocation Loc, clang::PragmaIntroducerKind Introducer) {
- if (pragmaOnce) {
- return;
- }
- clang::SourceManager& SM = getCompilerInstance().getSourceManager();
- if (!SM.isWrittenInMainFile(Loc)) {
- return;
- }
- clang::Preprocessor& PP = getCompilerInstance().getPreprocessor();
+ if (pragmaOnce) return;
+ auto &SM = getCompilerInstance().getSourceManager();
+ if (!SM.isWrittenInMainFile(Loc)) return;
+ clang::Preprocessor &PP = getCompilerInstance().getPreprocessor();
clang::Token tok;
PP.Lex(tok);
StringRef Text(SM.getCharacterData(tok.getLocation()), tok.getLength());
- if (Text == "once") {
+ if (Text == sOnce) {
pragmaOnce = true;
pragmaOnceLoc = tok.getEndLoc();
}
}
-bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
- StringRef refName = "cudaLaunchKernel";
- const auto* launchKernel = Result.Nodes.getNodeAs(refName);
- if (!launchKernel) {
- return false;
- }
- const clang::Expr* calleeExpr = launchKernel->getCallee();
- if (!calleeExpr) {
- return false;
- }
- const clang::FunctionDecl *caleeDecl = launchKernel->getDirectCallee();
- if (!caleeDecl) {
- return false;
- }
- const clang::CallExpr* config = launchKernel->getConfig();
- if (!config) {
- return false;
- }
+bool HipifyAction::cudaLaunchKernel(const mat::MatchFinder::MatchResult &Result) {
+ auto *launchKernel = Result.Nodes.getNodeAs(sCudaLaunchKernel);
+ if (!launchKernel) return false;
+ auto *calleeExpr = launchKernel->getCallee();
+ if (!calleeExpr) return false;
+ auto *caleeDecl = launchKernel->getDirectCallee();
+ if (!caleeDecl) return false;
+ auto *config = launchKernel->getConfig();
+ if (!config) return false;
clang::SmallString<40> XStr;
llvm::raw_svector_ostream OS(XStr);
clang::LangOptions DefaultLangOptions;
- clang::SourceManager* SM = Result.SourceManager;
- OS << sHipLaunchKernelGGL;
+ auto *SM = Result.SourceManager;
+ OS << sHipLaunchKernelGGL << "(";
if (caleeDecl->isTemplateInstantiation()) OS << sHIP_KERNEL_NAME << "(";
OS << readSourceText(*SM, calleeExpr->getSourceRange());
if (caleeDecl->isTemplateInstantiation()) OS << ")";
@@ -381,29 +364,24 @@ bool HipifyAction::cudaLaunchKernel(const clang::ast_matchers::MatchFinder::Matc
ct::Replacement Rep(*SM, launchStart, length, OS.str());
clang::FullSourceLoc fullSL(launchStart, *SM);
insertReplacement(Rep, fullSL);
- hipCounter counter = {"hipLaunchKernelGGL", "", ConvTypes::CONV_KERNEL_LAUNCH, ApiTypes::API_RUNTIME};
- Statistics::current().incrementCounter(counter, refName.str());
+ hipCounter counter = {sHipLaunchKernelGGL, "", ConvTypes::CONV_KERNEL_LAUNCH, ApiTypes::API_RUNTIME};
+ Statistics::current().incrementCounter(counter, sCudaLaunchKernel.str());
return true;
}
-bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
- StringRef refName = "cudaSharedIncompleteArrayVar";
- auto* sharedVar = Result.Nodes.getNodeAs(refName);
- if (!sharedVar) {
- return false;
- }
+bool HipifyAction::cudaSharedIncompleteArrayVar(const mat::MatchFinder::MatchResult &Result) {
+ auto *sharedVar = Result.Nodes.getNodeAs(sCudaSharedIncompleteArrayVar);
+ if (!sharedVar) return false;
// Example: extern __shared__ uint sRadix1[];
- if (!sharedVar->hasExternalFormalLinkage()) {
- return false;
- }
+ if (!sharedVar->hasExternalFormalLinkage()) return false;
clang::QualType QT = sharedVar->getType();
std::string typeName;
if (QT->isIncompleteArrayType()) {
- const clang::ArrayType* AT = QT.getTypePtr()->getAsArrayTypeUnsafe();
+ const clang::ArrayType *AT = QT.getTypePtr()->getAsArrayTypeUnsafe();
QT = AT->getElementType();
if (QT.getTypePtr()->isBuiltinType()) {
QT = QT.getCanonicalType();
- const auto* BT = clang::dyn_cast(QT);
+ auto *BT = clang::dyn_cast(QT);
if (BT) {
clang::LangOptions LO;
LO.CUDA = true;
@@ -417,7 +395,7 @@ bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::Match
if (!typeName.empty()) {
clang::SourceLocation slStart = sharedVar->getOuterLocStart();
clang::SourceLocation slEnd = llcompat::getEndLoc(sharedVar->getTypeSourceInfo()->getTypeLoc());
- clang::SourceManager* SM = Result.SourceManager;
+ auto *SM = Result.SourceManager;
size_t repLength = SM->getCharacterData(slEnd) - SM->getCharacterData(slStart) + 1;
std::string varName = sharedVar->getNameAsString();
std::string repName = sHIP_DYNAMIC_SHARED + "(" + typeName + ", " + varName + ")";
@@ -425,33 +403,57 @@ bool HipifyAction::cudaSharedIncompleteArrayVar(const clang::ast_matchers::Match
clang::FullSourceLoc fullSL(slStart, *SM);
insertReplacement(Rep, fullSL);
hipCounter counter = {sHIP_DYNAMIC_SHARED, "", ConvTypes::CONV_EXTERN_SHARED, ApiTypes::API_RUNTIME};
- Statistics::current().incrementCounter(counter, refName.str());
+ Statistics::current().incrementCounter(counter, sCudaSharedIncompleteArrayVar.str());
return true;
}
return false;
}
-bool HipifyAction::cudaDeviceFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
- if (const clang::CallExpr *call = Result.Nodes.getNodeAs("cudaDeviceFuncCall")) {
- const clang::FunctionDecl *funcDcl = call->getDirectCallee();
- if (!funcDcl) {
- return false;
- }
+bool HipifyAction::cudaDeviceFuncCall(const mat::MatchFinder::MatchResult &Result) {
+ if (const clang::CallExpr *call = Result.Nodes.getNodeAs(sCudaDeviceFuncCall)) {
+ auto *funcDcl = call->getDirectCallee();
+ if (!funcDcl) return false;
FindAndReplace(funcDcl->getDeclName().getAsString(), llcompat::getBeginLoc(call), CUDA_DEVICE_FUNC_MAP, false);
return true;
}
return false;
}
-bool HipifyAction::cudaHostFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
- if (const clang::CallExpr * call = Result.Nodes.getNodeAs("cudaHostFuncCall")) {
- if (!call->getNumArgs()) {
- return false;
- }
- const clang::FunctionDecl* funcDcl = call->getDirectCallee();
- if (!funcDcl) {
- return false;
+bool HipifyAction::cubNamespacePrefix(const mat::MatchFinder::MatchResult &Result) {
+ if (auto *decl = Result.Nodes.getNodeAs(sCubNamespacePrefix)) {
+ clang::QualType QT = decl->getUnderlyingType();
+ auto *t = QT.getTypePtr();
+ if (!t) return false;
+ const clang::ElaboratedType *et = t->getAs();
+ if (!et) return false;
+ const clang::NestedNameSpecifier *nns = et->getQualifier();
+ if (!nns) return false;
+ const clang::NamespaceDecl *nsd = nns->getAsNamespace();
+ if (!nsd) return false;
+ const clang::TypeSourceInfo *si = decl->getTypeSourceInfo();
+ const clang::TypeLoc tloc = si->getTypeLoc();
+ const clang::SourceRange sr = tloc.getSourceRange();
+ clang::SourceLocation sl(sr.getBegin());
+ clang::SourceLocation end(sr.getEnd());
+ auto &SM = getCompilerInstance().getSourceManager();
+ size_t length = SM.getCharacterData(end) - SM.getCharacterData(sl);
+ StringRef sfull = StringRef(SM.getCharacterData(sl), length);
+ std::string name = nsd->getDeclName().getAsString();
+ size_t offset = sfull.find(name);
+ if (offset > 0) {
+ sl = sl.getLocWithOffset(offset);
}
+ FindAndReplace(name, sl, CUDA_CUB_TYPE_NAME_MAP);
+ return true;
+ }
+ return false;
+}
+
+bool HipifyAction::cudaHostFuncCall(const mat::MatchFinder::MatchResult &Result) {
+ if (auto *call = Result.Nodes.getNodeAs(sCudaHostFuncCall)) {
+ if (!call->getNumArgs()) return false;
+ auto *funcDcl = call->getDirectCallee();
+ if (!funcDcl) return false;
std::string sName = funcDcl->getDeclName().getAsString();
unsigned int argNum = 0;
bool b_reinterpret = (ReinterpretFunctions.find(sName) != ReinterpretFunctions.end()) ? true : false;
@@ -465,7 +467,7 @@ bool HipifyAction::cudaHostFuncCall(const clang::ast_matchers::MatchFinder::Matc
clang::SmallString<40> XStr;
llvm::raw_svector_ostream OS(XStr);
clang::SourceRange sr = call->getArg(argNum)->getSourceRange();
- clang::SourceManager* SM = Result.SourceManager;
+ auto *SM = Result.SourceManager;
OS << (b_reinterpret ? s_reinterpret_cast : sHIP_SYMBOL) << "(" << readSourceText(*SM, sr) << ")";
clang::SourceRange replacementRange = getWriteRange(*SM, { sr.getBegin(), sr.getEnd() });
clang::SourceLocation s = replacementRange.getBegin();
@@ -480,7 +482,7 @@ bool HipifyAction::cudaHostFuncCall(const clang::ast_matchers::MatchFinder::Matc
return false;
}
-void HipifyAction::insertReplacement(const ct::Replacement& rep, const clang::FullSourceLoc& fullSL) {
+void HipifyAction::insertReplacement(const ct::Replacement &rep, const clang::FullSourceLoc &fullSL) {
llcompat::insertReplacement(*replacements, rep);
if (PrintStats) {
rep.getLength();
@@ -489,10 +491,10 @@ void HipifyAction::insertReplacement(const ct::Replacement& rep, const clang::Fu
}
}
-std::unique_ptr HipifyAction::CreateASTConsumer(clang::CompilerInstance& CI, llvm::StringRef) {
- Finder.reset(new clang::ast_matchers::MatchFinder);
+std::unique_ptr HipifyAction::CreateASTConsumer(clang::CompilerInstance &CI, StringRef) {
+ Finder.reset(new mat::MatchFinder);
// Replace the <<<...>>> language extension with a hip kernel launch
- Finder->addMatcher(mat::cudaKernelCallExpr(mat::isExpansionInMainFile()).bind("cudaLaunchKernel"), this);
+ Finder->addMatcher(mat::cudaKernelCallExpr(mat::isExpansionInMainFile()).bind(sCudaLaunchKernel), this);
Finder->addMatcher(
mat::varDecl(
mat::isExpansionInMainFile(),
@@ -500,7 +502,7 @@ std::unique_ptr HipifyAction::CreateASTConsumer(clang::Compi
mat::hasAttr(clang::attr::CUDAShared),
mat::hasType(mat::incompleteArrayType())
)
- ).bind("cudaSharedIncompleteArrayVar"),
+ ).bind(sCudaSharedIncompleteArrayVar),
this
);
Finder->addMatcher(
@@ -520,7 +522,7 @@ std::unique_ptr HipifyAction::CreateASTConsumer(clang::Compi
)
)
)
- ).bind("cudaHostFuncCall"),
+ ).bind(sCudaHostFuncCall),
this
);
Finder->addMatcher(
@@ -535,7 +537,22 @@ std::unique_ptr HipifyAction::CreateASTConsumer(clang::Compi
mat::unless(mat::hasAttr(clang::attr::CUDAHost))
)
)
- ).bind("cudaDeviceFuncCall"),
+ ).bind(sCudaDeviceFuncCall),
+ this
+ );
+ Finder->addMatcher(
+ mat::typedefDecl(
+ mat::isExpansionInMainFile(),
+ mat::hasType(
+ mat::elaboratedType(
+ mat::hasQualifier(
+ mat::specifiesNamespace(
+ mat::hasName(sCub)
+ )
+ )
+ )
+ )
+ ).bind(sCubNamespacePrefix),
this
);
// Ownership is transferred to the caller.
@@ -543,10 +560,8 @@ std::unique_ptr HipifyAction::CreateASTConsumer(clang::Compi
}
void HipifyAction::Ifndef(clang::SourceLocation Loc, const clang::Token &MacroNameTok, const clang::MacroDefinition &MD) {
- clang::SourceManager& SM = getCompilerInstance().getSourceManager();
- if (!SM.isWrittenInMainFile(Loc)) {
- return;
- }
+ auto &SM = getCompilerInstance().getSourceManager();
+ if (!SM.isWrittenInMainFile(Loc)) return;
StringRef Text(SM.getCharacterData(MacroNameTok.getLocation()), MacroNameTok.getLength());
Ifndefs.insert(std::make_pair(Text.str(), MacroNameTok.getEndLoc()));
}
@@ -559,12 +574,12 @@ void HipifyAction::EndSourceFileAction() {
// one copy of the hip include into every file.
bool placeForIncludeCalculated = false;
clang::SourceLocation sl, controllingMacroLoc;
- clang::SourceManager& SM = getCompilerInstance().getSourceManager();
- clang::Preprocessor& PP = getCompilerInstance().getPreprocessor();
- clang::HeaderSearch& HS = PP.getHeaderSearchInfo();
- clang::ExternalPreprocessorSource* EPL = HS.getExternalLookup();
- const clang::FileEntry* FE = SM.getFileEntryForID(SM.getMainFileID());
- const clang::IdentifierInfo* controllingMacro = HS.getFileInfo(FE).getControllingMacro(EPL);
+ auto &SM = getCompilerInstance().getSourceManager();
+ clang::Preprocessor &PP = getCompilerInstance().getPreprocessor();
+ clang::HeaderSearch &HS = PP.getHeaderSearchInfo();
+ clang::ExternalPreprocessorSource *EPL = HS.getExternalLookup();
+ const clang::FileEntry *FE = SM.getFileEntryForID(SM.getMainFileID());
+ const clang::IdentifierInfo *controllingMacro = HS.getFileInfo(FE).getControllingMacro(EPL);
if (controllingMacro) {
auto found = Ifndefs.find(controllingMacro->getName().str());
if (found != Ifndefs.end()) {
@@ -573,19 +588,13 @@ void HipifyAction::EndSourceFileAction() {
}
}
if (pragmaOnce) {
- if (placeForIncludeCalculated) {
- sl = pragmaOnceLoc < controllingMacroLoc ? pragmaOnceLoc : controllingMacroLoc;
- } else {
- sl = pragmaOnceLoc;
- }
+ if (placeForIncludeCalculated) sl = pragmaOnceLoc < controllingMacroLoc ? pragmaOnceLoc : controllingMacroLoc;
+ else sl = pragmaOnceLoc;
placeForIncludeCalculated = true;
}
if (!placeForIncludeCalculated) {
- if (firstHeader) {
- sl = firstHeaderLoc;
- } else {
- sl = SM.getLocForStartOfFile(SM.getMainFileID());
- }
+ if (firstHeader) sl = firstHeaderLoc;
+ else sl = SM.getLocForStartOfFile(SM.getMainFileID());
}
clang::FullSourceLoc fullSL(sl, SM);
ct::Replacement Rep(SM, sl, 0, "\n#include \n");
@@ -600,15 +609,15 @@ namespace {
* A silly little class to proxy PPCallbacks back to the HipifyAction class.
*/
class PPCallbackProxy : public clang::PPCallbacks {
- HipifyAction& hipifyAction;
+ HipifyAction &hipifyAction;
public:
- explicit PPCallbackProxy(HipifyAction& action): hipifyAction(action) {}
+ explicit PPCallbackProxy(HipifyAction &action): hipifyAction(action) {}
- void InclusionDirective(clang::SourceLocation hash_loc, const clang::Token& include_token,
+ void InclusionDirective(clang::SourceLocation hash_loc, const clang::Token &include_token,
StringRef file_name, bool is_angled, clang::CharSourceRange filename_range,
- const clang::FileEntry* file, StringRef search_path, StringRef relative_path,
- const clang::Module* imported
+ const clang::FileEntry *file, StringRef search_path, StringRef relative_path,
+ const clang::Module *imported
#if LLVM_VERSION_MAJOR > 6
, clang::SrcMgr::CharacteristicKind FileType
#endif
@@ -632,10 +641,10 @@ bool HipifyAction::BeginInvocation(clang::CompilerInstance &CI) {
}
void HipifyAction::ExecuteAction() {
- clang::Preprocessor& PP = getCompilerInstance().getPreprocessor();
- clang::SourceManager& SM = getCompilerInstance().getSourceManager();
+ clang::Preprocessor &PP = getCompilerInstance().getPreprocessor();
+ auto &SM = getCompilerInstance().getSourceManager();
// Start lexing the specified input file.
- const llvm::MemoryBuffer* FromFile = SM.getBuffer(SM.getMainFileID());
+ const llvm::MemoryBuffer *FromFile = SM.getBuffer(SM.getMainFileID());
clang::Lexer RawLex(SM.getMainFileID(), FromFile, SM, PP.getLangOpts());
RawLex.SetKeepWhitespaceMode(true);
// Perform a token-level rewrite of CUDA identifiers to hip ones. The raw-mode lexer gives us enough
@@ -653,9 +662,10 @@ void HipifyAction::ExecuteAction() {
clang::ASTFrontendAction::ExecuteAction();
}
-void HipifyAction::run(const clang::ast_matchers::MatchFinder::MatchResult& Result) {
+void HipifyAction::run(const mat::MatchFinder::MatchResult &Result) {
if (cudaLaunchKernel(Result)) return;
if (cudaSharedIncompleteArrayVar(Result)) return;
if (cudaHostFuncCall(Result)) return;
if (cudaDeviceFuncCall(Result)) return;
+ if (cubNamespacePrefix(Result)) return;
}
diff --git a/projects/hip/hipify-clang/src/HipifyAction.h b/projects/hip/hipify-clang/src/HipifyAction.h
index 3c85604ced..78532449ee 100644
--- a/projects/hip/hipify-clang/src/HipifyAction.h
+++ b/projects/hip/hipify-clang/src/HipifyAction.h
@@ -31,17 +31,18 @@ THE SOFTWARE.
#include "Statistics.h"
namespace ct = clang::tooling;
+namespace mat = clang::ast_matchers;
using namespace llvm;
/**
* A FrontendAction that hipifies CUDA programs.
*/
class HipifyAction : public clang::ASTFrontendAction,
- public clang::ast_matchers::MatchFinder::MatchCallback {
+ public mat::MatchFinder::MatchCallback {
private:
- ct::Replacements* replacements;
+ ct::Replacements *replacements;
std::map Ifndefs;
- std::unique_ptr Finder;
+ std::unique_ptr Finder;
// CUDA implicitly adds its runtime header. We rewrite explicitly-provided CUDA includes with equivalent
// ones, and track - using this flag - if the result led to us including the hip runtime header. If it did
// not, we insert it at the top of the file when we finish processing it.
@@ -67,11 +68,11 @@ public:
explicit HipifyAction(ct::Replacements *replacements): clang::ASTFrontendAction(),
replacements(replacements) {}
// MatchCallback listeners
- bool cudaBuiltin(const clang::ast_matchers::MatchFinder::MatchResult& Result);
- bool cudaLaunchKernel(const clang::ast_matchers::MatchFinder::MatchResult& Result);
- bool cudaSharedIncompleteArrayVar(const clang::ast_matchers::MatchFinder::MatchResult& Result);
- bool cudaDeviceFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result);
- bool cudaHostFuncCall(const clang::ast_matchers::MatchFinder::MatchResult& Result);
+ bool cudaLaunchKernel(const mat::MatchFinder::MatchResult &Result);
+ bool cudaSharedIncompleteArrayVar(const mat::MatchFinder::MatchResult &Result);
+ bool cudaDeviceFuncCall(const mat::MatchFinder::MatchResult &Result);
+ bool cudaHostFuncCall(const mat::MatchFinder::MatchResult &Result);
+ bool cubNamespacePrefix(const mat::MatchFinder::MatchResult &Result);
// Called by the preprocessor for each include directive during the non-raw lexing pass.
void InclusionDirective(clang::SourceLocation hash_loc,
const clang::Token &include_token,
@@ -90,7 +91,7 @@ public:
protected:
// Add a Replacement for the current file. These will all be applied after executing the FrontendAction.
- void insertReplacement(const ct::Replacement& rep, const clang::FullSourceLoc& fullSL);
+ void insertReplacement(const ct::Replacement &rep, const clang::FullSourceLoc &fullSL);
// FrontendAction entry point.
void ExecuteAction() override;
// Callback before starting processing a single input; used by hipify-clang for setting Preprocessor options.
@@ -98,8 +99,8 @@ protected:
// Called at the start of each new file to process.
void EndSourceFileAction() override;
// MatchCallback API entry point. Called by the AST visitor while searching the AST for things we registered an interest for.
- void run(const clang::ast_matchers::MatchFinder::MatchResult& Result) override;
- std::unique_ptr CreateASTConsumer(clang::CompilerInstance &CI, llvm::StringRef InFile) override;
- bool Exclude(const hipCounter & hipToken);
- void FindAndReplace(llvm::StringRef name, clang::SourceLocation sl, const std::map& repMap, bool bReplace = true);
+ void run(const mat::MatchFinder::MatchResult &Result) override;
+ std::unique_ptr CreateASTConsumer(clang::CompilerInstance &CI, StringRef InFile) override;
+ bool Exclude(const hipCounter &hipToken);
+ void FindAndReplace(StringRef name, clang::SourceLocation sl, const std::map &repMap, bool bReplace = true);
};
diff --git a/projects/hip/hipify-clang/src/Statistics.cpp b/projects/hip/hipify-clang/src/Statistics.cpp
index 9751763be3..70a75ac0cd 100644
--- a/projects/hip/hipify-clang/src/Statistics.cpp
+++ b/projects/hip/hipify-clang/src/Statistics.cpp
@@ -129,6 +129,7 @@ const char *apiTypes[NUM_API_TYPES] = {
"API_RAND",
"API_DNN",
"API_FFT",
+ "API_CUB",
"API_SPARSE",
"API_CAFFE2"
};
diff --git a/projects/hip/hipify-clang/src/Statistics.h b/projects/hip/hipify-clang/src/Statistics.h
index 051f680fb1..9b9889d0e5 100644
--- a/projects/hip/hipify-clang/src/Statistics.h
+++ b/projects/hip/hipify-clang/src/Statistics.h
@@ -134,6 +134,7 @@ enum ApiTypes {
API_DNN,
API_FFT,
API_SPARSE,
+ API_CUB,
API_CAFFE2,
API_LAST
};
diff --git a/projects/hip/include/hip/hcc_detail/library_types.h b/projects/hip/include/hip/hcc_detail/library_types.h
index 74bae81321..6fcd0dc3d1 100644
--- a/projects/hip/include/hip/hcc_detail/library_types.h
+++ b/projects/hip/include/hip/hcc_detail/library_types.h
@@ -32,11 +32,10 @@ typedef enum hipDataType {
HIP_C_64F = 5
} hipDataType;
-typedef enum libraryPropertyType {
- MAJOR_VERSION,
- MINOR_VERSION,
- PATCH_LEVEL
-} libraryPropertyType;
+typedef enum hipLibraryPropertyType {
+ HIP_LIBRARY_MAJOR_VERSION,
+ HIP_LIBRARY_MINOR_VERSION,
+ HIP_LIBRARY_PATCH_LEVEL
+} hipLibraryPropertyType;
#endif
-
diff --git a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h
index 02f83b329e..5684596e13 100644
--- a/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h
+++ b/projects/hip/include/hip/nvcc_detail/hip_runtime_api.h
@@ -59,6 +59,12 @@ typedef enum hipMemcpyKind {
#define HIP_C_32F CUDA_C_32F
#define HIP_C_64F CUDA_C_64F
+// hipLibraryPropertyType
+#define hipLibraryPropertyType libraryPropertyType
+#define HIP_LIBRARY_MAJOR_VERSION MAJOR_VERSION
+#define HIP_LIBRARY_MINOR_VERSION MINOR_VERSION
+#define HIP_LIBRARY_PATCH_LEVEL PATCH_LEVEL
+
// hipTextureAddressMode
#define hipTextureAddressMode cudaTextureAddressMode
#define hipAddressModeWrap cudaAddressModeWrap
diff --git a/projects/hip/tests/hipify-clang/lit.cfg b/projects/hip/tests/hipify-clang/lit.cfg
index 64f82e57fa..89e17249f0 100644
--- a/projects/hip/tests/hipify-clang/lit.cfg
+++ b/projects/hip/tests/hipify-clang/lit.cfg
@@ -12,18 +12,31 @@ import lit.util
site_cfg = lit_config.params.get('site_config', None)
lit_config.load_config(config, site_cfg)
-print(str("========================================"))
+config.excludes = ['cmdparser.hpp']
+config.excludes.append('spatial_batch_norm_op.h')
+config.excludes.append('common_cudnn.h')
+
+delimiter = "===============================================================";
+print(delimiter)
print("CUDA " + config.cuda_version + " - will be used for testing")
print("LLVM " + config.llvm_version + " - will be used for testing")
print(platform.machine() + " - Platform architecture")
print(platform.system() + " " + platform.release() + " - Platform OS")
print(str(config.pointer_size * 8) + " - hipify-clang binary bitness")
print(str(struct.calcsize("P") * 8) + " - python " + str(platform.python_version()) + " binary bitness")
-print(str("========================================"))
-
-config.excludes = ['cmdparser.hpp']
-config.excludes.append('spatial_batch_norm_op.h')
-config.excludes.append('common_cudnn.h')
+print(delimiter)
+warns = None
+if not config.cuda_dnn_root:
+ config.excludes.append('cudnn_convolution_forward.cu')
+ config.excludes.append('cudnn_softmax.cu')
+ print("WARN: cuDNN tests are excluded due to unset CUDA_DNN_ROOT_DIR")
+ warns = True
+if not config.cuda_cub_root:
+ config.excludes.append('cub_01.cu')
+ print("WARN: CUB tests are excluded due to unset CUDA_CUB_ROOT_DIR")
+ warns = True
+if warns:
+ print(delimiter)
if config.cuda_version_major == 7 and config.cuda_version_minor == 0:
config.excludes.append('headers_test_09.cu')
@@ -84,17 +97,32 @@ clang_arguments = "-v"
if sys.platform in ['win32']:
run_test_ext = ".bat"
hipify_path += "/" + config.build_type
+ # CUDA SDK ROOT
clang_arguments += " -isystem'%s'/common/inc"
else:
run_test_ext = ".sh"
+ # CUDA SDK ROOT
clang_arguments += " -isystem'%s'/samples/common/inc"
-clang_arguments += " -I'%s'/include"
if config.pointer_size == 8:
clang_arguments += " -D__LP64__"
-hipify_arguments = "--cuda-path='%s'"
+# cuDNN ROOT
+if config.cuda_dnn_root:
+ clang_arguments += " -I'%s'/include"
+# CUB ROOT
+if config.cuda_cub_root:
+ clang_arguments += " -I'%s'"
-config.substitutions.append(("%clang_args", clang_arguments % (config.cuda_sdk_root, config.cuda_dnn_root)))
-config.substitutions.append(("%hipify_args", hipify_arguments % (config.cuda_root)))
+if config.cuda_dnn_root and config.cuda_cub_root:
+ config.substitutions.append(("%clang_args", clang_arguments % (config.cuda_sdk_root, config.cuda_dnn_root, config.cuda_cub_root)))
+elif config.cuda_dnn_root:
+ config.substitutions.append(("%clang_args", clang_arguments % (config.cuda_sdk_root, config.cuda_dnn_root)))
+elif config.cuda_cub_root:
+ config.substitutions.append(("%clang_args", clang_arguments % (config.cuda_sdk_root, config.cuda_cub_root)))
+else:
+ config.substitutions.append(("%clang_args", clang_arguments % config.cuda_sdk_root))
+
+hipify_arguments = "--cuda-path='%s'"
+config.substitutions.append(("%hipify_args", hipify_arguments % config.cuda_root))
config.substitutions.append(("hipify", '"' + hipify_path + "/hipify-clang" + '"'))
config.substitutions.append(("%run_test", '"' + config.test_source_root + "/run_test" + run_test_ext + '"'))
diff --git a/projects/hip/tests/hipify-clang/lit.site.cfg.in b/projects/hip/tests/hipify-clang/lit.site.cfg.in
index 6ef4dc007a..3c17567903 100644
--- a/projects/hip/tests/hipify-clang/lit.site.cfg.in
+++ b/projects/hip/tests/hipify-clang/lit.site.cfg.in
@@ -8,6 +8,7 @@ config.llvm_tools_dir = "@LLVM_TOOLS_BINARY_DIR@"
config.obj_root = "@CMAKE_CURRENT_BINARY_DIR@"
config.cuda_root = "@CUDA_TOOLKIT_ROOT_DIR@"
config.cuda_dnn_root = "@CUDA_DNN_ROOT_DIR@"
+config.cuda_cub_root = "@CUDA_CUB_ROOT_DIR@"
config.cuda_version_major = int("@CUDA_VERSION_MAJOR@")
config.cuda_version_minor = int("@CUDA_VERSION_MINOR@")
config.cuda_version = "@CUDA_VERSION@"
diff --git a/projects/hip/tests/hipify-clang/unit_tests/libraries/CUB/cub_01.cu b/projects/hip/tests/hipify-clang/unit_tests/libraries/CUB/cub_01.cu
new file mode 100644
index 0000000000..4646015e74
--- /dev/null
+++ b/projects/hip/tests/hipify-clang/unit_tests/libraries/CUB/cub_01.cu
@@ -0,0 +1,60 @@
+// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
+// CHECK: #include
+#include
+// CHECK: #include
+#include
+// CHECK: #include
+#include
+
+#include
+
+// TODO:
+// using namespace cub;
+
+template
+__global__ void sort(const T* data_in, T* data_out){
+ // CHECK: typedef ::hipcub::BlockRadixSort BlockRadixSortT;
+ typedef ::cub::BlockRadixSort BlockRadixSortT;
+ __shared__ typename BlockRadixSortT::TempStorage tmp_sort;
+ double items[4];
+ int i0 = 4 * (blockIdx.x * blockDim.x + threadIdx.x);
+ for (int i = 0; i < 4; ++i){
+ items[i] = data_in[i0 + i];
+ }
+ BlockRadixSortT(tmp_sort).Sort(items);
+ for (int i = 0; i < 4; ++i){
+ data_out[i0 + i] = items[i];
+ }
+}
+
+int main(){
+ double* d_gpu = NULL;
+ double* result_gpu = NULL;
+ double* data_sorted = new double[4096];
+ // Allocate memory on the GPU
+ // CHECK: hipMalloc(&d_gpu, 4096 * sizeof(double));
+ cudaMalloc(&d_gpu, 4096 * sizeof(double));
+ // CHECK: hipMalloc(&result_gpu, 4096 * sizeof(double));
+ cudaMalloc(&result_gpu, 4096 * sizeof(double));
+ // CHECK: hiprandGenerator_t gen;
+ curandGenerator_t gen;
+ // Create generator
+ // CHECK: hiprandCreateGenerator(&gen, HIPRAND_RNG_PSEUDO_DEFAULT);
+ curandCreateGenerator(&gen, CURAND_RNG_PSEUDO_DEFAULT);
+ // Fill array with random numbers
+ // CHECK: hiprandGenerateNormalDouble(gen, d_gpu, 4096, 0.0, 1.0);
+ curandGenerateNormalDouble(gen, d_gpu, 4096, 0.0, 1.0);
+ // Destroy generator
+ // CHECK: hiprandDestroyGenerator(gen);
+ curandDestroyGenerator(gen);
+ // Sort data
+ // CHECK: hipLaunchKernelGGL(HIP_KERNEL_NAME(sort), dim3(1), dim3(1024), 0, 0, d_gpu, result_gpu);
+ sort<<<1, 1024>>>(d_gpu, result_gpu);
+ // CHECK: hipMemcpy(data_sorted, result_gpu, 4096 * sizeof(double), hipMemcpyDeviceToHost);
+ cudaMemcpy(data_sorted, result_gpu, 4096 * sizeof(double), cudaMemcpyDeviceToHost);
+ // Write the sorted data to standard out
+ for (int i = 0; i < 4096; ++i){
+ std::cout << data_sorted[i] << ", ";
+ }
+ std::cout << std::endl;
+}