[HIP] Implement ext_oneapi_queue_priority

Author: al42and

source/adapters/hip/queue.cpp

@@ -38,8 +38,8 @@ hipStream_t ur_queue_handle_t_::getNextComputeStream(uint32_t *StreamToken) {
       // The second check is done after mutex is locked so other threads can not
       // change NumComputeStreams after that
       if (NumComputeStreams < ComputeStreams.size()) {
-        UR_CHECK_ERROR(hipStreamCreateWithFlags(
-            &ComputeStreams[NumComputeStreams++], Flags));
+        UR_CHECK_ERROR(hipStreamCreateWithPriority(
+            &ComputeStreams[NumComputeStreams++], Flags, Priority));
       }
     }
     Token = ComputeStreamIdx++;
@@ -97,8 +97,8 @@ hipStream_t ur_queue_handle_t_::getNextTransferStream() {
     // The second check is done after mutex is locked so other threads can not
     // change NumTransferStreams after that
     if (NumTransferStreams < TransferStreams.size()) {
-      UR_CHECK_ERROR(hipStreamCreateWithFlags(
-          &TransferStreams[NumTransferStreams++], Flags));
+      UR_CHECK_ERROR(hipStreamCreateWithPriority(
+          &TransferStreams[NumTransferStreams++], Flags, Priority));
     }
   }
   uint32_t Stream_i = TransferStreamIdx++ % TransferStreams.size();
@@ -119,6 +119,19 @@ urQueueCreate(ur_context_handle_t hContext, ur_device_handle_t hDevice,
     }
 
     unsigned int Flags = 0;
+    ur_queue_flags_t URFlags = 0;
+    int Priority = 0; // Not guaranteed, but, in ROCm 5.7, 0 is the default
+
+    if (pProps && pProps->stype == UR_STRUCTURE_TYPE_QUEUE_PROPERTIES) {
+      URFlags = pProps->flags;
+      if (URFlags & UR_QUEUE_FLAG_PRIORITY_HIGH) {
+        ScopedContext Active(hContext->getDevice());
+        UR_CHECK_ERROR(hipDeviceGetStreamPriorityRange(nullptr, &Priority));
+      } else if (URFlags & UR_QUEUE_FLAG_PRIORITY_LOW) {
+        ScopedContext Active(hContext->getDevice());
+        UR_CHECK_ERROR(hipDeviceGetStreamPriorityRange(&Priority, nullptr));
+      }
+    }
 
     const bool IsOutOfOrder =
         pProps ? pProps->flags & UR_QUEUE_FLAG_OUT_OF_ORDER_EXEC_MODE_ENABLE
@@ -131,7 +144,7 @@ urQueueCreate(ur_context_handle_t hContext, ur_device_handle_t hDevice,
 
     QueueImpl = std::unique_ptr<ur_queue_handle_t_>(new ur_queue_handle_t_{
         std::move(ComputeHipStreams), std::move(TransferHipStreams), hContext,
-        hDevice, Flags, pProps ? pProps->flags : 0});
+        hDevice, Flags, pProps ? pProps->flags : 0, Priority});
 
     *phQueue = QueueImpl.release();
 

source/adapters/hip/queue.hpp

@@ -44,6 +44,7 @@ struct ur_queue_handle_t_ {
   unsigned int LastSyncTransferStreams;
   unsigned int Flags;
   ur_queue_flags_t URFlags;
+  int Priority;
   // When ComputeStreamSyncMutex and ComputeStreamMutex both need to be
   // locked at the same time, ComputeStreamSyncMutex should be locked first
   // to avoid deadlocks
@@ -56,7 +57,7 @@ struct ur_queue_handle_t_ {
   ur_queue_handle_t_(std::vector<native_type> &&ComputeStreams,
                      std::vector<native_type> &&TransferStreams,
                      ur_context_handle_t Context, ur_device_handle_t Device,
-                     unsigned int Flags, ur_queue_flags_t URFlags,
+                     unsigned int Flags, ur_queue_flags_t URFlags, int Priority,
                      bool BackendOwns = true)
       : ComputeStreams{std::move(ComputeStreams)}, TransferStreams{std::move(
                                                        TransferStreams)},
@@ -66,7 +67,7 @@ struct ur_queue_handle_t_ {
         Device{Device}, RefCount{1}, EventCount{0}, ComputeStreamIdx{0},
         TransferStreamIdx{0}, NumComputeStreams{0}, NumTransferStreams{0},
         LastSyncComputeStreams{0}, LastSyncTransferStreams{0}, Flags(Flags),
-        URFlags(URFlags), HasOwnership{BackendOwns} {
+        URFlags(URFlags), Priority(Priority), HasOwnership{BackendOwns} {
     urContextRetain(Context);
     urDeviceRetain(Device);
   }