Skip to content

Commit ac364f2

Browse files
AlexeySachkovAlexeySachkov
authored
[SYCL][NFC] Optmize handler.hpp compilation for device [1/N] (#15674)
Outlined runtime checks related to `-fsycl-id-queries-fit-in-int` into a separate header file which is only used in host compilation.
1 parent 65849fd commit ac364f2

File tree

3 files changed

+127
-86
lines changed

3 files changed

+127
-86
lines changed

‎sycl/include/sycl/detail/id_queries_fit_in_int.hpp‎

Lines changed: 110 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,110 @@
1+
//==-------------------- id_queries_fit_in_int.hpp -------------------------==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
//
9+
// Our SYCL implementation has a special mode (introduced for performance
10+
// reasons) in which it assume that all result of all id queries (i.e. global
11+
// sizes, work-group sizes, local id, global id, etc.) fit within MAX_INT.
12+
//
13+
// This header contains corresponding helper functions related to this mode.
14+
//
15+
//===----------------------------------------------------------------------===//
16+
17+
#pragma once
18+
19+
// We only use those helpers to throw an exception if user selected a range that
20+
// would violate the assumption. That can only happen on host and therefore to
21+
// optimize our headers, the helpers below are only available for host
22+
// compilation.
23+
#ifndef __SYCL_DEVICE_ONLY__
24+
25+
#include <sycl/exception.hpp>
26+
27+
#include <limits>
28+
#include <type_traits>
29+
30+
namespace sycl {
31+
inline namespace _V1 {
32+
namespace detail {
33+
34+
#if __SYCL_ID_QUERIES_FIT_IN_INT__
35+
template <typename T> struct NotIntMsg;
36+
37+
template <int Dims> struct NotIntMsg<range<Dims>> {
38+
constexpr static const char *Msg =
39+
"Provided range is out of integer limits. Pass "
40+
"`-fno-sycl-id-queries-fit-in-int' to disable range check.";
41+
};
42+
43+
template <int Dims> struct NotIntMsg<id<Dims>> {
44+
constexpr static const char *Msg =
45+
"Provided offset is out of integer limits. Pass "
46+
"`-fno-sycl-id-queries-fit-in-int' to disable offset check.";
47+
};
48+
49+
template <typename T, typename ValT>
50+
typename std::enable_if_t<std::is_same<ValT, size_t>::value ||
51+
std::is_same<ValT, unsigned long long>::value>
52+
checkValueRangeImpl(ValT V) {
53+
static constexpr size_t Limit =
54+
static_cast<size_t>((std::numeric_limits<int>::max)());
55+
if (V > Limit)
56+
throw sycl::exception(make_error_code(errc::nd_range), NotIntMsg<T>::Msg);
57+
}
58+
#endif
59+
60+
template <int Dims, typename T>
61+
typename std::enable_if_t<std::is_same_v<T, range<Dims>> ||
62+
std::is_same_v<T, id<Dims>>>
63+
checkValueRange([[maybe_unused]] const T &V) {
64+
#if __SYCL_ID_QUERIES_FIT_IN_INT__
65+
for (size_t Dim = 0; Dim < Dims; ++Dim)
66+
checkValueRangeImpl<T>(V[Dim]);
67+
68+
{
69+
unsigned long long Product = 1;
70+
for (size_t Dim = 0; Dim < Dims; ++Dim) {
71+
Product *= V[Dim];
72+
// check value now to prevent product overflow in the end
73+
checkValueRangeImpl<T>(Product);
74+
}
75+
}
76+
#endif
77+
}
78+
79+
template <int Dims>
80+
void checkValueRange([[maybe_unused]] const range<Dims> &R,
81+
[[maybe_unused]] const id<Dims> &O) {
82+
#if __SYCL_ID_QUERIES_FIT_IN_INT__
83+
checkValueRange<Dims>(R);
84+
checkValueRange<Dims>(O);
85+
86+
for (size_t Dim = 0; Dim < Dims; ++Dim) {
87+
unsigned long long Sum = R[Dim] + O[Dim];
88+
89+
checkValueRangeImpl<range<Dims>>(Sum);
90+
}
91+
#endif
92+
}
93+
94+
template <int Dims, typename T>
95+
typename std::enable_if_t<std::is_same_v<T, nd_range<Dims>>>
96+
checkValueRange([[maybe_unused]] const T &V) {
97+
#if __SYCL_ID_QUERIES_FIT_IN_INT__
98+
checkValueRange<Dims>(V.get_global_range());
99+
checkValueRange<Dims>(V.get_local_range());
100+
checkValueRange<Dims>(V.get_offset());
101+
102+
checkValueRange<Dims>(V.get_global_range(), V.get_offset());
103+
#endif
104+
}
105+
106+
} // namespace detail
107+
} // namespace _V1
108+
} // namespace sycl
109+
110+
#endif

‎sycl/include/sycl/handler.hpp‎

Lines changed: 16 additions & 86 deletions
Original file line numberDiff line numberDiff line change
@@ -16,6 +16,7 @@
1616
#include <sycl/detail/common.hpp>
1717
#include <sycl/detail/defines_elementary.hpp>
1818
#include <sycl/detail/export.hpp>
19+
#include <sycl/detail/id_queries_fit_in_int.hpp>
1920
#include <sycl/detail/impl_utils.hpp>
2021
#include <sycl/detail/kernel_desc.hpp>
2122
#include <sycl/detail/reduction_forward.hpp>
@@ -227,22 +228,6 @@ __SYCL_EXPORT void *getValueFromDynamicParameter(
227228
ext::oneapi::experimental::detail::dynamic_parameter_base
228229
&DynamicParamBase);
229230

230-
#if __SYCL_ID_QUERIES_FIT_IN_INT__
231-
template <typename T> struct NotIntMsg;
232-
233-
template <int Dims> struct NotIntMsg<range<Dims>> {
234-
constexpr static const char *Msg =
235-
"Provided range is out of integer limits. Pass "
236-
"`-fno-sycl-id-queries-fit-in-int' to disable range check.";
237-
};
238-
239-
template <int Dims> struct NotIntMsg<id<Dims>> {
240-
constexpr static const char *Msg =
241-
"Provided offset is out of integer limits. Pass "
242-
"`-fno-sycl-id-queries-fit-in-int' to disable offset check.";
243-
};
244-
#endif
245-
246231
// Helper for merging properties with ones defined in an optional kernel functor
247232
// getter.
248233
template <typename KernelType, typename PropertiesT, typename Cond = void>
@@ -265,70 +250,6 @@ struct GetMergedKernelProperties<
265250
PropertiesT, get_method_properties>;
266251
};
267252

