[SYCL] Diagnose local accessor use in single_task or parallel_for(range) (#8581)

* According to [local
accessors](https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:accessor.local)
of the SYCL specification, a local accessor must not be used in a SYCL
kernel function that is invoked via single_task or via the simple form
of parallel_for that takes a range parameter.
* Add test.

---------

Co-authored-by: Alexey Bader <[email protected]>

Author: mmoadeli

sycl/include/sycl/handler.hpp

@@ -354,6 +354,32 @@ class __SYCL_EXPORT handler {
                                 PI_ERROR_INVALID_OPERATION);
   }
 
+  constexpr static int AccessTargetMask = 0x7ff;
+  /// According to section 4.7.6.11. of the SYCL specification, a local accessor
+  /// must not be used in a SYCL kernel function that is invoked via single_task
+  /// or via the simple form of parallel_for that takes a range parameter.
+  template <typename KernelName, typename KernelType>
+  void throwOnLocalAccessorMisuse() const {
+    using NameT =
+        typename detail::get_kernel_name_t<KernelName, KernelType>::name;
+    using KI = sycl::detail::KernelInfo<NameT>;
+
+    auto *KernelArgs = &KI::getParamDesc(0);
+
+    for (unsigned I = 0; I < KI::getNumParams(); ++I) {
+      const detail::kernel_param_kind_t &Kind = KernelArgs[I].kind;
+      const access::target AccTarget =
+          static_cast<access::target>(KernelArgs[I].info & AccessTargetMask);
+      if ((Kind == detail::kernel_param_kind_t::kind_accessor) &&
+          (AccTarget == target::local))
+        throw sycl::exception(
+            make_error_code(errc::kernel_argument),
+            "A local accessor must not be used in a SYCL kernel function "
+            "that is invoked via single_task or via the simple form of "
+            "parallel_for that takes a range parameter.");
+    }
+  }
+
   /// Extracts and prepares kernel arguments from the lambda using integration
   /// header.
   void
@@ -906,6 +932,7 @@ class __SYCL_EXPORT handler {
   void parallel_for_lambda_impl(range<Dims> NumWorkItems,
                                 KernelType KernelFunc) {
     throwIfActionIsCreated();
+    throwOnLocalAccessorMisuse<KernelName, KernelType>();
     using LambdaArgType = sycl::detail::lambda_arg_type<KernelType, item<Dims>>;
 
     // If 1D kernel argument is an integral type, convert it to sycl::item<1>
@@ -1387,6 +1414,7 @@ class __SYCL_EXPORT handler {
                 ext::oneapi::experimental::detail::empty_properties_t>
   void single_task_lambda_impl(_KERNELFUNCPARAM(KernelFunc)) {
     throwIfActionIsCreated();
+    throwOnLocalAccessorMisuse<KernelName, KernelType>();
     // TODO: Properties may change the kernel function, so in order to avoid
     //       conflicts they should be included in the name.
     using NameT =

sycl/source/handler.cpp

@@ -491,7 +491,8 @@ void handler::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind,
   case kernel_param_kind_t::kind_accessor: {
     // For args kind of accessor Size is information about accessor.
     // The first 11 bits of Size encodes the accessor target.
-    const access::target AccTarget = static_cast<access::target>(Size & 0x7ff);
+    const access::target AccTarget =
+        static_cast<access::target>(Size & AccessTargetMask);
     switch (AccTarget) {
     case access::target::device:
     case access::target::constant_buffer: {
@@ -616,7 +617,7 @@ void handler::extractArgsAndReqsFromLambda(
       // For args kind of accessor Size is information about accessor.
       // The first 11 bits of Size encodes the accessor target.
       const access::target AccTarget =
-          static_cast<access::target>(Size & 0x7ff);
+          static_cast<access::target>(Size & AccessTargetMask);
       if ((AccTarget == access::target::device ||
            AccTarget == access::target::constant_buffer) ||
           (AccTarget == access::target::image ||