@@ -1646,7 +1646,7 @@ class __SYCL_EXPORT handler {
16461646 kernel_single_task_wrapper<NameT, KernelType, PropertiesT>(KernelFunc);
16471647#ifndef __SYCL_DEVICE_ONLY__
16481648 throwIfActionIsCreated ();
1649- throwOnLocalAccessorMisuse <KernelName, KernelType>();
1649+ throwOnKernelParameterMisuse <KernelName, KernelType>();
16501650 verifyUsedKernelBundleInternal (
16511651 detail::string_view{detail::getKernelName<NameT>()});
16521652 // No need to check if range is out of INT_MAX limits as it's compile-time
@@ -3565,7 +3565,7 @@ class __SYCL_EXPORT handler {
35653565 // / must not be used in a SYCL kernel function that is invoked via single_task
35663566 // / or via the simple form of parallel_for that takes a range parameter.
35673567 template <typename KernelName, typename KernelType>
3568- void throwOnLocalAccessorMisuse () const {
3568+ void throwOnKernelParameterMisuse () const {
35693569 using NameT =
35703570 typename detail::get_kernel_name_t <KernelName, KernelType>::name;
35713571 for (unsigned I = 0 ; I < detail::getKernelNumParams<NameT>(); ++I) {
@@ -3581,6 +3581,12 @@ class __SYCL_EXPORT handler {
35813581 " A local accessor must not be used in a SYCL kernel function "
35823582 " that is invoked via single_task or via the simple form of "
35833583 " parallel_for that takes a range parameter." );
3584+ if (Kind == detail::kernel_param_kind_t ::kind_work_group_memory)
3585+ throw sycl::exception (
3586+ make_error_code (errc::kernel_argument),
3587+ " A local accessor must not be used in a SYCL kernel function "
3588+ " that is invoked via single_task or via the simple form of "
3589+ " parallel_for that takes a range parameter." );
35843590 }
35853591 }
35863592
0 commit comments