Reintroduce unittest for range update errors

Author: Ewan Crawford

sycl/doc/design/CommandGraph.md

@@ -291,18 +291,17 @@ from a `std::unique_ptr` to a `std::shared_ptr` so that multiple nodes and the
 the overhead of having to allocate and free copies of the CG when a new active
 CG is selected.
 
-The `dynamic_command_group_impl` class contains weak pointers to the nodes which
-have been created with it, so that when a new active CG is selected it can
-propagate the change to those nodes. The `node_impl` class also contains a
-reference to the dynamic command-group that created it, so that when the graph
-is finalized each node can use the list of kernels in its dynamic command-group
-as part of the `urCommandBufferAppendKernelLaunchExp` call to pass the possible
-alternative kernels.
+The `dynamic_command_group_impl` class contains a list of weak pointers to the
+nodes which have been created with it, so that when a new active CG is selected
+it can propagate the change to those nodes. The `dynamic_parameter_impl` class
+also contains a list of weak pointers, but to the `dynamic_command_group_impl`
+instances of any dynamic command-groups where they are used. This allows
+updating the dynamic parameter to propagate to dynamic command-group nodes.
 
 The `sycl::detail::CGExecKernel` class has been added to, so that if the
 object was created from an element in the dynamic command-group list, the class
 stores a vector of weak pointers to the other alternative command-groups created
-from the same dynamic command-group object. This allows the DPC++ scheduler to
+from the same dynamic command-group object. This allows the SYCL runtime to
 access the list of alternative kernels when calling the UR API to append a
 kernel command to a command-buffer.
 

sycl/source/detail/graph_impl.cpp

@@ -562,7 +562,6 @@ graph_impl::add(std::shared_ptr<dynamic_command_group_impl> &DynCGImpl,
 
   // Track the dynamic command-group used inside the node object
   DynCGImpl->MNodes.push_back(NodeImpl);
-  NodeImpl->MDynCG = DynCGImpl;
 
   return NodeImpl;
 }

sycl/source/detail/graph_impl.hpp

@@ -99,8 +99,6 @@ class node_impl : public std::enable_shared_from_this<node_impl> {
   /// Stores the executable graph impl associated with this node if it is a
   /// subgraph node.
   std::shared_ptr<exec_graph_impl> MSubGraphImpl;
-  /// Dynamic command-group object used in node, if any.
-  std::shared_ptr<dynamic_command_group_impl> MDynCG;
 
   /// Used for tracking visited status during cycle checks.
   bool MVisited = false;
@@ -158,7 +156,7 @@ class node_impl : public std::enable_shared_from_this<node_impl> {
       : enable_shared_from_this(Other), MSuccessors(Other.MSuccessors),
         MPredecessors(Other.MPredecessors), MCGType(Other.MCGType),
         MNodeType(Other.MNodeType), MCommandGroup(Other.getCGCopy()),
-        MSubGraphImpl(Other.MSubGraphImpl), MDynCG(Other.MDynCG) {}
+        MSubGraphImpl(Other.MSubGraphImpl) {}
 
   /// Copy-assignment operator. This will perform a deep-copy of the
   /// command group object associated with this node.
@@ -170,7 +168,6 @@ class node_impl : public std::enable_shared_from_this<node_impl> {
       MNodeType = Other.MNodeType;
       MCommandGroup = Other.getCGCopy();
       MSubGraphImpl = Other.MSubGraphImpl;
-      MDynCG = Other.MDynCG;
     }
     return *this;
   }
@@ -420,7 +417,7 @@ class node_impl : public std::enable_shared_from_this<node_impl> {
       throw sycl::exception(sycl::errc::invalid,
                             "Cannot update execution range of a node with an "
                             "execution range of different dimensions than what "
-                            "the node was originall created with.");
+                            "the node was original created with.");
     }
 
     NDRDesc = sycl::detail::NDRDescT{ExecutionRange};
