[SYCL] Restrict nd_range parallel_for to nd_item (#11067)

According to the SYCL 2020 specification:

> The function object that represents the SYCL kernel function must take
one of: 1) a single SYCL nd_item parameter, 2) a single generic
parameter (template parameter or auto) that will be treated as an
nd_item parameter, 3) any other type converted from SYCL nd_item,
representing the currently executing work-item within the range
specified by the nd_range parameter.

However, the current implementation allows also sycl::item and any
argument types convertible from a sycl::item. This commit addresses this
discrepancy by disallowing anything other than sycl::nd_item and all
types convertible from sycl::nd_item in this kind of parallel_for. Since
this may break existing code, this new restriction is guarded by the
SYCL2020_CONFORMANT_APIS preprocessor macro.

---------

Signed-off-by: Larsen, Steffen <[email protected]>
Co-authored-by: Alexey Bader <[email protected]>

Author: steffenlarsen

sycl/include/sycl/handler.hpp

@@ -1191,10 +1191,19 @@ class __SYCL_EXPORT handler {
     verifyUsedKernelBundle(detail::KernelInfo<NameT>::getName());
     using LambdaArgType =
         sycl::detail::lambda_arg_type<KernelType, nd_item<Dims>>;
+#ifdef SYCL2020_CONFORMANT_APIS
+    static_assert(
+        std::is_convertible_v<sycl::nd_item<Dims>, LambdaArgType>,
+        "Kernel argument of a sycl::parallel_for with sycl::nd_range "
+        "must be either sycl::nd_item or be convertible from sycl::nd_item");
+    using TransformedArgType = sycl::nd_item<Dims>;
+#else
     // 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;
+#endif
+
     (void)ExecutionRange;
     kernel_parallel_for_wrapper<NameT, TransformedArgType, KernelType,
                                 PropertiesT>(KernelFunc);

sycl/test/basic_tests/handler/parallel_for_arg_restrictions.cpp

@@ -0,0 +1,103 @@
+// RUN: %clangxx -fsycl -fsyntax-only -ferror-limit=0 -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning,error %s
+// RUN: %clangxx -fsycl -fsyntax-only -ferror-limit=0 -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning,error %s -DSYCL2020_CONFORMANT_APIS
+
+// Tests the requirements on the first argument in a kernel lambda.
+// TODO: Some of the behavior is currently guarded behind
+//       SYCL2020_CONFORMANT_APIS. Remove the definition when this behavior has
+//       been promoted.
+
+#include <sycl/sycl.hpp>
+
+template <int Dims> struct ConvertibleFromItem {
+  ConvertibleFromItem(sycl::item<Dims>) {}
+};
+
+template <int Dims> struct ConvertibleFromNDItem {
+  ConvertibleFromNDItem(sycl::nd_item<Dims>) {}
+};
+
+int main() {
+  sycl::queue Q;
+
+// TODO: Remove this when the guarded behavior is promoted.
+#ifdef SYCL2020_CONFORMANT_APIS
+  // ND-range parallel_for with item.
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error@sycl/handler.hpp:* {{Kernel argument of a sycl::parallel_for with sycl::nd_range must be either sycl::nd_item or be convertible from sycl::nd_item}}
+    CGH.parallel_for(sycl::nd_range{sycl::range{1}, sycl::range{1}},
+                     [=](sycl::item<1>) {});
+  });
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error@sycl/handler.hpp:* {{Kernel argument of a sycl::parallel_for with sycl::nd_range must be either sycl::nd_item or be convertible from sycl::nd_item}}
+    CGH.parallel_for(sycl::nd_range{sycl::range{1, 1}, sycl::range{1, 1}},
+                     [=](sycl::item<2>) {});
+  });
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error@sycl/handler.hpp:* {{Kernel argument of a sycl::parallel_for with sycl::nd_range must be either sycl::nd_item or be convertible from sycl::nd_item}}
+    CGH.parallel_for(sycl::nd_range{sycl::range{1, 1, 1}, sycl::range{1, 1, 1}},
+                     [=](sycl::item<3>) {});
+  });
+
+  // ND-range parallel_for with id.
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error@sycl/handler.hpp:* {{Kernel argument of a sycl::parallel_for with sycl::nd_range must be either sycl::nd_item or be convertible from sycl::nd_item}}
+    CGH.parallel_for(sycl::nd_range{sycl::range{1}, sycl::range{1}},
+                     [=](sycl::id<1>) {});
+  });
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error@sycl/handler.hpp:* {{Kernel argument of a sycl::parallel_for with sycl::nd_range must be either sycl::nd_item or be convertible from sycl::nd_item}}
+    CGH.parallel_for(sycl::nd_range{sycl::range{1, 1}, sycl::range{1, 1}},
+                     [=](sycl::id<2>) {});
+  });
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error@sycl/handler.hpp:* {{Kernel argument of a sycl::parallel_for with sycl::nd_range must be either sycl::nd_item or be convertible from sycl::nd_item}}
+    CGH.parallel_for(sycl::nd_range{sycl::range{1, 1, 1}, sycl::range{1, 1, 1}},
+                     [=](sycl::id<3>) {});
+  });
+
+  // ND-range parallel_for with argument that is convertible from item.
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error@sycl/handler.hpp:* {{Kernel argument of a sycl::parallel_for with sycl::nd_range must be either sycl::nd_item or be convertible from sycl::nd_item}}
+    CGH.parallel_for(sycl::nd_range{sycl::range{1}, sycl::range{1}},
+                     [=](ConvertibleFromItem<1>) {});
+  });
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error@sycl/handler.hpp:* {{Kernel argument of a sycl::parallel_for with sycl::nd_range must be either sycl::nd_item or be convertible from sycl::nd_item}}
+    CGH.parallel_for(sycl::nd_range{sycl::range{1, 1}, sycl::range{1, 1}},
+                     [=](ConvertibleFromItem<2>) {});
+  });
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error@sycl/handler.hpp:* {{Kernel argument of a sycl::parallel_for with sycl::nd_range must be either sycl::nd_item or be convertible from sycl::nd_item}}
+    CGH.parallel_for(sycl::nd_range{sycl::range{1, 1, 1}, sycl::range{1, 1, 1}},
+                     [=](ConvertibleFromItem<3>) {});
+  });
+#endif // SYCL2020_CONFORMANT_APIS
+
+  // Range parallel_for with nd_item.
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error@sycl/handler.hpp:* {{Kernel argument cannot have a sycl::nd_item type in sycl::parallel_for with sycl::range}}
+    CGH.parallel_for(sycl::range{1}, [=](sycl::nd_item<1>) {});
+  });
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error@sycl/handler.hpp:* {{Kernel argument cannot have a sycl::nd_item type in sycl::parallel_for with sycl::range}}
+    CGH.parallel_for(sycl::range{1, 1}, [=](sycl::nd_item<2>) {});
+  });
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error@sycl/handler.hpp:* {{Kernel argument cannot have a sycl::nd_item type in sycl::parallel_for with sycl::range}}
+    CGH.parallel_for(sycl::range{1, 1, 1}, [=](sycl::nd_item<3>) {});
+  });
+
+  // Range parallel_for with argument that is convertible from nd_item.
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error@sycl/handler.hpp:* {{Kernel argument cannot have a sycl::nd_item type in sycl::parallel_for with sycl::range}}
+    CGH.parallel_for(sycl::range{1}, [=](ConvertibleFromNDItem<1>) {});
+  });
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error@sycl/handler.hpp:* {{Kernel argument cannot have a sycl::nd_item type in sycl::parallel_for with sycl::range}}
+    CGH.parallel_for(sycl::range{1, 1}, [=](ConvertibleFromNDItem<2>) {});
+  });
+  Q.submit([&](sycl::handler &CGH) {
+    // expected-error@sycl/handler.hpp:* {{Kernel argument cannot have a sycl::nd_item type in sycl::parallel_for with sycl::range}}
+    CGH.parallel_for(sycl::range{1, 1, 1}, [=](ConvertibleFromNDItem<3>) {});
+  });
+}