268-
#if __SYCL_ID_QUERIES_FIT_IN_INT__
269-
template <typename T, typename ValT>
270-
typename std::enable_if_t<std::is_same<ValT, size_t>::value ||
271-
std::is_same<ValT, unsigned long long>::value>
272-
checkValueRangeImpl(ValT V) {
273-
static constexpr size_t Limit =
274-
static_cast<size_t>((std::numeric_limits<int>::max)());
275-
if (V > Limit)
276-
throw sycl::exception(make_error_code(errc::nd_range), NotIntMsg<T>::Msg);
277-
}
278-
#endif
279-
280-
template <int Dims, typename T>
281-
typename std::enable_if_t<std::is_same_v<T, range<Dims>> ||
282-
std::is_same_v<T, id<Dims>>>
283-
checkValueRange(const T &V) {
284-
#if __SYCL_ID_QUERIES_FIT_IN_INT__
285-
for (size_t Dim = 0; Dim < Dims; ++Dim)
286-
checkValueRangeImpl<T>(V[Dim]);
287-
288-
{
289-
unsigned long long Product = 1;
290-
for (size_t Dim = 0; Dim < Dims; ++Dim) {
291-
Product *= V[Dim];
292-
// check value now to prevent product overflow in the end
293-
checkValueRangeImpl<T>(Product);
294-
}
295-
}
296-
#else
297-
(void)V;
298-
#endif
299-
}
300-
301-
template <int Dims>
302-
void checkValueRange(const range<Dims> &R, const id<Dims> &O) {
303-
#if __SYCL_ID_QUERIES_FIT_IN_INT__
304-
checkValueRange<Dims>(R);
305-
checkValueRange<Dims>(O);
306-
307-
for (size_t Dim = 0; Dim < Dims; ++Dim) {
308-
unsigned long long Sum = R[Dim] + O[Dim];
309-
310-
checkValueRangeImpl<range<Dims>>(Sum);
311-
}
312-
#else
313-
(void)R;
314-
(void)O;
315-
#endif
316-
}
317-
318-
template <int Dims, typename T>
319-
typename std::enable_if_t<std::is_same_v<T, nd_range<Dims>>>
320-
checkValueRange(const T &V) {
321-
#if __SYCL_ID_QUERIES_FIT_IN_INT__
322-
checkValueRange<Dims>(V.get_global_range());
323-
checkValueRange<Dims>(V.get_local_range());
324-
checkValueRange<Dims>(V.get_offset());
325-
326-
checkValueRange<Dims>(V.get_global_range(), V.get_offset());
327-
#else
328-
(void)V;
329-
#endif
330-
}
331-
332253
template <int Dims> class RoundedRangeIDGenerator {
333254
id<Dims> Id;
334255
id<Dims> InitId;
@@ -1353,8 +1274,10 @@ class __SYCL_EXPORT handler {
13531274
/// \param Kernel is a SYCL kernel function.
13541275
/// \param Properties is the properties.
13551276
template <int Dims, typename PropertiesT>
1356-
void parallel_for_impl(range<Dims> NumWorkItems, PropertiesT Props,
1357-
kernel Kernel) {
1277+
void parallel_for_impl([[maybe_unused]] range<Dims> NumWorkItems,
1278+
[[maybe_unused]] PropertiesT Props,
1279+
[[maybe_unused]] kernel Kernel) {
1280+
#ifndef __SYCL_DEVICE_ONLY__
13581281
throwIfActionIsCreated();
13591282
MKernel = detail::getSyclObjImpl(std::move(Kernel));
13601283
detail::checkValueRange<Dims>(NumWorkItems);
@@ -1364,6 +1287,7 @@ class __SYCL_EXPORT handler {
13641287
setNDRangeUsed(false);
13651288
extractArgsAndReqs();
13661289
MKernelName = getKernelName();
1290+
#endif
13671291
}
13681292

13691293
/// Defines and invokes a SYCL kernel function for the specified range and
@@ -1376,8 +1300,10 @@ class __SYCL_EXPORT handler {
13761300
/// \param Properties is the properties.
13771301
/// \param Kernel is a SYCL kernel function.
13781302
template <int Dims, typename PropertiesT>
1379-
void parallel_for_impl(nd_range<Dims> NDRange, PropertiesT Props,
1380-
kernel Kernel) {
1303+
void parallel_for_impl([[maybe_unused]] nd_range<Dims> NDRange,
1304+
[[maybe_unused]] PropertiesT Props,
1305+
[[maybe_unused]] kernel Kernel) {
1306+
#ifndef __SYCL_DEVICE_ONLY__
13811307
throwIfActionIsCreated();
13821308
MKernel = detail::getSyclObjImpl(std::move(Kernel));
13831309
detail::checkValueRange<Dims>(NDRange);
@@ -1387,6 +1313,7 @@ class __SYCL_EXPORT handler {
13871313
setNDRangeUsed(true);
13881314
extractArgsAndReqs();
13891315
MKernelName = getKernelName();
1316+
#endif
13901317
}
13911318

13921319
/// Hierarchical kernel invocation method of a kernel defined as a lambda
@@ -2136,8 +2063,10 @@ class __SYCL_EXPORT handler {
21362063
/// \param Kernel is a SYCL kernel function.
21372064
template <int Dims>
21382065
__SYCL2020_DEPRECATED("offsets are deprecated in SYCL 2020")
2139-
void parallel_for(range<Dims> NumWorkItems, id<Dims> WorkItemOffset,
2140-
kernel Kernel) {
2066+
void parallel_for([[maybe_unused]] range<Dims> NumWorkItems,
2067+
[[maybe_unused]] id<Dims> WorkItemOffset,
2068+
[[maybe_unused]] kernel Kernel) {
2069+
#ifndef __SYCL_DEVICE_ONLY__
21412070
throwIfActionIsCreated();
21422071
MKernel = detail::getSyclObjImpl(std::move(Kernel));
21432072
detail::checkValueRange<Dims>(NumWorkItems, WorkItemOffset);
@@ -2146,6 +2075,7 @@ class __SYCL_EXPORT handler {
21462075
setNDRangeUsed(false);
21472076
extractArgsAndReqs();
21482077
MKernelName = getKernelName();
2078+
#endif
21492079
}
21502080

21512081
/// Defines and invokes a SYCL kernel function for the specified range and

‎sycl/test/include_deps/sycl_detail_core.hpp.cpp‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -145,6 +145,7 @@
145145
// CHECK-NEXT: CL/cl_version.h
146146
// CHECK-NEXT: CL/cl_platform.h
147147
// CHECK-NEXT: CL/cl_ext.h
148+
// CHECK-NEXT: detail/id_queries_fit_in_int.hpp
148149
// CHECK-NEXT: detail/reduction_forward.hpp
149150
// CHECK-NEXT: detail/ur.hpp
150151
// CHECK-NEXT: ur_api_funcs.def

0 commit comments

Comments
 (0)