[SYCL] Fix support for classes implicitly converted from items in parallel_for (#5118)

Section 4.9.4.2.2 of the SYCL 2020 specification says that the function object that represents the SYCL kernel function must take any type implicitly converted from SYCL item, representing the currently executing work-item within the range specified by the range parameter. This PR adds that support.

E2E test: intel/llvm-test-suite#607

Author: aobolensk

sycl/include/CL/sycl/handler.hpp

@@ -921,6 +921,14 @@ class __SYCL_EXPORT handler {
            AccessMode == access::mode::discard_read_write;
   }
 
+  template <int Dims, typename LambdaArgType> struct TransformUserItemType {
+    using type = typename std::conditional<
+        std::is_convertible<nd_item<Dims>, LambdaArgType>::value, nd_item<Dims>,
+        typename std::conditional<
+            std::is_convertible<item<Dims>, LambdaArgType>::value, item<Dims>,
+            LambdaArgType>::type>::type;
+  };
+
   /// Defines and invokes a SYCL kernel function for the specified range.
   ///
   /// The SYCL kernel function is defined as a lambda function or a named
@@ -939,10 +947,12 @@ class __SYCL_EXPORT handler {
     using LambdaArgType = sycl::detail::lambda_arg_type<KernelType, item<Dims>>;
 
     // If 1D kernel argument is an integral type, convert it to sycl::item<1>
-    using TransformedArgType =
-        typename std::conditional<std::is_integral<LambdaArgType>::value &&
-                                      Dims == 1,
-                                  item<Dims>, LambdaArgType>::type;
+    // If user type is convertible from sycl::item/sycl::nd_item, use
+    // sycl::item/sycl::nd_item to transport item information
+    using TransformedArgType = typename std::conditional<
+        std::is_integral<LambdaArgType>::value && Dims == 1, item<Dims>,
+        typename TransformUserItemType<Dims, LambdaArgType>::type>::type;
+
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
 
@@ -1564,12 +1574,17 @@ class __SYCL_EXPORT handler {
     verifyUsedKernelBundle(detail::KernelInfo<NameT>::getName());
     using LambdaArgType =
         sycl::detail::lambda_arg_type<KernelType, nd_item<Dims>>;
+    // If user type is convertible from sycl::item/sycl::nd_item, use
+    // sycl::item/sycl::nd_item to transport item information
+    using TransformedArgType =
+        typename TransformUserItemType<Dims, LambdaArgType>::type;
     (void)ExecutionRange;
-    kernel_parallel_for_wrapper<NameT, LambdaArgType>(KernelFunc);
+    kernel_parallel_for_wrapper<NameT, TransformedArgType>(KernelFunc);
 #ifndef __SYCL_DEVICE_ONLY__
     detail::checkValueRange<Dims>(ExecutionRange);
     MNDRDesc.set(std::move(ExecutionRange));
-    StoreLambda<NameT, KernelType, Dims, LambdaArgType>(std::move(KernelFunc));
+    StoreLambda<NameT, KernelType, Dims, TransformedArgType>(
+        std::move(KernelFunc));
     setType(detail::CG::Kernel);
 #endif
   }

sycl/test/basic_tests/parallel_for_user_types.cpp

@@ -0,0 +1,54 @@
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %RUN_ON_HOST %t.out
+
+// This test performs basic check of supporting user defined class that are
+// implicitly converted from sycl::item/sycl::nd_item in parallel_for.
+
+#include <CL/sycl.hpp>
+#include <iostream>
+
+template <int Dimensions> class item_wrapper {
+public:
+  item_wrapper(sycl::item<Dimensions> it) : m_item(it) {}
+
+private:
+  sycl::item<Dimensions> m_item;
+};
+
+template <int Dimensions> class nd_item_wrapper {
+public:
+  nd_item_wrapper(sycl::nd_item<Dimensions> it) : m_item(it) {}
+
+private:
+  sycl::nd_item<Dimensions> m_item;
+};
+
+template <int Dimensions, typename T> class item_wrapper2 {
+public:
+  item_wrapper2(sycl::item<Dimensions> it) : m_item(it), m_value(T()) {}
+
+private:
+  sycl::item<Dimensions> m_item;
+  T m_value;
+};
+
+template <int Dimensions, typename T> class nd_item_wrapper2 {
+public:
+  nd_item_wrapper2(sycl::nd_item<Dimensions> it) : m_item(it), m_value(T()) {}
+
+private:
+  sycl::nd_item<Dimensions> m_item;
+  T m_value;
+};
+
+int main() {
+  sycl::queue q;
+
+  q.parallel_for(sycl::range<1>{1}, [=](item_wrapper<1> item) {});
+  q.parallel_for(sycl::nd_range<1>{1, 1}, [=](nd_item_wrapper<1> item) {});
+  q.parallel_for(sycl::range<1>{1}, [=](item_wrapper2<1, int> item) {});
+  q.parallel_for(sycl::nd_range<1>{1, 1},
+                 [=](nd_item_wrapper2<1, int> item) {});
+
+  return 0;
+}