[SYCL] Add support for sycl::ext::oneapi::property::queue::use_priori…

…ty (#7523)

Adds a new queue property (hint) for setting its priority in backends
that support it.
Currently supported in Level Zero backend only and ignored for others.
Test: intel/llvm-test-suite#1414

Signed-off-by: Sergey V Maslov <sergey.v.maslov@intel.com>

sycl/doc/extensions/supported/sycl_ext_oneapi_queue_priority.asciidoc

@@ -0,0 +1,107 @@
+= sycl_ext_oneapi_queue_priority
+
+:source-highlighter: coderay
+:coderay-linenums-mode: table
+
+// This section needs to be after the document title.
+:doctype: book
+:toc2:
+:toc: left
+:encoding: utf-8
+:lang: en
+:dpcpp: pass:[DPC++]
+
+// Set the default source code type in this document to C++,
+// for syntax highlighting purposes.  This is needed because
+// docbook uses c++ and html5 uses cpp.
+:language: {basebackend@docbook:c++:cpp}
+
+
+== Notice
+
+[%hardbreaks]
+Copyright (C) 2022-2022 Intel Corporation.  All rights reserved.
+
+Khronos(R) is a registered trademark and SYCL(TM) and SPIR(TM) are trademarks
+of The Khronos Group Inc.  OpenCL(TM) is a trademark of Apple Inc. used by
+permission by Khronos.
+
+== Contact
+
+To report problems with this extension, please open a new issue at:
+
+https://github.com/intel/llvm/issues
+
+
+== Dependencies
+
+This extension is written against the SYCL 2020 revision 6 specification.  All
+references below to the "core SYCL specification" or to section numbers in the
+SYCL specification refer to that revision.
+
+== Status
+
+This extension is implemented and fully supported by {dpcpp}.
+[NOTE]
+====
+Although {dpcpp} supports this extension on all backends, it is currently used
+only on Level Zero. Other backends ignore the properties defined in this specification.
+====
+
+== Overview
+
+Introduce SYCL queue properties specifying the desired priority of a queue.
+These priorities are a hint and may be ignored if not supported by
+underlying backends. 
+
+== Specification
+
+=== Feature test macro
+
+This extension provides a feature-test macro as described in the core SYCL
+specification.  An implementation supporting this extension must predefine
+the macro `SYCL_EXT_ONEAPI_QUEUE_PRIORITY` to one of the values defined
+in the table below.  Applications can test for the existence of this macro
+to determine if the implementation supports this feature, or applications
+can test the macro's value to determine which of the extension's features
+the implementation supports.
+
+[%header,cols="1,5"]
+|===
+|Value
+|Description
+
+|1
+|Initial version of this extension.
+|===
+
+=== API of the extension
+
+This extension adds support for new properties for SYCL queue constructors
+taking properties list:
+
+```c++
+namespace sycl::ext::oneapi::property::queue {
+
+  class priority_normal {
+  public:
+    priority_normal() = default;
+  };
+  class priority_low {
+  public:
+    priority_low() = default;
+  };
+  class priority_high {
+  public:
+    priority_high() = default;
+  };
+
+} // namespace
+```
+The new properties hint the SYCL runtime that the queue gets the specified
+priority for execution if supported by underlying target runtimes. These
+properties are hints and may safely be ignored by an implementation.
+
+It is illegal to specify multiple differrent priority hints for the same queue.
+Doing so causes the `queue` constructor to throw a synchronous `exception` with
+the `errc::invalid` error code.

sycl/include/sycl/detail/pi.h

@@ -56,6 +56,8 @@
 // 11.16 Add PI_EXT_INTEL_DEVICE_INFO_MEMORY_CLOCK_RATE and
 // PI_EXT_INTEL_DEVICE_INFO_MEMORY_BUS_WIDTH as an extension for
 // piDeviceGetInfo.
+// 11.17 Added new PI_EXT_ONEAPI_QUEUE_PRIORITY_LOW and
+// PI_EXT_ONEAPI_QUEUE_PRIORITY_HIGH queue properties.
 
 #define _PI_H_VERSION_MAJOR 11
 #define _PI_H_VERSION_MINOR 16
@@ -580,6 +582,8 @@ constexpr pi_queue_properties PI_QUEUE_PROFILING_ENABLE = (1 << 1);
 constexpr pi_queue_properties PI_QUEUE_ON_DEVICE = (1 << 2);
 constexpr pi_queue_properties PI_QUEUE_ON_DEVICE_DEFAULT = (1 << 3);
 constexpr pi_queue_properties PI_EXT_ONEAPI_QUEUE_DISCARD_EVENTS = (1 << 4);
+constexpr pi_queue_properties PI_EXT_ONEAPI_QUEUE_PRIORITY_LOW = (1 << 5);
+constexpr pi_queue_properties PI_EXT_ONEAPI_QUEUE_PRIORITY_HIGH = (1 << 6);
 
 using pi_result = _pi_result;
 using pi_platform_info = _pi_platform_info;

sycl/include/sycl/detail/properties_traits.def

@@ -4,10 +4,16 @@ __SYCL_PARAM_TRAITS_SPEC(sycl::property::buffer::context_bound)
 __SYCL_PARAM_TRAITS_SPEC(sycl::property::image::use_host_ptr)
 __SYCL_PARAM_TRAITS_SPEC(sycl::property::image::use_mutex)
 __SYCL_PARAM_TRAITS_SPEC(sycl::property::image::context_bound)
-__SYCL_PARAM_TRAITS_SPEC(sycl::ext::oneapi::property::buffer::use_pinned_host_memory)
+__SYCL_PARAM_TRAITS_SPEC(
+    sycl::ext::oneapi::property::buffer::use_pinned_host_memory)
 __SYCL_PARAM_TRAITS_SPEC(sycl::property::noinit)
 __SYCL_PARAM_TRAITS_SPEC(sycl::property::no_init)
-__SYCL_PARAM_TRAITS_SPEC(sycl::property::context::cuda::use_primary_context) // Deprecated
-__SYCL_PARAM_TRAITS_SPEC(sycl::ext::oneapi::cuda::property::context::use_primary_context)
+__SYCL_PARAM_TRAITS_SPEC(
+    sycl::property::context::cuda::use_primary_context) // Deprecated
+__SYCL_PARAM_TRAITS_SPEC(
+    sycl::ext::oneapi::cuda::property::context::use_primary_context)
 __SYCL_PARAM_TRAITS_SPEC(sycl::property::queue::in_order)
 __SYCL_PARAM_TRAITS_SPEC(sycl::property::reduction::initialize_to_identity)
+__SYCL_PARAM_TRAITS_SPEC(sycl::ext::oneapi::property::queue::priority_low)
+__SYCL_PARAM_TRAITS_SPEC(sycl::ext::oneapi::property::queue::priority_high)
+__SYCL_PARAM_TRAITS_SPEC(sycl::ext::oneapi::property::queue::priority_normal)

sycl/include/sycl/detail/property_helper.hpp

@@ -40,8 +40,11 @@ enum DataLessPropKind {
   FusionNoBarrier = 13,
   FusionEnable = 14,
   FusionForce = 15,
+  QueuePriorityNormal = 16,
+  QueuePriorityLow = 17,
+  QueuePriorityHigh = 18,
   // Indicates the last known dataless property.
-  LastKnownDataLessPropKind = 15,
+  LastKnownDataLessPropKind = 18,
   // Exceeding 32 may cause ABI breaking change on some of OSes.
   DataLessPropKindSize = 32
 };

sycl/include/sycl/feature_test.hpp.in

@@ -37,6 +37,7 @@ __SYCL_INLINE_VER_NAMESPACE(_V1) {
 #define SYCL_EXT_ONEAPI_ASSERT 1
 #define SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS 1
 #define SYCL_EXT_ONEAPI_DISCARD_QUEUE_EVENTS 1
+#define SYCL_EXT_ONEAPI_QUEUE_PRIORITY 1
 #define SYCL_EXT_ONEAPI_ENQUEUE_BARRIER 1
 #define SYCL_EXT_ONEAPI_FREE_FUNCTION_QUERIES 1
 #define SYCL_EXT_ONEAPI_GROUP_ALGORITHMS 1

sycl/include/sycl/properties/queue_properties.hpp

@@ -28,6 +28,15 @@ namespace property {
 namespace queue {
 class discard_events
     : public ::sycl::detail::DataLessProperty<::sycl::detail::DiscardEvents> {};
+
+class priority_normal
+    : public sycl::detail::DataLessProperty<sycl::detail::QueuePriorityNormal> {
+};
+class priority_low
+    : public sycl::detail::DataLessProperty<sycl::detail::QueuePriorityLow> {};
+class priority_high
+    : public sycl::detail::DataLessProperty<sycl::detail::QueuePriorityHigh> {};
+
 } // namespace queue
 } // namespace property
 
@@ -67,6 +76,15 @@ template <>
 struct is_property_of<ext::oneapi::property::queue::discard_events, queue>
     : std::true_type {};
 template <>
+struct is_property_of<ext::oneapi::property::queue::priority_normal, queue>
+    : std::true_type {};
+template <>
+struct is_property_of<ext::oneapi::property::queue::priority_low, queue>
+    : std::true_type {};
+template <>
+struct is_property_of<ext::oneapi::property::queue::priority_high, queue>
+    : std::true_type {};
+template <>
 struct is_property_of<property::queue::cuda::use_default_stream, queue>
     : std::true_type {};
 template <>

sycl/plugins/level_zero/pi_level_zero.cpp

@@ -1002,6 +1002,14 @@ bool _pi_queue::isDiscardEvents() const {
   return ((this->Properties & PI_EXT_ONEAPI_QUEUE_DISCARD_EVENTS) != 0);
 }
 
+bool _pi_queue::isPriorityLow() const {
+  return ((this->Properties & PI_EXT_ONEAPI_QUEUE_PRIORITY_LOW) != 0);
+}
+
+bool _pi_queue::isPriorityHigh() const {
+  return ((this->Properties & PI_EXT_ONEAPI_QUEUE_PRIORITY_HIGH) != 0);
+}
+
 pi_result
 _pi_queue::resetCommandList(pi_command_list_ptr_t CommandList,
                             bool MakeAvailable,
@@ -1829,16 +1837,24 @@ _pi_queue::pi_queue_group_t::getZeQueue(uint32_t *QueueGroupOrdinal) {
   ZeCommandQueueDesc.ordinal = *QueueGroupOrdinal;
   ZeCommandQueueDesc.index = QueueIndex;
   ZeCommandQueueDesc.mode = ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS;
+  const char *Priority = "Normal";
+  if (Queue->isPriorityLow()) {
+    ZeCommandQueueDesc.priority = ZE_COMMAND_QUEUE_PRIORITY_PRIORITY_LOW;
+    Priority = "Low";
+  } else if (Queue->isPriorityHigh()) {
+    ZeCommandQueueDesc.priority = ZE_COMMAND_QUEUE_PRIORITY_PRIORITY_HIGH;
+    Priority = "High";
+  }
 
   // Evaluate performance of explicit usage for "0" index.
   if (QueueIndex != 0) {
     ZeCommandQueueDesc.flags = ZE_COMMAND_QUEUE_FLAG_EXPLICIT_ONLY;
   }
 
   zePrint("[getZeQueue]: create queue ordinal = %d, index = %d "
-          "(round robin in [%d, %d])\n",
+          "(round robin in [%d, %d]) priority = %s\n",
           ZeCommandQueueDesc.ordinal, ZeCommandQueueDesc.index, LowerIndex,
-          UpperIndex);
+          UpperIndex, Priority);
 
   auto ZeResult = ZE_CALL_NOCHECK(
       zeCommandQueueCreate, (Queue->Context->ZeContext, Queue->Device->ZeDevice,
@@ -1864,16 +1880,24 @@ pi_command_list_ptr_t &_pi_queue::pi_queue_group_t::getImmCmdList() {
   ZeCommandQueueDesc.ordinal = QueueOrdinal;
   ZeCommandQueueDesc.index = QueueIndex;
   ZeCommandQueueDesc.mode = ZE_COMMAND_QUEUE_MODE_ASYNCHRONOUS;
+  const char *Priority = "Normal";
+  if (Queue->isPriorityLow()) {
+    ZeCommandQueueDesc.priority = ZE_COMMAND_QUEUE_PRIORITY_PRIORITY_LOW;
+    Priority = "Low";
+  } else if (Queue->isPriorityHigh()) {
+    ZeCommandQueueDesc.priority = ZE_COMMAND_QUEUE_PRIORITY_PRIORITY_HIGH;
+    Priority = "High";
+  }
 
   // Evaluate performance of explicit usage for "0" index.
   if (QueueIndex != 0) {
     ZeCommandQueueDesc.flags = ZE_COMMAND_QUEUE_FLAG_EXPLICIT_ONLY;
   }
 
   zePrint("[getZeQueue]: create queue ordinal = %d, index = %d "
-          "(round robin in [%d, %d])\n",
+          "(round robin in [%d, %d]) priority = %s\n",
           ZeCommandQueueDesc.ordinal, ZeCommandQueueDesc.index, LowerIndex,
-          UpperIndex);
+          UpperIndex, Priority);
 
   ze_command_list_handle_t ZeCommandList;
   ZE_CALL_NOCHECK(zeCommandListCreateImmediate,
@@ -3532,7 +3556,9 @@ pi_result piQueueCreate(pi_context Context, pi_device Device,
   PI_ASSERT(!(Properties & ~(PI_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE |
                              PI_QUEUE_PROFILING_ENABLE | PI_QUEUE_ON_DEVICE |
                              PI_QUEUE_ON_DEVICE_DEFAULT |
-                             PI_EXT_ONEAPI_QUEUE_DISCARD_EVENTS)),
+                             PI_EXT_ONEAPI_QUEUE_DISCARD_EVENTS |
+                             PI_EXT_ONEAPI_QUEUE_PRIORITY_LOW |
+                             PI_EXT_ONEAPI_QUEUE_PRIORITY_HIGH)),
             PI_ERROR_INVALID_VALUE);
 
   PI_ASSERT(Context, PI_ERROR_INVALID_CONTEXT);

sycl/plugins/level_zero/pi_level_zero.hpp

@@ -432,9 +432,9 @@ using pi_command_list_ptr_t = pi_command_list_map_t::iterator;
 struct _pi_context : _pi_object {
   _pi_context(ze_context_handle_t ZeContext, pi_uint32 NumDevices,
               const pi_device *Devs, bool OwnZeContext)
-      : ZeContext{ZeContext},
-        OwnZeContext{OwnZeContext}, Devices{Devs, Devs + NumDevices},
-        SingleRootDevice(getRootDevice()), ZeCommandListInit{nullptr} {
+      : ZeContext{ZeContext}, OwnZeContext{OwnZeContext},
+        Devices{Devs, Devs + NumDevices}, SingleRootDevice(getRootDevice()),
+        ZeCommandListInit{nullptr} {
     // NOTE: one must additionally call initialize() to complete
     // PI context creation.
   }
@@ -788,6 +788,10 @@ struct _pi_queue : _pi_object {
   // Returns true if the queue has discard events property.
   bool isDiscardEvents() const;
 
+  // Returns true if the queue has explicit priority set by user.
+  bool isPriorityLow() const;
+  bool isPriorityHigh() const;
+
   // adjust the queue's batch size, knowing that the current command list
   // is being closed with a full batch.
   // For copy commands, IsCopy is set to 'true'.
@@ -1366,9 +1370,9 @@ struct _pi_program : _pi_object {
 
   // Construct a program in IL or Native state.
   _pi_program(state St, pi_context Context, const void *Input, size_t Length)
-      : Context{Context},
-        OwnZeModule{true}, State{St}, Code{new uint8_t[Length]},
-        CodeLength{Length}, ZeModule{nullptr}, ZeBuildLog{nullptr} {
+      : Context{Context}, OwnZeModule{true}, State{St},
+        Code{new uint8_t[Length]}, CodeLength{Length}, ZeModule{nullptr},
+        ZeBuildLog{nullptr} {
     std::memcpy(Code.get(), Input, Length);
   }
 

sycl/source/detail/queue_impl.hpp

@@ -312,6 +312,31 @@ class queue_impl {
       // queue property.
       CreationFlags |= PI_EXT_ONEAPI_QUEUE_DISCARD_EVENTS;
     }
+    // Track that priority settings are not ambiguous.
+    bool PrioritySeen = false;
+    if (MPropList
+            .has_property<ext::oneapi::property::queue::priority_normal>()) {
+      // Normal is the default priority, don't pass anything.
+      PrioritySeen = true;
+    }
+    if (MPropList.has_property<ext::oneapi::property::queue::priority_low>()) {
+      if (PrioritySeen) {
+        throw sycl::exception(
+            make_error_code(errc::invalid),
+            "Queue cannot be constructed with different priorities.");
+      }
+      CreationFlags |= PI_EXT_ONEAPI_QUEUE_PRIORITY_LOW;
+      PrioritySeen = true;
+    }
+    if (MPropList.has_property<ext::oneapi::property::queue::priority_high>()) {
+      if (PrioritySeen) {
+        throw sycl::exception(
+            make_error_code(errc::invalid),
+            "Queue cannot be constructed with different priorities.");
+      }
+      CreationFlags |= PI_EXT_ONEAPI_QUEUE_PRIORITY_HIGH;
+      PrioritySeen = true;
+    }
     RT::PiQueue Queue{};
     RT::PiContext Context = MContext->getHandleRef();
     RT::PiDevice Device = MDevice->getHandleRef();