From 353da49e96353c90977d604d382f18015e90259a Mon Sep 17 00:00:00 2001
From: Evgeny Mankov <evgeny.mankov@gmail.com>
Date: Mon, 12 Aug 2024 21:37:07 +0100
Subject: [PATCH] [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<class T>`, 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<std::string, std::vector<ArgCastStruct>> 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 <hip/hip_runtime.h>
+
+template <typename T>
+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<unsigned int>(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;
+}