[llvm-branch-commits] [llvm] [libsycl] Add parallel_for feature (PR #189068)

[Yury Plyakhin via llvm-branch-commits]

================
@@ -199,18 +178,169 @@ class _LIBSYCL_EXPORT queue {
         "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)...);
+  }
+
   /// Blocks the calling thread until all commands previously submitted to this
   /// queue have completed. Synchronous errors are reported through SYCL
   /// exceptions.
   void wait();
 
 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,
----------------
YuriPlyakhin wrote:

static_assert might be better, since condition is constexpr.

https://github.com/llvm/llvm-project/pull/189068