[llvm-branch-commits] [llvm] [libsycl] Add parallel_for feature (PR #189068)
Sergey Semenov via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Fri Apr 24 06:25:48 PDT 2026
================
@@ -179,35 +159,191 @@ class _LIBSYCL_EXPORT queue {
/// \param depEvent is an event that specifies the kernel dependency.
/// \param kernelFunc is the kernel functor or lambda.
/// \return an event that represents the status of the submitted kernel.
- template <typename KernelName, typename KernelType>
+ template <typename KernelName = detail::AutoName, typename KernelType>
event single_task(event depEvent, const KernelType &kernelFunc) {
return single_task<KernelName, KernelType>({depEvent}, kernelFunc);
}
/// Defines and invokes a SYCL kernel function as a lambda expression or a
/// named function object type.
///
- /// \param depEvents is a collection of events that specify the kernel
+ /// \param depEvents is a collection of events which specify the kernel
/// dependencies.
/// \param kernelFunc is the kernel functor or lambda.
/// \return an event that represents the status of the submitted kernel.
- template <typename KernelName, typename KernelType>
+ template <typename KernelName = detail::AutoName, typename KernelType>
event single_task(const std::vector<event> &depEvents,
const KernelType &kernelFunc) {
static_assert(
- detail::CheckFunctionSignature<std::remove_reference_t<KernelType>,
- void()>::value,
+ (detail::CheckFunctionSignature<std::remove_reference_t<KernelType>,
+ void()>::value),
"sycl::queue::single_task() requires a kernel instead of a command "
"group");
setKernelParameters(depEvents);
- submitSingleTask<KernelName, KernelType>(kernelFunc);
+ using NameT =
+ typename detail::get_kernel_name_t<KernelName, KernelType>::name;
+ submitSingleTask<NameT, KernelType>(kernelFunc);
return getLastEvent();
}
+ /// Defines and invokes a SYCL kernel function as a lambda expression or a
+ /// named function object type, for the specified range.
+ ///
+ /// \param numWorkItems specifies the global work space of the kernel.
+ /// \param rest acts as-if: const KernelType &KernelFunc".
+ // TODO: Rest will represent reduction types once it is supported.
+ template <typename KernelName = detail::AutoName, typename... Rest>
+ event parallel_for(range<1> numWorkItems, Rest &&...rest) {
+ return parallel_for<KernelName>(numWorkItems, {},
+ std::forward<Rest>(rest)...);
+ }
+
+ /// Defines and invokes a SYCL kernel function as a lambda expression or a
+ /// named function object type, for the specified range.
+ ///
+ /// \param numWorkItems specifies the global work space of the kernel.
+ /// \param rest acts as-if: const KernelType &KernelFunc".
+ // TODO: Rest will represent reduction types once it is supported.
+ template <typename KernelName = detail::AutoName, typename... Rest>
+ event parallel_for(range<2> numWorkItems, Rest &&...rest) {
+ return parallel_for<KernelName>(numWorkItems, {},
+ std::forward<Rest>(rest)...);
+ }
+
+ /// Defines and invokes a SYCL kernel function as a lambda expression or a
+ /// named function object type, for the specified range.
+ ///
+ /// \param numWorkItems specifies the global work space of the kernel.
+ /// \param rest acts as-if: const KernelType &KernelFunc".
+ // TODO: Rest will represent reduction types once it is supported.
+ template <typename KernelName = detail::AutoName, typename... Rest>
+ event parallel_for(range<3> numWorkItems, Rest &&...rest) {
+ return parallel_for<KernelName>(numWorkItems, {},
+ std::forward<Rest>(rest)...);
+ }
+
+ /// Defines and invokes a SYCL kernel function as a lambda expression or a
+ /// named function object type, for the specified range.
+ ///
+ /// \param numWorkItems specifies the global work space of the kernel.
+ /// \param depEvent adds a requirement that the action represented by depEvent
+ /// must complete before executing this kernel.
+ /// \param rest acts as-if: const KernelType &KernelFunc".
+ // TODO: Rest will represent reduction types once it is supported.
+ template <typename KernelName = detail::AutoName, typename... Rest>
+ event parallel_for(range<1> numWorkItems, event depEvent, Rest &&...rest) {
+ return parallel_for<KernelName>(numWorkItems, {depEvent},
+ std::forward<Rest>(rest)...);
+ }
+
+ /// Defines and invokes a SYCL kernel function as a lambda expression or a
+ /// named function object type, for the specified range.
+ ///
+ /// \param numWorkItems specifies the global work space of the kernel.
+ /// \param depEvent adds a requirement that the action represented by depEvent
+ /// must complete before executing this kernel.
+ /// \param rest acts as-if: const KernelType &KernelFunc".
+ // TODO: Rest will represent reduction types once it is supported.
+ template <typename KernelName = detail::AutoName, typename... Rest>
+ event parallel_for(range<2> numWorkItems, event depEvent, Rest &&...rest) {
+ return parallel_for<KernelName>(numWorkItems, {depEvent},
+ std::forward<Rest>(rest)...);
+ }
+
+ /// Defines and invokes a SYCL kernel function as a lambda expression or a
+ /// named function object type, for the specified range.
+ ///
+ /// \param numWorkItems specifies the global work space of the kernel.
+ /// \param depEvent adds a requirement that the action represented by depEvent
+ /// must complete before executing this kernel.
+ /// \param rest acts as-if: const KernelType &KernelFunc".
+ // TODO: Rest will represent reduction types once it is supported.
+ template <typename KernelName = detail::AutoName, typename... Rest>
+ event parallel_for(range<3> numWorkItems, event depEvent, Rest &&...rest) {
+ return parallel_for<KernelName>(numWorkItems, {depEvent},
+ std::forward<Rest>(rest)...);
+ }
+
+ /// Defines and invokes a SYCL kernel function as a lambda expression or a
+ /// named function object type, for the specified range.
+ ///
+ /// \param numWorkItems specifies the global work space of the kernel
+ /// \param depEvents is a vector of events that specifies the kernel
+ /// dependencies.
+ /// \param rest acts as-if: const KernelType &KernelFunc".
+ // TODO: Rest will represent reduction types once it is supported.
+ template <typename KernelName = detail::AutoName, typename... Rest>
+ event parallel_for(range<1> numWorkItems, const std::vector<event> &depEvents,
+ Rest &&...rest) {
+ return parallelForImpl<KernelName>(numWorkItems, depEvents,
+ std::forward<Rest>(rest)...);
+ }
+
+ /// Defines and invokes a SYCL kernel function as a lambda expression or a
+ /// named function object type, for the specified range.
+ ///
+ /// \param numWorkItems specifies the global work space of the kernel
+ /// \param depEvents is a vector of events that specifies the kernel
+ /// dependencies.
+ /// \param rest acts as-if: const KernelType &KernelFunc".
+ // TODO: Rest will represent reduction types once it is supported.
+ template <typename KernelName = detail::AutoName, typename... Rest>
+ event parallel_for(range<2> numWorkItems, const std::vector<event> &depEvents,
+ Rest &&...rest) {
+ return parallelForImpl<KernelName>(numWorkItems, depEvents,
+ std::forward<Rest>(rest)...);
+ }
+
+ /// Defines and invokes a SYCL kernel function as a lambda expression or a
+ /// named function object type, for the specified range.
+ ///
+ /// \param numWorkItems specifies the global work space of the kernel
+ /// \param depEvents is a vector of events that specifies the kernel
+ /// dependencies.
+ /// \param rest acts as-if: const KernelType &KernelFunc".
+ // TODO: Rest will represent reduction types once it is supported.
+ template <typename KernelName = detail::AutoName, typename... Rest>
+ event parallel_for(range<3> numWorkItems, const std::vector<event> &depEvents,
+ Rest &&...rest) {
+ return parallelForImpl<KernelName>(numWorkItems, depEvents,
+ std::forward<Rest>(rest)...);
+ }
+
private:
- // Name of this function is defined by compiler. It generates call to this
- // function in the host implementation of KernelFunc in submitSingleTask.
+ 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();
+ }
+
+ /// Name of this function is defined by compiler. It generates call to this
+ /// function in the host implementation of KernelFunc in submitSingleTask or
+ /// submitParallelFor.
+ /// \param KernelName a name of the kernel being invoked.
+ /// \param args kernel arguments for kernel invocation.
----------------
sergey-semenov wrote:
```suggestion
/// \param KernelName the name of the kernel being invoked.
/// \param args the kernel arguments for the kernel invocation.
```
https://github.com/llvm/llvm-project/pull/189068
More information about the llvm-branch-commits
mailing list