[SYCL] Fix unexpected acceptance of id argument in nd_range parallel_for (#1348)

The kernel callable being invoked from an nd_range parallel_for is accepting an id argument, while it should be nd_item.

After my analysis, I found we check arguments' type for kernel_parallel_for instead of parallel_for. But that check is useless, because the compiler can still find a candidate for kernel_parallel_for with nd_range and id which is a wrong combination.

In my solution, parallel_for with nd_range calls kernel_parallel_for_nd_range(...) which is only available for nd_item.

Signed-off-by: Bing1 Yu <[email protected]>

Author: yubingex007-a11y

sycl/include/CL/sycl/handler.hpp

@@ -428,7 +428,8 @@ class handler {
 
   template <typename KernelName, typename KernelType, int Dims,
             EnableIfNDItem<KernelType, Dims> = 0>
-  __attribute__((sycl_kernel)) void kernel_parallel_for(KernelType KernelFunc) {
+  __attribute__((sycl_kernel)) void
+  kernel_parallel_for_nd_range(KernelType KernelFunc) {
     KernelFunc(detail::Builder::getNDItem<Dims>());
   }
 
@@ -618,7 +619,7 @@ class handler {
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
 #ifdef __SYCL_DEVICE_ONLY__
-    kernel_parallel_for<NameT, KernelType, Dims>(KernelFunc);
+    kernel_parallel_for_nd_range<NameT, KernelType, Dims>(KernelFunc);
 #else
     MNDRDesc.set(std::move(ExecutionRange));
     StoreLambda<NameT, KernelType, Dims>(std::move(KernelFunc));
@@ -855,7 +856,7 @@ class handler {
     using NameT =
         typename detail::get_kernel_name_t<KernelName, KernelType>::name;
 #ifdef __SYCL_DEVICE_ONLY__
-    kernel_parallel_for<NameT, KernelType, Dims>(KernelFunc);
+    kernel_parallel_for_nd_range<NameT, KernelType, Dims>(KernelFunc);
 #else
     MNDRDesc.set(std::move(NDRange));
     MKernel = detail::getSyclObjImpl(std::move(Kernel));

sycl/test/basic_tests/parallel_for_range_host.cpp

@@ -98,8 +98,8 @@ int main() {
   // parallel_for, 30 global, 1(implicit) local -> pass.
   try {
     Q.submit([&](handler &CGH) {
-        CGH.parallel_for<class g>(range<1>(30),
-            [=](nd_item<1>) {});
+      CGH.parallel_for<class g>(range<1>(30),
+                                [=](id<1>) {});
     });
     Q.wait_and_throw();
   } catch (nd_range_error) {

sycl/test/ordered_queue/oq_kernels.cpp

@@ -53,8 +53,8 @@ int main() {
     });
 
     nd_range<1> NDR(range<1>{N}, range<1>{2});
-    q.parallel_for<class NDFoo>(NDR, [=](id<1> ID) {
-      auto i = ID[0];
+    q.parallel_for<class NDFoo>(NDR, [=](nd_item<1> Item) {
+      auto i = Item.get_global_id(0);
       A[i]++;
     });