From e988e5e44898392ccc311cff93aa801be27a476a Mon Sep 17 00:00:00 2001
From: Julia Jiang <julia.jiang@amd.com>
Date: Tue, 30 Jul 2024 16:02:09 -0400
Subject: [PATCH] SWDEV-436608 - Un-deprecate hipHostAlloc()

Change-Id: I71393e6f536ba84b4e172acf54ba4f72350e2ae8
---
 hipamd/src/hip_internal.hpp |  1 +
 hipamd/src/hip_memory.cpp   | 95 +++++++++++++++++++++----------------
 2 files changed, 56 insertions(+), 40 deletions(-)

diff --git a/hipamd/src/hip_internal.hpp b/hipamd/src/hip_internal.hpp
index a83320c634..1e27dff5da 100644
--- a/hipamd/src/hip_internal.hpp
+++ b/hipamd/src/hip_internal.hpp
@@ -636,6 +636,7 @@ public:
   extern int ihipGetDevice();
 
   extern hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags);
+  extern hipError_t ihipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags);
   extern amd::Memory* getMemoryObject(const void* ptr, size_t& offset, size_t size = 0);
   extern amd::Memory* getMemoryObjectWithOffset(const void* ptr, const size_t size = 0);
   extern void getStreamPerThread(hipStream_t& stream);
diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp
index 90b01f43a3..95b9edd999 100644
--- a/hipamd/src/hip_memory.cpp
+++ b/hipamd/src/hip_memory.cpp
@@ -340,6 +340,56 @@ hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
   return hipSuccess;
 }
 
+// ================================================================================================
+hipError_t ihipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
+{
+  if (ptr == nullptr) {
+    return hipErrorInvalidValue;
+  }
+  if (sizeBytes == 0) {
+    *ptr = nullptr;
+    return hipSuccess;
+  }
+
+  *ptr = nullptr;
+  const unsigned int coherentFlags = hipHostMallocCoherent | hipHostMallocNonCoherent;
+
+  // can't have both Coherent and NonCoherent flags set at the same time
+  if ((flags & coherentFlags) == coherentFlags) {
+    LogPrintfError(
+        "Cannot have both coherent and non-coherent flags "
+        "at the same time, flags: %u coherent flags: %u",
+        flags, coherentFlags);
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+
+  unsigned int ihipFlags = CL_MEM_SVM_FINE_GRAIN_BUFFER;
+  if (flags == 0 ||
+      flags & (hipHostMallocCoherent | hipHostMallocMapped | hipHostMallocNumaUser) ||
+      (!(flags & hipHostMallocNonCoherent) && HIP_HOST_COHERENT)) {
+    ihipFlags |= CL_MEM_SVM_ATOMICS;
+  }
+
+  if (flags & hipHostMallocNumaUser) {
+    ihipFlags |= CL_MEM_FOLLOW_USER_NUMA_POLICY;
+  }
+
+  if (flags & hipHostMallocNonCoherent) {
+    ihipFlags &= ~CL_MEM_SVM_ATOMICS;
+  }
+
+  hipError_t status = ihipMalloc(ptr, sizeBytes, ihipFlags);
+
+  if ((status == hipSuccess) && ((*ptr) != nullptr)) {
+    size_t offset = 0; // This is ignored
+    amd::Memory* svmMem = getMemoryObject(*ptr, offset);
+    // Save the HIP memory flags so that they can be accessed later
+    svmMem->getUserData().flags = flags;
+  }
+
+  HIP_RETURN(status, *ptr);
+}
+
 // ================================================================================================
 bool IsHtoHMemcpyValid(void* dst, const void* src, hipMemcpyKind kind) {
   size_t sOffset = 0;
@@ -625,44 +675,7 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
   if (ptr == nullptr) {
     HIP_RETURN(hipErrorInvalidValue);
   }
-  *ptr = nullptr;
-
-  const unsigned int coherentFlags = hipHostMallocCoherent | hipHostMallocNonCoherent;
-
-  // can't have both Coherent and NonCoherent flags set at the same time
-  if ((flags & coherentFlags) == coherentFlags) {
-    LogPrintfError(
-        "Cannot have both coherent and non-coherent flags "
-        "at the same time, flags: %u coherent flags: %u",
-        flags, coherentFlags);
-    HIP_RETURN(hipErrorInvalidValue);
-  }
-
-  unsigned int ihipFlags = CL_MEM_SVM_FINE_GRAIN_BUFFER;
-  if (flags == 0 ||
-      flags & (hipHostMallocCoherent | hipHostMallocMapped | hipHostMallocNumaUser) ||
-      (!(flags & hipHostMallocNonCoherent) && HIP_HOST_COHERENT)) {
-    ihipFlags |= CL_MEM_SVM_ATOMICS;
-  }
-
-  if (flags & hipHostMallocNumaUser) {
-    ihipFlags |= CL_MEM_FOLLOW_USER_NUMA_POLICY;
-  }
-
-  if (flags & hipHostMallocNonCoherent) {
-    ihipFlags &= ~CL_MEM_SVM_ATOMICS;
-  }
-
-  hipError_t status = ihipMalloc(ptr, sizeBytes, ihipFlags);
-
-  if ((status == hipSuccess) && ((*ptr) != nullptr)) {
-    size_t offset = 0; // This is ignored
-    amd::Memory* svmMem = getMemoryObject(*ptr, offset);
-    // Save the HIP memory flags so that they can be accessed later
-    svmMem->getUserData().flags = flags;
-  }
-
-  HIP_RETURN_DURATION(status, *ptr);
+  HIP_RETURN(ihipHostMalloc(ptr, sizeBytes, flags));
 }
 
 hipError_t hipFree(void* ptr) {
@@ -1297,11 +1310,13 @@ hipError_t hipHostUnregister(void* hostPtr) {
   HIP_RETURN(ihipHostUnregister(hostPtr));
 }
 
-// Deprecated function:
 hipError_t hipHostAlloc(void** ptr, size_t sizeBytes, unsigned int flags) {
   HIP_INIT_API(hipHostAlloc, ptr, sizeBytes, flags);
   CHECK_STREAM_CAPTURE_SUPPORTED();
-  HIP_RETURN(ihipMalloc(ptr, sizeBytes, flags), (ptr != nullptr)? *ptr : nullptr);
+  if (ptr == nullptr) {
+    HIP_RETURN(hipErrorInvalidValue);
+  }
+  HIP_RETURN(ihipHostMalloc(ptr, sizeBytes, flags));
 };
 
 inline hipError_t ihipMemcpySymbol_validate(const void* symbol, size_t sizeBytes,