From 0331004a440f709d77e5d83e567319b0b8e7dc52 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Fri, 11 Oct 2019 15:18:08 +0300 Subject: [PATCH 1/8] [HIP][fix] Prefix libraryPropertyType to fix build of rocFFT and TensorFlow [ROCm/hip commit: 337b7ce06a120aa2140bc0479f9693c7162693aa] --- projects/hip/include/hip/hcc_detail/library_types.h | 11 +++++------ .../hip/include/hip/nvcc_detail/hip_runtime_api.h | 6 ++++++ 2 files changed, 11 insertions(+), 6 deletions(-) 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..3d17f1d867 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 +// hipDataType +#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 From a4b08bcfa35adb60b08d156e4cce881285151422 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Fri, 11 Oct 2019 15:20:58 +0300 Subject: [PATCH 2/8] [HIP] Fix typo in a comment [ROCm/hip commit: 28c23a7b1a70da536c5b0f08e5aa5a49a5a38bc5] --- projects/hip/include/hip/nvcc_detail/hip_runtime_api.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) 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 3d17f1d867..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,7 +59,7 @@ typedef enum hipMemcpyKind { #define HIP_C_32F CUDA_C_32F #define HIP_C_64F CUDA_C_64F -// hipDataType +// hipLibraryPropertyType #define hipLibraryPropertyType libraryPropertyType #define HIP_LIBRARY_MAJOR_VERSION MAJOR_VERSION #define HIP_LIBRARY_MINOR_VERSION MINOR_VERSION From 3deccebb070f35ac9e369937a3ad5e353f0049f7 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 14 Oct 2019 11:55:55 +0300 Subject: [PATCH 3/8] [HIPIFY][CUB] Initial support (hipify-clang only) + Add one matcher (will be more) + Update Maps and Statistics + Add cub_01.cu unit test + Update lit harness to support standalone CUB + Update README.md + Update hipify-perl (only CUB header is supported for now) [IMPORTANT] clang (and hipify-clang) works correctly only with official NVLabs version on GitHub. Compilation of CUB from official CUDA release has conflicts with THRUST. Thus, to compile CUB sources, option "-I" should be specified to the cloned CUB from NVLAB on GitHub. [ROCm/hip commit: 277d3b83698a906aeebdd7393ea142fa76f4087c] --- projects/hip/bin/hipify-perl | 1 + projects/hip/hipify-clang/README.md | 151 ++++++++++-------- projects/hip/hipify-clang/src/CUDA2HIP.cpp | 2 + projects/hip/hipify-clang/src/CUDA2HIP.h | 2 + .../hip/hipify-clang/src/HipifyAction.cpp | 57 +++++++ projects/hip/hipify-clang/src/HipifyAction.h | 1 + projects/hip/hipify-clang/src/Statistics.cpp | 1 + projects/hip/hipify-clang/src/Statistics.h | 1 + projects/hip/tests/hipify-clang/lit.cfg | 8 +- .../hip/tests/hipify-clang/lit.site.cfg.in | 1 + 10 files changed, 155 insertions(+), 70 deletions(-) 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/README.md b/projects/hip/hipify-clang/README.md index ab52c288c8..07466dbe62 100644 --- a/projects/hip/hipify-clang/README.md +++ b/projects/hip/hipify-clang/README.md @@ -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. @@ -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.12.3 - 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/HipifyAction.cpp b/projects/hip/hipify-clang/src/HipifyAction.cpp index 5c5663985a..c6d10f0cdc 100644 --- a/projects/hip/hipify-clang/src/HipifyAction.cpp +++ b/projects/hip/hipify-clang/src/HipifyAction.cpp @@ -443,6 +443,47 @@ bool HipifyAction::cudaDeviceFuncCall(const clang::ast_matchers::MatchFinder::Ma return false; } +bool HipifyAction::cubNamespacePrefix(const clang::ast_matchers::MatchFinder::MatchResult& Result) { + if (const clang::TypedefNameDecl *decl = Result.Nodes.getNodeAs("cubNamespacePrefix")) { + if (!decl) { + return false; + } + clang::QualType QT = decl->getUnderlyingType(); + const clang::Type* 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()); + clang::SourceManager& 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 clang::ast_matchers::MatchFinder::MatchResult& Result) { if (const clang::CallExpr * call = Result.Nodes.getNodeAs("cudaHostFuncCall")) { if (!call->getNumArgs()) { @@ -538,6 +579,21 @@ std::unique_ptr HipifyAction::CreateASTConsumer(clang::Compi ).bind("cudaDeviceFuncCall"), this ); + Finder->addMatcher( + mat::typedefDecl( + mat::isExpansionInMainFile(), + mat::hasType( + mat::elaboratedType( + mat::hasQualifier( + mat::specifiesNamespace( + mat::hasName("cub") + ) + ) + ) + ) + ).bind("cubNamespacePrefix"), + this + ); // Ownership is transferred to the caller. return Finder->newASTConsumer(); } @@ -658,4 +714,5 @@ void HipifyAction::run(const clang::ast_matchers::MatchFinder::MatchResult& Resu 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..a24404deee 100644 --- a/projects/hip/hipify-clang/src/HipifyAction.h +++ b/projects/hip/hipify-clang/src/HipifyAction.h @@ -72,6 +72,7 @@ public: 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 cubNamespacePrefix(const clang::ast_matchers::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, 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/tests/hipify-clang/lit.cfg b/projects/hip/tests/hipify-clang/lit.cfg index 64f82e57fa..1d092a4327 100644 --- a/projects/hip/tests/hipify-clang/lit.cfg +++ b/projects/hip/tests/hipify-clang/lit.cfg @@ -84,17 +84,23 @@ 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" +# cuDNN ROOT clang_arguments += " -I'%s'/include" if config.pointer_size == 8: clang_arguments += " -D__LP64__" +# CUB ROOT +clang_arguments += " -I'%s'" + hipify_arguments = "--cuda-path='%s'" -config.substitutions.append(("%clang_args", clang_arguments % (config.cuda_sdk_root, config.cuda_dnn_root))) +config.substitutions.append(("%clang_args", clang_arguments % (config.cuda_sdk_root, config.cuda_dnn_root, config.cuda_cub_root))) 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@" From d174cad0aa433c3083f130624c8164a2aac595aa Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 14 Oct 2019 12:03:20 +0300 Subject: [PATCH 4/8] [HIPIFY][CUB] Add missing unit test [ROCm/hip commit: 92fb5a3a94dd74c1ab6cb75bba286728693671be] --- .../unit_tests/libraries/CUB/cub_01.cu | 60 +++++++++++++++++++ 1 file changed, 60 insertions(+) create mode 100644 projects/hip/tests/hipify-clang/unit_tests/libraries/CUB/cub_01.cu 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; +} From b3f8782fb95bc82aa23132cd702eb21ea5eea385 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 14 Oct 2019 16:32:57 +0300 Subject: [PATCH 5/8] [HIPIFY][build fix] Add missing new source file [ROCm/hip commit: 2994fb342927cd494447293ba9fa25677d65e60e] --- .../src/CUDA2HIP_CUB_API_types.cpp | 28 +++++++++++++++++++ 1 file changed, 28 insertions(+) create mode 100644 projects/hip/hipify-clang/src/CUDA2HIP_CUB_API_types.cpp 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}}, +}; From 8456642b3d672d1c60ea9c2a0caf67adb788e11c Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 14 Oct 2019 18:08:29 +0300 Subject: [PATCH 6/8] [HIPIFY] HipifyAction code clean-up [ROCm/hip commit: c0f62157f917ef9aab4bd23798fe4135cf8a2929] --- .../hip/hipify-clang/src/HipifyAction.cpp | 301 ++++++++---------- projects/hip/hipify-clang/src/HipifyAction.h | 28 +- 2 files changed, 141 insertions(+), 188 deletions(-) diff --git a/projects/hip/hipify-clang/src/HipifyAction.cpp b/projects/hip/hipify-clang/src/HipifyAction.cpp index c6d10f0cdc..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,52 +403,39 @@ 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::cubNamespacePrefix(const clang::ast_matchers::MatchFinder::MatchResult& Result) { - if (const clang::TypedefNameDecl *decl = Result.Nodes.getNodeAs("cubNamespacePrefix")) { - if (!decl) { - return false; - } +bool HipifyAction::cubNamespacePrefix(const mat::MatchFinder::MatchResult &Result) { + if (auto *decl = Result.Nodes.getNodeAs(sCubNamespacePrefix)) { clang::QualType QT = decl->getUnderlyingType(); - const clang::Type* t = QT.getTypePtr(); - if (!t) { - return false; - } - const clang::ElaboratedType* et = t->getAs(); - if (!et) { - return false; - } + 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; - } + if (!nns) return false; const clang::NamespaceDecl *nsd = nns->getAsNamespace(); - if (!nsd) { - return false; - } + 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()); - clang::SourceManager& SM = getCompilerInstance().getSourceManager(); + 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(); @@ -484,15 +449,11 @@ bool HipifyAction::cubNamespacePrefix(const clang::ast_matchers::MatchFinder::Ma 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::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; @@ -506,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(); @@ -521,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(); @@ -530,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(), @@ -541,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( @@ -561,7 +522,7 @@ std::unique_ptr HipifyAction::CreateASTConsumer(clang::Compi ) ) ) - ).bind("cudaHostFuncCall"), + ).bind(sCudaHostFuncCall), this ); Finder->addMatcher( @@ -576,7 +537,7 @@ std::unique_ptr HipifyAction::CreateASTConsumer(clang::Compi mat::unless(mat::hasAttr(clang::attr::CUDAHost)) ) ) - ).bind("cudaDeviceFuncCall"), + ).bind(sCudaDeviceFuncCall), this ); Finder->addMatcher( @@ -586,12 +547,12 @@ std::unique_ptr HipifyAction::CreateASTConsumer(clang::Compi mat::elaboratedType( mat::hasQualifier( mat::specifiesNamespace( - mat::hasName("cub") + mat::hasName(sCub) ) ) ) ) - ).bind("cubNamespacePrefix"), + ).bind(sCubNamespacePrefix), this ); // Ownership is transferred to the caller. @@ -599,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())); } @@ -615,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()) { @@ -629,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"); @@ -656,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 @@ -688,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 @@ -709,7 +662,7 @@ 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; diff --git a/projects/hip/hipify-clang/src/HipifyAction.h b/projects/hip/hipify-clang/src/HipifyAction.h index a24404deee..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,12 +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 cubNamespacePrefix(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, @@ -91,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. @@ -99,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); }; From e49e3bfffad3955426341425e3da5aec573fbc21 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 15 Oct 2019 11:26:03 +0300 Subject: [PATCH 7/8] [HIPIFY][cmake] Make CMakeLists.txt compatible with default cmake 3.5.2 for Ubuntu 16.04 + Update README.md accordingly [ROCm/hip commit: d57c9a844bda3d18c477cd5e78c4673d54495faa] --- projects/hip/hipify-clang/CMakeLists.txt | 10 +++------- projects/hip/hipify-clang/README.md | 6 +++--- 2 files changed, 6 insertions(+), 10 deletions(-) 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 07466dbe62..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 @@ -196,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`: @@ -345,7 +345,7 @@ 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.5, Visual Studio 2017 (15.5.2) - 2019 (16.3.4). +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`: From 44ca50dfd06edc296c42fc96a95207bae7007f58 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 15 Oct 2019 14:20:23 +0300 Subject: [PATCH 8/8] [HIPIFY][tests] Exclude tests for the libs, which are not defined in cmake command line + affects cuDNN and CUB tests, paths to libraries of which are defined by CUDA_DNN_ROOT_DIR and CUDA_CUB_ROOT_DIR + Warn about excluding and why, for instance: "WARN: cuDNN tests are excluded due to unset CUDA_DNN_ROOT_DIR" [ROCm/hip commit: 3444834e3a465c750f2be76588fd4c90e964797e] --- projects/hip/tests/hipify-clang/lit.cfg | 46 ++++++++++++++++++------- 1 file changed, 34 insertions(+), 12 deletions(-) diff --git a/projects/hip/tests/hipify-clang/lit.cfg b/projects/hip/tests/hipify-clang/lit.cfg index 1d092a4327..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') @@ -90,17 +103,26 @@ else: run_test_ext = ".sh" # CUDA SDK ROOT clang_arguments += " -isystem'%s'/samples/common/inc" -# cuDNN ROOT -clang_arguments += " -I'%s'/include" if config.pointer_size == 8: clang_arguments += " -D__LP64__" +# cuDNN ROOT +if config.cuda_dnn_root: + clang_arguments += " -I'%s'/include" # CUB ROOT -clang_arguments += " -I'%s'" +if config.cuda_cub_root: + clang_arguments += " -I'%s'" + +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(("%clang_args", clang_arguments % (config.cuda_sdk_root, config.cuda_dnn_root, config.cuda_cub_root))) -config.substitutions.append(("%hipify_args", hipify_arguments % (config.cuda_root))) +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 + '"'))