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

Kseniya Tikhomirova via llvm-commits llvm-commits at lists.llvm.org
Thu Jul 16 03:09:07 PDT 2026


================
@@ -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() !=
----------------
KseniyaTikhomirova wrote:

I don't get it. I check if local range is 0 on line 352, since condition OR following it - next check should be executed only if local range is not 0. 
Could you please clarify what exact scenario you found incorrect here?

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


More information about the llvm-commits mailing list