[SYCL] Fix support for classes implicitly converted from items in parallel_for

sycl/include/CL/sycl/handler.hpp

@@ -917,6 +917,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
@@ -935,10 +943,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;
 
@@ -1560,12 +1570,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;
+}