From c986ec04772464fdb5be6701241878d8377224e0 Mon Sep 17 00:00:00 2001 From: Ewan Crawford Date: Mon, 21 Aug 2023 14:12:40 +0100 Subject: [PATCH] [SYCL][Doc] Update spec example Update the buffer example in the spec to respect the buffer restrictions we currently have. * Use lifetime properties in constructor * Use `set_write_back(false)` on buffers used, and use host accessors to check the result. * Make buffer lifetimes exceed that of the graph object. Addresses Lukas feedback on upstream PR https://github.com/intel/llvm/pull/10473#issuecomment-1682160365 --- .../sycl_ext_oneapi_graph.asciidoc | 44 +++++++++++++------ 1 file changed, 30 insertions(+), 14 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index f50498409d445..296c27b825cb3 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -476,7 +476,7 @@ node. [source,c++] ---- -namespace sycl::ext::oneapi::experimental::property::node +namespace sycl::ext::oneapi::experimental::property::node { class depends_on { public: template @@ -1376,14 +1376,24 @@ submitted in its entirety for execution via [source, c++] ---- using namespace sycl; + namespace sycl_ext = sycl::ext::oneapi::experimental; + queue q{default_selector{}}; - // New object representing graph of command-groups - ext::oneapi::experimental::command_graph graph(q.get_context(), q.get_device()); + // Lifetime of buffers must exceed the lifetime of graphs they are used in. + buffer bufferA{dataA.data(), range<1>{elements}}; + bufferA.set_write_back(false); + buffer bufferB{dataB.data(), range<1>{elements}}; + bufferB.set_write_back(false); + buffer bufferC{dataC.data(), range<1>{elements}}; + bufferC.set_write_back(false); + { - buffer bufferA{dataA.data(), range<1>{elements}}; - buffer bufferB{dataB.data(), range<1>{elements}}; - buffer bufferC{dataC.data(), range<1>{elements}}; + // New object representing graph of command-groups + sycl_ext::command_graph graph(q.get_context(), q.get_device(), + {sycl_ext::property::graph::assume_buffer_outlives_graph{}, + sycl_ext::property::graph::assume_data_outlives_buffer{}}); + // `q` will be put in the recording state where commands are recorded to // `graph` rather than submitted for execution immediately. @@ -1433,17 +1443,23 @@ submitted in its entirety for execution via // queue `q` will be returned to the executing state where commands are // submitted immediately for extension. graph.end_recording(); - } - // Finalize the modifiable graph to create an executable graph that can be - // submitted for execution. - auto exec_graph = graph.finalize(); + // Finalize the modifiable graph to create an executable graph that can be + // submitted for execution. + auto exec_graph = graph.finalize(); - // Execute graph - q.submit([&](handler& cgh) { - cgh.ext_oneapi_graph(exec_graph); - }); + // Execute graph + q.submit([&](handler& cgh) { + cgh.ext_oneapi_graph(exec_graph); + }).wait(); + } + + // Check output using host accessors + host_accessor hostAccA(bufferA); + host_accessor hostAccB(bufferB); + host_accessor hostAccC(bufferC); + ... ---- == Future Direction [[future-direction]]