[SYCL] Extend no-handler submission path to support kernel properties. (intel#20524)

Author: uditagarwal97

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

@@ -152,15 +152,12 @@ template <typename KernelName = sycl::detail::auto_name, typename KernelType>
 void single_task(queue Q, const KernelType &KernelObj,
  const sycl::detail::code_location &CodeLoc =
  sycl::detail::code_location::current()) {
- // TODO The handler-less path does not support kernel function properties
- // and kernel functions with the kernel_handler type argument yet.
- if constexpr (!(ext::oneapi::experimental::detail::
- HasKernelPropertiesGetMethod<
- const KernelType &>::value) &&
- !(detail::KernelLambdaHasKernelHandlerArgT<KernelType,
+ // TODO The handler-less path does not support kernel functions with the
+ // kernel_handler type argument yet.
+ if constexpr (!(detail::KernelLambdaHasKernelHandlerArgT<KernelType,
  void>::value)) {
  detail::submit_kernel_direct_single_task<KernelName>(
- std::move(Q), empty_properties_t{}, KernelObj, CodeLoc);
+ std::move(Q), KernelObj, empty_properties_t{}, CodeLoc);
  } else {
  submit(
  std::move(Q),
@@ -271,17 +268,13 @@ template <typename KernelName = sycl::detail::auto_name, int Dimensions,
  typename KernelType, typename... ReductionsT>
 void nd_launch(queue Q, nd_range<Dimensions> Range, const KernelType &KernelObj,
  ReductionsT &&...Reductions) {
- // TODO The handler-less path does not support reductions, kernel
- // function properties and kernel functions with the kernel_handler
- // type argument yet.
+ // TODO The handler-less path does not support reductions, and
+ // kernel functions with the kernel_handler type argument yet.
  if constexpr (sizeof...(ReductionsT) == 0 &&
- !(ext::oneapi::experimental::detail::
- HasKernelPropertiesGetMethod<
- const KernelType &>::value) &&
  !(detail::KernelLambdaHasKernelHandlerArgT<
  KernelType, sycl::nd_item<Dimensions>>::value)) {
- detail::submit_kernel_direct_parallel_for<KernelName>(
- std::move(Q), empty_properties_t{}, Range, KernelObj);
+ detail::submit_kernel_direct_parallel_for<KernelName>(std::move(Q), Range,
+   KernelObj);
  } else {
  submit(std::move(Q), [&](handler &CGH) {
  nd_launch<KernelName>(CGH, Range, KernelObj,
@@ -308,13 +301,25 @@ template <typename KernelName = sycl::detail::auto_name, int Dimensions,
  typename Properties, typename KernelType, typename... ReductionsT>
 void nd_launch(queue Q, launch_config<nd_range<Dimensions>, Properties> Config,
  const KernelType &KernelObj, ReductionsT &&...Reductions) {
- // TODO This overload of the nd_launch function takes the kernel function
- // properties, which are not yet supported for the handler-less path,
- // so it only supports handler based submission for now
- submit(std::move(Q), [&](handler &CGH) {
- nd_launch<KernelName>(CGH, Config, KernelObj,
- std::forward<ReductionsT>(Reductions)...);
- });
+ // TODO The handler-less path does not support reductions, and
+ // kernel functions with the kernel_handler type argument yet.
+ if constexpr (sizeof...(ReductionsT) == 0 &&
+ !(detail::KernelLambdaHasKernelHandlerArgT<
+ KernelType, sycl::nd_item<Dimensions>>::value)) {
+
+ ext::oneapi::experimental::detail::LaunchConfigAccess<nd_range<Dimensions>,
+ Properties>
+ LaunchConfigAccess(Config);
+
+ detail::submit_kernel_direct_parallel_for<KernelName>(
+ std::move(Q), LaunchConfigAccess.getRange(), KernelObj,
+ LaunchConfigAccess.getProperties());
+ } else {
+ submit(std::move(Q), [&](handler &CGH) {
+ nd_launch<KernelName>(CGH, Config, KernelObj,
+ std::forward<ReductionsT>(Reductions)...);
+ });
+ }
 }
 
 template <int Dimensions, typename... ArgsT>

sycl/include/sycl/khr/free_function_commands.hpp

@@ -157,16 +157,12 @@ template <typename KernelType, typename = typename std::enable_if_t<
 void launch_grouped(const queue &q, range<1> r, range<1> size, KernelType &&k,
  const sycl::detail::code_location &codeLoc =
  sycl::detail::code_location::current()) {
- // TODO The handler-less path does not support kernel function properties
- // and kernel functions with the kernel_handler type argument yet.
- if constexpr (!(ext::oneapi::experimental::detail::
- HasKernelPropertiesGetMethod<
- const KernelType &>::value) &&
- !(detail::KernelLambdaHasKernelHandlerArgT<
+ // TODO The handler-less path does not support kernel functions with the
+ // kernel_handler type argument yet.
+ if constexpr (!(detail::KernelLambdaHasKernelHandlerArgT<
  KernelType, sycl::nd_item<1>>::value)) {
- detail::submit_kernel_direct_parallel_for(
- q, ext::oneapi::experimental::empty_properties_t{},
- nd_range<1>(r, size), std::forward<KernelType>(k));
+ detail::submit_kernel_direct_parallel_for(q, nd_range<1>(r, size),
+ std::forward<KernelType>(k));
  } else {
  submit(
  q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); },
@@ -178,16 +174,12 @@ template <typename KernelType, typename = typename std::enable_if_t<
 void launch_grouped(const queue &q, range<2> r, range<2> size, KernelType &&k,
  const sycl::detail::code_location &codeLoc =
  sycl::detail::code_location::current()) {
- // TODO The handler-less path does not support kernel function properties
- // and kernel functions with the kernel_handler type argument yet.
- if constexpr (!(ext::oneapi::experimental::detail::
- HasKernelPropertiesGetMethod<
- const KernelType &>::value) &&
- !(detail::KernelLambdaHasKernelHandlerArgT<
+ // TODO The handler-less path does not support kernel functions with the
+ // kernel_handler type argument yet.
+ if constexpr (!(detail::KernelLambdaHasKernelHandlerArgT<
  KernelType, sycl::nd_item<2>>::value)) {
- detail::submit_kernel_direct_parallel_for(
- q, ext::oneapi::experimental::empty_properties_t{},
- nd_range<2>(r, size), std::forward<KernelType>(k));
+ detail::submit_kernel_direct_parallel_for(q, nd_range<2>(r, size),
+ std::forward<KernelType>(k));
  } else {
  submit(
  q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); },
@@ -199,16 +191,12 @@ template <typename KernelType, typename = typename std::enable_if_t<
 void launch_grouped(const queue &q, range<3> r, range<3> size, KernelType &&k,
  const sycl::detail::code_location &codeLoc =
  sycl::detail::code_location::current()) {
- // TODO The handler-less path does not support kernel function properties
- // and kernel functions with the kernel_handler type argument yet.
- if constexpr (!(ext::oneapi::experimental::detail::
- HasKernelPropertiesGetMethod<
- const KernelType &>::value) &&
- !(detail::KernelLambdaHasKernelHandlerArgT<
+ // TODO The handler-less path does not support kernel functions with the
+ // kernel_handler type argument yet.
+ if constexpr (!(detail::KernelLambdaHasKernelHandlerArgT<
  KernelType, sycl::nd_item<3>>::value)) {
- detail::submit_kernel_direct_parallel_for(
- q, ext::oneapi::experimental::empty_properties_t{},
- nd_range<3>(r, size), std::forward<KernelType>(k));
+ detail::submit_kernel_direct_parallel_for(q, nd_range<3>(r, size),
+ std::forward<KernelType>(k));
  } else {
  submit(
  q, [&](handler &h) { launch_grouped<KernelType>(h, r, size, k); },
@@ -324,16 +312,13 @@ template <typename KernelType, typename = typename std::enable_if_t<
 void launch_task(const sycl::queue &q, KernelType &&k,
  const sycl::detail::code_location &codeLoc =
  sycl::detail::code_location::current()) {
- // TODO The handler-less path does not support kernel function properties
- // and kernel functions with the kernel_handler type argument yet.
- if constexpr (!(ext::oneapi::experimental::detail::
- HasKernelPropertiesGetMethod<
- const KernelType &>::value) &&
- !(detail::KernelLambdaHasKernelHandlerArgT<KernelType,
+ // TODO The handler-less path does not support kernel functions with the
+ // kernel_handler type argument yet.
+ if constexpr (!(detail::KernelLambdaHasKernelHandlerArgT<KernelType,
  void>::value)) {
  detail::submit_kernel_direct_single_task(
- q, ext::oneapi::experimental::empty_properties_t{},
- std::forward<KernelType>(k), codeLoc);
+ q, std::forward<KernelType>(k),
+ ext::oneapi::experimental::empty_properties_t{}, codeLoc);
  } else {
  submit(q, [&](handler &h) { launch_task<KernelType>(h, k); }, codeLoc);
  }

sycl/include/sycl/queue.hpp

@@ -68,13 +68,15 @@ event __SYCL_EXPORT submit_kernel_direct_with_event_impl(
  const queue &Queue, const nd_range<Dims> &Range,
  detail::HostKernelRefBase &HostKernel,
  detail::DeviceKernelInfo *DeviceKernelInfo,
+ const detail::KernelPropertyHolderStructTy &Props,
  const detail::code_location &CodeLoc, bool IsTopCodeLoc);
 
 template <int Dims>
 void __SYCL_EXPORT submit_kernel_direct_without_event_impl(
  const queue &Queue, const nd_range<Dims> &Range,
  detail::HostKernelRefBase &HostKernel,
  detail::DeviceKernelInfo *DeviceKernelInfo,
+ const detail::KernelPropertyHolderStructTy &Props,
  const detail::code_location &CodeLoc, bool IsTopCodeLoc);
 
 namespace detail {
@@ -159,16 +161,14 @@ class __SYCL_EXPORT SubmissionInfo {
 
 template <detail::WrapAs WrapAs, typename LambdaArgType,
  typename KernelName = detail::auto_name, bool EventNeeded = false,
- typename PropertiesT, typename KernelTypeUniversalRef, int Dims>
+ typename PropertiesT = ext::oneapi::experimental::empty_properties_t,
+ typename KernelTypeUniversalRef, int Dims>
 auto submit_kernel_direct(
- const queue &Queue, [[maybe_unused]] PropertiesT Props,
- const nd_range<Dims> &Range, KernelTypeUniversalRef &&KernelFunc,
+ const queue &Queue, const nd_range<Dims> &Range,
+ KernelTypeUniversalRef &&KernelFunc,
+ const PropertiesT &ExtraProps =
+ ext::oneapi::experimental::empty_properties_t{},
  const detail::code_location &CodeLoc = detail::code_location::current()) {
- // TODO Properties not supported yet
- static_assert(
- std::is_same_v<PropertiesT,
- ext::oneapi::experimental::empty_properties_t>,
- "Setting properties not supported yet for no-CGH kernel submit.");
  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
 
  using KernelType =
@@ -210,22 +210,42 @@ auto submit_kernel_direct(
  "-fsycl-host-compiler-options='/std:c++latest' "
  "might also help.");
 
+ detail::KernelPropertyHolderStructTy ParsedProperties;
+ if constexpr (ext::oneapi::experimental::detail::HasKernelPropertiesGetMethod<
+ const KernelType &>::value) {
+ // Merge properties via get() and manually specified properties.
+ // get() method is used for specifying kernel properties but properties
+ // passed via launch_config (ExtraProps) should be kernel launch properties.
+ // They are mutually exclusive, so there should not be any conflict when
+ // merging properties. merge_properties() throws if there's a conflict.
+ auto MergedProps =
+ sycl::ext::oneapi::experimental::detail::merge_properties(
+ ExtraProps,
+ KernelFunc.get(ext::oneapi::experimental::properties_tag{}));
+
+ ParsedProperties = extractKernelProperties(MergedProps);
+ } else {
+ ParsedProperties = extractKernelProperties(ExtraProps);
+ }
+
  if constexpr (EventNeeded) {
  return submit_kernel_direct_with_event_impl(
- Queue, Range, HostKernel, DeviceKernelInfoPtr,
+ Queue, Range, HostKernel, DeviceKernelInfoPtr, ParsedProperties,
  TlsCodeLocCapture.query(), TlsCodeLocCapture.isToplevel());
  } else {
  submit_kernel_direct_without_event_impl(
- Queue, Range, HostKernel, DeviceKernelInfoPtr,
+ Queue, Range, HostKernel, DeviceKernelInfoPtr, ParsedProperties,
  TlsCodeLocCapture.query(), TlsCodeLocCapture.isToplevel());
  }
 }
 
 template <typename KernelName = detail::auto_name, bool EventNeeded = false,
- typename PropertiesT, typename KernelTypeUniversalRef, int Dims>
+ typename PropertiesT = ext::oneapi::experimental::empty_properties_t,
+ typename KernelTypeUniversalRef, int Dims>
 auto submit_kernel_direct_parallel_for(
- const queue &Queue, PropertiesT Props, const nd_range<Dims> &Range,
+ const queue &Queue, const nd_range<Dims> &Range,
  KernelTypeUniversalRef &&KernelFunc,
+ const PropertiesT &Props = ext::oneapi::experimental::empty_properties_t{},
  const detail::code_location &CodeLoc = detail::code_location::current()) {
 
  using KernelType =
@@ -246,21 +266,23 @@ auto submit_kernel_direct_parallel_for(
  return submit_kernel_direct<detail::WrapAs::parallel_for, TransformedArgType,
  KernelName, EventNeeded, PropertiesT,
  KernelTypeUniversalRef, Dims>(
- Queue, Props, Range, std::forward<KernelTypeUniversalRef>(KernelFunc),
+ Queue, Range, std::forward<KernelTypeUniversalRef>(KernelFunc), Props,
  CodeLoc);
 }
 
 template <typename KernelName = detail::auto_name, bool EventNeeded = false,
- typename PropertiesT, typename KernelTypeUniversalRef>
+ typename PropertiesT = ext::oneapi::experimental::empty_properties_t,
+ typename KernelTypeUniversalRef>
 auto submit_kernel_direct_single_task(
- const queue &Queue, PropertiesT Props, KernelTypeUniversalRef &&KernelFunc,
+ const queue &Queue, KernelTypeUniversalRef &&KernelFunc,
+ const PropertiesT &Props = ext::oneapi::experimental::empty_properties_t{},
  const detail::code_location &CodeLoc = detail::code_location::current()) {
 
  return submit_kernel_direct<detail::WrapAs::single_task, void, KernelName,
  EventNeeded, PropertiesT, KernelTypeUniversalRef,
  1>(
- Queue, Props, nd_range<1>{1, 1},
- std::forward<KernelTypeUniversalRef>(KernelFunc), CodeLoc);
+ Queue, nd_range<1>{1, 1},
+ std::forward<KernelTypeUniversalRef>(KernelFunc), Props, CodeLoc);
 }
 
 } // namespace detail
@@ -2775,18 +2797,12 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
 
  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
 
- // TODO The handler-less path does not support kernel
- // function properties and kernel functions with the kernel_handler
- // type argument yet.
- if constexpr (
- std::is_same_v<PropertiesT,
- ext::oneapi::experimental::empty_properties_t> &&
- !(ext::oneapi::experimental::detail::HasKernelPropertiesGetMethod<
- const KernelType &>::value) &&
- !(detail::KernelLambdaHasKernelHandlerArgT<KernelType, void>::value)) {
+ // TODO The handler-less path does not support kernel functions
+ // with the kernel_handler type argument yet.
+ if constexpr (!(detail::KernelLambdaHasKernelHandlerArgT<KernelType,
+ void>::value)) {
  return detail::submit_kernel_direct_single_task<KernelName, true>(
- *this, ext::oneapi::experimental::empty_properties_t{}, KernelFunc,
- TlsCodeLocCapture.query());
+ *this, KernelFunc, Properties, TlsCodeLocCapture.query());
  } else {
  return submit(
  [&](handler &CGH) {
@@ -3323,11 +3339,22 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
  RestT &&...Rest) {
  constexpr detail::code_location CodeLoc = getCodeLocation<KernelName>();
  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
- return submit(
- [&](handler &CGH) {
- CGH.template parallel_for<KernelName>(Range, Properties, Rest...);
- },
- TlsCodeLocCapture.query());
+ using KernelType = std::tuple_element_t<0, std::tuple<RestT...>>;
+
+ // TODO The handler-less path does not support reductions, and
+ // kernel functions with the kernel_handler type argument yet.
+ if constexpr (sizeof...(RestT) == 1 &&
+ !(detail::KernelLambdaHasKernelHandlerArgT<
+ KernelType, sycl::nd_item<Dims>>::value)) {
+
+ return detail::submit_kernel_direct_parallel_for<KernelName, true>(
+ *this, Range, Rest..., Properties, TlsCodeLocCapture.query());
+ } else
+ return submit(
+ [&](handler &CGH) {
+ CGH.template parallel_for<KernelName>(Range, Properties, Rest...);
+ },
+ TlsCodeLocCapture.query());
  }
 
  /// parallel_for version with a kernel represented as a lambda + nd_range that
@@ -3344,18 +3371,15 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
  detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
  using KernelType = std::tuple_element_t<0, std::tuple<RestT...>>;
 
- // TODO The handler-less path does not support reductions, kernel
- // function properties and kernel functions with the kernel_handler
- // type argument yet.
+ // TODO The handler-less path does not support reductions, and
+ // kernel functions with the kernel_handler type argument yet.
  if constexpr (sizeof...(RestT) == 1 &&
- !(ext::oneapi::experimental::detail::
- HasKernelPropertiesGetMethod<
- const KernelType &>::value) &&
  !(detail::KernelLambdaHasKernelHandlerArgT<
  KernelType, sycl::nd_item<Dims>>::value)) {
  return detail::submit_kernel_direct_parallel_for<KernelName, true>(
- *this, ext::oneapi::experimental::empty_properties_t{}, Range,
- Rest..., TlsCodeLocCapture.query());
+ *this, Range, Rest...,
+ ext::oneapi::experimental::empty_properties_t{},
+ TlsCodeLocCapture.query());
  } else {
  return submit(
  [&](handler &CGH) {

sycl/source/detail/queue_impl.cpp

@@ -567,13 +567,18 @@ EventImplPtr queue_impl::submit_command_to_graph(
 EventImplPtr queue_impl::submit_kernel_direct_impl(
  const NDRDescT &NDRDesc, detail::HostKernelRefBase &HostKernel,
  detail::DeviceKernelInfo *DeviceKernelInfo, bool CallerNeedsEvent,
+ const detail::KernelPropertyHolderStructTy &Props,
  const detail::code_location &CodeLoc, bool IsTopCodeLoc) {
 
  KernelData KData;
 
  KData.setDeviceKernelInfoPtr(DeviceKernelInfo);
  KData.setNDRDesc(NDRDesc);
 
+ // Validate and set kernel launch properties.
+ KData.validateAndSetKernelLaunchProperties(Props, hasCommandGraph(),
+ getDeviceImpl());
+
  auto SubmitKernelFunc = [&](detail::CG::StorageInitHelper &CGData,
  bool SchedulerBypass) -> EventImplPtr {
  if (SchedulerBypass) {