[SYCL] Fix performance regression in parallel_for with using id class

sycl/include/CL/sycl/detail/stl_type_traits.hpp

@@ -78,6 +78,12 @@ struct is_output_iterator<T, output_iterator_requirements<T>> {
   static constexpr bool value = true;
 };
 
+template <typename T, typename U>
+inline constexpr bool is_same_v = std::is_same<T, U>::value;
+
+template <typename T, typename U>
+inline constexpr bool is_convertible_v = std::is_convertible<T, U>::value;
+
 } // namespace detail
 } // namespace sycl
 } // __SYCL_INLINE_NAMESPACE(cl)

sycl/include/CL/sycl/handler.hpp

@@ -917,11 +917,14 @@ class __SYCL_EXPORT handler {
   }
 
   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;
+    using type = typename std::conditional_t<
+        detail::is_same_v<id<Dims>, LambdaArgType>, LambdaArgType,
+        typename std::conditional_t<
+            detail::is_convertible_v<nd_item<Dims>, LambdaArgType>,
+            nd_item<Dims>,
+            typename std::conditional_t<
+                detail::is_convertible_v<item<Dims>, LambdaArgType>, item<Dims>,
+                LambdaArgType>>>;
   };
 
   /// Defines and invokes a SYCL kernel function for the specified range.

sycl/test/basic_tests/parallel_for_type_check.cpp

@@ -0,0 +1,30 @@
+// RUN: %clangxx -fsycl -fsycl-device-only -D__SYCL_INTERNAL_API -O0 -c -emit-llvm -S -o - %s | FileCheck %s
+
+// This test performs basic type check for sycl::id that is used in result type.
+
+#include <CL/sycl.hpp>
+#include <iostream>
+
+int main() {
+  sycl::queue q;
+
+  // Initialize data array
+  const int sz = 16;
+  int data[sz] = {0};
+  for (int i = 0; i < sz; ++i) {
+    data[i] = i;
+  }
+
+  // Check user defined sycl::item wrapper
+  sycl::buffer<int> data_buf(data, sz);
+  q.submit([&](sycl::handler &h) {
+    auto buf_acc = data_buf.get_access<sycl::access::mode::read_write>(h);
+    h.parallel_for(
+        sycl::range<1>{sz},
+        // CHECK: cl{{.*}}sycl{{.*}}detail{{.*}}RoundedRangeKernel{{.*}}id{{.*}}main{{.*}}handler
+        [=](sycl::id<1> item) { buf_acc[item] += 1; });
+  });
+  q.wait();
+
+  return 0;
+}