[HIPIFY][fix] cudaMallocHost -> hipHostMalloc

+ Added the last argument `hipHostMallocDefault` as `unsigned int flags`, because there is no default value for it
emankov · Feb 9, 2024 · 2383fb9 · 2383fb9
1 parent 54f9d5c
commit 2383fb9
Show file tree

Hide file tree

Showing 3 changed files with 13 additions and 4 deletions.
diff --git 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
@@ -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
@@ -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);