@@ -280,14 +280,12 @@ dependencies in one of three ways. Firstly, through buffer accessors that
280
280
represent data dependencies between two command groups captured as nodes.
281
281
Secondly, by using the `handler::depends_on()` mechanism inside a command group
282
282
captured as a node. However, for an event passed to `handler::depends_on()` to
283
- create an edge, it must be an event returned from a queue
284
- submission captured by the same graph. Passing events from other sources (other
285
- graph submissions, regular SYCL submissions) will not create edges in the graph,
286
- but will create runtime dependencies for a graph node on those other events.
287
- `handler::depends_on()` can be used to express edges when a user is working with
288
- USM memory rather than SYCL buffers. Thirdly, for a graph recorded with an
289
- in-order queue, an edge is added automatically between two sequential command
290
- groups submitted to the in-order queue.
283
+ create an edge, it must be an event returned from a queue submission captured by
284
+ the same graph. Passing events from other sources (other graph submissions,
285
+ regular SYCL submissions) will not create edges in the graph, but will create
286
+ runtime dependencies for a graph node on those other events. Thirdly, for a
287
+ graph recorded with an in-order queue, an edge is added automatically between
288
+ two sequential command groups submitted to the in-order queue.
291
289
292
290
|===
293
291
@@ -478,7 +476,7 @@ Exceptions:
478
476
479
477
|===
480
478
481
- ==== Dynamic Parameters
479
+ ==== Dynamic Parameters [[dynamic-parameters]]
482
480
483
481
[source,c++]
484
482
----
@@ -569,9 +567,9 @@ class depends_on {
569
567
template<typename... NodeTN>
570
568
depends_on(NodeTN... nodes);
571
569
572
- depends_on(const event& ev );
570
+ depends_on(const event& depEvent );
573
571
574
- depends_on(const std::vector<event>& events );
572
+ depends_on(const std::vector<event>& depEvents );
575
573
};
576
574
}
577
575
----
@@ -580,7 +578,7 @@ The API for explicitly adding nodes to a `command_graph` includes a
580
578
`property_list` parameter. This extension defines the `depends_on` property to
581
579
be passed here. `depends_on` may be used in two ways:
582
580
583
- * Passing nodes from the same command_graph which will create dependencies and
581
+ * Passing nodes from the same ` command_graph` which will create dependencies and
584
582
graph edges between those nodes and the node being added.
585
583
586
584
* Passing SYCL events will create runtime dependencies for execution of the
@@ -723,14 +721,16 @@ Updates to a graph will be scheduled after any in-flight executions of the same
723
721
graph and will not affect previous submissions of the same graph. The user is
724
722
not required to wait on any previous submissions of a graph before updating it.
725
723
726
- The only type of nodes that are currently supported for updating in a graph are
727
- kernel execution nodes.
728
-
729
724
The aspects of a kernel execution node that can be configured during update are:
730
725
731
726
* Parameters to the kernel.
732
727
* Execution ND-Range of the kernel.
733
728
729
+ All node types may have the following aspects configured during update:
730
+
731
+ * Dependent events which were specified using <<dyanmic-events, Dynamic
732
+ Events>>.
733
+
734
734
To update an executable graph, the `property::graph::updatable` property must
735
735
have been set when the graph was created during finalization. Otherwise, an
736
736
exception will be thrown if a user tries to update an executable graph. This
@@ -777,7 +777,7 @@ will maintain the graphs data dependencies.
777
777
===== Node Event Dependency Update
778
778
779
779
Event dependencies for nodes can be updated using <<dynamic-events, Dynamic
780
- Events>> in a similar usage to Dynamic Parameters.
780
+ Events>> in a similar usage to <<dynamic-parameters, Dynamic Parameters>> .
781
781
782
782
Event updates are performed using a `dynamic_event` instance and calling
783
783
`dynamic_event::update()` to update all the associated event dependencies of
@@ -1125,6 +1125,20 @@ event get_event(const node& n);
1125
1125
only for the next execution of the graph. This event can be used as a dependency
1126
1126
in the same way as normal SYCL events.
1127
1127
1128
+ Constraints:
1129
+
1130
+ * This member function is only available when the `command_graph` state is
1131
+ `graph_state::executable`.
1132
+
1133
+ Parameters:
1134
+
1135
+ * `n` - The node to get the associated event for.
1136
+
1137
+ Exceptions:
1138
+
1139
+ * Throws synchronously with error code `invalid` if `n` is not a node within the
1140
+ graph.
1141
+
1128
1142
|===
1129
1143
1130
1144
Table {counter: tableNumber}. Member functions of the `command_graph` class for
@@ -1608,8 +1622,8 @@ different from the one with which the dynamic_parameter was created.
1608
1622
1609
1623
[source,c++]
1610
1624
----
1611
- namespace ext::oneapi::experimental{
1612
- class dynamic_event{
1625
+ namespace ext::oneapi::experimental {
1626
+ class dynamic_event {
1613
1627
dynamic_event();
1614
1628
dynamic_event(const event& ev);
1615
1629
@@ -1620,7 +1634,7 @@ namespace ext::oneapi::experimental{
1620
1634
1621
1635
Dynamic events represent SYCL events from outside of a given `command_graph`
1622
1636
(either obtained from normal SYCL operations or from another `command_graph`)
1623
- that nodes in that graph may depend on. Dynamic events also allow for these
1637
+ that nodes in that graph may depend on. The `dynamic_event` object enables these
1624
1638
dependent events to be updated between graph executions.
1625
1639
1626
1640
Dynamic events can be used to add dependencies to a graph node in the same way
@@ -2219,6 +2233,65 @@ node nodeA = myGraph.add([&](handler& cgh) {
2219
2233
dynParamAccessor.update(bufferB.get_access());
2220
2234
----
2221
2235
2236
+ === Dynamic Event Update
2237
+
2238
+ Example which shows create a dependency between two graphs using dynamic events
2239
+ which are then updated between executions.
2240
+
2241
+ [source,c++]
2242
+ ----
2243
+ ...
2244
+
2245
+ using namespace sycl;
2246
+ namespace sycl_ext = sycl::ext::oneapi::experimental;
2247
+
2248
+ queue myQueue;
2249
+ auto myContext = myQueue.get_context();
2250
+ auto myDevice = myQueue.get_device();
2251
+
2252
+ // Create the first graph
2253
+ sycl_ext::command_graph graphA(myContext, myDevice);
2254
+
2255
+ // Add a node to graphA
2256
+ sycl_ext::node nodeA = graphA.add(...);
2257
+
2258
+ sycl_ext::command_graph<sycl_ext::graph_state::executable> execGraphA =
2259
+ graphA.finalize();
2260
+
2261
+ // Get an event for the next execution of nodeA in execGraphA
2262
+ event eventA = execGraph.get_event(nodeA);
2263
+
2264
+ // Create the second graph
2265
+ sycl_ext::command_graph graphB(myContext, myDevice);
2266
+
2267
+ // Create the dynamic event which will wrap eventA and use it as a dependency
2268
+ // for a node in graphB
2269
+ sycl_ext::dynamic_event dynEvent{eventA};
2270
+
2271
+ auto nodeB = graphB.add([&](handler& CGH){
2272
+ CGH.depends_on(dynEvent);
2273
+ CGH.parallel_for(...);
2274
+ });
2275
+
2276
+ auto execGraphB = graphB.finalize({sycl_ext::property::graph::updatable});
2277
+
2278
+ // Submit both graphs for execution
2279
+ myQueue.submit(execGraphA);
2280
+ myQueue.submit(execGraphB);
2281
+ myQueue.wait_and_throw();
2282
+
2283
+ // Update the dynamic event with a new event frome execGraphA to reflect the
2284
+ // next execution
2285
+ dynEvent.update(execGraphA.get_event(nodeA));
2286
+ // Update execGraphB with the affected node to reflect those changes
2287
+ execGraphB.update(nodeB);
2288
+
2289
+ // Both graphs can now be executed again with updated dependencies between them.
2290
+ myQueue.submit(execGraphA);
2291
+ myQueue.submit(execGraphB);
2292
+ myQueue.wait_and_throw();
2293
+ ----
2294
+
2222
2295
== Future Direction [[future-direction]]
2223
2296
2224
2297
This section contains both features of the specification which have been
0 commit comments