diff --git a/projects/hip/src/hip_module.cpp b/projects/hip/src/hip_module.cpp
index a17a99fe17..2c268e1617 100644
--- a/projects/hip/src/hip_module.cpp
+++ b/projects/hip/src/hip_module.cpp
@@ -5,6 +5,26 @@
 #include "hcc_detail/trace_helper.h"
 #include <fstream>
 
+hsa_status_t FindRegions(hsa_region_t region, void *data){
+  hsa_region_segment_t segment_id;
+  hsa_region_get_info(region, HSA_REGION_INFO_SEGMENT, &segment_id);
+
+  if(segment_id != HSA_REGION_SEGMENT_GLOBAL){
+    return HSA_STATUS_SUCCESS;
+  }
+
+  hsa_region_global_flag_t flags;
+  hsa_region_get_info(region, HSA_REGION_INFO_GLOBAL_FLAGS, &flags);
+
+  hsa_region_t *reg = (hsa_region_t*)data;
+
+  if(flags & HSA_REGION_GLOBAL_FLAG_FINE_GRAINED){
+    *reg = region;
+  }
+
+  return HSA_STATUS_SUCCESS;
+}
+
 hipError_t hipModuleLoad(hipModule *module, const char *fname){
     HIP_INIT_API(fname);
     hipError_t ret = hipSuccess;
@@ -14,24 +34,25 @@ hipError_t hipModuleLoad(hipModule *module, const char *fname){
     }else{
         int deviceId = ctx->getDevice()->_deviceId;
         ihipDevice_t *currentDevice = ihipGetDevice(deviceId);
-        hc::accelerator acc = currentDevice->_acc;
         std::ifstream in(fname, std::ios::binary | std::ios::ate);
         if(!in){
             std::cout<<"Couldn't read file "<<fname<<std::endl;
         }else{
             size_t size = std::string::size_type(in.tellg());
             void *p = NULL;
-            hsa_amd_memory_pool_t *pool = (hsa_amd_memory_pool_t*)acc.get_hsa_am_system_region();
-            hsa_status_t status = hsa_amd_memory_pool_allocate(*pool, size, 0, (void**)&p);
-            assert(status = HSA_STATUS_SUCCESS);
+            hsa_agent_t agent = currentDevice->_hsaAgent;
+            hsa_region_t sysRegion;
+            hsa_status_t status = hsa_agent_iterate_regions(agent, FindRegions, &sysRegion);
+            assert(status == HSA_STATUS_SUCCESS);
+            status = hsa_memory_allocate(sysRegion, size, (void**)&p);
             char *ptr = (char*)p;
             if(!ptr){
                 std::cout<<"Error: failed to allocate memory for code object"<<std::endl;
             }
+            hsa_code_object_t obj;
             in.seekg(0, std::ios::beg);
             std::copy(std::istreambuf_iterator<char>(in),
                       std::istreambuf_iterator<char>(), ptr);
-            hsa_code_object_t obj;
             status = hsa_code_object_deserialize(ptr, size, NULL, &obj);
             *module = obj.handle;
             assert(status == HSA_STATUS_SUCCESS);
@@ -49,22 +70,21 @@ hipError_t hipModuleGetFunction(hipFunction *func, hipModule hmod, const char *n
     }else{
       int deviceId = ctx->getDevice()->_deviceId;
       ihipDevice_t *currentDevice = ihipGetDevice(deviceId);
-      hc::accelerator acc = currentDevice->_acc;
-      hsa_agent_t *gpuAgent = (hsa_agent_t*)acc.get_hsa_agent();
+      hsa_agent_t gpuAgent = (hsa_agent_t)currentDevice->_hsaAgent;
 
-      assert(gpuAgent != NULL);
       hsa_status_t status;
       hsa_executable_symbol_t kernel_symbol;
       hsa_executable_t executable;
       status = hsa_executable_create(HSA_PROFILE_FULL, HSA_EXECUTABLE_STATE_UNFROZEN, NULL, &executable);
+      
       assert(status == HSA_STATUS_SUCCESS);
       hsa_code_object_t obj;
       obj.handle = hmod;
-      status = hsa_executable_load_code_object(executable, *gpuAgent, obj, NULL);
+      status = hsa_executable_load_code_object(executable, gpuAgent, obj, NULL);
       assert(status == HSA_STATUS_SUCCESS);
       status = hsa_executable_freeze(executable, NULL);
       assert(status == HSA_STATUS_SUCCESS);
-      status = hsa_executable_get_symbol(executable, NULL, name, *gpuAgent, 0, &kernel_symbol);
+      status = hsa_executable_get_symbol(executable, NULL, name, gpuAgent, 0, &kernel_symbol);
       assert(status == HSA_STATUS_SUCCESS);
       status = hsa_executable_symbol_get_info(kernel_symbol,
                                    HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_OBJECT,