[llvm] draft (PR #211025)
Kseniya Tikhomirova via llvm-commits
llvm-commits at lists.llvm.org
Tue Jul 21 08:51:32 PDT 2026
https://github.com/KseniyaTikhomirova created https://github.com/llvm/llvm-project/pull/211025
None
>From bccecb61c330819b3dfef2a5cd0e43d08c3dbabf 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] draft
Signed-off-by: Tikhomirova, Kseniya <kseniya.tikhomirova at intel.com>
---
.../sycl/__impl/detail/kernel_submission.hpp | 119 +++++++++++++++
libsycl/include/sycl/__impl/handler.hpp | 122 +++++++++++++++
libsycl/include/sycl/__impl/queue.hpp | 139 ++++++++++--------
libsycl/src/CMakeLists.txt | 1 +
libsycl/src/detail/queue_impl.cpp | 37 +++++
libsycl/src/detail/queue_impl.hpp | 5 +
libsycl/src/handler.cpp | 33 +++++
libsycl/src/queue.cpp | 4 +
.../test/basic/handler_depends_on_memcpy.cpp | 49 ++++++
.../handler_parallel_for_arg_restrictions.cpp | 34 +++++
.../handler_parallel_for_generic_lambda.cpp | 41 ++++++
.../basic/handler_unnamed_lambda_functor.cpp | 23 +++
libsycl/test/basic/submit_fn_ptr_handler.cpp | 24 +++
libsycl/unittests/CMakeLists.txt | 1 +
libsycl/unittests/handler/CMakeLists.txt | 3 +
libsycl/unittests/handler/memcpy.cpp | 101 +++++++++++++
16 files changed, 672 insertions(+), 64 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/handler.cpp
create mode 100644 libsycl/test/basic/handler_depends_on_memcpy.cpp
create mode 100644 libsycl/test/basic/handler_parallel_for_arg_restrictions.cpp
create mode 100644 libsycl/test/basic/handler_parallel_for_generic_lambda.cpp
create mode 100644 libsycl/test/basic/handler_unnamed_lambda_functor.cpp
create mode 100644 libsycl/test/basic/submit_fn_ptr_handler.cpp
create mode 100644 libsycl/unittests/handler/CMakeLists.txt
create mode 100644 libsycl/unittests/handler/memcpy.cpp
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..3490eaccc192a
--- /dev/null
+++ b/libsycl/include/sycl/__impl/detail/kernel_submission.hpp
@@ -0,0 +1,119 @@
+//===----------------------------------------------------------------------===//
+//
+// 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 <tuple>
+#include <utility>
+
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+
+namespace detail {
+
+template <typename KernelName, typename SubmitKernelImpl, typename... Args>
+void syclKernelLaunch(SubmitKernelImpl &&SubmitKernelImplFn,
+ const char *KernelNameStr, Args &&...KernelArgs) {
+ static_assert(
+ sizeof...(KernelArgs) == 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(KernelArgs...));
+ SubmitKernelImplFn(detail::getDeviceKernelInfo<KernelName>(KernelNameStr),
+ &FirstArg, sizeof(FirstArg));
+}
+
+template <typename KernelType>
+void submitSingleTask(const KernelType &KernelFunc) {
+ KernelFunc();
+}
+
+template <typename ElementType, typename KernelType>
+void submitParallelFor(const KernelType &KernelFunc) {
+ KernelFunc(detail::Builder::getElement(detail::declptr<ElementType>()));
+}
+
+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) {
+ detail::submitSingleTask(KernelFunc);
+ }
+
+ template <typename KernelName, typename ElementType, typename KernelType>
+#ifdef SYCL_LANGUAGE_VERSION
+ [[clang::sycl_kernel_entry_point(KernelName)]]
+#endif
+ void submitParallelFor(const KernelType &KernelFunc) {
+ detail::submitParallelFor<ElementType>(KernelFunc);
+ }
+
+ 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)->template submitKernelFromLaunch<KN>(
+ KernelName, FirstArg);
+ }
+
+ template <typename KernelName, int Dims, typename... Rest>
+ decltype(auto) 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...>>;
+ using LambdaArgType = sycl::detail::lambda_arg_type<KernelType, item<Dims>>;
+ 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 either sycl::item or be convertible from sycl::item");
+
+ using TranformedLambdaArgType = std::conditional_t<
+ std::is_convertible_v<item<Dims>, LambdaArgType>, item<Dims>,
+ std::conditional_t<
+ std::is_convertible_v<item<Dims, false>, LambdaArgType>,
+ item<Dims, false>, LambdaArgType>>;
+
+ static_cast<DerivedT *>(this)->prepareParallelForRange(numWorkItems);
+
+ using NameT =
+ typename detail::get_kernel_name_t<KernelName, KernelType>::name;
+ submitParallelFor<NameT, TranformedLambdaArgType, KernelType>(
+ std::forward<Rest>(rest)...);
+
+ return static_cast<DerivedT *>(this)->completeParallelForSubmission();
+ }
+};
+
+} // namespace detail
+
+_LIBSYCL_END_NAMESPACE_SYCL
+
+#endif // _LIBSYCL___IMPL_DETAIL_KERNEL_SUBMISSION_HPP
diff --git a/libsycl/include/sycl/__impl/handler.hpp b/libsycl/include/sycl/__impl/handler.hpp
new file mode 100644
index 0000000000000..5686aa26fb429
--- /dev/null
+++ b/libsycl/include/sycl/__impl/handler.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
+//
+//===----------------------------------------------------------------------===//
+///
+/// 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 <vector>
+
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+
+namespace detail {
+class QueueImpl;
+}
+
+class _LIBSYCL_EXPORT handler : private detail::KernelSubmissionBase<handler> {
+public:
+ handler(detail::QueueImpl &Queue) : MQueue(Queue) {}
+
+ handler(const handler &) = delete;
+ handler(handler &&) = delete;
+ handler &operator=(const handler &) = delete;
+ handler &operator=(handler &&) = delete;
+
+ ~handler() = default;
+
+ void depends_on(event depEvent) {
+ return depends_on(std::vector<event>{depEvent});
+ }
+
+ void depends_on(const std::vector<event> &depEvents) {
+ MDepEvents.insert(MDepEvents.end(), depEvents.begin(), depEvents.end());
+ }
+
+ template <typename KernelName = detail::AutoName, typename KernelType>
+ void single_task(const KernelType &kernelFunc) {
+ submitSingleTask<KernelName, KernelType>(kernelFunc);
+ }
+
+ template <typename KernelName = detail::AutoName, typename... Rest>
+ void parallel_for(range<1> numWorkItems, Rest &&...rest) {
+ return parallelForImpl<KernelName>(numWorkItems,
+ std::forward<Rest>(rest)...);
+ }
+
+ template <typename KernelName = detail::AutoName, typename... Rest>
+ void parallel_for(range<2> numWorkItems, Rest &&...rest) {
+ return parallelForImpl<KernelName>(numWorkItems,
+ std::forward<Rest>(rest)...);
+ }
+
+ template <typename KernelName = detail::AutoName, typename... Rest>
+ void parallel_for(range<3> numWorkItems, Rest &&...rest) {
+ return parallelForImpl<KernelName>(numWorkItems,
+ std::forward<Rest>(rest)...);
+ }
+
+ void memcpy(void *dest, const void *src, std::size_t numBytes);
+
+private:
+ std::shared_ptr<detail::EventImpl> finalize();
+
+ void submitKernelImpl(detail::DeviceKernelInfo &KernelInfo, void *ArgData,
+ size_t ArgSize);
+
+ template <typename KN, typename ArgT>
+ void submitKernelFromLaunch(const char *KernelName, ArgT &FirstArg) {
+ constexpr size_t ArgSize = sizeof(FirstArg);
+ MArgData.resize(ArgSize);
+ std::memcpy(MArgData.data(), &FirstArg, ArgSize);
+ submitKernelImpl(detail::getDeviceKernelInfo<KN>(KernelName),
+ MArgData.data(), ArgSize);
+ MArgData.clear();
+ }
+
+ template <int Dims> void prepareParallelForRange(range<Dims> numWorkItems) {
+ MKernelRange = detail::UnifiedRangeView(numWorkItems);
+ }
+
+ void completeParallelForSubmission() {}
+
+ // Queue, this handler is attached to.
+ sycl::detail::QueueImpl &MQueue;
+
+ // Any command submission data.
+ std::vector<event> MDepEvents;
+ std::function<std::shared_ptr<detail::EventImpl>()> MCGF;
+
+ // Kernel specific data to be passed via a few libsycl calls.
+ std::vector<char> MArgData;
+ detail::UnifiedRangeView MKernelRange{};
+
+ 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/queue.hpp b/libsycl/include/sycl/__impl/queue.hpp
index f8fd0d265ecb4..1dda89ceb1bbb 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;
@@ -385,82 +419,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, 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...>>;
- using LambdaArgType = sycl::detail::lambda_arg_type<KernelType, item<Dims>>;
- 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 either sycl::item or be convertible from sycl::item");
- using TranformedLambdaArgType = std::conditional_t<
- std::is_convertible_v<item<Dims>, LambdaArgType>, item<Dims>,
- std::conditional_t<
- std::is_convertible_v<item<Dims, false>, LambdaArgType>,
- item<Dims, false>, LambdaArgType>>;
-
- using NameT =
- typename detail::get_kernel_name_t<KernelName, KernelType>::name;
- submitParallelFor<NameT, TranformedLambdaArgType, KernelType>(rest...);
- return getLastEvent();
+ return detail::KernelSubmissionBase<queue>::template parallelForImpl<
+ KernelName>(numWorkItems, std::forward<Rest>(rest)...);
}
- /// 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.");
+ template <int Dims> void prepareParallelForRange(range<Dims>) {}
- auto FirstArg = std::get<0>(std::tie(args...));
+ event completeParallelForSubmission() { return getLastEvent(); }
+
+ template <typename KN, typename ArgT>
+ void submitKernelFromLaunch(const char *KernelName, ArgT &FirstArg) {
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.
/// \param Events a collection of events representing dependencies of the
/// kernel to submit.
@@ -480,11 +487,15 @@ 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 sycl::handler;
+ 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/queue_impl.cpp b/libsycl/src/detail/queue_impl.cpp
index 0b7cae80ddf96..1ff92281f2a50 100644
--- a/libsycl/src/detail/queue_impl.cpp
+++ b/libsycl/src/detail/queue_impl.cpp
@@ -19,6 +19,25 @@ _LIBSYCL_BEGIN_NAMESPACE_SYCL
namespace detail {
+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;
+ }
+
+ ~NestedCallsTracker() { NestedCallsDetectorRef = false; }
+
+private:
+ // Cache the TLS location to decrease amount of TLS accesses.
+ bool &NestedCallsDetectorRef = NestedCallsDetector;
+};
+
static void setKernelLaunchArgs(const detail::UnifiedRangeView &Range,
ol_kernel_launch_size_args_t &ArgsToSet) {
assert(Range.MDims < 4 && "Invalid dimensions.");
@@ -208,5 +227,23 @@ EventImplPtr QueueImpl::createEvent(std::vector<EventImplPtr> &&Deps) {
return EventImpl::createEventWithHandle(NewEvent, MDevice.getPlatformImpl(),
std::move(Deps));
}
+
+EventImplPtr QueueImpl::submitWithHandler(const TypelessCGF &CGF) {
+ // detail::handler_impl HandlerImplVal(*this);
+ // handler Handler(HandlerImplVal);
+ handler Handler(*this);
+ // struct HandlerSetter {
+ // handler *Handler;
+ // HandlerSetter(handler *H) : Handler(H) { MCurrentSubmitInfo.Handler = H;
+ // } ~HandlerSetter() { MCurrentSubmitInfo.Handler = nullptr; }
+ // } Setter(&Handler);
+
+ {
+ NestedCallsTracker tracker;
+ CGF(Handler);
+ }
+
+ return Handler.finalize();
+}
} // namespace detail
_LIBSYCL_END_NAMESPACE_SYCL
diff --git a/libsycl/src/detail/queue_impl.hpp b/libsycl/src/detail/queue_impl.hpp
index e062546de5c5f..13f573405e391 100644
--- a/libsycl/src/detail/queue_impl.hpp
+++ b/libsycl/src/detail/queue_impl.hpp
@@ -117,6 +117,8 @@ 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);
+ EventImplPtr submitWithHandler(const TypelessCGF &CGF);
+
private:
void handleEventDependencies(const std::vector<EventImplPtr> &Dep);
EventImplPtr createEvent(std::vector<EventImplPtr> &&Deps = {});
@@ -131,9 +133,12 @@ class QueueImpl : public std::enable_shared_from_this<QueueImpl> {
// Submit data.
struct KernelSubmitInfo {
+ KernelSubmitInfo() : Handler(nullptr) {}
+
EventImplPtr LastEvent;
ol_kernel_launch_size_args_t Range;
std::vector<EventImplPtr> DepEvents;
+ handler *Handler;
};
inline static thread_local KernelSubmitInfo MCurrentSubmitInfo = {};
};
diff --git a/libsycl/src/handler.cpp b/libsycl/src/handler.cpp
new file mode 100644
index 0000000000000..8cf02eb060be1
--- /dev/null
+++ b/libsycl/src/handler.cpp
@@ -0,0 +1,33 @@
+//===----------------------------------------------------------------------===//
+//
+// 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/queue_impl.hpp>
+#include <sycl/__impl/handler.hpp>
+
+_LIBSYCL_BEGIN_NAMESPACE_SYCL
+
+void handler::submitKernelImpl(detail::DeviceKernelInfo &KernelInfo,
+ void *ArgData, size_t ArgSize) {
+ MCGF = [this, &KernelInfo, ArgData, ArgSize]() {
+ auto EventsImpl = detail::getSyclObjImpls(MDepEvents);
+ MQueue.setKernelParameters(std::move(EventsImpl), MKernelRange);
+ MQueue.submitKernelImpl(KernelInfo, ArgData, ArgSize);
+ return MQueue.getLastEvent();
+ };
+}
+
+void handler::memcpy(void *dest, const void *src, std::size_t numBytes) {
+ MCGF = [this, dest, src, numBytes]() {
+ return MQueue.memcpy(dest, src, numBytes,
+ detail::getSyclObjImpls(MDepEvents));
+ };
+}
+
+std::shared_ptr<detail::EventImpl> handler::finalize() { return MCGF(); }
+
+_LIBSYCL_END_NAMESPACE_SYCL
diff --git a/libsycl/src/queue.cpp b/libsycl/src/queue.cpp
index 6d15537636936..7823a15b0586e 100644
--- a/libsycl/src/queue.cpp
+++ b/libsycl/src/queue.cpp
@@ -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_depends_on_memcpy.cpp b/libsycl/test/basic/handler_depends_on_memcpy.cpp
new file mode 100644
index 0000000000000..97b918bebbde9
--- /dev/null
+++ b/libsycl/test/basic/handler_depends_on_memcpy.cpp
@@ -0,0 +1,49 @@
+// REQUIRES: any-device
+// RUN: %clangxx -fsycl %s -o %t.out
+// RUN: %t.out
+
+// Adapted from upstream SYCL tests that validate empty command-group
+// dependency behavior and handler memory operations.
+
+#include <sycl/sycl.hpp>
+
+#include <cassert>
+
+int main() {
+ sycl::queue q;
+
+ int *src = sycl::malloc_shared<int>(4, q);
+ int *dst = sycl::malloc_shared<int>(4, q);
+ assert(src && dst);
+ src[0] = 1;
+ src[1] = 2;
+ src[2] = 3;
+ src[3] = 4;
+ dst[0] = 0;
+ dst[1] = 0;
+ dst[2] = 0;
+ dst[3] = 0;
+
+ // Exercise handler::memcpy in a submitted command group.
+ auto memcpy_event = q.submit(
+ [&](sycl::handler &cgh) { cgh.memcpy(dst, src, 4 * sizeof(int)); });
+
+ auto *shared_value = sycl::malloc_shared<int>(1, q);
+ assert(shared_value);
+ *shared_value = 0;
+
+ auto kernel_event = q.submit([&](sycl::handler &cgh) {
+ cgh.depends_on(memcpy_event);
+ cgh.single_task<class DependsOnKernel>([=]() { *shared_value = dst[0]; });
+ });
+
+ kernel_event.wait();
+
+ assert(dst[0] == 1 && dst[1] == 2 && dst[2] == 3 && dst[3] == 4);
+ assert(*shared_value == 1);
+
+ sycl::free(src, q);
+ sycl::free(dst, q);
+ sycl::free(shared_value, q);
+ return 0;
+}
\ No newline at end of file
diff --git a/libsycl/test/basic/handler_parallel_for_arg_restrictions.cpp b/libsycl/test/basic/handler_parallel_for_arg_restrictions.cpp
new file mode 100644
index 0000000000000..e8cdbe2ff8c2d
--- /dev/null
+++ b/libsycl/test/basic/handler_parallel_for_arg_restrictions.cpp
@@ -0,0 +1,34 @@
+// RUN: %clangxx -fsycl -fsyntax-only %s
+// expected-no-diagnostics
+
+// Adapted from upstream SYCL test/basic_tests/handler/
+// parallel_for_arg_restrictions.cpp. This version keeps signatures that
+// should compile for current handler::parallel_for(range) support.
+
+#include <sycl/sycl.hpp>
+
+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::range{1}, [=](int) {});
+ });
+
+ return 0;
+}
diff --git a/libsycl/test/basic/handler_parallel_for_generic_lambda.cpp b/libsycl/test/basic/handler_parallel_for_generic_lambda.cpp
new file mode 100644
index 0000000000000..e6d89214a5dc7
--- /dev/null
+++ b/libsycl/test/basic/handler_parallel_for_generic_lambda.cpp
@@ -0,0 +1,41 @@
+// RUN: %clangxx -fsycl -fsyntax-only %s
+// expected-no-diagnostics
+
+// Adapted from upstream SYCL test/basic_tests/handler/
+// handler_generic_lambda_interface.cpp. This version keeps only
+// handler.parallel_for(range) coverage that is currently supported.
+
+#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});
+
+ 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_unnamed_lambda_functor.cpp b/libsycl/test/basic/handler_unnamed_lambda_functor.cpp
new file mode 100644
index 0000000000000..9ab6c4658ee69
--- /dev/null
+++ b/libsycl/test/basic/handler_unnamed_lambda_functor.cpp
@@ -0,0 +1,23 @@
+// RUN: %clangxx -fsycl -fsycl-device-only -std=c++17 -fsyntax-only %s
+
+// Adapted from upstream SYCL
+// test/basic_tests/handler/unnamed-lambda-functor.cpp. This version keeps only
+// operations supported by current libsycl.
+
+#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/submit_fn_ptr_handler.cpp b/libsycl/test/basic/submit_fn_ptr_handler.cpp
new file mode 100644
index 0000000000000..c43a5deef4830
--- /dev/null
+++ b/libsycl/test/basic/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 21cdbf32570e3..10ef115b56341 100644
--- a/libsycl/unittests/CMakeLists.txt
+++ b/libsycl/unittests/CMakeLists.txt
@@ -7,6 +7,7 @@ add_custom_target(check-sycl-unittests)
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/handler/CMakeLists.txt b/libsycl/unittests/handler/CMakeLists.txt
new file mode 100644
index 0000000000000..86d99f31f0202
--- /dev/null
+++ b/libsycl/unittests/handler/CMakeLists.txt
@@ -0,0 +1,3 @@
+add_sycl_unittest(HandlerTests
+ memcpy.cpp
+)
diff --git a/libsycl/unittests/handler/memcpy.cpp b/libsycl/unittests/handler/memcpy.cpp
new file mode 100644
index 0000000000000..4662bc063eb0e
--- /dev/null
+++ b/libsycl/unittests/handler/memcpy.cpp
@@ -0,0 +1,101 @@
+#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();
+
+ EXPECT_CALL(Mock.get(), olGetMemInfo(Src, OL_MEM_INFO_DEVICE,
+ sizeof(ol_device_handle_t), _))
+ .WillRepeatedly([&](const void *Ptr, ol_mem_info_t PropName,
+ size_t PropSize, void *PropValue) -> ol_result_t {
+ EXPECT_EQ(Ptr, static_cast<const void *>(Src));
+ std::ignore = PropName;
+ std::ignore = PropSize;
+ *(static_cast<ol_device_handle_t *>(PropValue)) = OLDev;
+ return OL_SUCCESS;
+ });
+
+ EXPECT_CALL(Mock.get(), olGetMemInfo(Dst, OL_MEM_INFO_DEVICE,
+ sizeof(ol_device_handle_t), _))
+ .WillRepeatedly([&](const void *Ptr, ol_mem_info_t PropName,
+ size_t PropSize, void *PropValue) -> ol_result_t {
+ EXPECT_EQ(Ptr, static_cast<const void *>(Dst));
+ std::ignore = PropName;
+ std::ignore = PropSize;
+ *(static_cast<ol_device_handle_t *>(PropValue)) = OLDev;
+ return OL_SUCCESS;
+ });
+
+ 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();
+
+ EXPECT_CALL(Mock.get(), olGetMemInfo(_, OL_MEM_INFO_DEVICE,
+ sizeof(ol_device_handle_t), _))
+ .Times(4)
+ .WillRepeatedly([&](const void *Ptr, ol_mem_info_t PropName,
+ size_t PropSize, void *PropValue) -> ol_result_t {
+ EXPECT_TRUE(Ptr == static_cast<const void *>(SrcA) ||
+ Ptr == static_cast<const void *>(Mid) ||
+ Ptr == static_cast<const void *>(Dst));
+ std::ignore = PropName;
+ std::ignore = PropSize;
+ *(static_cast<ol_device_handle_t *>(PropValue)) = OLDev;
+ return OL_SUCCESS;
+ });
+
+ 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();
+}
More information about the llvm-commits
mailing list