[SYCL][Graph] Bugfix buffer_ordering: make buffers outlives graph (in…

…tel#11380) (#332)

`buffer_ordering.cpp` did not meet the Graph specification. Buffers must
outlive the Graph that uses them.
The buffer-ordering test has been updated to fulfill this property.

sycl/test-e2e/Graph/Inputs/buffer_ordering.cpp

@@ -15,84 +15,86 @@ int main() {
 
  queue Queue{{sycl::ext::intel::property::queue::no_immediate_command_list{}}};
 
- exp_ext::command_graph Graph{
- Queue.get_context(),
- Queue.get_device(),
- {exp_ext::property::graph::assume_buffer_outlives_graph{}}};
-
  const size_t N = 10;
  std::vector<int> Arr(N, 0);
 
  buffer<int> Buf{N};
  Buf.set_write_back(false);
 
- // Buffer elements set to 3
- Queue.submit([&](handler &CGH) {
- auto Acc = Buf.get_access(CGH);
- CGH.parallel_for(range<1>{N}, [=](id<1> idx) {
- size_t i = idx;
- Acc[i] = 3;
+ {
+ exp_ext::command_graph Graph{
+ Queue.get_context(),
+ Queue.get_device(),
+ {exp_ext::property::graph::assume_buffer_outlives_graph{}}};
+
+ // Buffer elements set to 3
+ Queue.submit([&](handler &CGH) {
+ auto Acc = Buf.get_access(CGH);
+ CGH.parallel_for(range<1>{N}, [=](id<1> idx) {
+ size_t i = idx;
+ Acc[i] = 3;
+ });
  });
- });
 
- add_node(Graph, Queue, [&](handler &CGH) {
- auto Acc = Buf.get_access(CGH);
- CGH.parallel_for(range<1>{N}, [=](id<1> idx) {
- size_t i = idx;
- Acc[i] += 2;
+ add_node(Graph, Queue, [&](handler &CGH) {
+ auto Acc = Buf.get_access(CGH);
+ CGH.parallel_for(range<1>{N}, [=](id<1> idx) {
+ size_t i = idx;
+ Acc[i] += 2;
+ });
  });
- });
 
- for (size_t i = 0; i < N; i++) {
- assert(Arr[i] == 0);
- }
-
- // Buffer elements set to 4
- Queue.submit([&](handler &CGH) {
- auto Acc = Buf.get_access(CGH);
- CGH.parallel_for(range<1>{N}, [=](id<1> idx) {
- size_t i = idx;
- Acc[i] += 1;
+ for (size_t i = 0; i < N; i++) {
+ assert(Arr[i] == 0);
+ }
+
+ // Buffer elements set to 4
+ Queue.submit([&](handler &CGH) {
+ auto Acc = Buf.get_access(CGH);
+ CGH.parallel_for(range<1>{N}, [=](id<1> idx) {
+ size_t i = idx;
+ Acc[i] += 1;
+ });
  });
- });
 
- auto ExecGraph = Graph.finalize();
+  auto ExecGraph = Graph.finalize();
 
- for (size_t i = 0; i < N; i++) {
- assert(Arr[i] == 0);
- }
+  for (size_t i = 0; i < N; i++) {
+  assert(Arr[i] == 0);
+  }
 
- // Buffer elements set to 8
- Queue.submit([&](handler &CGH) {
- auto Acc = Buf.get_access(CGH);
- CGH.parallel_for(range<1>{N}, [=](id<1> idx) {
- size_t i = idx;
- Acc[i] *= 2;
+ // Buffer elements set to 8
+ Queue.submit([&](handler &CGH) {
+ auto Acc = Buf.get_access(CGH);
+ CGH.parallel_for(range<1>{N}, [=](id<1> idx) {
+ size_t i = idx;
+ Acc[i] *= 2;
+ });
  });
- });
-
- // Buffer elements set to 10
- auto Event =
- Queue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(ExecGraph); });
-
- // Buffer elements set to 20
- Queue.submit([&](handler &CGH) {
- auto Acc = Buf.get_access(CGH);
- CGH.parallel_for(range<1>{N}, [=](id<1> idx) {
- size_t i = idx;
- Acc[i] *= 2;
+
+ // Buffer elements set to 10
+  auto Event =
+  Queue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(ExecGraph); });
+
+ // Buffer elements set to 20
+  Queue.submit([&](handler &CGH) {
+  auto Acc = Buf.get_access(CGH);
+  CGH.parallel_for(range<1>{N}, [=](id<1> idx) {
+  size_t i = idx;
+  Acc[i] *= 2;
+ });
  });
- });
 
- Event.wait();
- // Buffer elements set to 22
- Queue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(ExecGraph); });
+  Event.wait();
+  // Buffer elements set to 22
+  Queue.submit([&](handler &CGH) { CGH.ext_oneapi_graph(ExecGraph); });
 
- Queue.submit([&](handler &CGH) {
- auto Acc = Buf.get_access(CGH);
- CGH.copy(Acc, Arr.data());
- });
- Queue.wait();
+ Queue.submit([&](handler &CGH) {
+ auto Acc = Buf.get_access(CGH);
+ CGH.copy(Acc, Arr.data());
+ });
+ Queue.wait();
+ }
 
  for (size_t i = 0; i < N; i++) {
  assert(Arr[i] == 22);