From 8753616c9727ebf74f4141b93fdcb45e35cb0753 Mon Sep 17 00:00:00 2001
From: foreman <dl.swbuild@amd.com>
Date: Fri, 5 Apr 2019 11:58:25 -0400
Subject: [PATCH] P4 to Git Change 1766264 by cpaquot@cpaquot-ocl-lc-lnx on
 2019/04/05 11:38:18

	SWDEV-145570 - [HIP] Use a context with all devices in system for host register
	hipHostRegister and hipMemcpy 0x10 and 0x20 fail in mGPU systems because
	we only register the memory on the current device. But in HIP, the registering
	needs to happen on all devices.

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_context.cpp#17 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_internal.hpp#26 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_memory.cpp#50 edit
---
 api/hip/hip_context.cpp  |  8 ++++++++
 api/hip/hip_internal.hpp |  1 +
 api/hip/hip_memory.cpp   | 12 +++++++-----
 3 files changed, 16 insertions(+), 5 deletions(-)

diff --git a/api/hip/hip_context.cpp b/api/hip/hip_context.cpp
index 94699fa38d..02b57f0316 100644
--- a/api/hip/hip_context.cpp
+++ b/api/hip/hip_context.cpp
@@ -34,6 +34,7 @@ thread_local amd::Context* g_context = nullptr;
 thread_local std::stack<amd::Context*> g_ctxtStack;
 thread_local hipError_t g_lastError = hipSuccess;
 std::once_flag g_ihipInitialized;
+amd::Context* host_context = nullptr;
 
 std::map<amd::Context*, amd::HostQueue*> g_nullStreams;
 
@@ -56,6 +57,13 @@ void init() {
       g_devices.push_back(context);
     }
   }
+
+  host_context = new amd::Context(devices, amd::Context::Info());
+  if (!host_context) return;
+
+  if (host_context && CL_SUCCESS != host_context->create(nullptr)) {
+    host_context->release();
+  }
 }
 
 amd::Context* getCurrentContext() {
diff --git a/api/hip/hip_internal.hpp b/api/hip/hip_internal.hpp
index ceff284dda..f7266ba82f 100644
--- a/api/hip/hip_internal.hpp
+++ b/api/hip/hip_internal.hpp
@@ -65,6 +65,7 @@ namespace hip {
   extern std::once_flag g_ihipInitialized;
   extern thread_local amd::Context* g_context;
   extern thread_local hipError_t g_lastError;
+  extern amd::Context* host_context;
 
   extern void init();
 
diff --git a/api/hip/hip_memory.cpp b/api/hip/hip_memory.cpp
index 2c8c30b588..75f45414f0 100644
--- a/api/hip/hip_memory.cpp
+++ b/api/hip/hip_memory.cpp
@@ -62,11 +62,14 @@ hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
     return hipErrorInvalidValue;
   }
 
-  if (hip::getCurrentContext()->devices()[0]->info().maxMemAllocSize_ < sizeBytes) {
+  amd::Context* amdContext = ((flags & CL_MEM_SVM_FINE_GRAIN_BUFFER) != 0)?
+    hip::host_context : hip::getCurrentContext();
+
+  if (amdContext->devices()[0]->info().maxMemAllocSize_ < sizeBytes) {
     return hipErrorMemoryAllocation;
   }
 
-  *ptr = amd::SvmBuffer::malloc(*hip::getCurrentContext(), flags, sizeBytes, hip::getCurrentContext()->devices()[0]->info().memBaseAddrAlign_);
+  *ptr = amd::SvmBuffer::malloc(*amdContext, flags, sizeBytes, amdContext->devices()[0]->info().memBaseAddrAlign_);
   if (*ptr == nullptr) {
     return hipErrorOutOfMemory;
   }
@@ -518,8 +521,7 @@ hipError_t hipHostGetFlags(unsigned int* flagsPtr, void* hostPtr) {
 hipError_t hipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags) {
   HIP_INIT_API(hostPtr, sizeBytes, flags);
   if(hostPtr != nullptr) {
-    amd::Context *amdContext = hip::getCurrentContext();
-    amd::Memory* mem = new (*amdContext) amd::Buffer(*amdContext, CL_MEM_USE_HOST_PTR, sizeBytes);
+    amd::Memory* mem = new (*hip::host_context) amd::Buffer(*hip::host_context, CL_MEM_USE_HOST_PTR, sizeBytes);
 
     if (!mem->create(hostPtr)) {
       mem->release();
@@ -538,7 +540,7 @@ hipError_t hipHostUnregister(void* hostPtr) {
   if (amd::SvmBuffer::malloced(hostPtr)) {
     hip::syncStreams();
     hip::getNullStream()->finish();
-    amd::SvmBuffer::free(*hip::getCurrentContext(), hostPtr);
+    amd::SvmBuffer::free(*hip::host_context, hostPtr);
     HIP_RETURN(hipSuccess);
   } else {
     size_t offset = 0;