From 748d98a7b97da0beed785f1dcfb3164c623a91d3 Mon Sep 17 00:00:00 2001 From: Pablo Reble Date: Tue, 21 Mar 2023 14:14:19 -0700 Subject: [PATCH 1/5] add implementation for empty node --- .../sycl/ext/oneapi/experimental/graph.hpp | 2 +- sycl/source/detail/graph_impl.cpp | 20 +++++++++++++++++++ sycl/source/detail/graph_impl.hpp | 10 +++++++++- 3 files changed, 30 insertions(+), 2 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp index e695499bfd6bf..21920b1a6e234 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp @@ -55,7 +55,7 @@ class __SYCL_EXPORT command_graph { command_graph(const property_list &propList = {}); // Adding empty node with [0..n] predecessors: - node add(const std::vector &dep = {}); + node add(const std::vector &dep = {}) { return add_impl(dep); } // Adding device node: template node add(T cgf, const std::vector &dep = {}) { diff --git a/sycl/source/detail/graph_impl.cpp b/sycl/source/detail/graph_impl.cpp index 337486d6096fc..bb142953a2530 100644 --- a/sycl/source/detail/graph_impl.cpp +++ b/sycl/source/detail/graph_impl.cpp @@ -85,6 +85,26 @@ bool check_for_arg(const sycl::detail::ArgDesc &Arg, return SuccessorAddedDep; } +std::shared_ptr +graph_impl::add(const std::shared_ptr &Impl, + const std::vector> &Dep) { + const std::shared_ptr &NodeImpl = + std::make_shared(Impl); + + // TODO: Encapsulate in separate function to avoid duplication + if (!Dep.empty()) { + for (auto N : Dep) { + N->register_successor(NodeImpl); // register successor + this->remove_root(NodeImpl); // remove receiver from root node + // list + } + } else { + this->add_root(NodeImpl); + } + + return NodeImpl; +} + std::shared_ptr graph_impl::add(const std::shared_ptr &Impl, std::function CGF, diff --git a/sycl/source/detail/graph_impl.hpp b/sycl/source/detail/graph_impl.hpp index f913cae633954..6bfc274c35176 100644 --- a/sycl/source/detail/graph_impl.hpp +++ b/sycl/source/detail/graph_impl.hpp @@ -85,6 +85,9 @@ struct node_impl { sycl::event get_event(void) const { return MEvent; } + node_impl(const std::shared_ptr &Graph) + : MScheduled(false), MGraph(Graph) {} + node_impl( const std::shared_ptr &Graph, std::shared_ptr Kernel, @@ -119,7 +122,8 @@ struct node_impl { if (!Next->MScheduled) Next->topology_sort(Schedule); } - Schedule.push_front(std::shared_ptr(this)); + if (MKernel != nullptr) + Schedule.push_front(std::shared_ptr(this)); } bool has_arg(const sycl::detail::ArgDesc &Arg) { @@ -169,6 +173,10 @@ struct graph_impl { const std::vector &Args, const std::vector> &Dep = {}); + std::shared_ptr + add(const std::shared_ptr &Impl, + const std::vector> &Dep = {}); + graph_impl() : MFirst(true) {} /// Add a queue to the set of queues which are currently recording to this From dd322e7b2ef980ebe0c318575d65ca6f100c1e68 Mon Sep 17 00:00:00 2001 From: Pablo Reble Date: Tue, 28 Mar 2023 14:22:37 -0700 Subject: [PATCH 2/5] Adding test for empty node --- sycl/test/graph/graph-explicit-empty.cpp | 45 ++++++++++++++++++++++++ 1 file changed, 45 insertions(+) create mode 100644 sycl/test/graph/graph-explicit-empty.cpp diff --git a/sycl/test/graph/graph-explicit-empty.cpp b/sycl/test/graph/graph-explicit-empty.cpp new file mode 100644 index 0000000000000..5eca9ba2c282c --- /dev/null +++ b/sycl/test/graph/graph-explicit-empty.cpp @@ -0,0 +1,45 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out +#include +#include + +#include + +int main() { + + sycl::property_list properties{ + sycl::property::queue::in_order{}, + sycl::ext::oneapi::property::queue::lazy_execution{}}; + + sycl::queue q{sycl::gpu_selector_v, properties}; + + sycl::ext::oneapi::experimental::command_graph g; + + const size_t n = 10; + float *arr = sycl::malloc_device(n, q); + + auto start = g.add(); + + auto init = g.add([&](sycl::handler &h) { + h.parallel_for(sycl::range<1>{n}, [=](sycl::id<1> idx) { + size_t i = idx; + arr[i] = 0; + }); + }, {start}); + + auto empty = g.add({init}); + + g.add([&](sycl::handler &h) { + h.parallel_for(sycl::range<1>{n}, [=](sycl::id<1> idx) { + size_t i = idx; + arr[i] = 1; + }); + }, {empty}); + + auto executable_graph = g.finalize(q.get_context()); + + q.submit([&](sycl::handler &h) { h.ext_oneapi_graph(executable_graph); }).wait(); + + sycl::free(arr, q); + + return 0; +} From 40d79ff204709a37a2e8a3ea32ef2cd14c7d91d2 Mon Sep 17 00:00:00 2001 From: Pablo Reble Date: Tue, 28 Mar 2023 14:23:53 -0700 Subject: [PATCH 3/5] Adding missing function --- sycl/include/sycl/ext/oneapi/experimental/graph.hpp | 2 ++ sycl/source/detail/graph_impl.cpp | 12 ++++++++++++ 2 files changed, 14 insertions(+) diff --git a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp index 21920b1a6e234..f3fef118d3aa4 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/graph.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/graph.hpp @@ -110,6 +110,8 @@ class __SYCL_EXPORT command_graph { node add_impl(std::function cgf, const std::vector &dep); + node add_impl(const std::vector &dep); + template friend decltype(Obj::impl) sycl::detail::getSyclObjImpl(const Obj &SyclObject); diff --git a/sycl/source/detail/graph_impl.cpp b/sycl/source/detail/graph_impl.cpp index bb142953a2530..f1a5d07acdf3e 100644 --- a/sycl/source/detail/graph_impl.cpp +++ b/sycl/source/detail/graph_impl.cpp @@ -213,6 +213,18 @@ command_graph::command_graph( const sycl::property_list &) : impl(std::make_shared()) {} +template <> +node command_graph::add_impl( + const std::vector &Deps) { + std::vector> DepImpls; + for (auto &D : Deps) { + DepImpls.push_back(sycl::detail::getSyclObjImpl(D)); + } + + std::shared_ptr NodeImpl = impl->add(impl, DepImpls); + return sycl::detail::createSyclObjFromImpl(NodeImpl); +} + template <> node command_graph::add_impl( std::function CGF, const std::vector &Deps) { From b140df5da85c66e661dd59b479a0ec7968722d57 Mon Sep 17 00:00:00 2001 From: Pablo Reble Date: Thu, 30 Mar 2023 11:48:28 -0500 Subject: [PATCH 4/5] Update sycl/test/graph/graph-explicit-empty.cpp Co-authored-by: Ben Tracy --- sycl/test/graph/graph-explicit-empty.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/sycl/test/graph/graph-explicit-empty.cpp b/sycl/test/graph/graph-explicit-empty.cpp index 5eca9ba2c282c..d48b606fac3d6 100644 --- a/sycl/test/graph/graph-explicit-empty.cpp +++ b/sycl/test/graph/graph-explicit-empty.cpp @@ -1,5 +1,4 @@ // RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out -#include #include #include From f198c613bb5fe51b79364e203ea13507be3ed799 Mon Sep 17 00:00:00 2001 From: Pablo Reble Date: Fri, 31 Mar 2023 12:23:49 -0500 Subject: [PATCH 5/5] Apply suggestions from code review --- sycl/test/graph/graph-explicit-empty.cpp | 7 +++++++ 1 file changed, 7 insertions(+) diff --git a/sycl/test/graph/graph-explicit-empty.cpp b/sycl/test/graph/graph-explicit-empty.cpp index d48b606fac3d6..7917a98fdf2f4 100644 --- a/sycl/test/graph/graph-explicit-empty.cpp +++ b/sycl/test/graph/graph-explicit-empty.cpp @@ -14,6 +14,7 @@ int main() { sycl::ext::oneapi::experimental::command_graph g; const size_t n = 10; + float *h_arr = sycl::malloc_host(n, q); float *arr = sycl::malloc_device(n, q); auto start = g.add(); @@ -38,6 +39,12 @@ int main() { q.submit([&](sycl::handler &h) { h.ext_oneapi_graph(executable_graph); }).wait(); + q.memcpy(&(h_arr[0]), arr, n).wait(); + + for (int i = 0; i < n; i++) + assert(h_arr[i] == 1.0f); + + sycl::free(h_arr, q); sycl::free(arr, q); return 0;