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;