AI/rocm-systems
2
0
Förgrening 0
Kod Ärenden Pull-förfrågningar Actions Packages Projekt Släpp Wiki Aktiviteter

Merge 'master' into 'amd-master'

Bläddra i källkod 
Change-Id: I49daa87c2f32c5b6a4591409dc18d50c9ffdfce8
This commit is contained in:
Jenkins
2019-10-16 05:10:43 -04:00
förälder 0402ffdfd3 dd665f446c
incheckning aad41d36ea
15 ändrade filer med 407 tillägg och 258 borttagningar
Download Patch File Download Diff File Expand all files Collapse all files
bin/hipify-perl
+1
Visa fil
@@ -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;
hipify-clang/CMakeLists.txt
+3 -7
Visa fil
@@ -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)
hipify-clang/README.md
+84 -71
Visa fil
@@ -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`
## <a name="build-and-install"></a> 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
```
### <a name="windows"></a >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
..
hipify-clang/src/CUDA2HIP.cpp
+2
Visa fil
@@ -67,6 +67,8 @@ const std::map <llvm::StringRef, hipCounter> 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}},
hipify-clang/src/CUDA2HIP.h
+2
Visa fil
@@ -67,6 +67,8 @@ extern const std::map<llvm::StringRef, hipCounter> CUDA_CAFFE2_TYPE_NAME_MAP;
extern const std::map<llvm::StringRef, hipCounter> CUDA_CAFFE2_FUNCTION_MAP;
// Maps the names of CUDA Device functions to the corresponding HIP functions
extern const std::map<llvm::StringRef, hipCounter> CUDA_DEVICE_FUNC_MAP;
// Maps the names of CUDA CUB API types to the corresponding HIP types
extern const std::map<llvm::StringRef, hipCounter> CUDA_CUB_TYPE_NAME_MAP;
/**
* The union of all the above maps, except includes.
hipify-clang/src/CUDA2HIP_CUB_API_types.cpp
+28
Visa fil
@@ -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<llvm::StringRef, hipCounter> CUDA_CUB_TYPE_NAME_MAP{
{"cub", {"hipcub", "", CONV_TYPE, API_CUB}},
};
hipify-clang/src/HipifyAction.cpp
+161 -151
Visa fil
@@ -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 void*>";
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<std::string> DeviceSymbolFunctions0 {
{sCudaMemcpyToSymbol},
@@ -78,7 +88,7 @@ std::set<std::string> 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<llvm::StringRef, hipCounter>& repMap, bool bReplace) {
const std::map<StringRef, hipCounter> &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<clang::CXXDefaultArgExpr>(arg)) {
return "0";
} else {
return readSourceText(SM, arg->getSourceRange());
}
std::string stringifyZeroDefaultedArg(clang::SourceManager &SM, const clang::Expr *arg) {
if (clang::isa<clang::CXXDefaultArgExpr>(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<clang::CUDAKernelCallExpr>(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<clang::CUDAKernelCallExpr>(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<clang::VarDecl>(refName);
if (!sharedVar) {
return false;
}
bool HipifyAction::cudaSharedIncompleteArrayVar(const mat::MatchFinder::MatchResult &Result) {
auto *sharedVar = Result.Nodes.getNodeAs<clang::VarDecl>(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<clang::BuiltinType>(QT);
auto *BT = clang::dyn_cast<clang::BuiltinType>(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<clang::CallExpr>("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<clang::CallExpr>(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<clang::CallExpr>("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<clang::TypedefNameDecl>(sCubNamespacePrefix)) {
clang::QualType QT = decl->getUnderlyingType();
auto *t = QT.getTypePtr();
if (!t) return false;
const clang::ElaboratedType *et = t->getAs<clang::ElaboratedType>();
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<clang::CallExpr>(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<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::CompilerInstance& CI, llvm::StringRef) {
Finder.reset(new clang::ast_matchers::MatchFinder);
std::unique_ptr<clang::ASTConsumer> 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<clang::ASTConsumer> 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<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::Compi
)
)
)
).bind("cudaHostFuncCall"),
).bind(sCudaHostFuncCall),
this
);
Finder->addMatcher(
@@ -535,7 +537,22 @@ std::unique_ptr<clang::ASTConsumer> 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<clang::ASTConsumer> 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 <hip/hip_runtime.h>\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;
}
hipify-clang/src/HipifyAction.h
+14 -13
Visa fil
@@ -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<std::string, clang::SourceLocation> Ifndefs;
std::unique_ptr<clang::ast_matchers::MatchFinder> Finder;
std::unique_ptr<mat::MatchFinder> 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<clang::ASTConsumer> CreateASTConsumer(clang::CompilerInstance &CI, llvm::StringRef InFile) override;
bool Exclude(const hipCounter & hipToken);
void FindAndReplace(llvm::StringRef name, clang::SourceLocation sl, const std::map<llvm::StringRef, hipCounter>& repMap, bool bReplace = true);
void run(const mat::MatchFinder::MatchResult &Result) override;
std::unique_ptr<clang::ASTConsumer> CreateASTConsumer(clang::CompilerInstance &CI, StringRef InFile) override;
bool Exclude(const hipCounter &hipToken);
void FindAndReplace(StringRef name, clang::SourceLocation sl, const std::map<StringRef, hipCounter> &repMap, bool bReplace = true);
};
hipify-clang/src/Statistics.cpp
+1
Visa fil
@@ -129,6 +129,7 @@ const char *apiTypes[NUM_API_TYPES] = {
"API_RAND",
"API_DNN",
"API_FFT",
"API_CUB",
"API_SPARSE",
"API_CAFFE2"
};
hipify-clang/src/Statistics.h
+1
Visa fil
@@ -134,6 +134,7 @@ enum ApiTypes {
API_DNN,
API_FFT,
API_SPARSE,
API_CUB,
API_CAFFE2,
API_LAST
};
include/hip/hcc_detail/library_types.h
+5 -6
Visa fil
@@ -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
include/hip/nvcc_detail/hip_runtime_api.h
+6
Visa fil
@@ -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
tests/hipify-clang/lit.cfg
+38 -10
Visa fil
@@ -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 + '"'))
tests/hipify-clang/lit.site.cfg.in
+1
Visa fil
@@ -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@"
tests/hipify-clang/unit_tests/libraries/CUB/cub_01.cu
+60
Visa fil
@@ -0,0 +1,60 @@
// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args
// CHECK: #include <hip/hip_runtime.h>
#include <iostream>
// CHECK: #include <hiprand.h>
#include <curand.h>
// CHECK: #include <hipcub/hipcub.hpp>
#include <cub/cub.cuh>
#include <iostream>
// TODO:
// using namespace cub;
template <typename T>
__global__ void sort(const T* data_in, T* data_out){
// CHECK: typedef ::hipcub::BlockRadixSort<T, 1024, 4> BlockRadixSortT;
typedef ::cub::BlockRadixSort<T, 1024, 4> 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;
}