From ed265e5015f037d040755e1a8f78acb1f913e356 Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Fri, 9 Aug 2024 14:48:24 +0100 Subject: [PATCH 1/6] Header device_launch_parameters present in CUDA should be skipped --- bin/hipify-perl | 1 + src/CUDA2HIP.cpp | 1 + 2 files changed, 2 insertions(+) diff --git a/bin/hipify-perl b/bin/hipify-perl index ef5d68f5..b25964e0 100755 --- a/bin/hipify-perl +++ b/bin/hipify-perl @@ -5635,6 +5635,7 @@ sub simpleSubstitutions { subst("cublas.h", "hipblas.h", "include_cuda_main_header"); subst("cuda.h", "hip\/hip_runtime.h", "include_cuda_main_header"); subst("cuda_runtime.h", "hip\/hip_runtime.h", "include_cuda_main_header"); + subst("device_launch_parameters.h", "", "include"); subst("cudnn.h", "hipDNN.h", "include_cuda_main_header"); subst("cufft.h", "hipfft\/hipfft.h", "include_cuda_main_header"); subst("curand.h", "hiprand\/hiprand.h", "include_cuda_main_header"); diff --git a/src/CUDA2HIP.cpp b/src/CUDA2HIP.cpp index 3aae6ba3..11d7395b 100644 --- a/src/CUDA2HIP.cpp +++ b/src/CUDA2HIP.cpp @@ -27,6 +27,7 @@ const std::map CUDA_INCLUDE_MAP { // CUDA includes {"cuda.h", {"hip/hip_runtime.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_DRIVER, 0}}, {"cuda_runtime.h", {"hip/hip_runtime.h", "", CONV_INCLUDE_CUDA_MAIN_H, API_RUNTIME, 0}}, + {"device_launch_parameters.h", {"", "", CONV_INCLUDE, API_RUNTIME, 0}}, {"cuda_runtime_api.h", {"hip/hip_runtime_api.h", "", CONV_INCLUDE, API_RUNTIME, 0}}, {"channel_descriptor.h", {"hip/channel_descriptor.h", "", CONV_INCLUDE, API_RUNTIME, 0}}, {"device_functions.h", {"hip/device_functions.h", "", CONV_INCLUDE, API_RUNTIME, 0}}, From 201ffcd7079d86647410bbabf556356e0cc77c1a Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Fri, 9 Aug 2024 22:07:18 +0100 Subject: [PATCH 2/6] [HIPIFY][tests][doc] Python 3.12.5 is supported + [IMP] In case of errors, similar to `ModuleNotFoundError: No module named 'setuptools'`, upgrade the `setuptools` package: `python -m pip install --upgrade pip setuptools` --- docs/hipify-clang.rst | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/docs/hipify-clang.rst b/docs/hipify-clang.rst index bb96d631..b15b5b2b 100644 --- a/docs/hipify-clang.rst +++ b/docs/hipify-clang.rst @@ -646,7 +646,7 @@ Minimum build system requirements for the above configurations: Recommended build system requirements: -* CMake 3.30.2, GNU C/C++ 13.2, Python 3.12.4. +* CMake 3.30.2, GNU C/C++ 13.2, Python 3.12.5. Here's how to build ``hipify-clang`` with testing support on ``Ubuntu 23.10.01``: @@ -690,7 +690,7 @@ The corresponding successful output is: -- - Binary path : /usr/llvm/18.1.8/dist/bin -- Linker detection: GNU ld -- ---- The below configuring for hipify-clang testing only ---- - -- Found Python: /usr/bin/python3.12 (found version "3.12.4") found components: Interpreter + -- Found Python: /usr/bin/python3.12 (found version "3.12.5") found components: Interpreter -- Found lit: /usr/local/bin/lit -- Found FileCheck: /GIT/LLVM/trunk/dist/FileCheck -- Initial CUDA to configure: @@ -726,7 +726,7 @@ The corresponding successful output is: x86_64 - Platform architecture Linux 6.5.0-15-generic - Platform OS 64 - hipify-clang binary bitness - 64 - python 3.12.4 binary bitness + 64 - python 3.12.5 binary bitness =============================================================== -- Testing: 106 tests, 12 threads -- Testing Time: 6.91s @@ -821,13 +821,13 @@ Tested configurations: - ``8.0.5 - 9.3.0`` - ``2019.16.11.38, 2022.17.10.5`` - ``3.30.0`` - - ``3.12.4`` + - ``3.12.5`` * - ``19.0.0git`` - ``7.0 - 12.5.1`` - ``8.0.5 - 9.3.0`` - ``2019.16.11.38, 2022.17.10.5`` - ``3.30.0`` - - ``3.12.4`` + - ``3.12.5`` :sup:`5` LLVM 14.x.x is the latest major release supporting Visual Studio 2017. @@ -886,7 +886,7 @@ The corresponding successful output is: -- - LLVM Include path : D:/LLVM/18.1.8/dist/include -- - Binary path : D:/LLVM/18.1.8/dist/bin -- ---- The below configuring for hipify-clang testing only ---- - -- Found Python: C:/Users/TT/AppData/Local/Programs/Python/Python312/python.exe (found version "3.12.4") found components: Interpreter + -- Found Python: C:/Users/TT/AppData/Local/Programs/Python/Python312/python.exe (found version "3.12.5") found components: Interpreter -- Found lit: C:/Users/TT/AppData/Local/Programs/Python/Python312/Scripts/lit.exe -- Found FileCheck: D:/LLVM/18.1.8/dist/bin/FileCheck.exe -- Initial CUDA to configure: From 94d9fe6b654f1992007df8304a8bf8df3a36acce Mon Sep 17 00:00:00 2001 From: Jatin Chaudhary Date: Mon, 12 Aug 2024 13:16:30 +0100 Subject: [PATCH 3/6] Generate hipify-perl instead of manually editing it --- bin/hipify-perl | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/bin/hipify-perl b/bin/hipify-perl index b25964e0..ac1393cf 100755 --- a/bin/hipify-perl +++ b/bin/hipify-perl @@ -5625,6 +5625,7 @@ sub simpleSubstitutions { subst("curand_precalc.h", "hiprand\/hiprand_kernel.h", "include"); subst("curand_uniform.h", "hiprand\/hiprand_kernel.h", "include"); subst("device_functions.h", "hip\/device_functions.h", "include"); + subst("device_launch_parameters.h", "", "include"); subst("driver_types.h", "hip\/driver_types.h", "include"); subst("library_types.h", "hip\/library_types.h", "include"); subst("math_constants.h", "hip\/hip_math_constants.h", "include"); @@ -5635,7 +5636,6 @@ sub simpleSubstitutions { subst("cublas.h", "hipblas.h", "include_cuda_main_header"); subst("cuda.h", "hip\/hip_runtime.h", "include_cuda_main_header"); subst("cuda_runtime.h", "hip\/hip_runtime.h", "include_cuda_main_header"); - subst("device_launch_parameters.h", "", "include"); subst("cudnn.h", "hipDNN.h", "include_cuda_main_header"); subst("cufft.h", "hipfft\/hipfft.h", "include_cuda_main_header"); subst("curand.h", "hiprand\/hiprand.h", "include_cuda_main_header"); From 353da49e96353c90977d604d382f18015e90259a Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Mon, 12 Aug 2024 21:37:07 +0100 Subject: [PATCH 4/6] [HIPIFY][SWDEV-475354][#1439][#1459][fix] Switched the `e_add_const_argument` matchers to using `getWriteRange` MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit [Synopsis] + The wrapping macro `checkErrors` on the function call `cudaMallocHost` leads to `Skipped some replacements` and the absence of any hipification. [IMP] + As for the particular `cudaMallocHost`, it has two versions: `C`-version with two arguments (the first argument is `void**`) and the templated `C++`-version with three arguments (the first argument is `T **`, where T is `template`, and the third argument is the default argument). + In the second case, clang reports three arguments (the third is empty) despite there is no explicit third argument in the function call. As a workaround solution, change the position of the adding argument in the matcher for `cudaMallocHost` from 2 to 3, which is treated as `add an argument at the end of the argument list`. [Solution] + Took into account possible macro expansions by starting to use the function getWriteRange for all `e_add_const_argument` matchers (including `cudaMallocHost`). + Provided the corresponding test `cudaMallocHost.cu` [ToDo] + Revise all the rest matchers and switch them to using `getWriteRange` AMAP. + Try to provide more detailed diagnostics for matchers (besides `Skipped some replacements`). --- src/HipifyAction.cpp | 4 ++- tests/unit_tests/samples/cudaMallocHost.cu | 41 ++++++++++++++++++++++ 2 files changed, 44 insertions(+), 1 deletion(-) create mode 100644 tests/unit_tests/samples/cudaMallocHost.cu diff --git a/src/HipifyAction.cpp b/src/HipifyAction.cpp index 889ac654..ff16b188 100644 --- a/src/HipifyAction.cpp +++ b/src/HipifyAction.cpp @@ -354,7 +354,7 @@ std::map> FuncArgCasts { { { { - {2, {e_add_const_argument, cw_None, "hipHostMallocDefault"}} + {3, {e_add_const_argument, cw_None, "hipHostMallocDefault"}} } } } @@ -2736,6 +2736,8 @@ bool HipifyAction::cudaHostFuncCall(const mat::MatchFinder::MatchResult &Result) OS << c.second.constValToAddOrReplace << ", "; else OS << ", " << c.second.constValToAddOrReplace; + clang::SourceRange replacementRange = getWriteRange(*Result.SourceManager, { s, s }); + s = replacementRange.getBegin(); break; } case e_add_var_argument: diff --git a/tests/unit_tests/samples/cudaMallocHost.cu b/tests/unit_tests/samples/cudaMallocHost.cu new file mode 100644 index 00000000..13f85ec8 --- /dev/null +++ b/tests/unit_tests/samples/cudaMallocHost.cu @@ -0,0 +1,41 @@ +// RUN: %run_test hipify "%s" "%t" %hipify_args %clang_args + +// CHECK: #include + +template +void check(T result, char const *const func, const char *const file, int const line) { + if (result) { + fprintf(stderr, "Error at %s:%d code=%d(%s) \" \n", file, line, static_cast(result), func); + exit(EXIT_FAILURE); + } +} + +#define checkErrors(val) check((val), #val, __FILE__, __LINE__) +#define num 1024 + +int main(int argc, const char *argv[]) { + int *input = nullptr; + int deviceCount = 0; + // CHECK: checkErrors(hipGetDeviceCount(&deviceCount)); + checkErrors(cudaGetDeviceCount(&deviceCount)); + printf("Device Count: %d\n", deviceCount); + // CHECK: hipDeviceProp_t deviceProp; + cudaDeviceProp deviceProp; + deviceProp.major = 0; + deviceProp.minor = 0; + int deviceID = 0; + // CHECK: checkErrors(hipGetDeviceProperties(&deviceProp, deviceID)); + checkErrors(cudaGetDeviceProperties(&deviceProp, deviceID)); + // CHECK: checkErrors(hipSetDevice(deviceID)); + checkErrors(cudaSetDevice(deviceID)); + // CHECK: checkErrors(hipHostMalloc(&input, sizeof(int) * num * 2, hipHostMallocDefault)); + checkErrors(cudaMallocHost(&input, sizeof(int) * num * 2)); + for (int i = 0; i < num * 2; ++i) { + input[i] = i; + } + // CHECK: checkErrors(hipHostFree(input)); + checkErrors(cudaFreeHost(input)); + // CHECK: checkErrors(hipDeviceSynchronize()); + checkErrors(cudaDeviceSynchronize()); + return 0; +} From 72911607a04ec1a7fd8a2bcb1129129f5454ebef Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Tue, 13 Aug 2024 22:39:34 +0100 Subject: [PATCH 5/6] [HIPIFY][SWDEV-475354][#1439][#1459][fix] Switched the rest of transforming matchers to using `getWriteRange` [ToDo] + Add more tests on different corner cases with all the transforming matchers + Try to provide more detailed diagnostics for matchers (besides `Skipped some replacements`) --- src/HipifyAction.cpp | 13 ++++++++----- .../transforming_matchers.cu} | 17 +++++++++++++++++ 2 files changed, 25 insertions(+), 5 deletions(-) rename tests/unit_tests/{samples/cudaMallocHost.cu => synthetic/transforming_matchers.cu} (67%) diff --git a/src/HipifyAction.cpp b/src/HipifyAction.cpp index ff16b188..9258bca9 100644 --- a/src/HipifyAction.cpp +++ b/src/HipifyAction.cpp @@ -2660,6 +2660,7 @@ bool HipifyAction::cudaHostFuncCall(const mat::MatchFinder::MatchResult &Result) auto it = FuncArgCasts.find(sName); if (it == FuncArgCasts.end()) return false; auto castStructs = it->second; + auto &SM = *Result.SourceManager; for (auto cc : castStructs) { if (cc.isToMIOpen != TranslateToMIOpen || cc.isToRoc != TranslateToRoc) continue; clang::LangOptions DefaultLangOptions; @@ -2668,17 +2669,17 @@ bool HipifyAction::cudaHostFuncCall(const mat::MatchFinder::MatchResult &Result) unsigned int argNum = c.first; clang::SmallString<40> XStr; llvm::raw_svector_ostream OS(XStr); - auto &SM = *Result.SourceManager; clang::SourceRange sr, replacementRange; clang::SourceLocation s, e; if (argNum < call->getNumArgs()) { sr = call->getArg(argNum)->getSourceRange(); replacementRange = getWriteRange(SM, { sr.getBegin(), sr.getEnd() }); - s = replacementRange.getBegin(); - e = replacementRange.getEnd(); } else { s = e = call->getEndLoc(); + replacementRange = getWriteRange(SM, { s, e }); } + s = replacementRange.getBegin(); + e = replacementRange.getEnd(); switch (c.second.castType) { case e_remove_argument: { @@ -2695,6 +2696,8 @@ bool HipifyAction::cudaHostFuncCall(const mat::MatchFinder::MatchResult &Result) s = prevComma->getLocation(); } } + replacementRange = getWriteRange(SM, { s, e }); + e = replacementRange.getEnd(); length = SM.getCharacterData(e) - SM.getCharacterData(s); break; } @@ -2727,6 +2730,8 @@ bool HipifyAction::cudaHostFuncCall(const mat::MatchFinder::MatchResult &Result) e = call->getArg(argNum + c.second.numberToMoveOrCopy)->getBeginLoc(); else e = call->getEndLoc(); + replacementRange = getWriteRange(SM, { s, e }); + e = replacementRange.getEnd(); length = SM.getCharacterData(e) - SM.getCharacterData(s); break; } @@ -2736,8 +2741,6 @@ bool HipifyAction::cudaHostFuncCall(const mat::MatchFinder::MatchResult &Result) OS << c.second.constValToAddOrReplace << ", "; else OS << ", " << c.second.constValToAddOrReplace; - clang::SourceRange replacementRange = getWriteRange(*Result.SourceManager, { s, s }); - s = replacementRange.getBegin(); break; } case e_add_var_argument: diff --git a/tests/unit_tests/samples/cudaMallocHost.cu b/tests/unit_tests/synthetic/transforming_matchers.cu similarity index 67% rename from tests/unit_tests/samples/cudaMallocHost.cu rename to tests/unit_tests/synthetic/transforming_matchers.cu index 13f85ec8..9ade8811 100644 --- a/tests/unit_tests/samples/cudaMallocHost.cu +++ b/tests/unit_tests/synthetic/transforming_matchers.cu @@ -15,6 +15,7 @@ void check(T result, char const *const func, const char *const file, int const l int main(int argc, const char *argv[]) { int *input = nullptr; + void *input_ptr = nullptr; int deviceCount = 0; // CHECK: checkErrors(hipGetDeviceCount(&deviceCount)); checkErrors(cudaGetDeviceCount(&deviceCount)); @@ -30,9 +31,25 @@ int main(int argc, const char *argv[]) { checkErrors(cudaSetDevice(deviceID)); // CHECK: checkErrors(hipHostMalloc(&input, sizeof(int) * num * 2, hipHostMallocDefault)); checkErrors(cudaMallocHost(&input, sizeof(int) * num * 2)); + // CHECK: checkErrors(hipHostMalloc(&input_ptr, sizeof(int) * num * 2, hipHostMallocDefault)); + checkErrors(cudaMallocHost(&input_ptr, sizeof(int) * num * 2)); for (int i = 0; i < num * 2; ++i) { input[i] = i; } + + int *value = 0; + int *value_2 = 0; + int iBlockSize = 0; + int iBlockSize_2 = 0; + size_t bytes = 0; + // CHECK: hipFunction_t function; + CUfunction function; + // CHECK: void* occupancyB2DSize; + CUoccupancyB2DSize occupancyB2DSize; + + // CHECK: checkErrors(hipModuleOccupancyMaxPotentialBlockSizeWithFlags(value, value_2, function, bytes, iBlockSize, iBlockSize_2)); + checkErrors(cuOccupancyMaxPotentialBlockSizeWithFlags(value, value_2, function, occupancyB2DSize, bytes, iBlockSize, iBlockSize_2)); + // CHECK: checkErrors(hipHostFree(input)); checkErrors(cudaFreeHost(input)); // CHECK: checkErrors(hipDeviceSynchronize()); From ae365a7dfc1d9471fd467af843dfefc6e3d10d9c Mon Sep 17 00:00:00 2001 From: Evgeny Mankov Date: Wed, 14 Aug 2024 14:35:12 +0100 Subject: [PATCH 6/6] [HIPIFY][doc][6.2.2] `CHANGELOG.md` update --- CHANGELOG.md | 1 + 1 file changed, 1 insertion(+) diff --git a/CHANGELOG.md b/CHANGELOG.md index fc0e50e2..7aa4a9b8 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -13,6 +13,7 @@ Documentation for HIPIFY is available at * Removed some post HIP 6.2 APIs from support * Added hipification support for HIP functions `hipSetValidDevices`, `hipMemcpy2DArrayToArray`, `hipMemcpyAtoA`, `hipMemcpyAtoD`, `hipMemcpyAtoA`, `hipMemcpyAtoHAsync`, and `hipMemcpyHtoAAsync` +* Fixed an issue with `Skipped some replacements` when hipification didn't occur at all ## HIPIFY for ROCm 6.2.1