Minor fixes

Author: uditagarwal97

sycl/include/sycl/detail/kernel_launch_helper.hpp

@@ -256,16 +256,18 @@ struct KernelWrapper<
   }
 }; // KernelWrapper struct
 
-// This struct is inherited by sycl::handler.
+// This class encapsulates everything related to parsing kernel launch
+// properties.
 class KernelLaunchPropertyWrapper {
 public:
   // This struct is used to store kernel launch properties.
   // std::optional is used to indicate that the property is not set.
   // In some code paths, kernel launch properties are set multiple times
   // for the same kernel, that is why using std::optional to avoid overriding
   // previously set properties.
+  // This struct is used to pass kernel launch properties across the ABI
+  // boundary.
   struct KernelLaunchPropertiesT {
-
     struct ScopeForwardProgressProperty {
       std::optional<sycl::ext::oneapi::experimental::forward_progress_guarantee>
           Guarantee;
@@ -282,28 +284,11 @@ class KernelLaunchPropertyWrapper {
     std::array<size_t, 3> MClusterSize = {0, 0, 0};
 
     // Forward progress guarantee properties for work_item, sub_group and
-    // work_group scopes.
-    // Indexed by ExecutionScope enum.
+    // work_group scopes. We need to store them for validation later.
     std::array<ScopeForwardProgressProperty, 3> MForwardProgressProperties;
-
-    KernelLaunchPropertiesT() = default;
-
-    // TODO: Do you even need this?
-    KernelLaunchPropertiesT(
-        ur_kernel_cache_config_t _CacheConfig, bool _IsCooperative,
-        uint32_t _WorkGroupMemorySize, bool _UsesClusterLaunch,
-        size_t _ClusterDims, std::array<size_t, 3> _ClusterSize,
-        std::array<ScopeForwardProgressProperty, 3> _ForwardProgressProperties)
-        : MCacheConfig(_CacheConfig), MIsCooperative(_IsCooperative),
-          MWorkGroupMemorySize(_WorkGroupMemorySize),
-          MUsesClusterLaunch(_UsesClusterLaunch), MClusterDims(_ClusterDims),
-          MClusterSize(_ClusterSize),
-          MForwardProgressProperties(_ForwardProgressProperties) {}
-  }; // struct KernelLaunchPropertiesT
+  };
 
   /// Process runtime kernel properties.
-  ///
-  /// Stores information about kernel properties into the handler.
   template <typename PropertiesT>
   static KernelLaunchPropertiesT
   processKernelLaunchProperties(PropertiesT Props) {
@@ -423,9 +408,6 @@ class KernelLaunchPropertyWrapper {
   }
 
   /// Process kernel properties.
-  ///
-  /// Stores information about kernel properties into the handler.
-  ///
   /// Note: it is important that this function *does not* depend on kernel
   /// name or kernel type, because then it will be instantiated for every
   /// kernel, even though body of those instantiated functions could be almost

sycl/source/detail/kernel_data.hpp

@@ -8,13 +8,12 @@
 
 #pragma once
 
+#include <detail/device_impl.hpp>
 #include <detail/device_kernel_info.hpp>
 #include <detail/graph/dynamic_impl.hpp>
 #include <detail/kernel_arg_desc.hpp>
 #include <detail/ndrange_desc.hpp>
 
-#include <detail/device_impl.hpp>
-
 #include <sycl/detail/kernel_desc.hpp>
 #include <sycl/detail/kernel_launch_helper.hpp>
 

sycl/source/handler.cpp

@@ -1758,9 +1758,46 @@ void handler::verifyDeviceHasProgressGuarantee(
     sycl::ext::oneapi::experimental::execution_scope threadScope,
     sycl::ext::oneapi::experimental::execution_scope coordinationScope) {
 
-  // FIXME!
-  detail::KernelLaunchPropertyWrapper::KernelLaunchPropertiesT Kprop;
-  setKernelLaunchProperties(Kprop);
+  using execution_scope = sycl::ext::oneapi::experimental::execution_scope;
+  using forward_progress =
+      sycl::ext::oneapi::experimental::forward_progress_guarantee;
+  const bool supported = impl->get_device().supportsForwardProgress(
+      guarantee, threadScope, coordinationScope);
+  if (threadScope == execution_scope::work_group) {
+    if (!supported) {
+      throw sycl::exception(
+          sycl::errc::feature_not_supported,
+          "Required progress guarantee for work groups is not "
+          "supported by this device.");
+    }
+    // If we are here, the device supports the guarantee required but there is a
+    // caveat in that if the guarantee required is a concurrent guarantee, then
+    // we most likely also need to enable cooperative launch of the kernel. That
+    // is, although the device supports the required guarantee, some setup work
+    // is needed to truly make the device provide that guarantee at runtime.
+    // Otherwise, we will get the default guarantee which is weaker than
+    // concurrent. Same reasoning applies for sub_group but not for work_item.
+    // TODO: Further design work is probably needed to reflect this behavior in
+    // Unified Runtime.
+    if (guarantee == forward_progress::concurrent)
+      setKernelIsCooperative(true);
+  } else if (threadScope == execution_scope::sub_group) {
+    if (!supported) {
+      throw sycl::exception(sycl::errc::feature_not_supported,
+                            "Required progress guarantee for sub groups is not "
+                            "supported by this device.");
+    }
+    // Same reasoning as above.
+    if (guarantee == forward_progress::concurrent)
+      setKernelIsCooperative(true);
+  } else { // threadScope is execution_scope::work_item otherwise undefined
+           // behavior
+    if (!supported) {
+      throw sycl::exception(sycl::errc::feature_not_supported,
+                            "Required progress guarantee for work items is not "
+                            "supported by this device.");
+    }
+  }
 }
 #endif