[SYCL] Fix enqueue functions taking both kernel and properties (#14743)

The current implementation of the enqueue free functions taking both a launch_config and a kernel do not properly process the properties. This commit addresses this and adds a static assert about the properties passed to these only applying to the launch of the kernel and not how the compiler handles compiling the kernel. --------- Signed-off-by: Larsen, Steffen <[email protected]> Co-authored-by: Sergey Semenov <[email protected]>
intel · Jul 26, 2024 · 33325d4 · 33325d4
1 parent 6ac800b
commit 33325d4
Show file tree

Hide file tree

Showing 10 changed files with 289 additions and 35 deletions.
diff --git a/sycl/include/sycl/ext/intel/experimental/fpga_kernel_properties.hpp b/sycl/include/sycl/ext/intel/experimental/fpga_kernel_properties.hpp
@@ -104,6 +104,15 @@ struct is_property_key_of<
     : std::true_type {};
 
 namespace detail {
+template <intel::experimental::streaming_interface_options_enum option>
+struct HasCompileTimeEffect<
+    intel::experimental::streaming_interface_key::value_t<option>>
+    : std::true_type {};
+template <intel::experimental::register_map_interface_options_enum option>
+struct HasCompileTimeEffect<
+    intel::experimental::register_map_interface_key::value_t<option>>
+    : std::true_type {};
+
 template <intel::experimental::streaming_interface_options_enum Stall_Free>
 struct PropertyMetaInfo<
     intel::experimental::streaming_interface_key::value_t<Stall_Free>> {

diff --git a/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp b/sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp
@@ -34,6 +34,16 @@ template <typename RangeT>
 constexpr bool is_range_or_nd_range_v = is_range_or_nd_range<RangeT>::value;
 
 template <typename LCRangeT, typename LCPropertiesT> struct LaunchConfigAccess;
+
+// Checks that none of the properties in the property list has compile-time
+// effects on the kernel.
+template <typename T>
+struct NoPropertyHasCompileTimeKernelEffect : std::false_type {};
+template <typename... Ts>
+struct NoPropertyHasCompileTimeKernelEffect<properties_t<Ts...>> {
+  static constexpr bool value =
+      !(HasCompileTimeEffect<Ts>::value || ... || false);
+};
 } // namespace detail
 
 // Available only when Range is range or nd_range
@@ -42,6 +52,11 @@ template <
     typename = std::enable_if_t<
         ext::oneapi::experimental::detail::is_range_or_nd_range_v<RangeT>>>
 class launch_config {
+  static_assert(ext::oneapi::experimental::detail::
+                    NoPropertyHasCompileTimeKernelEffect<PropertiesT>::value,
+                "launch_config does not allow properties with compile-time "
+                "kernel effects.");
+
 public:
   launch_config(RangeT Range, PropertiesT Properties = {})
       : MRange{Range}, MProperties{Properties} {}
@@ -187,7 +202,8 @@ void parallel_for(handler &CGH,
                                                         Properties>
       ConfigAccess(Config);
   CGH.set_args<ArgsT...>(std::forward<ArgsT>(Args)...);
-  CGH.parallel_for(ConfigAccess.getRange(), KernelObj);
+  sycl::detail::HandlerAccess::parallelForImpl(
+      CGH, ConfigAccess.getRange(), ConfigAccess.getProperties(), KernelObj);
 }
 
 template <int Dimensions, typename Properties, typename... ArgsT>
@@ -263,7 +279,8 @@ void nd_launch(handler &CGH,
                                                         Properties>
       ConfigAccess(Config);
   CGH.set_args<ArgsT...>(std::forward<ArgsT>(Args)...);
-  CGH.parallel_for(ConfigAccess.getRange(), KernelObj);
+  sycl::detail::HandlerAccess::parallelForImpl(
+      CGH, ConfigAccess.getRange(), ConfigAccess.getProperties(), KernelObj);
 }
 
 template <int Dimensions, typename Properties, typename... ArgsT>

diff --git a/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp b/sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp
@@ -240,6 +240,19 @@ template <> struct is_property_key<work_item_progress_key> : std::true_type {};
 
 namespace detail {
 
+template <size_t... Dims>
+struct HasCompileTimeEffect<work_group_size_key::value_t<Dims...>>
+    : std::true_type {};
+template <size_t... Dims>
+struct HasCompileTimeEffect<work_group_size_hint_key::value_t<Dims...>>
+    : std::true_type {};
+template <uint32_t Size>
+struct HasCompileTimeEffect<sub_group_size_key::value_t<Size>>
+    : std::true_type {};
+template <sycl::aspect... Aspects>
+struct HasCompileTimeEffect<device_has_key::value_t<Aspects...>>
+    : std::true_type {};
+
 template <size_t Dim0, size_t... Dims>
 struct PropertyMetaInfo<work_group_size_key::value_t<Dim0, Dims...>> {
   static constexpr const char *name = "sycl-work-group-size";

diff --git a/sycl/include/sycl/ext/oneapi/properties/property.hpp b/sycl/include/sycl/ext/oneapi/properties/property.hpp
@@ -266,6 +266,8 @@ template <typename PropertyT> struct PropertyMetaInfo {
   static constexpr std::nullptr_t value = nullptr;
 };
 
+template <typename> struct HasCompileTimeEffect : std::false_type {};
+
 } // namespace detail
 
 template <typename T>

diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp
@@ -970,28 +970,11 @@ class __SYCL_EXPORT handler {
     }
   }
 
-  /// Process kernel properties.
+  /// Process runtime kernel properties.
   ///
   /// Stores information about kernel properties into the handler.
-  template <
-      typename KernelName,
-      typename PropertiesT = ext::oneapi::experimental::empty_properties_t>
-  void processProperties(PropertiesT Props) {
-    using KI = detail::KernelInfo<KernelName>;
-    static_assert(
-        ext::oneapi::experimental::is_property_list<PropertiesT>::value,
-        "Template type is not a property list.");
-    static_assert(
-        !PropertiesT::template has_property<
-            sycl::ext::intel::experimental::fp_control_key>() ||
-            (PropertiesT::template has_property<
-                 sycl::ext::intel::experimental::fp_control_key>() &&
-             KI::isESIMD()),
-        "Floating point control property is supported for ESIMD kernels only.");
-    static_assert(
-        !PropertiesT::template has_property<
-            sycl::ext::oneapi::experimental::indirectly_callable_key>(),
-        "indirectly_callable property cannot be applied to SYCL kernels");
+  template <typename PropertiesT>
+  void processLaunchProperties(PropertiesT Props) {
     if constexpr (PropertiesT::template has_property<
                       sycl::ext::intel::experimental::cache_config_key>()) {
       auto Config = Props.template get_property<
@@ -1042,6 +1025,32 @@ class __SYCL_EXPORT handler {
     checkAndSetClusterRange(Props);
   }
 
+  /// Process kernel properties.
+  ///
+  /// Stores information about kernel properties into the handler.
+  template <
+      typename KernelName,
+      typename PropertiesT = ext::oneapi::experimental::empty_properties_t>
+  void processProperties(PropertiesT Props) {
+    using KI = detail::KernelInfo<KernelName>;
+    static_assert(
+        ext::oneapi::experimental::is_property_list<PropertiesT>::value,
+        "Template type is not a property list.");
+    static_assert(
+        !PropertiesT::template has_property<
+            sycl::ext::intel::experimental::fp_control_key>() ||
+            (PropertiesT::template has_property<
+                 sycl::ext::intel::experimental::fp_control_key>() &&
+             KI::isESIMD()),
+        "Floating point control property is supported for ESIMD kernels only.");
+    static_assert(
+        !PropertiesT::template has_property<
+            sycl::ext::oneapi::experimental::indirectly_callable_key>(),
+        "indirectly_callable property cannot be applied to SYCL kernels");
+
+    processLaunchProperties(Props);
+  }
+
   /// Checks whether it is possible to copy the source shape to the destination
   /// shape(the shapes are described by the accessor ranges) by using
   /// copying by regions of memory and not copying element by element
@@ -1440,18 +1449,44 @@ class __SYCL_EXPORT handler {
   ///
   /// \param NumWorkItems is a range defining indexing space.
   /// \param Kernel is a SYCL kernel function.
-  template <int Dims>
-  void parallel_for_impl(range<Dims> NumWorkItems, kernel Kernel) {
+  /// \param Properties is the properties.
+  template <int Dims, typename PropertiesT>
+  void parallel_for_impl(range<Dims> NumWorkItems, PropertiesT Props,
+                         kernel Kernel) {
     throwIfActionIsCreated();
     MKernel = detail::getSyclObjImpl(std::move(Kernel));
     detail::checkValueRange<Dims>(NumWorkItems);
     setNDRangeDescriptor(std::move(NumWorkItems));
+    processLaunchProperties<PropertiesT>(Props);
     setType(detail::CGType::Kernel);
     setNDRangeUsed(false);
     extractArgsAndReqs();
     MKernelName = getKernelName();
   }
 
+  /// Defines and invokes a SYCL kernel function for the specified range and
+  /// offsets.
+  ///
+  /// The SYCL kernel function is defined as SYCL kernel object.
+  ///
+  /// \param NDRange is a ND-range defining global and local sizes as
+  /// well as offset.
+  /// \param Properties is the properties.
+  /// \param Kernel is a SYCL kernel function.
+  template <int Dims, typename PropertiesT>
+  void parallel_for_impl(nd_range<Dims> NDRange, PropertiesT Props,
+                         kernel Kernel) {
+    throwIfActionIsCreated();
+    MKernel = detail::getSyclObjImpl(std::move(Kernel));
+    detail::checkValueRange<Dims>(NDRange);
+    setNDRangeDescriptor(std::move(NDRange));
+    processLaunchProperties(Props);
+    setType(detail::CGType::Kernel);
+    setNDRangeUsed(true);
+    extractArgsAndReqs();
+    MKernelName = getKernelName();
+  }
+
   /// Hierarchical kernel invocation method of a kernel defined as a lambda
   /// encoding the body of each work-group to launch.
   ///
@@ -2163,15 +2198,18 @@ class __SYCL_EXPORT handler {
   }
 
   void parallel_for(range<1> NumWorkItems, kernel Kernel) {
-    parallel_for_impl(NumWorkItems, Kernel);
+    parallel_for_impl(NumWorkItems,
+                      ext::oneapi::experimental::empty_properties_t{}, Kernel);
   }
 
   void parallel_for(range<2> NumWorkItems, kernel Kernel) {
-    parallel_for_impl(NumWorkItems, Kernel);
+    parallel_for_impl(NumWorkItems,
+                      ext::oneapi::experimental::empty_properties_t{}, Kernel);
   }
 
   void parallel_for(range<3> NumWorkItems, kernel Kernel) {
-    parallel_for_impl(NumWorkItems, Kernel);
+    parallel_for_impl(NumWorkItems,
+                      ext::oneapi::experimental::empty_properties_t{}, Kernel);
   }
 
   /// Defines and invokes a SYCL kernel function for the specified range and
@@ -2205,14 +2243,8 @@ class __SYCL_EXPORT handler {
   /// well as offset.
   /// \param Kernel is a SYCL kernel function.
   template <int Dims> void parallel_for(nd_range<Dims> NDRange, kernel Kernel) {
-    throwIfActionIsCreated();
-    MKernel = detail::getSyclObjImpl(std::move(Kernel));
-    detail::checkValueRange<Dims>(NDRange);
-    setNDRangeDescriptor(std::move(NDRange));
-    setType(detail::CGType::Kernel);
-    setNDRangeUsed(true);
-    extractArgsAndReqs();
-    MKernelName = getKernelName();
+    parallel_for_impl(NDRange, ext::oneapi::experimental::empty_properties_t{},
+                      Kernel);
   }
 
   /// Defines and invokes a SYCL kernel function.
@@ -3741,6 +3773,12 @@ class HandlerAccess {
   static void internalProfilingTagImpl(handler &Handler) {
     Handler.internalProfilingTagImpl();
   }
+
+  template <typename RangeT, typename PropertiesT>
+  static void parallelForImpl(handler &Handler, RangeT Range, PropertiesT Props,
+                              kernel Kernel) {
+    Handler.parallel_for_impl(Range, Props, Kernel);
+  }
 };
 } // namespace detail
 

diff --git a/sycl/test/extensions/properties/kernel_properties_negative.cpp b/sycl/test/extensions/properties/kernel_properties_negative.cpp
@@ -0,0 +1,25 @@
+// RUN: %clangxx -ferror-limit=0 %fsycl-host-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s
+
+// Negative tests for kernel properties.
+
+#include <sycl/sycl.hpp>
+
+namespace oneapi = sycl::ext::oneapi::experimental;
+
+extern sycl::kernel TestKernel;
+
+int main() {
+  sycl::queue Q{};
+
+  oneapi::properties props1{oneapi::sub_group_size<8>};
+  oneapi::properties props2{
+      oneapi::sub_group_size<8>,
+      oneapi::work_group_progress<oneapi::forward_progress_guarantee::parallel,
+                                  oneapi::execution_scope::root_group>};
+
+  // expected-error-re@sycl/ext/oneapi/experimental/enqueue_functions.hpp:* {{static assertion failed due to requirement {{.*}} launch_config does not allow properties with compile-time kernel effects.}}
+  oneapi::launch_config<sycl::range<1>, decltype(props1)> LC1{{1}, props1};
+
+  // expected-error-re@sycl/ext/oneapi/experimental/enqueue_functions.hpp:* {{static assertion failed due to requirement {{.*}} launch_config does not allow properties with compile-time kernel effects.}}
+  oneapi::launch_config<sycl::range<1>, decltype(props2)> LC22{{1}, props2};
+}
diff --git a/sycl/test/extensions/properties/properties_kernel.cpp b/sycl/test/extensions/properties/properties_kernel.cpp
@@ -42,6 +42,15 @@ int main() {
   static_assert(is_property_key<sub_group_size_key>::value);
   static_assert(is_property_key<device_has_key>::value);
 
+  static_assert(sycl::ext::oneapi::experimental::detail::HasCompileTimeEffect<
+                work_group_size_key::value_t<1>>::value);
+  static_assert(sycl::ext::oneapi::experimental::detail::HasCompileTimeEffect<
+                work_group_size_hint_key::value_t<1>>::value);
+  static_assert(sycl::ext::oneapi::experimental::detail::HasCompileTimeEffect<
+                sub_group_size_key::value_t<28>>::value);
+  static_assert(sycl::ext::oneapi::experimental::detail::HasCompileTimeEffect<
+                device_has_key::value_t<aspect::fp64>>::value);
+
   static_assert(is_property_value<decltype(work_group_size<1>)>::value);
   static_assert(is_property_value<decltype(work_group_size<2, 2>)>::value);
   static_assert(is_property_value<decltype(work_group_size<3, 3, 3>)>::value);

diff --git a/sycl/test/extensions/properties/properties_kernel_fpga.cpp b/sycl/test/extensions/properties/properties_kernel_fpga.cpp
@@ -16,6 +16,17 @@ int main() {
   static_assert(oneapi::experimental::is_property_key<
                 intel::experimental::pipelined_key>::value);
 
+  // Check that oneapi::experimental::detail::HasCompileTimeEffect is
+  // correctly specialized
+  static_assert(oneapi::experimental::detail::HasCompileTimeEffect<
+                intel::experimental::register_map_interface_key::value_t<
+                    intel::experimental::register_map_interface_options_enum::
+                        wait_for_done_write>>::value);
+  static_assert(oneapi::experimental::detail::HasCompileTimeEffect<
+                intel::experimental::streaming_interface_key::value_t<
+                    intel::experimental::streaming_interface_options_enum::
+                        accept_downstream_stall>>::value);
+
   // Check that oneapi::experimental::is_property_value is correctly specialized
   static_assert(oneapi::experimental::is_property_value<
                 decltype(intel::experimental::streaming_interface<

diff --git a/sycl/unittests/Extensions/CMakeLists.txt b/sycl/unittests/Extensions/CMakeLists.txt
@@ -13,6 +13,7 @@ add_sycl_unittest(ExtensionsTests OBJECT
   EnqueueFunctionsEvents.cpp
   DiscardEvent.cpp
   ProfilingTag.cpp
+  KernelProperties.cpp
 )
 
 add_subdirectory(CommandGraph)