diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 742af8ec31c23..b98567b3cf94b 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -354,32 +354,6 @@ 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 - void throwOnLocalAccessorMisuse() const { - using NameT = - typename detail::get_kernel_name_t::name; - using KI = sycl::detail::KernelInfo; - - 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(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 @@ -932,7 +906,6 @@ class __SYCL_EXPORT handler { void parallel_for_lambda_impl(range NumWorkItems, KernelType KernelFunc) { throwIfActionIsCreated(); - throwOnLocalAccessorMisuse(); using LambdaArgType = sycl::detail::lambda_arg_type>; // If 1D kernel argument is an integral type, convert it to sycl::item<1> @@ -1414,7 +1387,6 @@ class __SYCL_EXPORT handler { ext::oneapi::experimental::detail::empty_properties_t> void single_task_lambda_impl(_KERNELFUNCPARAM(KernelFunc)) { throwIfActionIsCreated(); - throwOnLocalAccessorMisuse(); // TODO: Properties may change the kernel function, so in order to avoid // conflicts they should be included in the name. using NameT = diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 4cdae7f896264..7fbb962df6b5a 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -491,8 +491,7 @@ 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(Size & AccessTargetMask); + const access::target AccTarget = static_cast(Size & 0x7ff); switch (AccTarget) { case access::target::device: case access::target::constant_buffer: { @@ -617,7 +616,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(Size & AccessTargetMask); + static_cast(Size & 0x7ff); if ((AccTarget == access::target::device || AccTarget == access::target::constant_buffer) || (AccTarget == access::target::image ||