Skip to content

Commit 9ab27f1

Browse files
aelovikov-intelaelovikov-intel
authored
[NFC][SYCL] Inline _KERNELFUNCPARAM macro (#17579)
There is a single unique value for it since #16118, so there is no reason for having it.
1 parent 25e1b55 commit 9ab27f1

File tree

2 files changed

+45
-64
lines changed

2 files changed

+45
-64
lines changed

‎sycl/include/sycl/handler.hpp

+35-42
Original file line numberDiff line numberDiff line change
@@ -66,9 +66,6 @@
6666
// 41(!!!) includes of SYCL headers + 10 includes of standard headers.
6767
// 3300+ lines of code
6868

69-
#define _KERNELFUNCPARAMTYPE const KernelType &
70-
#define _KERNELFUNCPARAM(a) _KERNELFUNCPARAMTYPE a
71-
7269
#if defined(__SYCL_UNNAMED_LAMBDA__)
7370
// We can't use nested types (e.g. struct S defined inside main() routine) to
7471
// name kernels. At the same time, we have to provide a unique kernel name for
@@ -1251,7 +1248,7 @@ class __SYCL_EXPORT handler {
12511248
template <typename KernelName, typename KernelType, int Dims,
12521249
typename PropertiesT>
12531250
void parallel_for_impl(nd_range<Dims> ExecutionRange, PropertiesT Props,
1254-
_KERNELFUNCPARAM(KernelFunc)) {
1251+
const KernelType &KernelFunc) {
12551252
// TODO: Properties may change the kernel function, so in order to avoid
12561253
// conflicts they should be included in the name.
12571254
using NameT =
@@ -1345,7 +1342,7 @@ class __SYCL_EXPORT handler {
13451342
typename PropertiesT = ext::oneapi::experimental::empty_properties_t>
13461343
void parallel_for_work_group_lambda_impl(range<Dims> NumWorkGroups,
13471344
PropertiesT Props,
1348-
_KERNELFUNCPARAM(KernelFunc)) {
1345+
const KernelType &KernelFunc) {
13491346
// TODO: Properties may change the kernel function, so in order to avoid
13501347
// conflicts they should be included in the name.
13511348
using NameT =
@@ -1386,7 +1383,7 @@ class __SYCL_EXPORT handler {
13861383
void parallel_for_work_group_lambda_impl(range<Dims> NumWorkGroups,
13871384
range<Dims> WorkGroupSize,
13881385
PropertiesT Props,
1389-
_KERNELFUNCPARAM(KernelFunc)) {
1386+
const KernelType &KernelFunc) {
13901387
// TODO: Properties may change the kernel function, so in order to avoid
13911388
// conflicts they should be included in the name.
13921389
using NameT =
@@ -1434,7 +1431,7 @@ class __SYCL_EXPORT handler {
14341431
#endif
14351432

14361433
__SYCL_KERNEL_ATTR__ static void
1437-
kernel_single_task(_KERNELFUNCPARAM(KernelFunc)) {
1434+
kernel_single_task(const KernelType &KernelFunc) {
14381435
#ifdef __SYCL_DEVICE_ONLY__
14391436
KernelFunc();
14401437
#else
@@ -1453,7 +1450,7 @@ class __SYCL_EXPORT handler {
14531450
ext::oneapi::experimental::detail::PropertyMetaInfo<Props>::value...)]]
14541451
#endif
14551452
__SYCL_KERNEL_ATTR__ static void
1456-
kernel_single_task(_KERNELFUNCPARAM(KernelFunc), kernel_handler KH) {
1453+
kernel_single_task(const KernelType &KernelFunc, kernel_handler KH) {
14571454
#ifdef __SYCL_DEVICE_ONLY__
14581455
KernelFunc(KH);
14591456
#else
@@ -1472,7 +1469,7 @@ class __SYCL_EXPORT handler {
14721469
ext::oneapi::experimental::detail::PropertyMetaInfo<Props>::value...)]]
14731470
#endif
14741471
__SYCL_KERNEL_ATTR__ static void
1475-
kernel_parallel_for(_KERNELFUNCPARAM(KernelFunc)) {
1472+
kernel_parallel_for(const KernelType &KernelFunc) {
14761473
#ifdef __SYCL_DEVICE_ONLY__
14771474
KernelFunc(detail::Builder::getElement(detail::declptr<ElementType>()));
14781475
#else
@@ -1490,7 +1487,7 @@ class __SYCL_EXPORT handler {
14901487
ext::oneapi::experimental::detail::PropertyMetaInfo<Props>::value...)]]
14911488
#endif
14921489
__SYCL_KERNEL_ATTR__ static void
1493-
kernel_parallel_for(_KERNELFUNCPARAM(KernelFunc), kernel_handler KH) {
1490+
kernel_parallel_for(const KernelType &KernelFunc, kernel_handler KH) {
14941491
#ifdef __SYCL_DEVICE_ONLY__
14951492
KernelFunc(detail::Builder::getElement(detail::declptr<ElementType>()), KH);
14961493
#else
@@ -1509,7 +1506,7 @@ class __SYCL_EXPORT handler {
15091506
ext::oneapi::experimental::detail::PropertyMetaInfo<Props>::value...)]]
15101507
#endif
15111508
__SYCL_KERNEL_ATTR__ static void
1512-
kernel_parallel_for_work_group(_KERNELFUNCPARAM(KernelFunc)) {
1509+
kernel_parallel_for_work_group(const KernelType &KernelFunc) {
15131510
#ifdef __SYCL_DEVICE_ONLY__
15141511
KernelFunc(detail::Builder::getElement(detail::declptr<ElementType>()));
15151512
#else
@@ -1527,7 +1524,7 @@ class __SYCL_EXPORT handler {
15271524
ext::oneapi::experimental::detail::PropertyMetaInfo<Props>::value...)]]
15281525
#endif
15291526
__SYCL_KERNEL_ATTR__ static void
1530-
kernel_parallel_for_work_group(_KERNELFUNCPARAM(KernelFunc),
1527+
kernel_parallel_for_work_group(const KernelType &KernelFunc,
15311528
kernel_handler KH) {
15321529
#ifdef __SYCL_DEVICE_ONLY__
15331530
KernelFunc(detail::Builder::getElement(detail::declptr<ElementType>()), KH);
@@ -1589,7 +1586,7 @@ class __SYCL_EXPORT handler {
15891586
// Couldn't think of a better way to achieve both.
15901587
template <typename KernelName, typename KernelType, typename PropertiesT,
15911588
bool HasKernelHandlerArg, typename FuncTy>
1592-
void unpack(_KERNELFUNCPARAM(KernelFunc), FuncTy Lambda) {
1589+
void unpack(const KernelType &KernelFunc, FuncTy Lambda) {
15931590
#ifdef __SYCL_DEVICE_ONLY__
15941591
detail::CheckDeviceCopyable<KernelType>();
15951592
#endif // __SYCL_DEVICE_ONLY__
@@ -1600,8 +1597,7 @@ class __SYCL_EXPORT handler {
16001597
#ifndef __SYCL_DEVICE_ONLY__
16011598
// If there are properties provided by get method then process them.
16021599
if constexpr (ext::oneapi::experimental::detail::
1603-
HasKernelPropertiesGetMethod<
1604-
_KERNELFUNCPARAMTYPE>::value) {
1600+
HasKernelPropertiesGetMethod<const KernelType &>::value) {
16051601
processProperties<detail::isKernelESIMD<KernelName>()>(
16061602
KernelFunc.get(ext::oneapi::experimental::properties_tag{}));
16071603
}
@@ -1620,7 +1616,7 @@ class __SYCL_EXPORT handler {
16201616
template <
16211617
typename KernelName, typename KernelType,
16221618
typename PropertiesT = ext::oneapi::experimental::empty_properties_t>
1623-
void kernel_single_task_wrapper(_KERNELFUNCPARAM(KernelFunc)) {
1619+
void kernel_single_task_wrapper(const KernelType &KernelFunc) {
16241620
unpack<KernelName, KernelType, PropertiesT,
16251621
detail::KernelLambdaHasKernelHandlerArgT<KernelType>::value>(
16261622
KernelFunc, [&](auto Unpacker, auto... args) {
@@ -1632,7 +1628,7 @@ class __SYCL_EXPORT handler {
16321628
template <
16331629
typename KernelName, typename ElementType, typename KernelType,
16341630
typename PropertiesT = ext::oneapi::experimental::empty_properties_t>
1635-
void kernel_parallel_for_wrapper(_KERNELFUNCPARAM(KernelFunc)) {
1631+
void kernel_parallel_for_wrapper(const KernelType &KernelFunc) {
16361632
unpack<KernelName, KernelType, PropertiesT,
16371633
detail::KernelLambdaHasKernelHandlerArgT<KernelType,
16381634
ElementType>::value>(
@@ -1645,7 +1641,7 @@ class __SYCL_EXPORT handler {
16451641
template <
16461642
typename KernelName, typename ElementType, typename KernelType,
16471643
typename PropertiesT = ext::oneapi::experimental::empty_properties_t>
1648-
void kernel_parallel_for_work_group_wrapper(_KERNELFUNCPARAM(KernelFunc)) {
1644+
void kernel_parallel_for_work_group_wrapper(const KernelType &KernelFunc) {
16491645
unpack<KernelName, KernelType, PropertiesT,
16501646
detail::KernelLambdaHasKernelHandlerArgT<KernelType,
16511647
ElementType>::value>(
@@ -1666,7 +1662,7 @@ class __SYCL_EXPORT handler {
16661662
typename KernelName, typename KernelType,
16671663
typename PropertiesT = ext::oneapi::experimental::empty_properties_t>
16681664
void single_task_lambda_impl(PropertiesT Props,
1669-
_KERNELFUNCPARAM(KernelFunc)) {
1665+
const KernelType &KernelFunc) {
16701666
(void)Props;
16711667
// TODO: Properties may change the kernel function, so in order to avoid
16721668
// conflicts they should be included in the name.
@@ -1892,27 +1888,27 @@ class __SYCL_EXPORT handler {
18921888
///
18931889
/// \param KernelFunc is a SYCL kernel function.
18941890
template <typename KernelName = detail::auto_name, typename KernelType>
1895-
void single_task(_KERNELFUNCPARAM(KernelFunc)) {
1891+
void single_task(const KernelType &KernelFunc) {
18961892
single_task_lambda_impl<KernelName>(
18971893
ext::oneapi::experimental::empty_properties_t{}, KernelFunc);
18981894
}
18991895

19001896
template <typename KernelName = detail::auto_name, typename KernelType>
1901-
void parallel_for(range<1> NumWorkItems, _KERNELFUNCPARAM(KernelFunc)) {
1897+
void parallel_for(range<1> NumWorkItems, const KernelType &KernelFunc) {
19021898
parallel_for_lambda_impl<KernelName>(
19031899
NumWorkItems, ext::oneapi::experimental::empty_properties_t{},
19041900
std::move(KernelFunc));
19051901
}
19061902

19071903
template <typename KernelName = detail::auto_name, typename KernelType>
1908-
void parallel_for(range<2> NumWorkItems, _KERNELFUNCPARAM(KernelFunc)) {
1904+
void parallel_for(range<2> NumWorkItems, const KernelType &KernelFunc) {
19091905
parallel_for_lambda_impl<KernelName>(
19101906
NumWorkItems, ext::oneapi::experimental::empty_properties_t{},
19111907
std::move(KernelFunc));
19121908
}
19131909

19141910
template <typename KernelName = detail::auto_name, typename KernelType>
1915-
void parallel_for(range<3> NumWorkItems, _KERNELFUNCPARAM(KernelFunc)) {
1911+
void parallel_for(range<3> NumWorkItems, const KernelType &KernelFunc) {
19161912
parallel_for_lambda_impl<KernelName>(
19171913
NumWorkItems, ext::oneapi::experimental::empty_properties_t{},
19181914
std::move(KernelFunc));
@@ -1955,7 +1951,7 @@ class __SYCL_EXPORT handler {
19551951
int Dims>
19561952
__SYCL2020_DEPRECATED("offsets are deprecated in SYCL2020")
19571953
void parallel_for(range<Dims> NumWorkItems, id<Dims> WorkItemOffset,
1958-
_KERNELFUNCPARAM(KernelFunc)) {
1954+
const KernelType &KernelFunc) {
19591955
using NameT =
19601956
typename detail::get_kernel_name_t<KernelName, KernelType>::name;
19611957
using LambdaArgType = sycl::detail::lambda_arg_type<KernelType, item<Dims>>;
@@ -1990,7 +1986,7 @@ class __SYCL_EXPORT handler {
19901986
template <typename KernelName = detail::auto_name, typename KernelType,
19911987
int Dims>
19921988
void parallel_for_work_group(range<Dims> NumWorkGroups,
1993-
_KERNELFUNCPARAM(KernelFunc)) {
1989+
const KernelType &KernelFunc) {
19941990
parallel_for_work_group_lambda_impl<KernelName>(
19951991
NumWorkGroups, ext::oneapi::experimental::empty_properties_t{},
19961992
KernelFunc);
@@ -2012,7 +2008,7 @@ class __SYCL_EXPORT handler {
20122008
int Dims>
20132009
void parallel_for_work_group(range<Dims> NumWorkGroups,
20142010
range<Dims> WorkGroupSize,
2015-
_KERNELFUNCPARAM(KernelFunc)) {
2011+
const KernelType &KernelFunc) {
20162012
parallel_for_work_group_lambda_impl<KernelName>(
20172013
NumWorkGroups, WorkGroupSize,
20182014
ext::oneapi::experimental::empty_properties_t{}, KernelFunc);
@@ -2096,7 +2092,7 @@ class __SYCL_EXPORT handler {
20962092
/// \param KernelFunc is a lambda that is used if device, queue is bound to,
20972093
/// is a host device.
20982094
template <typename KernelName = detail::auto_name, typename KernelType>
2099-
void single_task(kernel Kernel, _KERNELFUNCPARAM(KernelFunc)) {
2095+
void single_task(kernel Kernel, const KernelType &KernelFunc) {
21002096
// Ignore any set kernel bundles and use the one associated with the kernel
21012097
setHandlerKernelBundle(Kernel);
21022098
using NameT =
@@ -2132,7 +2128,7 @@ class __SYCL_EXPORT handler {
21322128
template <typename KernelName = detail::auto_name, typename KernelType,
21332129
int Dims>
21342130
void parallel_for(kernel Kernel, range<Dims> NumWorkItems,
2135-
_KERNELFUNCPARAM(KernelFunc)) {
2131+
const KernelType &KernelFunc) {
21362132
// Ignore any set kernel bundles and use the one associated with the kernel
21372133
setHandlerKernelBundle(Kernel);
21382134
using NameT =
@@ -2171,7 +2167,7 @@ class __SYCL_EXPORT handler {
21712167
int Dims>
21722168
__SYCL2020_DEPRECATED("offsets are deprecated in SYCL 2020")
21732169
void parallel_for(kernel Kernel, range<Dims> NumWorkItems,
2174-
id<Dims> WorkItemOffset, _KERNELFUNCPARAM(KernelFunc)) {
2170+
id<Dims> WorkItemOffset, const KernelType &KernelFunc) {
21752171
using NameT =
21762172
typename detail::get_kernel_name_t<KernelName, KernelType>::name;
21772173
using LambdaArgType = sycl::detail::lambda_arg_type<KernelType, item<Dims>>;
@@ -2210,7 +2206,7 @@ class __SYCL_EXPORT handler {
22102206
template <typename KernelName = detail::auto_name, typename KernelType,
22112207
int Dims>
22122208
void parallel_for(kernel Kernel, nd_range<Dims> NDRange,
2213-
_KERNELFUNCPARAM(KernelFunc)) {
2209+
const KernelType &KernelFunc) {
22142210
using NameT =
22152211
typename detail::get_kernel_name_t<KernelName, KernelType>::name;
22162212
using LambdaArgType =
@@ -2253,7 +2249,7 @@ class __SYCL_EXPORT handler {
22532249
template <typename KernelName = detail::auto_name, typename KernelType,
22542250
int Dims>
22552251
void parallel_for_work_group(kernel Kernel, range<Dims> NumWorkGroups,
2256-
_KERNELFUNCPARAM(KernelFunc)) {
2252+
const KernelType &KernelFunc) {
22572253
using NameT =
22582254
typename detail::get_kernel_name_t<KernelName, KernelType>::name;
22592255
using LambdaArgType =
@@ -2294,7 +2290,7 @@ class __SYCL_EXPORT handler {
22942290
int Dims>
22952291
void parallel_for_work_group(kernel Kernel, range<Dims> NumWorkGroups,
22962292
range<Dims> WorkGroupSize,
2297-
_KERNELFUNCPARAM(KernelFunc)) {
2293+
const KernelType &KernelFunc) {
22982294
using NameT =
22992295
typename detail::get_kernel_name_t<KernelName, KernelType>::name;
23002296
using LambdaArgType =
@@ -2327,7 +2323,7 @@ class __SYCL_EXPORT handler {
23272323
"member function instead.")
23282324
std::enable_if_t<ext::oneapi::experimental::is_property_list<
23292325
PropertiesT>::value> single_task(PropertiesT Props,
2330-
_KERNELFUNCPARAM(KernelFunc)) {
2326+
const KernelType &KernelFunc) {
23312327
single_task_lambda_impl<KernelName, KernelType, PropertiesT>(Props,
23322328
KernelFunc);
23332329
}
@@ -2341,7 +2337,7 @@ class __SYCL_EXPORT handler {
23412337
std::enable_if_t<ext::oneapi::experimental::is_property_list<
23422338
PropertiesT>::value> parallel_for(range<1> NumWorkItems,
23432339
PropertiesT Props,
2344-
_KERNELFUNCPARAM(KernelFunc)) {
2340+
const KernelType &KernelFunc) {
23452341
parallel_for_lambda_impl<KernelName, KernelType, 1, PropertiesT>(
23462342
NumWorkItems, Props, std::move(KernelFunc));
23472343
}
@@ -2355,7 +2351,7 @@ class __SYCL_EXPORT handler {
23552351
std::enable_if_t<ext::oneapi::experimental::is_property_list<
23562352
PropertiesT>::value> parallel_for(range<2> NumWorkItems,
23572353
PropertiesT Props,
2358-
_KERNELFUNCPARAM(KernelFunc)) {
2354+
const KernelType &KernelFunc) {
23592355
parallel_for_lambda_impl<KernelName, KernelType, 2, PropertiesT>(
23602356
NumWorkItems, Props, std::move(KernelFunc));
23612357
}
@@ -2369,7 +2365,7 @@ class __SYCL_EXPORT handler {
23692365
std::enable_if_t<ext::oneapi::experimental::is_property_list<
23702366
PropertiesT>::value> parallel_for(range<3> NumWorkItems,
23712367
PropertiesT Props,
2372-
_KERNELFUNCPARAM(KernelFunc)) {
2368+
const KernelType &KernelFunc) {
23732369
parallel_for_lambda_impl<KernelName, KernelType, 3, PropertiesT>(
23742370
NumWorkItems, Props, std::move(KernelFunc));
23752371
}
@@ -2383,7 +2379,7 @@ class __SYCL_EXPORT handler {
23832379
std::enable_if_t<ext::oneapi::experimental::is_property_list<
23842380
PropertiesT>::value> parallel_for(nd_range<Dims> Range,
23852381
PropertiesT Properties,
2386-
_KERNELFUNCPARAM(KernelFunc)) {
2382+
const KernelType &KernelFunc) {
23872383
parallel_for_impl<KernelName>(Range, Properties, std::move(KernelFunc));
23882384
}
23892385

@@ -2511,7 +2507,7 @@ class __SYCL_EXPORT handler {
25112507
"get(sycl::ext::oneapi::experimental::properties_tag) "
25122508
"member function instead.")
25132509
void parallel_for_work_group(range<Dims> NumWorkGroups, PropertiesT Props,
2514-
_KERNELFUNCPARAM(KernelFunc)) {
2510+
const KernelType &KernelFunc) {
25152511
parallel_for_work_group_lambda_impl<KernelName, KernelType, Dims,
25162512
PropertiesT>(NumWorkGroups, Props,
25172513
KernelFunc);
@@ -2525,15 +2521,12 @@ class __SYCL_EXPORT handler {
25252521
"member function instead.")
25262522
void parallel_for_work_group(range<Dims> NumWorkGroups,
25272523
range<Dims> WorkGroupSize, PropertiesT Props,
2528-
_KERNELFUNCPARAM(KernelFunc)) {
2524+
const KernelType &KernelFunc) {
25292525
parallel_for_work_group_lambda_impl<KernelName, KernelType, Dims,
25302526
PropertiesT>(
25312527
NumWorkGroups, WorkGroupSize, Props, KernelFunc);
25322528
}
25332529

2534-
// Clean up KERNELFUNC macro.
2535-
#undef _KERNELFUNCPARAM
2536-
25372530
// Explicit copy operations API
25382531

25392532
/// Copies the content of memory object accessed by Src into the memory

0 commit comments

Comments
 (0)