Skip to content

Commit

Permalink
Use array for clUpdateMutableCommandsKHR. (#1045)
Browse files Browse the repository at this point in the history 
Proposal to pass the update configs to `clUpdateMutableCommandsKHR` as
an array, rather than pointer changed linked list.

See #1041 for
motivation.
  • Loading branch information
@EwanC
EwanC authored Jul 16, 2024
1 parent fed48e7 commit c6cceb1
Show file tree
Hide file tree
Showing 3 changed files with 80 additions and 107 deletions.
60 changes: 23 additions & 37 deletions api/cl_khr_command_buffer_mutable_dispatch.asciidoc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -6,7 +6,7 @@ include::{generated}/meta/{refprefix}cl_khr_command_buffer_mutable_dispatch.txt[
=== Other Extension Metadata

*Last Modified Date*::
2022-08-31
2024-06-19
*IP Status*::
No known IP claims.
*Contributors*::
Expand Down Expand Up @@ -43,32 +43,15 @@ in a new command-buffer.

=== Interactions With Other Extensions

The {cl_command_buffer_structure_type_khr_TYPE} type has been added to this
extension for the purpose of allowing expansion of mutable functionality in
future extensions layered on top of
{cl_khr_command_buffer_mutable_dispatch_EXT}.
Any parameter that is a structure containing a `void* next` member *must*
have a value of `next` that is either `NULL`, or is a pointer to a valid
structure defined by {cl_khr_command_buffer_mutable_dispatch_EXT} or an
extension layered on top.
To be a valid structure in the pointer chain the first member of the
structure *must* be a {cl_command_buffer_structure_type_khr_TYPE} identifier
for the structure being iterated through, and the second member a `void*
next` pointer to the next structure in the chain.

[NOTE]
====
This approach is based on structure pointer chains in Vulkan, for more
details see the "`Valid Usage for Structure Pointer Chains`" section of the
Vulkan specification.
====

This is designed so that another extension layered on
{cl_khr_command_buffer_mutable_dispatch_EXT} could allow modification of
commands recorded to a command-buffer other than kernel execution commands.
As all command recording entry-points return a {cl_mutable_command_khr_TYPE}
handle, and aspects like which {cl_mem_TYPE} object a command uses could
also be updated between enqueues of the command-buffer.
The {clUpdateMutableCommandsKHR} entry-point has been designed for the purpose
of allowing expansion of mutable functionality in future extensions layered on
top of `cl_khr_command_buffer_mutable_dispatch`.

A new extension can define its own structure type to specify the update
configuration it requires, with a matching
{cl_command_buffer_update_type_khr_TYPE} value. This new structure type can
then be passed to {clUpdateMutableCommandsKHR} where it is reinterpreted from a
void pointer using {cl_command_buffer_update_type_khr_TYPE}.

=== New Commands

Expand All @@ -79,8 +62,7 @@ also be updated between enqueues of the command-buffer.

* {cl_mutable_dispatch_fields_khr_TYPE}
* {cl_mutable_command_info_khr_TYPE}
* {cl_command_buffer_structure_type_khr_TYPE}
* {cl_mutable_base_config_khr_TYPE}
* {cl_command_buffer_update_type_khr_TYPE}
* {cl_mutable_dispatch_asserts_khr_TYPE}
* {cl_mutable_dispatch_config_khr_TYPE}
* {cl_mutable_dispatch_exec_info_khr_TYPE}
Expand Down Expand Up @@ -115,8 +97,7 @@ also be updated between enqueues of the command-buffer.
** {CL_COMMAND_BUFFER_MUTABLE_KHR}
* {cl_command_buffer_properties_khr_TYPE}
** {CL_COMMAND_BUFFER_MUTABLE_DISPATCH_ASSERTS_KHR}
* {cl_command_buffer_structure_type_khr_TYPE}
** {CL_STRUCTURE_TYPE_MUTABLE_BASE_CONFIG_KHR}
* {cl_command_buffer_update_type_khr_TYPE}
** {CL_STRUCTURE_TYPE_MUTABLE_DISPATCH_CONFIG_KHR}
* New Error Codes
** {CL_INVALID_MUTABLE_COMMAND_KHR}
Expand Down Expand Up @@ -274,8 +255,6 @@ kernel void vector_addition(global int* tile1, global int* tile2,
cl_mutable_dispatch_arg_khr arg_2{2, sizeof(cl_mem), &output_buffer};
cl_mutable_dispatch_arg_khr args[] = {arg_0, arg_1, arg_2};
cl_mutable_dispatch_config_khr dispatch_config{
CL_STRUCTURE_TYPE_MUTABLE_DISPATCH_CONFIG_KHR,
nullptr,
command_handle,
3 /* num_args */,
0 /* num_svm_arg */,
Expand All @@ -287,12 +266,16 @@ kernel void vector_addition(global int* tile1, global int* tile2,
nullptr /* global_work_offset */,
nullptr /* global_work_size */,
nullptr /* local_work_size */};
cl_mutable_base_config_khr mutable_config{
CL_STRUCTURE_TYPE_MUTABLE_BASE_CONFIG_KHR, nullptr, 1,
&dispatch_config};
// Update the command buffer with the mutable configuration
error = clUpdateMutableCommandsKHR(command_buffer, &mutable_config);
cl_uint num_configs = 1;
cl_command_buffer_update_type_khr config_types[1] = {
CL_STRUCTURE_TYPE_MUTABLE_DISPATCH_CONFIG_KHR
};
const void* configs[1] = {&dispatch_config};
error = clUpdateMutableCommandsKHR(command_buffer, num_configs,
config_types, configs);
CL_CHECK(error);
}
Expand Down Expand Up @@ -374,3 +357,6 @@ may be a introduced as a stand alone extension.
* Revision 0.9.1, 2023-11-07
** Add type {cl_mutable_dispatch_asserts_khr_TYPE} and its possible values
(provisional).
* Revision 0.9.2, 2024-06-19
** Change {clUpdateMutableCommandsKHR} API to pass configs as an array rather
than linked list (provisional).
93 changes: 45 additions & 48 deletions api/opencl_runtime_layer.asciidoc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -15647,7 +15647,7 @@ endif::cl_khr_command_buffer_multi_device[]

ifdef::cl_khr_command_buffer_mutable_dispatch[]
[[mutable-commands]]
=== Mutable Commands:
=== Mutable Commands

A generic {cl_mutable_command_khr_TYPE} handle is called a _mutable-command_
object as it can be returned from any command recording entry-point in the
Expand All @@ -15658,11 +15658,10 @@ modified through the fields of {cl_mutable_dispatch_config_khr_TYPE}.

Mutable-command handles are updated between enqueues using entry-point
{clUpdateMutableCommandsKHR}.
To enable performant usage, all aspects of mutation are encapsulated inside
a single {cl_mutable_base_config_khr_TYPE} parameter.
This means that the runtime has access to all the information about how the
command-buffer will change, allowing the command-buffer to be rebuilt as
efficiently as possible.
To enable performant usage, all aspects of mutation can be passed in a single
call using an array. This means that the runtime has access to all the
information about how the command-buffer will change, allowing the
command-buffer to be rebuilt as efficiently as possible.
Any modifications to the arguments or execution info of a mutable-dispatch
handle using {cl_mutable_dispatch_arg_khr_TYPE} or
{cl_mutable_dispatch_exec_info_khr_TYPE} have no affect on the original
Expand Down Expand Up @@ -15703,8 +15702,13 @@ include::{generated}/api/protos/clUpdateMutableCommandsKHR.txt[]
include::{generated}/api/version-notes/clUpdateMutableCommandsKHR.asciidoc[]

* _command_buffer_ refers to a valid command-buffer object.
* _mutable_config_ is a pointer to a {cl_mutable_base_config_khr_TYPE}
structure defining updates to make to mutable-commands.
* _num_configs_ Number of elements in the _config_types_ and _config_ arrays.
* _config_types_ An array of length _num_configs_ with each element identifying
the type of each config in _configs_ at the same array index.
* _configs_ An array of length _num_configs_ containing structs which define how a
mutable-command handle in _command_buffer_ is to be updated, each of which is
interpreted using _config_types_ at the same index with the mapping defined
in the <<update-config-mapping, Mutable Command Update Structs>> section.

// refError

Expand All @@ -15718,16 +15722,13 @@ one of the errors below is returned:
* {CL_INVALID_OPERATION} if _command_buffer_ has not been finalized.
* {CL_INVALID_OPERATION} if _command_buffer_ was not created with the
{CL_COMMAND_BUFFER_MUTABLE_KHR} flag.
* {CL_INVALID_VALUE} if the _type_ member of _mutable_config_ is not
{CL_STRUCTURE_TYPE_MUTABLE_BASE_CONFIG_KHR}.
* {CL_INVALID_VALUE} if the _mutable_dispatch_list_ member of
_mutable_config_ is `NULL` and _num_mutable_dispatch_ > 0, or
_mutable_dispatch_list_ is not `NULL` and _num_mutable_dispatch_ is 0.
* {CL_INVALID_VALUE} if the _next_ member of _mutable_config_ is not
`NULL` and any iteration of the structure pointer chain does not contain
valid _type_ and _next_ members.
* {CL_INVALID_VALUE} if _mutable_config_ is `NULL`, or if both _next_ and
_mutable_dispatch_list_ members of _mutable_config_ are `NULL`.
* {CL_INVALID_VALUE} if _config_types_ is `NULL` and _num_configs_ > 0, or
_config_types_ is not `NULL` and _num_configs_ is 0.
* {CL_INVALID_VALUE} if _configs_ is `NULL` and _num_configs_ > 0, or
_configs_ is not `NULL` and _num_configs_ is 0.
* {CL_INVALID_VALUE} if any element of _config_types_ is not a valid
{cl_command_buffer_update_type_khr_TYPE} enum.
* {CL_INVALID_VALUE} if any element of _configs_ is NULL.
* {CL_OUT_OF_RESOURCES} if there is a failure to allocate resources
required by the OpenCL implementation on the device.
* {CL_OUT_OF_HOST_MEMORY} if there is a failure to allocate resources
Expand All @@ -15753,19 +15754,17 @@ parameters are updated so that the new number of work-groups exceeds the
number when the ND-range command was recorded, the behavior is undefined.
====

If the _mutable_dispatch_list_ member of _mutable_config_ is non-`NULL`,
then errors defined by {clEnqueueNDRangeKernel}, {clSetKernelExecInfo},
{clSetKernelArg}, and {clSetKernelArgSVMPointer} are returned by
{clUpdateMutableCommandsKHR} if any of the array elements are set to an
invalid value.
Additionally, the following errors are returned if any
{cl_mutable_dispatch_config_khr_TYPE} element of the array violates the
defined conditions:
If _configs_ is non-`NULL`, then for any {cl_mutable_dispatch_config_khr_TYPE}
element of the array the errors defined by {clEnqueueNDRangeKernel},
{clSetKernelExecInfo}, {clSetKernelArg}, and {clSetKernelArgSVMPointer} are
returned by {clUpdateMutableCommandsKHR} if any of the struct elements are set
to an invalid value. Additionally, the following errors are returned if any
{cl_mutable_dispatch_config_khr_TYPE} element of the array violates the defined
conditions:

* {CL_INVALID_MUTABLE_COMMAND_KHR} if _command_ is not a valid mutable
command object, or created from _command_buffer_.
* {CL_INVALID_VALUE} if _type_ is not
{CL_STRUCTURE_TYPE_MUTABLE_DISPATCH_CONFIG_KHR}.
command object returned from {clCommandNDRangeKernelKHR}, or created from
_command_buffer_.
* {CL_INVALID_OPERATION} if the values of _local_work_size_ and/or
_global_work_size_ result in a change to work-group uniformity.
* {CL_INVALID_OPERATION} if the _work_dim_ is different from the
Expand Down Expand Up @@ -15793,24 +15792,25 @@ defined conditions:
0, or _exec_info_list_ is not `NULL` and _num_exec_infos_ is 0.
--

[open,refpage='cl_mutable_base_config_khr',desc='DESC',type='structs']
--
The {cl_mutable_base_config_khr_TYPE} structure encapsulates all aspects of
mutation and is defined as:
[[mutable-commands-update-structs]]
==== Mutable Command Update Structs

include::{generated}/api/structs/cl_mutable_base_config_khr.txt[]
The following table defines the mapping of
{cl_command_buffer_update_type_khr_TYPE} values to the structs they define
reinterpreting a void pointer as when passed to {clUpdateMutableCommandsKHR}.

* _type_ is the type of this structure, and must be
{CL_STRUCTURE_TYPE_MUTABLE_BASE_CONFIG_KHR_anchor}
* _next_ is `NULL` or a pointer to an extending structure.
* _num_mutable_dispatch_ is the number of mutable-dispatch objects to
configure in this enqueue of the command-buffer.
* _mutable_dispatch_list_ is an array containing _num_mutable_dispatch_
elements describing the configurations of mutable kernel execution
commands in the command-buffer.
For a description of struct members making up each array element see
{cl_mutable_dispatch_config_khr_TYPE}.
--
[[update-config-mapping]]
[cols=",,",options="header",]
|====
| Enum Value | Struct Type | Entry Point

| {CL_STRUCTURE_TYPE_MUTABLE_DISPATCH_CONFIG_KHR_anchor}
| {cl_mutable_dispatch_config_khr_TYPE}
| {clCommandNDRangeKernelKHR}

|====

==== Kernel Command Update Structs

[open,refpage='cl_mutable_dispatch_config_khr',desc='Set kernel configuration of a mutable clCommandNDRangeKernelKHR command',type='structs']
--
Expand All @@ -15820,9 +15820,6 @@ The {cl_mutable_dispatch_arg_khr_TYPE} structure is passed to

include::{generated}/api/structs/cl_mutable_dispatch_config_khr.txt[]

* _type_ is the type of this structure, and must be
{CL_STRUCTURE_TYPE_MUTABLE_DISPATCH_CONFIG_KHR_anchor}.
* _next_ is `NULL` or a pointer to an extending structure.
* _command_ is a mutable-command object returned by
{clCommandNDRangeKernelKHR} representing a kernel execution as part of a
command-buffer.
Expand Down
34 changes: 12 additions & 22 deletions xml/cl.xml
View file
Original file line number Diff line number Diff line change
Expand Up @@ -249,7 +249,7 @@ server's OpenCL/api-docs repository.
<type category="define">typedef struct _cl_mutable_command_khr* <name>cl_mutable_command_khr</name>;</type>
<type category="define">typedef <type>cl_bitfield</type> <name>cl_mutable_dispatch_fields_khr</name>;</type>
<type category="define">typedef <type>cl_uint</type> <name>cl_mutable_command_info_khr</name>;</type>
<type category="define">typedef <type>cl_uint</type> <name>cl_command_buffer_structure_type_khr</name>;</type>
<type category="define">typedef <type>cl_uint</type> <name>cl_command_buffer_update_type_khr</name>;</type>
<type category="define">typedef <type>cl_bitfield</type> <name>cl_device_fp_atomic_capabilities_ext</name>;</type>
<type category="define">typedef <type>cl_uint</type> <name>cl_image_requirements_info_ext</name>;</type>
<type category="define">typedef <type>cl_bitfield</type> <name>cl_platform_command_buffer_capabilities_khr</name>;</type>
Expand Down Expand Up @@ -370,8 +370,6 @@ server's OpenCL/api-docs repository.
<member>const <type>void</type>* <name>param_value</name></member>
</type>
<type category="struct" name="cl_mutable_dispatch_config_khr">
<member><type>cl_command_buffer_structure_type_khr</type> <name>type</name></member>
<member>const <type>void</type>* <name>next</name></member>
<member><type>cl_mutable_command_khr</type> <name>command</name></member>
<member><type>cl_uint</type> <name>num_args</name></member>
<member><type>cl_uint</type> <name>num_svm_args</name></member>
Expand All @@ -384,13 +382,6 @@ server's OpenCL/api-docs repository.
<member>const <type>size_t</type>* <name>global_work_size</name></member>
<member>const <type>size_t</type>* <name>local_work_size</name></member>
</type>

<type category="struct" name="cl_mutable_base_config_khr">
<member><type>cl_command_buffer_structure_type_khr</type> <name>type</name></member>
<member>const <type>void</type>* <name>next</name></member>
<member><type>cl_uint</type> <name>num_mutable_dispatch</name></member>
<member>const <type>cl_mutable_dispatch_config_khr</type>* <name>mutable_dispatch_list</name></member>
</type>
</types>

<!-- SECTION: OpenCL enumerant (token) definitions. -->
Expand Down Expand Up @@ -1370,10 +1361,9 @@ server's OpenCL/api-docs repository.
<enum bitpos="0" name="CL_MUTABLE_DISPATCH_ASSERT_NO_ADDITIONAL_WORK_GROUPS_KHR"/>
</enums>

<enums name="cl_command_buffer_structure_type_khr" vendor="Khronos">
<enum value="0" name="CL_STRUCTURE_TYPE_MUTABLE_BASE_CONFIG_KHR"/>
<enum value="1" name="CL_STRUCTURE_TYPE_MUTABLE_DISPATCH_CONFIG_KHR"/>
<unused start="2" end="2" comment="Used by future command-buffer extensions"/>
<enums name="cl_command_buffer_update_type_khr" vendor="Khronos">
<enum value="0" name="CL_STRUCTURE_TYPE_MUTABLE_DISPATCH_CONFIG_KHR"/>
<unused start="1" end="1" comment="Used by future command-buffer extensions"/>
</enums>

<enums name="cl_device_fp_atomic_capabilities_ext" vendor="EXT" type="bitmask">
Expand Down Expand Up @@ -3280,9 +3270,11 @@ server's OpenCL/api-docs repository.
<param><type>size_t</type>* <name>param_value_size_ret</name></param>
</command>
<command>
<proto><type>cl_int</type> <name>clUpdateMutableCommandsKHR</name></proto>
<param><type>cl_command_buffer_khr</type> <name>command_buffer</name></param>
<param>const <type>cl_mutable_base_config_khr</type>* <name>mutable_config</name></param>
<proto><type>cl_int</type> <name>clUpdateMutableCommandsKHR</name></proto>
<param><type>cl_command_buffer_khr</type> <name>command_buffer</name></param>
<param><type>cl_uint</type> <name>num_configs</name></param>
<param>const <type>cl_command_buffer_update_type_khr</type>* <name>config_types</name></param>
<param>const <type>void</type>** <name>configs</name></param>
</command>
<command>
<proto><type>cl_int</type> <name>clGetMutableCommandInfoKHR</name></proto>
Expand Down Expand Up @@ -7324,18 +7316,17 @@ server's OpenCL/api-docs repository.
<enum name="CL_QUEUE_JOB_SLOT_ARM"/>
</require>
</extension>
<extension name="cl_khr_command_buffer_mutable_dispatch" revision="0.9.1" supported="opencl" depends="cl_khr_command_buffer" ratified="opencl" provisional="true" comment="requires cl_khr_command_buffer 0.9.0 or later">
<extension name="cl_khr_command_buffer_mutable_dispatch" revision="0.9.2" supported="opencl" depends="cl_khr_command_buffer" ratified="opencl" provisional="true" comment="requires cl_khr_command_buffer 0.9.0 or later">
<require>
<type name="CL/cl.h"/>
</require>
<require>
<type name="cl_command_buffer_structure_type_khr"/>
<type name="cl_command_buffer_update_type_khr"/>
<type name="cl_mutable_dispatch_fields_khr"/>
<type name="cl_mutable_command_info_khr"/>
<type name="cl_mutable_dispatch_arg_khr"/>
<type name="cl_mutable_dispatch_exec_info_khr"/>
<type name="cl_mutable_dispatch_config_khr"/>
<type name="cl_mutable_base_config_khr"/>
<type name="cl_mutable_dispatch_asserts_khr"/>
</require>
<require comment="cl_command_buffer_flags_khr - bitfield">
Expand Down Expand Up @@ -7369,8 +7360,7 @@ server's OpenCL/api-docs repository.
<enum name="CL_MUTABLE_DISPATCH_GLOBAL_WORK_SIZE_KHR"/>
<enum name="CL_MUTABLE_DISPATCH_LOCAL_WORK_SIZE_KHR"/>
</require>
<require comment="cl_command_buffer_structure_type_khr">
<enum name="CL_STRUCTURE_TYPE_MUTABLE_BASE_CONFIG_KHR"/>
<require comment="cl_command_buffer_update_type_khr">
<enum name="CL_STRUCTURE_TYPE_MUTABLE_DISPATCH_CONFIG_KHR"/>
</require>
<require comment="cl_command_buffer_properties_khr">
Expand Down

0 comments on commit c6cceb1

Please sign in to comment.