merge amd-staging into amd-mainline

Change-Id: I73d6f0da42b32668a92759b75a744fb9fce7e833

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
 

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");

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:

src/CUDA2HIP.cpp

@@ -27,6 +27,7 @@ const std::map <llvm::StringRef, hipCounter> 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}},

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"}}
         }
       }
     }
@@ -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;
           }

tests/unit_tests/synthetic/transforming_matchers.cu

@@ -0,0 +1,58 @@
+// 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;
+  void *input_ptr = 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));
+  // 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());
+  checkErrors(cudaDeviceSynchronize());
+  return 0;
+}