From 372f06fc3dc0f85b55af5bd82a8b1ca78c4c4ebd Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio=20Mestre?= Date: Tue, 17 Dec 2024 17:51:15 +0000 Subject: [PATCH 1/6] Add local memory parameter update functionality to sycl graphs. Updates the sycl graph specification to add the dynamic_accessor, dynamic_local_accessor and dynamic_work_group_memory classes. This adds the required functionality to support updating local memory parameters to sycl graph kernel nodes. Additionally, it also moves the accessor update functionality from the dynamic_parameter class to the new dynamic_accessor class. This improves the cohesion of the API and removes the need to use placeholder accessors when updating buffer arguments in sycl graphs. --- .../sycl_ext_oneapi_graph.asciidoc | 776 ++++++++++++++++-- 1 file changed, 712 insertions(+), 64 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 1cdcb2860432d..563f750fe9f3f 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,6 +481,55 @@ 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, +with each registration associating one or more node arguments to the class instance. +Registration happens inside the command-group that the node +represents, and is done when the dynamic parameter is set as a parameter 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. + +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. + +===== The dynamic_parameter class [[dynamic-parameter-class]] + +The type of the underlying object a dynamic parameter represents is set at +compile time using a template parameter. This underlying type can be 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]. + [source,c++] ---- namespace ext::oneapi::experimental{ @@ -485,29 +543,6 @@ 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 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. - -See <> for more information -about updating node parameters. - The `dynamic_parameter` class provides the {crs}[common reference semantics]. Table {counter: tableNumber}. Member functions of the `dynamic_parameter` class. @@ -551,6 +586,548 @@ Parameters: |=== +===== 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: + 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(handler &cgh); +}; +} +---- + +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}. 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 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. This 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. This 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(handler &cgh); +---- + +| Returns a `sycl::accessor` to use in the current command-group. +The accessor is registered with this `dynamic_accessor` instance +and will be updated whenever `dynamic_accessor::update()` is called. + +Parameters: + +* `cgh` - The kernel handler that represents the current submission. + +Returns: + +An instance of `sycl::accessor` to use in the current command-group. + +|=== + +===== The dynamic_local_accessor class [[dynamic-local-accessor-class]] + +[source,c++] +---- +namespace ext::oneapi::experimental{ +template +class dynamic_local_accessor { +public: + dynamic_local_accessor( + command_graph graph, + range allocationSize, + const property_list &propList = {}); + + void update(range newAllocationSize); + + local_accessor get(handler &cgh); +}; +} +---- + +The `dynamic_local_accessor` class provides the {crs}[common reference semantics]. + +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); +---- +|Available only when `(Dimensions > 0)`. + +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(handler &cgh); +---- + +| Returns a `sycl::local_accessor` to use in the current command-group. +The local accessor is registered with this `dynamic_local_accessor` +instance and will be updated whenever `dynamic_local_accessor::update()` +is called. + +Parameters: + +* `cgh` - The kernel handler that represents the current submission. + +Returns: + +An instance of `sycl::local_accessor` to use in the current command-group. + +|=== + +===== 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: + dynamic_work_group_memory( + command_graph graph, + size_t num); + + void update(size_t newNum); + + work_group_memory get(handler &cgh); +}; +} +---- + +The `dynamic_work_group_memory` class provides the {crs}[common reference semantics]. + +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 number of `DataT` elements in the dynamic work group memory. + +| +[source,c++] +---- +void update(size_t newNum); +---- + +|Updates the number of `DataT` elements that this dynamic work group memory will +allocate memory for in all graph nodes where it is registered. This 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 `newNum` is set to its current value in any registered nodes. + +Parameters: + +* `newNum` - The new number of `DataT` elements that this dynamic work group memory +will allocate memory for in all the graph nodes where it is registered. + +| +[source,c++] +---- +work_group_memory get(handler &cgh); +---- + +| Returns a `work_group_memory` object to use in the current command-group. +The work group memory is registered with this `dynamic_work_group_memory` +instance and will be updated whenever `dynamic_work_group_memory::update()` +is called. + +Parameters: + +* `cgh` - The kernel handler that represents the current submission. + +Returns: + +An instance of `ext::oneapi::experimental::work_group_memory` to use +in the current command-group. + +|=== + ==== Dynamic Command Groups [source,c++] @@ -856,16 +1433,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 +can then be registered with nodes in that graph when passed to calls to +`set_arg()/set_args()`. + +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 +1455,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 +1487,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 +1550,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 +1785,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++] @@ -1866,62 +2446,130 @@ Exceptions: | [source,c++] ---- -template void -handler::require(ext::oneapi::experimental::dynamic_parameter< - accessor> - dynamicParamAcc) +template +void handler::set_arg(int argIndex, + ext::oneapi::experimental::dynamic_parameter &dynamicParam); ---- -|Requires access to a memory object associated with an accessor contained in a -dynamic parameter. +|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: -* `dynamicParamAcc` - The dynamic parameter which contains the accessor that is -required. +* `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. +a command-group submitted to a queue which 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. +different from the one with which the dynamic_parameter was created. | [source,c++] ---- -template +template +void handler::set_arg( + int argIndex, ext::oneapi::experimental::dynamic_accessor + &dynamicAccessor); +---- + +|Sets a memory object argument to a kernel based on the value represented +by a dynamic accessor, and registers that dynamic accessor with the graph +node encapsulating the submission of the command-group that calls this function. + +Parameters: + +* `argIndex` - The index of the kernel argument. + +* `dynamicAccessor` - The dynamic accessor that represents the memory +object argument. + +Exceptions: + +* Throws synchronously with error code `invalid` if this function is called from +a command-group submitted to a queue which 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 `dynamicAccessor` was created. + +| +[source,c++] +---- +template void handler::set_arg(int argIndex, - ext::oneapi::experimental::dynamic_parameter &dynamicParam); + ext::oneapi::experimental::dynamic_local_accessor + &dynamicLocalAccessor); ---- -|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. +|Sets a local memory argument to a kernel based on a dynamic local accessor, and +registers that dynamic accessor 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. +* `dynamicLocalAccessor` - The dynamic local accessor that represents the local +memory 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. +a command-group submitted to a queue which 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. +different from the one with which `dynamicLocalAccessor` was created. + +| +[source,c++] +---- +template +void handler::set_arg(int argIndex, + ext::oneapi::experimental::dynamic_work_group_memory + &dynamicWorkGroupMemory); +---- + +|Sets a local memory argument to a kernel based on a `dynamic_work_group_memory` +object, and registers that dynamic work group memory with the graph node +encapsulating the submission of the command-group that calls this function. + +Parameters: + +* `argIndex` - The index of the kernel argument. + +* `dynamicWorkGroupMemory` - The dynamic work group memory that represents the +local memory argument. + +Exceptions: + +* Throws synchronously with error code `invalid` if this function is called from +a command-group submitted to a queue which 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 `dynamicWorkGroupMemory` was created. |=== From c4f6fe2e36095a0f52af23c66f20baf7b46008a9 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio=20Mestre?= Date: Fri, 14 Feb 2025 12:33:20 +0000 Subject: [PATCH 2/6] Address Review comments: - Specify new usage for dynamic parameters after compiler support is added. - Update get() function to be used only in device code. - Remove set_arg() overloads - Clarify template parameters limitations and add static_asserts to class definition. - Fix wording of work_group_memory parameters. --- .../sycl_ext_oneapi_graph.asciidoc | 388 +++++++++--------- 1 file changed, 201 insertions(+), 187 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 563f750fe9f3f..acb9a99693d27 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -503,18 +503,17 @@ 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, -with each registration associating one or more node arguments to the class instance. -Registration happens inside the command-group that the node -represents, and is done when the dynamic parameter is set as a parameter 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. +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 @@ -522,13 +521,25 @@ nodes can then be passed to an executable graph to update it with new values. See <> for more information about updating node parameters. -===== The dynamic_parameter class [[dynamic-parameter-class]] +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. -The type of the underlying object a dynamic parameter represents is set at -compile time using a template parameter. This underlying type can be 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 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++] ---- @@ -545,6 +556,24 @@ public: The `dynamic_parameter` class provides the {crs}[common reference semantics]. + +Table {counter: tableNumber}. Template parameters of the `dynamic_parameter` class. +[cols="2a,a"] +|=== +|Template Parameter|Description + +| +ValueT +| +The type of the underlying object that a dynamic parameter represents is set at +compile time using the `ValueT` parameter. This underlying type can be 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]. +| + +|=== + Table {counter: tableNumber}. Member functions of the `dynamic_parameter` class. [cols="2a,a"] |=== @@ -584,9 +613,30 @@ 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]] +===== The `dynamic_accessor` Class [[dynamic-accessor-class]] [source,c++] ---- @@ -598,6 +648,9 @@ template class dynamic_accessor { public: + + static_assert(AccessTarget == target::device); + template dynamic_accessor( command_graph &graph, @@ -661,7 +714,7 @@ public: range accessRange, id accessOffset); - accessor get(handler &cgh); + accessor get(); }; } ---- @@ -670,6 +723,39 @@ 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"] |=== @@ -860,7 +946,10 @@ void update(buffer &newBufferRef); ---- |Updates the buffer that this dynamic accessor provides access to in all graph nodes -where it is registered. This new value will be reflected immediately in the modifiable +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 @@ -884,7 +973,10 @@ void update( ---- |Updates the buffer that this dynamic accessor provides access to in all graph nodes -where it is registered. This new value will be reflected immediately in the modifiable +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 @@ -915,7 +1007,7 @@ void update( ---- |Updates the buffer that this dynamic accessor provides access to in all graph nodes -where it is registered. This new value will be reflected immediately in the modifiable +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 @@ -939,24 +1031,26 @@ Throws an exception with the `errc::invalid` error code if the sum of | [source,c++] ---- -accessor get(handler &cgh); +accessor get(); ---- -| Returns a `sycl::accessor` to use in the current command-group. -The accessor is registered with this `dynamic_accessor` instance -and will be updated whenever `dynamic_accessor::update()` is called. +| 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. -Parameters: +Returns: -* `cgh` - The kernel handler that represents the current submission. +An instance of `sycl::accessor` to use in the current kernel. -Returns: +Exceptions: -An instance of `sycl::accessor` to use in the current command-group. +Throws an exception with the `errc::invalid` error code if used +in host code. |=== -===== The dynamic_local_accessor class [[dynamic-local-accessor-class]] +===== The `dynamic_local_accessor` Class [[dynamic-local-accessor-class]] [source,c++] ---- @@ -964,6 +1058,9 @@ namespace ext::oneapi::experimental{ template class dynamic_local_accessor { public: + + static_assert(Dimensions > 0 && Dimensions <= 3); + dynamic_local_accessor( command_graph graph, range allocationSize, @@ -971,13 +1068,33 @@ public: void update(range newAllocationSize); - local_accessor get(handler &cgh); + 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"] |=== @@ -991,9 +1108,7 @@ dynamic_local_accessor( command_graph graph, range allocationSize); ---- -|Available only when `(Dimensions > 0)`. - -Constructs a dynamic local accessor object that can be registered with command graph. +|Constructs a dynamic local accessor object that can be registered with command graph. Parameters: @@ -1024,25 +1139,27 @@ use in all the graph nodes where it is registered. | [source,c++] ---- -local_accessor get(handler &cgh); +local_accessor get(); ---- -| Returns a `sycl::local_accessor` to use in the current command-group. -The local accessor is registered with this `dynamic_local_accessor` -instance and will be updated whenever `dynamic_local_accessor::update()` -is called. +| +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. -Parameters: +Returns: -* `cgh` - The kernel handler that represents the current submission. +An instance of `sycl::local_accessor` to use in the current kernel. -Returns: +Exceptions: -An instance of `sycl::local_accessor` to use in the current command-group. +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 [[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] @@ -1054,19 +1171,42 @@ 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(handler &cgh); + 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"] |=== @@ -1084,7 +1224,9 @@ dynamic_work_group_memory( Parameters: * `graph` - Graph which will contain the nodes that use the dynamic work group memory. -* `num` - The number of `DataT` elements in 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++] @@ -1092,9 +1234,9 @@ Parameters: void update(size_t newNum); ---- -|Updates the number of `DataT` elements that this dynamic work group memory will -allocate memory for in all graph nodes where it is registered. This new value will -be reflected immediately in the modifiable graph which contains the registered nodes. +|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. @@ -1103,28 +1245,28 @@ It is not an error if `newNum` is set to its current value in any registered nod Parameters: -* `newNum` - The new number of `DataT` elements that this dynamic work group memory -will allocate memory for in all the graph nodes where it is registered. +* `newNum` - The new size for the first dimension of the unbounded array `DataT`. | [source,c++] ---- -work_group_memory get(handler &cgh); +work_group_memory get(); ---- -| Returns a `work_group_memory` object to use in the current command-group. +| +Returns a `work_group_memory` object to use in the current kernel. The work group memory is registered with this `dynamic_work_group_memory` -instance and will be updated whenever `dynamic_work_group_memory::update()` -is called. +and will be updated whenever `dynamic_work_group_memory::update()` is called. +It is an error to use this function in host code. -Parameters: +Returns: -* `cgh` - The kernel handler that represents the current submission. +An instance of `work_group_memory` to use in the current kernel. -Returns: +Exceptions: -An instance of `ext::oneapi::experimental::work_group_memory` to use -in the current command-group. +Throws an exception with the `errc::invalid` error code if used +in host code. |=== @@ -1435,8 +1577,8 @@ of the node that requires updating, different API's should be used: Parameters to individual nodes in a graph in the `executable` state can be updated between graph executions using <>. A dynamic parameter is created with a modifiable state graph. Dynamic parameters -can then be registered with nodes in that graph when passed to calls to -`set_arg()/set_args()`. +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 @@ -2443,134 +2585,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::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 which 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. - -| -[source,c++] ----- -template -void handler::set_arg( - int argIndex, ext::oneapi::experimental::dynamic_accessor - &dynamicAccessor); ----- - -|Sets a memory object argument to a kernel based on the value represented -by a dynamic accessor, and registers that dynamic accessor with the graph -node encapsulating the submission of the command-group that calls this function. - -Parameters: - -* `argIndex` - The index of the kernel argument. - -* `dynamicAccessor` - The dynamic accessor that represents the memory -object argument. - -Exceptions: - -* Throws synchronously with error code `invalid` if this function is called from -a command-group submitted to a queue which 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 `dynamicAccessor` was created. - -| -[source,c++] ----- -template -void handler::set_arg(int argIndex, - ext::oneapi::experimental::dynamic_local_accessor - &dynamicLocalAccessor); ----- - -|Sets a local memory argument to a kernel based on a dynamic local accessor, and -registers that dynamic accessor with the graph node encapsulating the submission -of the command-group that calls this function. - -Parameters: - -* `argIndex` - The index of the kernel argument. - -* `dynamicLocalAccessor` - The dynamic local accessor that represents the local -memory argument. - -Exceptions: - -* Throws synchronously with error code `invalid` if this function is called from -a command-group submitted to a queue which 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 `dynamicLocalAccessor` was created. - -| -[source,c++] ----- -template -void handler::set_arg(int argIndex, - ext::oneapi::experimental::dynamic_work_group_memory - &dynamicWorkGroupMemory); ----- - -|Sets a local memory argument to a kernel based on a `dynamic_work_group_memory` -object, and registers that dynamic work group memory with the graph node -encapsulating the submission of the command-group that calls this function. - -Parameters: - -* `argIndex` - The index of the kernel argument. - -* `dynamicWorkGroupMemory` - The dynamic work group memory that represents the -local memory argument. - -Exceptions: - -* Throws synchronously with error code `invalid` if this function is called from -a command-group submitted to a queue which 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 `dynamicWorkGroupMemory` was created. - |=== === Thread Safety From 7361cd4784c7dc56a90559f882a35bc84c8d8dc4 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio=20Mestre?= Date: Fri, 14 Feb 2025 18:41:36 +0000 Subject: [PATCH 3/6] Update the usage guide with the new dynamic classes --- sycl/doc/syclgraph/SYCLGraphUsageGuide.md | 114 ++++++++++++++++++++-- 1 file changed, 104 insertions(+), 10 deletions(-) diff --git a/sycl/doc/syclgraph/SYCLGraphUsageGuide.md b/sycl/doc/syclgraph/SYCLGraphUsageGuide.md index a2ca77258bdd8..a510ca0450891 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 Queue; +sycl_ext::command_graph Graph{Queue.get_context(), Queue.get_device()}; + +const size_t Size = 1024; +int *ptrX = malloc_shared(Size, Queue); +int *ptrY = malloc_shared(Size, Queue); + +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}); +Queue.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. +Queue.ext_oneapi_graph(execGraph); +Queue.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 From e9201c97d41ef263c674a9cb05fa315307e8a31e Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio=20Mestre?= Date: Mon, 24 Feb 2025 18:26:03 +0000 Subject: [PATCH 4/6] Fix typos in usage guide examples --- sycl/doc/syclgraph/SYCLGraphUsageGuide.md | 18 +++++++++--------- 1 file changed, 9 insertions(+), 9 deletions(-) diff --git a/sycl/doc/syclgraph/SYCLGraphUsageGuide.md b/sycl/doc/syclgraph/SYCLGraphUsageGuide.md index a510ca0450891..0f6c952ee2695 100644 --- a/sycl/doc/syclgraph/SYCLGraphUsageGuide.md +++ b/sycl/doc/syclgraph/SYCLGraphUsageGuide.md @@ -438,17 +438,17 @@ function and lambdas: using namespace sycl; namespace sycl_ext = sycl::ext::oneapi::experimental; -queue Queue; -sycl_ext::command_graph Graph{Queue.get_context(), Queue.get_device()}; +queue myQueue; +sycl_ext::command_graph myGraph{myQueue.get_context(), myQueue.get_device()}; -const size_t Size = 1024; -int *ptrX = malloc_shared(Size, Queue); -int *ptrY = malloc_shared(Size, Queue); +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) { + cgh.parallel_for(range<1>(size), [=](item<1> Id) { // Get the USM pointer to use in the kernel. auto Ptr = dynParam.get(); @@ -458,7 +458,7 @@ sycl_ext::node kernelNode = myGraph.add([&](handler& cgh) { }); auto execGraph = myGraph.finalize({sycl_ext::property::graph::updatable}); -Queue.ext_oneapi_graph(execGraph); +myQueue.ext_oneapi_graph(execGraph); // Change ptrX argument to ptrY. dynParam.update(ptrY); @@ -467,8 +467,8 @@ dynParam.update(ptrY); execGraph.update(kernelNode); // Execute the graph again. -Queue.ext_oneapi_graph(execGraph); -Queue.wait(); +myQueue.ext_oneapi_graph(execGraph); +myQueue.wait(); sycl::free(ptrX, myQueue); sycl::free(ptrY, myQueue); From 795949bdfce468f3584060e3b5b968de5c75bc72 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?F=C3=A1bio=20Mestre?= Date: Wed, 12 Mar 2025 15:44:08 +0000 Subject: [PATCH 5/6] Update restriction for the dynamic_parameter class Updates the dynamic parameter class to require underlying type to be device copyable. --- .../experimental/sycl_ext_oneapi_graph.asciidoc | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index acb9a99693d27..03cc6895c70c2 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -547,6 +547,9 @@ 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); @@ -565,11 +568,10 @@ Table {counter: tableNumber}. Template parameters of the `dynamic_parameter` cla | 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 can be 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]. +compile time using the `ValueT` parameter. This underlying type must be {device_copyable}[device +copyable]. | |=== From 7b7f955f7e4d6ebab4562fcf40bfd74dfa85a7f1 Mon Sep 17 00:00:00 2001 From: Fabio Mestre Date: Mon, 19 May 2025 14:45:21 +0100 Subject: [PATCH 6/6] Make get() methods const --- .../experimental/sycl_ext_oneapi_graph.asciidoc | 17 +++++++++-------- 1 file changed, 9 insertions(+), 8 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 03cc6895c70c2..aa5e02e7a66c5 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -618,7 +618,7 @@ Parameters: | [source,c++] ---- -ValueT& get(); +ValueT& get() const; ---- | Returns a reference to the underlying value that @@ -716,7 +716,7 @@ public: range accessRange, id accessOffset); - accessor get(); + accessor get() const; }; } ---- @@ -1033,7 +1033,7 @@ Throws an exception with the `errc::invalid` error code if the sum of | [source,c++] ---- -accessor get(); +accessor get() const; ---- | Returns a `sycl::accessor` to use in the current kernel. @@ -1070,7 +1070,7 @@ public: void update(range newAllocationSize); - local_accessor get(); + local_accessor get() const; }; } ---- @@ -1108,7 +1108,8 @@ Table {counter: tableNumber}. Member functions of the `dynamic_local_accessor` c template dynamic_local_accessor( command_graph graph, - range allocationSize); + range allocationSize, + const property_list &propList = {}); ---- |Constructs a dynamic local accessor object that can be registered with command graph. @@ -1141,7 +1142,7 @@ use in all the graph nodes where it is registered. | [source,c++] ---- -local_accessor get(); +local_accessor get() const; ---- | @@ -1185,7 +1186,7 @@ public: void update(size_t newNum); - work_group_memory get(); + work_group_memory get() const; }; } ---- @@ -1252,7 +1253,7 @@ Parameters: | [source,c++] ---- -work_group_memory get(); +work_group_memory get() const; ---- |