[SYCL] Implement forward_progress extension API

sycl/include/sycl/detail/pi.h

@@ -437,7 +437,12 @@ typedef enum {
   // The number of max registers per block (device specific)
   PI_EXT_CODEPLAY_DEVICE_INFO_MAX_REGISTERS_PER_WORK_GROUP = 0x20009,
   PI_EXT_INTEL_DEVICE_INFO_ESIMD_SUPPORT = 0x2000A,
-
+  PI_EXT_ONEAPI_DEVICE_INFO_WORK_GROUP_PROGRESS_AT_ROOT_GROUP_LEVEL = 0x2000B,
+  PI_EXT_ONEAPI_DEVICE_INFO_SUB_GROUP_PROGRESS_AT_ROOT_GROUP_LEVEL = 0x2000C,
+  PI_EXT_ONEAPI_DEVICE_INFO_SUB_GROUP_PROGRESS_AT_WORK_GROUP_LEVEL = 0x2000D,
+  PI_EXT_ONEAPI_DEVICE_INFO_WORK_ITEM_PROGRESS_AT_ROOT_GROUP_LEVEL = 0x2000E,
+  PI_EXT_ONEAPI_DEVICE_INFO_WORK_ITEM_PROGRESS_AT_WORK_GROUP_LEVEL = 0x2000F,
+  PI_EXT_ONEAPI_DEVICE_INFO_WORK_ITEM_PROGRESS_AT_SUB_GROUP_LEVEL = 0x20010,
   // Bindless images, mipmaps, interop
   PI_EXT_ONEAPI_DEVICE_INFO_BINDLESS_IMAGES_SUPPORT = 0x20100,
   PI_EXT_ONEAPI_DEVICE_INFO_BINDLESS_IMAGES_SHARED_USM_SUPPORT = 0x20101,

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

@@ -0,0 +1,30 @@
+//==------- forward_progress.hpp - sycl_ext_oneapi_forward_progress -------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+namespace sycl {
+inline namespace _V1 {
+namespace ext {
+namespace oneapi {
+namespace experimental {
+
+enum class forward_progress_guarantee { concurrent, parallel, weakly_parallel };
+
+enum class execution_scope {
+  work_item,
+  sub_group,
+  work_group,
+  root_group,
+};
+
+} // namespace experimental
+} // namespace oneapi
+} // namespace ext
+} // namespace _V1
+} // namespace sycl

sycl/include/sycl/ext/oneapi/kernel_properties/properties.hpp

@@ -8,17 +8,16 @@
 
 #pragma once
 
-#include <sycl/aspects.hpp>                              // for aspect
+#include <array>                                             // for array
+#include <stddef.h>                                          // for size_t
+#include <stdint.h>                                          // for uint32_T
+#include <sycl/aspects.hpp>                                  // for aspect
+#include <sycl/ext/oneapi/experimental/forward_progress.hpp> // for forward_progress_guarantee enum
 #include <sycl/ext/oneapi/properties/property.hpp>       // for PropKind
 #include <sycl/ext/oneapi/properties/property_utils.hpp> // for SizeListToStr
 #include <sycl/ext/oneapi/properties/property_value.hpp> // for property_value
