Skip to content

Commit d7a5ec0

Browse files
mmoadelibaderMahmoud Moadelisteffenlarsen
authored
[SYCL] Disallow local accessor misuse (#9269)
* According to the [specification](https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:accessor.local) 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. * Modifies existing misuse in tests. * This PR was once reviewed and merged in, but it was reverted as needed AccessTargetMask as an exported symbol on Windows to be added. --------- Co-authored-by: Alexey Bader <[email protected]> Co-authored-by: Mahmoud Moadeli <[email protected]> Co-authored-by: Steffen Larsen <[email protected]>
1 parent 854ab7e commit d7a5ec0

File tree

13 files changed

+215
-140
lines changed

13 files changed

+215
-140
lines changed

‎sycl/include/sycl/handler.hpp

Lines changed: 28 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -365,6 +365,32 @@ class __SYCL_EXPORT handler {
365365
PI_ERROR_INVALID_OPERATION);
366366
}
367367

368+
constexpr static int AccessTargetMask = 0x7ff;
369+
/// According to section 4.7.6.11. of the SYCL specification, a local accessor
370+
/// must not be used in a SYCL kernel function that is invoked via single_task
371+
/// or via the simple form of parallel_for that takes a range parameter.
372+
template <typename KernelName, typename KernelType>
373+
void throwOnLocalAccessorMisuse() const {
374+
using NameT =
375+
typename detail::get_kernel_name_t<KernelName, KernelType>::name;
376+
using KI = sycl::detail::KernelInfo<NameT>;
377+
378+
auto *KernelArgs = &KI::getParamDesc(0);
379+
380+
for (unsigned I = 0; I < KI::getNumParams(); ++I) {
381+
const detail::kernel_param_kind_t &Kind = KernelArgs[I].kind;
382+
const access::target AccTarget =
383+
static_cast<access::target>(KernelArgs[I].info & AccessTargetMask);
384+
if ((Kind == detail::kernel_param_kind_t::kind_accessor) &&
385+
(AccTarget == target::local))
386+
throw sycl::exception(
387+
make_error_code(errc::kernel_argument),
388+
"A local accessor must not be used in a SYCL kernel function "
389+
"that is invoked via single_task or via the simple form of "
390+
"parallel_for that takes a range parameter.");
391+
}
392+
}
393+
368394
/// Extracts and prepares kernel arguments from the lambda using integration
369395
/// header.
370396
void
@@ -940,6 +966,7 @@ class __SYCL_EXPORT handler {
940966
void parallel_for_lambda_impl(range<Dims> NumWorkItems, PropertiesT Props,
941967
KernelType KernelFunc) {
942968
throwIfActionIsCreated();
969+
throwOnLocalAccessorMisuse<KernelName, KernelType>();
943970
using LambdaArgType = sycl::detail::lambda_arg_type<KernelType, item<Dims>>;
944971

945972
// If 1D kernel argument is an integral type, convert it to sycl::item<1>
@@ -1435,6 +1462,7 @@ class __SYCL_EXPORT handler {
14351462
void single_task_lambda_impl(PropertiesT Props,
14361463
_KERNELFUNCPARAM(KernelFunc)) {
14371464
throwIfActionIsCreated();
1465+
throwOnLocalAccessorMisuse<KernelName, KernelType>();
14381466
// TODO: Properties may change the kernel function, so in order to avoid
14391467
// conflicts they should be included in the name.
14401468
using NameT =

‎sycl/source/handler.cpp

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -512,7 +512,8 @@ void handler::processArg(void *Ptr, const detail::kernel_param_kind_t &Kind,
512512
case kernel_param_kind_t::kind_accessor: {
513513
// For args kind of accessor Size is information about accessor.
514514
// The first 11 bits of Size encodes the accessor target.
515-
const access::target AccTarget = static_cast<access::target>(Size & 0x7ff);
515+
const access::target AccTarget =
516+
static_cast<access::target>(Size & AccessTargetMask);
516517
switch (AccTarget) {
517518
case access::target::device:
518519
case access::target::constant_buffer: {
@@ -637,7 +638,7 @@ void handler::extractArgsAndReqsFromLambda(
637638
// For args kind of accessor Size is information about accessor.
638639
// The first 11 bits of Size encodes the accessor target.
639640
const access::target AccTarget =
640-
static_cast<access::target>(Size & 0x7ff);
641+
static_cast<access::target>(Size & AccessTargetMask);
641642
if ((AccTarget == access::target::device ||
642643
AccTarget == access::target::constant_buffer) ||
643644
(AccTarget == access::target::image ||

‎sycl/test-e2e/Basic/accessor/accessor.cpp

Lines changed: 59 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -107,7 +107,8 @@ template <typename T> void TestAccSizeFuncs(const std::vector<T> &vec) {
107107
q.submit([&](sycl::handler &cgh) {
108108
sycl::accessor accRes(bufRes, cgh);
109109
sycl::local_accessor<T, 1> locAcc(vec.size(), cgh);
110-
cgh.single_task([=]() { test(accRes, locAcc); });
110+
cgh.parallel_for(sycl::nd_range<1>{1, 1},
111+
[=](sycl::nd_item<1>) { test(accRes, locAcc); });
111112
});
112113
q.wait();
113114
}
@@ -120,7 +121,7 @@ template <typename GlobAcc, typename LocAcc>
120121
void testLocalAccItersImpl(sycl::handler &cgh, GlobAcc &globAcc, LocAcc &locAcc,
121122
bool testConstIter) {
122123
if (testConstIter) {
123-
cgh.single_task([=]() {
124+
cgh.parallel_for(sycl::nd_range<1>{1, 1}, [=](sycl::nd_item<1>) {
124125
size_t Idx = 0;
125126
for (auto &It : locAcc) {
126127
It = globAcc[Idx++];
@@ -133,7 +134,7 @@ void testLocalAccItersImpl(sycl::handler &cgh, GlobAcc &globAcc, LocAcc &locAcc,
133134
globAcc[Idx--] += *It;
134135
});
135136
} else {
136-
cgh.single_task([=]() {
137+
cgh.parallel_for(sycl::nd_range<1>{1, 1}, [=](sycl::nd_item<1>) {
137138
size_t Idx = 0;
138139
for (auto It = locAcc.begin(); It != locAcc.end(); It++)
139140
*It = globAcc[Idx++] * 2;
@@ -991,10 +992,11 @@ int main() {
991992
sycl::accessor acc1(buf1, cgh);
992993
sycl::accessor acc2(buf2, cgh);
993994
acc1.swap(acc2);
994-
cgh.single_task([=]() {
995-
acc1[15] = 4;
996-
acc2[7] = 4;
997-
});
995+
cgh.parallel_for<class swap1>(sycl::nd_range<1>{1, 1},
996+
[=](sycl::nd_item<1>) {
997+
acc1[15] = 4;
998+
acc2[7] = 4;
999+
});
9981000
});
9991001
}
10001002
assert(vec1[7] == 4 && vec2[15] == 4);
@@ -1012,10 +1014,11 @@ int main() {
10121014
sycl::accessor acc2(buf2, cgh);
10131015
sycl::local_accessor<int, 1> locAcc1(8, cgh), locAcc2(16, cgh);
10141016
locAcc1.swap(locAcc2);
1015-
cgh.single_task([=]() {
1016-
acc1[0] = locAcc1.size();
1017-
acc2[0] = locAcc2.size();
1018-
});
1017+
cgh.parallel_for<class swap2>(sycl::nd_range<1>{1, 1},
1018+
[=](sycl::nd_item<1>) {
1019+
acc1[0] = locAcc1.size();
1020+
acc2[0] = locAcc2.size();
1021+
});
10191022
});
10201023
}
10211024
assert(size1 == 16 && size2 == 8);
@@ -1082,20 +1085,55 @@ int main() {
10821085
// Explicit block to prompt copy-back to Data
10831086
{
10841087
sycl::buffer<int, 1> DataBuffer(&Data, sycl::range<1>(1));
1085-
10861088
Queue.submit([&](sycl::handler &CGH) {
10871089
sycl::accessor<int, 0> Acc(DataBuffer, CGH);
10881090
sycl::local_accessor<int, 0> LocalAcc(CGH);
1089-
CGH.single_task<class local_acc_0_dim_assignment>([=]() {
1090-
LocalAcc = 64;
1091-
Acc = LocalAcc;
1092-
});
1091+
CGH.parallel_for<class copyblock>(sycl::nd_range<1>{1, 1},
1092+
[=](sycl::nd_item<1>) {
1093+
LocalAcc = 64;
1094+
Acc = LocalAcc;
1095+
});
10931096
});
10941097
}
10951098

10961099
assert(Data == 64);
10971100
}
10981101

1102+
// Throws exception on local_accessors used in single_task
1103+
{
1104+
constexpr static int size = 1;
1105+
sycl::queue Queue;
1106+
1107+
try {
1108+
Queue.submit([&](sycl::handler &cgh) {
1109+
auto local_acc = sycl::local_accessor<int, 1>({size}, cgh);
1110+
cgh.single_task<class local_acc_exception>([=]() { (void)local_acc; });
1111+
});
1112+
assert(0 && "local accessor must not be used in single task.");
1113+
} catch (sycl::exception e) {
1114+
std::cout << "SYCL exception caught: " << e.what() << std::endl;
1115+
}
1116+
}
1117+
1118+
// Throws exception on local_accessors used in parallel_for taking a range
1119+
// parameter.
1120+
{
1121+
constexpr static int size = 1;
1122+
sycl::queue Queue;
1123+
1124+
try {
1125+
Queue.submit([&](sycl::handler &cgh) {
1126+
auto local_acc = sycl::local_accessor<int, 1>({size}, cgh);
1127+
cgh.parallel_for<class parallel_for_exception>(
1128+
sycl::range<1>{size}, [=](sycl::id<1> ID) { (void)local_acc; });
1129+
});
1130+
assert(0 &&
1131+
"local accessor must not be used in parallel for with range.");
1132+
} catch (sycl::exception e) {
1133+
std::cout << "SYCL exception caught: " << e.what() << std::endl;
1134+
}
1135+
}
1136+
10991137
// local_accessor::operator& and local_accessor::operator[] with const DataT
11001138
{
11011139
using AccT_zero = sycl::local_accessor<const int, 0>;
@@ -1106,10 +1144,11 @@ int main() {
11061144
queue.submit([&](sycl::handler &cgh) {
11071145
AccT_zero acc_zero(cgh);
11081146
AccT_non_zero acc_non_zero(sycl::range<1>(5), cgh);
1109-
cgh.single_task([=] {
1110-
const int &ref_zero = acc_zero;
1111-
const int &ref_non_zero = acc_non_zero[0];
1112-
});
1147+
cgh.parallel_for<class local_acc_const_type>(
1148+
sycl::nd_range<1>{1, 1}, [=](sycl::nd_item<1> ID) {
1149+
const int &ref_zero = acc_zero;
1150+
const int &ref_non_zero = acc_non_zero[0];
1151+
});
11131152
});
11141153
}
11151154
}

‎sycl/test-e2e/Basic/multi_ptr.hpp

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -104,11 +104,11 @@ template <typename T, access::decorated IsDecorated> void testMultPtr() {
104104
local_accessor<T, 1> localAccessor(numOfItems, cgh);
105105

106106
cgh.parallel_for<class testMultPtrKernel<
107-
T, IsDecorated>>(range<1>{10}, [=](id<1> wiID) {
107+
T, IsDecorated>>(nd_range<1>{10, 10}, [=](nd_item<1> wiID) {
108108
T private_data[10];
109109
for (size_t i = 0; i < 10; ++i)
110110
private_data[i] = 0;
111-
localAccessor[wiID] = 0;
111+
localAccessor[wiID.get_local_id()] = 0;
112112

113113
auto ptr_1 =
114114
multi_ptr<T, access::address_space::global_space, IsDecorated>(
@@ -166,8 +166,8 @@ template <typename T, access::decorated IsDecorated> void testMultPtr() {
166166
global_ptr<void, IsDecorated> ptr_12 =
167167
global_ptr<void, IsDecorated>(ptr_11);
168168

169-
innerFunc<T, IsDecorated>(wiID.get(0), ptr_1, ptr_2, ptr_3, ptr_4,
170-
ptr_5, local_ptr, priv_ptr);
169+
innerFunc<T, IsDecorated>(wiID.get_local_id().get(0), ptr_1, ptr_2,
170+
ptr_3, ptr_4, ptr_5, local_ptr, priv_ptr);
171171
});
172172
});
173173
}
@@ -201,8 +201,8 @@ void testMultPtrArrowOperator() {
201201
access::placeholder::false_t>
202202
accessorData_3(bufferData_3, cgh);
203203

204-
cgh.single_task<class testMultPtrArrowOperatorKernel<T, IsDecorated>>(
205-
[=]() {
204+
cgh.parallel_for<class testMultPtrArrowOperatorKernel<T, IsDecorated>>(
205+
sycl::nd_range<1>{1, 1}, [=](sycl::nd_item<1>) {
206206
point<T> private_val = 0;
207207

208208
auto ptr_1 =

0 commit comments

Comments
 (0)