Skip to content

Commit 42a9485

Browse files
authored
[SYCL][Graph] Enable host-task update in graphs (#16853)
- Update spec wording to allow updating host-task function in graphs - Support host-tasks in dynamic command-groups - Support host-tasks in whole graph update - Add E2E tests for both scenarios - Fix passing incorrect accessors to graph update command after update which could cause issues if the graph was destroyed while the scheduler still had graph execution commands alive
1 parent ba6cc2c commit 42a9485

12 files changed

+894
-111
lines changed

sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc

+65-27
Original file line numberDiff line numberDiff line change
@@ -551,7 +551,7 @@ Parameters:
551551

552552
|===
553553

554-
==== Dynamic Command Groups
554+
==== Dynamic Command Groups [[dynamic-command-groups]]
555555

556556
[source,c++]
557557
----
@@ -570,12 +570,13 @@ public:
570570
Dynamic command-groups can be added as nodes to a graph. They provide a
571571
mechanism that allows updating the command-group function of a node after the
572572
graph is finalized. There is always one command-group function in the dynamic
573-
command-group that is set as active, this is the kernel which will execute for
574-
the node when the graph is finalized into an executable state `command_graph`,
575-
and all the other command-group functions in `cgfList` will be ignored. The
576-
executable `command_graph` node can then be updated to a different kernel in
577-
`cgfList`, by selecting a new active index on the dynamic command-group object
578-
and calling the `update(node& node)` method on the executable `command_graph`.
573+
command-group that is set as active, this is the command-group which will
574+
execute for the node when the graph is finalized into an executable state
575+
`command_graph`, and all the other command-group functions in `cgfList` will be
576+
ignored. The executable `command_graph` node can then be updated to a different
577+
kernel in `cgfList`, by selecting a new active index on the dynamic
578+
command-group object and calling the `update(node& node)` method on the
579+
executable `command_graph`.
579580

580581
The `dynamic_command_group` class provides the {crs}[common reference semantics].
581582

@@ -584,9 +585,13 @@ about updating command-groups.
584585

585586
===== Limitations
586587

587-
Dynamic command-groups can only contain kernel operations. Trying to construct
588-
a dynamic command-group with functions that contain other operations will
589-
result in an error.
588+
Dynamic command-groups can only contain the following operations:
589+
590+
* Kernel operations
591+
* <<host-tasks, Host-tasks>>
592+
593+
Trying to construct a dynamic command-group with functions that contain other
594+
operations will result in an error.
590595

591596
All the command-group functions in a dynamic command-group must have identical dependencies.
592597
It is not allowed for a dynamic command-group to have command-group functions that would
@@ -625,10 +630,13 @@ Exceptions:
625630
property for more information.
626631

627632
* Throws with error code `invalid` if the `dynamic_command_group` is created with
628-
command-group functions that are not kernel executions.
633+
command-group functions that are not kernel executions or host-tasks.
629634

630635
* Throws with error code `invalid` if `cgfList` is empty.
631636

637+
* Throws with error code `invalid` if the types of all command-groups in
638+
`cgfList` do not match.
639+
632640
|
633641
[source,c++]
634642
----
@@ -829,32 +837,54 @@ possible.
829837

830838
===== Supported Features
831839

832-
The only types of nodes that are currently able to be updated in a graph are
833-
kernel execution nodes.
840+
The only types of nodes that are currently able to be updated in a graph are:
834841

835-
There are two different API's that can be used to update a graph:
842+
* Kernel executions
843+
* <<host-tasks, Host-tasks>>
844+
845+
There are two different APIs that can be used to update a graph:
836846

837847
* <<individual-node-update, Individual Node Update>> which allows updating
838848
individual nodes of a command-graph.
839849
* <<whole-graph-update, Whole Graph Update>> which allows updating the
840850
entirety of the graph simultaneously by using another graph as a
841851
reference.
842852

843-
The aspects of a kernel execution node that can be changed during update are
844-
different depending on the API used to perform the update:
853+
The following table illustrates the aspects of each supported node type that can be changed
854+
depending on the API used to perform the update.
855+
856+
Table {counter: tableNumber}. Graph update capabilites for supported node types.
857+
[cols="1,2a,2a"]
858+
|===
859+
|Node Type|<<individual-node-update, Individual Node Update>>|<<whole-graph-update, Whole Graph Update>>
860+
861+
|`node_type::kernel`
862+
|
863+
864+
* Kernel function
865+
* Kernel parameters
866+
* ND-range
867+
868+
|
869+
* Kernel parameters
870+
* ND-range
845871

846-
* For the <<individual-node-update, Individual Node Update>> API it's possible to update
847-
the kernel function, the parameters to the kernel, and the ND-range.
848-
* For the <<whole-graph-update, Whole Graph Update>> API, only the parameters of the kernel
849-
and the ND-range can be updated.
872+
|`node_type::host_task`
873+
|
874+
* Host-task function
875+
|
876+
* Host-task function
877+
878+
|===
850879

851880
===== Individual Node Update [[individual-node-update]]
852881

853-
Individual nodes of an executable graph can be updated directly. Depending on the attribute
854-
of the node that requires updating, different API's should be used:
882+
Individual nodes of an executable graph can be updated directly. Depending on the attribute or `node_type` of the node that requires updating, different API's should be used:
855883

856884
====== Parameter Updates
857885

886+
_Supported Node Types: Kernel_
887+
858888
Parameters to individual nodes in a graph in the `executable` state can be
859889
updated between graph executions using dynamic parameters. A `dynamic_parameter`
860890
object is created with a modifiable state graph and an initial value for the
@@ -884,6 +914,8 @@ will maintain the graphs data dependencies.
884914

885915
====== Execution Range Updates
886916

917+
_Supported Node Types: Kernel_
918+
887919
Another configuration that can be updated is the execution range of the
888920
kernel, this can be set through `node::update_nd_range()` or
889921
`node::update_range()` but does not require any prior registration.
@@ -897,10 +929,13 @@ code may be defined as operating in a different dimension.
897929

898930
====== Command Group Updates
899931

900-
The command-groups of a kernel node can be updated using dynamic command-groups.
901-
Dynamic command-groups allow replacing the command-group function of a kernel
902-
node with a different one. This effectively allows updating the kernel function
903-
and/or the kernel execution range.
932+
_Supported Node Types: Kernel, Host-task_
933+
934+
The command-groups of a kernel node can be updated using
935+
<<dynamic-command-groups, Dynamic Command-Groups>>. Dynamic command-groups allow
936+
replacing the command-group function of a kernel node with a different one. This
937+
effectively allows updating the kernel function and/or the kernel execution
938+
range.
904939

905940
Command-group updates are performed by creating an instance of the
906941
`dynamic_command_group` class. A dynamic command-group is created with a modifiable
@@ -1972,7 +2007,7 @@ Any code like this should be moved to a separate host-task and added to the
19722007
graph via the recording or explicit APIs in order to be compatible with this
19732008
extension.
19742009

1975-
=== Host Tasks
2010+
=== Host Tasks [[host-tasks]]
19762011

19772012
:host-task: https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#subsec:interfaces.hosttasks
19782013

@@ -1992,6 +2027,9 @@ auto node = graph.add([&](sycl::handler& cgh){
19922027
});
19932028
----
19942029

2030+
Host-tasks can be updated using <<executable-graph-update, Executable Graph Update>>.
2031+
2032+
19952033
=== Queue Behavior In Recording Mode
19962034

19972035
When a queue is placed in recording mode via a call to `command_graph::begin_recording`,

0 commit comments

Comments
 (0)