[SYCL][Graph] Update UR tag for L0 kernel binary update

sycl/cmake/modules/UnifiedRuntimeTag.cmake

@@ -1,5 +1,7 @@
-# commit f2af85f35957601dd91f81e8fad39fea413ccbf2
-# Author: Yang Zhao <[email protected]>
-# Date:   Wed Nov 27 00:20:29 2024 +0800
-#     [DeviceSanitizer] Support "-fsanitize-ignorelist=" to disable sanitizing on some of kernels (#2055)
-set(UNIFIED_RUNTIME_TAG f2af85f35957601dd91f81e8fad39fea413ccbf2)
+# commit 0a90db9b2c36960c9b28ce18557ca15760724c4d
+# Merge: c4d9fdb4 6e0bdeb9
+# Author: Callum Fare <[email protected]>
+# Date:   Wed Nov 27 12:16:44 2024 +0000
+#     Merge pull request #2369 from Bensuo/ben/kernel-binary-update-l0
+#     [CMDBUF] Implement kernel binary update for L0 adapter
+set(UNIFIED_RUNTIME_TAG 0a90db9b2c36960c9b28ce18557ca15760724c4d)

sycl/include/sycl/ext/oneapi/experimental/graph.hpp

@@ -444,7 +444,8 @@ class command_graph : public detail::modifiable_command_graph {
   /// Constructor.
   /// @param SyclQueue Queue to use for the graph device and context.
   /// @param PropList Optional list of properties to pass.
-  command_graph(const queue &SyclQueue, const property_list &PropList = {})
+  explicit command_graph(const queue &SyclQueue,
+                         const property_list &PropList = {})
       : modifiable_command_graph(SyclQueue, PropList) {}
 
 private:

sycl/source/detail/device_impl.cpp

@@ -723,7 +723,8 @@ bool device_impl::has(aspect Aspect) const {
         UR_DEVICE_COMMAND_BUFFER_UPDATE_CAPABILITY_FLAG_KERNEL_ARGUMENTS |
         UR_DEVICE_COMMAND_BUFFER_UPDATE_CAPABILITY_FLAG_LOCAL_WORK_SIZE |
         UR_DEVICE_COMMAND_BUFFER_UPDATE_CAPABILITY_FLAG_GLOBAL_WORK_SIZE |
-        UR_DEVICE_COMMAND_BUFFER_UPDATE_CAPABILITY_FLAG_GLOBAL_WORK_OFFSET;
+        UR_DEVICE_COMMAND_BUFFER_UPDATE_CAPABILITY_FLAG_GLOBAL_WORK_OFFSET |
+        UR_DEVICE_COMMAND_BUFFER_UPDATE_CAPABILITY_FLAG_KERNEL_HANDLE;
 
     return has(aspect::ext_oneapi_limited_graph) &&
            (UpdateCapabilities & RequiredCapabilities) == RequiredCapabilities;

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

@@ -5,9 +5,6 @@
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 
-// XFAIL: level_zero
-// XFAIL-TRACKER: OFNAAO-307
-
 // Tests using dynamic command-group objects with buffer accessors
 
 #include "../graph_common.hpp"
@@ -23,14 +20,14 @@ int main() {
       Queue.get_device(),
       {exp_ext::property::graph::assume_buffer_outlives_graph{}}};
 
-  int PatternA = 42;
+  const int PatternA = 42;
   auto CGFA = [&](handler &CGH) {
     CGH.require(Acc);
     CGH.parallel_for(Size,
                      [=](item<1> Item) { Acc[Item.get_id()] = PatternA; });
   };
 
-  int PatternB = 0xA;
+  const int PatternB = 0xA;
   auto CGFB = [&](handler &CGH) {
     CGH.require(Acc);
     CGH.parallel_for(Size,

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

@@ -5,9 +5,6 @@
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 
-// XFAIL: level_zero
-// XFAIL-TRACKER: OFNAAO-307
-
 // Tests adding a dynamic command-group node to a graph using buffer
 // accessors for the node edges.
 

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

@@ -5,9 +5,6 @@
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 
-// XFAIL: level_zero
-// XFAIL-TRACKER: OFNAAO-307
-
 // Tests adding a dynamic command-group node to a graph using buffer
 // accessors for the node edges, but where different command-groups
 // use different buffers that create identical edges.

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

@@ -1,13 +1,11 @@
 // RUN: %{build} -o %t.out
-// RUN: %{run} %t.out %S/../Inputs/Kernels/update_with_indices_accessor.spv
+// RUN: %{run} %t.out %S/../Inputs/Kernels/dyn_cgf_accessor.spv
 // Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG
-// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out %S/../Inputs/Kernels/update_with_indices_accessor.spv 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
+// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out %S/../Inputs/Kernels/dyn_cgf_accessor.spv 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 // Extra run to check for immediate-command-list in Level Zero
-// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out %S/../Inputs/Kernels/update_with_indices_accessor.spv 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
+// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out %S/../Inputs/Kernels/dyn_cgf_accessor.spv 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 
 // REQUIRES: level_zero
-// XFAIL: level_zero
-// XFAIL-TRACKER: OFNAAO-307
 
 // Tests updating an accessor argument to a graph node created from SPIR-V
 // using dynamic command-groups.
@@ -23,8 +21,12 @@ int main(int, char **argv) {
         return bundle.ext_oneapi_get_kernel(name);
       };
 
-  kernel kernel = getKernel(
-      KernelBundle, "_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_");
+  kernel kernelA = getKernel(
+      KernelBundle,
+      "_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlNS0_4itemILi1ELb1EEEE_");
+  kernel kernelB = getKernel(
+      KernelBundle,
+      "_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE0_clES2_EUlNS0_4itemILi1ELb1EEEE_");
 
   exp_ext::command_graph Graph{
       Queue.get_context(),
@@ -36,22 +38,27 @@ int main(int, char **argv) {
   BufA.set_write_back(false);
   BufB.set_write_back(false);
 
+  int PatternA = 42;
+  int PatternB = 0xA;
+
   auto AccA = BufA.get_access();
   auto AccB = BufB.get_access();
 
   auto CGFA = [&](handler &CGH) {
     CGH.require(AccA);
     CGH.set_arg(0, AccA);
-    CGH.single_task(kernel);
+    CGH.set_arg(2, PatternA);
+    CGH.parallel_for(sycl::range<1>(Size), kernelA);
   };
 
   auto CGFB = [&](handler &CGH) {
     CGH.require(AccB);
     CGH.set_arg(0, AccB);
-    CGH.single_task(kernel);
+    CGH.set_arg(2, PatternB);
+    CGH.parallel_for(sycl::range<1>(Size), kernelB);
   };
 
-  auto DynamicCG = exp_ext::dynamic_command_group(Queue, {CGFA, CGFB});
+  auto DynamicCG = exp_ext::dynamic_command_group(Graph, {CGFA, CGFB});
   auto DynamicCGNode = Graph.add(DynamicCG);
   auto ExecGraph = Graph.finalize(exp_ext::property::graph::updatable{});
 
@@ -62,8 +69,8 @@ int main(int, char **argv) {
   Queue.copy(BufA.get_access(), HostDataA.data()).wait();
   Queue.copy(BufB.get_access(), HostDataB.data()).wait();
   for (size_t i = 0; i < Size; i++) {
-    assert(HostDataA[i] == i);
-    assert(HostDataB[i] == 0);
+    assert(check_value(i, PatternA, HostDataA[i], "HostDataA"));
+    assert(check_value(i, 0, HostDataB[i], "HostDataB"));
   }
 
   DynamicCG.set_active_cgf(1);
@@ -74,8 +81,8 @@ int main(int, char **argv) {
   Queue.copy(BufA.get_access(), HostDataA.data()).wait();
   Queue.copy(BufB.get_access(), HostDataB.data()).wait();
   for (size_t i = 0; i < Size; i++) {
-    assert(HostDataA[i] == i);
-    assert(HostDataB[i] == i);
+    assert(check_value(i, PatternA, HostDataA[i], "HostDataA"));
+    assert(check_value(i, PatternB, HostDataB[i], "HostDataB"));
   }
   return 0;
 }

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

@@ -1,12 +1,9 @@
 // RUN: %{build} -o %t.out
 // RUN: env SYCL_UR_TRACE=2 %{run} %t.out | FileCheck %s
 // Extra run to check for leaks in Level Zero using UR_L0_LEAKS_DEBUG
-// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
+// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=0 env SYCL_UR_TRACE=2 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 // Extra run to check for immediate-command-list in Level Zero
-// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
-
-// XFAIL: level_zero
-// XFAIL-TRACKER: OFNAAO-307
+// RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 env SYCL_UR_TRACE=2 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 
 // Tests updating a dynamic command-group with command-groups containing a
 // different number of arguments.

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

@@ -5,9 +5,6 @@
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 
-// XFAIL: level_zero
-// XFAIL-TRACKER: OFNAAO-307
-
 // Tests adding a dynamic command-group node to a graph using graph limited
 // events for dependencies.
 

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

@@ -5,9 +5,6 @@
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 
-// XFAIL: level_zero
-// XFAIL-TRACKER: OFNAAO-307
-
 // Tests updating a dynamic command-group node where the dynamic command-groups
 // have different ranges/nd-ranges
 

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

@@ -5,9 +5,6 @@
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 
-// XFAIL: level_zero
-// XFAIL-TRACKER: OFNAAO-307
-
 // Tests updating a dynamic command-group node where the dynamic command-groups
 // have different range/nd-range dimensions
 

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

@@ -5,9 +5,6 @@
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 
-// XFAIL: level_zero
-// XFAIL-TRACKER: OFNAAO-307
-
 // Tests updating kernel code using dynamic command-groups that have different
 // parameters in each command-group.
 

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

@@ -5,9 +5,6 @@
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 
-// XFAIL: level_zero
-// XFAIL-TRACKER: OFNAAO-307
-
 // Tests using the same dynamic command-group in more than one graph node.
 
 #include "../graph_common.hpp"

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

@@ -5,9 +5,6 @@
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 
-// XFAIL: level_zero
-// XFAIL-TRACKER: OFNAAO-307
-
 // Tests updating usm kernel code using dynamic command-groups
 
 #include "../graph_common.hpp"

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

@@ -5,9 +5,6 @@
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 
-// XFAIL: level_zero
-// XFAIL-TRACKER: OFNAAO-307
-
 // Tests using a dynamic command-group object with dynamic parameters inside it
 
 #include "../graph_common.hpp"

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

@@ -5,9 +5,6 @@
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 
-// XFAIL: level_zero
-// XFAIL-TRACKER: OFNAAO-307
-
 // Tests using a dynamic command-group object with dynamic parameters of
 // different types
 

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

@@ -5,9 +5,6 @@
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 
-// XFAIL: level_zero
-// XFAIL-TRACKER: OFNAAO-307
-
 // Tests using a dynamic command-group object where some but not all the
 // command-groups use dynamic parameters.
 

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

@@ -5,9 +5,6 @@
 // Extra run to check for immediate-command-list in Level Zero
 // RUN: %if level_zero %{env SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS=1 %{l0_leak_check} %{run} %t.out 2>&1 | FileCheck %s --implicit-check-not=LEAK %}
 
-// XFAIL: level_zero
-// XFAIL-TRACKER: OFNAAO-307
-
 // Tests interaction of whole graph update and dynamic command-groups
 
 #include "../graph_common.hpp"