Skip to content

Commit 6ea1490

Browse files
uditagarwal97uditagarwal97
authored
[SYCL][NFC] Move KernelWrapper outside the sycl::handler (#19511)
`KernelWrapper` is also used in the no-handler submission path, so moving this struct outside the handler class into the detail:: namespace. This PR also moves functions like `kernel_single_task` that are called by `KernelWrapper` outside the `sycl::handler` class.
1 parent 9d225d0 commit 6ea1490

File tree

7 files changed

+329
-290
lines changed

7 files changed

+329
-290
lines changed

‎sycl/doc/design/CompileTimeProperties.md

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -362,11 +362,11 @@ string if it is not already a string.
362362
[9]: <https://llvm.org/doxygen/classllvm_1_1Function.html#ae7b919df259dce5480774e656791c079>
363363

364364
**NOTE**: The intention is to replace the existing member functions like
365-
`handler::kernel_single_task()` with wrapper classes like
365+
`detail::KernelWrapperHelperFuncs::kernel_single_task()` with wrapper classes like
366366
`KernelSingleTaskWrapper`. We believe this will not cause problems for the
367367
device compiler front-end because it recognizes kernel functions via the
368368
`[[clang::sycl_kernel]]` attribute, not by the name
369-
`handler::kernel_single_task()`.
369+
`detail::KernelWrapperHelperFuncs::kernel_single_task()`.
370370

371371

372372
## Properties on a non-global variable type

‎sycl/include/sycl/detail/kernel_launch_helper.hpp

Lines changed: 268 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,268 @@
1+
//==-------- kernel_launch_helper.hpp --- SYCL kernel launch utilities ----==//
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+
#pragma once
10+
11+
#include <sycl/detail/cg_types.hpp>
12+
#include <sycl/detail/helpers.hpp>
13+
#include <sycl/ext/intel/experimental/fp_control_kernel_properties.hpp>
14+
#include <sycl/ext/intel/experimental/kernel_execution_properties.hpp>
15+
#include <sycl/ext/oneapi/experimental/virtual_functions.hpp>
16+
#include <sycl/ext/oneapi/kernel_properties/properties.hpp>
17+
#include <sycl/ext/oneapi/work_group_scratch_memory.hpp>
18+
#include <sycl/kernel_handler.hpp>
19+
20+
#include <assert.h>
21+
#include <type_traits>
22+
23+
namespace sycl {
24+
inline namespace _V1 {
25+
namespace detail {
26+
27+
enum class WrapAs { single_task, parallel_for, parallel_for_work_group };
28+
29+
// Helper for merging properties with ones defined in an optional kernel functor
30+
// getter.
31+
template <typename KernelType, typename PropertiesT, typename Cond = void>
32+
struct GetMergedKernelProperties {
33+
using type = PropertiesT;
34+
};
35+
template <typename KernelType, typename PropertiesT>
36+
struct GetMergedKernelProperties<
37+
KernelType, PropertiesT,
38+
std::enable_if_t<ext::oneapi::experimental::detail::
39+
HasKernelPropertiesGetMethod<KernelType>::value>> {
40+
using get_method_properties =
41+
typename ext::oneapi::experimental::detail::HasKernelPropertiesGetMethod<
42+
KernelType>::properties_t;
43+
static_assert(
44+
ext::oneapi::experimental::is_property_list<get_method_properties>::value,
45+
"get(sycl::ext::oneapi::experimental::properties_tag) member in kernel "
46+
"functor class must return a valid property list.");
47+
using type = ext::oneapi::experimental::detail::merged_properties_t<
48+
PropertiesT, get_method_properties>;
49+
};
50+
51+
struct KernelWrapperHelperFuncs {
52+
53+
#ifdef SYCL_LANGUAGE_VERSION
54+
#ifndef __INTEL_SYCL_USE_INTEGRATION_HEADERS
55+
#define __SYCL_KERNEL_ATTR__ [[clang::sycl_kernel_entry_point(KernelName)]]
56+
#else
57+
#define __SYCL_KERNEL_ATTR__ [[clang::sycl_kernel]]
58+
#endif // __INTEL_SYCL_USE_INTEGRATION_HEADERS
59+
#else
60+
#define __SYCL_KERNEL_ATTR__
61+
#endif // SYCL_LANGUAGE_VERSION
62+
63+
// NOTE: the name of this function - "kernel_single_task" - is used by the
64+
// Front End to determine kernel invocation kind.
65+
template <typename KernelName, typename KernelType, typename... Props>
66+
#ifdef __SYCL_DEVICE_ONLY__
67+
[[__sycl_detail__::add_ir_attributes_function(
68+
"sycl-single-task",
69+
ext::oneapi::experimental::detail::PropertyMetaInfo<Props>::name...,
70+
nullptr,
71+
ext::oneapi::experimental::detail::PropertyMetaInfo<Props>::value...)]]
72+
#endif
73+
74+
__SYCL_KERNEL_ATTR__ static void
75+
kernel_single_task(const KernelType &KernelFunc) {
76+
#ifdef __SYCL_DEVICE_ONLY__
77+
KernelFunc();
78+
#else
79+
(void)KernelFunc;
80+
#endif
81+
}
82+
83+
// NOTE: the name of this function - "kernel_single_task" - is used by the
84+
// Front End to determine kernel invocation kind.
85+
template <typename KernelName, typename KernelType, typename... Props>
86+
#ifdef __SYCL_DEVICE_ONLY__
87+
[[__sycl_detail__::add_ir_attributes_function(
88+
"sycl-single-task",
89+
ext::oneapi::experimental::detail::PropertyMetaInfo<Props>::name...,
90+
nullptr,
91+
ext::oneapi::experimental::detail::PropertyMetaInfo<Props>::value...)]]
92+
#endif
93+
__SYCL_KERNEL_ATTR__ static void
94+
kernel_single_task(const KernelType &KernelFunc, kernel_handler KH) {
95+
#ifdef __SYCL_DEVICE_ONLY__
96+
KernelFunc(KH);
97+
#else
98+
(void)KernelFunc;
99+
(void)KH;
100+
#endif
101+
}
102+
103+
// NOTE: the name of these functions - "kernel_parallel_for" - are used by the
104+
// Front End to determine kernel invocation kind.
105+
template <typename KernelName, typename ElementType, typename KernelType,
106+
typename... Props>
107+
#ifdef __SYCL_DEVICE_ONLY__
108+
[[__sycl_detail__::add_ir_attributes_function(
109+
ext::oneapi::experimental::detail::PropertyMetaInfo<Props>::name...,
110+
ext::oneapi::experimental::detail::PropertyMetaInfo<Props>::value...)]]
111+
#endif
112+
__SYCL_KERNEL_ATTR__ static void
113+
kernel_parallel_for(const KernelType &KernelFunc) {
114+
#ifdef __SYCL_DEVICE_ONLY__
115+
KernelFunc(detail::Builder::getElement(detail::declptr<ElementType>()));
116+
#else
117+
(void)KernelFunc;
118+
#endif
119+
}
120+
121+
// NOTE: the name of these functions - "kernel_parallel_for" - are used by the
122+
// Front End to determine kernel invocation kind.
123+
template <typename KernelName, typename ElementType, typename KernelType,
124+
typename... Props>
125+
#ifdef __SYCL_DEVICE_ONLY__
126+
[[__sycl_detail__::add_ir_attributes_function(
127+
ext::oneapi::experimental::detail::PropertyMetaInfo<Props>::name...,
128+
ext::oneapi::experimental::detail::PropertyMetaInfo<Props>::value...)]]
129+
#endif
130+
__SYCL_KERNEL_ATTR__ static void
131+
kernel_parallel_for(const KernelType &KernelFunc, kernel_handler KH) {
132+
#ifdef __SYCL_DEVICE_ONLY__
133+
KernelFunc(detail::Builder::getElement(detail::declptr<ElementType>()), KH);
134+
#else
135+
(void)KernelFunc;
136+
(void)KH;
137+
#endif
138+
}
139+
140+
// NOTE: the name of this function - "kernel_parallel_for_work_group" - is
141+
// used by the Front End to determine kernel invocation kind.
142+
template <typename KernelName, typename ElementType, typename KernelType,
143+
typename... Props>
144+
#ifdef __SYCL_DEVICE_ONLY__
145+
[[__sycl_detail__::add_ir_attributes_function(
146+
ext::oneapi::experimental::detail::PropertyMetaInfo<Props>::name...,
147+
ext::oneapi::experimental::detail::PropertyMetaInfo<Props>::value...)]]
148+
#endif
149+
__SYCL_KERNEL_ATTR__ static void
150+
kernel_parallel_for_work_group(const KernelType &KernelFunc) {
151+
#ifdef __SYCL_DEVICE_ONLY__
152+
KernelFunc(detail::Builder::getElement(detail::declptr<ElementType>()));
153+
#else
154+
(void)KernelFunc;
155+
#endif
156+
}
157+
158+
// NOTE: the name of this function - "kernel_parallel_for_work_group" - is
159+
// used by the Front End to determine kernel invocation kind.
160+
template <typename KernelName, typename ElementType, typename KernelType,
161+
typename... Props>
162+
#ifdef __SYCL_DEVICE_ONLY__
163+
[[__sycl_detail__::add_ir_attributes_function(
164+
ext::oneapi::experimental::detail::PropertyMetaInfo<Props>::name...,
165+
ext::oneapi::experimental::detail::PropertyMetaInfo<Props>::value...)]]
166+
#endif
167+
__SYCL_KERNEL_ATTR__ static void
168+
kernel_parallel_for_work_group(const KernelType &KernelFunc,
169+
kernel_handler KH) {
170+
#ifdef __SYCL_DEVICE_ONLY__
171+
KernelFunc(detail::Builder::getElement(detail::declptr<ElementType>()), KH);
172+
#else
173+
(void)KernelFunc;
174+
(void)KH;
175+
#endif
176+
}
177+
}; // KernelWrapperSingletonFunc
178+
179+
// The KernelWrapper below has two purposes.
180+
//
181+
// First, from SYCL 2020, Table 129 (Member functions of the `handler ` class)
182+
// > The callable ... can optionally take a `kernel_handler` ... in
183+
// > which case the SYCL runtime will construct an instance of
184+
// > `kernel_handler` and pass it to the callable.
185+
//
186+
// Note: "..." due to slight wording variability between
187+
// single_task/parallel_for (e.g. only parameter vs last). This helper class
188+
// calls `kernel_*` entry points (both hardcoded names known to FE and special
189+
// device-specific entry point attributes) with proper arguments (with/without
190+
// `kernel_handler` argument, depending on the signature of the SYCL kernel
191+
// function).
192+
//
193+
// Second, it performs a few checks and some properties processing (including
194+
// the one provided via `sycl_ext_oneapi_kernel_properties` extension by
195+
// embedding them into the kernel's type).
196+
197+
template <WrapAs WrapAsVal, typename KernelName, typename KernelType,
198+
typename ElementType, typename PropertyProcessor,
199+
typename PropertiesT = ext::oneapi::experimental::empty_properties_t,
200+
typename MergedPropertiesT = typename detail::
201+
GetMergedKernelProperties<KernelType, PropertiesT>::type>
202+
struct KernelWrapper;
203+
template <WrapAs WrapAsVal, typename KernelName, typename KernelType,
204+
typename ElementType, typename PropertyProcessor,
205+
typename PropertiesT, typename... MergedProps>
206+
struct KernelWrapper<
207+
WrapAsVal, KernelName, KernelType, ElementType, PropertyProcessor,
208+
PropertiesT,
209+
ext::oneapi::experimental::detail::properties_t<MergedProps...>>
210+
: public KernelWrapperHelperFuncs {
211+
212+
static void wrap([[maybe_unused]] PropertyProcessor h,
213+
[[maybe_unused]] const KernelType &KernelFunc) {
214+
#ifdef __SYCL_DEVICE_ONLY__
215+
detail::CheckDeviceCopyable<KernelType>();
216+
#else
217+
// If there are properties provided by get method then process them.
218+
if constexpr (ext::oneapi::experimental::detail::
219+
HasKernelPropertiesGetMethod<const KernelType &>::value) {
220+
221+
// TODO: decouple property processing from KernelWrapper.
222+
h->template processProperties<detail::isKernelESIMD<KernelName>()>(
223+
KernelFunc.get(ext::oneapi::experimental::properties_tag{}));
224+
}
225+
#endif
226+
// Note: the static_assert below need to be run on both the host and the
227+
// device ends to avoid test issues, so don't put it into the #ifdef
228+
// __SYCL_DEVICE_ONLY__ directive above print out diagnostic message if
229+
// the kernel functor has a get(properties_tag) member, but it's not const
230+
static_assert(
231+
(ext::oneapi::experimental::detail::HasKernelPropertiesGetMethod<
232+
const KernelType &>::value) ||
233+
!(ext::oneapi::experimental::detail::HasKernelPropertiesGetMethod<
234+
KernelType>::value),
235+
"get(sycl::ext::oneapi::experimental::properties_tag) member in "
236+
"kernel functor class must be declared as a const member function");
237+
auto L = [&](auto &&...args) {
238+
if constexpr (WrapAsVal == WrapAs::single_task) {
239+
kernel_single_task<KernelName, KernelType, MergedProps...>(
240+
std::forward<decltype(args)>(args)...);
241+
} else if constexpr (WrapAsVal == WrapAs::parallel_for) {
242+
kernel_parallel_for<KernelName, ElementType, KernelType,
243+
MergedProps...>(
244+
std::forward<decltype(args)>(args)...);
245+
} else if constexpr (WrapAsVal == WrapAs::parallel_for_work_group) {
246+
kernel_parallel_for_work_group<KernelName, ElementType, KernelType,
247+
MergedProps...>(
248+
std::forward<decltype(args)>(args)...);
249+
} else {
250+
// Always false, but template-dependent. Can't compare `WrapAsVal`
251+
// with itself because of `-Wtautological-compare` warning.
252+
static_assert(!std::is_same_v<KernelName, KernelName>,
253+
"Unexpected WrapAsVal");
254+
}
255+
};
256+
if constexpr (detail::KernelLambdaHasKernelHandlerArgT<
257+
KernelType, ElementType>::value) {
258+
kernel_handler KH;
259+
L(KernelFunc, KH);
260+
} else {
261+
L(KernelFunc);
262+
}
263+
}
264+
}; // KernelWrapper struct
265+
266+
} // namespace detail
267+
} // namespace _V1
268+
} // namespace sycl

0 commit comments

Comments
 (0)