From 600d5d7c06f622d6f895fe79fe950f8d0453b4a8 Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 23 Jan 2018 23:43:36 +0300 Subject: [PATCH] [HIPIFY][fix] CUDA and cuBLAS main headers correct handling --- hipify-clang/src/CUDA2HipMap.cpp | 8 +++--- hipify-clang/src/HipifyAction.cpp | 35 +++++++++++++++++++-------- hipify-clang/src/HipifyAction.h | 5 ++-- tests/hipify-clang/headers_test_01.cu | 4 ++- tests/hipify-clang/headers_test_02.cu | 10 +++++--- tests/hipify-clang/headers_test_06.cu | 8 ++++++ tests/hipify-clang/headers_test_07.cu | 8 ++++++ tests/hipify-clang/headers_test_08.cu | 14 +++++++++++ 8 files changed, 71 insertions(+), 21 deletions(-) create mode 100644 tests/hipify-clang/headers_test_06.cu create mode 100644 tests/hipify-clang/headers_test_07.cu create mode 100644 tests/hipify-clang/headers_test_08.cu diff --git a/hipify-clang/src/CUDA2HipMap.cpp b/hipify-clang/src/CUDA2HipMap.cpp index 1f9c6287ed..6bcc0b38c4 100644 --- a/hipify-clang/src/CUDA2HipMap.cpp +++ b/hipify-clang/src/CUDA2HipMap.cpp @@ -379,12 +379,12 @@ const std::map CUDA_INCLUDE_MAP{ {"vector_types.h", {"hip/hip_vector_types.h", CONV_INCLUDE, API_RUNTIME}}, // CUBLAS includes - {"cublas.h", {"hipblas.h", CONV_INCLUDE, API_BLAS}}, - {"cublas_v2.h", {"hipblas.h", CONV_INCLUDE, API_BLAS}}, + {"cublas.h", {"hipblas.h", CONV_INCLUDE_CUDA_MAIN_H, API_BLAS}}, + {"cublas_v2.h", {"hipblas.h", CONV_INCLUDE_CUDA_MAIN_H, API_BLAS}}, // CURAND includes - {"curand.h", {"hiprand.h", CONV_INCLUDE, API_RAND}}, - {"curand_kernel.h", {"hiprand_kernel.h", CONV_INCLUDE, API_RAND}}, + {"curand.h", {"hiprand.h", CONV_INCLUDE, API_RAND}}, + {"curand_kernel.h", {"hiprand_kernel.h", CONV_INCLUDE, API_RAND}}, // HIP includes // TODO: uncomment this when hip/cudacommon.h will be renamed to hip/hipcommon.h diff --git a/hipify-clang/src/HipifyAction.cpp b/hipify-clang/src/HipifyAction.cpp index 7cd5b3d402..9ac0eacbd1 100644 --- a/hipify-clang/src/HipifyAction.cpp +++ b/hipify-clang/src/HipifyAction.cpp @@ -149,23 +149,38 @@ void HipifyAction::InclusionDirective(clang::SourceLocation hash_loc, 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()) { - if (!firstNotMainHeader) { - firstNotMainHeader = true; - firstNotMainHeaderLoc = hash_loc; - } return; } // Special-casing to avoid duplication of the hip_runtime include. bool secondMainInclude = false; - if (found->second.hipName == "hip/hip_runtime.h") { - if (insertedRuntimeHeader) { - secondMainInclude = true; + if (found->second.countType == CONV_INCLUDE_CUDA_MAIN_H) { + switch (found->second.countApiType) { + case API_DRIVER: + case API_RUNTIME: + if (insertedRuntimeHeader) { + secondMainInclude = true; + break; + } + insertedRuntimeHeader = true; + break; + case API_BLAS: + if (insertedBLASHeader) { + secondMainInclude = true; + break; + } + insertedBLASHeader = true; + break; + default: + break; } - insertedRuntimeHeader = true; } Statistics::current().incrementCounter(found->second, file_name.str()); @@ -356,8 +371,8 @@ void HipifyAction::EndSourceFileAction() { clang::SourceLocation sl; if (pragmaOnce) { sl = pragmaOnceLoc; - } else if (firstNotMainHeader) { - sl = firstNotMainHeaderLoc; + } else if (firstHeader) { + sl = firstHeaderLoc; } else { sl = SM.getLocForStartOfFile(SM.getMainFileID()); } diff --git a/hipify-clang/src/HipifyAction.h b/hipify-clang/src/HipifyAction.h index a269a37117..42622c1e01 100644 --- a/hipify-clang/src/HipifyAction.h +++ b/hipify-clang/src/HipifyAction.h @@ -23,9 +23,10 @@ private: // not, we insert it at the top of the file when we finish processing it. // This approach means we do the best it's possible to do w.r.t preserving the user's include order. bool insertedRuntimeHeader = false; - bool firstNotMainHeader = false; + bool insertedBLASHeader = false; + bool firstHeader = false; bool pragmaOnce = false; - clang::SourceLocation firstNotMainHeaderLoc; + clang::SourceLocation firstHeaderLoc; clang::SourceLocation pragmaOnceLoc; /** diff --git a/tests/hipify-clang/headers_test_01.cu b/tests/hipify-clang/headers_test_01.cu index c39ef80d8f..3747c339e8 100644 --- a/tests/hipify-clang/headers_test_01.cu +++ b/tests/hipify-clang/headers_test_01.cu @@ -1,6 +1,8 @@ // RUN: %run_test hipify "%s" "%t" %cuda_args // CHECK: #include +// CHECK-NOT: #include +// CHECK: #include #include -// CHECK-NOT: #include #include +#include diff --git a/tests/hipify-clang/headers_test_02.cu b/tests/hipify-clang/headers_test_02.cu index 90d412f797..57308efd59 100644 --- a/tests/hipify-clang/headers_test_02.cu +++ b/tests/hipify-clang/headers_test_02.cu @@ -1,6 +1,8 @@ // RUN: %run_test hipify "%s" "%t" %cuda_args -// CHECK: #include -#include -// CHECK-NOT: #include -#include +// CHECK: #include "hip/hip_runtime.h" +// CHECK-NOT: #include "cuda_runtime.h" +// CHECK: #include +#include "cuda.h" +#include "cuda_runtime.h" +#include diff --git a/tests/hipify-clang/headers_test_06.cu b/tests/hipify-clang/headers_test_06.cu new file mode 100644 index 0000000000..bce73c42df --- /dev/null +++ b/tests/hipify-clang/headers_test_06.cu @@ -0,0 +1,8 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +// CHECK: #include +// CHECK-NOT: #include +// CHECK: #include +#include +#include +#include diff --git a/tests/hipify-clang/headers_test_07.cu b/tests/hipify-clang/headers_test_07.cu new file mode 100644 index 0000000000..4237e1eb72 --- /dev/null +++ b/tests/hipify-clang/headers_test_07.cu @@ -0,0 +1,8 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +// CHECK: #include "hipblas.h" +// CHECK-NOT: #include "cublas.h" +// CHECK: #include +#include "cublas_v2.h" +#include "cublas.h" +#include diff --git a/tests/hipify-clang/headers_test_08.cu b/tests/hipify-clang/headers_test_08.cu new file mode 100644 index 0000000000..ad54871bd8 --- /dev/null +++ b/tests/hipify-clang/headers_test_08.cu @@ -0,0 +1,14 @@ +// RUN: %run_test hipify "%s" "%t" %cuda_args + +// CHECK: #include +// CHECK-NOT: #include +// CHECK: #include +// CHECK: #include "hipblas.h" +// CHECK-NOT: #include "cublas.h" +// CHECK: #include +#include +#include +#include +#include "cublas_v2.h" +#include "cublas.h" +#include