Skip to content
Merged
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
56 changes: 37 additions & 19 deletions sycl/doc/extensions/proposed/sycl_ext_oneapi_graph.asciidoc
Original file line number Diff line number Diff line change
Expand Up @@ -262,6 +262,18 @@ enum class queue_state {
recording
};

namespace property {
namespace node {

class depends_on {
public:
template<typename... NodeTN>
depends_on(NodeTN... nodes);
};

} // namespace node
} // namespace property

class node {};

// State of a graph
Expand All @@ -280,13 +292,13 @@ public:
command_graph(const property_list &propList = {});
command_graph<graph_state::executable> finalize(const context &syclContext) const;

node add(const std::vector<node>& dep = {});
node add(const property_list& propList = {});

template<typename T>
node add(T cgf, const std::vector<node>& dep = {});
node add(T cgf, const property_list& propList = {});

node add_malloc_device(void *&data, size_t numBytes, const std::vector<node>& dep = {});
node add_free(void *data, const std::vector<node>& dep = {});
node add_malloc_device(void *&data, size_t numBytes, const property_list& propList = {});
node add_free(void *data, const property_list& propList = {});

void make_edge(node sender, node receiver);
};
Expand Down Expand Up @@ -411,7 +423,7 @@ Table 6. Member functions of the `command_graph` class.
[source,c++]
----
using namespace ext::oneapi::experimental;
node add(const std::vector<node>& dep = {});
node add(const property_list& propList = {});
----
|This creates an empty node which is associated to no task. Its intended use is
either a connection point inside a graph between groups of nodes, and can
Expand All @@ -425,7 +437,8 @@ Preconditions:

Parameters:

* `dep` - Nodes the created node will be dependent on.
* `propList` - Zero or more properties can be provided to the constructed node
via an instance of `property_list`.

Returns: The empty node which has been added to the graph.

Expand All @@ -439,7 +452,7 @@ Exceptions:
----
using namespace ext::oneapi::experimental;
template<typename T>
node add(T cgf, const std::vector<node>& dep = {});
node add(T cgf, const property_list& propList = {});
----
|This function adds a command group function object to a graph. The function
object can contain single or multiple commands such as a host task which is
Expand All @@ -455,7 +468,8 @@ Parameters:

* `cgf` - Command group function object to be added as a node

* `dep` - Nodes the created node will be dependent on.
* `propList` - Zero or more properties can be provided to the constructed node
via an instance of `property_list`.

Returns: The command-group function object node which has been added to the graph.

Expand Down Expand Up @@ -532,7 +546,7 @@ Table 7. Member functions of the `command_graph` class (memory operations).
[source,c++]
----
using namespace ext::oneapi::experimental;
node add_malloc_device(void *&data, size_t numBytes, const std::vector<node>& dep = {});
node add_malloc_device(void *&data, size_t numBytes, const property_list& propList = {});
----
|Adding a node that encapsulates a memory allocation operation.

Expand All @@ -547,7 +561,8 @@ Parameters:

* `numBytes` - Size in bytes to allocate.

* `dep` - Nodes the created node will be dependent on.
* `propList` - Zero or more properties can be provided to the constructed node
via an instance of `property_list`.

Returns: The memory allocation node which has been added to the graph

Expand All @@ -560,7 +575,7 @@ Exceptions:
[source,c++]
----
using namespace ext::oneapi::experimental;
node add_free(void *data, const std::vector<node>& dep = {});
node add_free(void *data, const property_list& propList = {});
----
|Adding a node that encapsulates a memory free operation.

Expand All @@ -573,7 +588,8 @@ Parameters:

* `data` - Address of memory to free.

* `dep` - Nodes the created node will be dependent on.
* `propList` - Zero or more properties can be provided to the constructed node
via an instance of `property_list`.

Returns: The memory freeing node which has been added to the graph.

Expand Down Expand Up @@ -843,14 +859,16 @@ implementation is currently under development.
#include <sycl/ext/oneapi/experimental/graph.hpp>

int main() {
using namespace sycl::ext::oneapi::experimental;

const size_t n = 10;
float alpha = 1.0f;
float beta = 2.0f;
float gamma = 3.0f;

sycl::queue q;

sycl::ext::oneapi::experimental::command_graph g;
command_graph g;

float *x , *y, *z;

Expand Down Expand Up @@ -880,21 +898,21 @@ int main() {
y[i] = 2.0f;
z[i] = 3.0f;
});
}, {node_x, node_y, node_z});
}, { property::node::depends_on(node_x, node_y, node_z)});

auto node_a = g.add([&](sycl::handler &h) {
h.parallel_for(sycl::range<1>{n}, [=](sycl::id<1> it) {
const size_t i = it[0];
x[i] = alpha * x[i] + beta * y[i];
});
}, {node_i});
}, { property::node::depends_on(node_i)});

auto node_b = g.add([&](sycl::handler &h) {
h.parallel_for(sycl::range<1>{n}, [=](sycl::id<1> it) {
const size_t i = it[0];
z[i] = gamma * z[i] + beta * y[i];
});
}, {node_i});
}, { property::node::depends_on(node_i)});

auto node_c = g.add(
[&](sycl::handler &h) {
Expand All @@ -905,10 +923,10 @@ int main() {
sum += x[i] * z[i];
});
},
{node_a, node_b});
{ property::node::depends_on(node_a, node_b)});

auto node_fx = g.add_free(x, {node_c});
auto node_fy = g.add_free(y, {node_b});
auto node_fx = g.add_free(x, {property::node::depends_on(node_c)});
auto node_fy = g.add_free(y, {property::node::depends_on(node_b)});

auto exec = g.finalize(q.get_context());

Expand Down