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

- Fix potential implicit conversion from queue to command_graph when using dynamic command groups
- Remove XFAIL for dyn_cgf* tests on L0
- Fix dyn_cgf_accessor_spv test and add new prebuilt spirv binary for it

Author: Bensuo

sycl/cmake/modules/FetchUnifiedRuntime.cmake

@@ -116,7 +116,7 @@ if(SYCL_UR_USE_FETCH_CONTENT)
       CACHE PATH "Path to external '${name}' adapter source dir" FORCE)
   endfunction()
 
-  set(UNIFIED_RUNTIME_REPO "https://github.com/oneapi-src/unified-runtime.git")
+  set(UNIFIED_RUNTIME_REPO "https://github.com/bensuo/unified-runtime.git")
   include(${CMAKE_CURRENT_SOURCE_DIR}/cmake/modules/UnifiedRuntimeTag.cmake)
 
   set(UMF_BUILD_EXAMPLES OFF CACHE INTERNAL "EXAMPLES")

sycl/cmake/modules/UnifiedRuntimeTag.cmake

@@ -4,4 +4,4 @@
 # Date:   Wed Nov 20 14:49:17 2024 +0000
 #     Merge pull request #2258 from aarongreig/aaron/tryUseExtensionSubgroupInfo
 #     Use extension version of clGetKernelSubGroupInfo when necessary.
-set(UNIFIED_RUNTIME_TAG 9937d029c7fdcbf101e89f8515f640c145e059c5)
+set(UNIFIED_RUNTIME_TAG b24eb57d99e5180ce851dc726984c7dc117e36b9)

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/test-e2e/Graph/Inputs/Kernels/dyn_cgf_accessor.spv

@@ -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.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_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, but where different command-groups
 // use different buffers that create identical edges.

sycl/test-e2e/Graph/Update/dyn_cgf_accessor_deps2.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_accessor_spv.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_different_arg_nums.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.