Skip to content

Commit 08f8656

Browse files
s-kanaevs-kanaev
authored
[SYCL] Throw exception if range/offset of kernel execution exceeds INT_MAX (#1713)
The exception will only be thrown if `__SYCL_ID_QUERIES_FIT_IN_INT__` macro is defined. Signed-off-by: Sergey Kanaev <sergey.kanaev@intel.com>
1 parent 72edbd9 commit 08f8656

File tree

2 files changed

+251
-2
lines changed

2 files changed

+251
-2
lines changed

‎sycl/include/CL/sycl/handler.hpp

Lines changed: 74 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -25,6 +25,7 @@
2525

2626
#include <algorithm>
2727
#include <functional>
28+
#include <limits>
2829
#include <memory>
2930
#include <type_traits>
3031

@@ -128,6 +129,39 @@ struct check_fn_signature<F, RetT(Args...)> {
128129

129130
__SYCL_EXPORT device getDeviceFromHandler(handler &);
130131

132+
#if defined(__SYCL_ID_QUERIES_FIT_IN_INT__)
133+
template <typename T> struct NotIntMsg;
134+
135+
// TODO reword for "`fsycl-id-queries-fit-in-int' optimization flag." when
136+
// implemented
137+
template <int Dims> struct NotIntMsg<range<Dims>> {
138+
constexpr static char *Msg = "Provided range is out of integer limits. "
139+
"Pass `-U__SYCL_ID_QUERIES_FIT_IN_INT__' to "
140+
"disable range check.";
141+
};
142+
143+
template <int Dims> struct NotIntMsg<id<Dims>> {
144+
constexpr static char *Msg = "Provided offset is out of integer limits. "
145+
"Pass `-U__SYCL_ID_QUERIES_FIT_IN_INT__' to "
146+
"disable offset check.";
147+
};
148+
#endif
149+
150+
template <int Dims, typename T>
151+
typename std::enable_if<std::is_same<T, range<Dims>>::value ||
152+
std::is_same<T, id<Dims>>::value>::type
153+
checkValueRange(const T &V) {
154+
#if defined(__SYCL_ID_QUERIES_FIT_IN_INT__)
155+
static constexpr size_t Limit =
156+
static_cast<size_t>((std::numeric_limits<int>::max)());
157+
for (size_t Dim = 0; Dim < Dims; ++Dim)
158+
if (V[Dim] > Limit)
159+
throw runtime_error(NotIntMsg<T>::Msg, PI_INVALID_VALUE);
160+
#else
161+
(void)V;
162+
#endif
163+
}
164+
131165
} // namespace detail
132166

133167
namespace intel {
@@ -770,6 +804,8 @@ class __SYCL_EXPORT handler {
770804
#ifdef __SYCL_DEVICE_ONLY__
771805
kernel_single_task<NameT>(KernelFunc);
772806
#else
807+
// No need to check if range is out of INT_MAX limits as it's compile-time
808+
// known constant.
773809
MNDRDesc.set(range<1>{1});
774810

775811
StoreLambda<NameT, KernelType, /*Dims*/ 0, void>(KernelFunc);
@@ -798,6 +834,7 @@ class __SYCL_EXPORT handler {
798834
(void)NumWorkItems;
799835
kernel_parallel_for<NameT, KernelType, Dims>(KernelFunc);
800836
#else
837+
detail::checkValueRange<Dims>(NumWorkItems);
801838
MNDRDesc.set(std::move(NumWorkItems));
802839
StoreLambda<NameT, KernelType, Dims>(std::move(KernelFunc));
803840
MCGType = detail::CG::KERNEL;
@@ -810,6 +847,8 @@ class __SYCL_EXPORT handler {
810847
/// named function object type.
811848
template <typename FuncT> void run_on_host_intel(FuncT Func) {
812849
throwIfActionIsCreated();
850+
// No need to check if range is out of INT_MAX limits as it's compile-time
851+
// known constant
813852
MNDRDesc.set(range<1>{1});
814853

815854
MArgs = std::move(MAssociatedAccesors);
@@ -856,6 +895,8 @@ class __SYCL_EXPORT handler {
856895
(void)WorkItemOffset;
857896
kernel_parallel_for<NameT, KernelType, Dims>(KernelFunc);
858897
#else
898+
detail::checkValueRange<Dims>(NumWorkItems);
899+
detail::checkValueRange<Dims>(WorkItemOffset);
859900
MNDRDesc.set(std::move(NumWorkItems), std::move(WorkItemOffset));
860901
StoreLambda<NameT, KernelType, Dims>(std::move(KernelFunc));
861902
MCGType = detail::CG::KERNEL;
@@ -884,6 +925,9 @@ class __SYCL_EXPORT handler {
884925
(void)ExecutionRange;
885926
kernel_parallel_for_nd_range<NameT, KernelType, Dims>(KernelFunc);
886927
#else
928+
detail::checkValueRange<Dims>(ExecutionRange.get_global_range());
929+
detail::checkValueRange<Dims>(ExecutionRange.get_local_range());
930+
detail::checkValueRange<Dims>(ExecutionRange.get_offset());
887931
MNDRDesc.set(std::move(ExecutionRange));
888932
StoreLambda<NameT, KernelType, Dims>(std::move(KernelFunc));
889933
MCGType = detail::CG::KERNEL;
@@ -1053,6 +1097,7 @@ class __SYCL_EXPORT handler {
10531097
(void)NumWorkGroups;
10541098
kernel_parallel_for_work_group<NameT, KernelType, Dims>(KernelFunc);
10551099
#else
1100+
detail::checkValueRange<Dims>(NumWorkGroups);
10561101
MNDRDesc.setNumWorkGroups(NumWorkGroups);
10571102
StoreLambda<NameT, KernelType, Dims>(std::move(KernelFunc));
10581103
MCGType = detail::CG::KERNEL;
@@ -1084,7 +1129,12 @@ class __SYCL_EXPORT handler {
10841129
(void)WorkGroupSize;
10851130
kernel_parallel_for_work_group<NameT, KernelType, Dims>(KernelFunc);
10861131
#else
1087-
MNDRDesc.set(nd_range<Dims>(NumWorkGroups * WorkGroupSize, WorkGroupSize));
1132+
nd_range<Dims> ExecRange =
1133+
nd_range<Dims>(NumWorkGroups * WorkGroupSize, WorkGroupSize);
1134+
detail::checkValueRange<Dims>(ExecRange.get_global_range());
1135+
detail::checkValueRange<Dims>(ExecRange.get_local_range());
1136+
detail::checkValueRange<Dims>(ExecRange.get_offset());
1137+
MNDRDesc.set(std::move(ExecRange));
10881138
StoreLambda<NameT, KernelType, Dims>(std::move(KernelFunc));
10891139
MCGType = detail::CG::KERNEL;
10901140
#endif // __SYCL_DEVICE_ONLY__
@@ -1099,6 +1149,8 @@ class __SYCL_EXPORT handler {
10991149
void single_task(kernel Kernel) {
11001150
throwIfActionIsCreated();
11011151
verifyKernelInvoc(Kernel);
1152+
// No need to check if range is out of INT_MAX limits as it's compile-time
1153+
// known constant
11021154
MNDRDesc.set(range<1>{1});
11031155
MKernel = detail::getSyclObjImpl(std::move(Kernel));
11041156
MCGType = detail::CG::KERNEL;
@@ -1117,6 +1169,7 @@ class __SYCL_EXPORT handler {
11171169
throwIfActionIsCreated();
11181170
verifyKernelInvoc(Kenrel);
11191171
MKernel = detail::getSyclObjImpl(std::move(Kenrel));
1172+
detail::checkValueRange<Dims>(NumWorkItems);
11201173
MNDRDesc.set(std::move(NumWorkItems));
11211174
MCGType = detail::CG::KERNEL;
11221175
extractArgsAndReqs();
@@ -1136,6 +1189,8 @@ class __SYCL_EXPORT handler {
11361189
throwIfActionIsCreated();
11371190
verifyKernelInvoc(Kernel);
11381191
MKernel = detail::getSyclObjImpl(std::move(Kernel));
1192+
detail::checkValueRange<Dims>(NumWorkItems);
1193+
detail::checkValueRange<Dims>(WorkItemOffset);
11391194
MNDRDesc.set(std::move(NumWorkItems), std::move(WorkItemOffset));
11401195
MCGType = detail::CG::KERNEL;
11411196
extractArgsAndReqs();
@@ -1153,6 +1208,9 @@ class __SYCL_EXPORT handler {
11531208
throwIfActionIsCreated();
11541209
verifyKernelInvoc(Kernel);
11551210
MKernel = detail::getSyclObjImpl(std::move(Kernel));
1211+
detail::checkValueRange<Dims>(NDRange.get_global_range());
1212+
detail::checkValueRange<Dims>(NDRange.get_local_range());
1213+
detail::checkValueRange<Dims>(NDRange.get_offset());
11561214
MNDRDesc.set(std::move(NDRange));
11571215
MCGType = detail::CG::KERNEL;
11581216
extractArgsAndReqs();
@@ -1173,6 +1231,8 @@ class __SYCL_EXPORT handler {
11731231
(void)Kernel;
11741232
kernel_single_task<NameT>(KernelFunc);
11751233
#else
1234+
// No need to check if range is out of INT_MAX limits as it's compile-time
1235+
// known constant
11761236
MNDRDesc.set(range<1>{1});
11771237
MKernel = detail::getSyclObjImpl(std::move(Kernel));
11781238
MCGType = detail::CG::KERNEL;
@@ -1211,6 +1271,7 @@ class __SYCL_EXPORT handler {
12111271
(void)NumWorkItems;
12121272
kernel_parallel_for<NameT, KernelType, Dims>(KernelFunc);
12131273
#else
1274+
detail::checkValueRange<Dims>(NumWorkItems);
12141275
MNDRDesc.set(std::move(NumWorkItems));
12151276
MKernel = detail::getSyclObjImpl(std::move(Kernel));
12161277
MCGType = detail::CG::KERNEL;
@@ -1243,6 +1304,8 @@ class __SYCL_EXPORT handler {
12431304
(void)WorkItemOffset;
12441305
kernel_parallel_for<NameT, KernelType, Dims>(KernelFunc);
12451306
#else
1307+
detail::checkValueRange<Dims>(NumWorkItems);
1308+
detail::checkValueRange<Dims>(WorkItemOffset);
12461309
MNDRDesc.set(std::move(NumWorkItems), std::move(WorkItemOffset));
12471310
MKernel = detail::getSyclObjImpl(std::move(Kernel));
12481311
MCGType = detail::CG::KERNEL;
@@ -1274,6 +1337,9 @@ class __SYCL_EXPORT handler {
12741337
(void)NDRange;
12751338
kernel_parallel_for_nd_range<NameT, KernelType, Dims>(KernelFunc);
12761339
#else
1340+
detail::checkValueRange<Dims>(NDRange.get_global_range());
1341+
detail::checkValueRange<Dims>(NDRange.get_local_range());
1342+
detail::checkValueRange<Dims>(NDRange.get_offset());
12771343
MNDRDesc.set(std::move(NDRange));
12781344
MKernel = detail::getSyclObjImpl(std::move(Kernel));
12791345
MCGType = detail::CG::KERNEL;
@@ -1309,6 +1375,7 @@ class __SYCL_EXPORT handler {
13091375
(void)NumWorkGroups;
13101376
kernel_parallel_for_work_group<NameT, KernelType, Dims>(KernelFunc);
13111377
#else
1378+
detail::checkValueRange<Dims>(NumWorkGroups);
13121379
MNDRDesc.setNumWorkGroups(NumWorkGroups);
13131380
MKernel = detail::getSyclObjImpl(std::move(Kernel));
13141381
StoreLambda<NameT, KernelType, Dims>(std::move(KernelFunc));
@@ -1345,7 +1412,12 @@ class __SYCL_EXPORT handler {
13451412
(void)WorkGroupSize;
13461413
kernel_parallel_for_work_group<NameT, KernelType, Dims>(KernelFunc);
13471414
#else
1348-
MNDRDesc.set(nd_range<Dims>(NumWorkGroups * WorkGroupSize, WorkGroupSize));
1415+
nd_range<Dims> ExecRange =
1416+
nd_range<Dims>(NumWorkGroups * WorkGroupSize, WorkGroupSize);
1417+
detail::checkValueRange<Dims>(ExecRange.get_global_range());
1418+
detail::checkValueRange<Dims>(ExecRange.get_local_range());
1419+
detail::checkValueRange<Dims>(ExecRange.get_offset());
1420+
MNDRDesc.set(std::move(ExecRange));
13491421
MKernel = detail::getSyclObjImpl(std::move(Kernel));
13501422
StoreLambda<NameT, KernelType, Dims>(std::move(KernelFunc));
13511423
MCGType = detail::CG::KERNEL;

‎sycl/test/basic_tests/range_offset_fit_in_int.cpp

Lines changed: 177 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,177 @@
1+
// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple -D__SYCL_ID_QUERIES_FIT_IN_INT__=1 %s -o %t.out
2+
// RUN: %CPU_RUN_PLACEHOLDER %t.out
3+
4+
#include <CL/sycl.hpp>
5+
#include <climits>
6+
7+
namespace S = cl::sycl;
8+
9+
void checkRangeException(S::runtime_error &E) {
10+
constexpr char Msg[] = "Provided range is out of integer limits. "
11+
"Pass `-U__SYCL_ID_QUERIES_FIT_IN_INT__' to "
12+
"disable range check.";
13+
14+
std::cerr << E.what() << std::endl;
15+
16+
assert(std::string(E.what()).find(Msg) == 0 && "Unexpected message");
17+
}
18+
19+
void checkOffsetException(S::runtime_error &E) {
20+
constexpr char Msg[] = "Provided offset is out of integer limits. "
21+
"Pass `-U__SYCL_ID_QUERIES_FIT_IN_INT__' to "
22+
"disable offset check.";
23+
24+
std::cerr << E.what() << std::endl;
25+
26+
assert(std::string(E.what()).find(Msg) == 0 && "Unexpected message");
27+
}
28+
29+
void test() {
30+
auto EH = [](S::exception_list EL) {
31+
for (const std::exception_ptr &E : EL) {
32+
throw E;
33+
}
34+
};
35+
36+
S::queue Queue(EH);
37+
38+
static constexpr size_t OutOfLimitsSize = static_cast<size_t>(INT_MAX) + 1;
39+
40+
S::range<1> RangeOutOfLimits{OutOfLimitsSize};
41+
S::range<1> RangeInLimits{1};
42+
S::id<1> OffsetOutOfLimits{OutOfLimitsSize};
43+
S::id<1> OffsetInLimits{1};
44+
S::nd_range<1> NDRange_ROL_LIL_OIL{RangeOutOfLimits, RangeInLimits,
45+
OffsetInLimits};
46+
S::nd_range<1> NDRange_RIL_LOL_OIL{RangeInLimits, RangeOutOfLimits,
47+
OffsetInLimits};
48+
S::nd_range<1> NDRange_RIL_LIL_OOL{RangeInLimits, RangeInLimits,
49+
OffsetOutOfLimits};
50+
S::nd_range<1> NDRange_RIL_LIL_OIL(RangeInLimits, RangeInLimits,
51+
OffsetInLimits);
52+
53+
int Data = 0;
54+
S::buffer<int, 1> Buf{&Data, 1};
55+
56+
try {
57+
Queue.submit([&](S::handler &CGH) {
58+
auto Acc = Buf.get_access<cl::sycl::access::mode::read_write>(CGH);
59+
60+
CGH.parallel_for<class PF_ROL>(RangeOutOfLimits,
61+
[=](S::id<1> Id) { Acc[0] += 1; });
62+
});
63+
64+
assert(false && "Exception expected");
65+
} catch (S::runtime_error &E) {
66+
checkRangeException(E);
67+
} catch (...) {
68+
assert(false && "Unexpected exception catched");
69+
}
70+
71+
try {
72+
Queue.submit([&](S::handler &CGH) {
73+
auto Acc = Buf.get_access<cl::sycl::access::mode::read_write>(CGH);
74+
75+
CGH.parallel_for<class PF_RIL>(RangeInLimits,
76+
[Acc](S::id<1> Id) { Acc[0] += 1; });
77+
});
78+
} catch (...) {
79+
assert(false && "Unexpected exception catched");
80+
}
81+
82+
try {
83+
Queue.submit([&](S::handler &CGH) {
84+
auto Acc = Buf.get_access<cl::sycl::access::mode::read_write>(CGH);
85+
86+
CGH.parallel_for<class PF_ROL_OIL>(RangeOutOfLimits, OffsetInLimits,
87+
[Acc](S::id<1> Id) { Acc[0] += 1; });
88+
});
89+
90+
assert(false && "Exception expected");
91+
} catch (S::runtime_error &E) {
92+
checkRangeException(E);
93+
} catch (...) {
94+
assert(false && "Unexpected exception catched");
95+
}
96+
97+
try {
98+
Queue.submit([&](S::handler &CGH) {
99+
auto Acc = Buf.get_access<cl::sycl::access::mode::read_write>(CGH);
100+
101+
CGH.parallel_for<class PF_RIL_OOL>(RangeInLimits, OffsetOutOfLimits,
102+
[Acc](S::id<1> Id) { Acc[0] += 1; });
103+
});
104+
105+
assert(false && "Exception expected");
106+
} catch (S::runtime_error &E) {
107+
checkOffsetException(E);
108+
} catch (...) {
109+
assert(false && "Unexpected exception catched");
110+
}
111+
112+
try {
113+
Queue.submit([&](S::handler &CGH) {
114+
auto Acc = Buf.get_access<cl::sycl::access::mode::read_write>(CGH);
115+
116+
CGH.parallel_for<class PF_RIL_OIL>(RangeInLimits, OffsetInLimits,
117+
[Acc](S::id<1> Id) { Acc[0] += 1; });
118+
});
119+
} catch (...) {
120+
assert(false && "Unexpected exception catched");
121+
}
122+
123+
try {
124+
Queue.submit([&](S::handler &CGH) {
125+
auto Acc = Buf.get_access<cl::sycl::access::mode::read_write>(CGH);
126+
127+
CGH.parallel_for<class PF_ND_GOL_LIL_OIL>(
128+
NDRange_ROL_LIL_OIL, [Acc](S::nd_item<1> Id) { Acc[0] += 1; });
129+
});
130+
} catch (S::runtime_error &E) {
131+
checkRangeException(E);
132+
} catch (...) {
133+
assert(false && "Unexpected exception catched");
134+
}
135+
136+
try {
137+
Queue.submit([&](S::handler &CGH) {
138+
auto Acc = Buf.get_access<cl::sycl::access::mode::read_write>(CGH);
139+
140+
CGH.parallel_for<class PF_ND_GIL_LOL_OIL>(
141+
NDRange_RIL_LOL_OIL, [Acc](S::nd_item<1> Id) { Acc[0] += 1; });
142+
});
143+
} catch (S::runtime_error &E) {
144+
checkRangeException(E);
145+
} catch (...) {
146+
assert(false && "Unexpected exception catched");
147+
}
148+
149+
try {
150+
Queue.submit([&](S::handler &CGH) {
151+
auto Acc = Buf.get_access<cl::sycl::access::mode::read_write>(CGH);
152+
153+
CGH.parallel_for<class PF_ND_GIL_LIL_OOL>(
154+
NDRange_RIL_LIL_OOL, [Acc](S::nd_item<1> Id) { Acc[0] += 1; });
155+
});
156+
} catch (S::runtime_error &E) {
157+
checkOffsetException(E);
158+
} catch (...) {
159+
assert(false && "Unexpected exception catched");
160+
}
161+
162+
try {
163+
Queue.submit([&](S::handler &CGH) {
164+
auto Acc = Buf.get_access<cl::sycl::access::mode::read_write>(CGH);
165+
166+
CGH.parallel_for<class PF_ND_GIL_LIL_OIL>(
167+
NDRange_RIL_LIL_OIL, [Acc](S::nd_item<1> Id) { Acc[0] += 1; });
168+
});
169+
} catch (...) {
170+
assert(false && "Unexpected exception catched");
171+
}
172+
}
173+
174+
int main(void) {
175+
test();
176+
return 0;
177+
}

0 commit comments

Comments
 (0)