From 2383fb90c3c59567daab813aa4f1d9ef10bc4072 Mon Sep 17 00:00:00 2001
From: Evgeny Mankov <evgeny.mankov@gmail.com>
Date: Fri, 9 Feb 2024 15:43:41 +0000
Subject: [PATCH] [HIPIFY][fix] `cudaMallocHost` -> `hipHostMalloc`

+ Added the last argument `hipHostMallocDefault` as `unsigned int flags`, because there is no default value for it
---
 src/HipifyAction.cpp                            | 11 ++++++++++-
 tests/unit_tests/device/device_symbols.cu       |  4 ++--
 tests/unit_tests/synthetic/runtime_functions.cu |  2 +-
 3 files changed, 13 insertions(+), 4 deletions(-)

diff --git a/src/HipifyAction.cpp b/src/HipifyAction.cpp
index 9afa929d..305bd679 100644
--- a/src/HipifyAction.cpp
+++ b/src/HipifyAction.cpp
@@ -212,6 +212,7 @@ const std::string sCusparseSpMM_bufferSize = "cusparseSpMM_bufferSize";
 const std::string sCusparseSpSM_analysis = "cusparseSpSM_analysis";
 const std::string sCusparseSpSM_solve = "cusparseSpSM_solve";
 const std::string sCusparseXcsrgeam2Nnz = "cusparseXcsrgeam2Nnz";
+const std::string sCudaMallocHost = "cudaMallocHost";
 
 // CUDA_OVERLOADED
 const std::string sCudaEventCreate = "cudaEventCreate";
@@ -269,6 +270,13 @@ std::map<std::string, hipify::FuncOverloadsStruct> FuncOverloads {
 };
 
 std::map<std::string, ArgCastStruct> FuncArgCasts {
+  {sCudaMallocHost,
+    {
+      {
+        {2, {e_add_const_argument, cw_None, "hipHostMallocDefault"}}
+      }
+    }
+  },
   {sCudaMemcpyToSymbol,
     {
       {
@@ -2574,7 +2582,8 @@ std::unique_ptr<clang::ASTConsumer> HipifyAction::CreateASTConsumer(clang::Compi
             sCusparseSpMM_bufferSize,
             sCusparseSpSM_analysis,
             sCusparseSpSM_solve,
-            sCusparseXcsrgeam2Nnz
+            sCusparseXcsrgeam2Nnz,
+            sCudaMallocHost
           )
         )
       )
diff --git a/tests/unit_tests/device/device_symbols.cu b/tests/unit_tests/device/device_symbols.cu
index 4139478c..0edb313d 100644
--- a/tests/unit_tests/device/device_symbols.cu
+++ b/tests/unit_tests/device/device_symbols.cu
@@ -57,9 +57,9 @@ int main() {
     }
     // CHECK: hipMalloc((void**)&Ad, SIZE);
     cudaMalloc((void**)&Ad, SIZE);
-    // CHECK: hipHostMalloc((void**)&Am, SIZE);
+    // CHECK: hipHostMalloc((void**)&Am, SIZE, hipHostMallocDefault);
     cudaMallocHost((void**)&Am, SIZE);
-    // CHECK: hipHostMalloc((void**)&Cm, SIZE);
+    // CHECK: hipHostMalloc((void**)&Cm, SIZE, hipHostMallocDefault);
     cudaMallocHost((void**)&Cm, SIZE);
     for (int i = 0; i < NUM; ++i) {
         Am[i] = -1 * i;
diff --git a/tests/unit_tests/synthetic/runtime_functions.cu b/tests/unit_tests/synthetic/runtime_functions.cu
index fa90300b..4e9438da 100644
--- a/tests/unit_tests/synthetic/runtime_functions.cu
+++ b/tests/unit_tests/synthetic/runtime_functions.cu
@@ -496,7 +496,7 @@ int main() {
 
   // CUDA: extern __host__ cudaError_t CUDARTAPI cudaMallocHost(void **ptr, size_t size);
   // HIP: hipError_t hipHostMalloc(void** ptr, size_t size, unsigned int flags);
-  // CHECK: result = hipHostMalloc(&deviceptr, bytes);
+  // CHECK: result = hipHostMalloc(&deviceptr, bytes, hipHostMallocDefault);
   result = cudaMallocHost(&deviceptr, bytes);
 
   // CUDA: extern __host__ __cudart_builtin__ cudaError_t CUDARTAPI cudaMallocManaged(void **devPtr, size_t size, unsigned int flags = cudaMemAttachGlobal);