[SYCL] RoundedRangeKernel should inherit properties from the original (#18900)

Fix issue #16839.

---------

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

Author: HPS-1

sycl/include/sycl/handler.hpp

@@ -349,6 +349,16 @@ class RoundedRangeKernel {
       KernelFunc(item);
     }
   }
+
+  // Copy the properties_tag getter from the original kernel to propagate
+  // property(s)
+  template <
+      typename T = KernelType,
+      typename = std::enable_if_t<ext::oneapi::experimental::detail::
+                                      HasKernelPropertiesGetMethod<T>::value>>
+  auto get(ext::oneapi::experimental::properties_tag) const {
+    return KernelFunc.get(ext::oneapi::experimental::properties_tag{});
+  }
 };
 
 template <typename TransformedArgType, int Dims, typename KernelType>
@@ -364,6 +374,16 @@ class RoundedRangeKernelWithKH {
       KernelFunc(item, KH);
     }
   }
+
+  // Copy the properties_tag getter from the original kernel to propagate
+  // property(s)
+  template <
+      typename T = KernelType,
+      typename = std::enable_if_t<ext::oneapi::experimental::detail::
+                                      HasKernelPropertiesGetMethod<T>::value>>
+  auto get(ext::oneapi::experimental::properties_tag) const {
+    return KernelFunc.get(ext::oneapi::experimental::properties_tag{});
+  }
 };
 
 using std::enable_if_t;

sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf-2.cpp

@@ -4,10 +4,7 @@
 // kernels when different work-items perform calls to different virtual
 // functions using the same object.
 //
-// TODO: Currently using the -Wno-deprecated-declarations flag due to issue
-// https://github.com/intel/llvm/issues/16839. Remove the flag as well as the
-// variable 'props' once the issue is resolved.
-// RUN: %{build} -o %t.out -Wno-deprecated-declarations %helper-includes
+// RUN: %{build} -o %t.out %helper-includes
 // RUN: %{run} %t.out
 
 #include <sycl/detail/core.hpp>
@@ -76,7 +73,6 @@ int main() try {
   auto *DeviceStorage = sycl::malloc_shared<storage_t>(1, q);
   sycl::range R{1024};
 
-  constexpr oneapi::properties props{oneapi::assume_indirect_calls};
   for (size_t TestCase = 0; TestCase < 2; ++TestCase) {
     std::vector<int> HostData(R.size());
     std::iota(HostData.begin(), HostData.end(), 0);
@@ -91,7 +87,7 @@ int main() try {
 
     q.submit([&](sycl::handler &CGH) {
       sycl::accessor DataAcc(DataStorage, CGH, sycl::read_write);
-      CGH.parallel_for(R, props, KernelFunctor(DeviceStorage, DataAcc));
+      CGH.parallel_for(R, KernelFunctor(DeviceStorage, DataAcc));
     });
 
     BaseOp *Ptr = HostStorage.construct</* ret type = */ BaseOp>(TestCase);

sycl/test-e2e/VirtualFunctions/misc/range-non-uniform-vf.cpp

@@ -4,10 +4,7 @@
 // kernels when different work-items perform a virtual function calls using
 // different objects.
 //
-// TODO: Currently using the -Wno-deprecated-declarations flag due to issue
-// https://github.com/intel/llvm/issues/16839. Remove the flag as well as the
-// variable 'props' once the issue is resolved.
-// RUN: %{build} -o %t.out -Wno-deprecated-declarations %helper-includes
+// RUN: %{build} -o %t.out %helper-includes
 // RUN: %{run} %t.out
 
 #include <sycl/builtins.hpp>
@@ -71,7 +68,6 @@ int main() try {
   auto *DeviceStorage = sycl::malloc_shared<storage_t>(3, q);
   sycl::range R{1024};
 
-  constexpr oneapi::properties props{oneapi::assume_indirect_calls};
   {
     std::vector<float> HostData(R.size());
     for (size_t I = 1; I < HostData.size(); ++I)
@@ -89,7 +85,7 @@ int main() try {
 
     q.submit([&](sycl::handler &CGH) {
       sycl::accessor DataAcc(DataStorage, CGH, sycl::read_write);
-      CGH.parallel_for(R, props, KernelFunctor(DeviceStorage, DataAcc));
+      CGH.parallel_for(R, KernelFunctor(DeviceStorage, DataAcc));
     });
 
     BaseOp *Ptr[] = {HostStorage[0].construct</* ret type = */ BaseOp>(0),

sycl/test-e2e/VirtualFunctions/misc/range-uniform-vf.cpp

@@ -4,10 +4,7 @@
 // kernels when every work-item calls the same virtual function on the same
 // object.
 //
-// TODO: Currently using the -Wno-deprecated-declarations flag due to issue
-// https://github.com/intel/llvm/issues/16839. Remove the flag as well as the
-// variable 'props' once the issue is resolved.
-// RUN: %{build} -o %t.out -Wno-deprecated-declarations %helper-includes
+// RUN: %{build} -o %t.out %helper-includes
 // RUN: %{run} %t.out
 
 #include <sycl/builtins.hpp>
@@ -69,7 +66,6 @@ int main() try {
   auto *DeviceStorage = sycl::malloc_shared<storage_t>(1, q);
   sycl::range R{1024};
 
-  constexpr oneapi::properties props{oneapi::assume_indirect_calls};
   for (unsigned TestCase = 0; TestCase < 3; ++TestCase) {
     std::vector<float> HostData(R.size());
     for (size_t I = 1; I < HostData.size(); ++I)
@@ -85,7 +81,7 @@ int main() try {
 
     q.submit([&](sycl::handler &CGH) {
       sycl::accessor DataAcc(DataStorage, CGH, sycl::read_write);
-      CGH.parallel_for(R, props, KernelFunctor(DeviceStorage, DataAcc));
+      CGH.parallel_for(R, KernelFunctor(DeviceStorage, DataAcc));
     });
 
     auto *Ptr = HostStorage.construct</* ret type = */ BaseOp>(TestCase);