[llvm] [libsycl] add initial sycl::handler impl (PR #216754)
Kseniya Tikhomirova via llvm-commits
llvm-commits at lists.llvm.org
Tue Aug 25 05:03:19 PDT 2026
https://github.com/KseniyaTikhomirova updated https://github.com/llvm/llvm-project/pull/216754
>From 61f6e2babf8f0c563e3f4dbb80bc7de0f0e13b3c Mon Sep 17 00:00:00 2001
From: "Tikhomirova, Kseniya" <kseniya.tikhomirova at intel.com>
Date: Thu, 2 Jul 2026 06:01:02 -0700
Subject: [PATCH 1/2] add base handler impl
Signed-off-by: Tikhomirova, Kseniya <kseniya.tikhomirova at intel.com>
---
.../sycl/__impl/detail/kernel_submission.hpp | 122 ++++++++++++
.../sycl/__impl/detail/unified_range_view.hpp | 1 +
libsycl/include/sycl/__impl/handler.hpp | 157 ++++++++++++++++
.../include/sycl/__impl/info/desc_base.hpp | 2 +
libsycl/include/sycl/__impl/property_list.hpp | 2 +
libsycl/include/sycl/__impl/queue.hpp | 173 ++++++++----------
libsycl/src/CMakeLists.txt | 1 +
libsycl/src/detail/handler_impl.hpp | 62 +++++++
libsycl/src/detail/offload/offload_utils.cpp | 47 +++++
libsycl/src/detail/offload/offload_utils.hpp | 5 +
libsycl/src/detail/queue_impl.cpp | 92 +++++-----
libsycl/src/detail/queue_impl.hpp | 34 +++-
libsycl/src/handler.cpp | 61 ++++++
libsycl/src/queue.cpp | 10 +-
.../test/basic/handler/handler_depends_on.cpp | 82 +++++++++
.../handler_parallel_for_arg_restrictions.cpp | 50 +++++
.../handler_parallel_for_generic_lambda.cpp | 43 +++++
...parallel_for_nd_range_invalid_arg_type.cpp | 17 ++
.../handler/handler_parallel_for_runtime.cpp | 149 +++++++++++++++
.../handler_unnamed_lambda_functor.cpp | 19 ++
.../basic/handler/submit_fn_ptr_handler.cpp | 24 +++
libsycl/unittests/CMakeLists.txt | 1 +
libsycl/unittests/event/event.cpp | 4 +-
libsycl/unittests/handler/CMakeLists.txt | 4 +
libsycl/unittests/handler/memcpy.cpp | 70 +++++++
libsycl/unittests/handler/semantics.cpp | 164 +++++++++++++++++
libsycl/unittests/handler/test_helpers.hpp | 37 ++++
.../unittests/queue/sycl_kernel_launch.cpp | 12 +-
28 files changed, 1282 insertions(+), 163 deletions(-)
create mode 100644 libsycl/include/sycl/__impl/detail/kernel_submission.hpp
create mode 100644 libsycl/include/sycl/__impl/handler.hpp
create mode 100644 libsycl/src/detail/handler_impl.hpp
create mode 100644 libsycl/src/handler.cpp
create mode 100644 libsycl/test/basic/handler/handler_depends_on.cpp
create mode 100644 libsycl/test/basic/handler/handler_parallel_for_arg_restrictions.cpp
create mode 100644 libsycl/test/basic/handler/handler_parallel_for_generic_lambda.cpp
create mode 100644 libsycl/test/basic/handler/handler_parallel_for_nd_range_invalid_arg_type.cpp
create mode 100644 libsycl/test/basic/handler/handler_parallel_for_runtime.cpp
create mode 100644 libsycl/test/basic/handler/handler_unnamed_lambda_functor.cpp
create mode 100644 libsycl/test/basic/handler/submit_fn_ptr_handler.cpp
create mode 100644 libsycl/unittests/handler/CMakeLists.txt
create mode 100644 libsycl/unittests/handler/memcpy.cpp
create mode 100644 libsycl/unittests/handler/semantics.cpp
create mode 100644 libsycl/unittests/handler/test_helpers.hpp
diff --git a/libsycl/include/sycl/__impl/detail/kernel_submission.hpp b/libsycl/include/sycl/__impl/detail/kernel_submission.hpp
new file mode 100644
index 0000000000000..b7381c46bf820
--- /dev/null
+++ b/libsycl/include/sycl/__impl/detail/kernel_submission.hpp
@@ -0,0 +1,122 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// This file contains helpers for kernel submission entry points.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef _LIBSYCL___IMPL_DETAIL_KERNEL_SUBMISSION_HPP
+#define _LIBSYCL___IMPL_DETAIL_KERNEL_SUBMISSION_HPP
+
+#include <sycl/__impl/detail/config.hpp>
+#include <sycl/__impl/detail/get_device_kernel_info.hpp>
+#include <sycl/__impl/detail/kernel_arg_helpers.hpp>
+#include <sycl/__impl/exception.hpp>
+#include <sycl/__impl/index_space_classes.hpp>
+#include <sycl/__impl/nd_item.hpp>
+#include <sycl/__impl/nd_range.hpp>
+
+#include <tuple>
+#include <utility>
+
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+
+namespace detail {
+
+template <int Dims>
+void checkNDRangeAndThrow(const sycl::nd_range<Dims> executionRange) {
+ if (executionRange.get_global_range() != range<Dims>{} &&
+ (executionRange.get_local_range().size() == 0 ||
+ executionRange.get_global_range() % executionRange.get_local_range() !=
+ range<Dims>{}))
+ throw sycl::exception(sycl::make_error_code(sycl::errc::nd_range),
+ "Invalid nd_range submission: global size must be "
+ "evenly divisible by local size.");
+}
+
+template <typename DerivedT> class KernelSubmissionBase {
+protected:
+ template <typename KernelName, typename KernelType>
+#ifdef SYCL_LANGUAGE_VERSION
+ [[clang::sycl_kernel_entry_point(KernelName)]]
+#endif
+ void submitSingleTask(const KernelType &KernelFunc) {
+ KernelFunc();
+ }
+
+ template <typename KernelName, typename ElementType, typename KernelType>
+#ifdef SYCL_LANGUAGE_VERSION
+ [[clang::sycl_kernel_entry_point(KernelName)]]
+#endif
+ void submitParallelFor(const KernelType &KernelFunc) {
+ KernelFunc(detail::Builder::getElement(detail::declptr<ElementType>()));
+ }
+
+ template <typename KN, typename... Args>
+ void sycl_kernel_launch(const char *KernelName, Args &&...args) {
+ static_assert(
+ sizeof...(args) == 1,
+ "sycl_kernel_launch expects only 2 arguments now: name of kernel and "
+ "callable object passed to kernel invocation by the user.");
+
+ auto FirstArg = std::get<0>(std::tie(args...));
+ static_cast<DerivedT *>(this)->submitKernelImpl(
+ detail::getDeviceKernelInfo<KN>(KernelName), &FirstArg,
+ sizeof(FirstArg));
+ }
+
+ template <typename KernelName, int Dims, template <int> class Range,
+ typename... Rest>
+ void parallelForImpl(Range<Dims> numWorkItems, Rest &&...rest) {
+ if constexpr (sizeof...(Rest) != 1)
+ throw sycl::exception(errc::feature_not_supported,
+ "Reductions are not supported");
+
+ using KernelType =
+ std::decay_t<detail::nth_type_t<sizeof...(Rest) - 1, Rest...>>;
+ constexpr bool IsNdRangeSubmission =
+ std::is_same_v<Range<Dims>, nd_range<Dims>>;
+ using SuggestedArgType =
+ std::conditional_t<IsNdRangeSubmission, nd_item<Dims>, item<Dims>>;
+ using LambdaArgType =
+ sycl::detail::lambda_arg_type<KernelType, SuggestedArgType>;
+
+ if constexpr (IsNdRangeSubmission) {
+ static_assert(
+ std::is_convertible_v<sycl::nd_item<Dims>, LambdaArgType>,
+ "Kernel argument of a sycl::parallel_for with sycl::nd_range "
+ "must be sycl::nd_item or be convertible from sycl::nd_item");
+ } else {
+ static_assert(
+ std::is_convertible_v<sycl::item<Dims>, LambdaArgType> ||
+ std::is_convertible_v<sycl::item<Dims, false>, LambdaArgType>,
+ "Kernel argument of a sycl::parallel_for with sycl::range "
+ "must be sycl::item or be convertible from sycl::item");
+ }
+
+ using TransformedLambdaArgType = std::conditional_t<
+ IsNdRangeSubmission, nd_item<Dims>,
+ std::conditional_t<
+ std::is_convertible_v<sycl::item<Dims>, LambdaArgType>, item<Dims>,
+ std::conditional_t<
+ std::is_convertible_v<sycl::item<Dims, false>, LambdaArgType>,
+ item<Dims, false>, LambdaArgType>>>;
+
+ using NameT =
+ typename detail::get_kernel_name_t<KernelName, KernelType>::name;
+ return submitParallelFor<NameT, TransformedLambdaArgType, KernelType>(
+ std::forward<Rest>(rest)...);
+ }
+};
+
+} // namespace detail
+
+_LIBSYCL_END_NAMESPACE_SYCL
+
+#endif // _LIBSYCL___IMPL_DETAIL_KERNEL_SUBMISSION_HPP
diff --git a/libsycl/include/sycl/__impl/detail/unified_range_view.hpp b/libsycl/include/sycl/__impl/detail/unified_range_view.hpp
index be53d0e6331b6..ed8b789b47b88 100644
--- a/libsycl/include/sycl/__impl/detail/unified_range_view.hpp
+++ b/libsycl/include/sycl/__impl/detail/unified_range_view.hpp
@@ -17,6 +17,7 @@
#include <sycl/__impl/detail/config.hpp>
#include <sycl/__impl/index_space_classes.hpp>
+#include <sycl/__impl/nd_range.hpp>
_LIBSYCL_BEGIN_NAMESPACE_SYCL
diff --git a/libsycl/include/sycl/__impl/handler.hpp b/libsycl/include/sycl/__impl/handler.hpp
new file mode 100644
index 0000000000000..88f8b479ba850
--- /dev/null
+++ b/libsycl/include/sycl/__impl/handler.hpp
@@ -0,0 +1,157 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+///
+/// This file contains the declaration of the SYCL handler class, which provides
+/// the interface for the commands that can be executed inside the command group
+/// scope.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef _LIBSYCL___IMPL_HANDLER_HPP
+#define _LIBSYCL___IMPL_HANDLER_HPP
+
+#include <sycl/__impl/detail/config.hpp>
+#include <sycl/__impl/detail/get_device_kernel_info.hpp>
+#include <sycl/__impl/detail/kernel_arg_helpers.hpp>
+#include <sycl/__impl/detail/kernel_submission.hpp>
+#include <sycl/__impl/detail/unified_range_view.hpp>
+#include <sycl/__impl/event.hpp>
+#include <sycl/__impl/exception.hpp>
+#include <sycl/__impl/index_space_classes.hpp>
+
+#include <array>
+#include <cstring>
+#include <memory>
+#include <type_traits>
+#include <vector>
+
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+
+namespace detail {
+class HandlerImpl;
+class QueueImpl;
+} // namespace detail
+
+class _LIBSYCL_EXPORT handler : private detail::KernelSubmissionBase<handler> {
+public:
+ handler(const handler &) = delete;
+ handler(handler &&) = delete;
+ handler &operator=(const handler &) = delete;
+ handler &operator=(handler &&) = delete;
+
+ ~handler() = default;
+
+ /// Adds an event dependency to this command group.
+ ///
+ /// \param depEvent is the event that must complete before this command
+ /// group is executed.
+ void depends_on(event depEvent) {
+ return depends_on(std::vector<event>{depEvent});
+ }
+
+ /// Adds event dependencies to this command group.
+ ///
+ /// \param depEvents are the events that must complete before this command
+ /// group is executed.
+ void depends_on(const std::vector<event> &depEvents) {
+ MDepEvents.insert(MDepEvents.end(), depEvents.begin(), depEvents.end());
+ }
+
+ /// Defines a single-task kernel for this command group.
+ ///
+ /// \param kernelFunc is the kernel functor or lambda to execute.
+ template <typename KernelName = detail::AutoName, typename KernelType>
+ void single_task(const KernelType &kernelFunc) {
+ setKernelRange();
+ submitSingleTask<KernelName, KernelType>(kernelFunc);
+ }
+
+ /// Defines a one-dimensional range kernel for this command group.
+ ///
+ /// \param numWorkItems is the number of work-items in the kernel range.
+ /// \param rest contains the kernel functor or lambda and its optional
+ /// submission arguments.
+ template <typename KernelName = detail::AutoName, typename... Rest>
+ void parallel_for(range<1> numWorkItems, Rest &&...rest) {
+ return parallelForImpl<KernelName>(numWorkItems,
+ std::forward<Rest>(rest)...);
+ }
+
+ /// Defines a two-dimensional range kernel for this command group.
+ ///
+ /// \param numWorkItems is the number of work-items in the kernel range.
+ /// \param rest contains the kernel functor or lambda and its optional
+ /// submission arguments.
+ template <typename KernelName = detail::AutoName, typename... Rest>
+ void parallel_for(range<2> numWorkItems, Rest &&...rest) {
+ return parallelForImpl<KernelName>(numWorkItems,
+ std::forward<Rest>(rest)...);
+ }
+
+ /// Defines a three-dimensional range kernel for this command group.
+ ///
+ /// \param numWorkItems is the number of work-items in the kernel range.
+ /// \param rest contains the kernel functor or lambda and its optional
+ /// submission arguments.
+ template <typename KernelName = detail::AutoName, typename... Rest>
+ void parallel_for(range<3> numWorkItems, Rest &&...rest) {
+ return parallelForImpl<KernelName>(numWorkItems,
+ std::forward<Rest>(rest)...);
+ }
+
+ /// Defines an nd-range kernel for this command group.
+ ///
+ /// \param executionRange is the global and local range of the kernel.
+ /// \param rest contains the kernel functor or lambda and its optional
+ /// submission arguments.
+ template <typename KernelName = detail::AutoName, int Dims, typename... Rest>
+ void parallel_for(nd_range<Dims> executionRange, Rest &&...rest) {
+ detail::checkNDRangeAndThrow(executionRange);
+
+ return parallelForImpl<KernelName>(executionRange,
+ std::forward<Rest>(rest)...);
+ }
+
+ /// Defines a memory copy operation for this command group.
+ ///
+ /// \param dest is the destination memory address.
+ /// \param src is the source memory address.
+ /// \param numBytes is the number of bytes to copy.
+ void memcpy(void *dest, const void *src, std::size_t numBytes);
+
+private:
+ template <typename KernelName, int Dims, template <int> class Range,
+ typename... Rest>
+ void parallelForImpl(Range<Dims> numWorkItems, Rest &&...rest) {
+ setKernelRange(numWorkItems);
+
+ detail::KernelSubmissionBase<handler>::template parallelForImpl<KernelName>(
+ numWorkItems, std::forward<Rest>(rest)...);
+ }
+
+ std::shared_ptr<detail::EventImpl> finalize();
+
+ void submitKernelImpl(detail::DeviceKernelInfo &KernelInfo, void *ArgData,
+ size_t ArgSize);
+
+ void setKernelRange(const detail::UnifiedRangeView &Range = {});
+
+ handler(detail::HandlerImpl &HandlerImplVal) : MImpl(HandlerImplVal) {}
+
+ std::vector<event> MDepEvents;
+
+ detail::HandlerImpl &MImpl;
+
+ friend sycl::detail::ImplUtils;
+ friend class detail::QueueImpl;
+ friend class detail::KernelSubmissionBase<handler>;
+};
+
+_LIBSYCL_END_NAMESPACE_SYCL
+
+#endif // _LIBSYCL___IMPL_HANDLER_HPP
diff --git a/libsycl/include/sycl/__impl/info/desc_base.hpp b/libsycl/include/sycl/__impl/info/desc_base.hpp
index 0fc4284d60b68..8f4a7ad8959fb 100644
--- a/libsycl/include/sycl/__impl/info/desc_base.hpp
+++ b/libsycl/include/sycl/__impl/info/desc_base.hpp
@@ -16,6 +16,8 @@
#include <sycl/__impl/detail/config.hpp>
+#include <type_traits>
+
_LIBSYCL_BEGIN_NAMESPACE_SYCL
namespace detail {
diff --git a/libsycl/include/sycl/__impl/property_list.hpp b/libsycl/include/sycl/__impl/property_list.hpp
index 8b1da2c24e6dc..d12101ec76180 100644
--- a/libsycl/include/sycl/__impl/property_list.hpp
+++ b/libsycl/include/sycl/__impl/property_list.hpp
@@ -17,6 +17,8 @@
#ifndef _LIBSYCL___IMPL_PROPERTY_LIST_HPP
#define _LIBSYCL___IMPL_PROPERTY_LIST_HPP
+#include <sycl/__impl/detail/config.hpp>
+
_LIBSYCL_BEGIN_NAMESPACE_SYCL
/// Collection of properties for SYCL objects. Supported properties are defined
diff --git a/libsycl/include/sycl/__impl/queue.hpp b/libsycl/include/sycl/__impl/queue.hpp
index 32b057104a165..a7921ff678195 100644
--- a/libsycl/include/sycl/__impl/queue.hpp
+++ b/libsycl/include/sycl/__impl/queue.hpp
@@ -18,11 +18,13 @@
#include <sycl/__impl/async_handler.hpp>
#include <sycl/__impl/device.hpp>
#include <sycl/__impl/event.hpp>
+#include <sycl/__impl/handler.hpp>
#include <sycl/__impl/property_list.hpp>
#include <sycl/__impl/detail/config.hpp>
#include <sycl/__impl/detail/get_device_kernel_info.hpp>
#include <sycl/__impl/detail/kernel_arg_helpers.hpp>
+#include <sycl/__impl/detail/kernel_submission.hpp>
#include <sycl/__impl/detail/obj_utils.hpp>
#include <sycl/__impl/detail/unified_range_view.hpp>
#include <sycl/__impl/exception.hpp>
@@ -58,8 +60,40 @@ struct CheckFunctionCallOperator<F, RetT(Args...)> {
};
} // namespace detail
+class TypelessCGF {
+ // SYCL 2020 command group function object is a type which is callable with
+ // operator() that takes a reference to a command group handler, that defines
+ // a command group which can be submitted by a queue. The function object can
+ // be a named type, lambda expression or std::function.
+ template <typename T> struct Invoker {
+ static void call(const void *Object, handler &CGH) {
+ (*const_cast<T *>(static_cast<const T *>(Object)))(CGH);
+ }
+ };
+ const void *Object;
+ using InvokerTy = void (*)(const void *, handler &);
+ const InvokerTy InvokerF;
+
+public:
+ template <class T>
+ TypelessCGF(T &&F)
+ // NOTE: Even if `F` is a pointer to a function, `&F` is a pointer to a
+ // pointer to a function and as such can be casted to `void *` (pointer to
+ // a function cannot be casted).
+ : Object(static_cast<const void *>(&F)),
+ InvokerF(&Invoker<std::remove_reference_t<T>>::call) {}
+ ~TypelessCGF() = default;
+
+ TypelessCGF(const TypelessCGF &) = delete;
+ TypelessCGF(TypelessCGF &&) = delete;
+ TypelessCGF &operator=(const TypelessCGF &) = delete;
+ TypelessCGF &operator=(TypelessCGF &&) = delete;
+
+ void operator()(handler &CGH) const { InvokerF(Object, CGH); }
+};
+
// SYCL 2020 4.6.5. Queue class.
-class _LIBSYCL_EXPORT queue {
+class _LIBSYCL_EXPORT queue : private detail::KernelSubmissionBase<queue> {
public:
queue(const queue &rhs) = default;
queue(queue &&rhs) = default;
@@ -215,7 +249,7 @@ class _LIBSYCL_EXPORT queue {
void()>::value,
"Invalid kernel function signature.");
- setKernelParameters(depEvents);
+ setKernelLaunchParams(depEvents);
using NameT =
typename detail::get_kernel_name_t<KernelName, KernelType>::name;
submitSingleTask<NameT, KernelType>(kernelFunc);
@@ -365,13 +399,7 @@ class _LIBSYCL_EXPORT queue {
template <typename KernelName = detail::AutoName, int Dims, typename... Rest>
event parallel_for(nd_range<Dims> executionRange,
const std::vector<event> &depEvents, Rest &&...rest) {
- if (executionRange.get_global_range() != range<Dims>{} &&
- (executionRange.get_local_range().size() == 0 ||
- executionRange.get_global_range() % executionRange.get_local_range() !=
- range<Dims>{}))
- throw sycl::exception(sycl::make_error_code(sycl::errc::nd_range),
- "Invalid nd_range submission: global size must be "
- "evenly divisible by local size.");
+ detail::checkNDRangeAndThrow(executionRange);
return parallelForImpl<KernelName>(executionRange, depEvents,
std::forward<Rest>(rest)...);
}
@@ -413,109 +441,55 @@ class _LIBSYCL_EXPORT queue {
event memcpy(void *dest, const void *src, std::size_t numBytes,
const std::vector<event> &depEvents);
+ /// Immediately calls the command group function object.
+ ///
+ /// The command group may submit no more than one command to this queue for
+ /// execution on the associated device.
+ ///
+ /// \param cgf command group function object.
+ /// \return an event that represents the status of the submitted command.
+ template <typename T>
+ std::enable_if_t<std::is_invocable_r_v<void, T, handler &>, event>
+ submit(T cgf) {
+ return submitWithHandler(cgf);
+ }
+
+ /// Immediately calls the command group function object.
+ ///
+ /// The command group may submit no more than one command to this queue for
+ /// execution on the associated device. On a kernel error, this command group
+ /// may be scheduled for execution on \p secondaryQueue.
+ ///
+ /// \param cgf command group function object.
+ /// \param secondaryQueue queue used as a fallback for kernel errors. Unused
+ /// (See SYCL 2020 3.9.10. Fallback mechanism).
+ /// \return an event that represents the status of the submitted command.
+ template <typename T>
+ std::enable_if_t<std::is_invocable_r_v<void, T, handler &>, event>
+ submit(T cgf, [[maybe_unused]] queue &secondaryQueue) {
+ return submitWithHandler(cgf);
+ }
+
private:
template <typename KernelName, int Dims, template <int> class Range,
typename... Rest>
event parallelForImpl(Range<Dims> numWorkItems,
const std::vector<event> &depEvents, Rest &&...rest) {
- if constexpr (sizeof...(Rest) != 1)
- throw sycl::exception(errc::feature_not_supported,
- "Reductions are not supported");
- setKernelParameters(depEvents, numWorkItems);
-
- using KernelType =
- std::decay_t<detail::nth_type_t<sizeof...(Rest) - 1, Rest...>>;
- constexpr bool IsNdRangeSubmission =
- std::is_same_v<Range<Dims>, nd_range<Dims>>;
- using SuggestedArgType =
- std::conditional_t<IsNdRangeSubmission, nd_item<Dims>, item<Dims>>;
- using LambdaArgType =
- sycl::detail::lambda_arg_type<KernelType, SuggestedArgType>;
-
- if constexpr (IsNdRangeSubmission) {
- static_assert(
- std::is_convertible_v<sycl::nd_item<Dims>, LambdaArgType>,
- "Kernel argument of a sycl::parallel_for with sycl::nd_range "
- "must be sycl::nd_item or be convertible from sycl::nd_item");
- } else {
- static_assert(
- std::is_convertible_v<sycl::item<Dims>, LambdaArgType> ||
- std::is_convertible_v<sycl::item<Dims, false>, LambdaArgType>,
- "Kernel argument of a sycl::parallel_for with sycl::range "
- "must be sycl::item or be convertible from sycl::item");
- }
+ setKernelLaunchParams(depEvents, numWorkItems);
- using TransformedLambdaArgType = std::conditional_t<
- IsNdRangeSubmission, nd_item<Dims>,
- std::conditional_t<
- std::is_convertible_v<sycl::item<Dims>, LambdaArgType>, item<Dims>,
- std::conditional_t<
- std::is_convertible_v<sycl::item<Dims, false>, LambdaArgType>,
- item<Dims, false>, LambdaArgType>>>;
-
- using NameT =
- typename detail::get_kernel_name_t<KernelName, KernelType>::name;
- submitParallelFor<NameT, TransformedLambdaArgType, KernelType>(rest...);
+ detail::KernelSubmissionBase<queue>::template parallelForImpl<KernelName>(
+ numWorkItems, std::forward<Rest>(rest)...);
return getLastEvent();
}
- /// Name of this function is defined by compiler. It generates a call to this
- /// function in the host implementation of KernelFunc in submitSingleTask or
- /// submitParallelFor.
- /// \param KernelName the name of the kernel being invoked.
- /// \param args the kernel arguments for the kernel invocation.
- template <typename KN, typename... Args>
- void sycl_kernel_launch(const char *KernelName, Args &&...args) {
- static_assert(
- sizeof...(args) == 1,
- "sycl_kernel_launch expects only 2 arguments now: name of kernel and "
- "callable object passed to kernel invocation by the user.");
-
- auto FirstArg = std::get<0>(std::tie(args...));
- submitKernelImpl(detail::getDeviceKernelInfo<KN>(KernelName), &FirstArg,
- sizeof(FirstArg));
- }
-
- /// The sycl_kernel_entry_point attribute facilitates the generation of an
- /// offload kernel entry point function with parameters corresponding to the
- /// (potentially decomposed) kernel arguments and a body that executes the
- /// kernel (after reconstructing the arguments if required).
-#ifdef SYCL_LANGUAGE_VERSION
-# define _LIBSYCL_ENTRY_POINT_ATTR__(KernelName) \
- [[clang::sycl_kernel_entry_point(KernelName)]]
-#else
-# define _LIBSYCL_ENTRY_POINT_ATTR__(KernelName)
-#endif // SYCL_LANGUAGE_VERSION
-
- /// Specifies the parameters and body of the generated offload kernel entry
- /// point for single_task invocations. On host, the compiler generates a call
- /// to sycl_kernel_launch instead of the KernelFunc invocation.
- template <typename KernelName, typename KernelType>
- _LIBSYCL_ENTRY_POINT_ATTR__(KernelName)
- void submitSingleTask(const KernelType &KernelFunc) {
- KernelFunc();
- }
-
- /// Specifies the parameters and body of the generated offload kernel entry
- /// point for parallel_for invocations. On host, the compiler generates a call
- /// to sycl_kernel_launch instead of the KernelFunc invocation.
- template <typename KernelName, typename ElementType, typename KernelType>
- _LIBSYCL_ENTRY_POINT_ATTR__(KernelName)
- void submitParallelFor(const KernelType &KernelFunc) {
- KernelFunc(detail::Builder::getElement(detail::declptr<ElementType>()));
- }
-#undef _LIBSYCL_ENTRY_POINT_ATTR__
-
- /// Passes kernel parameters to the runtime.
+ /// Passes kernel dependencies and execution range to the runtime.
/// \param Events a collection of events representing dependencies of the
/// kernel to submit.
/// \param Range a unified view of the kernel execution range.
- void setKernelParameters(const std::vector<event> &Events,
- const detail::UnifiedRangeView &Range = {});
+ void setKernelLaunchParams(const std::vector<event> &Events,
+ const detail::UnifiedRangeView &Range = {});
/// Passes kernel arguments to runtime.
- /// If all the dependencies can be handled by the backend, the kernel is
- /// submitted to it directly in this call.
/// \param KernelInfo the information for the kernel being invoked.
/// \param ArgData a pointer to the kernel argument.
/// \param ArgSize the size of the kernel argument.
@@ -525,11 +499,14 @@ class _LIBSYCL_EXPORT queue {
/// \return an event representing last kernel invocation.
event getLastEvent();
+ event submitWithHandler(const TypelessCGF &CGF);
+
queue(const std::shared_ptr<detail::QueueImpl> &Impl) : impl(Impl) {}
std::shared_ptr<detail::QueueImpl> impl;
friend sycl::detail::ImplUtils;
friend sycl::detail::MockQueue;
+ friend class detail::KernelSubmissionBase<queue>;
}; // class queue
_LIBSYCL_END_NAMESPACE_SYCL
diff --git a/libsycl/src/CMakeLists.txt b/libsycl/src/CMakeLists.txt
index c591a162db2f8..870fc06b6334e 100644
--- a/libsycl/src/CMakeLists.txt
+++ b/libsycl/src/CMakeLists.txt
@@ -90,6 +90,7 @@ set(LIBSYCL_SOURCES
"exception_list.cpp"
"device.cpp"
"device_selector.cpp"
+ "handler.cpp"
"platform.cpp"
"queue.cpp"
"usm_functions.cpp"
diff --git a/libsycl/src/detail/handler_impl.hpp b/libsycl/src/detail/handler_impl.hpp
new file mode 100644
index 0000000000000..a6a8dbaa737b0
--- /dev/null
+++ b/libsycl/src/detail/handler_impl.hpp
@@ -0,0 +1,62 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// This file contains the declaration of the HandlerImpl, which stores
+/// command group function, argument data, and kernel execution range in
+/// liboffload style for deferred submission.
+///
+//===----------------------------------------------------------------------===//
+
+#ifndef _LIBSYCL_HANDLER_IMPL
+#define _LIBSYCL_HANDLER_IMPL
+
+#include <sycl/__impl/detail/config.hpp>
+
+#include <OffloadAPI.h>
+
+#include <functional>
+#include <memory>
+#include <vector>
+
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+
+namespace detail {
+
+class EventImpl;
+class QueueImpl;
+
+/// Stores the deferred command group state for a sycl::handler submission.
+struct HandlerImpl {
+ HandlerImpl(QueueImpl &Queue) : MQueue(Queue) {}
+
+ HandlerImpl(const HandlerImpl &) = delete;
+ HandlerImpl(HandlerImpl &&) = delete;
+ HandlerImpl &operator=(const HandlerImpl &) = delete;
+ HandlerImpl &operator=(HandlerImpl &&) = delete;
+
+ ~HandlerImpl() = default;
+
+ // Queue this handler is attached to.
+ QueueImpl &MQueue;
+
+ /// The command group function to execute at finalize time.
+ std::function<std::shared_ptr<EventImpl>()> MCGF;
+
+ /// Captured kernel argument data.
+ std::vector<char> MArgData;
+
+ /// Kernel execution range in liboffload format, set by setKernelRange().
+ ol_kernel_launch_size_args_t MRange = {};
+};
+
+} // namespace detail
+
+_LIBSYCL_END_NAMESPACE_SYCL
+
+#endif // _LIBSYCL_HANDLER_IMPL
diff --git a/libsycl/src/detail/offload/offload_utils.cpp b/libsycl/src/detail/offload/offload_utils.cpp
index 594c41f9e965d..1376c11bb955d 100644
--- a/libsycl/src/detail/offload/offload_utils.cpp
+++ b/libsycl/src/detail/offload/offload_utils.cpp
@@ -8,6 +8,9 @@
#include <detail/offload/offload_utils.hpp>
+#include <cassert>
+#include <limits>
+
_LIBSYCL_BEGIN_NAMESPACE_SYCL
namespace detail {
@@ -104,5 +107,49 @@ ol_alloc_type_t getOlAllocType(usm::alloc USMKind) {
}
}
+ol_kernel_launch_size_args_t convertToOlRange(const UnifiedRangeView &Range) {
+ assert(Range.MDims < 4 && "Invalid dimensions.");
+
+ uint32_t GlobalSize[3] = {1, 1, 1};
+ if (Range.MGlobalSize) {
+ for (size_t I = 0; I < Range.MDims; ++I) {
+ assert(Range.MGlobalSize[I] <= std::numeric_limits<uint32_t>::max());
+ GlobalSize[I] = static_cast<uint32_t>(Range.MGlobalSize[I]);
+ }
+ }
+
+ uint32_t GroupSize[3] = {1, 1, 1};
+ if (Range.MLocalSize) {
+ for (size_t I = 0; I < Range.MDims; ++I) {
+ assert(Range.MLocalSize[I] <= std::numeric_limits<uint32_t>::max() &&
+ Range.MLocalSize[I] != 0);
+ GroupSize[I] = static_cast<uint32_t>(Range.MLocalSize[I]);
+ }
+ }
+
+ // We have the following mapping between dimensions with SPIR-V builtins:
+ // 1D: id[0] -> x
+ // 2D: id[0] -> y, id[1] -> x
+ // 3D: id[0] -> z, id[1] -> y, id[2] -> x
+ // So in order to ensure the correctness we update all the kernel
+ // parameters accordingly.
+ if (Range.MDims > 1) {
+ // TODO: Offset is not supported in liboffload so just ignore it for now.
+ std::swap(GlobalSize[0], GlobalSize[Range.MDims - 1]);
+ std::swap(GroupSize[0], GroupSize[Range.MDims - 1]);
+ }
+
+ ol_kernel_launch_size_args_t olRange = {};
+ olRange.Dimensions = Range.MDims;
+ olRange.NumGroups.x = GlobalSize[0] / GroupSize[0];
+ olRange.NumGroups.y = GlobalSize[1] / GroupSize[1];
+ olRange.NumGroups.z = GlobalSize[2] / GroupSize[2];
+ olRange.GroupSize.x = GroupSize[0];
+ olRange.GroupSize.y = GroupSize[1];
+ olRange.GroupSize.z = GroupSize[2];
+ olRange.DynSharedMemory = 0;
+ return olRange;
+}
+
} // namespace detail
_LIBSYCL_END_NAMESPACE_SYCL
diff --git a/libsycl/src/detail/offload/offload_utils.hpp b/libsycl/src/detail/offload/offload_utils.hpp
index 12ddb13d350e4..f565ca86aef4d 100644
--- a/libsycl/src/detail/offload/offload_utils.hpp
+++ b/libsycl/src/detail/offload/offload_utils.hpp
@@ -17,6 +17,7 @@
#include <sycl/__impl/backend.hpp>
#include <sycl/__impl/detail/config.hpp>
+#include <sycl/__impl/detail/unified_range_view.hpp>
#include <sycl/__impl/exception.hpp>
#include <sycl/__impl/info/device_type.hpp>
#include <sycl/__impl/usm_alloc_type.hpp>
@@ -139,6 +140,10 @@ constexpr To map_info_desc(typename info_ol_mapping<To>::template M<Ts>... ms) {
.value;
}
+/// Converts a UnifiedRangeView into the liboffload
+/// ol_kernel_launch_size_args_t format.
+ol_kernel_launch_size_args_t convertToOlRange(const UnifiedRangeView &Range);
+
} // namespace detail
_LIBSYCL_END_NAMESPACE_SYCL
diff --git a/libsycl/src/detail/queue_impl.cpp b/libsycl/src/detail/queue_impl.cpp
index bd76f2aaa2a6e..7fd0b0ef08eec 100644
--- a/libsycl/src/detail/queue_impl.cpp
+++ b/libsycl/src/detail/queue_impl.cpp
@@ -12,6 +12,7 @@
#include <detail/device_impl.hpp>
#include <detail/event_impl.hpp>
#include <detail/global_objects.hpp>
+#include <detail/handler_impl.hpp>
#include <detail/program_manager.hpp>
#include <algorithm>
@@ -20,47 +21,24 @@ _LIBSYCL_BEGIN_NAMESPACE_SYCL
namespace detail {
-static void setKernelLaunchArgs(const detail::UnifiedRangeView &Range,
- ol_kernel_launch_size_args_t &ArgsToSet) {
- assert(Range.MDims < 4 && "Invalid dimensions.");
- uint32_t GlobalSize[3] = {1, 1, 1};
- if (Range.MGlobalSize) {
- for (size_t I = 0; I < Range.MDims; ++I) {
- assert(Range.MGlobalSize[I] <= std::numeric_limits<uint32_t>::max());
- GlobalSize[I] = static_cast<uint32_t>(Range.MGlobalSize[I]);
- }
+thread_local bool NestedCallsDetector = false;
+class NestedCallsTracker {
+public:
+ NestedCallsTracker() {
+ if (NestedCallsDetectorRef)
+ throw sycl::exception(
+ make_error_code(errc::invalid),
+ "Calls to sycl::queue::submit cannot be nested. Command group "
+ "function objects should use the sycl::handler API instead.");
+ NestedCallsDetectorRef = true;
}
- uint32_t GroupSize[3] = {1, 1, 1};
- if (Range.MLocalSize) {
- for (size_t I = 0; I < Range.MDims; ++I) {
- assert(Range.MLocalSize[I] <= std::numeric_limits<uint32_t>::max() &&
- Range.MLocalSize[I] != 0);
- GroupSize[I] = static_cast<uint32_t>(Range.MLocalSize[I]);
- }
- }
+ ~NestedCallsTracker() { NestedCallsDetectorRef = false; }
- // We have the following mapping between dimensions with SPIR-V builtins:
- // 1D: id[0] -> x
- // 2D: id[0] -> y, id[1] -> x
- // 3D: id[0] -> z, id[1] -> y, id[2] -> x
- // So in order to ensure the correctness we update all the kernel
- // parameters accordingly.
- if (Range.MDims > 1) {
- // TODO: Offset is not supported in liboffload so just ignore it for now.
- std::swap(GlobalSize[0], GlobalSize[Range.MDims - 1]);
- std::swap(GroupSize[0], GroupSize[Range.MDims - 1]);
- }
-
- ArgsToSet.Dimensions = Range.MDims;
- ArgsToSet.NumGroups.x = GlobalSize[0] / GroupSize[0];
- ArgsToSet.NumGroups.y = GlobalSize[1] / GroupSize[1];
- ArgsToSet.NumGroups.z = GlobalSize[2] / GroupSize[2];
- ArgsToSet.GroupSize.x = GroupSize[0];
- ArgsToSet.GroupSize.y = GroupSize[1];
- ArgsToSet.GroupSize.z = GroupSize[2];
- ArgsToSet.DynSharedMemory = 0;
-}
+private:
+ // Cache the TLS location to decrease amount of TLS accesses.
+ bool &NestedCallsDetectorRef = NestedCallsDetector;
+};
QueueImpl::QueueImpl(DeviceImpl &deviceImpl, const async_handler &asyncHandler,
const property_list &propList, PrivateTag)
@@ -113,17 +91,19 @@ static void checkEventsPlatformMatch(const std::vector<EventImplPtr> &Events,
}
}
-void QueueImpl::setKernelParameters(std::vector<EventImplPtr> &&Events,
- const detail::UnifiedRangeView &Range) {
+void QueueImpl::setKernelLaunchParams(std::vector<EventImplPtr> &&Events,
+ const detail::UnifiedRangeView &Range) {
checkEventsPlatformMatch(Events, MDevice.getPlatformImpl());
+ MCurrentSubmitInfo.DepEvents = std::move(Events);
+ MCurrentSubmitInfo.Range = convertToOlRange(Range);
+}
- // It is done at the beginning of a new submission to ensure that we can still
- // submit a kernel properly if the previous submission throws.
- MCurrentSubmitInfo.DepEvents.clear();
- MCurrentSubmitInfo.Range = {};
-
+void QueueImpl::setKernelLaunchParams(
+ std::vector<EventImplPtr> &&Events,
+ const ol_kernel_launch_size_args_t &Range) {
+ checkEventsPlatformMatch(Events, MDevice.getPlatformImpl());
MCurrentSubmitInfo.DepEvents = std::move(Events);
- setKernelLaunchArgs(Range, MCurrentSubmitInfo.Range);
+ MCurrentSubmitInfo.Range = Range;
}
void QueueImpl::submitKernelImpl(DeviceKernelInfo &KernelInfo, void *ArgData,
@@ -143,6 +123,7 @@ void QueueImpl::submitKernelImpl(DeviceKernelInfo &KernelInfo, void *ArgData,
auto Result =
olLaunchKernel(MOffloadQueue, MDevice.getOLHandle(), Kernel,
&MCurrentSubmitInfo.Range, NULL, 1, ArgPtrs, ArgSizes);
+
if (isFailed(Result))
throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
std::string("Kernel submission (") +
@@ -177,8 +158,7 @@ QueueImpl::memcpy(void *Dest, const void *Src, std::size_t NumBytes,
const std::vector<EventImplPtr> &DepEvents) {
checkEventsPlatformMatch(DepEvents, MDevice.getPlatformImpl());
if (NumBytes == 0) {
- handleEventDependencies(DepEvents);
- return createEvent();
+ return submitWait(DepEvents);
}
if (!Dest || !Src) {
@@ -224,5 +204,21 @@ EventImplPtr QueueImpl::createEvent(std::vector<EventImplPtr> &&Deps) {
return EventImpl::createEventWithHandle(NewEvent, MDevice.getPlatformImpl(),
std::move(Deps));
}
+
+EventImplPtr QueueImpl::submitWithHandler(const TypelessCGF &CGF) {
+ detail::HandlerImpl HandlerImplVal(*this);
+ handler Handler(HandlerImplVal);
+ {
+ NestedCallsTracker tracker;
+ CGF(Handler);
+ }
+
+ return Handler.finalize();
+}
+
+EventImplPtr QueueImpl::submitWait(const std::vector<EventImplPtr> &DepEvents) {
+ handleEventDependencies(DepEvents);
+ return createEvent(std::vector<EventImplPtr>(DepEvents));
+}
} // namespace detail
_LIBSYCL_END_NAMESPACE_SYCL
diff --git a/libsycl/src/detail/queue_impl.hpp b/libsycl/src/detail/queue_impl.hpp
index e062546de5c5f..c04476454fd31 100644
--- a/libsycl/src/detail/queue_impl.hpp
+++ b/libsycl/src/detail/queue_impl.hpp
@@ -80,8 +80,8 @@ class QueueImpl : public std::enable_shared_from_this<QueueImpl> {
void throwAsynchronous();
/// Enqueues a kernel to liboffload.
- /// Kernel parameters like dependencies and range must be passed in advance by
- /// calling setKernelParameters.
+ /// Kernel dependencies and range must be passed in advance by calling
+ /// setKernelLaunchParams.
/// \param KernelInfo a kernel info that is uniform between different
/// submissions of the same kernel.
/// \param ArgData a pointer to kernel argument.
@@ -97,12 +97,19 @@ class QueueImpl : public std::enable_shared_from_this<QueueImpl> {
return MCurrentSubmitInfo.LastEvent;
}
- /// Sets kernel parameters to be used in the next submitKernelImpl call.
- /// Must be called prior to a submitKernelImpl call.
+ /// Sets event dependencies and execution range for the next kernel
+ /// submission.
/// \param Events a collection of events that the kernel depends on.
/// \param Range a unified range view of the execution range.
- void setKernelParameters(std::vector<EventImplPtr> &&Events,
- const detail::UnifiedRangeView &Range);
+ void setKernelLaunchParams(std::vector<EventImplPtr> &&Events,
+ const detail::UnifiedRangeView &Range);
+
+ /// Sets event dependencies and execution range for the next kernel
+ /// submission.
+ /// \param Events a collection of events that the kernel depends on.
+ /// \param Range a pre-converted liboffload kernel launch size args struct.
+ void setKernelLaunchParams(std::vector<EventImplPtr> &&Events,
+ const ol_kernel_launch_size_args_t &Range);
/// \return the async_handler associated with this queue.
const async_handler &getAsyncHandler() const { return MAsyncHandler; }
@@ -117,6 +124,19 @@ class QueueImpl : public std::enable_shared_from_this<QueueImpl> {
EventImplPtr memcpy(void *Dest, const void *Src, std::size_t NumBytes,
const std::vector<EventImplPtr> &DepEvents);
+ /// Submits a command group function to this queue.
+ ///
+ /// \param CGF is the command group function to invoke.
+ /// \return an event impl object representing the submitted operation.
+ EventImplPtr submitWithHandler(const TypelessCGF &CGF);
+
+ /// Submits a dependency-only wait operation to this queue.
+ ///
+ /// \param DepEvents are the events that must complete before the wait
+ /// operation completes.
+ /// \return an event impl object representing the wait operation.
+ EventImplPtr submitWait(const std::vector<EventImplPtr> &DepEvents);
+
private:
void handleEventDependencies(const std::vector<EventImplPtr> &Dep);
EventImplPtr createEvent(std::vector<EventImplPtr> &&Deps = {});
@@ -131,6 +151,8 @@ class QueueImpl : public std::enable_shared_from_this<QueueImpl> {
// Submit data.
struct KernelSubmitInfo {
+ KernelSubmitInfo() {}
+
EventImplPtr LastEvent;
ol_kernel_launch_size_args_t Range;
std::vector<EventImplPtr> DepEvents;
diff --git a/libsycl/src/handler.cpp b/libsycl/src/handler.cpp
new file mode 100644
index 0000000000000..aa59834123374
--- /dev/null
+++ b/libsycl/src/handler.cpp
@@ -0,0 +1,61 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <detail/handler_impl.hpp>
+#include <detail/offload/offload_utils.hpp>
+#include <detail/queue_impl.hpp>
+#include <sycl/__impl/handler.hpp>
+
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+
+static void checkSingleCommand(
+ const std::function<std::shared_ptr<detail::EventImpl>()> &CGF) {
+ if (CGF) {
+ throw sycl::exception(
+ sycl::make_error_code(sycl::errc::invalid),
+ "Attempt to set multiple actions for the command group");
+ }
+}
+
+void handler::submitKernelImpl(detail::DeviceKernelInfo &KernelInfo,
+ void *ArgData, size_t ArgSize) {
+ checkSingleCommand(MImpl.MCGF);
+ MImpl.MArgData.resize(ArgSize);
+ std::memcpy(MImpl.MArgData.data(), ArgData, ArgSize);
+ MImpl.MCGF = [this, &KernelInfo]() {
+ auto EventsImpl = detail::getSyclObjImpls(MDepEvents);
+ MImpl.MQueue.setKernelLaunchParams(std::move(EventsImpl), MImpl.MRange);
+ MImpl.MQueue.submitKernelImpl(KernelInfo, MImpl.MArgData.data(),
+ MImpl.MArgData.size());
+ return MImpl.MQueue.getLastEvent();
+ };
+}
+
+void handler::setKernelRange(const detail::UnifiedRangeView &Range) {
+ MImpl.MRange = convertToOlRange(Range);
+}
+
+void handler::memcpy(void *dest, const void *src, std::size_t numBytes) {
+ checkSingleCommand(MImpl.MCGF);
+ MImpl.MCGF = [this, dest, src, numBytes]() {
+ return MImpl.MQueue.memcpy(dest, src, numBytes,
+ detail::getSyclObjImpls(MDepEvents));
+ };
+}
+
+std::shared_ptr<detail::EventImpl> handler::finalize() {
+ if (!MImpl.MCGF) {
+ auto EventsImpl = detail::getSyclObjImpls(MDepEvents);
+ return MImpl.MQueue.submitWait(EventsImpl);
+ }
+
+ auto Event = MImpl.MCGF();
+ return Event;
+}
+
+_LIBSYCL_END_NAMESPACE_SYCL
diff --git a/libsycl/src/queue.cpp b/libsycl/src/queue.cpp
index 6d15537636936..9117063cb11ac 100644
--- a/libsycl/src/queue.cpp
+++ b/libsycl/src/queue.cpp
@@ -51,9 +51,9 @@ event queue::getLastEvent() {
return detail::createSyclObjFromImpl<event>(impl->getLastEvent());
}
-void queue::setKernelParameters(const std::vector<event> &Events,
- const detail::UnifiedRangeView &Range) {
- return impl->setKernelParameters(detail::getSyclObjImpls(Events), Range);
+void queue::setKernelLaunchParams(const std::vector<event> &Events,
+ const detail::UnifiedRangeView &Range) {
+ return impl->setKernelLaunchParams(detail::getSyclObjImpls(Events), Range);
}
void queue::submitKernelImpl(detail::DeviceKernelInfo &KernelInfo,
@@ -61,4 +61,8 @@ void queue::submitKernelImpl(detail::DeviceKernelInfo &KernelInfo,
impl->submitKernelImpl(KernelInfo, ArgData, ArgSize);
}
+event queue::submitWithHandler(const TypelessCGF &CGF) {
+ return detail::createSyclObjFromImpl<event>(impl->submitWithHandler(CGF));
+}
+
_LIBSYCL_END_NAMESPACE_SYCL
diff --git a/libsycl/test/basic/handler/handler_depends_on.cpp b/libsycl/test/basic/handler/handler_depends_on.cpp
new file mode 100644
index 0000000000000..1869d6973701f
--- /dev/null
+++ b/libsycl/test/basic/handler/handler_depends_on.cpp
@@ -0,0 +1,82 @@
+// REQUIRES: any-device
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %t.out
+
+#include <sycl/sycl.hpp>
+
+#include <algorithm>
+#include <cassert>
+#include <numeric>
+
+namespace {
+
+void testMemcpyDependency() {
+ constexpr size_t N = 4;
+ sycl::queue Queue;
+
+ int *Src = sycl::malloc_shared<int>(N, Queue);
+ int *Dst = sycl::malloc_shared<int>(N, Queue);
+ std::iota(Src, Src + N, 1);
+ std::fill(Dst, Dst + N, 0);
+
+ auto MemCpyEvent = Queue.submit(
+ [&](sycl::handler &CGH) { CGH.memcpy(Dst, Src, N * sizeof(int)); });
+
+ auto *DstElementCopy = sycl::malloc_shared<int>(1, Queue);
+ *DstElementCopy = 0;
+
+ auto KernelEvent = Queue.submit([&](sycl::handler &CGH) {
+ CGH.depends_on(MemCpyEvent);
+ CGH.single_task<class DependsOnMemcpyKernel>(
+ [=]() { *DstElementCopy = Dst[0]; });
+ });
+
+ KernelEvent.wait();
+
+ assert(Dst[0] == 1 && Dst[1] == 2 && Dst[2] == 3 && Dst[3] == 4);
+ assert(*DstElementCopy == 1);
+
+ sycl::free(Src, Queue);
+ sycl::free(Dst, Queue);
+ sycl::free(DstElementCopy, Queue);
+}
+
+void testNDRangeDependency() {
+ constexpr size_t N = 16;
+ sycl::queue Queue;
+
+ int *Data = sycl::malloc_shared<int>(N, Queue);
+ int *Token = sycl::malloc_shared<int>(1, Queue);
+
+ std::fill(Data, Data + N, 0);
+ *Token = 0;
+
+ auto InitEvent =
+ Queue.single_task<class DependsOnNDRangeInit>([=]() { *Token = 9; });
+
+ auto KernelEvent = Queue.submit([&](sycl::handler &CGH) {
+ CGH.depends_on(InitEvent);
+ CGH.parallel_for<class DependsOnNDRangeKernel>(
+ sycl::nd_range<1>{sycl::range<1>{N}, sycl::range<1>{4}},
+ [=](sycl::nd_item<1> Item) {
+ const size_t I = Item.get_global_id(0);
+ Data[I] = static_cast<int>(I) + *Token;
+ });
+ });
+
+ KernelEvent.wait();
+
+ for (size_t I = 0; I < N; ++I)
+ assert(Data[I] == static_cast<int>(I) + 9);
+
+ sycl::free(Data, Queue);
+ sycl::free(Token, Queue);
+}
+
+} // namespace
+
+int main() {
+ testMemcpyDependency();
+ testNDRangeDependency();
+ return 0;
+}
\ No newline at end of file
diff --git a/libsycl/test/basic/handler/handler_parallel_for_arg_restrictions.cpp b/libsycl/test/basic/handler/handler_parallel_for_arg_restrictions.cpp
new file mode 100644
index 0000000000000..b33087e1c5b9a
--- /dev/null
+++ b/libsycl/test/basic/handler/handler_parallel_for_arg_restrictions.cpp
@@ -0,0 +1,50 @@
+// RUN: %clangxx -fsycl -fsyntax-only %s
+// expected-no-diagnostics
+
+#include <sycl/sycl.hpp>
+
+template <int Dims> struct ConvertibleFromNDItem {
+ ConvertibleFromNDItem(sycl::nd_item<Dims>) {}
+};
+
+int main() {
+ sycl::queue Q;
+
+ Q.submit([&](sycl::handler &CGH) {
+ CGH.parallel_for(sycl::range{1}, [=](sycl::item<1>) {});
+ });
+
+ Q.submit([&](sycl::handler &CGH) {
+ CGH.parallel_for(sycl::range{1, 1}, [=](sycl::item<2>) {});
+ });
+
+ Q.submit([&](sycl::handler &CGH) {
+ CGH.parallel_for(sycl::range{1, 1, 1}, [=](sycl::item<3>) {});
+ });
+
+ Q.submit([&](sycl::handler &CGH) {
+ CGH.parallel_for(sycl::range{1}, [=](sycl::id<1>) {});
+ });
+
+ Q.submit([&](sycl::handler &CGH) {
+ CGH.parallel_for(sycl::nd_range{sycl::range{4}, sycl::range{2}},
+ [=](sycl::nd_item<1>) {});
+ });
+
+ Q.submit([&](sycl::handler &CGH) {
+ CGH.parallel_for(sycl::nd_range{sycl::range{2, 4}, sycl::range{1, 2}},
+ [=](sycl::nd_item<2>) {});
+ });
+
+ Q.submit([&](sycl::handler &CGH) {
+ CGH.parallel_for(sycl::nd_range{sycl::range{2, 2, 4}, sycl::range{1, 1, 2}},
+ [=](sycl::nd_item<3>) {});
+ });
+
+ Q.submit([&](sycl::handler &CGH) {
+ CGH.parallel_for(sycl::nd_range{sycl::range{4}, sycl::range{2}},
+ [=](ConvertibleFromNDItem<1>) {});
+ });
+
+ return 0;
+}
diff --git a/libsycl/test/basic/handler/handler_parallel_for_generic_lambda.cpp b/libsycl/test/basic/handler/handler_parallel_for_generic_lambda.cpp
new file mode 100644
index 0000000000000..da34fab3b6741
--- /dev/null
+++ b/libsycl/test/basic/handler/handler_parallel_for_generic_lambda.cpp
@@ -0,0 +1,43 @@
+// RUN: %clangxx -fsycl -fsyntax-only %s
+// expected-no-diagnostics
+
+#include <sycl/sycl.hpp>
+
+#include <type_traits>
+
+template <typename KernelName, typename ExpectedType, typename Range>
+void test_parallel_for(Range r) {
+ sycl::queue q;
+ q.submit([&](sycl::handler &cgh) {
+ cgh.parallel_for<KernelName>(r, [=](auto item) {
+ static_assert(std::is_same<decltype(item), ExpectedType>::value,
+ "Argument type is unexpected");
+ });
+ });
+}
+
+int main() {
+ test_parallel_for<class Item1Name, sycl::item<1>>(sycl::range<1>{1});
+ test_parallel_for<class Item2Name, sycl::item<2>>(sycl::range<2>{1, 1});
+ test_parallel_for<class Item3Name, sycl::item<3>>(sycl::range<3>{1, 1, 1});
+ test_parallel_for<class NDItem1Name, sycl::nd_item<1>>(
+ sycl::nd_range<1>{sycl::range<1>{1}, sycl::range<1>{1}});
+ test_parallel_for<class NDItem2Name, sycl::nd_item<2>>(
+ sycl::nd_range<2>{sycl::range<2>{2, 2}, sycl::range<2>{1, 1}});
+ test_parallel_for<class NDItem3Name, sycl::nd_item<3>>(
+ sycl::nd_range<3>{sycl::range<3>{2, 2, 2}, sycl::range<3>{1, 1, 1}});
+
+ sycl::queue q;
+ q.submit([&](sycl::handler &cgh) {
+ cgh.parallel_for<class GenericInitList1>(sycl::range{1}, [=](auto &) {});
+ });
+ q.submit([&](sycl::handler &cgh) {
+ cgh.parallel_for<class GenericInitList2>(sycl::range{1, 1}, [=](auto &) {});
+ });
+ q.submit([&](sycl::handler &cgh) {
+ cgh.parallel_for<class GenericInitList3>(sycl::range{1, 1, 1},
+ [=](auto &) {});
+ });
+
+ return 0;
+}
diff --git a/libsycl/test/basic/handler/handler_parallel_for_nd_range_invalid_arg_type.cpp b/libsycl/test/basic/handler/handler_parallel_for_nd_range_invalid_arg_type.cpp
new file mode 100644
index 0000000000000..6a810a21fdf94
--- /dev/null
+++ b/libsycl/test/basic/handler/handler_parallel_for_nd_range_invalid_arg_type.cpp
@@ -0,0 +1,17 @@
+// RUN: not %clangxx -fsycl -fsyntax-only %s 2>&1 | FileCheck %s
+
+#include <sycl/sycl.hpp>
+
+int main() {
+ sycl::queue q;
+
+ q.submit([&](sycl::handler &cgh) {
+ cgh.parallel_for<class HandlerNDRangeInvalidArgType>(
+ sycl::nd_range<1>{sycl::range<1>{4}, sycl::range<1>{2}},
+ [=](sycl::item<1>) {});
+ });
+
+ return 0;
+}
+
+// CHECK: must be sycl::nd_item or be convertible from sycl::nd_item
diff --git a/libsycl/test/basic/handler/handler_parallel_for_runtime.cpp b/libsycl/test/basic/handler/handler_parallel_for_runtime.cpp
new file mode 100644
index 0000000000000..8f6b31c332d9b
--- /dev/null
+++ b/libsycl/test/basic/handler/handler_parallel_for_runtime.cpp
@@ -0,0 +1,149 @@
+// REQUIRES: any-device
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %t.out
+
+#include <sycl/sycl.hpp>
+
+#include <cassert>
+
+void test1D(sycl::queue &q) {
+ constexpr size_t N = 16;
+ constexpr size_t LocalSize = 4;
+ int *data = sycl::malloc_shared<int>(N, q);
+ assert(data);
+
+ for (size_t i = 0; i < N; ++i)
+ data[i] = 0;
+
+ q.submit([&](sycl::handler &cgh) {
+ cgh.parallel_for<class HandlerParallelForRuntime>(
+ sycl::range<1>{N},
+ [=](sycl::item<1> it) { data[it[0]] = static_cast<int>(it[0]) + 7; });
+ }).wait();
+
+ for (size_t i = 0; i < N; ++i)
+ assert(data[i] == static_cast<int>(i) + 7);
+
+ for (size_t i = 0; i < N; ++i)
+ data[i] = 0;
+
+ q.submit([&](sycl::handler &cgh) {
+ cgh.parallel_for<class HandlerParallelForNDRangeRuntime>(
+ sycl::nd_range<1>{sycl::range<1>{N}, sycl::range<1>{LocalSize}},
+ [=](sycl::nd_item<1> it) {
+ const size_t idx = it.get_global_id(0);
+ data[idx] = static_cast<int>(idx) + 11;
+ });
+ }).wait();
+
+ for (size_t i = 0; i < N; ++i)
+ assert(data[i] == static_cast<int>(i) + 11);
+
+ sycl::free(data, q);
+}
+
+void test2D(sycl::queue &q) {
+ constexpr size_t G0 = 4;
+ constexpr size_t G1 = 6;
+ constexpr size_t L0 = 2;
+ constexpr size_t L1 = 3;
+ int *data = sycl::malloc_shared<int>(G0 * G1, q);
+ assert(data);
+
+ for (size_t i = 0; i < G0 * G1; ++i)
+ data[i] = -1;
+
+ q.submit([&](sycl::handler &cgh) {
+ cgh.parallel_for<class HandlerParallelFor2DRuntime>(
+ sycl::range<2>{G0, G1}, [=](sycl::item<2> it) {
+ const size_t i = it.get_id(0);
+ const size_t j = it.get_id(1);
+ data[i * G1 + j] = static_cast<int>(i * 100 + j) + 7;
+ });
+ }).wait();
+
+ for (size_t i = 0; i < G0; ++i)
+ for (size_t j = 0; j < G1; ++j)
+ assert(data[i * G1 + j] == static_cast<int>(i * 100 + j) + 7);
+
+ for (size_t i = 0; i < G0 * G1; ++i)
+ data[i] = -1;
+
+ q.submit([&](sycl::handler &cgh) {
+ cgh.parallel_for<class HandlerNDRange2DRuntime>(
+ sycl::nd_range<2>{sycl::range<2>{G0, G1}, sycl::range<2>{L0, L1}},
+ [=](sycl::nd_item<2> it) {
+ const size_t i = it.get_global_id(0);
+ const size_t j = it.get_global_id(1);
+ data[i * G1 + j] = static_cast<int>(i * 100 + j);
+ });
+ }).wait();
+
+ for (size_t i = 0; i < G0; ++i)
+ for (size_t j = 0; j < G1; ++j)
+ assert(data[i * G1 + j] == static_cast<int>(i * 100 + j));
+
+ sycl::free(data, q);
+}
+
+void test3D(sycl::queue &q) {
+ constexpr size_t G0 = 2;
+ constexpr size_t G1 = 3;
+ constexpr size_t G2 = 4;
+ constexpr size_t L0 = 1;
+ constexpr size_t L1 = 3;
+ constexpr size_t L2 = 2;
+ int *data = sycl::malloc_shared<int>(G0 * G1 * G2, q);
+ assert(data);
+
+ for (size_t i = 0; i < G0 * G1 * G2; ++i)
+ data[i] = -1;
+
+ q.submit([&](sycl::handler &cgh) {
+ cgh.parallel_for<class HandlerParallelFor3DRuntime>(
+ sycl::range<3>{G0, G1, G2}, [=](sycl::item<3> it) {
+ const size_t i = it.get_id(0);
+ const size_t j = it.get_id(1);
+ const size_t k = it.get_id(2);
+ data[(i * G1 + j) * G2 + k] =
+ static_cast<int>(i * 100 + j * 10 + k) + 7;
+ });
+ }).wait();
+
+ for (size_t i = 0; i < G0; ++i)
+ for (size_t j = 0; j < G1; ++j)
+ for (size_t k = 0; k < G2; ++k)
+ assert(data[(i * G1 + j) * G2 + k] ==
+ static_cast<int>(i * 100 + j * 10 + k) + 7);
+
+ for (size_t i = 0; i < G0 * G1 * G2; ++i)
+ data[i] = -1;
+
+ q.submit([&](sycl::handler &cgh) {
+ cgh.parallel_for<class HandlerNDRange3DRuntime>(
+ sycl::nd_range<3>{sycl::range<3>{G0, G1, G2},
+ sycl::range<3>{L0, L1, L2}},
+ [=](sycl::nd_item<3> it) {
+ const size_t i = it.get_global_id(0);
+ const size_t j = it.get_global_id(1);
+ const size_t k = it.get_global_id(2);
+ data[(i * G1 + j) * G2 + k] = static_cast<int>(i * 100 + j * 10 + k);
+ });
+ }).wait();
+
+ for (size_t i = 0; i < G0; ++i)
+ for (size_t j = 0; j < G1; ++j)
+ for (size_t k = 0; k < G2; ++k)
+ assert(data[(i * G1 + j) * G2 + k] ==
+ static_cast<int>(i * 100 + j * 10 + k));
+
+ sycl::free(data, q);
+}
+
+int main() {
+ sycl::queue q;
+ test1D(q);
+ test2D(q);
+ test3D(q);
+ return 0;
+}
diff --git a/libsycl/test/basic/handler/handler_unnamed_lambda_functor.cpp b/libsycl/test/basic/handler/handler_unnamed_lambda_functor.cpp
new file mode 100644
index 0000000000000..ae9c0790f9a99
--- /dev/null
+++ b/libsycl/test/basic/handler/handler_unnamed_lambda_functor.cpp
@@ -0,0 +1,19 @@
+// RUN: %clangxx -fsycl -fsycl-device-only -std=c++17 -fsyntax-only %s
+
+#include <sycl/sycl.hpp>
+
+struct SingleTaskKernel {
+ void operator()() const {}
+};
+
+int main() {
+ sycl::queue q;
+
+ q.single_task<class QueueSingleTaskNamed>(SingleTaskKernel{});
+
+ q.submit([&](sycl::handler &cgh) {
+ cgh.single_task<class HandlerSingleTaskNamed>(SingleTaskKernel{});
+ });
+
+ return 0;
+}
diff --git a/libsycl/test/basic/handler/submit_fn_ptr_handler.cpp b/libsycl/test/basic/handler/submit_fn_ptr_handler.cpp
new file mode 100644
index 0000000000000..c43a5deef4830
--- /dev/null
+++ b/libsycl/test/basic/handler/submit_fn_ptr_handler.cpp
@@ -0,0 +1,24 @@
+// REQUIRES: any-device
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %t.out
+
+#include <sycl/sycl.hpp>
+
+#include <cassert>
+
+int *p = nullptr;
+
+void foo(sycl::handler &cgh) {
+ auto *copy = p;
+ cgh.single_task([=]() { *copy = 42; });
+}
+
+int main() {
+ sycl::queue q;
+ p = sycl::malloc_shared<int>(1, q);
+ *p = 0;
+ q.submit(foo).wait();
+ assert(*p == 42);
+ sycl::free(p, q);
+ return 0;
+}
diff --git a/libsycl/unittests/CMakeLists.txt b/libsycl/unittests/CMakeLists.txt
index 597651abd0bfc..a54229b9e3b66 100644
--- a/libsycl/unittests/CMakeLists.txt
+++ b/libsycl/unittests/CMakeLists.txt
@@ -7,6 +7,7 @@ set_target_properties(LibsyclUnitTests PROPERTIES FOLDER "libsycl/Tests/Unit")
add_subdirectory(mock)
add_subdirectory(device_selector)
add_subdirectory(event)
+add_subdirectory(handler)
add_subdirectory(platform)
add_subdirectory(program_manager)
add_subdirectory(queue)
diff --git a/libsycl/unittests/event/event.cpp b/libsycl/unittests/event/event.cpp
index 327c25415fc95..d8543663e9f6f 100644
--- a/libsycl/unittests/event/event.cpp
+++ b/libsycl/unittests/event/event.cpp
@@ -42,7 +42,7 @@ createEventImplWithHandle(detail::PlatformImpl &PlatformImpl,
class sycl::detail::MockQueue : public sycl::queue {
public:
using sycl::queue::getLastEvent;
- using sycl::queue::setKernelParameters;
+ using sycl::queue::setKernelLaunchParams;
using sycl::queue::sycl_kernel_launch;
};
@@ -121,7 +121,7 @@ TEST(Event, GetWaitListWithSetKernelParametersAndLaunch) {
std::vector<event> DepEvents = {Dep1, Dep2};
- Q.setKernelParameters(DepEvents);
+ Q.setKernelLaunchParams(DepEvents, {});
Q.sycl_kernel_launch<class TestKernelWithDeps>(TestKernelWithDeps, Data);
event KernelEvent = Q.getLastEvent();
diff --git a/libsycl/unittests/handler/CMakeLists.txt b/libsycl/unittests/handler/CMakeLists.txt
new file mode 100644
index 0000000000000..ad2b633ecd28b
--- /dev/null
+++ b/libsycl/unittests/handler/CMakeLists.txt
@@ -0,0 +1,4 @@
+add_sycl_unittest(HandlerTests
+ memcpy.cpp
+ semantics.cpp
+)
diff --git a/libsycl/unittests/handler/memcpy.cpp b/libsycl/unittests/handler/memcpy.cpp
new file mode 100644
index 0000000000000..24f04376ff79c
--- /dev/null
+++ b/libsycl/unittests/handler/memcpy.cpp
@@ -0,0 +1,70 @@
+#include "test_helpers.hpp"
+#include <mock/helpers.hpp>
+
+#include <detail/device_impl.hpp>
+
+#include <sycl/__impl/device.hpp>
+#include <sycl/__impl/queue.hpp>
+
+#include <gmock/gmock.h>
+#include <gtest/gtest.h>
+
+using namespace sycl;
+using namespace ::testing;
+
+TEST(Handler, MemcpyViaSubmit) {
+ constexpr int NumBytes = 64;
+
+ mock::MockWrapper Mock;
+ queue Q;
+
+ int Src[NumBytes / sizeof(int)] = {};
+ int Dst[NumBytes / sizeof(int)] = {};
+
+ ol_device_handle_t OLDev =
+ detail::getSyclObjImpl(Q.get_device())->getOLHandle();
+
+ sycl::unittests::expectDeviceMemoryInfo(Mock, {Src, Dst}, OLDev, 2);
+
+ EXPECT_CALL(Mock.get(), olMemcpy(_, Dst, OLDev, Src, OLDev, NumBytes))
+ .Times(1);
+ EXPECT_CALL(Mock.get(), olCreateEvent(_, _, _)).Times(1);
+
+ auto E = Q.submit([&](handler &CGH) { CGH.memcpy(Dst, Src, NumBytes); });
+
+ EXPECT_CALL(Mock.get(), olSyncEvent(_)).Times(1);
+ E.wait();
+}
+
+TEST(Handler, DependsOnWithMemcpy) {
+ constexpr int NumBytes = 32;
+
+ mock::MockWrapper Mock;
+ queue Q;
+
+ int SrcA[NumBytes / sizeof(int)] = {};
+ int Mid[NumBytes / sizeof(int)] = {};
+ int Dst[NumBytes / sizeof(int)] = {};
+
+ ol_device_handle_t OLDev =
+ detail::getSyclObjImpl(Q.get_device())->getOLHandle();
+
+ sycl::unittests::expectDeviceMemoryInfo(Mock, {SrcA, Mid, Dst}, OLDev, 4);
+
+ EXPECT_CALL(Mock.get(), olMemcpy(_, Mid, OLDev, SrcA, OLDev, NumBytes))
+ .Times(1);
+ EXPECT_CALL(Mock.get(), olMemcpy(_, Dst, OLDev, Mid, OLDev, NumBytes))
+ .Times(1);
+ EXPECT_CALL(Mock.get(), olCreateEvent(_, _, _)).Times(2);
+
+ auto First = Q.memcpy(Mid, SrcA, NumBytes);
+
+ EXPECT_CALL(Mock.get(), olWaitEvents(_, _, 1)).Times(1);
+ auto Second = Q.submit([&](handler &CGH) {
+ CGH.depends_on(First);
+ CGH.memcpy(Dst, Mid, NumBytes);
+ });
+
+ EXPECT_CALL(Mock.get(), olSyncEvent(_)).Times(1);
+ Second.wait();
+}
diff --git a/libsycl/unittests/handler/semantics.cpp b/libsycl/unittests/handler/semantics.cpp
new file mode 100644
index 0000000000000..3bd5d2bba4360
--- /dev/null
+++ b/libsycl/unittests/handler/semantics.cpp
@@ -0,0 +1,164 @@
+#include "test_helpers.hpp"
+#include <mock/helpers.hpp>
+
+#include <detail/device_impl.hpp>
+
+#include <sycl/__impl/device.hpp>
+#include <sycl/__impl/queue.hpp>
+
+#include <gmock/gmock.h>
+#include <gtest/gtest.h>
+
+#include <string>
+
+using namespace sycl;
+using namespace ::testing;
+
+TEST(Handler, MultipleActionsRejected) {
+ mock::MockWrapper Mock;
+ queue Q;
+ int Src = 1;
+ int Dst = 0;
+
+ bool Thrown = false;
+ try {
+ Q.submit([&](handler &CGH) {
+ CGH.memcpy(&Dst, &Src, sizeof(int));
+ CGH.memcpy(&Dst, &Src, sizeof(int));
+ });
+ } catch (const sycl::exception &E) {
+ Thrown = true;
+ EXPECT_NE(std::string(E.what()).find("multiple actions"),
+ std::string::npos);
+ }
+
+ EXPECT_TRUE(Thrown);
+}
+
+TEST(Handler, DependsOnOnlyCommandGroup) {
+ mock::MockWrapper Mock;
+ queue Q;
+
+ constexpr int NumBytes = sizeof(int);
+ int Src = 1;
+ int Dst = 0;
+ ol_device_handle_t OLDev =
+ detail::getSyclObjImpl(Q.get_device())->getOLHandle();
+
+ sycl::unittests::expectDeviceMemoryInfo(Mock, {&Src, &Dst}, OLDev, 2);
+ EXPECT_CALL(Mock.get(), olMemcpy(_, &Dst, OLDev, &Src, OLDev, NumBytes))
+ .Times(1);
+ EXPECT_CALL(Mock.get(), olCreateEvent(_, _, _)).Times(2);
+
+ auto DepEvent = Q.memcpy(&Dst, &Src, NumBytes);
+
+ EXPECT_CALL(Mock.get(), olWaitEvents(_, _, 1)).Times(1);
+ auto E = Q.submit([&](handler &CGH) { CGH.depends_on(DepEvent); });
+
+ EXPECT_CALL(Mock.get(), olSyncEvent(_)).Times(1);
+ E.wait();
+}
+
+TEST(Handler, DependsOnVectorOverload) {
+ constexpr int NumBytes = 16;
+
+ mock::MockWrapper Mock;
+ queue Q;
+
+ int SrcA[NumBytes / sizeof(int)] = {};
+ int Mid[NumBytes / sizeof(int)] = {};
+ int Dst[NumBytes / sizeof(int)] = {};
+
+ ol_device_handle_t OLDev =
+ detail::getSyclObjImpl(Q.get_device())->getOLHandle();
+
+ sycl::unittests::expectDeviceMemoryInfo(Mock, {SrcA, Mid, Dst}, OLDev, 4);
+
+ EXPECT_CALL(Mock.get(), olMemcpy(_, Mid, OLDev, SrcA, OLDev, NumBytes))
+ .Times(1);
+ EXPECT_CALL(Mock.get(), olMemcpy(_, Dst, OLDev, Mid, OLDev, NumBytes))
+ .Times(1);
+
+ // Two events for the two memcpys, one for the dependency-only command group.
+ EXPECT_CALL(Mock.get(), olCreateEvent(_, _, _)).Times(3);
+
+ auto E1 = Q.memcpy(Mid, SrcA, NumBytes);
+ auto E2 = Q.memcpy(Dst, Mid, NumBytes);
+
+ EXPECT_CALL(Mock.get(), olWaitEvents(_, _, 2)).Times(1);
+ auto E = Q.submit(
+ [&](handler &CGH) { CGH.depends_on(std::vector<event>{E1, E2}); });
+
+ EXPECT_CALL(Mock.get(), olSyncEvent(_)).Times(1);
+ E.wait();
+}
+
+TEST(Handler, EmptyCommandGroupNoDependencies) {
+ mock::MockWrapper Mock;
+ queue Q;
+
+ EXPECT_CALL(Mock.get(), olWaitEvents(_, _, _)).Times(0);
+ EXPECT_CALL(Mock.get(), olCreateEvent(_, _, _)).Times(1);
+
+ auto E = Q.submit([&](handler &CGH) { (void)CGH; });
+
+ EXPECT_CALL(Mock.get(), olSyncEvent(_)).Times(1);
+ E.wait();
+}
+
+TEST(Queue, SubmitCannotBeNested) {
+ mock::MockWrapper Mock;
+ queue Q;
+
+ bool Thrown = false;
+ try {
+ Q.submit([&](handler &CGH) {
+ (void)CGH;
+ Q.submit([&](handler &Nested) {
+ Nested.single_task<class UTNestedSubmitKernel>([]() {});
+ });
+ });
+ } catch (const sycl::exception &E) {
+ Thrown = true;
+ EXPECT_NE(std::string(E.what()).find("cannot be nested"),
+ std::string::npos);
+ }
+
+ EXPECT_TRUE(Thrown);
+}
+
+TEST(Handler, ParallelForNDRangeRejectsNonDivisibleRange) {
+ mock::MockWrapper Mock;
+ queue Q;
+
+ bool Thrown = false;
+ try {
+ Q.submit([&](handler &CGH) {
+ CGH.parallel_for<class UTNDRangeInvalidNonDivisible>(
+ nd_range<1>{range<1>{10}, range<1>{3}}, [=](nd_item<1>) {});
+ });
+ } catch (const sycl::exception &E) {
+ Thrown = true;
+ EXPECT_EQ(E.code(), sycl::errc::nd_range);
+ }
+
+ EXPECT_TRUE(Thrown);
+}
+
+TEST(Handler, ParallelForNDRangeRejectsZeroLocalRange) {
+ mock::MockWrapper Mock;
+ queue Q;
+
+ bool Thrown = false;
+ try {
+ Q.submit([&](handler &CGH) {
+ CGH.parallel_for<class UTNDRangeInvalidZeroLocal>(
+ nd_range<1>{range<1>{8}, range<1>{0}}, [=](nd_item<1>) {});
+ });
+ } catch (const sycl::exception &E) {
+ Thrown = true;
+ EXPECT_EQ(E.code(), sycl::errc::nd_range);
+ }
+
+ EXPECT_TRUE(Thrown);
+}
diff --git a/libsycl/unittests/handler/test_helpers.hpp b/libsycl/unittests/handler/test_helpers.hpp
new file mode 100644
index 0000000000000..ad43374df79cd
--- /dev/null
+++ b/libsycl/unittests/handler/test_helpers.hpp
@@ -0,0 +1,37 @@
+#ifndef LIBSYCL_UNITTESTS_HANDLER_TEST_HELPERS_HPP
+#define LIBSYCL_UNITTESTS_HANDLER_TEST_HELPERS_HPP
+
+#include <mock/helpers.hpp>
+
+#include <sycl/__impl/detail/config.hpp>
+
+#include <gmock/gmock.h>
+
+#include <algorithm>
+#include <initializer_list>
+#include <vector>
+
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+namespace unittests {
+
+inline void expectDeviceMemoryInfo(mock::MockWrapper &Mock,
+ const std::vector<const void *> ExpectedPtrs,
+ ol_device_handle_t Device, int Count) {
+ EXPECT_CALL(Mock.get(),
+ olGetMemInfo(::testing::_, OL_MEM_INFO_DEVICE,
+ sizeof(ol_device_handle_t), ::testing::_))
+ .Times(Count)
+ .WillRepeatedly([ExpectedPtrs, Device](const void *Ptr, ol_mem_info_t,
+ size_t,
+ void *PropValue) -> ol_result_t {
+ EXPECT_NE(std::find(ExpectedPtrs.begin(), ExpectedPtrs.end(), Ptr),
+ ExpectedPtrs.end());
+ *(static_cast<ol_device_handle_t *>(PropValue)) = Device;
+ return OL_SUCCESS;
+ });
+}
+
+} // namespace unittests
+_LIBSYCL_END_NAMESPACE_SYCL
+
+#endif
diff --git a/libsycl/unittests/queue/sycl_kernel_launch.cpp b/libsycl/unittests/queue/sycl_kernel_launch.cpp
index 686f0860309d8..fd2f1602ed880 100644
--- a/libsycl/unittests/queue/sycl_kernel_launch.cpp
+++ b/libsycl/unittests/queue/sycl_kernel_launch.cpp
@@ -21,7 +21,7 @@ using namespace ::testing;
class sycl::detail::MockQueue : public sycl::queue {
public:
- using sycl::queue::setKernelParameters;
+ using sycl::queue::setKernelLaunchParams;
using sycl::queue::sycl_kernel_launch;
};
@@ -97,7 +97,7 @@ static ol_kernel_launch_size_args_t captureKernelLaunchArgs(
struct DimSwapParam {
// Name used by the test runner to identify the case.
const char *Description;
- // Calls setKernelParameters on Q with the appropriate range.
+ // Calls kernel dependency/range setup on Q with the appropriate range.
std::function<void(sycl::detail::MockQueue &)> SetParams;
// Expected fields of ol_kernel_launch_size_args_t after the swap.
uint32_t ExpDims;
@@ -132,7 +132,7 @@ INSTANTIATE_TEST_SUITE_P(
[](sycl::detail::MockQueue &Q) {
sycl::nd_range<1> NDR(sycl::range<1>{8},
sycl::range<1>{2});
- Q.setKernelParameters({}, NDR);
+ Q.setKernelLaunchParams(std::vector<sycl::event>{}, NDR);
},
/*Dims=*/1, /*NG=*/4, 1, 1, /*GS=*/2, 1, 1},
// 2D nd_range: swap [0]<->[1].
@@ -141,7 +141,7 @@ INSTANTIATE_TEST_SUITE_P(
[](sycl::detail::MockQueue &Q) {
sycl::nd_range<2> NDR(sycl::range<2>{4, 6},
sycl::range<2>{2, 3});
- Q.setKernelParameters({}, NDR);
+ Q.setKernelLaunchParams(std::vector<sycl::event>{}, NDR);
},
/*Dims=*/2, /*NG=*/2, 2, 1, /*GS=*/3, 2, 1},
// 3D nd_range: swap [0]<->[2].
@@ -151,7 +151,7 @@ INSTANTIATE_TEST_SUITE_P(
[](sycl::detail::MockQueue &Q) {
sycl::nd_range<3> NDR(sycl::range<3>{2, 4, 6},
sycl::range<3>{1, 2, 3});
- Q.setKernelParameters({}, NDR);
+ Q.setKernelLaunchParams(std::vector<sycl::event>{}, NDR);
},
/*Dims=*/3, /*NG=*/2, 2, 2, /*GS=*/3, 2, 1},
// 2D range (no local): swap [0]<->[1], GroupSize stays {1,1,1}.
@@ -159,7 +159,7 @@ INSTANTIATE_TEST_SUITE_P(
DimSwapParam{"2D_Range",
[](sycl::detail::MockQueue &Q) {
sycl::range<2> R(4, 6);
- Q.setKernelParameters({}, R);
+ Q.setKernelLaunchParams(std::vector<sycl::event>{}, R);
},
/*Dims=*/2, /*NG=*/6, 4, 1, /*GS=*/1, 1, 1}),
[](const ::testing::TestParamInfo<DimSwapParam> &Info) {
>From 4b1cbf396dd2f5aa21d2fd3d38d63458b8dc34f6 Mon Sep 17 00:00:00 2001
From: "Tikhomirova, Kseniya" <kseniya.tikhomirova at intel.com>
Date: Tue, 25 Aug 2026 05:03:00 -0700
Subject: [PATCH 2/2] fix comments
Signed-off-by: Tikhomirova, Kseniya <kseniya.tikhomirova at intel.com>
---
libsycl/include/sycl/__impl/queue.hpp | 31 ++++++++++---------
libsycl/src/detail/queue_impl.cpp | 4 +--
.../test/basic/handler/handler_depends_on.cpp | 6 +---
3 files changed, 18 insertions(+), 23 deletions(-)
diff --git a/libsycl/include/sycl/__impl/queue.hpp b/libsycl/include/sycl/__impl/queue.hpp
index 3c473bebde694..e55cf3f1cf8cb 100644
--- a/libsycl/include/sycl/__impl/queue.hpp
+++ b/libsycl/include/sycl/__impl/queue.hpp
@@ -61,25 +61,12 @@ struct CheckFunctionCallOperator<F, RetT(Args...)> {
} // namespace detail
class TypelessCGF {
- // SYCL 2020 command group function object is a type which is callable with
- // operator() that takes a reference to a command group handler, that defines
- // a command group which can be submitted by a queue. The function object can
- // be a named type, lambda expression or std::function.
- template <typename T> struct Invoker {
- static void call(const void *Object, handler &CGH) {
- (*const_cast<T *>(static_cast<const T *>(Object)))(CGH);
- }
- };
- const void *Object;
- using InvokerTy = void (*)(const void *, handler &);
- const InvokerTy InvokerF;
-
public:
template <class T>
TypelessCGF(T &&F)
// NOTE: Even if `F` is a pointer to a function, `&F` is a pointer to a
- // pointer to a function and as such can be casted to `void *` (pointer to
- // a function cannot be casted).
+ // pointer to a function and as such can be cast to `void *` (pointer to
+ // a function cannot be cast).
: Object(static_cast<const void *>(&F)),
InvokerF(&Invoker<std::remove_reference_t<T>>::call) {}
~TypelessCGF() = default;
@@ -90,6 +77,20 @@ class TypelessCGF {
TypelessCGF &operator=(TypelessCGF &&) = delete;
void operator()(handler &CGH) const { InvokerF(Object, CGH); }
+
+private:
+ // SYCL 2020 command group function object is a type that is callable with
+ // operator(), takes a reference to a command group handler, and defines
+ // a command group that can be submitted by a queue. The function object can
+ // be a named type, lambda expression or std::function.
+ template <typename T> struct Invoker {
+ static void call(const void *Object, handler &CGH) {
+ (*const_cast<T *>(static_cast<const T *>(Object)))(CGH);
+ }
+ };
+ const void *Object;
+ using InvokerTy = void (*)(const void *, handler &);
+ const InvokerTy InvokerF;
};
// SYCL 2020 4.6.5. Queue class.
diff --git a/libsycl/src/detail/queue_impl.cpp b/libsycl/src/detail/queue_impl.cpp
index 9e60c7d863cdc..a156b4c8eee23 100644
--- a/libsycl/src/detail/queue_impl.cpp
+++ b/libsycl/src/detail/queue_impl.cpp
@@ -93,9 +93,7 @@ static void checkEventsPlatformMatch(const std::vector<EventImplPtr> &Events,
void QueueImpl::setKernelLaunchParams(std::vector<EventImplPtr> &&Events,
const detail::UnifiedRangeView &Range) {
- checkEventsPlatformMatch(Events, MDevice.getPlatformImpl());
- MCurrentSubmitInfo.DepEvents = std::move(Events);
- MCurrentSubmitInfo.Range = convertToOlRange(Range);
+ setKernelLaunchParams(std::move(Events), convertToOlRange(Range));
}
void QueueImpl::setKernelLaunchParams(
diff --git a/libsycl/test/basic/handler/handler_depends_on.cpp b/libsycl/test/basic/handler/handler_depends_on.cpp
index 1869d6973701f..3fc8958d30581 100644
--- a/libsycl/test/basic/handler/handler_depends_on.cpp
+++ b/libsycl/test/basic/handler/handler_depends_on.cpp
@@ -8,8 +8,6 @@
#include <cassert>
#include <numeric>
-namespace {
-
void testMemcpyDependency() {
constexpr size_t N = 4;
sycl::queue Queue;
@@ -73,10 +71,8 @@ void testNDRangeDependency() {
sycl::free(Token, Queue);
}
-} // namespace
-
int main() {
testMemcpyDependency();
testNDRangeDependency();
return 0;
-}
\ No newline at end of file
+}
More information about the llvm-commits
mailing list