[UR][CUDA][HIP] Fix command buffer update with shared kernel handles (#19479)

- Fix issue where updating multiple nodes with the same UR kernel handle
would give incorrect results due to arg caching.
- Add SYCL E2E test based on the example in the issue that reported
this.

Addresses issue reported in #19450

Author: Bensuo

sycl/test-e2e/Graph/Update/independent_kernels_update.cpp

@@ -0,0 +1,58 @@
+// RUN: %{build} -o %t.out
+// RUN: %{run} %t.out
+
+// REQUIRES: aspect-usm_shared_allocations
+
+// Modified example from issue #19450 which identified an issue with updating
+// multiple kernel nodes which share the same kernel.
+
+#include "../graph_common.hpp"
+
+int main() {
+  sycl::queue q;
+
+  static constexpr size_t R = 10;
+  static constexpr size_t I = 5;
+  int *output = sycl::malloc_shared<int>(I, q);
+  std::fill(output, output + I, 0);
+
+  std::unique_ptr<sycl::ext::oneapi::experimental::command_graph<
+      sycl::ext::oneapi::experimental::graph_state::executable>>
+      graph;
+  for (int r = 0; r < R; ++r) {
+
+    sycl::ext::oneapi::experimental::command_graph<
+        sycl::ext::oneapi::experimental::graph_state::modifiable>
+        modifiable_graph(q.get_context(), q.get_device());
+    for (size_t i = 1; i < I; i++) {
+      sycl::range global = {i, i, i};
+      sycl::range local = {i, i, i};
+      modifiable_graph.add([=](sycl::handler &h) {
+        h.parallel_for<class test>(sycl::nd_range{global, local},
+                                   [=](sycl::nd_item<3> it) noexcept {
+                                     if (it.get_group().leader()) {
+                                       output[i]++;
+                                     }
+                                   });
+      });
+    }
+
+    if (r == 0) {
+      const auto instance = modifiable_graph.finalize(
+          sycl::ext::oneapi::experimental::property::graph::updatable{});
+      graph = std::make_unique<sycl::ext::oneapi::experimental::command_graph<
+          sycl::ext::oneapi::experimental::graph_state::executable>>(
+          std::move(instance));
+    } else {
+      graph->update(modifiable_graph);
+    }
+    q.ext_oneapi_graph(*graph).wait();
+  }
+
+  q.wait();
+  std::array<int, I> Ref{0, R, R, R, R};
+
+  for (int i = 0; i < I; ++i) {
+    assert(output[i] == Ref[i]);
+  }
+}

unified-runtime/source/adapters/cuda/command_buffer.cpp

@@ -1347,14 +1347,12 @@ UR_APIEXPORT ur_result_t UR_APICALL urCommandBufferUpdateKernelLaunchExp(
     UR_CHECK_ERROR(validateCommandDesc(hCommandBuffer, pUpdateKernelLaunch[i]));
   }
 
-  // Store changes in config struct in command handle object
+  // Store changes in config struct in command handle object and propagate
+  // changes to CUDA graph
   for (uint32_t i = 0; i < numKernelUpdates; i++) {
     UR_CHECK_ERROR(updateCommand(pUpdateKernelLaunch[i]));
     UR_CHECK_ERROR(updateKernelArguments(pUpdateKernelLaunch[i]));
-  }
 
-  // Propagate changes to CUDA driver API
-  for (uint32_t i = 0; i < numKernelUpdates; i++) {
     const auto &UpdateCommandDesc = pUpdateKernelLaunch[i];
 
     // If no work-size is provided make sure we pass nullptr to setKernelParams

unified-runtime/source/adapters/hip/command_buffer.cpp

@@ -984,14 +984,12 @@ UR_APIEXPORT ur_result_t UR_APICALL urCommandBufferUpdateKernelLaunchExp(
     UR_CHECK_ERROR(validateCommandDesc(hCommandBuffer, pUpdateKernelLaunch[i]));
   }
 
-  // Store changes in config struct in command handle object
+  // Store changes in config struct in command handle object and propagate
+  // changes to HIP Graph.
   for (uint32_t i = 0; i < numKernelUpdates; i++) {
     UR_CHECK_ERROR(updateCommand(pUpdateKernelLaunch[i]));
     UR_CHECK_ERROR(updateKernelArguments(pUpdateKernelLaunch[i]));
-  }
 
-  // Propagate changes to HIP driver API
-  for (uint32_t i = 0; i < numKernelUpdates; i++) {
     const auto &UpdateCommandDesc = pUpdateKernelLaunch[i];
 
     // If no worksize is provided make sure we pass nullptr to setKernelParams