[vdi] Add device assertion support.

- Once device assertion occurs, abort the host execution as well.
- TODO: This's the initial support. As we need to drain hostcall queue
  to ensure device assertion message being flushed out, hostcall
  listener needs an interface to explicitly drain its queue.

Change-Id: I8a04400aa7109bfd054ae5777c41a4abbf0db4a9

rocclr/device/rocm/rocdevice.cpp

@@ -1901,6 +1901,13 @@ bool Device::SetClockMode(const cl_set_device_clock_mode_input_amd setClockModeI
   return result;
 }
 
+static void callbackQueue(hsa_status_t status, hsa_queue_t* queue, void* data) {
+  if (status != HSA_STATUS_SUCCESS && status != HSA_STATUS_INFO_BREAK) {
+    // Abort on device exceptions.
+    abort();
+  }
+}
+
 hsa_queue_t* Device::acquireQueue(uint32_t queue_size_hint, bool coop_queue) {
   assert(queuePool_.size() <= GPU_MAX_HW_QUEUES);
   ClPrint(amd::LOG_INFO, amd::LOG_QUEUE, "number of allocated hardware queues: %d, maximum: %d",
@@ -1938,7 +1945,7 @@ hsa_queue_t* Device::acquireQueue(uint32_t queue_size_hint, bool coop_queue) {
     queue_type = HSA_QUEUE_TYPE_COOPERATIVE;
   }
 
-  while (hsa_queue_create(_bkendDevice, queue_size, queue_type, nullptr, nullptr,
+  while (hsa_queue_create(_bkendDevice, queue_size, queue_type, callbackQueue, this,
                           std::numeric_limits<uint>::max(), std::numeric_limits<uint>::max(),
                           &queue) != HSA_STATUS_SUCCESS) {
     queue_size >>= 1;

rocclr/device/rocm/rocvirtual.cpp

@@ -1841,6 +1841,13 @@ 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) {
+    // Abort on device exceptions.
+    abort();
+  }
+}
+
 bool VirtualGPU::createSchedulerParam()
 {
   if (nullptr != schedulerParam_) {
@@ -1856,7 +1863,7 @@ 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,
-        nullptr, nullptr, std::numeric_limits<uint>::max(), std::numeric_limits<uint>::max(),
+        callbackQueue, this, std::numeric_limits<uint>::max(), std::numeric_limits<uint>::max(),
         &schedulerQueue_)) {
       break;
     }