diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc index 4ed2abdf0e880..5c486e65687c2 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_graph.asciidoc @@ -246,7 +246,7 @@ to define an edge between existing nodes, or using a Edges can also be created when explicitly adding nodes to the graph through existing SYCL mechanisms for expressing dependencies. Data dependencies from -buffer accessors to existing nodes in the graph are captured as an edge. Using +accessors to existing nodes in the graph are captured as an edge. Using `handler::depends_on()` will also create a graph edge when passed an event returned from a queue submission captured by a queue recording to the same graph. |=== @@ -328,6 +328,11 @@ class assume_buffer_outlives_graph { public: assume_buffer_outlives_graph() = default; }; + +class updatable { + public: + updatable() = default; +}; } // namespace graph namespace node { @@ -355,6 +360,87 @@ struct graphs_support; } // namespace device } // namespace info +enum class dynamic_parameter_type { + usm, + accessor, + scalar, + raw +}; + +template +class dynamic_parameter; + +template <> +class dynamic_parameter { +public: + dynamic_parameter(command_graph graph, + std::string label = ""); + + void register(void* param, node node); + + void register(void* param); + + void register(int argIndex, node node); + + void update(void* newValue, + command_graph execGraph); + + const std::string& get_label() const; +}; + +template <> +class dynamic_parameter { +public: + dynamic_parameter(command_graph graph, + std::string label = ""); + + template + void register(T param, node node); + + template + void register(T param); + + void register(int argIndex, node node); + + template + void update(T newValue, + command_graph execGraph); + + const std::string& get_label() const; +}; + + +template <> +class dynamic_parameter { +public: + dynamic_parameter(command_graph graph, + size_t size, std::string label = ""); + + void register(int argIndex, node node); + + template + void update(T newValue, + command_graph execGraph); + + size_t get_size() const; + const std::string& get_label() const; +}; + +template <> +class dynamic_parameter { +public: + dynamic_parameter(command_graph graph, + size_t size, std::string label = ""); + + void register(int argIndex, node node); + + void update(void* newValue, + command_graph execGraph); + + size_t get_size() const; + const std::string& get_label() const; +}; + class node {}; // State of a graph @@ -399,7 +485,11 @@ public: template<> class command_graph { public: - command_graph() = delete; + command_graph() = delete; + + template + void update_nd_range(const std::vector& nodes, + nd_range executionRange); }; } // namespace ext::oneapi::experimental @@ -464,12 +554,425 @@ edges. The `node` class provides the {crs}[common reference semantics]. +==== Dynamic Parameters + +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. They are updated on an +executable graph which contains the nodes associated with the dynamic parameter +object. + +The type of the underlying object a dynamic parameter represents is set at +compile time using a template parameter. This underlying memory object 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:../proposed/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. It is valid for a node argument to be registered with more than one +dynamic parameter instance, the semantics of this are that the last update to +any of the dynamic parameters on an executable graph is the new node argument +on the next submissions of that executable graph. + +Registration is done either using the argument index (known when argument +originally set using `set_arg()`/`set_args()`), or by matching the argument +value. Matching for USM typed dynamic parameters is done by direct comparison +of the pointer values. Matching for accessors is done by verifying the +underlying managed data is the same. Raw and scalar typed dynamic parameters +can't be registered by matching however, as the chances of unintentional +collisions in the comparison are too high. + +When registration is done by matching the node argument value, it must +be the original argument value. Attempting to register the dynamic parameter +again to the same node argument by matching against an updated value will not +work. This is because registration operates on a modifiable graph, while update +operates on an executable graph. For example, + [source,c++] ---- -namespace sycl::ext::oneapi::experimental { - class node {}; -} +// Update a node with a single USM parameter 'PtrA', to 'PtrB'. +dynamic_parameter DynParam(Graph); +DynParam.register(PtrA, Node); +DynParam.update(PtrB, ExecGraph); + +// Exception thrown here as no match against original value. +DynParam.register(PtrB, Node); + +// No-op as dynamic parameter already registered with 'PtrA'. +DynParam.register(PtrA, Node); +---- + +Dynamic parameters can optionally have a user-specified string label associated +with them. This can provide meaningful names to help track what the parameters +are when updating them. These labels are not used by the SYCL runtime in any +way. + +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. +[cols="2a,a"] +|=== +|Member Function|Description + +| +[source,c++] +---- +dynamic_parameter(command_graph graph, + std::string label = ""); ---- +|Constructs a dynamic parameter object that can be registered with node +arguments. + +Preconditions: + +* This constructor is only available when the templated `dynamic_parameter_type` + is `dynamic_parameter_type::usm` or `dynamic_parameter_type::accessor`. + +Parameters: + +* `graph` - Graph containing the nodes which will have dynamic parameters. +* `label` - An optional string label which can be used to help identify the +parameter which has been registered as dynamic. + +| +[source,c++] +---- +dynamic_parameter(command_graph graph, + size_t size, std::string label = ""); +---- +|Constructs a dynamic parameter object representing defined number of bytes +that can be registered with node arguments. + +Preconditions: + +* This constructor is only available when the templated `dynamic_parameter_type` + is `dynamic_parameter_type::scalar` or `dynamic_parameter_type::raw`. + +Parameters: + +* `graph` - Graph containing the nodes which will have dynamic parameters. +* `size` - Size in bytes of the parameter. +* `label` - An optional string label which can be used to help identify the +parameter which has been registered as dynamic. + +| +[source,c++] +---- +const std::string& get_label() const; +---- +|Returns the label associated with the dynamic parameter. + +| +[source,c++] +---- +size_t get_size() const; +---- +|Returns the size associated with the scalar dynamic parameter. + +Preconditions: + +* This constructor is only available when the templated `dynamic_parameter_type` + is `dynamic_parameter_type::scalar` or `dynamic_parameter_type::raw`. + +|=== + +Table {counter: tableNumber}. Member functions of the `dynamic_parameter` class +for registration. +[cols="2a,a"] +|=== +|Member Function|Description + +| +[source,c++] +---- +void register(void* param, node node); +---- +|Associate the dynamic parameter with USM kernel parameter for `node`. All USM +kernel arguments matching `param` are bound to the dynamic parameter. If `param` +has already been registered with the dynamic parameter for `node`, this +operation is a no-op. + +Preconditions: + +* This method is only available when the templated `dynamic_parameter_type` + of the `dynamic_parameter` instance is `dynamic_parameter_type::usm`. + +Parameters: + +* `param` - USM pointer used as a node argument. +* `node` - Kernel command to look for an argument matching `param`. + +Exceptions: + +* Throws with error code `invalid` if `param` does not match any arguments in + `node`. +* Throws with error code `invalid` if the type of `node` is not a kernel + execution. +* Throws with error code `invalid` if `node` is not a member of the graph that + this dynamic parameter was created with. + +| +[source,c++] +---- +void register(void* param); +---- +|Associate the dynamic parameter with all the USM kernel parameters which match +`param` in the graph used on creation of the dynamic parameter. For every kernel +node in the graph, the dynamic parameter binds to any kernel arguments matching +`param`. It is not an error if the dynamic parameter fails to bind to a single +node, only if no nodes in the graph have an argument matching `param`. + +Preconditions: + +* This method is only available when the templated `dynamic_parameter_type` + of the `dynamic_parameter` instance is `dynamic_parameter_type::usm`. + +Parameters: + +* `param` - USM pointer used as a node argument. + +Exceptions: + +* Throws with error code `invalid` if no node arguments are matched to `param` + in the graph. + +| +[source,c++] +---- +template +void register(T param, node node); +---- +|Associate an accessor dynamic parameter with an accessor kernel parameter of +`node`. All accessor kernel arguments matching `param` are bound to the dynamic +parameter. If `param` has already been registered the dynamic parameter for +`node`, this operation is a no-op. + +Preconditions: + +* This method is only available when the templated `dynamic_parameter_type` + of the `dynamic_parameter` instance is `dynamic_parameter_type::accessor`. + +Parameters: + +* `param` - Accessor used as a node argument. +* `node` - Kernel command to look for an argument matching `param`. + +Exceptions: + +* Throws with error code `invalid` if `param` does not match any arguments in + `node`. +* Throws with error code `invalid` if `param` is not an accessor. +* Throws with error code `invalid` if the type of `node` is not a kernel + execution. +* Throws with error code `invalid` if `node` is not a member of the graph that + this dynamic parameter was created with. + +| +[source,c++] +---- +template +void register(T param); +---- +|Associate the dynamic parameter with all the accessor kernel parameters which +match `param` in the graph used on creation of the dynamic parameter. For +every kernel node in the graph, the dynamic parameter binds to any kernel +arguments matching `param`. It is not an error if the dynamic parameter fails +to bind to a single node, only if no nodes in the graph have an argument +matching `param`. + +Preconditions: + +* This method is only available when the templated `dynamic_parameter_type` + of the `dynamic_parameter` instance is `dynamic_parameter_type::accessor`. + +Parameters: + +* `param` - Accessor used as a node argument. + +Exceptions: + +* Throws with error code `invalid` if no node arguments are matched to `param` + in the graph. + +| +[source,c++] +---- +template +void register(int argIndex, node node); +---- +|Associate the dynamic parameter with a kernel parameter of `node` via the +argument index. Although argument ordering is not defined for lambda captures, +the ordering is reliable for arguments set using `handler::set_arg()` or +`handler::set_args()`. If `argIndex` is already registered with the dynamic +parameter for `node`, this operation is a no-op. + +Parameters: + +* `argIndex` - Kernel argument index starting from zero to associate with + the dynamic parameter. +* `node` - Kernel command to associate with argument at index `argIndex`. + +Exceptions: + +* Throws with error code `invalid` if `argIndex` is not a valid argument index + for the `node` kernel. +* Throws with error code `invalid` if the type of the dynamic parameter does not + correspond to the type of the `node` argument at `argIndex`. +* Throws with error code `invalid` if the type of `node` is not a kernel + execution. +* Throws with error code `invalid` if `node` is not a member of the graph that + this dynamic parameter was created with. + +|=== + +Table {counter: tableNumber}. Member functions of the `dynamic_parameter` class +for update. +[cols="2a,a"] +|=== +|Member Function|Description + +| +[source,c++] +---- +void update(void* newValue, + command_graph execGraph); +---- + +|Updates USM parameters in `execGraph` nodes associated with the dynamic +parameter to `newValue`. This new value will not affect any prior submissions +of `execGraph` and will take affect only for future submissions. It is not an +error if `newValue` is set to the current parameter value in nodes of +`execGraph`, i.e. no update occurred. + +Preconditions: + +* This method is only available when the templated `dynamic_parameter_type` + of the `dynamic_parameter` instance is `dynamic_parameter_type::usm`. + +Parameters: + +* `newValue` - Value to the parameters being updated to. +* `execGraph` - The executable graph to update the node parameters for. + +Exceptions: + +* Throws synchronously with error code `invalid` if + `property::graph::updatable` was not set when `execGraph` was created. +* Throws synchronously with error code `invalid` if `newValue` is not a pointer + to a valid USM allocation. +* Throws synchronously with error code `invalid` if the graph set on creation of + the dynamic parameter is not the same graph as used to create `execGraph`. +* Throws synchronously with error code `invalid` if the dynamic parameter is + not associated with any nodes in `execGraph`. + +| +[source,c++] +---- +template +void update(T newValue, + command_graph execGraph); +---- + +|Updates accessor parameters in `execGraph` nodes associated with the dynamic +parameter to `newValue`. This new value will not affect any prior submissions +of `execGraph` and will take affect only for future submissions. It is not an +error if `newValue` is set to the current parameter value in nodes of `execGraph`, +i.e. no update occurred. + +Preconditions: + +* This method is only available when the templated `dynamic_parameter_type` + of the `dynamic_parameter` instance is `dynamic_parameter_type::accessor`. + +Parameters: + +* `newValue` - Value to the parameters being updated to. +* `execGraph` - The executable graph to update the node parameters for. + +Exceptions: + +* Throws synchronously with error code `invalid` if + `property::graph::updatable` was not set when `execGraph` was created. +* Throws with error code `invalid` if `newValue` is not a valid accessor. +* Throws synchronously with error code `invalid` if the graph set on creation of + the dynamic parameter is not the same graph as used to create `execGraph`. +* Throws synchronously with error code `invalid` if the dynamic parameter is + not associated with any nodes in `execGraph`. + +| +[source,c++] +---- +template +void update(T newValue, + command_graph execGraph); +---- + +|Updates scalar parameters in `execGraph` nodes associated with the dynamic +parameter to `newValue`. This new value will not affect any prior submissions +of `execGraph` and will take affect only for future submissions. It is not an +error if `newValue` is set to the current parameter value of nodes in +`execGraph`, i.e. no update occurred. + +Preconditions: + +* This method is only available when the templated `dynamic_parameter_type` + of the `dynamic_parameter` instance is `dynamic_parameter_type::scalar`. + +Parameters: + +* `newValue` - Value to the parameters being updated to. +* `execGraph` - The executable graph to update the node parameters for. + +Exceptions: + +* Throws synchronously with error code `invalid` if + `property::graph::updatable` was not set when `execGraph` was created. +* Throws with error code `invalid` if `newValue` does not have the same size + as specified on creation of the dynamic parameter. +* Throws synchronously with error code `invalid` if the graph set on creation of + the dynamic parameter is not the same graph as used to create `execGraph`. +* Throws synchronously with error code `invalid` if the dynamic parameter is + not associated with any nodes in `execGraph`. + +| +[source,c++] +---- +void update(void* newValue, + command_graph execGraph); +---- + +|Updates raw byte parameters in `execGraph` nodes associated with the dynamic +parameter to a new value represented by array `newValue` of the size set when +the dynamic parameter was created. This new value will not affect any prior submissions +of `execGraph` and will take affect only for future submissions. It is not an +error if `newValue` is set to the current parameter value in nodes of +`execGraph`, i.e. no update occurred. + +Preconditions: + +* This method is only available when the templated `dynamic_parameter_type` + of the `dynamic_parameter` instance is `dynamic_parameter_type::raw`. + +Parameters: + +* `newValue` - Byte array containing the value to update the parameters to. +* `execGraph` - The executable graph to update the node parameters for. + +Exceptions: + +* Throws synchronously with error code `invalid` if `newValue` is a null pointer. +* Throws synchronously with error code `invalid` if + `property::graph::updatable` was not set when `execGraph` was created. +* Throws synchronously with error code `invalid` if the graph set on creation of + the dynamic parameter is not the same graph as used to create `execGraph`. +* Throws synchronously with error code `invalid` if the dynamic parameter is + not associated with any nodes in `execGraph`. + +|=== ==== Depends-On Property @@ -548,6 +1051,80 @@ graph LR Modifiable -->|Finalize| Executable .... +==== Executable Graph Update + +A graph in the executable state can have the configuration of its nodes modified +using a concept called graph _update_. This avoids a user having to rebuild and +finalize a new executable graph when only the inputs & outputs to a graph +change between submissions. Updates to a node take effect from the next +submission of a graph and will not affect any previous submissions or in-flight +executions of the same graph. + +The only type of node that is currently supported for updating in a graph is +kernel execution nodes. + +The aspects of a kernel execution node that can be configured during update are: + +* Parameters to the kernel. +* ND-Range of the kernel. + +Two methods are provided by the API to the user for performing this update. +An API for updating the whole graph object, which is most useful when the +graph was recorded, and an individual node update API. + +To update an executable graph, the `property::graph::updatable` property must +have been set when the graph was created during finalization. Otherwise, an +exception will be thrown if a user tries to update an executable graph. This +guarantee allows the backend to provide a more optimized implementation, if +possible. + +The `property::graph::updatable` property also allows an executable graph to be +submitted for execution while a previous submission of the same executable +graph instance is still executing. This is because the ability to change the +graph inputs/outputs can remove the data race conditions that could otherwise +exist if the same executable graph was executed concurrently. + +===== Whole Graph Update + +===== Individual Node Update + +Memory parameters to individual nodes in a graph in the `executable` state +can be updated between graph executions using dynamic parameters. When a +`dynamic_parameter` instance is created with a modifiable state graph it is +empty, with no associated kernel arguments. By registering the +`dynamic_parameter` with a node using `dynamic_parameter::register()`, the +dynamic parameter is matched to a node argument, and the `dynamic_parameter` +can then be used to update that node argument in future. + +If a user doesn't have the individual `node` handles to register parameters +with, and the dynamic parameter type isn't raw or scalar, they can use the +`dynamic_parameter::register()` API that doesn't take a `node` argument. +This entry-point iterates through the modifiable graph nodes, and tries to +register the parameter against every kernel node in the graph. + +The update itself is then performed using a `dynamic_parameter` instance by +calling `dynamic_parameter::update()` to update all the parameters of nodes in +an executable graph to which the `dynamic_parameter` is registered. + +The other node configuration that can be updated is the ND-Range, this can +be set through `command_graph::update_nd_range()` but +does not require any prior registration. + +Since the structure of the graph became fixed when finalizing, updating +parameters on a node in a graph in the `executable` state will not change +the already defined dependencies between nodes. This is important to note +when updating buffer parameters to a node, since no edges will be automatically +created or removed based on this change. Care should be taken that updates +of buffer parameters do not change the 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. + ==== Graph Properties [[graph-properties]] ===== No-Cycle-Check Property @@ -775,8 +1352,9 @@ Preconditions: Parameters: -* `propList` - Optional parameter for passing properties. No finalization - properties are defined by this extension. +* `propList` - Optional parameter for passing properties. The only property + that is valid to pass here is `property::graph::updatable`, to enable the + returned executable graph to be <>. Returns: A new executable graph object which can be submitted to a queue. @@ -806,9 +1384,66 @@ Exceptions: * Throws synchronously with error code `invalid` if the path is invalid or the file extension is not supported or if the write operation failed. +| +[source,c++] +---- +std::vector get_nodes() const; +---- +|Returns a list of all the nodes present in the graph in the order that they +were added. + +| +[source,c++] +---- +std::vector get_root_nodes() const; +---- +|Returns a list of all nodes in the graph which have no dependencies in the +order they were added to the graph. + |=== -Table {counter: tableNumber}. Member functions of the `command_graph` class for queue recording. +Table {counter: tableNumber}. Member functions of the `command_graph` class for +graph update. +[cols="2a,a"] +|=== +|Member function|Description + +| +[source,c++] +---- +template +void update_nd_range(const std::vector& nodes, + nd_range executionRange); +---- + +| Updates the ND-Range for the nodes with a new value. This new value will +not affect any prior submissions of this graph and will take affect only +for future submissions. See <> +for more information about updating node parameters. + +Preconditions: + +* This member function is only available when the `command_graph` state is + `graph_state::executable`. + +Parameters: + +* `nodes` - The nodes in this graph which will have their ND-Range values updated. + +* `executionRange` - The new value for the ND-Range. + +Exceptions: + +* Throws synchronously with error code `invalid` if + `property::graph::updatable` was not set when the executable graph was created. +* Throws with error code `invalid` if any node in `nodes` is not part of the + graph. +* Throws with error code `nd_range` if `executionRange` is an invalid ND-Range + for any node in `nodes`. +|=== + +Table {counter: tableNumber}. Member functions of the `command_graph` class for +queue recording. [cols="2a,a"] |=== |Member function|Description @@ -844,6 +1479,7 @@ Exceptions: * Throws synchronously with error code `invalid` if `recordingQueue` is associated with a device or context that is different from the device and context used on creation of the graph. + | [source, c++] ---- @@ -1091,9 +1727,9 @@ handler::ext_oneapi_graph(command_graph& graph) ---- |Invokes the execution of a graph. Only one instance of `graph` will -execute at any time. If `graph` is submitted multiple times, dependencies -are automatically added by the runtime to prevent concurrent executions of -an identical graph. +execute at any time. If `graph` is submitted multiple times and +`property::graph::updatable_graph` is not set, dependencies are automatically +added by the runtime to prevent concurrent executions of an identical graph. Parameters: @@ -1548,6 +2184,97 @@ submitted in its entirety for execution via ... ---- +=== Dynamic Parameter Update + +Example showing a graph with two kernel nodes, the first created using a kernel +bundle with `handler::set_args()`, having its node arguments updated. + +[source,c++] +---- +... + +using namespace sycl; +namespace sycl_ext = sycl::ext::oneapi::experimental; + +queue myQueue; +auto myContext = myQueue.get_context(); +auto myDevice = myQueue.get_device(); + +// USM allocations for kernel input/output +const size_t n = 1024; +int *ptrX = malloc_shared(n, myQueue); +int *ptrY = malloc_device(n, myQueue); + +int *ptrZ = malloc_shared(n, myQueue); +int *ptrQ = malloc_device(n, myQueue); + +// Kernel loaded from kernel bundle +const std::vector builtinKernelIds = + myDevice.get_info(); +kernel_bundle myBundle = + get_kernel_bundle(myContext, { myDevice }, builtinKernelIds); +kernel builtinKernel = myBundle.get_kernel(builtinKernelIds[0]); + +// Graph containing a two kernels node +sycl_ext::command_graph myGraph(myContext, myDevice); + +// Create graph dynamic parameters +dynamic_parameter dynParamInput(myGraph); +dynamic_parameter dynParamOutput(myGraph); +dynamic_parameter dynParamScalar(myGraph); + +// First node uses ptrX as an input & output parameter, with operand +// mySclar as another argument. +int myScalar = 42; +node nodeA = myGraph.add([&](handler& cgh) { + cgh.set_args(ptrX, myScalar); + cgh.parallel_for(range {n}, builtinKernel); +}); + +// Register nodeA dynamic parameters +dynParamInput.register(0, nodeA); // Argument index 0 is ptrX +dynParamInput.register(1, nodeA); // Argument index 1 is myScalar + +// Second node uses ptrX as an inputs, and ptrY as input/output/ +node nodeB = myGraph.add([&](handler& cgh) { + cgh.parallel_for(sycl::range<1>{n}, [=](sycl::id<1> it) { + const size_t i = it[0]; + ptrY[i] += ptrX[i]; + }); + }, + {sycl_ext::property::node::depends_on(nodeA)}); + +// Register nodeB dynamic parameters +dynParamInput.register(ptrX, nodeB); +dynParamInput.register(ptrY, nodeB); + +// Create an executable graph with the updatable property. +auto execGraph = myGraph.finalize({sycl_ext::property::graph::updatable}); + +// Execute graph, then update without needing to wait for it to complete +myQueue.ext_oneapi_graph(execGraph); + +// Change ptrX argument to node A & B to ptrZ +dynParamInput.update(ptrZ, execGraph); + +// Change myScalar argument to node A to newScalar +int newScalar = 12; +dynParamScalar.update(newScalar, execGraph); + +// Change ptrY argument to node B to ptrQ +dynParamOutput.update(ptrQ, execGraph); + +// Execute graph again, without needing to wait on completion of first submission. +myQueue.ext_oneapi_graph(exec); +myQueue.wait(); + +sycl::free(ptrX, myQueue); +sycl::free(ptrY, myQueue); +sycl::free(ptrZ, myQueue); +sycl::free(ptrQ, myQueue); + +---- + == Future Direction [[future-direction]] This section contains both features of the specification which have been @@ -1752,6 +2479,14 @@ runtime. == Issues +=== Update More Command Types + +Support updating arguments to types of nodes other that kernel execution +commands. + +**RESOLVED** Should be added be for at least memory copy nodes, however +fully scope of support needs designed and implemented. + === Multi Device Graph Allow an executable graph to contain nodes targeting different devices. @@ -1816,8 +2551,8 @@ if used in application code. `sycl::ext::intel::property::queue::no_immediate_command_list` should be set on construction to any queues an executable graph is submitted to. -. Synchronization between multiple executions of the same command-buffer - must be handled in the host for level-zero backend, which may involve +. Synchronization between multiple executions of the same command-buffer + must be handled in the host for level-zero backend, which may involve extra latency for subsequent submissions. == Revision History