[SYCL] Address Issue about Unable to Specify Kernel Properties when Reductions are Present (#20491)

Basically to fix this
[issue](#16320). Note: seems that
other than certain variants (overloads) of `nd_launch` which is
explicitly mentioned in the issue, some variants of `parallel_for` with
reductions are also having the same problem. Furthermore, these
problematic `nd_launch` variants actually calls the problematic
`paralle_for` variants, so in the end only `parallel_for` variants are
modified.

Also there're some test revisions in this PR. These changes are
previously blocked by the aforementioned issue. And as of the revision
itself, the key idea here is to get rid of all
`parallel_for`/`single_task` overloads that take a property list as a
parameter (since these overloads are to be deprecated).

---------

Signed-off-by: Hu, Peisen <[email protected]>

Author: HPS-1

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

@@ -444,6 +444,24 @@ struct ConflictingProperties<max_linear_work_group_size_key, Properties> {
   }();
 };
 
+// If the kernel (last element in the parameter pack) has a get(properties_tag)
+// method, return the property list specified by this getter. Otherwise, return
+// an empty properety list.
+template <typename... RestT>
+auto RetrieveGetMethodPropertiesOrEmpty(RestT &&...Rest) {
+  // Note: the following trivial identity lambda is used to avoid the issue
+  // that line "const auto &KernelObj = (Rest, ...);" may result in a "left
+  // operand of comma operator has no effect" error for certain compiler(s)
+  auto Identity = [](const auto &x) -> decltype(auto) { return x; };
+  const auto &KernelObj = (Identity(Rest), ...);
+  if constexpr (ext::oneapi::experimental::detail::HasKernelPropertiesGetMethod<
+                    decltype(KernelObj)>::value) {
+    return KernelObj.get(ext::oneapi::experimental::properties_tag{});
+  } else {
+    return ext::oneapi::experimental::empty_properties_t{};
+  }
+}
+
 } // namespace detail
 } // namespace ext::oneapi::experimental
 } // namespace _V1

sycl/include/sycl/handler.hpp

