diff --git a/projects/clr/rocclr/device/rocm/rocdevice.cpp b/projects/clr/rocclr/device/rocm/rocdevice.cpp
index b3c64d1700..71fae55468 100644
--- a/projects/clr/rocclr/device/rocm/rocdevice.cpp
+++ b/projects/clr/rocclr/device/rocm/rocdevice.cpp
@@ -2957,38 +2957,6 @@ void Device::getHwEventTime(const amd::Event& event, uint64_t* start, uint64_t*
   }
 }
 
-// ================================================================================================
-static void callbackQueue(hsa_status_t status, hsa_queue_t* queue, void* data) {
-  if (status != HSA_STATUS_SUCCESS && status != HSA_STATUS_INFO_BREAK) {
-    Device* dev = reinterpret_cast<Device*>(data);
-    for (auto it : dev->vgpus()) {
-      roc::VirtualGPU* vgpu = reinterpret_cast<roc::VirtualGPU*>(it);
-      if (vgpu->gpu_queue() == queue) {
-        vgpu->AnalyzeAqlQueue();
-      }
-    }
-    // Abort on device exceptions.
-    const char* errorMsg = 0;
-    hsa_status_string(status, &errorMsg);
-    if (status == HSA_STATUS_ERROR_OUT_OF_RESOURCES) {
-      size_t global_available_mem = 0;
-      if (HSA_STATUS_SUCCESS != hsa_agent_get_info(dev->getBackendDevice(),
-                         static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_MEMORY_AVAIL),
-                         &global_available_mem)) {
-        LogError("HSA_AMD_AGENT_INFO_MEMORY_AVAIL query failed.");
-      }
-      ClPrint(amd::LOG_NONE, amd::LOG_ALWAYS,
-              "Callback: Queue %p Aborting with error : %s Code: 0x%x Available Free mem : %zu MB",
-              queue->base_address, errorMsg, status, global_available_mem/Mi);
-    } else {
-      ClPrint(amd::LOG_NONE, amd::LOG_ALWAYS,
-        "Callback: Queue %p aborting with error : %s code: 0x%x", queue->base_address,
-        errorMsg, status);
-    }
-    abort();
-  }
-}
-
 // ================================================================================================
 hsa_queue_t* Device::getQueueFromPool(const uint qIndex) {
   // Check if queue with refCount 0 is available to use
@@ -3630,6 +3598,39 @@ ProfilingSignal::~ProfilingSignal() {
   }
 }
 
+// ================================================================================================
+void callbackQueue(hsa_status_t status, hsa_queue_t* queue, void* data) {
+  if (status != HSA_STATUS_SUCCESS && status != HSA_STATUS_INFO_BREAK) {
+    Device* dev = reinterpret_cast<Device*>(data);
+    for (auto it : dev->vgpus()) {
+      roc::VirtualGPU* vgpu = reinterpret_cast<roc::VirtualGPU*>(it);
+      if (vgpu->gpu_queue() == queue) {
+        vgpu->AnalyzeAqlQueue();
+      }
+    }
+    // Abort on device exceptions.
+    const char* errorMsg = 0;
+    hsa_status_string(status, &errorMsg);
+    if (status == HSA_STATUS_ERROR_OUT_OF_RESOURCES) {
+      size_t global_available_mem = 0;
+      if (HSA_STATUS_SUCCESS != hsa_agent_get_info(dev->getBackendDevice(),
+                         static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_MEMORY_AVAIL),
+                         &global_available_mem)) {
+        LogError("HSA_AMD_AGENT_INFO_MEMORY_AVAIL query failed.");
+      }
+      ClPrint(amd::LOG_NONE, amd::LOG_ALWAYS,
+              "Callback: Queue %p Aborting with error : %s Code: 0x%x Available Free mem : %zu MB",
+              queue->base_address, errorMsg, status, global_available_mem/Mi);
+    } else {
+      ClPrint(amd::LOG_NONE, amd::LOG_ALWAYS,
+        "Callback: Queue %p aborting with error : %s code: 0x%x", queue->base_address,
+        errorMsg, status);
+    }
+    abort();
+  }
+}
+
+// ================================================================================================
 #if defined(__clang__)
 #if __has_feature(address_sanitizer)
 device::UriLocator* Device::createUriLocator() const {
diff --git a/projects/clr/rocclr/device/rocm/rocdevice.hpp b/projects/clr/rocclr/device/rocm/rocdevice.hpp
index d360f90e52..db602bbb66 100644
--- a/projects/clr/rocclr/device/rocm/rocdevice.hpp
+++ b/projects/clr/rocclr/device/rocm/rocdevice.hpp
@@ -701,6 +701,9 @@ class Device : public NullDevice {
 #endif
 #endif
 };                                // class roc::Device
+
+void callbackQueue(hsa_status_t status, hsa_queue_t* queue, void* data);
+
 }  // namespace amd::roc
 
 /**
diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.cpp b/projects/clr/rocclr/device/rocm/rocvirtual.cpp
index 492ef6996c..938c4fcc5d 100644
--- a/projects/clr/rocclr/device/rocm/rocvirtual.cpp
+++ b/projects/clr/rocclr/device/rocm/rocvirtual.cpp
@@ -2979,33 +2979,6 @@ void VirtualGPU::submitMigrateMemObjects(amd::MigrateMemObjectsCommand& vcmd) {
   profilingEnd(vcmd);
 }
 
-// ================================================================================================
-static void callbackQueue(hsa_status_t status, hsa_queue_t* queue, void* data) {
-  if (status != HSA_STATUS_SUCCESS && status != HSA_STATUS_INFO_BREAK) {
-    VirtualGPU* vgpu = reinterpret_cast<VirtualGPU*>(data);
-    vgpu->AnalyzeAqlQueue();
-    // Abort on device exceptions.
-    const char* errorMsg = 0;
-    hsa_status_string(status, &errorMsg);
-    if (status == HSA_STATUS_ERROR_OUT_OF_RESOURCES) {
-      size_t global_available_mem = 0;
-      if (HSA_STATUS_SUCCESS != hsa_agent_get_info(vgpu->gpu_device(),
-                         static_cast<hsa_agent_info_t>(HSA_AMD_AGENT_INFO_MEMORY_AVAIL),
-                         &global_available_mem)) {
-        LogError("HSA_AMD_AGENT_INFO_MEMORY_AVAIL query failed.");
-      }
-      ClPrint(amd::LOG_NONE, amd::LOG_ALWAYS,
-              "Callback: Queue %p Aborting with error : %s Code: 0x%x Available Free mem : %zu MB",
-              queue->base_address, errorMsg, status, global_available_mem/Mi);
-    } else {
-      ClPrint(amd::LOG_NONE, amd::LOG_ALWAYS,
-        "Callback: Queue %p aborting with error : %s code: 0x%x", queue->base_address,
-        errorMsg, status);
-    }
-    abort();
-  }
-}
-
 // ================================================================================================
 bool VirtualGPU::createSchedulerParam()
 {
@@ -3023,8 +2996,8 @@ bool VirtualGPU::createSchedulerParam()
 
     // The queue is written by multiple threads of the scheduler kernel
     if (HSA_STATUS_SUCCESS != hsa_queue_create(gpu_device(), 2048, HSA_QUEUE_TYPE_MULTI,
-        callbackQueue, this, std::numeric_limits<uint>::max(), std::numeric_limits<uint>::max(),
-        &schedulerQueue_)) {
+        callbackQueue, &roc_device_, std::numeric_limits<uint>::max(),
+        std::numeric_limits<uint>::max(), &schedulerQueue_)) {
       break;
     }