@@ -441,7 +438,7 @@ class node_impl : public std::enable_shared_from_this<node_impl> {
       throw sycl::exception(sycl::errc::invalid,
                             "Cannot update execution range of a node with an "
                             "execution range of different dimensions than what "
-                            "the node was originall created with.");
+                            "the node was original created with.");
     }
 
     NDRDesc = sycl::detail::NDRDescT{ExecutionRange};

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

@@ -14,9 +14,7 @@
 
 int main() {
   queue Queue{};
-  const size_t N = 1024;
-  std::vector<int> HostData(N, 0);
-  buffer Buf{HostData};
+  buffer<int> Buf{sycl::range<1>(Size)};
   Buf.set_write_back(false);
   auto Acc = Buf.get_access();
 
@@ -28,22 +26,25 @@ int main() {
   int PatternA = 42;
   auto CGFA = [&](handler &CGH) {
     CGH.require(Acc);
-    CGH.parallel_for(N, [=](item<1> Item) { Acc[Item.get_id()] = PatternA; });
+    CGH.parallel_for(Size,
+                     [=](item<1> Item) { Acc[Item.get_id()] = PatternA; });
   };
 
   int PatternB = 0xA;
   auto CGFB = [&](handler &CGH) {
     CGH.require(Acc);
-    CGH.parallel_for(N, [=](item<1> Item) { Acc[Item.get_id()] = PatternB; });
+    CGH.parallel_for(Size,
+                     [=](item<1> Item) { Acc[Item.get_id()] = PatternB; });
   };
 
   auto DynamicCG = exp_ext::dynamic_command_group(Graph, {CGFA, CGFB});
   auto DynamicCGNode = Graph.add(DynamicCG);
   auto ExecGraph = Graph.finalize(exp_ext::property::graph::updatable{});
 
   Queue.ext_oneapi_graph(ExecGraph).wait();
+  std::vector<int> HostData(Size, 0);
   Queue.copy(Acc, HostData.data()).wait();
-  for (size_t i = 0; i < N; i++) {
+  for (size_t i = 0; i < Size; i++) {
     assert(HostData[i] == PatternA);
   }
 
@@ -52,7 +53,7 @@ int main() {
   Queue.ext_oneapi_graph(ExecGraph).wait();
 
   Queue.copy(Acc, HostData.data()).wait();
-  for (size_t i = 0; i < N; i++) {
+  for (size_t i = 0; i < Size; i++) {
     assert(HostData[i] == PatternB);
   }
 

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

@@ -15,10 +15,8 @@
 
 int main() {
   queue Queue{};
-  const size_t N = 1024;
-  int *Ptr = (int *)sycl::malloc_device<int>(N, Queue);
-  std::vector<int> HostData(N, 0);
-  buffer Buf{HostData};
+  int *Ptr = (int *)sycl::malloc_device<int>(Size, Queue);
+  buffer<int> Buf{sycl::range<1>(Size)};
   Buf.set_write_back(false);
 
   exp_ext::command_graph Graph{
@@ -28,19 +26,21 @@ int main() {
 
   auto RootNode = Graph.add([&](handler &CGH) {
     auto Acc = Buf.get_access<access::mode::write>(CGH);
-    CGH.parallel_for(N, [=](item<1> Item) { Acc[Item.get_id()] = 1; });
+    CGH.parallel_for(Size, [=](item<1> Item) { Acc[Item.get_id()] = 1; });
   });
 
   int PatternA = 42;
   auto CGFA = [&](handler &CGH) {
     auto Acc = Buf.get_access<access::mode::read_write>(CGH);
-    CGH.parallel_for(N, [=](item<1> Item) { Acc[Item.get_id()] += PatternA; });
+    CGH.parallel_for(Size,
+                     [=](item<1> Item) { Acc[Item.get_id()] += PatternA; });
   };
 
   int PatternB = 0xA;
   auto CGFB = [&](handler &CGH) {
     auto Acc = Buf.get_access<access::mode::read_write>(CGH);
-    CGH.parallel_for(N, [=](item<1> Item) { Acc[Item.get_id()] += PatternB; });
+    CGH.parallel_for(Size,
+                     [=](item<1> Item) { Acc[Item.get_id()] += PatternB; });
   };
 
   auto DynamicCG = exp_ext::dynamic_command_group(Graph, {CGFA, CGFB});
@@ -49,23 +49,28 @@ int main() {
   auto LeafNode = Graph.add([&](handler &CGH) {
     auto Acc = Buf.get_access<access::mode::read>(CGH);
     CGH.parallel_for(
-        N, [=](item<1> Item) { Ptr[Item.get_id()] = Acc[Item.get_id()]; });
+        Size, [=](item<1> Item) { Ptr[Item.get_id()] = Acc[Item.get_id()]; });
   });
   auto ExecGraph = Graph.finalize(exp_ext::property::graph::updatable{});
 
   Queue.ext_oneapi_graph(ExecGraph).wait();
-  Queue.copy(Ptr, HostData.data(), N).wait();
-  for (size_t i = 0; i < N; i++) {
-    assert(HostData[i] == (PatternA + 1));
+
+  std::vector<int> HostData(Size, 0);
+  Queue.copy(Ptr, HostData.data(), Size).wait();
+
+  int Ref = PatternA + 1;
+  for (size_t i = 0; i < Size; i++) {
+    assert(HostData[i] == Ref);
   }
 
   DynamicCG.set_active_cgf(1);
   ExecGraph.update(DynamicCGNode);
 
   Queue.ext_oneapi_graph(ExecGraph).wait();
-  Queue.copy(Ptr, HostData.data(), N).wait();
-  for (size_t i = 0; i < N; i++) {
-    assert(HostData[i] == (PatternB + 1));
+  Queue.copy(Ptr, HostData.data(), Size).wait();
+  Ref = PatternB + 1;
+  for (size_t i = 0; i < Size; i++) {
+    assert(HostData[i] == Ref);
   }
 
   sycl::free(Ptr, Queue);

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

@@ -16,11 +16,9 @@
 
 int main() {
   queue Queue{};
-  const size_t N = 1024;
-  int *Ptr = (int *)sycl::malloc_device<int>(N, Queue);
-  std::vector<int> HostData(N, 0);
-  buffer<int> BufA{sycl::range<1>(N)};
-  buffer<int> BufB{sycl::range<1>(N)};
+  int *Ptr = (int *)sycl::malloc_device<int>(Size, Queue);
+  buffer<int> BufA{sycl::range<1>(Size)};
+  buffer<int> BufB{sycl::range<1>(Size)};
   BufA.set_write_back(false);
   BufB.set_write_back(false);
 
@@ -34,7 +32,7 @@ int main() {
   auto RootNode = Graph.add([&](handler &CGH) {
     auto AccA = BufA.get_access<access::mode::write>(CGH);
     auto AccB = BufB.get_access<access::mode::write>(CGH);
-    CGH.parallel_for(N, [=](item<1> Item) {
+    CGH.parallel_for(Size, [=](item<1> Item) {
       AccA[Item.get_id()] = InitA;
       AccB[Item.get_id()] = InitB;
     });
@@ -43,13 +41,15 @@ int main() {
   int PatternA = 42;
   auto CGFA = [&](handler &CGH) {
     auto AccA = BufA.get_access<access::mode::read_write>(CGH);
-    CGH.parallel_for(N, [=](item<1> Item) { AccA[Item.get_id()] += PatternA; });
+    CGH.parallel_for(Size,
+                     [=](item<1> Item) { AccA[Item.get_id()] += PatternA; });
   };
 
   int PatternB = 0xA;
   auto CGFB = [&](handler &CGH) {
     auto AccB = BufB.get_access<access::mode::read_write>(CGH);
-    CGH.parallel_for(N, [=](item<1> Item) { AccB[Item.get_id()] += PatternB; });
+    CGH.parallel_for(Size,
+                     [=](item<1> Item) { AccB[Item.get_id()] += PatternB; });
   };
 
   auto DynamicCG = exp_ext::dynamic_command_group(Graph, {CGFA, CGFB});
@@ -58,25 +58,28 @@ int main() {
   auto LeafNode = Graph.add([&](handler &CGH) {
     auto AccA = BufA.get_access<access::mode::read>(CGH);
     auto AccB = BufB.get_access<access::mode::read>(CGH);
-    CGH.parallel_for(N, [=](item<1> Item) {
+    CGH.parallel_for(Size, [=](item<1> Item) {
       Ptr[Item.get_id()] = AccA[Item.get_id()] + AccB[Item.get_id()];
     });
   });
   auto ExecGraph = Graph.finalize(exp_ext::property::graph::updatable{});
 
   Queue.ext_oneapi_graph(ExecGraph).wait();
-  Queue.copy(Ptr, HostData.data(), N).wait();
-  for (size_t i = 0; i < N; i++) {
+
+  std::vector<int> HostData(Size, 0);
+  Queue.copy(Ptr, HostData.data(), Size).wait();
+  for (size_t i = 0; i < Size; i++) {
     assert(HostData[i] == (InitA + InitB + PatternA));
   }
 
   DynamicCG.set_active_cgf(1);
   ExecGraph.update(DynamicCGNode);
 
   Queue.ext_oneapi_graph(ExecGraph).wait();
-  Queue.copy(Ptr, HostData.data(), N).wait();
-  for (size_t i = 0; i < N; i++) {
-    assert(HostData[i] == (InitA + InitB + PatternB));
+  Queue.copy(Ptr, HostData.data(), Size).wait();
+  int Ref = InitA + InitB + PatternB;
+  for (size_t i = 0; i < Size; i++) {
+    assert(HostData[i] == Ref);
   }
 
   sycl::free(Ptr, Queue);

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

@@ -26,17 +26,13 @@ int main(int, char **argv) {
   kernel kernel = getKernel(
       KernelBundle, "_ZTSZZ4mainENKUlRN4sycl3_V17handlerEE_clES2_EUlvE_");
 
-  const size_t N = 1024;
-
   exp_ext::command_graph Graph{
       Queue.get_context(),
       Queue.get_device(),
       {exp_ext::property::graph::assume_buffer_outlives_graph{}}};
-  std::vector<int> HostDataA(N, 0);
-  std::vector<int> HostDataB(N, 0);
 
-  buffer BufA{HostDataA};
-  buffer BufB{HostDataB};
+  buffer<int> BufA{sycl::range<1>(Size)};
+  buffer<int> BufB{sycl::range<1>(Size)};
   BufA.set_write_back(false);
   BufB.set_write_back(false);
 
@@ -60,9 +56,12 @@ int main(int, char **argv) {
   auto ExecGraph = Graph.finalize(exp_ext::property::graph::updatable{});
 
   Queue.ext_oneapi_graph(ExecGraph).wait();
+
+  std::vector<int> HostDataA(Size, 0);
+  std::vector<int> HostDataB(Size, 0);
   Queue.copy(BufA.get_access(), HostDataA.data()).wait();
   Queue.copy(BufB.get_access(), HostDataB.data()).wait();
-  for (size_t i = 0; i < N; i++) {
+  for (size_t i = 0; i < Size; i++) {
     assert(HostDataA[i] == i);
     assert(HostDataB[i] == 0);
   }
@@ -71,9 +70,10 @@ int main(int, char **argv) {
   ExecGraph.update(DynamicCGNode);
 
   Queue.ext_oneapi_graph(ExecGraph).wait();
+
   Queue.copy(BufA.get_access(), HostDataA.data()).wait();
   Queue.copy(BufB.get_access(), HostDataB.data()).wait();
-  for (size_t i = 0; i < N; i++) {
+  for (size_t i = 0; i < Size; i++) {
     assert(HostDataA[i] == i);
     assert(HostDataB[i] == i);
   }