Skip to content

Commit 43a54f6

Browse files
committed
[SYCL][NFC] Reduce branching in parallel_for_lambda_impl
- SYCL 2020 is the oldest standard we support, so check `SYCL >= 2020` is always true, dropped the corresponding check; - Outlined calls repeated for both rounded and non-rounded paths into a common section; - Outlined `(void)` casts into a common section to reduce size of `#ifdef` branches for improved readability;
1 parent fb8d4ab commit 43a54f6

File tree

1 file changed

+13
-20
lines changed

1 file changed

+13
-20
lines changed

sycl/include/sycl/handler.hpp

Lines changed: 13 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -1238,6 +1238,9 @@ class __SYCL_EXPORT handler {
12381238
typename PropertiesT = ext::oneapi::experimental::empty_properties_t>
12391239
void parallel_for_lambda_impl(range<Dims> UserRange, PropertiesT Props,
12401240
const KernelType &KernelFunc) {
1241+
(void)Props;
1242+
(void)UserRange;
1243+
(void)KernelFunc;
12411244
// TODO: Properties may change the kernel function, so in order to avoid
12421245
// conflicts they should be included in the name.
12431246
using NameT =
@@ -1251,6 +1254,13 @@ class __SYCL_EXPORT handler {
12511254
throw sycl::exception(make_error_code(errc::runtime),
12521255
"The total number of work-items in "
12531256
"a range must fit within size_t");
1257+
verifyUsedKernelBundleInternal(Info.Name);
1258+
processProperties<Info.IsESIMD, PropertiesT>(Props);
1259+
// Even if range rounding will be used, there will still be items/ids
1260+
// constructed in the range rounded kernel that use items/ids in the
1261+
// user range, which means that __SYCL_ASSUME_INT can still be violated.
1262+
// So check the bounds of the user range regardless.
1263+
detail::checkValueRange<Dims>(UserRange);
12541264
#endif
12551265

12561266
using LambdaArgType = sycl::detail::lambda_arg_type<KernelType, item<Dims>>;
@@ -1280,9 +1290,7 @@ class __SYCL_EXPORT handler {
12801290
"invocable with sycl::item and optionally sycl::kernel_handler");
12811291

12821292
// Range rounding can be disabled by the user.
1283-
// Range rounding is supported only for newer SYCL standards.
1284-
#if !defined(__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__) && \
1285-
SYCL_LANGUAGE_VERSION >= 202012L
1293+
#ifndef __SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__
12861294
auto [RoundedRange, HasRoundedRange] = getRoundedRange(UserRange);
12871295
if (HasRoundedRange) {
12881296
using NameWT = typename detail::get_kernel_wrapper_name_t<NameT>::name;
@@ -1300,23 +1308,13 @@ class __SYCL_EXPORT handler {
13001308
detail::KernelLaunchPropertyWrapper::parseProperties<KName>(this,
13011309
Wrapper);
13021310
#ifndef __SYCL_DEVICE_ONLY__
1303-
verifyUsedKernelBundleInternal(Info.Name);
1304-
// We are executing over the rounded range, but there are still
1305-
// items/ids that are are constructed in ther range rounded
1306-
// kernel use items/ids in the user range, which means that
1307-
// __SYCL_ASSUME_INT can still be violated. So check the bounds
1308-
// of the user range, instead of the rounded range.
1309-
detail::checkValueRange<Dims>(UserRange);
13101311
setNDRangeDescriptor(RoundedRange);
13111312
StoreLambda<KName, decltype(Wrapper), Dims, TransformedArgType>(
13121313
std::move(Wrapper));
13131314
#endif
13141315
} else
1315-
#endif // !__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__ &&
1316-
// SYCL_LANGUAGE_VERSION >= 202012L
1316+
#endif // !__SYCL_DISABLE_PARALLEL_FOR_RANGE_ROUNDING__
13171317
{
1318-
(void)UserRange;
1319-
(void)Props;
13201318
#ifndef __SYCL_FORCE_PARALLEL_FOR_RANGE_ROUNDING__
13211319
// If parallel_for range rounding is forced then only range rounded
13221320
// kernel is generated
@@ -1325,16 +1323,11 @@ class __SYCL_EXPORT handler {
13251323
detail::KernelLaunchPropertyWrapper::parseProperties<NameT>(this,
13261324
KernelFunc);
13271325
#ifndef __SYCL_DEVICE_ONLY__
1328-
verifyUsedKernelBundleInternal(Info.Name);
1329-
processProperties<Info.IsESIMD, PropertiesT>(Props);
1330-
detail::checkValueRange<Dims>(UserRange);
13311326
setNDRangeDescriptor(std::move(UserRange));
13321327
StoreLambda<NameT, KernelType, Dims, TransformedArgType>(
13331328
std::move(KernelFunc));
13341329
#endif
1335-
#else
1336-
(void)KernelFunc;
1337-
#endif // __SYCL_FORCE_PARALLEL_FOR_RANGE_ROUNDING__
1330+
#endif // !__SYCL_FORCE_PARALLEL_FOR_RANGE_ROUNDING__
13381331
}
13391332
}
13401333

0 commit comments

Comments
 (0)