diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 1cdcb2860432d..03cc6895c70c2 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -9,6 +9,7 @@ :encoding: utf-8 :lang: en :sectnums: +:sectnumlevels: 4 :dpcpp: pass:[DPC++] :blank: pass:[ +] @@ -224,6 +225,14 @@ Table {counter: tableNumber}. Terminology. | Edge | Dependency between commands as a happens-before relationship. +| Dynamic Parameter +| An instance of the `dynamic_parameter`, `dynamic_accessor`, `dynamic_local_accessor` or +`dynamic_work_group_memory` classes. Used to update the parameters of a graph node. + +| Dynamic Command Group +| An instance of the `dynamic_command_group` class. Used to update the command group +function of a graph node. + |=== ==== Explicit Graph Building API @@ -472,12 +481,75 @@ Exceptions: ==== Dynamic Parameters +The following classes provide a mechanism by which a parameter of a graph +node can be updated. They are arguments to a node's command-group which can +be updated by the user after the node has been added to a graph. The choice +of which class to use depends on the type of the argument that needs to be updated: + +[#parameter-update-classes] +- <>: used to update pointers +to a USM allocation, scalars passed by value, or instances of `raw_kernel_arg` +(as defined in the +link:../experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc[sycl_ext_oneapi_raw_kernel_arg] extension). + +- <>: used to update the accessors +to `sycl::buffer` arguments. + +- <>: used to update the +allocation size of `sycl::local_accessor` arguments. + +- <>: used to update +the number of elements in `work_group_memory` arguments (as defined in the +link:../experimental/sycl_ext_oneapi_work_group_memory.asciidoc[sycl_ext_oneapi_work_group_memory] +extension). + +For simplicity, in this document the classes mentioned above are referred to as +dynamic parameters. Any references to the term "dynamic parameter(s)" will apply +not only to the `dynamic_parameter` class, but also to the `dynamic_accessor`, +`dynamic_local_accessor` and `dynamic_work_group_memory` classes. + +Dynamic parameters can be registered with nodes in a modifiable graph. +Registration happens when a dynamic parameter is set as a parameter to +the kernel that the node represents. Dynamic parameters behave like regular +kernel parameters and can be set using the same mechanisms (e.g. using +lambda captures or `handler::set_arg()` if using free function kernels). It is +valid to use multiple dynamic parameters in the same node. + +After registration, updating the value of a dynamic parameter will be reflected +immediately in the modifiable graph which contains the node. These updated +nodes can then be passed to an executable graph to update it with new values. +See <> for more information +about updating node parameters. + +When writing kernels that use dynamic parameters, the underlying object that +the dynamic parameter represents can be obtained using the `get()` member +function of the dynamic parameter. This function must only be used within a +kernel and will throw an exception if used in host code. + +===== Dynamic Parameter Exceptions + +The following uses of dynamic parameters are invalid and will throw an +`errc::invalid` exception: + +* Registering a dynamic parameter in a command-group that was submitted +to a queue which is currently recording to a graph. +* Registering a dynamic parameter on a SYCL command-group submission that +is not associated with a graph. +* Registering a dynamic parameter with a Graph that is different from the +one passed to the dynamic parameter constructor. + + +===== The `dynamic_parameter` Class [[dynamic-parameter-class]] + [source,c++] ---- namespace ext::oneapi::experimental{ template class dynamic_parameter { public: + + static_assert(sycl::is_device_copyable_v); + dynamic_parameter(command_graph graph, const ValueT &initialValue); void update(const ValueT& newValue); @@ -485,30 +557,24 @@ public: } ---- -Dynamic parameters are arguments to a node's command-group which can be updated -by the user after the node has been added to a graph. Updating the value of a -dynamic parameter will be reflected in the modifiable graph which contains this -node. These updated nodes can then be passed to an executable graph to update -it with new values. +The `dynamic_parameter` class provides the {crs}[common reference semantics]. -The type of the underlying object a dynamic parameter represents is set at -compile time using a template parameter. This underlying type can be an -accessor, a pointer to a USM allocation, scalar passed by value, or a raw byte -representation of the argument. The raw byte representation is intended to -enable updating arguments set using -link:../experimental/sycl_ext_oneapi_raw_kernel_arg.asciidoc[sycl_ext_oneapi_raw_kernel_arg]. -Dynamic parameters are registered with nodes in a modifiable graph, with each -registration associating one or more node arguments to the dynamic parameter -instance. Registration happens inside the command-group that the node -represents, and is done when dynamic parameters are set as parameters to the -kernel using `handler::set_arg()`/`handler::set_args()`. It is valid for a node -argument to be registered with more than one dynamic parameter instance. +Table {counter: tableNumber}. Template parameters of the `dynamic_parameter` class. +[cols="2a,a"] +|=== +|Template Parameter|Description -See <> for more information -about updating node parameters. +| +ValueT +| +:device_copyable: https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec::device.copyable +The type of the underlying object that a dynamic parameter represents is set at +compile time using the `ValueT` parameter. This underlying type must be {device_copyable}[device +copyable]. +| -The `dynamic_parameter` class provides the {crs}[common reference semantics]. +|=== Table {counter: tableNumber}. Member functions of the `dynamic_parameter` class. [cols="2a,a"] @@ -549,6 +615,661 @@ Parameters: * `newValue` - Value to update the registered node parameters to. +| +[source,c++] +---- +ValueT& get(); +---- + +| Returns a reference to the underlying value that +this dynamic parameter represents. This value is registered with +this `dynamic_parameter` and will be updated whenever +`dynamic_parameter::update()` is called. It is an error to use +this function in host code. + +Returns: + +A reference to `ValueT` to use in the current kernel. + +Exceptions: + +Throws an exception with the `errc::invalid` error code if used +in host code. + +|=== + +===== The `dynamic_accessor` Class [[dynamic-accessor-class]] + +[source,c++] +---- +namespace ext::oneapi::experimental{ +template ? access_mode::read + : access_mode::read_write), + target AccessTarget = target::device> +class dynamic_accessor { +public: + + static_assert(AccessTarget == target::device); + + template + dynamic_accessor( + command_graph &graph, + buffer &bufferRef, + const property_list &propList = {}); + + template + dynamic_accessor( + command_graph &graph, + buffer &bufferRef, + const property_list &propList = {}); + + template + dynamic_accessor( + command_graph &graph, + buffer &bufferRef, + TagT tag, const property_list &propList = {}); + + template + dynamic_accessor( + command_graph &graph, + buffer &bufferRef, + range accessRange, + const property_list &propList = {}); + + template + dynamic_accessor( + command_graph &graph, + buffer &bufferRef, + range accessRange, TagT tag, + const property_list &propList = {}); + + template + dynamic_accessor( + command_graph &graph, + buffer &bufferRef, + range accessRange, + id accessOffset, + const property_list &propList = {}); + + template + dynamic_accessor( + command_graph &graph, + buffer &bufferRef, + range accessRange, + id accessOffset, + TagT tag, const property_list &propList = {}); + + template + void update( + buffer &newBufferRef); + + template + void update( + buffer &newBufferRef, + range accessRange); + + template + void update( + buffer &newBufferRef, + range accessRange, + id accessOffset); + + accessor get(); +}; +} +---- + +The `dynamic_accessor` class provides the {crs}[common reference semantics]. + +:acbt: https://www.khronos.org/registry/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:accessor.command.buffer.tags + +Table {counter: tableNumber}. Template parameters of the `dynamic_accessor` class. +[cols="2a,a"] +|=== +|Template Parameter|Description + +| +DataT +| +:data_types: https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#_data_type +Specifies the data type for the `accessor` represented by this `dynamic_accessor`. The +restrictions and behavior are analogous to those of the +{data_types}[`accessor` class]. + +| +Dimensions +| +The `Dimensions` must be 0, 1, 2 or 3. + +| +AccessMode +| +:access_modes: https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#_access_modes +Specifies the access mode for the `accessor` represented by this `dynamic_accessor`. +The restrictions and behavior are analogous to those of the {access_modes}[`accessor` class]. + +| +AccessTarget +| +The access target must be equal to `target::device`. +| + +|=== + +Table {counter: tableNumber}. Member functions of the `dynamic_accessor` class. +[cols="2a,a"] +|=== +|Member Function|Description + +| +[source,c++] +---- +template +dynamic_accessor( + command_graph &graph, + buffer &bufferRef, + const property_list &propList = {}) +---- +|Available only when `(Dimensions == 0)`. + +Constructs a dynamic accessor for accessing the first element of a buffer. +This object can be registered with command graph nodes. + +Parameters: + +* `graph` - Graph which will contain the nodes that use the dynamic accessor. +* `bufferRef` - The buffer to access. +* `propList` - List of properties for the underlying accessor. + +| +[source,c++] +---- +template +dynamic_accessor( + command_graph &graph, + buffer &bufferRef, + const property_list &propList = {}) +---- +|Available only when `(Dimensions > 0)`. + +Constructs a dynamic accessor for accessing a buffer. This object can be registered +with command graph nodes. + +Parameters: + +* `graph` - Graph which will contain the nodes that use the dynamic accessor. +* `bufferRef` - The buffer to access. +* `propList` - List of properties for the underlying accessor. + +| +[source,c++] +---- +template +dynamic_accessor( + command_graph &graph, + buffer &bufferRef, + TagT tag, const property_list &propList = {}) +---- +|Available only when `(Dimensions > 0)`. + +Constructs a dynamic accessor for accessing a buffer. This object can be registered +with command graph nodes. + +Parameters: + +* `graph` - Graph which will contain the nodes that use the dynamic accessor. +* `bufferRef` - The buffer to access. +* `tag` - A {acbt}[deduction tag] for the template arguments of the dynamic accessor. +* `propList` - List of properties for the underlying accessor. + +| +[source,c++] +---- +template +dynamic_accessor( + command_graph &graph, + buffer &bufferRef, + range accessRange, + const property_list &propList = {}) +---- +|Available only when `(Dimensions > 0)`. + +Constructs a dynamic accessor that is a ranged accessor, where the range starts +at the beginning of the buffer. This object can be registered with command graph nodes. + +Parameters: + +* `graph` - Graph which will contain the nodes that use the dynamic accessor. +* `bufferRef` - The buffer to access. +* `accessRange` - The access range for the ranged accessor. +* `propList` - List of properties for the underlying accessor. + +Exceptions: + +Throws an exception with the `errc::invalid` error code if `accessRange` exceeds +the range of `bufferRef` in any dimension. + +| +[source,c++] +---- +template +dynamic_accessor( + command_graph &graph, + buffer &bufferRef, + range accessRange, + TagT tag, const property_list &propList = {}) +---- +|Available only when `(Dimensions > 0)`. + +Constructs a dynamic accessor that is a ranged accessor, where the range starts +at the beginning of the buffer. This object can be registered with command graph nodes. + +Parameters: + +* `graph` - Graph which will contain the nodes that use the dynamic accessor. +* `bufferRef` - The buffer to access. +* `accessRange` - The access range for the ranged accessor. +* `tag` - A {acbt}[deduction tag] for the template arguments of the dynamic accessor. +* `propList` - List of properties for the underlying accessor. + +Exceptions: + +Throws an exception with the `errc::invalid` error code if `accessRange` exceeds +the range of `bufferRef` in any dimension. + +| +[source,c++] +---- +template +dynamic_accessor( + command_graph &graph, + buffer &bufferRef, + range accessRange, + id accessOffset, + const property_list &propList = {}) +---- +|Available only when `(Dimensions > 0)`. + +Constructs a dynamic accessor that is a ranged accessor, where the range starts +at an offset from the beginning of the buffer. This object can be registered +with command graph nodes. + +Parameters: + +* `graph` - Graph which will contain the nodes that use the dynamic accessor. +* `bufferRef` - The buffer to access. +* `accessRange` - The access range for the ranged accessor. +* `accessOffset` - The offset for the range. +* `propList` - List of properties for the underlying accessor. + +Exceptions: + +Throws an exception with the `errc::invalid` error code if the sum of +`accessRange` and `accessOffset` exceeds the range of `bufferRef` in any dimension. + +| +[source,c++] +---- +template +dynamic_accessor( + command_graph &graph, + buffer &bufferRef, + range accessRange, + id accessOffset, + TagT tag, const property_list &propList = {}) +---- +|Available only when `(Dimensions > 0)`. + +Constructs a dynamic accessor that is a ranged accessor, where the range starts +at an offset from the beginning of the buffer. This object can be registered +with command graph nodes. + +Parameters: + +* `graph` - Graph which will contain the nodes that use the dynamic accessor. +* `bufferRef` - The buffer to access. +* `accessRange` - The access range for the ranged accessor. +* `accessOffset` - The offset for the range. +* `tag` - A {acbt}[deduction tag] for the template arguments of the dynamic accessor. +* `propList` - List of properties for the underlying accessor. + +Exceptions: + +Throws an exception with the `errc::invalid` error code if the sum of +`accessRange` and `accessOffset` exceeds the range of `bufferRef` in any dimension. + +| +[source,c++] +---- +template +void update(buffer &newBufferRef); +---- + +|Updates the buffer that this dynamic accessor provides access to in all graph nodes +where it is registered. This call will eliminate any access range or access offset +attributes currently associated with the dynamic accessor. + +The new value will be reflected immediately in the modifiable +graph which contains the registered nodes. The new value will not be reflected in any +executable graphs created from that modifiable graph until `command_graph::update()` +is called passing the modified nodes, or a new executable graph is finalized from +the modifiable graph. + +It is not an error if `newBufferRef` is set to the current parameter value in any +registered nodes. + +Parameters: + +* `newBufferRef` - The new buffer that the dynamic accessor will provide access to +in all registered graph nodes. + +| +[source,c++] +---- +template +void update( + buffer &newBufferRef, + range newAccessRange); +---- + +|Updates the buffer that this dynamic accessor provides access to in all graph nodes +where it is registered. If an access offset attribute is currently associated with +this dynamic accessor, this call will eliminate it. + +The new value will be reflected immediately in the modifiable +graph which contains the registered nodes. The new value will not be reflected in any +executable graphs created from that modifiable graph until `command_graph::update()` +is called passing the modified nodes, or a new executable graph is finalized from +the modifiable graph. + +It is not an error if `newBufferRef` is set to the current parameter value in any +registered nodes. + +Parameters: + +* `newBufferRef` - The new buffer that the dynamic accessor will provide access to +in all registered graph nodes. +* `newAccessRange` - Access range for the accessor of the new buffer parameter. + +Exceptions: + +Throws an exception with the `errc::invalid` error code if `newAccessRange` exceeds +the range of `bufferRef` in any dimension. + +| +[source,c++] +---- +template +void update( + buffer &newBufferRef, + range newAccessRange, + id newAccessOffset); +---- + +|Updates the buffer that this dynamic accessor provides access to in all graph nodes +where it is registered. The new value will be reflected immediately in the modifiable +graph which contains the registered nodes. The new value will not be reflected in any +executable graphs created from that modifiable graph until `command_graph::update()` +is called passing the modified nodes, or a new executable graph is finalized from +the modifiable graph. + +It is not an error if `newBufferRef` is set to the current parameter value in any +registered nodes. + +Parameters: + +* `newBufferRef` - The new buffer that the dynamic accessor will provide access to +in all registered graph nodes. +* `newAccessRange` - Access range for the accessor of the new buffer parameter. +* `newAccessOffset` - Access offset for the range of the new buffer parameter. + +Exceptions: + +Throws an exception with the `errc::invalid` error code if the sum of +`newAccessRange` and `newAccessOffset` exceeds the range of `bufferRef` in any dimension. + +| +[source,c++] +---- +accessor get(); +---- + +| Returns a `sycl::accessor` to use in the current kernel. +The accessor is registered with this `dynamic_accessor` and +will be updated whenever `dynamic_accessor::update()` is called. +It is an error to use this function in host code. + +Returns: + +An instance of `sycl::accessor` to use in the current kernel. + +Exceptions: + +Throws an exception with the `errc::invalid` error code if used +in host code. + +|=== + +===== The `dynamic_local_accessor` Class [[dynamic-local-accessor-class]] + +[source,c++] +---- +namespace ext::oneapi::experimental{ +template +class dynamic_local_accessor { +public: + + static_assert(Dimensions > 0 && Dimensions <= 3); + + dynamic_local_accessor( + command_graph graph, + range allocationSize, + const property_list &propList = {}); + + void update(range newAllocationSize); + + local_accessor get(); +}; +} +---- + +The `dynamic_local_accessor` class provides the {crs}[common reference semantics]. + +Table {counter: tableNumber}. Template parameters of the `dynamic_local_accessor` class. +[cols="2a,a"] +|=== +|Template Parameter|Description + +| +DataT +| +:accessor_local: https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:accessor.local +The `DataT` parameter can be any C++ type that the device supports. Its +restrictions and behaviour are analogous to the ones for the +{accessor_local}[`local_accessor` class]. + +| +Dimensions +| + +The `Dimensions` must be 1, 2 or 3. +|=== + +Table {counter: tableNumber}. Member functions of the `dynamic_local_accessor` class. +[cols="2a,a"] +|=== +|Member Function|Description + +| +[source,c++] +---- +template +dynamic_local_accessor( + command_graph graph, + range allocationSize); +---- +|Constructs a dynamic local accessor object that can be registered with command graph. + +Parameters: + +* `graph` - Graph which will contain the nodes that use the dynamic local accessor. +* `allocationSize` - The size of the local accessor. + +| +[source,c++] +---- +void update(range newAllocationSize); +---- + +|Updates the allocation size of this dynamic local accessor in all graph nodes +where it is registered. This new size will be reflected immediately in the +modifiable graph which contains the registered nodes. The new size will not be +reflected in any executable graphs created from that modifiable graph until +`command_graph::update()` is called passing the modified nodes, or a new +executable graph is finalized from the modifiable graph. + +It is not an error if `newAllocationSize` is set to the current allocation size +of any registered nodes. + +Parameters: + +* `newAllocationSize` - The new allocation size that this dynamic local accessor will +use in all the graph nodes where it is registered. + +| +[source,c++] +---- +local_accessor get(); +---- + +| +Returns a `sycl::local_accessor` to use in the current kernel. +The accessor is registered with this `dynamic_local_accessor` and +will be updated whenever `dynamic_local_accessor::update()` is called. +It is an error to use this function in host code. + +Returns: + +An instance of `sycl::local_accessor` to use in the current kernel. + +Exceptions: + +Throws an exception with the `errc::invalid` error code if used +in host code. + +|=== + +===== The `dynamic_work_group_memory` Class [[dynamic-work-group-memory-class]] + +The `dynamic_work_group_memory` class is a wrapper for the +link:../experimental/sycl_ext_oneapi_work_group_memory.asciidoc[sycl_ext_oneapi_work_group_memory] +extension. It provides a mechanism to update the local memory size allocated in graph kernel nodes. + +[source,c++] +---- +namespace ext::oneapi::experimental{ +template +class dynamic_work_group_memory { +public: + + // Check that DataT is an unbounded array type. + static_assert(std::is_array_v && std::extent_v == 0); + + static_assert(std::is_same_v); + + dynamic_work_group_memory( + command_graph graph, + size_t num); + + void update(size_t newNum); + + work_group_memory get(); +}; +} +---- + +The `dynamic_work_group_memory` class provides the {crs}[common reference semantics]. + +Table {counter: tableNumber}. Template parameters of the `dynamic_work_group_memory` class. +[cols="2a,a"] +|=== +|Template Parameter|Description + +| +DataT +| +The type of the object created in device local memory. This type must be an unbounded +array of a type that is supported in device code. + +| +PropertyListT +| +The property list must be `empty_properties_t`. +|=== + +Table {counter: tableNumber}. Member functions of the `dynamic_work_group_memory` class. +[cols="2a,a"] +|=== +|Member Function|Description + +| +[source,c++] +---- +dynamic_work_group_memory( + command_graph graph, + size_t num); +---- +|Constructs a dynamic work group memory object that can be registered with command graph. + +Parameters: + +* `graph` - Graph which will contain the nodes that use the dynamic work group memory. +* `num` - The size of the first dimension of the unbounded array `DataT`. This +value is used when allocating local memory in all graph nodes where this object is +registered. + +| +[source,c++] +---- +void update(size_t newNum); +---- + +|Updates the size of the first dimension of the unbounded array `DataT`. The new value +will be used when allocating local memory in all graph nodes where this object is registered. The +update will be reflected immediately in the modifiable graph which contains the registered nodes. +The new value will not be reflected in any executable graphs created from that modifiable +graph until `command_graph::update()` is called passing the modified nodes, or a new +executable graph is finalized from the modifiable graph. + +It is not an error if `newNum` is set to its current value in any registered nodes. + +Parameters: + +* `newNum` - The new size for the first dimension of the unbounded array `DataT`. + +| +[source,c++] +---- +work_group_memory get(); +---- + +| +Returns a `work_group_memory` object to use in the current kernel. +The work group memory is registered with this `dynamic_work_group_memory` +and will be updated whenever `dynamic_work_group_memory::update()` is called. +It is an error to use this function in host code. + +Returns: + +An instance of `work_group_memory` to use in the current kernel. + +Exceptions: + +Throws an exception with the `errc::invalid` error code if used +in host code. + |=== ==== Dynamic Command Groups @@ -856,16 +1577,17 @@ of the node that requires updating, different API's should be used: ====== Parameter Updates Parameters to individual nodes in a graph in the `executable` state can be -updated between graph executions using dynamic parameters. A `dynamic_parameter` -object is created with a modifiable state graph and an initial value for the -parameter. Dynamic parameters can then be registered with nodes in that graph -when passed to calls to `set_arg()/set_args()`. - -Parameter updates are performed using a `dynamic_parameter` instance by calling -`dynamic_parameter::update()` to update all the parameters of nodes to which the -`dynamic_parameter` is registered. Updates will not affect any nodes which were -not registered, even if they use the same parameter value as a -`dynamic_parameter`. +updated between graph executions using <>. +A dynamic parameter is created with a modifiable state graph. Dynamic parameters +will be automatically registered with nodes in that graph when passed as +arguments to the kernel. + +Parameter updates are performed using an instance of one of the +<> by calling +their `update()` member function which updates all the nodes to which +they are registered. Updates will not affect any nodes which were not +registered, even if they use the same parameter value as a dynamic +parameter. Since the structure of the graph became fixed when finalizing, updating parameters on a node will not change the already defined dependencies between @@ -877,10 +1599,11 @@ behavior of a graph when executed. For example, if there are two nodes (NodeA and NodeB) which are connected by an edge due to a dependency on the same buffer, both nodes must have this buffer parameter updated to the new value. This maintains the correct -data dependency and prevents unexpected behavior. To achieve this, one -dynamic parameter for the buffer can be registered with all the nodes which -use the buffer as a parameter. Then a single `dynamic_parameter::update()` call -will maintain the graphs data dependencies. +data dependency and prevents unexpected behavior. To achieve this, the buffer +accessor should be registered for update with all the nodes which +use the buffer as a parameter. Since it is an accessor, the registration should +be done using an instance of the `dynamic_accessor` class. Then, a single `update()` +call will maintain the graph's data dependencies. ====== Execution Range Updates @@ -908,9 +1631,9 @@ state graph and a list of possible command-group functions. Command-group functi within a dynamic command-group can then be set to active by using the member function `dynamic_command_group::set_active_index()`. -Dynamic command-groups are compatible with dynamic parameters. This means that -dynamic parameters can be used in command-group functions that are part of -dynamic command-groups. Updates to such dynamic parameters will be reflected +Dynamic command-groups are compatible with <>. +This means that dynamic parameters can be used in command-group functions that are +part of dynamic command-groups. Updates to such dynamic parameters will be reflected in the command-group functions once they are activated. Note that the execution range is tied to the command-group, therefore updating @@ -971,10 +1694,10 @@ satisfy the conditions of topological identity results in undefined behaviour, as it may prevent the runtime from pairing nodes in the source and target graphs. -It is valid to use nodes that contain dynamic parameters in whole graph updates. -If a node containing a dynamic parameter is updated through the whole graph -update API, then any previous updates to the dynamic parameter will be reflected -in the new graph. +It is valid to use nodes that contain <> +in whole graph updates. If a node containing a dynamic parameter is updated +through the whole graph update API, then any previous updates to the +dynamic parameter will be reflected in the new graph. ==== Graph Properties [[graph-properties]] @@ -1206,7 +1929,8 @@ Exceptions: <> property for more information. * Throws with error code `invalid` if the type of the command-group is not a - kernel execution and a `dynamic_parameter` was registered inside `cgf`. + kernel execution and a `dynamic_parameter`, `dynamic_accessor`, + `dynamic_local_accessor` or `dynamic_work_group_memory` was registered inside `cgf`. | [source,c++] @@ -1863,66 +2587,6 @@ Exceptions: to a queue which is associated with a device or context that is different from the device and context used on creation of the graph. -| -[source,c++] ----- -template void -handler::require(ext::oneapi::experimental::dynamic_parameter< - accessor> - dynamicParamAcc) ----- - -|Requires access to a memory object associated with an accessor contained in a -dynamic parameter. - -Parameters: - -* `dynamicParamAcc` - The dynamic parameter which contains the accessor that is -required. - -Exceptions: - -* Throws synchronously with error code `invalid` if this function is called from -a command-group submitted to a queue with is currently recording to a graph. - -* Throws synchronously with error code `invalid` if this function is called from -a normal SYCL command-group submission. - -* Throws synchronously with error code `invalid` if the graph which will be -associated with the graph node resulting from this command-group submission is -different from the one with which `dynamicParameterAcc` was created. - -| -[source,c++] ----- -template -void handler::set_arg(int argIndex, - ext::oneapi::experimental::dynamic_parameter &dynamicParam); ----- - -|Sets an argument to a kernel based on the value inside a dynamic parameter, and -registers that dynamic parameter with the graph node encapsulating the -submission of the command-group that calls this function. - -Parameters: - -* `argIndex` - The index of the kernel argument. - -* `dynamicParam` - The dynamic parameter which contains the argument. - -Exceptions: - -* Throws synchronously with error code `invalid` if this function is called from -a command-group submitted to a queue with is currently recording to a graph. - -* Throws synchronously with error code `invalid` if this function is called from -a normal SYCL command-group submission. - -* Throws synchronously with error code `invalid` if the graph which will be -associated with the graph node resulting from this command-group submission is -different from the one with which the dynamic_parameter was created. - |=== === Thread Safety diff --git a/sycl/doc/syclgraph/SYCLGraphUsageGuide.md b/sycl/doc/syclgraph/SYCLGraphUsageGuide.md index a2ca77258bdd8..0f6c952ee2695 100644 --- a/sycl/doc/syclgraph/SYCLGraphUsageGuide.md +++ b/sycl/doc/syclgraph/SYCLGraphUsageGuide.md @@ -427,30 +427,124 @@ sycl::free(ptrX, myQueue); sycl::free(ptrY, myQueue); sycl::free(ptrZ, myQueue); sycl::free(ptrQ, myQueue); +``` + +Example snippet showing how to update USM memory parameters using the `parallel_for` +function and lambdas: + +```c++ +... + +using namespace sycl; +namespace sycl_ext = sycl::ext::oneapi::experimental; + +queue myQueue; +sycl_ext::command_graph myGraph{myQueue.get_context(), myQueue.get_device()}; + +const size_t size = 1024; +int *ptrX = malloc_shared(size, myQueue); +int *ptrY = malloc_shared(size, myQueue); + +sycl_ext::dynamic_parameter dynParam(myGraph, ptrX); + +sycl_ext::node kernelNode = myGraph.add([&](handler& cgh) { + cgh.parallel_for(range<1>(size), [=](item<1> Id) { + // Get the USM pointer to use in the kernel. + auto Ptr = dynParam.get(); + + auto LinID = Id.get_linear_id(); + Ptr[LinID] = 1; + }); +}); + +auto execGraph = myGraph.finalize({sycl_ext::property::graph::updatable}); +myQueue.ext_oneapi_graph(execGraph); + +// Change ptrX argument to ptrY. +dynParam.update(ptrY); + +// Update kernelNode in the executable graph with the new parameters +execGraph.update(kernelNode); +// Execute the graph again. +myQueue.ext_oneapi_graph(execGraph); +myQueue.wait(); + +sycl::free(ptrX, myQueue); +sycl::free(ptrY, myQueue); ``` -Example snippet showing how to use accessors with `dynamic_parameter` update: +Example snippet showing how to update accessors with the `dynamic_accessor` class +using a kernel bundle: ```c++ +... + sycl::buffer bufferA{...}; sycl::buffer bufferB{...}; -// Create graph dynamic parameter using a placeholder accessor, since the -// sycl::handler is not available here outside of the command-group scope. -sycl_ext::dynamic_parameter dynParamAccessor(myGraph, bufferA.get_access()); +// Create a dynamic accessor for bufferA. +sycl_ext::dynamic_accessor dynAccessor(myGraph, bufferA); sycl_ext::node kernelNode = myGraph.add([&](handler& cgh) { - // Require the accessor contained in the dynamic paramter - cgh.require(dynParamAccessor); - // Set the arg on the kernel using the dynamic parameter directly - cgh.set_args(dynParamAccessor); + // Set the dynamic accessor as an arg for the kernel. + cgh.set_args(dynAccessor); cgh.parallel_for(range {n}, builtinKernel); }); ... -// Update the dynamic parameter with a placeholder accessor from bufferB instead -dynParamAccessor.update(bufferB.get_access()); +// Update the dynamic accessor to access bufferB instead. +dynAccessor.update(bufferB); +``` + +### Dynamic Work Group Memory Update with Free Function Kernels + +Example snippet showing how to update work group memory using the +`dynamic_work_group_memory` class and usm memory using the +`dynamic_parameter` class while utilizing the free function kernel +extension: + +```cpp +using namespace sycl; +namespace sycl_ext = sycl::ext::oneapi::experimental; + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((exp_ext::nd_range_kernel<1>)) +void freeFunction(sycl_ext::dynamic_work_group_memory DynWorkGroupMemory, + sycl_ext::dynamic_parameter DynPtr) { + sycl_ext::work_group_memory LocalMem = DynWorkGroupMemory.get(); + int *&Ptr = DynPtr.get(); + + ... + // Use Ptr and LocalMem directly to perform calculations. +} + +void workGroupMemExample() { + queue Queue{}; + sycl_ext::command_graph Graph{Queue.get_context(), Queue.get_device()}; + int *PtrA = malloc_shared(Size, Queue); + int *PtrB = malloc_shared(Size, Queue); + + sycl_ext::dynamic_work_group_memory DynWorkGroupMemory{Graph, + LocalMemSize}; + sycl_ext::dynamic_parameter DynPtr(Graph, PtrA); + + // Free Function Kernel + kernel_id Kernel_id = sycl_ext::get_kernel_id(); + kernel Kernel = Bundle.get_kernel(Kernel_id); + + auto Node = Graph.add([&](handler &CGH) { + CGH.set_arg(0, DynWorkGroupMemory); + CGH.set_arg(1, DynPtr); + CGH.parallel_for(nd_range({Size}, {LocalMemSize}), Kernel); + }); + + ... + int NewLocalMemSize = 1024; + // Update the size of the work group memory in Node to NewLocalMemSize. + DynWorkGroupMemory.update(NewLocalMemSize); + // Update the memory used in Node to PtrB. + DynPtr.update(PtrB); +} ``` ### Dynamic Command Groups