-
-#include <array>       // for array
-#include <stddef.h>    // for size_t
-#include <stdint.h>    // for uint32_t
-#include <type_traits> // for true_type
-#include <utility>     // for declval
-
+#include <type_traits>                                   // for true_type
+#include <utility>                                       // for declval
 namespace sycl {
 inline namespace _V1 {
 namespace ext::oneapi::experimental {
@@ -55,7 +54,8 @@ struct sub_group_size_key
                                  std::integral_constant<uint32_t, Size>>;
 };
 
-struct device_has_key : detail::compile_time_property_key<detail::PropKind::DeviceHas> {
+struct device_has_key
+    : detail::compile_time_property_key<detail::PropKind::DeviceHas> {
   template <aspect... Aspects>
   using value_t = property_value<device_has_key,
                                  std::integral_constant<aspect, Aspects>...>;
@@ -126,6 +126,88 @@ inline constexpr sub_group_size_key::value_t<Size> sub_group_size;
 template <aspect... Aspects>
 inline constexpr device_has_key::value_t<Aspects...> device_has;
 
+struct work_group_progress_key
+    : detail::compile_time_property_key<detail::PropKind::WorkGroupProgress> {
+  template <forward_progress_guarantee Guarantee,
+            execution_scope CoordinationScope>
+  using value_t = property_value<
+      work_group_progress_key,
+      std::integral_constant<forward_progress_guarantee, Guarantee>,
+      std::integral_constant<execution_scope, CoordinationScope>>;
+};
+
+struct sub_group_progress_key
+    : detail::compile_time_property_key<detail::PropKind::SubGroupProgress> {
+  template <forward_progress_guarantee Guarantee,
+            execution_scope CoordinationScope>
+  using value_t = property_value<
+      sub_group_progress_key,
+      std::integral_constant<forward_progress_guarantee, Guarantee>,
+      std::integral_constant<execution_scope, CoordinationScope>>;
+};
+
+struct work_item_progress_key
+    : detail::compile_time_property_key<detail::PropKind::WorkItemProgress> {
+  template <forward_progress_guarantee Guarantee,
+            execution_scope CoordinationScope>
+  using value_t = property_value<
+      work_item_progress_key,
+      std::integral_constant<forward_progress_guarantee, Guarantee>,
+      std::integral_constant<execution_scope, CoordinationScope>>;
+};
+
+template <forward_progress_guarantee Guarantee,
+          execution_scope CoordinationScope>
+struct property_value<
+    work_group_progress_key,
+    std::integral_constant<forward_progress_guarantee, Guarantee>,
+    std::integral_constant<execution_scope, CoordinationScope>> {
+  using key_t = work_group_progress_key;
+  static constexpr forward_progress_guarantee guarantee = Guarantee;
+  static constexpr execution_scope coordinationScope = CoordinationScope;
+};
+
+template <forward_progress_guarantee Guarantee,
+          execution_scope CoordinationScope>
+struct property_value<
+    sub_group_progress_key,
+    std::integral_constant<forward_progress_guarantee, Guarantee>,
+    std::integral_constant<execution_scope, CoordinationScope>> {
+  using key_t = work_group_progress_key;
+  static constexpr forward_progress_guarantee guarantee = Guarantee;
+  static constexpr execution_scope coordinationScope = CoordinationScope;
+};
+
+template <forward_progress_guarantee Guarantee,
+          execution_scope CoordinationScope>
+struct property_value<
+    work_item_progress_key,
+    std::integral_constant<forward_progress_guarantee, Guarantee>,
+    std::integral_constant<execution_scope, CoordinationScope>> {
+  using key_t = work_group_progress_key;
+  static constexpr forward_progress_guarantee guarantee = Guarantee;
+  static constexpr execution_scope coordinationScope = CoordinationScope;
+};
+
+template <forward_progress_guarantee Guarantee,
+          execution_scope CoordinationScope>
+inline constexpr work_group_progress_key::value_t<Guarantee, CoordinationScope>
+    work_group_progress;
+
+template <forward_progress_guarantee Guarantee,
+          execution_scope CoordinationScope>
+inline constexpr sub_group_progress_key::value_t<Guarantee, CoordinationScope>
+    sub_group_progress;
+
+template <forward_progress_guarantee Guarantee,
+          execution_scope CoordinationScope>
+inline constexpr work_item_progress_key::value_t<Guarantee, CoordinationScope>
+    work_item_progress;
+
+template <> struct is_property_key<work_group_progress_key> : std::true_type {};
+template <> struct is_property_key<sub_group_progress_key> : std::true_type {};
+template <> struct is_property_key<work_item_progress_key> : std::true_type {};
+
 namespace detail {
 template <size_t Dim0, size_t... Dims>
 struct PropertyMetaInfo<work_group_size_key::value_t<Dim0, Dims...>> {

sycl/include/sycl/ext/oneapi/properties/property.hpp

@@ -200,8 +200,11 @@ enum PropKind : uint32_t {
   ContiguousMemory = 59,
   FullGroup = 60,
   Naive = 61,
+  WorkGroupProgress = 62,
+  SubGroupProgress = 63,
+  WorkItemProgress = 64,
   // PropKindSize must always be the last value.
-  PropKindSize = 62,
+  PropKindSize = 65,
 };
 
 struct property_key_base_tag {};

sycl/include/sycl/handler.hpp

@@ -939,6 +939,11 @@ class __SYCL_EXPORT handler {
     }
   }
 
+  void verifyDeviceHasProgressGuarantee(
+      sycl::ext::oneapi::experimental::forward_progress_guarantee guarantee,
+      sycl::ext::oneapi::experimental::execution_scope threadScope,
+      sycl::ext::oneapi::experimental::execution_scope coordinationScope);
+
   /// Process kernel properties.
   ///
   /// Stores information about kernel properties into the handler.
@@ -973,6 +978,36 @@ class __SYCL_EXPORT handler {
     constexpr bool UsesRootSync = PropertiesT::template has_property<
         sycl::ext::oneapi::experimental::use_root_sync_key>();
     setKernelIsCooperative(UsesRootSync);
+    if constexpr (PropertiesT::template has_property<
+                      sycl::ext::oneapi::experimental::
+                          work_group_progress_key>()) {
+      auto prop = Props.template get_property<
+          sycl::ext::oneapi::experimental::work_group_progress_key>();
+      verifyDeviceHasProgressGuarantee(
+          prop.guarantee,
+          sycl::ext::oneapi::experimental::execution_scope::work_group,
+          prop.coordinationScope);
+    }
+    if constexpr (PropertiesT::template has_property<
+                      sycl::ext::oneapi::experimental::
+                          sub_group_progress_key>()) {
+      auto prop = Props.template get_property<
+          sycl::ext::oneapi::experimental::sub_group_progress_key>();
+      verifyDeviceHasProgressGuarantee(
+          prop.guarantee,
+          sycl::ext::oneapi::experimental::execution_scope::sub_group,
+          prop.coordinationScope);
+    }
+    if constexpr (PropertiesT::template has_property<
+                      sycl::ext::oneapi::experimental::
+                          work_item_progress_key>()) {
+      auto prop = Props.template get_property<
+          sycl::ext::oneapi::experimental::work_item_progress_key>();
+      verifyDeviceHasProgressGuarantee(
+          prop.guarantee,
+          sycl::ext::oneapi::experimental::execution_scope::work_item,
+          prop.coordinationScope);
+    }
   }
 
   /// Checks whether it is possible to copy the source shape to the destination

sycl/include/sycl/info/ext_oneapi_device_traits.def

@@ -6,6 +6,45 @@ __SYCL_PARAM_TRAITS_SPEC(ext::oneapi::experimental,device, max_global_work_group
 __SYCL_PARAM_TRAITS_TEMPLATE_SPEC(ext::oneapi::experimental,device, max_work_groups<1>, id<1>, PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_1D)
 __SYCL_PARAM_TRAITS_TEMPLATE_SPEC(ext::oneapi::experimental,device, max_work_groups<2>, id<2>, PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_2D)
 __SYCL_PARAM_TRAITS_TEMPLATE_SPEC(ext::oneapi::experimental,device, max_work_groups<3>, id<3>, PI_EXT_ONEAPI_DEVICE_INFO_MAX_WORK_GROUPS_3D)
+
+// Forward progress guarantees
+__SYCL_PARAM_TRAITS_TEMPLATE_SPEC(
+    ext::oneapi::experimental, device,
+    work_group_progress_capabilities<
+        ext::oneapi::experimental::execution_scope::root_group>,
+    std::vector<ext::oneapi::experimental::forward_progress_guarantee>,
+    PI_EXT_ONEAPI_DEVICE_INFO_WORK_GROUP_PROGRESS_AT_ROOT_GROUP_LEVEL)
+__SYCL_PARAM_TRAITS_TEMPLATE_SPEC(
+    ext::oneapi::experimental, device,
+    sub_group_progress_capabilities<
+        ext::oneapi::experimental::execution_scope::root_group>,
+    std::vector<ext::oneapi::experimental::forward_progress_guarantee>,
+    PI_EXT_ONEAPI_DEVICE_INFO_SUB_GROUP_PROGRESS_AT_ROOT_GROUP_LEVEL)
+__SYCL_PARAM_TRAITS_TEMPLATE_SPEC(
+    ext::oneapi::experimental, device,
+    sub_group_progress_capabilities<
+        ext::oneapi::experimental::execution_scope::work_group>,
+    std::vector<ext::oneapi::experimental::forward_progress_guarantee>,
+    PI_EXT_ONEAPI_DEVICE_INFO_SUB_GROUP_PROGRESS_AT_WORK_GROUP_LEVEL)
+__SYCL_PARAM_TRAITS_TEMPLATE_SPEC(
+    ext::oneapi::experimental, device,
+    work_item_progress_capabilities<
+        ext::oneapi::experimental::execution_scope::root_group>,
+    std::vector<ext::oneapi::experimental::forward_progress_guarantee>,
+    PI_EXT_ONEAPI_DEVICE_INFO_WORK_ITEM_PROGRESS_AT_ROOT_GROUP_LEVEL)
+__SYCL_PARAM_TRAITS_TEMPLATE_SPEC(
+    ext::oneapi::experimental, device,
+    work_item_progress_capabilities<
+        ext::oneapi::experimental::execution_scope::work_group>,
+    std::vector<ext::oneapi::experimental::forward_progress_guarantee>,
+    PI_EXT_ONEAPI_DEVICE_INFO_WORK_ITEM_PROGRESS_AT_WORK_GROUP_LEVEL)
+__SYCL_PARAM_TRAITS_TEMPLATE_SPEC(
+    ext::oneapi::experimental, device,
+    work_item_progress_capabilities<
+        ext::oneapi::experimental::execution_scope::sub_group>,
+    std::vector<ext::oneapi::experimental::forward_progress_guarantee>,
+    PI_EXT_ONEAPI_DEVICE_INFO_WORK_ITEM_PROGRESS_AT_SUB_GROUP_LEVEL)
+
 __SYCL_PARAM_TRAITS_SPEC(ext::oneapi::experimental, device, architecture,
                          ext::oneapi::experimental::architecture,
                          PI_EXT_ONEAPI_DEVICE_INFO_IP_VERSION)

sycl/include/sycl/info/info_desc.hpp

@@ -17,6 +17,7 @@
 #include <sycl/aspects.hpp>
 #include <sycl/detail/type_traits.hpp>
 #include <sycl/ext/oneapi/experimental/device_architecture.hpp>
+#include <sycl/ext/oneapi/experimental/forward_progress.hpp>
 #include <sycl/ext/oneapi/matrix/query-types.hpp>
 
 #include <sycl/range.hpp>
@@ -195,6 +196,13 @@ template <typename T, T param> struct compatibility_param_traits {};
 
 namespace ext::oneapi::experimental::info::device {
 template <int Dimensions> struct max_work_groups;
+template <ext::oneapi::experimental::execution_scope CoordinationScope>
+struct work_group_progress_capabilities;
+template <ext::oneapi::experimental::execution_scope CoordinationScope>
+struct sub_group_progress_capabilities;
+template <ext::oneapi::experimental::execution_scope CoordinationScope>
+struct work_item_progress_capabilities;
+
 } // namespace ext::oneapi::experimental::info::device
 #include <sycl/info/ext_codeplay_device_traits.def>
 #include <sycl/info/ext_intel_device_traits.def>

sycl/include/sycl/kernel.hpp

@@ -8,6 +8,8 @@
 
 #pragma once
 
+#include <cstddef>                            // for size_t
+#include <memory>                             // for shared_ptr, hash, opera...
 #include <sycl/backend_types.hpp>             // for backend, backend_return_t
 #include <sycl/context.hpp>                   // for context
 #include <sycl/detail/defines_elementary.hpp> // for __SYCL2020_DEPRECATED
@@ -21,10 +23,7 @@
 #include <sycl/device.hpp>              // for device
 #include <sycl/kernel_bundle_enums.hpp> // for bundle_state
 #include <sycl/range.hpp>               // for range
-
-#include <cstddef> // for size_t
-#include <memory>  // for shared_ptr, hash, opera...
-#include <variant> // for hash
+#include <variant>                      // for hash
 
 namespace sycl {
 inline namespace _V1 {

sycl/include/sycl/sycl.hpp

@@ -91,6 +91,7 @@
 #include <sycl/ext/oneapi/experimental/cuda/barrier.hpp>
 #include <sycl/ext/oneapi/experimental/enqueue_functions.hpp>
 #include <sycl/ext/oneapi/experimental/fixed_size_group.hpp>
+#include <sycl/ext/oneapi/experimental/forward_progress.hpp>
 #include <sycl/ext/oneapi/experimental/group_load_store.hpp>
 #include <sycl/ext/oneapi/experimental/opportunistic_group.hpp>
 #include <sycl/ext/oneapi/experimental/prefetch.hpp>