From 0e4817fc7d09053ca7dc73a360c14220b6f2ab3d Mon Sep 17 00:00:00 2001
From: foreman <dl.constructicon@amd.com>
Date: Fri, 6 Apr 2018 18:08:18 -0400
Subject: [PATCH] P4 to Git Change 1537729 by cpaquot@cpaquot-ocl-lc-lnx on
 2018/04/06 17:59:29

	SWDEV-145570 - [HIP] Use Svm path for both hipMalloc and hipHostMalloc
	Make sure hipMemCpy uses SvmBuffer to fetch the cl_mem

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#8 edit


[ROCm/hip commit: e7f206d249c6b040027397f97bd3d09af21b432f]
---
 projects/hip/api/hip/hip_memory.cpp | 61 ++++++++++-------------------
 1 file changed, 21 insertions(+), 40 deletions(-)

diff --git a/projects/hip/api/hip/hip_memory.cpp b/projects/hip/api/hip/hip_memory.cpp
index f2dfcf4b4a..ef40bdbdf6 100644
--- a/projects/hip/api/hip/hip_memory.cpp
+++ b/projects/hip/api/hip/hip_memory.cpp
@@ -24,9 +24,8 @@ THE SOFTWARE.
 
 #include "hip_internal.hpp"
 
-hipError_t hipMalloc(void** ptr, size_t sizeBytes) {
-  HIP_INIT_API(ptr, sizeBytes);
-
+hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
+{
   if (sizeBytes == 0) {
     *ptr = nullptr;
     return hipSuccess;
@@ -39,36 +38,7 @@ hipError_t hipMalloc(void** ptr, size_t sizeBytes) {
     return hipErrorOutOfMemory;
   }
 
-  amd::Memory* mem = new (*g_context) amd::Buffer(*g_context, 0, sizeBytes);
-  if (!mem) {
-    return hipErrorOutOfMemory;
-  }
-
-  if (!mem->create(nullptr)) {
-    return hipErrorMemoryAllocation;
-  }
-
-  *ptr = reinterpret_cast<void*>(as_cl(mem));
-
-  return hipSuccess;
-}
-
-hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
-  HIP_INIT_API(ptr, sizeBytes, flags);
-
-  if (sizeBytes == 0) {
-    *ptr = nullptr;
-    return hipSuccess;
-  }
-  else if (!ptr) {
-    return hipErrorInvalidValue;
-  }
-
-  if (g_context->devices()[0]->info().maxMemAllocSize_ < sizeBytes) {
-    return hipErrorOutOfMemory;
-  }
-
-  *ptr = amd::SvmBuffer::malloc(*g_context, 0, sizeBytes, g_context->devices()[0]->info().memBaseAddrAlign_);
+  *ptr = amd::SvmBuffer::malloc(*g_context, flags, sizeBytes, g_context->devices()[0]->info().memBaseAddrAlign_);
   if (!*ptr) {
     return hipErrorOutOfMemory;
   }
@@ -76,16 +46,24 @@ hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
   return hipSuccess;
 }
 
+hipError_t hipMalloc(void** ptr, size_t sizeBytes) {
+  HIP_INIT_API(ptr, sizeBytes);
+
+  return ihipMalloc(ptr, sizeBytes, 0);
+}
+
+hipError_t hipHostMalloc(void** ptr, size_t sizeBytes, unsigned int flags) {
+  HIP_INIT_API(ptr, sizeBytes, flags);
+
+  return ihipMalloc(ptr, sizeBytes, CL_MEM_SVM_FINE_GRAIN_BUFFER);
+}
+
 hipError_t hipFree(void* ptr) {
   if (amd::SvmBuffer::malloced(ptr)) {
     amd::SvmBuffer::free(*g_context, ptr);
     return hipSuccess;
   }
-  if (!is_valid(reinterpret_cast<cl_mem>(ptr))) {
-    return hipErrorInvalidValue;
-  }
-  as_amd(reinterpret_cast<cl_mem>(ptr))->release();
-  return hipSuccess;
+  return hipErrorInvalidValue;
 }
 
 hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind kind) {
@@ -102,15 +80,18 @@ hipError_t hipMemcpy(void* dst, const void* src, size_t sizeBytes, hipMemcpyKind
 
   amd::Command* command;
   amd::Command::EventWaitList waitList;
+  amd::Memory* memory;
 
   switch (kind) {
   case hipMemcpyDeviceToHost:
+    memory = amd::SvmManager::FindSvmBuffer(src);
     command = new amd::ReadMemoryCommand(*queue, CL_COMMAND_READ_BUFFER, waitList,
-      *as_amd(reinterpret_cast<cl_mem>(const_cast<void*>(src)))->asBuffer(), 0, sizeBytes, dst);
+      *memory->asBuffer(), 0, sizeBytes, dst);
     break;
   case hipMemcpyHostToDevice:
+    memory = amd::SvmManager::FindSvmBuffer(dst);
     command = new amd::WriteMemoryCommand(*queue, CL_COMMAND_WRITE_BUFFER, waitList,
-      *as_amd(reinterpret_cast<cl_mem>(dst))->asBuffer(), 0, sizeBytes, src);
+      *memory->asBuffer(), 0, sizeBytes, src);
     break;
   default:
     assert(!"Shouldn't reach here");