diff --git a/hipamd/src/hip_context.cpp b/hipamd/src/hip_context.cpp
index 515e181343..307f8452f7 100644
--- a/hipamd/src/hip_context.cpp
+++ b/hipamd/src/hip_context.cpp
@@ -29,7 +29,7 @@ std::vector<hip::Device*> g_devices;
 
 namespace hip {
 thread_local TlsAggregator tls;
-Device* host_device = nullptr;
+amd::Context* host_context = nullptr;
 
 //init() is only to be called from the HIP_INIT macro only once
 bool init() {
@@ -74,7 +74,7 @@ bool init() {
   if (CL_SUCCESS != hContext->create(nullptr)) {
     hContext->release();
   }
-  host_device = new Device(hContext, -1);
+  host_context = hContext;
 
   PlatformState::instance().init();
   return true;
@@ -113,7 +113,7 @@ amd::HostQueue* getNullStream(amd::Context& ctx) {
   }
   // If it's a pure SVM allocation with system memory access, then it shouldn't matter which device
   // runtime selects by default
-  if (hip::host_device->asContext() == &ctx) {
+  if (hip::host_context == &ctx) {
     // Return current...
     return getNullStream();
   }
diff --git a/hipamd/src/hip_device_runtime.cpp b/hipamd/src/hip_device_runtime.cpp
index def9bcb0ef..0288c885f4 100644
--- a/hipamd/src/hip_device_runtime.cpp
+++ b/hipamd/src/hip_device_runtime.cpp
@@ -602,7 +602,7 @@ hipError_t hipSetDeviceFlags ( unsigned int  flags ) {
   switch (scheduleFlag) {
     case hipDeviceScheduleAuto:
       // Current behavior is different from the spec, due to MT usage in runtime
-      if (hip::host_device->devices().size() >= std::thread::hardware_concurrency()) {
+      if (hip::host_context->devices().size() >= std::thread::hardware_concurrency()) {
         device->SetActiveWait(false);
         break;
       }
diff --git a/hipamd/src/hip_hmm.cpp b/hipamd/src/hip_hmm.cpp
index 00687f84e4..be4c6cb4fe 100644
--- a/hipamd/src/hip_hmm.cpp
+++ b/hipamd/src/hip_hmm.cpp
@@ -233,8 +233,8 @@ hipError_t ihipMallocManaged(void** ptr, size_t size, unsigned int align) {
     return hipSuccess;
   }
 
-  assert((hip::host_device->asContext()!= nullptr) && "Current host context must be valid");
-  amd::Context& ctx = *hip::host_device->asContext();
+  assert((hip::host_context != nullptr) && "Current host context must be valid");
+  amd::Context& ctx = *hip::host_context;
 
   const amd::Device& dev = *ctx.devices()[0];
 
diff --git a/hipamd/src/hip_internal.hpp b/hipamd/src/hip_internal.hpp
index 2c3f7fb919..a416c1e7fb 100644
--- a/hipamd/src/hip_internal.hpp
+++ b/hipamd/src/hip_internal.hpp
@@ -513,7 +513,7 @@ namespace hip {
   extern thread_local TlsAggregator tls;
 
   /// Device representing the host - for pinned memory
-  extern Device* host_device;
+  extern amd::Context* host_context;
 
   extern bool init();
 
diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp
index da7ed30a55..2fdb28b537 100644
--- a/hipamd/src/hip_memory.cpp
+++ b/hipamd/src/hip_memory.cpp
@@ -274,7 +274,7 @@ hipError_t ihipMalloc(void** ptr, size_t sizeBytes, unsigned int flags)
 
   bool useHostDevice = (flags & CL_MEM_SVM_FINE_GRAIN_BUFFER) != 0;
   amd::Context* curDevContext = hip::getCurrentDevice()->asContext();
-  amd::Context* amdContext = useHostDevice ? hip::host_device->asContext() : curDevContext;
+  amd::Context* amdContext = useHostDevice ? hip::host_context : curDevContext;
 
   if (amdContext == nullptr) {
     return hipErrorOutOfMemory;
@@ -1155,7 +1155,7 @@ hipError_t ihipHostRegister(void* hostPtr, size_t sizeBytes, unsigned int flags)
   if (hostPtr == nullptr || sizeBytes == 0 || flags > 15) {
     return hipErrorInvalidValue;
   } else {
-    amd::Memory* mem = new (*hip::host_device->asContext()) amd::Buffer(*hip::host_device->asContext(),
+    amd::Memory* mem = new (*hip::host_context) amd::Buffer(*hip::host_context,
                             CL_MEM_USE_HOST_PTR | CL_MEM_SVM_ATOMICS, sizeBytes);
 
     constexpr bool sysMemAlloc = false;