[llvm] [libsycl] Implement nd_range kernel submissions (PR #206505)

[Sergey Semenov via llvm-commits]

================
@@ -332,6 +332,33 @@ class _LIBSYCL_EXPORT queue {
                                        std::forward<Rest>(rest)...);
   }
 
+  template <typename KernelName = detail::AutoName, int Dims, typename... Rest>
+  event parallel_for(nd_range<Dims> executionRange, Rest &&...rest) {
+    return parallel_for<KernelName, Dims, Rest...>(executionRange, {},
+                                                   std::forward<Rest>(rest)...);
+  }
+
+  template <typename KernelName = detail::AutoName, int Dims, typename... Rest>
+  event parallel_for(nd_range<Dims> executionRange, event depEvent,
+                     Rest &&...rest) {
+    return parallel_for<KernelName, Dims, Rest...>(executionRange, {depEvent},
+                                                   std::forward<Rest>(rest)...);
+  }
+
+  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() == range<Dims>{} ||
+         executionRange.get_global_range() % executionRange.get_local_range() !=
----------------
sergey-semenov wrote:

This check looks like it could lead to mod 0 for certain inputs.

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