@@ -1789,27 +1789,33 @@ class __SYCL_EXPORT handler {
   std::enable_if_t<detail::AreAllButLastReductions<RestT...>::value &&
                    (sizeof...(RestT) > 1)>
   parallel_for(range<1> Range, RestT &&...Rest) {
-    parallel_for<KernelName>(Range,
-                             ext::oneapi::experimental::empty_properties_t{},
-                             std::forward<RestT>(Rest)...);
+    parallel_for<KernelName>(
+        Range,
+        ext::oneapi::experimental::detail::RetrieveGetMethodPropertiesOrEmpty(
+            Rest...),
+        std::forward<RestT>(Rest)...);
   }
 
   template <typename KernelName = detail::auto_name, typename... RestT>
   std::enable_if_t<detail::AreAllButLastReductions<RestT...>::value &&
                    (sizeof...(RestT) > 1)>
   parallel_for(range<2> Range, RestT &&...Rest) {
-    parallel_for<KernelName>(Range,
-                             ext::oneapi::experimental::empty_properties_t{},
-                             std::forward<RestT>(Rest)...);
+    parallel_for<KernelName>(
+        Range,
+        ext::oneapi::experimental::detail::RetrieveGetMethodPropertiesOrEmpty(
+            Rest...),
+        std::forward<RestT>(Rest)...);
   }
 
   template <typename KernelName = detail::auto_name, typename... RestT>
   std::enable_if_t<detail::AreAllButLastReductions<RestT...>::value &&
                    (sizeof...(RestT) > 1)>
   parallel_for(range<3> Range, RestT &&...Rest) {
-    parallel_for<KernelName>(Range,
-                             ext::oneapi::experimental::empty_properties_t{},
-                             std::forward<RestT>(Rest)...);
+    parallel_for<KernelName>(
+        Range,
+        ext::oneapi::experimental::detail::RetrieveGetMethodPropertiesOrEmpty(
+            Rest...),
+        std::forward<RestT>(Rest)...);
   }
 
   template <typename KernelName = detail::auto_name, int Dims,
@@ -1834,7 +1840,20 @@ class __SYCL_EXPORT handler {
 
   template <typename KernelName = detail::auto_name, int Dims,
             typename... RestT>
-  std::enable_if_t<detail::AreAllButLastReductions<RestT...>::value>
+  std::enable_if_t<detail::AreAllButLastReductions<RestT...>::value &&
+                   (sizeof...(RestT) > 1)> // variant with reductions
+  parallel_for(nd_range<Dims> Range, RestT &&...Rest) {
+    parallel_for<KernelName>(
+        Range,
+        ext::oneapi::experimental::detail::RetrieveGetMethodPropertiesOrEmpty(
+            Rest...),
+        std::forward<RestT>(Rest)...);
+  }
+
+  template <typename KernelName = detail::auto_name, int Dims,
+            typename... RestT>
+  std::enable_if_t<detail::AreAllButLastReductions<RestT...>::value &&
+                   (sizeof...(RestT) == 1)> // variant without reductions
   parallel_for(nd_range<Dims> Range, RestT &&...Rest) {
     parallel_for<KernelName>(Range,
                              ext::oneapi::experimental::empty_properties_t{},

sycl/test-e2e/Properties/cache_config.cpp

@@ -1,9 +1,6 @@
 // REQUIRES: gpu, level_zero
 
-// TODO: Currently using the -Wno-deprecated-declarations flag due to issue
-// https://github.com/intel/llvm/issues/16320. Remove the flag once the issue is
-// resolved.
-// RUN: %{build} -o %t.out -Wno-deprecated-declarations
+// RUN: %{build} -o %t.out
 // RUN: env UR_L0_DEBUG=1 %{run} %t.out 2>&1 | FileCheck %s
 
 #include <numeric>
@@ -16,34 +13,39 @@ using namespace sycl::ext::intel::experimental;
 using namespace sycl::ext::oneapi::experimental;
 
 struct KernelFunctor {
-
-  KernelFunctor() {}
-
   void operator()() const {}
   auto get(properties_tag) const { return properties{cache_config(large_slm)}; }
 };
 
 struct KernelFunctorND {
-
-  KernelFunctorND() {}
-
   void operator()(nd_item<2> i) const {}
   auto get(properties_tag) const { return properties{cache_config(large_slm)}; }
 };
 
 struct NegativeKernelFunctor {
-
-  NegativeKernelFunctor() {}
-
   void operator()(nd_item<2> i) const {}
   auto get(properties_tag) const { return properties{}; }
 };
 
 struct RangeKernelFunctor {
+  void operator()(id<2> i) const {}
+  auto get(properties_tag) const { return properties{cache_config(large_slm)}; }
+};
+
+struct WorkGroupFunctor {
+  void operator()(group<1> g) const {
+    g.parallel_for_work_item([&](h_item<1>) {});
+  }
+  auto get(properties_tag) const { return properties{cache_config(large_slm)}; }
+};
 
-  RangeKernelFunctor() {}
+template <typename T1> struct ReductionKernelFunctor {
+  T1 mInput_values;
+  ReductionKernelFunctor(T1 &Input_values) : mInput_values(Input_values) {}
 
-  void operator()(id<2> i) const {}
+  template <typename sumT> void operator()(id<1> idx, sumT &sum) const {
+    sum += mInput_values[idx];
+  }
   auto get(properties_tag) const { return properties{cache_config(large_slm)}; }
 };
 
@@ -63,18 +65,16 @@ int main() {
   // CHECK: zeKernelSetCacheConfig
   std::cout << "parallel_for_work_group(range, func)" << std::endl;
   q.submit([&](handler &cgh) {
-    cgh.parallel_for_work_group<class hpar_range>(
-        range<1>(8), properties,
-        [=](group<1> g) { g.parallel_for_work_item([&](h_item<1> i) {}); });
+    cgh.parallel_for_work_group<class hpar_range>(range<1>(8),
+                                                  WorkGroupFunctor{});
   });
 
   // CHECK: parallel_for_work_group(range, range, func)
   // CHECK: zeKernelSetCacheConfig
   std::cout << "parallel_for_work_group(range, range, func)" << std::endl;
   q.submit([&](handler &cgh) {
     cgh.parallel_for_work_group<class hpar_range_range>(
-        range<1>(8), range<1>(4), properties,
-        [=](group<1> g) { g.parallel_for_work_item([&](h_item<1> i) {}); });
+        range<1>(8), range<1>(4), WorkGroupFunctor{});
   });
 
   buffer<int> values_buf{1024};
@@ -92,8 +92,8 @@ int main() {
   q.submit([&](handler &cgh) {
     auto input_values = values_buf.get_access<access_mode::read>(cgh);
     auto sum_reduction = reduction(sum_buf, cgh, plus<>());
-    cgh.parallel_for(range<1>{1024}, properties, sum_reduction,
-                     [=](id<1> idx, auto &sum) { sum += input_values[idx]; });
+    cgh.parallel_for(range<1>{1024}, sum_reduction,
+                     ReductionKernelFunctor(input_values));
   });
 
   // CHECK: KernelFunctor single_task
@@ -111,7 +111,7 @@ int main() {
   // CHECK-NOT: zeKernelSetCacheConfig
   std::cout << "negative parallel_for with sycl::nd_range" << std::endl;
   q.parallel_for(nd_range<2>{range<2>(4, 4), range<2>(2, 2)},
-                 [=](nd_item<2> i) {})
+                 NegativeKernelFunctor{})
       .wait();
 
   // CHECK: negative parallel_for with KernelFunctor