[SYCL][Graph] Fix access modes not being respected

- Fix access modes not being respected and creating unnecessary edges - Update printing E2E tests since output has changed - Add unit tests for access modes
reble · Mar 13, 2024 · d89a3ea · d89a3ea
1 parent 37cb495
commit d89a3ea
Show file tree

Hide file tree

Showing 9 changed files with 255 additions and 54 deletions.
diff --git a/sycl/source/detail/graph_impl.cpp b/sycl/source/detail/graph_impl.cpp
@@ -439,12 +439,12 @@ graph_impl::add(node_type NodeType,
     }
     // Look through the graph for nodes which share this requirement
     for (auto &Node : MNodeStorage) {
-      if (Node->hasRequirement(Req)) {
+      if (Node->hasRequirementDependency(Req)) {
         bool ShouldAddDep = true;
         // If any of this node's successors have this requirement then we skip
         // adding the current node as a dependency.
         for (auto &Succ : Node->MSuccessors) {
-          if (Succ.lock()->hasRequirement(Req)) {
+          if (Succ.lock()->hasRequirementDependency(Req)) {
             ShouldAddDep = false;
             break;
           }

diff --git a/sycl/source/detail/graph_impl.hpp b/sycl/source/detail/graph_impl.hpp
@@ -154,16 +154,38 @@ class node_impl {
         MCGType(Other.MCGType), MNodeType(Other.MNodeType),
         MCommandGroup(Other.getCGCopy()), MSubGraphImpl(Other.MSubGraphImpl) {}
 
-  /// Checks if this node has a given requirement.
-  /// @param Requirement Requirement to lookup.
-  /// @return True if \p Requirement is present in node, false otherwise.
-  bool hasRequirement(sycl::detail::AccessorImplHost *IncomingReq) {
+  /// Checks if this node should be a dependency of another node based on
+  /// accessor requirements. This is calculated using access modes if a
+  /// requirement to the same buffer is found inside this node.
+  /// @param IncomingReq Incoming requirement.
+  /// @return True if a dependency is needed, false if not.
+  bool hasRequirementDependency(sycl::detail::AccessorImplHost *IncomingReq) {
+    access_mode InMode = IncomingReq->MAccessMode;
+    switch (InMode) {
+    case access_mode::read:
+    case access_mode::read_write:
+    case access_mode::atomic:
+      break;
+    // These access modes don't care about existing buffer data, so we don't
+    // need a dependency.
+    case access_mode::write:
+    case access_mode::discard_read_write:
+    case access_mode::discard_write:
+      return false;
+    }
+
     for (sycl::detail::AccessorImplHost *CurrentReq :
          MCommandGroup->getRequirements()) {
       if (IncomingReq->MSYCLMemObj == CurrentReq->MSYCLMemObj) {
-        return true;
+        access_mode CurrentMode = CurrentReq->MAccessMode;
+        // Since we have an incoming read requirement, we only care
+        // about requirements on this node if they are write
+        if (CurrentMode != access_mode::read) {
+          return true;
+        }
       }
     }
+    // No dependency necessary
     return false;
   }
 

diff --git a/sycl/test-e2e/Graph/Explicit/debug_print_graph.cpp b/sycl/test-e2e/Graph/Explicit/debug_print_graph.cpp
@@ -12,22 +12,22 @@
 // CHECK-SAME: [style=bold, label="ID = 0x[[#NODE2]]\nTYPE = CGExecKernel \nNAME = _ZTSZZ11add_kernelsItEN4sycl3_V13ext6oneapi12experimental4nodeENS4_13command_graphILNS4_11graph_stateE0EEEmNS1_6bufferIT_Li1ENS1_6detail17aligned_allocatorINSt12remove_constISA_E4typeEEEvEESH_SH_ENKUlRNS1_7handlerEE0_clESJ_EUlNS1_4itemILi1ELb1EEEE_\n"];
 // CHECK-NEXT: "0x[[#NODE1]]" -> "0x[[#NODE2]]"
 // CHECK-NEXT: "0x[[#%x,NODE3:]]"
-// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE3]]\nTYPE = CGExecKernel \nNAME = _ZTSZZ11add_kernelsItEN4sycl3_V13ext6oneapi12experimental4nodeENS4_13command_graphILNS4_11graph_stateE0EEEmNS1_6bufferIT_Li1ENS1_6detail17aligned_allocatorINSt12remove_constISA_E4typeEEEvEESH_SH_ENKUlRNS1_7handlerEE1_clESJ_EUlNS1_4itemILi1ELb1EEEE_\n"];
-// CHECK-NEXT: "0x[[#NODE2]]" -> "0x[[#NODE3]]"
+// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE3]]\nTYPE = CGExecKernel \nNAME = _ZTSZZ11add_kernelsItEN4sycl3_V13ext6oneapi12experimental4nodeENS4_13command_graphILNS4_11graph_stateE0EEEmNS1_6bufferIT_Li1ENS1_6detail17aligned_allocatorINSt12remove_constISA_E4typeEEEvEESH_SH_ENKUlRNS1_7handlerEE2_clESJ_EUlNS1_4itemILi1ELb1EEEE_\n"];
+// CHECK-DAG: "0x[[#NODE2]]" -> "0x[[#NODE3]]"
+// CHECK-DAG: "0x[[#%x,NODE7:]]" -> "0x[[#NODE3]]"
 // CHECK-NEXT: "0x[[#%x,NODE4:]]"
-// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE4]]\nTYPE = CGExecKernel \nNAME = _ZTSZZ11add_kernelsItEN4sycl3_V13ext6oneapi12experimental4nodeENS4_13command_graphILNS4_11graph_stateE0EEEmNS1_6bufferIT_Li1ENS1_6detail17aligned_allocatorINSt12remove_constISA_E4typeEEEvEESH_SH_ENKUlRNS1_7handlerEE2_clESJ_EUlNS1_4itemILi1ELb1EEEE_\n"];
+// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE4]]\nTYPE = CGCopy Device-to-Device \n"];
 // CHECK-DAG: "0x[[#NODE3]]" -> "0x[[#NODE4]]"
-// CHECK-DAG: "0x[[#NODE2]]" -> "0x[[#NODE4]]"
+// CHECK-DAG: "0x[[#NODE1]]" -> "0x[[#NODE4]]"
 // CHECK-NEXT: "0x[[#%x,NODE5:]]"
-// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE5]]\nTYPE = CGCopy Device-to-Device \n"];
+// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE5]]\nTYPE = CGCopy Device-to-Host \n"];
 // CHECK-DAG: "0x[[#NODE3]]" -> "0x[[#NODE5]]"
-// CHECK-DAG: "0x[[#NODE4]]" -> "0x[[#NODE5]]
 // CHECK-NEXT: "0x[[#%x,NODE6:]]"
-// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE6]]\nTYPE = CGCopy Device-to-Host \n"];
-// CHECK-DAG: "0x[[#NODE4]]" -> "0x[[#NODE6]]"
-// CHECK-NEXT: "0x[[#%x,NODE7:]]"
-// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE7]]\nTYPE = None \n"];
-// CHECK-DAG: "0x[[#NODE6]]" -> "0x[[#NODE7]]"
+// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE6]]\nTYPE = None \n"];
+// CHECK-DAG: "0x[[#NODE5]]" -> "0x[[#NODE6]]"
+// CHECK-NEXT: "0x[[#NODE7]]"
+// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE7]]\nTYPE = CGExecKernel \nNAME = _ZTSZZ11add_kernelsItEN4sycl3_V13ext6oneapi12experimental4nodeENS4_13command_graphILNS4_11graph_stateE0EEEmNS1_6bufferIT_Li1ENS1_6detail17aligned_allocatorINSt12remove_constISA_E4typeEEEvEESH_SH_ENKUlRNS1_7handlerEE1_clESJ_EUlNS1_4itemILi1ELb1EEEE_\n"];
+// CHECK-DAG: "0x[[#NODE1]]" -> "0x[[#NODE7]]"
 
 #define GRAPH_E2E_EXPLICIT
 

diff --git a/sycl/test-e2e/Graph/Explicit/debug_print_graph_verbose.cpp b/sycl/test-e2e/Graph/Explicit/debug_print_graph_verbose.cpp
@@ -16,26 +16,26 @@
 // CHECK-SAME: 3) Type: STD_Layout Ptr: 0x[[#ADDR4]]\n4) Type: Accessor Ptr: 0x[[#%x,ADDR7:]]\n5) Type: STD_Layout Ptr: 0x[[#%x,ADDR8:]]\n6) Type: STD_Layout Ptr: 0x[[#%x,ADDR9:]]\n7) Type: STD_Layout Ptr: 0x[[#%x,ADDR10:]]\n"];
 // CHECK-NEXT: "0x[[#NODE1]]" -> "0x[[#NODE2]]"
 // CHECK-NEXT: "0x[[#%x,NODE3:]]"
-// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE3]]\nTYPE = CGExecKernel \nNAME = _ZTSZZ11add_kernelsItEN4sycl3_V13ext6oneapi12experimental4nodeENS4_13command_graphILNS4_11graph_stateE0EEEmNS1_6bufferIT_Li1ENS1_6detail17aligned_allocatorINSt12remove_constISA_E4typeEEEvEESH_SH_ENKUlRNS1_7handlerEE1_clESJ_EUlNS1_4itemILi1ELb1EEEE_\n
+// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE3]]\nTYPE = CGExecKernel \nNAME = _ZTSZZ11add_kernelsItEN4sycl3_V13ext6oneapi12experimental4nodeENS4_13command_graphILNS4_11graph_stateE0EEEmNS1_6bufferIT_Li1ENS1_6detail17aligned_allocatorINSt12remove_constISA_E4typeEEEvEESH_SH_ENKUlRNS1_7handlerEE2_clESJ_EUlNS1_4itemILi1ELb1EEEE_\n
 // CHECK-SAME: ARGS = \n0) Type: Accessor Ptr: 0x[[#%x,ADDR11:]]\n1) Type: STD_Layout Ptr: 0x[[#%x,ADDR12:]]\n2) Type: STD_Layout Ptr: 0x[[#%x,ADDR13:]]\n
 // CHECK-SAME: 3) Type: STD_Layout Ptr: 0x[[#ADDR11]]\n4) Type: Accessor Ptr: 0x[[#%x,ADDR14:]]\n5) Type: STD_Layout Ptr: 0x[[#%x,ADDR15:]]\n6) Type: STD_Layout Ptr: 0x[[#%x,ADDR16:]]\n7) Type: STD_Layout Ptr: 0x[[#%x,ADDR17:]]\n"];
 // CHECK-NEXT: "0x[[#NODE2]]" -> "0x[[#NODE3]]"
+// CHECK-NEXT: "0x[[#%x,NODE7:]]" -> "0x[[#NODE3]]"
 // CHECK-NEXT: "0x[[#%x,NODE4:]]"
-// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE4]]\nTYPE = CGExecKernel \nNAME = _ZTSZZ11add_kernelsItEN4sycl3_V13ext6oneapi12experimental4nodeENS4_13command_graphILNS4_11graph_stateE0EEEmNS1_6bufferIT_Li1ENS1_6detail17aligned_allocatorINSt12remove_constISA_E4typeEEEvEESH_SH_ENKUlRNS1_7handlerEE2_clESJ_EUlNS1_4itemILi1ELb1EEEE_\n
-// CHECK-SAME: ARGS = \n0) Type: Accessor Ptr: 0x[[#%x,ADDR18:]]\n1) Type: STD_Layout Ptr: 0x[[#%x,ADDR19:]]\n2) Type: STD_Layout Ptr: 0x[[#%x,ADDR20:]]\n
-// CHECK-SAME: 3) Type: STD_Layout Ptr: 0x[[#ADDR18]]\n4) Type: Accessor Ptr: 0x[[#%x,ADDR21:]]\n5) Type: STD_Layout Ptr: 0x[[#%x,ADDR22:]]\n6) Type: STD_Layout Ptr: 0x[[#%x,ADDR23:]]\n7) Type: STD_Layout Ptr: 0x[[#%x,ADDR24:]]\n"];
+// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE4]]\nTYPE = CGCopy Device-to-Device \nSrc: 0x[[#%x,ADDR18:]] Dst: 0x[[#%x,ADDR19:]]\n"];
 // CHECK-DAG: "0x[[#NODE3]]" -> "0x[[#NODE4]]"
-// CHECK-DAG: "0x[[#NODE2]]" -> "0x[[#NODE4]]"
+// CHECK-DAG: "0x[[#NODE1]]" -> "0x[[#NODE4]]
 // CHECK-NEXT: "0x[[#%x,NODE5:]]"
-// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE5]]\nTYPE = CGCopy Device-to-Device \nSrc: 0x[[#%x,ADDR25:]] Dst: 0x[[#%x,ADDR26:]]\n"];
+// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE5]]\nTYPE = CGCopy Device-to-Host \nSrc: 0x[[#%x,ADDR20:]] Dst: 0x[[#%x,ADDR21:]]\n"];
 // CHECK-DAG: "0x[[#NODE3]]" -> "0x[[#NODE5]]"
-// CHECK-DAG: "0x[[#NODE4]]" -> "0x[[#NODE5]]
 // CHECK-NEXT: "0x[[#%x,NODE6:]]"
-// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE6]]\nTYPE = CGCopy Device-to-Host \nSrc: 0x[[#%x,ADDR27:]] Dst: 0x[[#%x,ADDR28:]]\n"];
-// CHECK-DAG: "0x[[#NODE4]]" -> "0x[[#NODE6]]"
-// CHECK-NEXT: "0x[[#%x,NODE7:]]"
-// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE7]]\nTYPE = None \n"];
-// CHECK-DAG: "0x[[#NODE6]]" -> "0x[[#NODE7]]"
+// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE6]]\nTYPE = None \n"];
+// CHECK-DAG: "0x[[#NODE5]]" -> "0x[[#NODE6]]"
+// CHECK-NEXT: "0x[[#NODE7]]"
+// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE7]]\nTYPE = CGExecKernel \nNAME = _ZTSZZ11add_kernelsItEN4sycl3_V13ext6oneapi12experimental4nodeENS4_13command_graphILNS4_11graph_stateE0EEEmNS1_6bufferIT_Li1ENS1_6detail17aligned_allocatorINSt12remove_constISA_E4typeEEEvEESH_SH_ENKUlRNS1_7handlerEE1_clESJ_EUlNS1_4itemILi1ELb1EEEE_\n
+// CHECK-SAME: ARGS = \n0) Type: Accessor Ptr: 0x[[#%x,ADDR22:]]\n1) Type: STD_Layout Ptr: 0x[[#%x,ADDR23:]]\n2) Type: STD_Layout Ptr: 0x[[#%x,ADDR24:]]\n
+// CHECK-SAME: 3) Type: STD_Layout Ptr: 0x[[#ADDR22]]\n4) Type: Accessor Ptr: 0x[[#%x,ADDR25:]]\n5) Type: STD_Layout Ptr: 0x[[#%x,ADDR26:]]\n6) Type: STD_Layout Ptr: 0x[[#%x,ADDR27:]]\n7) Type: STD_Layout Ptr: 0x[[#%x,ADDR28:]]\n"];
+// CHECK-DAG: "0x[[#NODE1]]" -> "0x[[#NODE7]]"
 
 #define GRAPH_E2E_EXPLICIT
 

diff --git a/sycl/test-e2e/Graph/RecordReplay/debug_print_graph.cpp b/sycl/test-e2e/Graph/RecordReplay/debug_print_graph.cpp
@@ -12,22 +12,22 @@
 // CHECK-SAME: [style=bold, label="ID = 0x[[#NODE2]]\nTYPE = CGExecKernel \nNAME = _ZTSZZ11run_kernelsItEN4sycl3_V15eventENS1_5queueEmNS1_6bufferIT_Li1ENS1_6detail17aligned_allocatorINSt12remove_constIS5_E4typeEEEvEESC_SC_ENKUlRNS1_7handlerEE0_clESE_EUlNS1_4itemILi1ELb1EEEE_\n"];
 // CHECK-NEXT: "0x[[#NODE1]]" -> "0x[[#NODE2]]"
 // CHECK-NEXT: "0x[[#%x,NODE3:]]"
-// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE3]]\nTYPE = CGExecKernel \nNAME = _ZTSZZ11run_kernelsItEN4sycl3_V15eventENS1_5queueEmNS1_6bufferIT_Li1ENS1_6detail17aligned_allocatorINSt12remove_constIS5_E4typeEEEvEESC_SC_ENKUlRNS1_7handlerEE1_clESE_EUlNS1_4itemILi1ELb1EEEE_\n"];
-// CHECK-NEXT: "0x[[#NODE2]]" -> "0x[[#NODE3]]"
+// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE3]]\nTYPE = CGExecKernel \nNAME = _ZTSZZ11run_kernelsItEN4sycl3_V15eventENS1_5queueEmNS1_6bufferIT_Li1ENS1_6detail17aligned_allocatorINSt12remove_constIS5_E4typeEEEvEESC_SC_ENKUlRNS1_7handlerEE2_clESE_EUlNS1_4itemILi1ELb1EEEE_\n"];
+// CHECK-DAG: "0x[[#NODE2]]" -> "0x[[#NODE3]]"
+// CHECK-DAG: "0x[[#%x,NODE7:]]" -> "0x[[#NODE3]]"
 // CHECK-NEXT: "0x[[#%x,NODE4:]]"
-// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE4]]\nTYPE = CGExecKernel \nNAME = _ZTSZZ11run_kernelsItEN4sycl3_V15eventENS1_5queueEmNS1_6bufferIT_Li1ENS1_6detail17aligned_allocatorINSt12remove_constIS5_E4typeEEEvEESC_SC_ENKUlRNS1_7handlerEE2_clESE_EUlNS1_4itemILi1ELb1EEEE_\n"];
+// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE4]]\nTYPE = CGCopy Device-to-Device \n"];
 // CHECK-DAG: "0x[[#NODE3]]" -> "0x[[#NODE4]]"
-// CHECK-DAG: "0x[[#NODE2]]" -> "0x[[#NODE4]]"
+// CHECK-DAG: "0x[[#NODE1]]" -> "0x[[#NODE4]]"
 // CHECK-NEXT: "0x[[#%x,NODE5:]]"
-// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE5]]\nTYPE = CGCopy Device-to-Device \n"];
+// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE5]]\nTYPE = CGCopy Device-to-Host \n"];
 // CHECK-DAG: "0x[[#NODE3]]" -> "0x[[#NODE5]]"
-// CHECK-DAG: "0x[[#NODE4]]" -> "0x[[#NODE5]]
 // CHECK-NEXT: "0x[[#%x,NODE6:]]"
-// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE6]]\nTYPE = CGCopy Device-to-Host \n"];
-// CHECK-DAG: "0x[[#NODE4]]" -> "0x[[#NODE6]]"
-// CHECK-NEXT: "0x[[#%x,NODE7:]]"
-// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE7]]\nTYPE = None \n"];
-// CHECK-DAG: "0x[[#NODE6]]" -> "0x[[#NODE7]]"
+// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE6]]\nTYPE = None \n"];
+// CHECK-DAG: "0x[[#NODE5]]" -> "0x[[#NODE6]]"
+// CHECK-NEXT: "0x[[#NODE7]]"
+// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE7]]\nTYPE = CGExecKernel \nNAME = _ZTSZZ11run_kernelsItEN4sycl3_V15eventENS1_5queueEmNS1_6bufferIT_Li1ENS1_6detail17aligned_allocatorINSt12remove_constIS5_E4typeEEEvEESC_SC_ENKUlRNS1_7handlerEE1_clESE_EUlNS1_4itemILi1ELb1EEEE_\n"];
+// CHECK-DAG: "0x[[#NODE1]]" -> "0x[[#NODE7]]"
 
 #define GRAPH_E2E_RECORD_REPLAY
 

diff --git a/sycl/test-e2e/Graph/RecordReplay/debug_print_graph_verbose.cpp b/sycl/test-e2e/Graph/RecordReplay/debug_print_graph_verbose.cpp
@@ -17,26 +17,26 @@
 // CHECK-SAME: 3) Type: STD_Layout Ptr: 0x[[#ADDR4]]\n4) Type: Accessor Ptr: 0x[[#%x,ADDR7:]]\n5) Type: STD_Layout Ptr: 0x[[#%x,ADDR8:]]\n6) Type: STD_Layout Ptr: 0x[[#%x,ADDR9:]]\n7) Type: STD_Layout Ptr: 0x[[#%x,ADDR10:]]\n"];
 // CHECK-NEXT: "0x[[#NODE1]]" -> "0x[[#NODE2]]"
 // CHECK-NEXT: "0x[[#%x,NODE3:]]"
-// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE3]]\nTYPE = CGExecKernel \nNAME = _ZTSZZ11run_kernelsItEN4sycl3_V15eventENS1_5queueEmNS1_6bufferIT_Li1ENS1_6detail17aligned_allocatorINSt12remove_constIS5_E4typeEEEvEESC_SC_ENKUlRNS1_7handlerEE1_clESE_EUlNS1_4itemILi1ELb1EEEE_\n
+// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE3]]\nTYPE = CGExecKernel \nNAME = _ZTSZZ11run_kernelsItEN4sycl3_V15eventENS1_5queueEmNS1_6bufferIT_Li1ENS1_6detail17aligned_allocatorINSt12remove_constIS5_E4typeEEEvEESC_SC_ENKUlRNS1_7handlerEE2_clESE_EUlNS1_4itemILi1ELb1EEEE_\n
 // CHECK-SAME: ARGS = \n0) Type: Accessor Ptr: 0x[[#%x,ADDR11:]]\n1) Type: STD_Layout Ptr: 0x[[#%x,ADDR12:]]\n2) Type: STD_Layout Ptr: 0x[[#%x,ADDR13:]]\n
 // CHECK-SAME: 3) Type: STD_Layout Ptr: 0x[[#ADDR11]]\n4) Type: Accessor Ptr: 0x[[#%x,ADDR14:]]\n5) Type: STD_Layout Ptr: 0x[[#%x,ADDR15:]]\n6) Type: STD_Layout Ptr: 0x[[#%x,ADDR16:]]\n7) Type: STD_Layout Ptr: 0x[[#%x,ADDR17:]]\n"];
 // CHECK-NEXT: "0x[[#NODE2]]" -> "0x[[#NODE3]]"
+// CHECK-NEXT: "0x[[#%x,NODE7:]]" -> "0x[[#NODE3]]"
 // CHECK-NEXT: "0x[[#%x,NODE4:]]"
-// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE4]]\nTYPE = CGExecKernel \nNAME = _ZTSZZ11run_kernelsItEN4sycl3_V15eventENS1_5queueEmNS1_6bufferIT_Li1ENS1_6detail17aligned_allocatorINSt12remove_constIS5_E4typeEEEvEESC_SC_ENKUlRNS1_7handlerEE2_clESE_EUlNS1_4itemILi1ELb1EEEE_\n
-// CHECK-SAME: ARGS = \n0) Type: Accessor Ptr: 0x[[#%x,ADDR18:]]\n1) Type: STD_Layout Ptr: 0x[[#%x,ADDR19:]]\n2) Type: STD_Layout Ptr: 0x[[#%x,ADDR20:]]\n
-// CHECK-SAME: 3) Type: STD_Layout Ptr: 0x[[#ADDR18]]\n4) Type: Accessor Ptr: 0x[[#%x,ADDR21:]]\n5) Type: STD_Layout Ptr: 0x[[#%x,ADDR22:]]\n6) Type: STD_Layout Ptr: 0x[[#%x,ADDR23:]]\n7) Type: STD_Layout Ptr: 0x[[#%x,ADDR24:]]\n"];
+// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE4]]\nTYPE = CGCopy Device-to-Device \nSrc: 0x[[#%x,ADDR18:]] Dst: 0x[[#%x,ADDR19:]]\n"];
 // CHECK-DAG: "0x[[#NODE3]]" -> "0x[[#NODE4]]"
-// CHECK-DAG: "0x[[#NODE2]]" -> "0x[[#NODE4]]"
+// CHECK-DAG: "0x[[#NODE1]]" -> "0x[[#NODE4]]
 // CHECK-NEXT: "0x[[#%x,NODE5:]]"
-// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE5]]\nTYPE = CGCopy Device-to-Device \nSrc: 0x[[#%x,ADDR25:]] Dst: 0x[[#%x,ADDR26:]]\n"];
+// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE5]]\nTYPE = CGCopy Device-to-Host \nSrc: 0x[[#%x,ADDR20:]] Dst: 0x[[#%x,ADDR21:]]\n"];
 // CHECK-DAG: "0x[[#NODE3]]" -> "0x[[#NODE5]]"
-// CHECK-DAG: "0x[[#NODE4]]" -> "0x[[#NODE5]]
 // CHECK-NEXT: "0x[[#%x,NODE6:]]"
-// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE6]]\nTYPE = CGCopy Device-to-Host \nSrc: 0x[[#%x,ADDR27:]] Dst: 0x[[#%x,ADDR28:]]\n"];
-// CHECK-DAG: "0x[[#NODE4]]" -> "0x[[#NODE6]]"
-// CHECK-NEXT: "0x[[#%x,NODE7:]]"
-// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE7]]\nTYPE = None \n"];
-// CHECK-DAG: "0x[[#NODE6]]" -> "0x[[#NODE7]]"
+// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE6]]\nTYPE = None \n"];
+// CHECK-DAG: "0x[[#NODE5]]" -> "0x[[#NODE6]]"
+// CHECK-NEXT: "0x[[#NODE7]]"
+// CHECK-SAME: [style=bold, label="ID = 0x[[#NODE7]]\nTYPE = CGExecKernel \nNAME = _ZTSZZ11run_kernelsItEN4sycl3_V15eventENS1_5queueEmNS1_6bufferIT_Li1ENS1_6detail17aligned_allocatorINSt12remove_constIS5_E4typeEEEvEESC_SC_ENKUlRNS1_7handlerEE1_clESE_EUlNS1_4itemILi1ELb1EEEE_\n
+// CHECK-SAME: ARGS = \n0) Type: Accessor Ptr: 0x[[#%x,ADDR22:]]\n1) Type: STD_Layout Ptr: 0x[[#%x,ADDR23:]]\n2) Type: STD_Layout Ptr: 0x[[#%x,ADDR24:]]\n
+// CHECK-SAME: 3) Type: STD_Layout Ptr: 0x[[#ADDR22]]\n4) Type: Accessor Ptr: 0x[[#%x,ADDR25:]]\n5) Type: STD_Layout Ptr: 0x[[#%x,ADDR26:]]\n6) Type: STD_Layout Ptr: 0x[[#%x,ADDR27:]]\n7) Type: STD_Layout Ptr: 0x[[#%x,ADDR28:]]\n"];
+// CHECK-DAG: "0x[[#NODE1]]" -> "0x[[#NODE7]]"
 
 #define GRAPH_E2E_RECORD_REPLAY
 

diff --git a/sycl/unittests/Extensions/CommandGraph/CMakeLists.txt b/sycl/unittests/Extensions/CommandGraph/CMakeLists.txt
@@ -7,5 +7,6 @@ add_sycl_unittest(CommandGraphExtensionTests OBJECT
   InOrderQueue.cpp
   MultiThreaded.cpp
   Queries.cpp
+  Regressions.cpp
   Subgraph.cpp
 )