[SYCL][Graph] Throw an exception when unsupported features used in a graph (#10789)

This PR contains a set of changes that implement throwing an exception
when a feature unsupported by

[sycl_ext_oneapi_graph](https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_graph.asciidoc)
is used.

An error is thrown when an application uses the following features with
the graph extension:
* An unsupported extension.
* Querying the event returned from graph submission for profiling
information.
* Level Zero immediate command-list (see
#10467)
* Specialization constants, kernel bundles, or reductions are used in a
graph node.

## Authors

Co-authored-by: Pablo Reble <[email protected]>
Co-authored-by: Julian Miller <[email protected]>
Co-authored-by: Ben Tracy <[email protected]>
Co-authored-by: Ewan Crawford <[email protected]>
Co-authored-by: Maxime France-Pillois
<[email protected]>

Author: 4 people

sycl/include/sycl/ext/oneapi/experimental/graph.hpp

@@ -31,6 +31,44 @@ namespace oneapi {
 namespace experimental {
 
 namespace detail {
+// List of sycl features and extensions which are not supported by graphs. Used
+// for throwing errors when these features are used with graphs.
+enum class UnsupportedGraphFeatures {
+  sycl_reductions = 0,
+  sycl_specialization_constants = 1,
+  sycl_kernel_bundle = 2,
+  sycl_ext_oneapi_kernel_properties = 3,
+  sycl_ext_oneapi_enqueue_barrier = 4,
+  sycl_ext_oneapi_memcpy2d = 5,
+  sycl_ext_oneapi_device_global = 6,
+  sycl_ext_oneapi_bindless_images = 7
+};
+
+constexpr const char *
+UnsupportedFeatureToString(UnsupportedGraphFeatures Feature) {
+  using UGF = UnsupportedGraphFeatures;
+  switch (Feature) {
+  case UGF::sycl_reductions:
+    return "Reductions";
+  case UGF::sycl_specialization_constants:
+    return "Specialization Constants";
+  case UGF::sycl_kernel_bundle:
+    return "Kernel Bundles";
+  case UGF::sycl_ext_oneapi_kernel_properties:
+    return "sycl_ext_oneapi_kernel_properties";
+  case UGF::sycl_ext_oneapi_enqueue_barrier:
+    return "sycl_ext_oneapi_enqueue_barrier";
+  case UGF::sycl_ext_oneapi_memcpy2d:
+    return "sycl_ext_oneapi_memcpy2d";
+  case UGF::sycl_ext_oneapi_device_global:
+    return "sycl_ext_oneapi_device_global";
+  case UGF::sycl_ext_oneapi_bindless_images:
+    return "sycl_ext_oneapi_bindless_images";
+  default:
+    return {};
+  }
+}
+
 class node_impl;
 class graph_impl;
 class exec_graph_impl;

sycl/include/sycl/handler.hpp

@@ -135,7 +135,7 @@ class pipe;
 
 namespace ext::oneapi::experimental::detail {
 class graph_impl;
-}
+} // namespace ext::oneapi::experimental::detail
 namespace detail {
 
 class handler_impl;
@@ -1578,6 +1578,10 @@ class __SYCL_EXPORT handler {
   void set_specialization_constant(
       typename std::remove_reference_t<decltype(SpecName)>::value_type Value) {
 
+    throwIfGraphAssociated<
+        ext::oneapi::experimental::detail::UnsupportedGraphFeatures::
+            sycl_specialization_constants>();
+
     setStateSpecConstSet();
 
     std::shared_ptr<detail::kernel_bundle_impl> KernelBundleImplPtr =
@@ -1592,6 +1596,10 @@ class __SYCL_EXPORT handler {
   typename std::remove_reference_t<decltype(SpecName)>::value_type
   get_specialization_constant() const {
 
+    throwIfGraphAssociated<
+        ext::oneapi::experimental::detail::UnsupportedGraphFeatures::
+            sycl_specialization_constants>();
+
     if (isStateExplicitKernelBundle())
       throw sycl::exception(make_error_code(errc::invalid),
                             "Specialization constants cannot be read after "
@@ -2107,6 +2115,7 @@ class __SYCL_EXPORT handler {
   std::enable_if_t<
       ext::oneapi::experimental::is_property_list<PropertiesT>::value>
   single_task(PropertiesT Props, _KERNELFUNCPARAM(KernelFunc)) {
+    throwIfGraphAssociatedAndKernelProperties<PropertiesT>();
     single_task_lambda_impl<KernelName, KernelType, PropertiesT>(Props,
                                                                  KernelFunc);
   }
@@ -2117,6 +2126,7 @@ class __SYCL_EXPORT handler {
       ext::oneapi::experimental::is_property_list<PropertiesT>::value>
   parallel_for(range<1> NumWorkItems, PropertiesT Props,
                _KERNELFUNCPARAM(KernelFunc)) {
+    throwIfGraphAssociatedAndKernelProperties<PropertiesT>();
     parallel_for_lambda_impl<KernelName, KernelType, 1, PropertiesT>(
         NumWorkItems, Props, std::move(KernelFunc));
   }
@@ -2127,6 +2137,7 @@ class __SYCL_EXPORT handler {
       ext::oneapi::experimental::is_property_list<PropertiesT>::value>
   parallel_for(range<2> NumWorkItems, PropertiesT Props,
                _KERNELFUNCPARAM(KernelFunc)) {
+    throwIfGraphAssociatedAndKernelProperties<PropertiesT>();
     parallel_for_lambda_impl<KernelName, KernelType, 2, PropertiesT>(
         NumWorkItems, Props, std::move(KernelFunc));
   }
@@ -2137,6 +2148,7 @@ class __SYCL_EXPORT handler {
       ext::oneapi::experimental::is_property_list<PropertiesT>::value>
   parallel_for(range<3> NumWorkItems, PropertiesT Props,
                _KERNELFUNCPARAM(KernelFunc)) {
+    throwIfGraphAssociatedAndKernelProperties<PropertiesT>();
     parallel_for_lambda_impl<KernelName, KernelType, 3, PropertiesT>(
         NumWorkItems, Props, std::move(KernelFunc));
   }
@@ -2147,6 +2159,7 @@ class __SYCL_EXPORT handler {
       ext::oneapi::experimental::is_property_list<PropertiesT>::value>
   parallel_for(nd_range<Dims> Range, PropertiesT Properties,
                _KERNELFUNCPARAM(KernelFunc)) {
+    throwIfGraphAssociatedAndKernelProperties<PropertiesT>();
     parallel_for_impl<KernelName>(Range, Properties, std::move(KernelFunc));
   }
 
@@ -2159,6 +2172,9 @@ class __SYCL_EXPORT handler {
       detail::AreAllButLastReductions<RestT...>::value &&
       ext::oneapi::experimental::is_property_list<PropertiesT>::value>
   parallel_for(range<1> Range, PropertiesT Properties, RestT &&...Rest) {
+    throwIfGraphAssociated<ext::oneapi::experimental::detail::
+                               UnsupportedGraphFeatures::sycl_reductions>();
+    throwIfGraphAssociatedAndKernelProperties<PropertiesT>();
     detail::reduction_parallel_for<KernelName>(*this, Range, Properties,
                                                std::forward<RestT>(Rest)...);
   }
@@ -2170,6 +2186,9 @@ class __SYCL_EXPORT handler {
       detail::AreAllButLastReductions<RestT...>::value &&
       ext::oneapi::experimental::is_property_list<PropertiesT>::value>
   parallel_for(range<2> Range, PropertiesT Properties, RestT &&...Rest) {
+    throwIfGraphAssociated<ext::oneapi::experimental::detail::
+                               UnsupportedGraphFeatures::sycl_reductions>();
+    throwIfGraphAssociatedAndKernelProperties<PropertiesT>();
     detail::reduction_parallel_for<KernelName>(*this, Range, Properties,
                                                std::forward<RestT>(Rest)...);
   }
@@ -2181,6 +2200,9 @@ class __SYCL_EXPORT handler {
       detail::AreAllButLastReductions<RestT...>::value &&
       ext::oneapi::experimental::is_property_list<PropertiesT>::value>
   parallel_for(range<3> Range, PropertiesT Properties, RestT &&...Rest) {
+    throwIfGraphAssociated<ext::oneapi::experimental::detail::
+                               UnsupportedGraphFeatures::sycl_reductions>();
+    throwIfGraphAssociatedAndKernelProperties<PropertiesT>();
     detail::reduction_parallel_for<KernelName>(*this, Range, Properties,
                                                std::forward<RestT>(Rest)...);
   }
@@ -2216,6 +2238,8 @@ class __SYCL_EXPORT handler {
       detail::AreAllButLastReductions<RestT...>::value &&
       ext::oneapi::experimental::is_property_list<PropertiesT>::value>
   parallel_for(nd_range<Dims> Range, PropertiesT Properties, RestT &&...Rest) {
+    throwIfGraphAssociated<ext::oneapi::experimental::detail::
+                               UnsupportedGraphFeatures::sycl_reductions>();
     detail::reduction_parallel_for<KernelName>(*this, Range, Properties,
                                                std::forward<RestT>(Rest)...);
   }
@@ -2235,6 +2259,7 @@ class __SYCL_EXPORT handler {
             int Dims, typename PropertiesT>
   void parallel_for_work_group(range<Dims> NumWorkGroups, PropertiesT Props,
                                _KERNELFUNCPARAM(KernelFunc)) {
+    throwIfGraphAssociatedAndKernelProperties<PropertiesT>();
     parallel_for_work_group_lambda_impl<KernelName, KernelType, Dims,
                                         PropertiesT>(NumWorkGroups, Props,
                                                      KernelFunc);
@@ -2245,6 +2270,7 @@ class __SYCL_EXPORT handler {
   void parallel_for_work_group(range<Dims> NumWorkGroups,
                                range<Dims> WorkGroupSize, PropertiesT Props,
                                _KERNELFUNCPARAM(KernelFunc)) {
+    throwIfGraphAssociatedAndKernelProperties<PropertiesT>();
     parallel_for_work_group_lambda_impl<KernelName, KernelType, Dims,
                                         PropertiesT>(
         NumWorkGroups, WorkGroupSize, Props, KernelFunc);
@@ -2552,6 +2578,9 @@ class __SYCL_EXPORT handler {
   /// until all commands previously submitted to this queue have entered the
   /// complete state.
   void ext_oneapi_barrier() {
+    throwIfGraphAssociated<
+        ext::oneapi::experimental::detail::UnsupportedGraphFeatures::
+            sycl_ext_oneapi_enqueue_barrier>();
     throwIfActionIsCreated();
     setType(detail::CG::Barrier);
   }
@@ -2637,6 +2666,9 @@ class __SYCL_EXPORT handler {
             typename = std::enable_if_t<std::is_same_v<T, unsigned char>>>
   void ext_oneapi_memcpy2d(void *Dest, size_t DestPitch, const void *Src,
                            size_t SrcPitch, size_t Width, size_t Height) {
+    throwIfGraphAssociated<
+        ext::oneapi::experimental::detail::UnsupportedGraphFeatures::
+            sycl_ext_oneapi_memcpy2d>();
     throwIfActionIsCreated();
     if (Width > DestPitch)
       throw sycl::exception(sycl::make_error_code(errc::invalid),
@@ -2815,6 +2847,9 @@ class __SYCL_EXPORT handler {
   void memcpy(ext::oneapi::experimental::device_global<T, PropertyListT> &Dest,
               const void *Src, size_t NumBytes = sizeof(T),
               size_t DestOffset = 0) {
+    throwIfGraphAssociated<
+        ext::oneapi::experimental::detail::UnsupportedGraphFeatures::
+            sycl_ext_oneapi_device_global>();
     if (sizeof(T) < DestOffset + NumBytes)
       throw sycl::exception(make_error_code(errc::invalid),
                             "Copy to device_global is out of bounds.");
@@ -2847,6 +2882,9 @@ class __SYCL_EXPORT handler {
   memcpy(void *Dest,
          const ext::oneapi::experimental::device_global<T, PropertyListT> &Src,
          size_t NumBytes = sizeof(T), size_t SrcOffset = 0) {
+    throwIfGraphAssociated<
+        ext::oneapi::experimental::detail::UnsupportedGraphFeatures::
+            sycl_ext_oneapi_device_global>();
     if (sizeof(T) < SrcOffset + NumBytes)
       throw sycl::exception(make_error_code(errc::invalid),
                             "Copy from device_global is out of bounds.");
@@ -3368,8 +3406,34 @@ class __SYCL_EXPORT handler {
                             "handler::require() before it can be used.");
   }
 
+  template <typename PropertiesT>
+  std::enable_if_t<
+      ext::oneapi::experimental::is_property_list<PropertiesT>::value>
+  throwIfGraphAssociatedAndKernelProperties() const {
+    if (!std::is_same_v<PropertiesT,
+                        ext::oneapi::experimental::detail::empty_properties_t>)
+      throwIfGraphAssociated<
+          ext::oneapi::experimental::detail::UnsupportedGraphFeatures::
+              sycl_ext_oneapi_kernel_properties>();
+  }
+
   // Set value of the gpu cache configuration for the kernel.
   void setKernelCacheConfig(sycl::detail::pi::PiKernelCacheConfig);
+
+  template <
+      ext::oneapi::experimental::detail::UnsupportedGraphFeatures FeatureT>
+  void throwIfGraphAssociated() const {
+
+    if (getCommandGraph()) {
+      std::string FeatureString =
+          ext::oneapi::experimental::detail::UnsupportedFeatureToString(
+              FeatureT);
+      throw sycl::exception(sycl::make_error_code(errc::invalid),
+                            "The " + FeatureString +
+                                " feature is not yet available "
+                                "for use with the SYCL Graph extension.");
+    }
+  }
 };
 } // namespace _V1
 } // namespace sycl

sycl/include/sycl/reduction.hpp

@@ -1355,7 +1355,8 @@ struct NDRangeReduction<
               sycl::atomic_ref<int, memory_order::acq_rel, memory_scope::device,
                                access::address_space::global_space>(
                   NWorkGroupsFinished[0]);
-          DoReducePartialSumsInLastWG[0] = ++NFinished == NWorkGroups;
+          DoReducePartialSumsInLastWG[0] =
+              ++NFinished == static_cast<int>(NWorkGroups);
         }
 
         workGroupBarrier();

sycl/plugins/unified_runtime/ur/adapters/level_zero/command_buffer.cpp

@@ -663,6 +663,12 @@ UR_APIEXPORT ur_result_t UR_APICALL urCommandBufferEnqueueExp(
     ur_exp_command_buffer_handle_t CommandBuffer, ur_queue_handle_t Queue,
     uint32_t NumEventsInWaitList, const ur_event_handle_t *EventWaitList,
     ur_event_handle_t *Event) {
+  // There are issues with immediate command lists so return an error if the
+  // queue is in that mode.
+  if (Queue->UsingImmCmdLists) {
+    return UR_RESULT_ERROR_INVALID_QUEUE_PROPERTIES;
+  }
+
   std::scoped_lock<ur_shared_mutex> lock(Queue->Mutex);
   // Use compute engine rather than copy engine
   const auto UseCopyEngine = false;

sycl/source/detail/event_impl.cpp

@@ -278,6 +278,11 @@ void event_impl::checkProfilingPreconditions() const {
         "Profiling information is unavailable as the queue associated with "
         "the event does not have the 'enable_profiling' property.");
   }
+  if (MEventFromSubmitedExecCommandBuffer) {
+    throw sycl::exception(make_error_code(sycl::errc::invalid),
+                          "Profiling information is unavailable for events "
+                          "returned by a graph submission.");
+  }
 }
 
 template <>

sycl/source/detail/event_impl.hpp

@@ -282,6 +282,14 @@ class event_impl {
     return MGraph.lock();
   }
 
+  void setEventFromSubmitedExecCommandBuffer(bool value) {
+    MEventFromSubmitedExecCommandBuffer = value;
+  }
+
+  bool isEventFromSubmitedExecCommandBuffer() const {
+    return MEventFromSubmitedExecCommandBuffer;
+  }
+
 protected:
   // When instrumentation is enabled emits trace event for event wait begin and
   // returns the telemetry event generated for the wait
@@ -332,6 +340,8 @@ class event_impl {
   /// Store the command graph associated with this event, if any.
   /// This event is also be stored in the graph so a weak_ptr is used.
   std::weak_ptr<ext::oneapi::experimental::detail::graph_impl> MGraph;
+  /// Indicates that the event results from a command graph submission
+  bool MEventFromSubmitedExecCommandBuffer = false;
 
   // If this event represents a submission to a
   // sycl::detail::pi::PiExtCommandBuffer the sync point for that submission is

sycl/source/detail/fusion/fusion_wrapper_impl.cpp

@@ -27,6 +27,11 @@ bool fusion_wrapper_impl::is_in_fusion_mode() const {
 }
 
 void fusion_wrapper_impl::start_fusion() {
+  if (MQueue->getCommandGraph()) {
+    throw sycl::exception(sycl::make_error_code(errc::invalid),
+                          "SYCL kernel fusion can NOT be started "
+                          "on a queue that is in a recording state.");
+  }
   detail::Scheduler::getInstance().startFusion(MQueue);
 }
 

sycl/source/detail/graph_impl.cpp

@@ -372,6 +372,7 @@ exec_graph_impl::enqueue(const std::shared_ptr<sycl::detail::queue_impl> &Queue,
     auto NewEvent = std::make_shared<sycl::detail::event_impl>(Queue);
     NewEvent->setContextImpl(Queue->getContextImplPtr());
     NewEvent->setStateIncomplete();
+    NewEvent->setEventFromSubmitedExecCommandBuffer(true);
     return NewEvent;
   });
 
@@ -395,7 +396,14 @@ exec_graph_impl::enqueue(const std::shared_ptr<sycl::detail::queue_impl> &Queue,
               ->call_nocheck<
                   sycl::detail::PiApiKind::piextEnqueueCommandBuffer>(
                   CommandBuffer, Queue->getHandleRef(), 0, nullptr, OutEvent);
-      if (Res != pi_result::PI_SUCCESS) {
+      if (Res == pi_result::PI_ERROR_INVALID_QUEUE_PROPERTIES) {
+        throw sycl::exception(
+            make_error_code(errc::invalid),
+            "Graphs cannot be submitted to a queue which uses "
+            "immediate command lists. Use "
+            "sycl::ext::intel::property::queue::no_immediate_"
+            "command_list to disable them.");
+      } else if (Res != pi_result::PI_SUCCESS) {
         throw sycl::exception(
             errc::event,
             "Failed to enqueue event for command buffer submission");
@@ -509,6 +517,12 @@ modifiable_command_graph::finalize(const sycl::property_list &) const {
 bool modifiable_command_graph::begin_recording(queue &RecordingQueue) {
   auto QueueImpl = sycl::detail::getSyclObjImpl(RecordingQueue);
 
+  if (QueueImpl->is_in_fusion_mode()) {
+    throw sycl::exception(sycl::make_error_code(errc::invalid),
+                          "SYCL queue in kernel in fusion mode "
+                          "can NOT be recorded.");
+  }
+
   if (QueueImpl->get_context() != impl->getContext()) {
     throw sycl::exception(sycl::make_error_code(errc::invalid),
                           "begin_recording called for a queue whose context "