Skip to content

Revert "[SYCL] Diagnose local accessor use in single_task or parallel_for(range)" #8773

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 1 commit into from
Mar 24, 2023
Merged

Revert "[SYCL] Diagnose local accessor use in single_task or parallel_for(range)" #8773

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
28 changes: 0 additions & 28 deletions sycl/include/sycl/handler.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 <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
Expand Down Expand Up @@ -932,7 +906,6 @@ 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>
Expand Down Expand Up @@ -1414,7 +1387,6 @@ 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 =
Expand Down
5 changes: 2 additions & 3 deletions sycl/source/handler.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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<access::target>(Size & AccessTargetMask);
const access::target AccTarget = static_cast<access::target>(Size & 0x7ff);
switch (AccTarget) {
case access::target::device:
case access::target::constant_buffer: {
Expand Down Expand Up @@ -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<access::target>(Size & AccessTargetMask);
static_cast<access::target>(Size & 0x7ff);
if ((AccTarget == access::target::device ||
AccTarget == access::target::constant_buffer) ||
(AccTarget == access::target::image ||
Expand Down