Skip to content

Commit a4f8899

Browse files
slawekptakslawekptak
authored
[SYCL] Fallback path for handler-less kernel properties (#20283)
Add a fallback path (handler-based submission) for the handler-less kernel submission path, if kernel function properties are provided.
1 parent b0f6da7 commit a4f8899

File tree

3 files changed

+60
-42
lines changed

3 files changed

+60
-42
lines changed

‎sycl/include/sycl/ext/oneapi/experimental/enqueue_functions.hpp‎

Lines changed: 13 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -260,8 +260,12 @@ template <typename KernelName = sycl::detail::auto_name, int Dimensions,
260260
void nd_launch(queue Q, nd_range<Dimensions> Range, const KernelType &KernelObj,
261261
ReductionsT &&...Reductions) {
262262
#ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
263-
// TODO The handler-less path does not support reductions yet.
264-
if constexpr (sizeof...(ReductionsT) == 0) {
263+
// TODO The handler-less path does not support reductions and kernel function
264+
// properties yet.
265+
if constexpr (sizeof...(ReductionsT) == 0 &&
266+
!(ext::oneapi::experimental::detail::
267+
HasKernelPropertiesGetMethod<
268+
const KernelType &>::value)) {
265269
detail::submit_kernel_direct<KernelName>(std::move(Q), empty_properties_t{},
266270
Range, KernelObj);
267271
} else
@@ -292,23 +296,13 @@ template <typename KernelName = sycl::detail::auto_name, int Dimensions,
292296
typename Properties, typename KernelType, typename... ReductionsT>
293297
void nd_launch(queue Q, launch_config<nd_range<Dimensions>, Properties> Config,
294298
const KernelType &KernelObj, ReductionsT &&...Reductions) {
295-
#ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
296-
// TODO The handler-less path does not support reductions yet.
297-
if constexpr (sizeof...(ReductionsT) == 0) {
298-
ext::oneapi::experimental::detail::LaunchConfigAccess<nd_range<Dimensions>,
299-
Properties>
300-
ConfigAccess(Config);
301-
detail::submit_kernel_direct<KernelName>(
302-
std::move(Q), ConfigAccess.getProperties(), ConfigAccess.getRange(),
303-
KernelObj);
304-
} else
305-
#endif
306-
{
307-
submit(std::move(Q), [&](handler &CGH) {
308-
nd_launch<KernelName>(CGH, Config, KernelObj,
309-
std::forward<ReductionsT>(Reductions)...);
310-
});
311-
}
299+
// TODO This overload of the nd_launch function takes the kernel function
300+
// properties, which are not yet supported for the handler-less path,
301+
// so it only supports handler based submission for now
302+
submit(std::move(Q), [&](handler &CGH) {
303+
nd_launch<KernelName>(CGH, Config, KernelObj,
304+
std::forward<ReductionsT>(Reductions)...);
305+
});
312306
}
313307

314308
template <int Dimensions, typename... ArgsT>

‎sycl/include/sycl/khr/free_function_commands.hpp‎

Lines changed: 39 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -158,44 +158,62 @@ void launch_grouped(const queue &q, range<1> r, range<1> size, KernelType &&k,
158158
const sycl::detail::code_location &codeLoc =
159159
sycl::detail::code_location::current()) {
160160
#ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
161-
detail::submit_kernel_direct(
162-
q, ext::oneapi::experimental::empty_properties_t{}, nd_range<1>(r, size),
163-
std::forward<KernelType>(k));
164-
#else
165-
submit(
166-
q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); },
167-
codeLoc);
161+
// TODO The handler-less path does not support kernel function properties yet.
162+
if constexpr (!(ext::oneapi::experimental::detail::
163+
HasKernelPropertiesGetMethod<
164+
const KernelType &>::value)) {
165+
detail::submit_kernel_direct(
166+
q, ext::oneapi::experimental::empty_properties_t{},
167+
nd_range<1>(r, size), std::forward<KernelType>(k));
168+
} else
168169
#endif
170+
{
171+
submit(
172+
q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); },
173+
codeLoc);
174+
}
169175
}
170176
template <typename KernelType, typename = typename std::enable_if_t<
171177
enable_kernel_function_overload<KernelType>>>
172178
void launch_grouped(const queue &q, range<2> r, range<2> size, KernelType &&k,
173179
const sycl::detail::code_location &codeLoc =
174180
sycl::detail::code_location::current()) {
175181
#ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
176-
detail::submit_kernel_direct(
177-
q, ext::oneapi::experimental::empty_properties_t{}, nd_range<2>(r, size),
178-
std::forward<KernelType>(k));
179-
#else
180-
submit(
181-
q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); },
182-
codeLoc);
182+
// TODO The handler-less path does not support kernel function properties yet.
183+
if constexpr (!(ext::oneapi::experimental::detail::
184+
HasKernelPropertiesGetMethod<
185+
const KernelType &>::value)) {
186+
detail::submit_kernel_direct(
187+
q, ext::oneapi::experimental::empty_properties_t{},
188+
nd_range<2>(r, size), std::forward<KernelType>(k));
189+
} else
183190
#endif
191+
{
192+
submit(
193+
q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); },
194+
codeLoc);
195+
}
184196
}
185197
template <typename KernelType, typename = typename std::enable_if_t<
186198
enable_kernel_function_overload<KernelType>>>
187199
void launch_grouped(const queue &q, range<3> r, range<3> size, KernelType &&k,
188200
const sycl::detail::code_location &codeLoc =
189201
sycl::detail::code_location::current()) {
190202
#ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
191-
detail::submit_kernel_direct(
192-
q, ext::oneapi::experimental::empty_properties_t{}, nd_range<3>(r, size),
193-
std::forward<KernelType>(k));
194-
#else
195-
submit(
196-
q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); },
197-
codeLoc);
203+
// TODO The handler-less path does not support kernel function properties yet.
204+
if constexpr (!(ext::oneapi::experimental::detail::
205+
HasKernelPropertiesGetMethod<
206+
const KernelType &>::value)) {
207+
detail::submit_kernel_direct(
208+
q, ext::oneapi::experimental::empty_properties_t{},
209+
nd_range<3>(r, size), std::forward<KernelType>(k));
210+
} else
198211
#endif
212+
{
213+
submit(
214+
q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); },
215+
codeLoc);
216+
}
199217
}
200218

201219
template <typename... Args>

‎sycl/include/sycl/queue.hpp‎

Lines changed: 8 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -3276,8 +3276,14 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
32763276
constexpr detail::code_location CodeLoc = getCodeLocation<KernelName>();
32773277
detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
32783278
#ifdef __DPCPP_ENABLE_UNFINISHED_NO_CGH_SUBMIT
3279-
// TODO The handler-less path does not support reductions yet.
3280-
if constexpr (sizeof...(RestT) == 1) {
3279+
using KernelType = std::tuple_element_t<0, std::tuple<RestT...>>;
3280+
3281+
// TODO The handler-less path does not support reductions and kernel
3282+
// function properties yet.
3283+
if constexpr (sizeof...(RestT) == 1 &&
3284+
!(ext::oneapi::experimental::detail::
3285+
HasKernelPropertiesGetMethod<
3286+
const KernelType &>::value)) {
32813287
return detail::submit_kernel_direct<KernelName, true>(
32823288
*this, ext::oneapi::experimental::empty_properties_t{}, Range,
32833289
Rest...);

0 commit comments

Comments
 (0)