From a92bbec0383e22e350583cf27ea894550deb161f Mon Sep 17 00:00:00 2001 From: Steffen Larsen Date: Wed, 6 Sep 2023 18:18:32 +0100 Subject: [PATCH] [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 Co-authored-by: Alexey Bader --- sycl/include/sycl/handler.hpp | 9 ++ .../handler/parallel_for_arg_restrictions.cpp | 103 ++++++++++++++++++ 2 files changed, 112 insertions(+) create mode 100644 sycl/test/basic_tests/handler/parallel_for_arg_restrictions.cpp diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 5b262e221f494..d7b1f8017aeb4 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -1191,10 +1191,19 @@ class __SYCL_EXPORT handler { verifyUsedKernelBundle(detail::KernelInfo::getName()); using LambdaArgType = sycl::detail::lambda_arg_type>; +#ifdef SYCL2020_CONFORMANT_APIS + static_assert( + std::is_convertible_v, 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; +#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::type; +#endif + (void)ExecutionRange; kernel_parallel_for_wrapper(KernelFunc); diff --git a/sycl/test/basic_tests/handler/parallel_for_arg_restrictions.cpp b/sycl/test/basic_tests/handler/parallel_for_arg_restrictions.cpp new file mode 100644 index 0000000000000..031aa811c64ed --- /dev/null +++ b/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 + +template struct ConvertibleFromItem { + ConvertibleFromItem(sycl::item) {} +}; + +template struct ConvertibleFromNDItem { + ConvertibleFromNDItem(sycl::nd_item) {} +}; + +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>) {}); + }); +}