From d4700e59f0ae2bcd031c0e504e6d7967fd7c6c3a Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Thu, 27 Feb 2025 11:02:25 +0000
Subject: [PATCH 01/34] [NATIVECPU] faster enqueue for larger ranges

---
 .../source/adapters/native_cpu/enqueue.cpp    | 173 +++++++++++-------
 1 file changed, 104 insertions(+), 69 deletions(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index 6cd1f6af8e660..40232868ea65e 100644
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -63,6 +63,29 @@ static native_cpu::state getResizedState(const native_cpu::NDRDescT &ndr,
 }
 #endif
 
+using IndexT = std::array<size_t, 3>;
+using RangeT = native_cpu::NDRDescT::RangeT;
+
+static inline void execute_range(native_cpu::state &state,
+                                 const ur_kernel_handle_t_ &hKernel,
+                                 const std::vector<void *> &args, IndexT first,
+                                 IndexT lastPlusOne) {
+  for (size_t g2 = first[2]; g2 < lastPlusOne[2]; g2++) {
+    for (size_t g1 = first[1]; g1 < lastPlusOne[1]; g1++) {
+      for (size_t g0 = first[0]; g0 < lastPlusOne[0]; g0 += 1) {
+        state.update(g0, g1, g2);
+        hKernel._subhandler(args.data(), &state);
+      }
+    }
+  }
+}
+
+static inline void execute_range(native_cpu::state &state,
+                                 const ur_kernel_handle_t_ &hKernel,
+                                 IndexT first, IndexT lastPlusOne) {
+  execute_range(state, hKernel, hKernel.getArgs(), first, lastPlusOne);
+}
+
 UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
     ur_queue_handle_t hQueue, ur_kernel_handle_t hKernel, uint32_t workDim,
     const size_t *pGlobalWorkOffset, const size_t *pGlobalWorkSize,
@@ -158,89 +181,101 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
     size_t new_num_work_groups_0 = numParallelThreads;
     size_t itemsPerThread = ndr.GlobalSize[0] / numParallelThreads;
 
-    for (unsigned g2 = 0; g2 < numWG2; g2++) {
-      for (unsigned g1 = 0; g1 < numWG1; g1++) {
-        for (unsigned g0 = 0; g0 < new_num_work_groups_0; g0 += 1) {
-          futures.emplace_back(tp.schedule_task(
-              [ndr, itemsPerThread, &kernel = *kernel, g0, g1, g2](size_t) {
-                native_cpu::state resized_state =
-                    getResizedState(ndr, itemsPerThread);
-                resized_state.update(g0, g1, g2);
-                kernel._subhandler(kernel.getArgs().data(), &resized_state);
-              }));
-        }
-        // Peel the remaining work items. Since the local size is 1, we iterate
-        // over the work groups.
-        for (unsigned g0 = new_num_work_groups_0 * itemsPerThread; g0 < numWG0;
-             g0++) {
-          state.update(g0, g1, g2);
-          kernel->_subhandler(kernel->getArgs().data(), &state);
-        }
-      }
+    for (size_t t = 0; t < numParallelThreads;) {
+      IndexT first = {t, 0, 0};
+      IndexT last = {++t, numWG1, numWG2};
+      futures.emplace_back(tp.schedule_task(
+          [ndr, itemsPerThread, &kernel = *kernel, first, last](size_t) {
+            native_cpu::state resized_state =
+                getResizedState(ndr, itemsPerThread);
+            execute_range(resized_state, kernel, first, last);
+          }));
+    }
+
+    size_t start_wg0_remainder = new_num_work_groups_0 * itemsPerThread;
+    if (start_wg0_remainder < numWG0) {
+      // Peel the remaining work items. Since the local size is 1, we iterate
+      // over the work groups.
+      futures.emplace_back(
+          tp.schedule_task([state, &kernel = *kernel, start_wg0_remainder,
+                            numWG0, numWG1, numWG2](size_t) mutable {
+            IndexT first = {start_wg0_remainder, 0, 0};
+            IndexT last = {numWG0, numWG1, numWG2};
+            execute_range(state, kernel, first, last);
+          }));
     }
 
   } else {
     // We are running a parallel_for over an nd_range
+    const auto numWG0_per_thread = numWG0 / numParallelThreads;
 
-    if (numWG1 * numWG2 >= numParallelThreads) {
-      // Dimensions 1 and 2 have enough work, split them across the threadpool
-      for (unsigned g2 = 0; g2 < numWG2; g2++) {
-        for (unsigned g1 = 0; g1 < numWG1; g1++) {
-          futures.emplace_back(
-              tp.schedule_task([state, &kernel = *kernel, numWG0, g1, g2,
-                                numParallelThreads](size_t threadId) mutable {
-                for (unsigned g0 = 0; g0 < numWG0; g0++) {
-                  state.update(g0, g1, g2);
-                  kernel._subhandler(
-                      kernel.getArgs(numParallelThreads, threadId).data(),
-                      &state);
-                }
-              }));
-        }
+    if (numWG0_per_thread) {
+      for (size_t t = 0, WG0_start = 0; t < numParallelThreads; t++) {
+        IndexT first = {WG0_start, 0, 0};
+        WG0_start += numWG0_per_thread;
+        IndexT last = {WG0_start, numWG1, numWG2};
+        futures.emplace_back(
+            tp.schedule_task([state, numParallelThreads, &kernel = *kernel,
+                              first, last](size_t threadId) mutable {
+              execute_range(state, kernel,
+                            kernel.getArgs(numParallelThreads, threadId), first,
+                            last);
+            }));
       }
+      size_t start_wg0_remainder = numWG0_per_thread * numParallelThreads;
+      if (start_wg0_remainder < numWG0) {
+        IndexT first = {start_wg0_remainder, 0, 0};
+        IndexT last = {numWG0, numWG1, numWG2};
+        futures.emplace_back(
+            tp.schedule_task([state, numParallelThreads, &kernel = *kernel,
+                              first, last](size_t threadId) mutable {
+              execute_range(state, kernel,
+                            kernel.getArgs(numParallelThreads, threadId), first,
+                            last);
+            }));
+      }
+
     } else {
-      // Split dimension 0 across the threadpool
       // Here we try to create groups of workgroups in order to reduce
       // synchronization overhead
-      for (unsigned g2 = 0; g2 < numWG2; g2++) {
-        for (unsigned g1 = 0; g1 < numWG1; g1++) {
-          for (unsigned g0 = 0; g0 < numWG0; g0++) {
-            groups.push_back([state, g0, g1, g2, numParallelThreads](
-                                 size_t threadId,
-                                 ur_kernel_handle_t_ &kernel) mutable {
-              state.update(g0, g1, g2);
-              kernel._subhandler(
-                  kernel.getArgs(numParallelThreads, threadId).data(), &state);
-            });
-          }
-        }
-      }
-      auto numGroups = groups.size();
+
+      // todo: deal with overflow
+      auto numGroups = numWG2 * numWG1 * numWG0;
       auto groupsPerThread = numGroups / numParallelThreads;
+
+      IndexT first = {0, 0, 0};
+      size_t counter = 0;
       if (groupsPerThread) {
-        for (unsigned thread = 0; thread < numParallelThreads; thread++) {
-          futures.emplace_back(
-              tp.schedule_task([groups, thread, groupsPerThread,
-                                &kernel = *kernel](size_t threadId) {
-                for (unsigned i = 0; i < groupsPerThread; i++) {
-                  auto index = thread * groupsPerThread + i;
-                  groups[index](threadId, kernel);
-                }
-              }));
+        for (unsigned g2 = 0; g2 < numWG2; g2++) {
+          for (unsigned g1 = 0; g1 < numWG1; g1++) {
+            for (unsigned g0 = 0; g0 < numWG0; g0++) {
+              if (counter == 0)
+                first = {g0, g1, g2};
+              if (++counter == groupsPerThread) {
+                IndexT last = {g0 + 1, g1 + 1, g2 + 1};
+                futures.emplace_back(tp.schedule_task(
+                    [state, numParallelThreads, &kernel = *kernel, first,
+                     last](size_t threadId) mutable {
+                      execute_range(
+                          state, kernel,
+                          kernel.getArgs(numParallelThreads, threadId), first,
+                          last);
+                    }));
+                counter = 0;
+              }
+            }
+          }
         }
       }
-
-      // schedule the remaining tasks
-      auto remainder = numGroups % numParallelThreads;
-      if (remainder) {
+      if (numGroups % numParallelThreads) {
+        // we have a remainder
+        IndexT last = {numWG0, numWG1, numWG2};
         futures.emplace_back(
-            tp.schedule_task([groups, remainder,
-                              scheduled = numParallelThreads * groupsPerThread,
-                              &kernel = *kernel](size_t threadId) {
-              for (unsigned i = 0; i < remainder; i++) {
-                auto index = scheduled + i;
-                groups[index](threadId, kernel);
-              }
+            tp.schedule_task([state, numParallelThreads, &kernel = *kernel,
+                              first, last](size_t threadId) mutable {
+              execute_range(state, kernel,
+                            kernel.getArgs(numParallelThreads, threadId), first,
+                            last);
             }));
       }
     }

From 960b1d50353f6cca12f8cccf3f5ea9f52123edf5 Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Fri, 28 Feb 2025 10:00:03 +0000
Subject: [PATCH 02/34] [NATIVECPU] removed unused groups

---
 unified-runtime/source/adapters/native_cpu/enqueue.cpp | 1 -
 1 file changed, 1 deletion(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index 40232868ea65e..e3efa5e3c774d 100644
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -130,7 +130,6 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
   auto &tp = hQueue->getDevice()->tp;
   const size_t numParallelThreads = tp.num_threads();
   std::vector<std::future<void>> futures;
-  std::vector<std::function<void(size_t, ur_kernel_handle_t_ &)>> groups;
   auto numWG0 = ndr.GlobalSize[0] / ndr.LocalSize[0];
   auto numWG1 = ndr.GlobalSize[1] / ndr.LocalSize[1];
   auto numWG2 = ndr.GlobalSize[2] / ndr.LocalSize[2];

From 04bd48abee0f319f6909fd118c90b8405f6e46dc Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Thu, 27 Mar 2025 16:32:23 +0000
Subject: [PATCH 03/34] [NATIVECPU] added async memcpy

---
 .../source/adapters/native_cpu/enqueue.cpp    | 58 ++++++++++++++-----
 1 file changed, 43 insertions(+), 15 deletions(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index e3efa5e3c774d..6f33b23f80bd6 100644
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -299,26 +299,43 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
   return UR_RESULT_SUCCESS;
 }
 
-template <class T>
+template <class T, class I>
 static inline ur_result_t
 withTimingEvent(ur_command_t command_type, ur_queue_handle_t hQueue,
                 uint32_t numEventsInWaitList,
                 const ur_event_handle_t *phEventWaitList,
-                ur_event_handle_t *phEvent, T &&f) {
+                ur_event_handle_t *phEvent, T &&f, I &&inv) {
   urEventWait(numEventsInWaitList, phEventWaitList);
   ur_event_handle_t event = nullptr;
   if (phEvent) {
-    event = new ur_event_handle_t_(hQueue, command_type);
+    ur_event_handle_t event = new ur_event_handle_t_(hQueue, command_type);
     event->tick_start();
+    ur_result_t result = inv(std::forward<T>(f), event);
+    *phEvent = event;
+    return result;
   }
-
   ur_result_t result = f();
+  return result;
+}
 
-  if (phEvent) {
+struct BlockingWithEvent {
+  template <class T>
+  ur_result_t operator()(T &&op, ur_event_handle_t event) const {
+    ur_result_t result = op();
     event->tick_end();
-    *phEvent = event;
+    return result;
   }
-  return result;
+};
+
+template <class T>
+static inline ur_result_t
+withTimingEvent(ur_command_t command_type, ur_queue_handle_t hQueue,
+                uint32_t numEventsInWaitList,
+                const ur_event_handle_t *phEventWaitList,
+                ur_event_handle_t *phEvent, T &&f) {
+  return withTimingEvent(command_type, hQueue, numEventsInWaitList,
+                         phEventWaitList, phEvent, std::forward<T>(f),
+                         BlockingWithEvent());
 }
 
 UR_APIEXPORT ur_result_t UR_APICALL urEnqueueEventsWait(
@@ -654,18 +671,29 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMMemcpy(
     ur_queue_handle_t hQueue, bool blocking, void *pDst, const void *pSrc,
     size_t size, uint32_t numEventsInWaitList,
     const ur_event_handle_t *phEventWaitList, ur_event_handle_t *phEvent) {
-  std::ignore = blocking;
+  UR_ASSERT(hQueue, UR_RESULT_ERROR_INVALID_QUEUE);
+  UR_ASSERT(pDst, UR_RESULT_ERROR_INVALID_NULL_POINTER);
+  UR_ASSERT(pSrc, UR_RESULT_ERROR_INVALID_NULL_POINTER);
+
+  auto Inv = [blocking, hQueue, size](auto &&f, ur_event_handle_t event) {
+    if (blocking || size < 100)
+      return BlockingWithEvent()(f, event);
+    auto &tp = hQueue->getDevice()->tp;
+    std::vector<std::future<void>> futures;
+    futures.emplace_back(tp.schedule_task([f](size_t) { f(); }));
+    event->set_futures(futures);
+    event->set_callback([event]() { event->tick_end(); });
+    return UR_RESULT_SUCCESS;
+  };
+  // blocking op
   return withTimingEvent(
       UR_COMMAND_USM_MEMCPY, hQueue, numEventsInWaitList, phEventWaitList,
-      phEvent, [&]() {
-        UR_ASSERT(hQueue, UR_RESULT_ERROR_INVALID_QUEUE);
-        UR_ASSERT(pDst, UR_RESULT_ERROR_INVALID_NULL_POINTER);
-        UR_ASSERT(pSrc, UR_RESULT_ERROR_INVALID_NULL_POINTER);
-
+      phEvent,
+      [pDst, pSrc, size]() {
         memcpy(pDst, pSrc, size);
-
         return UR_RESULT_SUCCESS;
-      });
+      },
+      Inv);
 }
 
 UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMPrefetch(

From 7008b8ba778ea5e79d1c25b49ac1773b61d7265b Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Fri, 11 Apr 2025 09:30:55 +0100
Subject: [PATCH 04/34] [NATIVECPU] added non-blocking invoker, removed unused
 variable

---
 .../source/adapters/native_cpu/enqueue.cpp    | 22 +++++++++++++------
 1 file changed, 15 insertions(+), 7 deletions(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index 6f33b23f80bd6..67bb4cb00ac7d 100644
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -306,7 +306,6 @@ withTimingEvent(ur_command_t command_type, ur_queue_handle_t hQueue,
                 const ur_event_handle_t *phEventWaitList,
                 ur_event_handle_t *phEvent, T &&f, I &&inv) {
   urEventWait(numEventsInWaitList, phEventWaitList);
-  ur_event_handle_t event = nullptr;
   if (phEvent) {
     ur_event_handle_t event = new ur_event_handle_t_(hQueue, command_type);
     event->tick_start();
@@ -327,6 +326,20 @@ struct BlockingWithEvent {
   }
 };
 
+struct NonBlocking {
+  ur_queue_handle_t hQueue;
+  NonBlocking(ur_queue_handle_t hQueue) : hQueue(hQueue) {}
+  template <class T>
+  ur_result_t operator()(T &&op, ur_event_handle_t event) const {
+    auto &tp = hQueue->getDevice()->tp;
+    std::vector<std::future<void>> futures;
+    futures.emplace_back(tp.schedule_task([op](size_t) { op(); }));
+    event->set_futures(futures);
+    event->set_callback([event]() { event->tick_end(); });
+    return UR_RESULT_SUCCESS;
+  }
+};
+
 template <class T>
 static inline ur_result_t
 withTimingEvent(ur_command_t command_type, ur_queue_handle_t hQueue,
@@ -678,12 +691,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMMemcpy(
   auto Inv = [blocking, hQueue, size](auto &&f, ur_event_handle_t event) {
     if (blocking || size < 100)
       return BlockingWithEvent()(f, event);
-    auto &tp = hQueue->getDevice()->tp;
-    std::vector<std::future<void>> futures;
-    futures.emplace_back(tp.schedule_task([f](size_t) { f(); }));
-    event->set_futures(futures);
-    event->set_callback([event]() { event->tick_end(); });
-    return UR_RESULT_SUCCESS;
+    return NonBlocking(hQueue)(f, event);
   };
   // blocking op
   return withTimingEvent(

From 2f1b3fed3f33668f3f4445a8e2d43639e455511e Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Wed, 16 Apr 2025 10:16:28 +0100
Subject: [PATCH 05/34] [NATIVECPU] waiting for events in threads

---
 .../source/adapters/native_cpu/enqueue.cpp    | 48 +++++++++++++++----
 1 file changed, 38 insertions(+), 10 deletions(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index 6cd1f6af8e660..bd0181f91433b 100644
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -50,6 +50,16 @@ struct NDRDescT {
        << GlobalOffset[2] << "\n";
   }
 };
+
+namespace {
+struct WaitInfo {
+  std::vector<ur_event_handle_t> events;
+  WaitInfo() = default;
+  WaitInfo(uint32_t numEvents, const ur_event_handle_t *WaitList)
+      : events(WaitList, WaitList + numEvents) {}
+  void wait() const { urEventWait(events.size(), events.data()); }
+};
+} // namespace
 } // namespace native_cpu
 
 #ifdef NATIVECPU_USE_OCK
@@ -69,7 +79,6 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
     const size_t *pLocalWorkSize, uint32_t numEventsInWaitList,
     const ur_event_handle_t *phEventWaitList, ur_event_handle_t *phEvent) {
 
-  urEventWait(numEventsInWaitList, phEventWaitList);
   UR_ASSERT(hQueue, UR_RESULT_ERROR_INVALID_NULL_HANDLE);
   UR_ASSERT(hKernel, UR_RESULT_ERROR_INVALID_NULL_HANDLE);
   UR_ASSERT(pGlobalWorkOffset, UR_RESULT_ERROR_INVALID_NULL_POINTER);
@@ -123,6 +132,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
   kernel->updateMemPool(numParallelThreads);
 
 #ifndef NATIVECPU_USE_OCK
+  urEventWait(numEventsInWaitList, phEventWaitList);
   for (unsigned g2 = 0; g2 < numWG2; g2++) {
     for (unsigned g1 = 0; g1 < numWG1; g1++) {
       for (unsigned g0 = 0; g0 < numWG0; g0++) {
@@ -138,6 +148,11 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
     }
   }
 #else
+  native_cpu::WaitInfo *const InEvents =
+      (numEventsInWaitList && phEventWaitList)
+          ? new native_cpu::WaitInfo(numEventsInWaitList, phEventWaitList)
+          : nullptr;
+
   bool isLocalSizeOne =
       ndr.LocalSize[0] == 1 && ndr.LocalSize[1] == 1 && ndr.LocalSize[2] == 1;
   if (isLocalSizeOne && ndr.GlobalSize[0] > numParallelThreads &&
@@ -157,15 +172,19 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
 
     size_t new_num_work_groups_0 = numParallelThreads;
     size_t itemsPerThread = ndr.GlobalSize[0] / numParallelThreads;
+    bool doneWaiting = false;
 
     for (unsigned g2 = 0; g2 < numWG2; g2++) {
       for (unsigned g1 = 0; g1 < numWG1; g1++) {
         for (unsigned g0 = 0; g0 < new_num_work_groups_0; g0 += 1) {
-          futures.emplace_back(tp.schedule_task(
-              [ndr, itemsPerThread, &kernel = *kernel, g0, g1, g2](size_t) {
+          futures.emplace_back(
+              tp.schedule_task([ndr, itemsPerThread, &kernel = *kernel, g0, g1,
+                                g2, InEvents](size_t) {
                 native_cpu::state resized_state =
                     getResizedState(ndr, itemsPerThread);
                 resized_state.update(g0, g1, g2);
+                if (InEvents)
+                  InEvents->wait();
                 kernel._subhandler(kernel.getArgs().data(), &resized_state);
               }));
         }
@@ -174,6 +193,10 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
         for (unsigned g0 = new_num_work_groups_0 * itemsPerThread; g0 < numWG0;
              g0++) {
           state.update(g0, g1, g2);
+          if (InEvents && !doneWaiting) {
+            InEvents->wait();
+            doneWaiting = true;
+          }
           kernel->_subhandler(kernel->getArgs().data(), &state);
         }
       }
@@ -186,11 +209,13 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
       // Dimensions 1 and 2 have enough work, split them across the threadpool
       for (unsigned g2 = 0; g2 < numWG2; g2++) {
         for (unsigned g1 = 0; g1 < numWG1; g1++) {
-          futures.emplace_back(
-              tp.schedule_task([state, &kernel = *kernel, numWG0, g1, g2,
-                                numParallelThreads](size_t threadId) mutable {
+          futures.emplace_back(tp.schedule_task(
+              [state, &kernel = *kernel, numWG0, g1, g2, numParallelThreads,
+               InEvents](size_t threadId) mutable {
                 for (unsigned g0 = 0; g0 < numWG0; g0++) {
                   state.update(g0, g1, g2);
+                  if (InEvents)
+                    InEvents->wait();
                   kernel._subhandler(
                       kernel.getArgs(numParallelThreads, threadId).data(),
                       &state);
@@ -205,10 +230,12 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
       for (unsigned g2 = 0; g2 < numWG2; g2++) {
         for (unsigned g1 = 0; g1 < numWG1; g1++) {
           for (unsigned g0 = 0; g0 < numWG0; g0++) {
-            groups.push_back([state, g0, g1, g2, numParallelThreads](
-                                 size_t threadId,
-                                 ur_kernel_handle_t_ &kernel) mutable {
+            groups.push_back([state, g0, g1, g2, numParallelThreads,
+                              InEvents](size_t threadId,
+                                        ur_kernel_handle_t_ &kernel) mutable {
               state.update(g0, g1, g2);
+              if (InEvents)
+                InEvents->wait();
               kernel._subhandler(
                   kernel.getArgs(numParallelThreads, threadId).data(), &state);
             });
@@ -252,10 +279,11 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
   if (phEvent) {
     *phEvent = event;
   }
-  event->set_callback([kernel = std::move(kernel), hKernel, event]() {
+  event->set_callback([kernel = std::move(kernel), hKernel, event, InEvents]() {
     event->tick_end();
     // TODO: avoid calling clear() here.
     hKernel->_localArgInfo.clear();
+    delete InEvents;
   });
 
   if (hQueue->isInOrder()) {

From 8efb1e437658b81404112031fc097b6341ddc3c1 Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Tue, 22 Apr 2025 16:51:28 +0100
Subject: [PATCH 06/34] [NATIVECPU] ndrange enqueue with less work for main
 thread

---
 .../source/adapters/native_cpu/enqueue.cpp    | 87 ++++++-------------
 1 file changed, 27 insertions(+), 60 deletions(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index e3efa5e3c774d..bf63d53fee5b0 100644
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -206,25 +206,25 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
 
   } else {
     // We are running a parallel_for over an nd_range
-    const auto numWG0_per_thread = numWG0 / numParallelThreads;
 
-    if (numWG0_per_thread) {
-      for (size_t t = 0, WG0_start = 0; t < numParallelThreads; t++) {
-        IndexT first = {WG0_start, 0, 0};
-        WG0_start += numWG0_per_thread;
-        IndexT last = {WG0_start, numWG1, numWG2};
-        futures.emplace_back(
-            tp.schedule_task([state, numParallelThreads, &kernel = *kernel,
-                              first, last](size_t threadId) mutable {
-              execute_range(state, kernel,
-                            kernel.getArgs(numParallelThreads, threadId), first,
-                            last);
-            }));
-      }
-      size_t start_wg0_remainder = numWG0_per_thread * numParallelThreads;
-      if (start_wg0_remainder < numWG0) {
-        IndexT first = {start_wg0_remainder, 0, 0};
-        IndexT last = {numWG0, numWG1, numWG2};
+    const IndexT numWG = {numWG0, numWG1, numWG2};
+    IndexT groupsPerThread;
+    for (size_t t = 0; t < 3; t++)
+      groupsPerThread[t] = numWG[t] / numParallelThreads;
+    size_t dim = 0;
+    if (groupsPerThread[0] == 0) {
+      if (groupsPerThread[1])
+        dim = 1;
+      else if (groupsPerThread[2])
+        dim = 2;
+    }
+    IndexT first = {0, 0, 0}, last = numWG;
+    size_t wg_start = 0;
+    if (groupsPerThread[dim]) {
+      for (size_t t = 0; t < numParallelThreads; t++) {
+        first[dim] = wg_start;
+        wg_start += groupsPerThread[dim];
+        last[dim] = wg_start;
         futures.emplace_back(
             tp.schedule_task([state, numParallelThreads, &kernel = *kernel,
                               first, last](size_t threadId) mutable {
@@ -233,50 +233,17 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
                             last);
             }));
       }
-
-    } else {
-      // Here we try to create groups of workgroups in order to reduce
-      // synchronization overhead
-
-      // todo: deal with overflow
-      auto numGroups = numWG2 * numWG1 * numWG0;
-      auto groupsPerThread = numGroups / numParallelThreads;
-
-      IndexT first = {0, 0, 0};
-      size_t counter = 0;
-      if (groupsPerThread) {
-        for (unsigned g2 = 0; g2 < numWG2; g2++) {
-          for (unsigned g1 = 0; g1 < numWG1; g1++) {
-            for (unsigned g0 = 0; g0 < numWG0; g0++) {
-              if (counter == 0)
-                first = {g0, g1, g2};
-              if (++counter == groupsPerThread) {
-                IndexT last = {g0 + 1, g1 + 1, g2 + 1};
-                futures.emplace_back(tp.schedule_task(
-                    [state, numParallelThreads, &kernel = *kernel, first,
-                     last](size_t threadId) mutable {
-                      execute_range(
-                          state, kernel,
+    }
+    if (wg_start < numWG[dim]) {
+      first[dim] = wg_start;
+      last[dim] = numWG[dim];
+      futures.emplace_back(
+          tp.schedule_task([state, numParallelThreads, &kernel = *kernel, first,
+                            last](size_t threadId) mutable {
+            execute_range(state, kernel,
                           kernel.getArgs(numParallelThreads, threadId), first,
                           last);
-                    }));
-                counter = 0;
-              }
-            }
-          }
-        }
-      }
-      if (numGroups % numParallelThreads) {
-        // we have a remainder
-        IndexT last = {numWG0, numWG1, numWG2};
-        futures.emplace_back(
-            tp.schedule_task([state, numParallelThreads, &kernel = *kernel,
-                              first, last](size_t threadId) mutable {
-              execute_range(state, kernel,
-                            kernel.getArgs(numParallelThreads, threadId), first,
-                            last);
-            }));
-      }
+          }));
     }
   }
 

From 2c5218664cf9d4f8be7333ca576c09e0e3818edc Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Tue, 22 Apr 2025 17:02:57 +0100
Subject: [PATCH 07/34] [NATIVECPU] static_assert for pointer type

---
 unified-runtime/source/adapters/native_cpu/enqueue.cpp | 1 +
 1 file changed, 1 insertion(+)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index bd0181f91433b..cac2fa1dd2b89 100644
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -55,6 +55,7 @@ namespace {
 struct WaitInfo {
   std::vector<ur_event_handle_t> events;
   WaitInfo() = default;
+  static_assert(std::is_pointer_v<ur_event_handle_t>);
   WaitInfo(uint32_t numEvents, const ur_event_handle_t *WaitList)
       : events(WaitList, WaitList + numEvents) {}
   void wait() const { urEventWait(events.size(), events.data()); }

From 5348490068b7e91713b2f542550c58398ef70e91 Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Tue, 22 Apr 2025 17:59:40 +0100
Subject: [PATCH 08/34] [NATIVECPU] added anonymous namespace

---
 unified-runtime/source/adapters/native_cpu/enqueue.cpp | 2 ++
 1 file changed, 2 insertions(+)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index 7e16e1fd81e9a..439d52ef43d85 100644
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -284,6 +284,7 @@ withTimingEvent(ur_command_t command_type, ur_queue_handle_t hQueue,
   return result;
 }
 
+namespace {
 struct BlockingWithEvent {
   template <class T>
   ur_result_t operator()(T &&op, ur_event_handle_t event) const {
@@ -306,6 +307,7 @@ struct NonBlocking {
     return UR_RESULT_SUCCESS;
   }
 };
+} // namespace
 
 template <class T>
 static inline ur_result_t

From 1de1251079b23671624f9aa2870415b6bef1eb49 Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Tue, 22 Apr 2025 19:28:44 +0100
Subject: [PATCH 09/34] [NATIVECPU] separated out Invokers for enqueues

---
 .../source/adapters/native_cpu/enqueue.cpp    | 30 ++++++++++++-------
 1 file changed, 19 insertions(+), 11 deletions(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index 439d52ef43d85..23252c88be796 100644
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -276,7 +276,7 @@ withTimingEvent(ur_command_t command_type, ur_queue_handle_t hQueue,
   if (phEvent) {
     ur_event_handle_t event = new ur_event_handle_t_(hQueue, command_type);
     event->tick_start();
-    ur_result_t result = inv(std::forward<T>(f), event);
+    ur_result_t result = inv(std::forward<T>(f), event, hQueue);
     *phEvent = event;
     return result;
   }
@@ -287,7 +287,8 @@ withTimingEvent(ur_command_t command_type, ur_queue_handle_t hQueue,
 namespace {
 struct BlockingWithEvent {
   template <class T>
-  ur_result_t operator()(T &&op, ur_event_handle_t event) const {
+  ur_result_t operator()(T &&op, ur_event_handle_t event,
+                         ur_queue_handle_t) const {
     ur_result_t result = op();
     event->tick_end();
     return result;
@@ -295,10 +296,9 @@ struct BlockingWithEvent {
 };
 
 struct NonBlocking {
-  ur_queue_handle_t hQueue;
-  NonBlocking(ur_queue_handle_t hQueue) : hQueue(hQueue) {}
   template <class T>
-  ur_result_t operator()(T &&op, ur_event_handle_t event) const {
+  ur_result_t operator()(T &&op, ur_event_handle_t event,
+                         ur_queue_handle_t hQueue) const {
     auto &tp = hQueue->getDevice()->tp;
     std::vector<std::future<void>> futures;
     futures.emplace_back(tp.schedule_task([op](size_t) { op(); }));
@@ -307,6 +307,19 @@ struct NonBlocking {
     return UR_RESULT_SUCCESS;
   }
 };
+
+struct Invoker {
+  const bool blocking;
+  Invoker(bool blocking) : blocking(blocking) {}
+  template <class T>
+  ur_result_t operator()(T &&f, ur_event_handle_t event,
+                         ur_queue_handle_t hQueue) const {
+    if (blocking)
+      return BlockingWithEvent()(std::forward<T>(f), event, hQueue);
+    return NonBlocking()(std::forward<T>(f), event, hQueue);
+  };
+};
+
 } // namespace
 
 template <class T>
@@ -624,11 +637,6 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMMemcpy(
   UR_ASSERT(pDst, UR_RESULT_ERROR_INVALID_NULL_POINTER);
   UR_ASSERT(pSrc, UR_RESULT_ERROR_INVALID_NULL_POINTER);
 
-  auto Inv = [blocking, hQueue, size](auto &&f, ur_event_handle_t event) {
-    if (blocking || size < 100)
-      return BlockingWithEvent()(f, event);
-    return NonBlocking(hQueue)(f, event);
-  };
   return withTimingEvent(
       UR_COMMAND_USM_MEMCPY, hQueue, numEventsInWaitList, phEventWaitList,
       phEvent,
@@ -636,7 +644,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMMemcpy(
         memcpy(pDst, pSrc, size);
         return UR_RESULT_SUCCESS;
       },
-      Inv);
+      Invoker(blocking || size < 100));
 }
 
 UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMPrefetch(

From 9173f5ebc27c1202cb47a46d89e33a077126a4c3 Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Wed, 23 Apr 2025 09:51:53 +0100
Subject: [PATCH 10/34] [NATIVECPU] made more memops async

---
 .../source/adapters/native_cpu/enqueue.cpp    | 39 +++++++++++--------
 1 file changed, 22 insertions(+), 17 deletions(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index 23252c88be796..e016e6cc79367 100644
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -406,39 +406,43 @@ static inline ur_result_t enqueueMemBufferReadWriteRect_impl(
       });
 }
 
-static inline ur_result_t doCopy_impl(ur_queue_handle_t hQueue, void *DstPtr,
-                                      const void *SrcPtr, size_t Size,
-                                      uint32_t numEventsInWaitList,
-                                      const ur_event_handle_t *phEventWaitList,
-                                      ur_event_handle_t *phEvent,
-                                      ur_command_t command_type) {
-  return withTimingEvent(command_type, hQueue, numEventsInWaitList,
-                         phEventWaitList, phEvent, [&]() {
-                           if (SrcPtr != DstPtr && Size)
-                             memmove(DstPtr, SrcPtr, Size);
-                           return UR_RESULT_SUCCESS;
-                         });
+template <class T>
+static inline ur_result_t
+doCopy_impl(ur_queue_handle_t hQueue, void *DstPtr, const void *SrcPtr,
+            size_t Size, uint32_t numEventsInWaitList,
+            const ur_event_handle_t *phEventWaitList,
+            ur_event_handle_t *phEvent, ur_command_t command_type, T &&Inv) {
+  return withTimingEvent(
+      command_type, hQueue, numEventsInWaitList, phEventWaitList, phEvent,
+      [DstPtr, SrcPtr, Size]() {
+        if (SrcPtr != DstPtr && Size)
+          memmove(DstPtr, SrcPtr, Size);
+        return UR_RESULT_SUCCESS;
+      },
+      Inv);
 }
 
 UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemBufferRead(
-    ur_queue_handle_t hQueue, ur_mem_handle_t hBuffer, bool /*blockingRead*/,
+    ur_queue_handle_t hQueue, ur_mem_handle_t hBuffer, bool blockingRead,
     size_t offset, size_t size, void *pDst, uint32_t numEventsInWaitList,
     const ur_event_handle_t *phEventWaitList, ur_event_handle_t *phEvent) {
 
   void *FromPtr = /*Src*/ hBuffer->_mem + offset;
   auto res = doCopy_impl(hQueue, pDst, FromPtr, size, numEventsInWaitList,
-                         phEventWaitList, phEvent, UR_COMMAND_MEM_BUFFER_READ);
+                         phEventWaitList, phEvent, UR_COMMAND_MEM_BUFFER_READ,
+                         Invoker(blockingRead));
   return res;
 }
 
 UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemBufferWrite(
-    ur_queue_handle_t hQueue, ur_mem_handle_t hBuffer, bool /*blockingWrite*/,
+    ur_queue_handle_t hQueue, ur_mem_handle_t hBuffer, bool blockingWrite,
     size_t offset, size_t size, const void *pSrc, uint32_t numEventsInWaitList,
     const ur_event_handle_t *phEventWaitList, ur_event_handle_t *phEvent) {
 
   void *ToPtr = hBuffer->_mem + offset;
   auto res = doCopy_impl(hQueue, ToPtr, pSrc, size, numEventsInWaitList,
-                         phEventWaitList, phEvent, UR_COMMAND_MEM_BUFFER_WRITE);
+                         phEventWaitList, phEvent, UR_COMMAND_MEM_BUFFER_WRITE,
+                         Invoker(blockingWrite));
   return res;
 }
 
@@ -477,7 +481,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemBufferCopy(
   const void *SrcPtr = hBufferSrc->_mem + srcOffset;
   void *DstPtr = hBufferDst->_mem + dstOffset;
   return doCopy_impl(hQueue, DstPtr, SrcPtr, size, numEventsInWaitList,
-                     phEventWaitList, phEvent, UR_COMMAND_MEM_BUFFER_COPY);
+                     phEventWaitList, phEvent, UR_COMMAND_MEM_BUFFER_COPY,
+                     BlockingWithEvent() /*TODO: check blocking*/);
 }
 
 UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemBufferCopyRect(

From 7cd7caa8848d25065c676bdf25fc0526d5452f7d Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Wed, 23 Apr 2025 10:11:31 +0100
Subject: [PATCH 11/34] [NATIVECPU] memop pointer check outside worker lambda

---
 unified-runtime/source/adapters/native_cpu/enqueue.cpp | 8 ++++++--
 1 file changed, 6 insertions(+), 2 deletions(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index e016e6cc79367..baa97cb1af841 100644
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -412,11 +412,15 @@ doCopy_impl(ur_queue_handle_t hQueue, void *DstPtr, const void *SrcPtr,
             size_t Size, uint32_t numEventsInWaitList,
             const ur_event_handle_t *phEventWaitList,
             ur_event_handle_t *phEvent, ur_command_t command_type, T &&Inv) {
+  if (SrcPtr == DstPtr || Size == 0)
+    return withTimingEvent(
+        command_type, hQueue, numEventsInWaitList, phEventWaitList, phEvent,
+        []() { return UR_RESULT_SUCCESS; }, BlockingWithEvent());
+
   return withTimingEvent(
       command_type, hQueue, numEventsInWaitList, phEventWaitList, phEvent,
       [DstPtr, SrcPtr, Size]() {
-        if (SrcPtr != DstPtr && Size)
-          memmove(DstPtr, SrcPtr, Size);
+        memmove(DstPtr, SrcPtr, Size);
         return UR_RESULT_SUCCESS;
       },
       Inv);

From 32ecf0928b749a839b2b3f1a9745c5284cb67a1d Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Wed, 23 Apr 2025 15:18:07 +0100
Subject: [PATCH 12/34] [NATIVECPU] moved inEvents

---
 unified-runtime/source/adapters/native_cpu/enqueue.cpp | 10 +++++-----
 1 file changed, 5 insertions(+), 5 deletions(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index 75b1e63db2fd6..158f472a4382e 100644
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -132,6 +132,11 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
   auto kernel = std::make_unique<ur_kernel_handle_t_>(*hKernel);
   kernel->updateMemPool(numParallelThreads);
 
+  native_cpu::WaitInfo *const InEvents =
+      (numEventsInWaitList && phEventWaitList)
+          ? new native_cpu::WaitInfo(numEventsInWaitList, phEventWaitList)
+          : nullptr;
+
 #ifndef NATIVECPU_USE_OCK
   urEventWait(numEventsInWaitList, phEventWaitList);
   for (unsigned g2 = 0; g2 < numWG2; g2++) {
@@ -149,11 +154,6 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
     }
   }
 #else
-  native_cpu::WaitInfo *const InEvents =
-      (numEventsInWaitList && phEventWaitList)
-          ? new native_cpu::WaitInfo(numEventsInWaitList, phEventWaitList)
-          : nullptr;
-
   bool isLocalSizeOne =
       ndr.LocalSize[0] == 1 && ndr.LocalSize[1] == 1 && ndr.LocalSize[2] == 1;
   if (isLocalSizeOne && ndr.GlobalSize[0] > numParallelThreads &&

From 4b05062a0f2660b4a10984026eaec09f37f3e53e Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Mon, 28 Apr 2025 10:40:51 +0100
Subject: [PATCH 13/34] [NATIVECPU] use unique_ptr for WaitInfo

---
 .../source/adapters/native_cpu/enqueue.cpp    | 29 ++++++++++++-------
 1 file changed, 18 insertions(+), 11 deletions(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index 158f472a4382e..94fcf47b29be0 100644
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -54,12 +54,21 @@ struct NDRDescT {
 namespace {
 struct WaitInfo {
   std::vector<ur_event_handle_t> events;
-  WaitInfo() = default;
   static_assert(std::is_pointer_v<ur_event_handle_t>);
   WaitInfo(uint32_t numEvents, const ur_event_handle_t *WaitList)
       : events(WaitList, WaitList + numEvents) {}
   void wait() const { urEventWait(events.size(), events.data()); }
 };
+
+inline static std::unique_ptr<WaitInfo>
+getWaitInfo(uint32_t numEventsInWaitList,
+            const ur_event_handle_t *phEventWaitList) {
+  return (numEventsInWaitList && phEventWaitList)
+             ? std::make_unique<native_cpu::WaitInfo>(numEventsInWaitList,
+                                                      phEventWaitList)
+             : nullptr;
+}
+
 } // namespace
 } // namespace native_cpu
 
@@ -132,10 +141,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
   auto kernel = std::make_unique<ur_kernel_handle_t_>(*hKernel);
   kernel->updateMemPool(numParallelThreads);
 
-  native_cpu::WaitInfo *const InEvents =
-      (numEventsInWaitList && phEventWaitList)
-          ? new native_cpu::WaitInfo(numEventsInWaitList, phEventWaitList)
-          : nullptr;
+  auto InEvents = native_cpu::getWaitInfo(numEventsInWaitList, phEventWaitList);
 
 #ifndef NATIVECPU_USE_OCK
   urEventWait(numEventsInWaitList, phEventWaitList);
@@ -180,7 +186,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
         for (unsigned g0 = 0; g0 < new_num_work_groups_0; g0 += 1) {
           futures.emplace_back(
               tp.schedule_task([ndr, itemsPerThread, &kernel = *kernel, g0, g1,
-                                g2, InEvents](size_t) {
+                                g2, InEvents = InEvents.get()](size_t) {
                 native_cpu::state resized_state =
                     getResizedState(ndr, itemsPerThread);
                 resized_state.update(g0, g1, g2);
@@ -212,7 +218,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
         for (unsigned g1 = 0; g1 < numWG1; g1++) {
           futures.emplace_back(tp.schedule_task(
               [state, &kernel = *kernel, numWG0, g1, g2, numParallelThreads,
-               InEvents](size_t threadId) mutable {
+               InEvents = InEvents.get()](size_t threadId) mutable {
                 for (unsigned g0 = 0; g0 < numWG0; g0++) {
                   state.update(g0, g1, g2);
                   if (InEvents)
@@ -232,8 +238,9 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
         for (unsigned g1 = 0; g1 < numWG1; g1++) {
           for (unsigned g0 = 0; g0 < numWG0; g0++) {
             groups.push_back([state, g0, g1, g2, numParallelThreads,
-                              InEvents](size_t threadId,
-                                        ur_kernel_handle_t_ &kernel) mutable {
+                              InEvents = InEvents.get()](
+                                 size_t threadId,
+                                 ur_kernel_handle_t_ &kernel) mutable {
               state.update(g0, g1, g2);
               if (InEvents)
                 InEvents->wait();
@@ -280,11 +287,11 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
   if (phEvent) {
     *phEvent = event;
   }
-  event->set_callback([kernel = std::move(kernel), hKernel, event, InEvents]() {
+  event->set_callback([kernel = std::move(kernel), hKernel, event,
+                       InEvents = std::move(InEvents)]() {
     event->tick_end();
     // TODO: avoid calling clear() here.
     hKernel->_localArgInfo.clear();
-    delete InEvents;
   });
 
   if (hQueue->isInOrder()) {

From 2722cad08b7c0fbf2fe0925acc188f7e2c23452b Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Mon, 28 Apr 2025 12:53:51 +0100
Subject: [PATCH 14/34] [NATIVECPU] async memcopy

---
 .../source/adapters/native_cpu/enqueue.cpp    | 43 ++++++++++++++-----
 1 file changed, 33 insertions(+), 10 deletions(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index 055de9af4d6f8..bd57bdcea2db1 100644
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -294,14 +294,15 @@ withTimingEvent(ur_command_t command_type, ur_queue_handle_t hQueue,
                 uint32_t numEventsInWaitList,
                 const ur_event_handle_t *phEventWaitList,
                 ur_event_handle_t *phEvent, T &&f, I &&inv) {
-  urEventWait(numEventsInWaitList, phEventWaitList);
   if (phEvent) {
     ur_event_handle_t event = new ur_event_handle_t_(hQueue, command_type);
     event->tick_start();
-    ur_result_t result = inv(std::forward<T>(f), event, hQueue);
+    ur_result_t result = inv(std::forward<T>(f), event, hQueue,
+                             numEventsInWaitList, phEventWaitList);
     *phEvent = event;
     return result;
   }
+  urEventWait(numEventsInWaitList, phEventWaitList);
   ur_result_t result = f();
   return result;
 }
@@ -309,8 +310,10 @@ withTimingEvent(ur_command_t command_type, ur_queue_handle_t hQueue,
 namespace {
 struct BlockingWithEvent {
   template <class T>
-  ur_result_t operator()(T &&op, ur_event_handle_t event,
-                         ur_queue_handle_t) const {
+  ur_result_t operator()(T &&op, ur_event_handle_t event, ur_queue_handle_t,
+                         uint32_t numEventsInWaitList,
+                         const ur_event_handle_t *phEventWaitList) const {
+    urEventWait(numEventsInWaitList, phEventWaitList);
     ur_result_t result = op();
     event->tick_end();
     return result;
@@ -320,12 +323,24 @@ struct BlockingWithEvent {
 struct NonBlocking {
   template <class T>
   ur_result_t operator()(T &&op, ur_event_handle_t event,
-                         ur_queue_handle_t hQueue) const {
+                         ur_queue_handle_t hQueue, uint32_t numEventsInWaitList,
+                         const ur_event_handle_t *phEventWaitList) const {
     auto &tp = hQueue->getDevice()->tp;
     std::vector<std::future<void>> futures;
-    futures.emplace_back(tp.schedule_task([op](size_t) { op(); }));
+    native_cpu::WaitInfo *const InEvents =
+        (numEventsInWaitList && phEventWaitList)
+            ? new native_cpu::WaitInfo(numEventsInWaitList, phEventWaitList)
+            : nullptr;
+    futures.emplace_back(tp.schedule_task([op, InEvents](size_t) {
+      if (InEvents)
+        InEvents->wait();
+      op();
+    }));
     event->set_futures(futures);
-    event->set_callback([event]() { event->tick_end(); });
+    event->set_callback([event, InEvents]() {
+      event->tick_end();
+      delete InEvents;
+    });
     return UR_RESULT_SUCCESS;
   }
 };
@@ -335,10 +350,13 @@ struct Invoker {
   Invoker(bool blocking) : blocking(blocking) {}
   template <class T>
   ur_result_t operator()(T &&f, ur_event_handle_t event,
-                         ur_queue_handle_t hQueue) const {
+                         ur_queue_handle_t hQueue, uint32_t numEventsInWaitList,
+                         const ur_event_handle_t *phEventWaitList) const {
     if (blocking)
-      return BlockingWithEvent()(std::forward<T>(f), event, hQueue);
-    return NonBlocking()(std::forward<T>(f), event, hQueue);
+      return BlockingWithEvent()(std::forward<T>(f), event, hQueue,
+                                 numEventsInWaitList, phEventWaitList);
+    return NonBlocking()(std::forward<T>(f), event, hQueue, numEventsInWaitList,
+                         phEventWaitList);
   };
 };
 
@@ -668,6 +686,11 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMMemcpy(
   UR_ASSERT(pDst, UR_RESULT_ERROR_INVALID_NULL_POINTER);
   UR_ASSERT(pSrc, UR_RESULT_ERROR_INVALID_NULL_POINTER);
 
+  if (pSrc == pDst || size == 0)
+    return withTimingEvent(
+        UR_COMMAND_USM_MEMCPY, hQueue, numEventsInWaitList, phEventWaitList, phEvent,
+        []() { return UR_RESULT_SUCCESS; }, BlockingWithEvent());
+
   return withTimingEvent(
       UR_COMMAND_USM_MEMCPY, hQueue, numEventsInWaitList, phEventWaitList,
       phEvent,

From 22898b4a8774e6575b44c23715cf1f92349267b0 Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Mon, 28 Apr 2025 15:38:53 +0100
Subject: [PATCH 15/34] [NATIVECPU] code reuse for memcopies

---
 .../source/adapters/native_cpu/enqueue.cpp    | 21 ++++++-------------
 1 file changed, 6 insertions(+), 15 deletions(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index 526060075f31d..f77151f3f3200 100755
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -452,7 +452,8 @@ static inline ur_result_t enqueueMemBufferReadWriteRect_impl(
       });
 }
 
-template <class T>
+template <void *(copy_func)(void *dst, const void *src, size_t) = memmove,
+          class T>
 static inline ur_result_t
 doCopy_impl(ur_queue_handle_t hQueue, void *DstPtr, const void *SrcPtr,
             size_t Size, uint32_t numEventsInWaitList,
@@ -466,7 +467,7 @@ doCopy_impl(ur_queue_handle_t hQueue, void *DstPtr, const void *SrcPtr,
   return withTimingEvent(
       command_type, hQueue, numEventsInWaitList, phEventWaitList, phEvent,
       [DstPtr, SrcPtr, Size]() {
-        memmove(DstPtr, SrcPtr, Size);
+        copy_func(DstPtr, SrcPtr, Size);
         return UR_RESULT_SUCCESS;
       },
       Inv);
@@ -692,19 +693,9 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMMemcpy(
   UR_ASSERT(pDst, UR_RESULT_ERROR_INVALID_NULL_POINTER);
   UR_ASSERT(pSrc, UR_RESULT_ERROR_INVALID_NULL_POINTER);
 
-  if (pSrc == pDst || size == 0)
-    return withTimingEvent(
-        UR_COMMAND_USM_MEMCPY, hQueue, numEventsInWaitList, phEventWaitList,
-        phEvent, []() { return UR_RESULT_SUCCESS; }, BlockingWithEvent());
-
-  return withTimingEvent(
-      UR_COMMAND_USM_MEMCPY, hQueue, numEventsInWaitList, phEventWaitList,
-      phEvent,
-      [pDst, pSrc, size]() {
-        memcpy(pDst, pSrc, size);
-        return UR_RESULT_SUCCESS;
-      },
-      Invoker(blocking || size < 100));
+  return doCopy_impl<memcpy>(hQueue, pDst, pSrc, size, numEventsInWaitList,
+                             phEventWaitList, phEvent, UR_COMMAND_USM_MEMCPY,
+                             Invoker(blocking));
 }
 
 UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMPrefetch(

From 5d12b7ae3474cedef4b50b7719ed743668b7c88c Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Tue, 29 Apr 2025 09:58:45 +0100
Subject: [PATCH 16/34] [NATIVECPU] removed invoker

---
 .../source/adapters/native_cpu/enqueue.cpp    | 53 +++++++++----------
 1 file changed, 25 insertions(+), 28 deletions(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index f77151f3f3200..91b343a2eb51e 100755
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -350,22 +350,6 @@ struct NonBlocking {
     return UR_RESULT_SUCCESS;
   }
 };
-
-struct Invoker {
-  const bool blocking;
-  Invoker(bool blocking) : blocking(blocking) {}
-  template <class T>
-  ur_result_t operator()(T &&f, ur_event_handle_t event,
-                         ur_queue_handle_t hQueue, uint32_t numEventsInWaitList,
-                         const ur_event_handle_t *phEventWaitList) const {
-    if (blocking)
-      return BlockingWithEvent()(std::forward<T>(f), event, hQueue,
-                                 numEventsInWaitList, phEventWaitList);
-    return NonBlocking()(std::forward<T>(f), event, hQueue, numEventsInWaitList,
-                         phEventWaitList);
-  };
-};
-
 } // namespace
 
 template <class T>
@@ -379,6 +363,21 @@ withTimingEvent(ur_command_t command_type, ur_queue_handle_t hQueue,
                          BlockingWithEvent());
 }
 
+template <class T>
+static inline ur_result_t
+withTimingEvent(ur_command_t command_type, ur_queue_handle_t hQueue,
+                uint32_t numEventsInWaitList,
+                const ur_event_handle_t *phEventWaitList,
+                ur_event_handle_t *phEvent, T &&f, bool blocking) {
+  if (blocking)
+    return withTimingEvent(command_type, hQueue, numEventsInWaitList,
+                           phEventWaitList, phEvent, std::forward<T>(f),
+                           BlockingWithEvent());
+  return withTimingEvent(command_type, hQueue, numEventsInWaitList,
+                         phEventWaitList, phEvent, std::forward<T>(f),
+                         NonBlocking());
+}
+
 UR_APIEXPORT ur_result_t UR_APICALL urEnqueueEventsWait(
     ur_queue_handle_t hQueue, uint32_t numEventsInWaitList,
     const ur_event_handle_t *phEventWaitList, ur_event_handle_t *phEvent) {
@@ -452,13 +451,11 @@ static inline ur_result_t enqueueMemBufferReadWriteRect_impl(
       });
 }
 
-template <void *(copy_func)(void *dst, const void *src, size_t) = memmove,
-          class T>
-static inline ur_result_t
-doCopy_impl(ur_queue_handle_t hQueue, void *DstPtr, const void *SrcPtr,
-            size_t Size, uint32_t numEventsInWaitList,
-            const ur_event_handle_t *phEventWaitList,
-            ur_event_handle_t *phEvent, ur_command_t command_type, T &&Inv) {
+template <void *(copy_func)(void *dst, const void *src, size_t) = memmove>
+static inline ur_result_t doCopy_impl(
+    ur_queue_handle_t hQueue, void *DstPtr, const void *SrcPtr, size_t Size,
+    uint32_t numEventsInWaitList, const ur_event_handle_t *phEventWaitList,
+    ur_event_handle_t *phEvent, ur_command_t command_type, bool blocking) {
   if (SrcPtr == DstPtr || Size == 0)
     return withTimingEvent(
         command_type, hQueue, numEventsInWaitList, phEventWaitList, phEvent,
@@ -470,7 +467,7 @@ doCopy_impl(ur_queue_handle_t hQueue, void *DstPtr, const void *SrcPtr,
         copy_func(DstPtr, SrcPtr, Size);
         return UR_RESULT_SUCCESS;
       },
-      Inv);
+      blocking);
 }
 
 UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemBufferRead(
@@ -481,7 +478,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemBufferRead(
   void *FromPtr = /*Src*/ hBuffer->_mem + offset;
   auto res = doCopy_impl(hQueue, pDst, FromPtr, size, numEventsInWaitList,
                          phEventWaitList, phEvent, UR_COMMAND_MEM_BUFFER_READ,
-                         Invoker(blockingRead));
+                         blockingRead);
   return res;
 }
 
@@ -493,7 +490,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemBufferWrite(
   void *ToPtr = hBuffer->_mem + offset;
   auto res = doCopy_impl(hQueue, ToPtr, pSrc, size, numEventsInWaitList,
                          phEventWaitList, phEvent, UR_COMMAND_MEM_BUFFER_WRITE,
-                         Invoker(blockingWrite));
+                         blockingWrite);
   return res;
 }
 
@@ -533,7 +530,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemBufferCopy(
   void *DstPtr = hBufferDst->_mem + dstOffset;
   return doCopy_impl(hQueue, DstPtr, SrcPtr, size, numEventsInWaitList,
                      phEventWaitList, phEvent, UR_COMMAND_MEM_BUFFER_COPY,
-                     BlockingWithEvent() /*TODO: check blocking*/);
+                     true /*TODO: check blocking*/);
 }
 
 UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemBufferCopyRect(
@@ -695,7 +692,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMMemcpy(
 
   return doCopy_impl<memcpy>(hQueue, pDst, pSrc, size, numEventsInWaitList,
                              phEventWaitList, phEvent, UR_COMMAND_USM_MEMCPY,
-                             Invoker(blocking));
+                             blocking);
 }
 
 UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMPrefetch(

From 400ba0d6da60fe5e218297524ebf516800a421ee Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Tue, 29 Apr 2025 10:57:16 +0100
Subject: [PATCH 17/34] [NATIVECPU] removed unneeded function

---
 .../source/adapters/native_cpu/enqueue.cpp          | 13 +------------
 1 file changed, 1 insertion(+), 12 deletions(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index 91b343a2eb51e..6e688bca37e5b 100755
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -357,18 +357,7 @@ static inline ur_result_t
 withTimingEvent(ur_command_t command_type, ur_queue_handle_t hQueue,
                 uint32_t numEventsInWaitList,
                 const ur_event_handle_t *phEventWaitList,
-                ur_event_handle_t *phEvent, T &&f) {
-  return withTimingEvent(command_type, hQueue, numEventsInWaitList,
-                         phEventWaitList, phEvent, std::forward<T>(f),
-                         BlockingWithEvent());
-}
-
-template <class T>
-static inline ur_result_t
-withTimingEvent(ur_command_t command_type, ur_queue_handle_t hQueue,
-                uint32_t numEventsInWaitList,
-                const ur_event_handle_t *phEventWaitList,
-                ur_event_handle_t *phEvent, T &&f, bool blocking) {
+                ur_event_handle_t *phEvent, T &&f, bool blocking = true) {
   if (blocking)
     return withTimingEvent(command_type, hQueue, numEventsInWaitList,
                            phEventWaitList, phEvent, std::forward<T>(f),

From 40f7270fff3927e3b1821667fce0337bc0b67417 Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Tue, 29 Apr 2025 14:43:12 +0100
Subject: [PATCH 18/34] [NATIVECPU] async wait in noop copy

---
 unified-runtime/source/adapters/native_cpu/enqueue.cpp | 6 ++++--
 1 file changed, 4 insertions(+), 2 deletions(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index 6e688bca37e5b..3032c7e6e6381 100755
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -445,10 +445,12 @@ static inline ur_result_t doCopy_impl(
     ur_queue_handle_t hQueue, void *DstPtr, const void *SrcPtr, size_t Size,
     uint32_t numEventsInWaitList, const ur_event_handle_t *phEventWaitList,
     ur_event_handle_t *phEvent, ur_command_t command_type, bool blocking) {
-  if (SrcPtr == DstPtr || Size == 0)
+  if (SrcPtr == DstPtr || Size == 0) {
+    bool hasInEvents = numEventsInWaitList && phEventWaitList;
     return withTimingEvent(
         command_type, hQueue, numEventsInWaitList, phEventWaitList, phEvent,
-        []() { return UR_RESULT_SUCCESS; }, BlockingWithEvent());
+        []() { return UR_RESULT_SUCCESS; }, blocking || !hasInEvents);
+  }
 
   return withTimingEvent(
       command_type, hQueue, numEventsInWaitList, phEventWaitList, phEvent,

From bd161bcd058cc623c7ef4629b8ac98f43b4db15b Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Tue, 29 Apr 2025 16:00:49 +0100
Subject: [PATCH 19/34] [NATIVECPU] async membuffer ops

---
 .../source/adapters/native_cpu/enqueue.cpp          | 13 ++++++++-----
 1 file changed, 8 insertions(+), 5 deletions(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index 3032c7e6e6381..1292091649b3b 100755
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -339,7 +339,7 @@ struct NonBlocking {
     auto InEvents =
         native_cpu::getWaitInfo(numEventsInWaitList, phEventWaitList);
     futures.emplace_back(
-        tp.schedule_task([op, InEvents = InEvents.get()](size_t) {
+        tp.schedule_task([op, InEvents = InEvents.get()](size_t) mutable {
           if (InEvents)
             InEvents->wait();
           op();
@@ -395,7 +395,7 @@ UR_APIEXPORT ur_result_t urEnqueueEventsWaitWithBarrierExt(
 
 template <bool IsRead>
 static inline ur_result_t enqueueMemBufferReadWriteRect_impl(
-    ur_queue_handle_t hQueue, ur_mem_handle_t Buff, bool,
+    ur_queue_handle_t hQueue, ur_mem_handle_t Buff, bool blocking,
     ur_rect_offset_t BufferOffset, ur_rect_offset_t HostOffset,
     ur_rect_region_t region, size_t BufferRowPitch, size_t BufferSlicePitch,
     size_t HostRowPitch, size_t HostSlicePitch,
@@ -408,7 +408,9 @@ static inline ur_result_t enqueueMemBufferReadWriteRect_impl(
   else
     command_t = UR_COMMAND_MEM_BUFFER_WRITE_RECT;
   return withTimingEvent(
-      command_t, hQueue, NumEventsInWaitList, phEventWaitList, phEvent, [&]() {
+      command_t, hQueue, NumEventsInWaitList, phEventWaitList, phEvent,
+      [BufferRowPitch, region, BufferSlicePitch, HostRowPitch, HostSlicePitch,
+       BufferOffset, HostOffset, Buff, DstMem]() mutable {
         // TODO: blocking, check other constraints, performance optimizations
         //       More sharing with level_zero where possible
 
@@ -437,7 +439,8 @@ static inline ur_result_t enqueueMemBufferReadWriteRect_impl(
             }
 
         return UR_RESULT_SUCCESS;
-      });
+      },
+      blocking);
 }
 
 template <void *(copy_func)(void *dst, const void *src, size_t) = memmove>
@@ -532,7 +535,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemBufferCopyRect(
     uint32_t numEventsInWaitList, const ur_event_handle_t *phEventWaitList,
     ur_event_handle_t *phEvent) {
   return enqueueMemBufferReadWriteRect_impl<true /*read*/>(
-      hQueue, hBufferSrc, false /*todo: check blocking*/, srcOrigin,
+      hQueue, hBufferSrc, true /*todo: check blocking*/, srcOrigin,
       /*HostOffset*/ dstOrigin, region, srcRowPitch, srcSlicePitch, dstRowPitch,
       dstSlicePitch, hBufferDst->_mem, numEventsInWaitList, phEventWaitList,
       phEvent);

From 870754a2df5d85a89a94739d7acde9bbb83836bf Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Wed, 30 Apr 2025 17:20:59 +0100
Subject: [PATCH 20/34] [NATIVECPU] quick fix for in-order queues

---
 unified-runtime/source/adapters/native_cpu/enqueue.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index 1292091649b3b..e4926d78c20d3 100755
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -358,7 +358,7 @@ withTimingEvent(ur_command_t command_type, ur_queue_handle_t hQueue,
                 uint32_t numEventsInWaitList,
                 const ur_event_handle_t *phEventWaitList,
                 ur_event_handle_t *phEvent, T &&f, bool blocking = true) {
-  if (blocking)
+  if (blocking || hQueue->isInOrder())
     return withTimingEvent(command_type, hQueue, numEventsInWaitList,
                            phEventWaitList, phEvent, std::forward<T>(f),
                            BlockingWithEvent());

From e11f5966df87f502f32f9c5b25adda9d3129a472 Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Wed, 30 Apr 2025 19:21:54 +0100
Subject: [PATCH 21/34] [NATIVECPU] construct state inside thread

---
 .../source/adapters/native_cpu/enqueue.cpp    | 26 ++++++++++---------
 1 file changed, 14 insertions(+), 12 deletions(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index e4926d78c20d3..02b87c1b2ec14 100755
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -72,16 +72,18 @@ getWaitInfo(uint32_t numEventsInWaitList,
 } // namespace
 } // namespace native_cpu
 
-#ifdef NATIVECPU_USE_OCK
-static native_cpu::state getResizedState(const native_cpu::NDRDescT &ndr,
-                                         size_t itemsPerThread) {
+static inline native_cpu::state getResizedState(const native_cpu::NDRDescT &ndr,
+                                                size_t itemsPerThread) {
   native_cpu::state resized_state(
       ndr.GlobalSize[0], ndr.GlobalSize[1], ndr.GlobalSize[2], itemsPerThread,
       ndr.LocalSize[1], ndr.LocalSize[2], ndr.GlobalOffset[0],
       ndr.GlobalOffset[1], ndr.GlobalOffset[2]);
   return resized_state;
 }
-#endif
+
+static inline native_cpu::state getState(const native_cpu::NDRDescT &ndr) {
+  return getResizedState(ndr, ndr.LocalSize[0]);
+}
 
 using IndexT = std::array<size_t, 3>;
 using RangeT = native_cpu::NDRDescT::RangeT;
@@ -152,10 +154,6 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
   auto numWG0 = ndr.GlobalSize[0] / ndr.LocalSize[0];
   auto numWG1 = ndr.GlobalSize[1] / ndr.LocalSize[1];
   auto numWG2 = ndr.GlobalSize[2] / ndr.LocalSize[2];
-  native_cpu::state state(ndr.GlobalSize[0], ndr.GlobalSize[1],
-                          ndr.GlobalSize[2], ndr.LocalSize[0], ndr.LocalSize[1],
-                          ndr.LocalSize[2], ndr.GlobalOffset[0],
-                          ndr.GlobalOffset[1], ndr.GlobalOffset[2]);
   auto event = new ur_event_handle_t_(hQueue, UR_COMMAND_KERNEL_LAUNCH);
   event->tick_start();
 
@@ -166,6 +164,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
   auto InEvents = native_cpu::getWaitInfo(numEventsInWaitList, phEventWaitList);
 
 #ifndef NATIVECPU_USE_OCK
+  native_cpu::state state = getState(ndr);
   urEventWait(numEventsInWaitList, phEventWaitList);
   for (unsigned g2 = 0; g2 < numWG2; g2++) {
     for (unsigned g1 = 0; g1 < numWG1; g1++) {
@@ -221,12 +220,13 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
       // Peel the remaining work items. Since the local size is 1, we iterate
       // over the work groups.
       futures.emplace_back(tp.schedule_task(
-          [state, &kernel = *kernel, start_wg0_remainder, numWG0, numWG1,
-           numWG2, InEvents = InEvents.get()](size_t) mutable {
+          [ndr, &kernel = *kernel, start_wg0_remainder, numWG0, numWG1, numWG2,
+           InEvents = InEvents.get()](size_t) mutable {
             IndexT first = {start_wg0_remainder, 0, 0};
             IndexT last = {numWG0, numWG1, numWG2};
             if (InEvents)
               InEvents->wait();
+            native_cpu::state state = getState(ndr);
             execute_range(state, kernel, first, last);
           }));
     }
@@ -252,10 +252,11 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
         wg_start += groupsPerThread[dim];
         last[dim] = wg_start;
         futures.emplace_back(tp.schedule_task(
-            [state, numParallelThreads, &kernel = *kernel, first, last,
+            [ndr, numParallelThreads, &kernel = *kernel, first, last,
              InEvents = InEvents.get()](size_t threadId) mutable {
               if (InEvents)
                 InEvents->wait();
+              native_cpu::state state = getState(ndr);
               execute_range(state, kernel,
                             kernel.getArgs(numParallelThreads, threadId), first,
                             last);
@@ -266,10 +267,11 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
       first[dim] = wg_start;
       last[dim] = numWG[dim];
       futures.emplace_back(tp.schedule_task(
-          [state, numParallelThreads, &kernel = *kernel, first, last,
+          [ndr, numParallelThreads, &kernel = *kernel, first, last,
            InEvents = InEvents.get()](size_t threadId) mutable {
             if (InEvents)
               InEvents->wait();
+            native_cpu::state state = getState(ndr);
             execute_range(state, kernel,
                           kernel.getArgs(numParallelThreads, threadId), first,
                           last);

From b4069d1c8d0ff8d4c58b449dfd0572c17e657a82 Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Thu, 1 May 2025 14:00:19 +0100
Subject: [PATCH 22/34] [NATIVECPU] update comments

---
 unified-runtime/source/adapters/native_cpu/enqueue.cpp | 7 ++++---
 1 file changed, 4 insertions(+), 3 deletions(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index 02b87c1b2ec14..7b36f8985a0c9 100755
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -413,7 +413,7 @@ static inline ur_result_t enqueueMemBufferReadWriteRect_impl(
       command_t, hQueue, NumEventsInWaitList, phEventWaitList, phEvent,
       [BufferRowPitch, region, BufferSlicePitch, HostRowPitch, HostSlicePitch,
        BufferOffset, HostOffset, Buff, DstMem]() mutable {
-        // TODO: blocking, check other constraints, performance optimizations
+        // TODO: check other constraints, performance optimizations
         //       More sharing with level_zero where possible
 
         if (BufferRowPitch == 0)
@@ -526,7 +526,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemBufferCopy(
   void *DstPtr = hBufferDst->_mem + dstOffset;
   return doCopy_impl(hQueue, DstPtr, SrcPtr, size, numEventsInWaitList,
                      phEventWaitList, phEvent, UR_COMMAND_MEM_BUFFER_COPY,
-                     true /*TODO: check blocking*/);
+                     true /*TODO: check false for non-blocking*/);
 }
 
 UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemBufferCopyRect(
@@ -537,7 +537,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemBufferCopyRect(
     uint32_t numEventsInWaitList, const ur_event_handle_t *phEventWaitList,
     ur_event_handle_t *phEvent) {
   return enqueueMemBufferReadWriteRect_impl<true /*read*/>(
-      hQueue, hBufferSrc, true /*todo: check blocking*/, srcOrigin,
+      hQueue, hBufferSrc, true /*todo: check false for non-blocking*/,
+      srcOrigin,
       /*HostOffset*/ dstOrigin, region, srcRowPitch, srcSlicePitch, dstRowPitch,
       dstSlicePitch, hBufferDst->_mem, numEventsInWaitList, phEventWaitList,
       phEvent);

From dfc67d8a7342f0906cf7effedeac35a96bc08e31 Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Fri, 2 May 2025 14:39:31 +0100
Subject: [PATCH 23/34] [NATIVECPU] removed nullptr check for pHEventWaitList

---
 unified-runtime/source/adapters/native_cpu/enqueue.cpp | 7 +++----
 1 file changed, 3 insertions(+), 4 deletions(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index 7b36f8985a0c9..c14a97954d8e8 100755
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -63,10 +63,9 @@ struct WaitInfo {
 inline static std::unique_ptr<WaitInfo>
 getWaitInfo(uint32_t numEventsInWaitList,
             const ur_event_handle_t *phEventWaitList) {
-  return (numEventsInWaitList && phEventWaitList)
-             ? std::make_unique<native_cpu::WaitInfo>(numEventsInWaitList,
-                                                      phEventWaitList)
-             : nullptr;
+  return (numEventsInWaitList) ? std::make_unique<native_cpu::WaitInfo>(
+                                     numEventsInWaitList, phEventWaitList)
+                               : nullptr;
 }
 
 } // namespace

From 3074b16b826e738ba98e6b86e05c9aa37df0484b Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Tue, 6 May 2025 14:48:06 +0100
Subject: [PATCH 24/34] [NATIVECPU] removed unneeded mutable

---
 .../source/adapters/native_cpu/enqueue.cpp         | 14 +++++++-------
 1 file changed, 7 insertions(+), 7 deletions(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index c14a97954d8e8..73c19b86b7a1e 100755
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -218,9 +218,9 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
     if (start_wg0_remainder < numWG0) {
       // Peel the remaining work items. Since the local size is 1, we iterate
       // over the work groups.
-      futures.emplace_back(tp.schedule_task(
-          [ndr, &kernel = *kernel, start_wg0_remainder, numWG0, numWG1, numWG2,
-           InEvents = InEvents.get()](size_t) mutable {
+      futures.emplace_back(
+          tp.schedule_task([ndr, &kernel = *kernel, start_wg0_remainder, numWG0,
+                            numWG1, numWG2, InEvents = InEvents.get()](size_t) {
             IndexT first = {start_wg0_remainder, 0, 0};
             IndexT last = {numWG0, numWG1, numWG2};
             if (InEvents)
@@ -252,7 +252,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
         last[dim] = wg_start;
         futures.emplace_back(tp.schedule_task(
             [ndr, numParallelThreads, &kernel = *kernel, first, last,
-             InEvents = InEvents.get()](size_t threadId) mutable {
+             InEvents = InEvents.get()](size_t threadId) {
               if (InEvents)
                 InEvents->wait();
               native_cpu::state state = getState(ndr);
@@ -265,9 +265,9 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
     if (wg_start < numWG[dim]) {
       first[dim] = wg_start;
       last[dim] = numWG[dim];
-      futures.emplace_back(tp.schedule_task(
-          [ndr, numParallelThreads, &kernel = *kernel, first, last,
-           InEvents = InEvents.get()](size_t threadId) mutable {
+      futures.emplace_back(
+          tp.schedule_task([ndr, numParallelThreads, &kernel = *kernel, first,
+                            last, InEvents = InEvents.get()](size_t threadId) {
             if (InEvents)
               InEvents->wait();
             native_cpu::state state = getState(ndr);

From 070f0cfe36743c31811746a1202af60f18df59ce Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Tue, 6 May 2025 15:26:39 +0100
Subject: [PATCH 25/34] [NATIVECPU] moved lambda code from
 enqueueMemBufferReadWriteRect_impl into seperate function to be able to
 remove mutable

---
 .../source/adapters/native_cpu/enqueue.cpp    | 68 +++++++++++--------
 1 file changed, 39 insertions(+), 29 deletions(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index 73c19b86b7a1e..3be5c115dc594 100755
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -340,7 +340,7 @@ struct NonBlocking {
     auto InEvents =
         native_cpu::getWaitInfo(numEventsInWaitList, phEventWaitList);
     futures.emplace_back(
-        tp.schedule_task([op, InEvents = InEvents.get()](size_t) mutable {
+        tp.schedule_task([op, InEvents = InEvents.get()](size_t) {
           if (InEvents)
             InEvents->wait();
           op();
@@ -394,6 +394,40 @@ UR_APIEXPORT ur_result_t urEnqueueEventsWaitWithBarrierExt(
                                         phEventWaitList, phEvent);
 }
 
+template <bool IsRead>
+static inline void MemBufferReadWriteRect_impl(
+    ur_mem_handle_t Buff, ur_rect_offset_t BufferOffset,
+    ur_rect_offset_t HostOffset, ur_rect_region_t region, size_t BufferRowPitch,
+    size_t BufferSlicePitch, size_t HostRowPitch, size_t HostSlicePitch,
+    typename std::conditional<IsRead, void *, const void *>::type DstMem) {
+  // TODO: check other constraints, performance optimizations
+  //       More sharing with level_zero where possible
+
+  if (BufferRowPitch == 0)
+    BufferRowPitch = region.width;
+  if (BufferSlicePitch == 0)
+    BufferSlicePitch = BufferRowPitch * region.height;
+  if (HostRowPitch == 0)
+    HostRowPitch = region.width;
+  if (HostSlicePitch == 0)
+    HostSlicePitch = HostRowPitch * region.height;
+  for (size_t w = 0; w < region.width; w++)
+    for (size_t h = 0; h < region.height; h++)
+      for (size_t d = 0; d < region.depth; d++) {
+        size_t buff_orign = (d + BufferOffset.z) * BufferSlicePitch +
+                            (h + BufferOffset.y) * BufferRowPitch + w +
+                            BufferOffset.x;
+        size_t host_origin = (d + HostOffset.z) * HostSlicePitch +
+                             (h + HostOffset.y) * HostRowPitch + w +
+                             HostOffset.x;
+        int8_t &buff_mem = ur_cast<int8_t *>(Buff->_mem)[buff_orign];
+        if constexpr (IsRead)
+          ur_cast<int8_t *>(DstMem)[host_origin] = buff_mem;
+        else
+          buff_mem = ur_cast<const int8_t *>(DstMem)[host_origin];
+      }
+}
+
 template <bool IsRead>
 static inline ur_result_t enqueueMemBufferReadWriteRect_impl(
     ur_queue_handle_t hQueue, ur_mem_handle_t Buff, bool blocking,
@@ -411,34 +445,10 @@ static inline ur_result_t enqueueMemBufferReadWriteRect_impl(
   return withTimingEvent(
       command_t, hQueue, NumEventsInWaitList, phEventWaitList, phEvent,
       [BufferRowPitch, region, BufferSlicePitch, HostRowPitch, HostSlicePitch,
-       BufferOffset, HostOffset, Buff, DstMem]() mutable {
-        // TODO: check other constraints, performance optimizations
-        //       More sharing with level_zero where possible
-
-        if (BufferRowPitch == 0)
-          BufferRowPitch = region.width;
-        if (BufferSlicePitch == 0)
-          BufferSlicePitch = BufferRowPitch * region.height;
-        if (HostRowPitch == 0)
-          HostRowPitch = region.width;
-        if (HostSlicePitch == 0)
-          HostSlicePitch = HostRowPitch * region.height;
-        for (size_t w = 0; w < region.width; w++)
-          for (size_t h = 0; h < region.height; h++)
-            for (size_t d = 0; d < region.depth; d++) {
-              size_t buff_orign = (d + BufferOffset.z) * BufferSlicePitch +
-                                  (h + BufferOffset.y) * BufferRowPitch + w +
-                                  BufferOffset.x;
-              size_t host_origin = (d + HostOffset.z) * HostSlicePitch +
-                                   (h + HostOffset.y) * HostRowPitch + w +
-                                   HostOffset.x;
-              int8_t &buff_mem = ur_cast<int8_t *>(Buff->_mem)[buff_orign];
-              if constexpr (IsRead)
-                ur_cast<int8_t *>(DstMem)[host_origin] = buff_mem;
-              else
-                buff_mem = ur_cast<const int8_t *>(DstMem)[host_origin];
-            }
-
+       BufferOffset, HostOffset, Buff, DstMem]() {
+        MemBufferReadWriteRect_impl<IsRead>(
+            Buff, BufferOffset, HostOffset, region, BufferRowPitch,
+            BufferSlicePitch, HostRowPitch, HostSlicePitch, DstMem);
         return UR_RESULT_SUCCESS;
       },
       blocking);

From 3207ffacd7674ce1966bfb21ca689a522ffa83cd Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Wed, 7 May 2025 17:03:47 +0100
Subject: [PATCH 26/34] [NATIVECPU] simplified event generation

---
 .../source/adapters/native_cpu/enqueue.cpp    | 62 +++++--------------
 1 file changed, 14 insertions(+), 48 deletions(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index 3be5c115dc594..7dffab6f3e681 100755
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -298,74 +298,40 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
   return UR_RESULT_SUCCESS;
 }
 
-template <class T, class I>
+template <class T>
 static inline ur_result_t
 withTimingEvent(ur_command_t command_type, ur_queue_handle_t hQueue,
                 uint32_t numEventsInWaitList,
                 const ur_event_handle_t *phEventWaitList,
-                ur_event_handle_t *phEvent, T &&f, I &&inv) {
+                ur_event_handle_t *phEvent, T &&f, bool blocking = true) {
   if (phEvent) {
     ur_event_handle_t event = new ur_event_handle_t_(hQueue, command_type);
-    event->tick_start();
-    ur_result_t result = inv(std::forward<T>(f), event, hQueue,
-                             numEventsInWaitList, phEventWaitList);
     *phEvent = event;
-    return result;
-  }
-  urEventWait(numEventsInWaitList, phEventWaitList);
-  ur_result_t result = f();
-  return result;
-}
-
-namespace {
-struct BlockingWithEvent {
-  template <class T>
-  ur_result_t operator()(T &&op, ur_event_handle_t event, ur_queue_handle_t,
-                         uint32_t numEventsInWaitList,
-                         const ur_event_handle_t *phEventWaitList) const {
-    urEventWait(numEventsInWaitList, phEventWaitList);
-    ur_result_t result = op();
-    event->tick_end();
-    return result;
-  }
-};
-
-struct NonBlocking {
-  template <class T>
-  ur_result_t operator()(T &&op, ur_event_handle_t event,
-                         ur_queue_handle_t hQueue, uint32_t numEventsInWaitList,
-                         const ur_event_handle_t *phEventWaitList) const {
+    event->tick_start();
+    if (blocking || hQueue->isInOrder()) {
+      urEventWait(numEventsInWaitList, phEventWaitList);
+      ur_result_t result = f();
+      event->tick_end();
+      return result;
+    }
     auto &tp = hQueue->getDevice()->tp;
     std::vector<std::future<void>> futures;
     auto InEvents =
         native_cpu::getWaitInfo(numEventsInWaitList, phEventWaitList);
     futures.emplace_back(
-        tp.schedule_task([op, InEvents = InEvents.get()](size_t) {
+        tp.schedule_task([f, InEvents = InEvents.get()](size_t) {
           if (InEvents)
             InEvents->wait();
-          op();
+          f();
         }));
     event->set_futures(futures);
     event->set_callback(
         [event, InEvents = std::move(InEvents)]() { event->tick_end(); });
     return UR_RESULT_SUCCESS;
   }
-};
-} // namespace
-
-template <class T>
-static inline ur_result_t
-withTimingEvent(ur_command_t command_type, ur_queue_handle_t hQueue,
-                uint32_t numEventsInWaitList,
-                const ur_event_handle_t *phEventWaitList,
-                ur_event_handle_t *phEvent, T &&f, bool blocking = true) {
-  if (blocking || hQueue->isInOrder())
-    return withTimingEvent(command_type, hQueue, numEventsInWaitList,
-                           phEventWaitList, phEvent, std::forward<T>(f),
-                           BlockingWithEvent());
-  return withTimingEvent(command_type, hQueue, numEventsInWaitList,
-                         phEventWaitList, phEvent, std::forward<T>(f),
-                         NonBlocking());
+  urEventWait(numEventsInWaitList, phEventWaitList);
+  ur_result_t result = f();
+  return result;
 }
 
 UR_APIEXPORT ur_result_t UR_APICALL urEnqueueEventsWait(

From 941932bd307323be10c6c746cba3a389cdac0dec Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Fri, 9 May 2025 09:06:45 +0100
Subject: [PATCH 27/34] [NATIVECPU] replaced function pointer template
 parameter

---
 .../source/adapters/native_cpu/enqueue.cpp         | 14 +++++++++-----
 1 file changed, 9 insertions(+), 5 deletions(-)
 mode change 100755 => 100644 unified-runtime/source/adapters/native_cpu/enqueue.cpp

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
old mode 100755
new mode 100644
index 7dffab6f3e681..bf3766d486519
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -420,7 +420,7 @@ static inline ur_result_t enqueueMemBufferReadWriteRect_impl(
       blocking);
 }
 
-template <void *(copy_func)(void *dst, const void *src, size_t) = memmove>
+template <bool AllowPartialOverlap = true>
 static inline ur_result_t doCopy_impl(
     ur_queue_handle_t hQueue, void *DstPtr, const void *SrcPtr, size_t Size,
     uint32_t numEventsInWaitList, const ur_event_handle_t *phEventWaitList,
@@ -435,7 +435,11 @@ static inline ur_result_t doCopy_impl(
   return withTimingEvent(
       command_type, hQueue, numEventsInWaitList, phEventWaitList, phEvent,
       [DstPtr, SrcPtr, Size]() {
-        copy_func(DstPtr, SrcPtr, Size);
+        if constexpr (AllowPartialOverlap) {
+          memmove(DstPtr, SrcPtr, Size);
+        } else {
+          memcpy(DstPtr, SrcPtr, Size);
+        }
         return UR_RESULT_SUCCESS;
       },
       blocking);
@@ -662,9 +666,9 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMMemcpy(
   UR_ASSERT(pDst, UR_RESULT_ERROR_INVALID_NULL_POINTER);
   UR_ASSERT(pSrc, UR_RESULT_ERROR_INVALID_NULL_POINTER);
 
-  return doCopy_impl<memcpy>(hQueue, pDst, pSrc, size, numEventsInWaitList,
-                             phEventWaitList, phEvent, UR_COMMAND_USM_MEMCPY,
-                             blocking);
+  return doCopy_impl<false /*use memcpy*/>(
+      hQueue, pDst, pSrc, size, numEventsInWaitList, phEventWaitList, phEvent,
+      UR_COMMAND_USM_MEMCPY, blocking);
 }
 
 UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMPrefetch(

From 153277991cce79687688a91d956764da235fae07 Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Tue, 13 May 2025 16:43:01 +0100
Subject: [PATCH 28/34] [NATIVECPU] simplified WaitInfo

---
 .../source/adapters/native_cpu/enqueue.cpp    | 76 +++++++++----------
 1 file changed, 38 insertions(+), 38 deletions(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index 0146924f344ab..526a7a8ed33a0 100644
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -52,20 +52,27 @@ struct NDRDescT {
 };
 
 namespace {
-struct WaitInfo {
-  std::vector<ur_event_handle_t> events;
+class WaitInfo {
+  std::vector<ur_event_handle_t> *const events;
   static_assert(std::is_pointer_v<ur_event_handle_t>);
+
+public:
   WaitInfo(uint32_t numEvents, const ur_event_handle_t *WaitList)
-      : events(WaitList, WaitList + numEvents) {}
-  void wait() const { urEventWait(events.size(), events.data()); }
+      : events(numEvents ? new std::vector<ur_event_handle_t>(
+                               WaitList, WaitList + numEvents)
+                         : nullptr) {}
+  void wait() const {
+    if (events)
+      urEventWait(events->size(), events->data());
+  }
+  std::unique_ptr<std::vector<ur_event_handle_t>> getUniquePtr() {
+    return std::unique_ptr<std::vector<ur_event_handle_t>>(events);
+  }
 };
 
-inline static std::unique_ptr<WaitInfo>
-getWaitInfo(uint32_t numEventsInWaitList,
-            const ur_event_handle_t *phEventWaitList) {
-  return (numEventsInWaitList) ? std::make_unique<native_cpu::WaitInfo>(
-                                     numEventsInWaitList, phEventWaitList)
-                               : nullptr;
+inline static WaitInfo getWaitInfo(uint32_t numEventsInWaitList,
+                                   const ur_event_handle_t *phEventWaitList) {
+  return native_cpu::WaitInfo(numEventsInWaitList, phEventWaitList);
 }
 
 } // namespace
@@ -203,15 +210,13 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
     for (size_t t = 0; t < numParallelThreads;) {
       IndexT first = {t, 0, 0};
       IndexT last = {++t, numWG1, numWG2};
-      futures.emplace_back(
-          tp.schedule_task([ndr, itemsPerThread, &kernel = *kernel, first, last,
-                            InEvents = InEvents.get()](size_t) {
-            native_cpu::state resized_state =
-                getResizedState(ndr, itemsPerThread);
-            if (InEvents)
-              InEvents->wait();
-            execute_range(resized_state, kernel, first, last);
-          }));
+      futures.emplace_back(tp.schedule_task([ndr, itemsPerThread,
+                                             &kernel = *kernel, first, last,
+                                             InEvents](size_t) {
+        native_cpu::state resized_state = getResizedState(ndr, itemsPerThread);
+        InEvents.wait();
+        execute_range(resized_state, kernel, first, last);
+      }));
     }
 
     size_t start_wg0_remainder = new_num_work_groups_0 * itemsPerThread;
@@ -220,11 +225,10 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
       // over the work groups.
       futures.emplace_back(
           tp.schedule_task([ndr, &kernel = *kernel, start_wg0_remainder, numWG0,
-                            numWG1, numWG2, InEvents = InEvents.get()](size_t) {
+                            numWG1, numWG2, InEvents](size_t) {
             IndexT first = {start_wg0_remainder, 0, 0};
             IndexT last = {numWG0, numWG1, numWG2};
-            if (InEvents)
-              InEvents->wait();
+            InEvents.wait();
             native_cpu::state state = getState(ndr);
             execute_range(state, kernel, first, last);
           }));
@@ -250,11 +254,10 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
         first[dim] = wg_start;
         wg_start += groupsPerThread[dim];
         last[dim] = wg_start;
-        futures.emplace_back(tp.schedule_task(
-            [ndr, numParallelThreads, &kernel = *kernel, first, last,
-             InEvents = InEvents.get()](size_t threadId) {
-              if (InEvents)
-                InEvents->wait();
+        futures.emplace_back(
+            tp.schedule_task([ndr, numParallelThreads, &kernel = *kernel, first,
+                              last, InEvents](size_t threadId) {
+              InEvents.wait();
               native_cpu::state state = getState(ndr);
               execute_range(state, kernel,
                             kernel.getArgs(numParallelThreads, threadId), first,
@@ -267,9 +270,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
       last[dim] = numWG[dim];
       futures.emplace_back(
           tp.schedule_task([ndr, numParallelThreads, &kernel = *kernel, first,
-                            last, InEvents = InEvents.get()](size_t threadId) {
-            if (InEvents)
-              InEvents->wait();
+                            last, InEvents](size_t threadId) {
+            InEvents.wait();
             native_cpu::state state = getState(ndr);
             execute_range(state, kernel,
                           kernel.getArgs(numParallelThreads, threadId), first,
@@ -285,7 +287,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
     *phEvent = event;
   }
   event->set_callback([kernel = std::move(kernel), hKernel, event,
-                       InEvents = std::move(InEvents)]() {
+                       InEvents = InEvents.getUniquePtr()]() {
     event->tick_end();
     // TODO: avoid calling clear() here.
     hKernel->_localArgInfo.clear();
@@ -318,15 +320,13 @@ withTimingEvent(ur_command_t command_type, ur_queue_handle_t hQueue,
     std::vector<std::future<void>> futures;
     auto InEvents =
         native_cpu::getWaitInfo(numEventsInWaitList, phEventWaitList);
-    futures.emplace_back(
-        tp.schedule_task([f, InEvents = InEvents.get()](size_t) {
-          if (InEvents)
-            InEvents->wait();
-          f();
-        }));
+    futures.emplace_back(tp.schedule_task([f, InEvents](size_t) {
+      InEvents.wait();
+      f();
+    }));
     event->set_futures(futures);
     event->set_callback(
-        [event, InEvents = std::move(InEvents)]() { event->tick_end(); });
+        [event, InEvents = InEvents.getUniquePtr()]() { event->tick_end(); });
     return UR_RESULT_SUCCESS;
   }
   urEventWait(numEventsInWaitList, phEventWaitList);

From 6fcea0f0d655849d4881f70e339a506eba178cee Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Wed, 4 Jun 2025 17:16:39 +0100
Subject: [PATCH 29/34] [NATIVECPU] launch ranges with number of work items
 that is multiple of vector width

---
 .../source/adapters/native_cpu/enqueue.cpp    | 37 +++++++++++++++----
 1 file changed, 29 insertions(+), 8 deletions(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index 57df579b88086..bff27211596f8 100644
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -197,8 +197,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
 #else
   bool isLocalSizeOne =
       ndr.LocalSize[0] == 1 && ndr.LocalSize[1] == 1 && ndr.LocalSize[2] == 1;
-  if (isLocalSizeOne && ndr.GlobalSize[0] > numParallelThreads &&
-      !kernel->hasLocalArgs()) {
+  if (isLocalSizeOne && !kernel->hasLocalArgs()) {
     // If the local size is one, we make the assumption that we are running a
     // parallel_for over a sycl::range.
     // Todo: we could add more compiler checks and
@@ -212,10 +211,33 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
     // divide the global range by the number of threads, set that as the local
     // size and peel everything else.
 
+    // The number of items per kernel invocation should ideally be at least a
+    // multiple of the applied vector width, which we currently assume to be 8.
+    // TODO: Encode this and other kernel capabilities in the binary so we can
+    // use actual values to efficiently enqueue kernels instead of relying on
+    // assumptions.
+    const size_t itemsPerKernelInvocation = 8;
+
     size_t new_num_work_groups_0 = numParallelThreads;
     size_t itemsPerThread = ndr.GlobalSize[0] / numParallelThreads;
+    if (itemsPerThread < itemsPerKernelInvocation) {
+      if (itemsPerKernelInvocation <= numWG0)
+        itemsPerThread = itemsPerKernelInvocation;
+      else if (itemsPerThread == 0)
+        itemsPerThread = numWG0;
+    } else if (itemsPerThread > itemsPerKernelInvocation) {
+      // Launch kernel with number of items that is the next multiple of the
+      // vector width.
+      const size_t nextMult = (itemsPerThread + itemsPerKernelInvocation - 1) /
+                              itemsPerKernelInvocation *
+                              itemsPerKernelInvocation;
+      if (nextMult < numWG0)
+        itemsPerThread = nextMult;
+    }
 
-    for (size_t t = 0; t < numParallelThreads;) {
+    size_t wg0_index = 0;
+    for (size_t t = 0; (wg0_index + itemsPerThread) <= numWG0;
+         wg0_index += itemsPerThread) {
       IndexT first = {t, 0, 0};
       IndexT last = {++t, numWG1, numWG2};
       futures.emplace_back(tp.schedule_task([ndr, itemsPerThread,
@@ -227,14 +249,13 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
       }));
     }
 
-    size_t start_wg0_remainder = new_num_work_groups_0 * itemsPerThread;
-    if (start_wg0_remainder < numWG0) {
+    if (wg0_index < numWG0) {
       // Peel the remaining work items. Since the local size is 1, we iterate
       // over the work groups.
       futures.emplace_back(
-          tp.schedule_task([ndr, &kernel = *kernel, start_wg0_remainder, numWG0,
-                            numWG1, numWG2, InEvents](size_t) {
-            IndexT first = {start_wg0_remainder, 0, 0};
+          tp.schedule_task([ndr, &kernel = *kernel, wg0_index, numWG0, numWG1,
+                            numWG2, InEvents](size_t) {
+            IndexT first = {wg0_index, 0, 0};
             IndexT last = {numWG0, numWG1, numWG2};
             InEvents.wait();
             native_cpu::state state = getState(ndr);

From ddb908f5e1a6e2322f6061efb7e22b08a347847a Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Thu, 5 Jun 2025 13:20:46 +0100
Subject: [PATCH 30/34] [NATIVECPU] removed unused local

---
 unified-runtime/source/adapters/native_cpu/enqueue.cpp | 1 -
 1 file changed, 1 deletion(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index bff27211596f8..3b5d1b8433e34 100644
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -218,7 +218,6 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
     // assumptions.
     const size_t itemsPerKernelInvocation = 8;
 
-    size_t new_num_work_groups_0 = numParallelThreads;
     size_t itemsPerThread = ndr.GlobalSize[0] / numParallelThreads;
     if (itemsPerThread < itemsPerKernelInvocation) {
       if (itemsPerKernelInvocation <= numWG0)

From 1d629032915211e8e16a08ba948eba94e27a0deb Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Mon, 9 Jun 2025 18:44:44 +0100
Subject: [PATCH 31/34] [NATIVECPU] removed reference captures in enqueue
 lambdas

---
 .../source/adapters/native_cpu/enqueue.cpp      | 17 ++++++++---------
 1 file changed, 8 insertions(+), 9 deletions(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index 3b5d1b8433e34..8fff9391fcb0e 100644
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -369,7 +369,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueEventsWait(
   // TODO: the wait here should be async
   return withTimingEvent(UR_COMMAND_EVENTS_WAIT, hQueue, numEventsInWaitList,
                          phEventWaitList, phEvent,
-                         [&]() { return UR_RESULT_SUCCESS; });
+                         []() { return UR_RESULT_SUCCESS; });
 }
 
 UR_APIEXPORT ur_result_t UR_APICALL urEnqueueEventsWaitWithBarrier(
@@ -377,7 +377,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueEventsWaitWithBarrier(
     const ur_event_handle_t *phEventWaitList, ur_event_handle_t *phEvent) {
   return withTimingEvent(UR_COMMAND_EVENTS_WAIT_WITH_BARRIER, hQueue,
                          numEventsInWaitList, phEventWaitList, phEvent,
-                         [&]() { return UR_RESULT_SUCCESS; });
+                         []() { return UR_RESULT_SUCCESS; });
 }
 
 UR_APIEXPORT ur_result_t urEnqueueEventsWaitWithBarrierExt(
@@ -556,12 +556,10 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemBufferFill(
     size_t patternSize, size_t offset, size_t size,
     uint32_t numEventsInWaitList, const ur_event_handle_t *phEventWaitList,
     ur_event_handle_t *phEvent) {
-
+  UR_ASSERT(hQueue, UR_RESULT_ERROR_INVALID_NULL_HANDLE);
   return withTimingEvent(
       UR_COMMAND_MEM_BUFFER_FILL, hQueue, numEventsInWaitList, phEventWaitList,
-      phEvent, [&]() {
-        UR_ASSERT(hQueue, UR_RESULT_ERROR_INVALID_NULL_HANDLE);
-
+      phEvent, [hBuffer, offset, size, patternSize, pPattern]() {
         // TODO: error checking
         // TODO: handle async
         void *startingPtr = hBuffer->_mem + offset;
@@ -615,7 +613,8 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemBufferMap(
     ur_event_handle_t *phEvent, void **ppRetMap) {
 
   return withTimingEvent(UR_COMMAND_MEM_BUFFER_MAP, hQueue, numEventsInWaitList,
-                         phEventWaitList, phEvent, [&]() {
+                         phEventWaitList, phEvent,
+                         [ppRetMap, hBuffer, offset]() {
                            *ppRetMap = hBuffer->_mem + offset;
                            return UR_RESULT_SUCCESS;
                          });
@@ -627,7 +626,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueMemUnmap(
     ur_event_handle_t *phEvent) {
   return withTimingEvent(UR_COMMAND_MEM_UNMAP, hQueue, numEventsInWaitList,
                          phEventWaitList, phEvent,
-                         [&]() { return UR_RESULT_SUCCESS; });
+                         []() { return UR_RESULT_SUCCESS; });
 }
 
 UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMFill(
@@ -636,7 +635,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueUSMFill(
     const ur_event_handle_t *phEventWaitList, ur_event_handle_t *phEvent) {
   return withTimingEvent(
       UR_COMMAND_USM_FILL, hQueue, numEventsInWaitList, phEventWaitList,
-      phEvent, [&]() {
+      phEvent, [ptr, pPattern, patternSize, size]() {
         UR_ASSERT(ptr, UR_RESULT_ERROR_INVALID_NULL_POINTER);
         UR_ASSERT(pPattern, UR_RESULT_ERROR_INVALID_NULL_POINTER);
         UR_ASSERT(patternSize != 0, UR_RESULT_ERROR_INVALID_SIZE)

From 5e0b99db40a6b28dea8ac68449a06735a8b1b605 Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Thu, 24 Jul 2025 15:06:34 +0100
Subject: [PATCH 32/34] [NATIVECPU] removed unneeded capture

---
 unified-runtime/source/adapters/native_cpu/enqueue.cpp | 9 +++------
 1 file changed, 3 insertions(+), 6 deletions(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index f524b6c3f85ce..156993f89f3a0 100644
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -180,9 +180,6 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
     futures.emplace_back(
         tp.schedule_task([ndr, InEvents, &kernel = *kernel, rangeStart,
                           rangeEnd = rangeEnd[3], numWG0, numWG1,
-#ifndef NATIVECPU_USE_OCK
-                          localSize = ndr.LocalSize,
-#endif
                           numParallelThreads](size_t threadId) mutable {
           auto state = getState(ndr);
           InEvents.wait();
@@ -194,9 +191,9 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
             kernel._subhandler(
                 kernel.getArgs(numParallelThreads, threadId).data(), &state);
 #else
-            for (size_t local2 = 0; local2 < localSize[2]; ++local2) {
-              for (size_t local1 = 0; local1 < localSize[1]; ++local1) {
-                for (size_t local0 = 0; local0 < localSize[0]; ++local0) {
+            for (size_t local2 = 0; local2 < ndr.LocalSize[2]; ++local2) {
+              for (size_t local1 = 0; local1 < ndr.LocalSize[1]; ++local1) {
+                for (size_t local0 = 0; local0 < ndr.LocalSize[0]; ++local0) {
                   state.update(g0, g1, g2, local0, local1, local2);
                   kernel._subhandler(
                       kernel.getArgs(numParallelThreads, threadId).data(),

From f05bba10e28ef28ef443e85cb598e62d9f929b2c Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Thu, 24 Jul 2025 15:15:39 +0100
Subject: [PATCH 33/34] [NATIVECPU] removed mutable from task lambda

---
 unified-runtime/source/adapters/native_cpu/enqueue.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index 156993f89f3a0..f98ab905c47c4 100644
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -180,7 +180,7 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
     futures.emplace_back(
         tp.schedule_task([ndr, InEvents, &kernel = *kernel, rangeStart,
                           rangeEnd = rangeEnd[3], numWG0, numWG1,
-                          numParallelThreads](size_t threadId) mutable {
+                          numParallelThreads](size_t threadId) {
           auto state = getState(ndr);
           InEvents.wait();
           for (size_t g0 = rangeStart[0], g1 = rangeStart[1],

From 58ffb890266e71a4636a76c23552af12a93287c1 Mon Sep 17 00:00:00 2001
From: Uwe Dolinsky <uwe@codeplay.com>
Date: Thu, 24 Jul 2025 16:31:59 +0100
Subject: [PATCH 34/34] [NATIVECPU] clang-format

---
 unified-runtime/source/adapters/native_cpu/enqueue.cpp | 7 +++----
 1 file changed, 3 insertions(+), 4 deletions(-)

diff --git a/unified-runtime/source/adapters/native_cpu/enqueue.cpp b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
index f98ab905c47c4..86da10bbffef7 100644
--- a/unified-runtime/source/adapters/native_cpu/enqueue.cpp
+++ b/unified-runtime/source/adapters/native_cpu/enqueue.cpp
@@ -177,10 +177,9 @@ UR_APIEXPORT ur_result_t UR_APICALL urEnqueueKernelLaunch(
     rangeEnd[0] = rangeEnd[3] % numWG0;
     rangeEnd[1] = (rangeEnd[3] / numWG0) % numWG1;
     rangeEnd[2] = rangeEnd[3] / (numWG0 * numWG1);
-    futures.emplace_back(
-        tp.schedule_task([ndr, InEvents, &kernel = *kernel, rangeStart,
-                          rangeEnd = rangeEnd[3], numWG0, numWG1,
-                          numParallelThreads](size_t threadId) {
+    futures.emplace_back(tp.schedule_task(
+        [ndr, InEvents, &kernel = *kernel, rangeStart, rangeEnd = rangeEnd[3],
+         numWG0, numWG1, numParallelThreads](size_t threadId) {
           auto state = getState(ndr);
           InEvents.wait();
           for (size_t g0 = rangeStart[0], g1 = rangeStart[1],