Skip to content

Commit

Permalink
Use array for clUpdateMutableCommandsKHR.
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 committed Jan 19, 2024
1 parent b34543b commit 8b0b63a
Show file tree
Hide file tree
Showing 2 changed files with 63 additions and 115 deletions.
144 changes: 51 additions & 93 deletions ext/cl_khr_command_buffer_mutable_dispatch.asciidoc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -20,6 +20,7 @@ commands between command-buffer enqueues.
|====
| *Date* | *Version* | *Description*
| 2022-08-31 | 0.9.0 | First assigned version (provisional).
| 2023-01-15 | 0.9.1 | Change {clUpdateMutableCommandsKHR} API to pass configs as an array rather than linked list (provisional).
|====

include::provisional_notice.asciidoc[]
Expand Down Expand Up @@ -61,30 +62,15 @@ without having to re-record the entire command sequence 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`.
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` 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.
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`.

[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` 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.
A new extension can define its own structure type to specify the update
configuration it requires, with a matching {cl_update_config_type_khr_TYPE}
value. This new structure type can then be passed to {clUpdateMutableCommandsKHR}
where it is reinterpreted from a void pointer using
{cl_update_config_type_khr_TYPE}.

=== New Types

Expand All @@ -101,8 +87,9 @@ typedef cl_bitfield cl_mutable_dispatch_fields_khr;
// For querying mutable-command objects with clGetMutableCommandInfoKHR
typedef cl_uint cl_mutable_command_info_khr;
// Identifies the type of a structure to allow structure pointer chains
typedef cl_uint cl_command_buffer_structure_type_khr;
// Identifies the structure type to reinterpret a void pointer element in
// clUpdateMutableCommandsKHR config array parameter as.
typedef cl_uint cl_update_config_type_khr;
----

Struct type for setting kernel arguments normally passed using {clSetKernelArg}
Expand All @@ -127,11 +114,6 @@ configuration of a mutable {clCommandNDRangeKernelKHR} command:

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

_type_ Type of this structure, must be
{CL_STRUCTURE_TYPE_MUTABLE_DISPATCH_CONFIG_KHR}.

_next_ Is `NULL` or a pointer to an extending structure.

_command_ A mutable-command object returned by {clCommandNDRangeKernelKHR}
representing a kernel execution as part of a command-buffer.

Expand Down Expand Up @@ -179,38 +161,16 @@ describe the number of work-items that make up a work-group that will execute
the kernel. If _local_work_size_ is `NULL` then the number of local work-items
in the dispatch is not changed. See {clEnqueueNDRangeKernel} for valid usage.

[[cl_mutable_base_config_khr]]
[source,opencl]
----
typedef struct _cl_mutable_base_config_khr {
cl_command_buffer_structure_type_khr type,
const void* next,
cl_uint num_mutable_dispatch,
const cl_mutable_dispatch_config_khr* mutable_dispatch_list
} cl_mutable_base_config_khr;
----

_type_ Type of this structure, must be
{CL_STRUCTURE_TYPE_MUTABLE_BASE_CONFIG_KHR}

_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}.

=== New API Functions

Mutable-handle entry points from <<mutable-commands, Section 5.X.5>>:
[source,opencl]
----
cl_int clUpdateMutableCommandsKHR(
cl_command_buffer_khr command_buffer,
const cl_mutable_base_config_khr* mutable_config);
cl_uint num_configs,
const cl_update_config_type_khr* config_types,
const void** configs);
cl_int clGetMutableCommandInfoKHR(
cl_mutable_command_khr command,
Expand Down Expand Up @@ -258,14 +218,12 @@ CL_MUTABLE_COMMAND_COMMAND_TYPE_KHR 0x12AD
CL_COMMAND_BUFFER_MUTABLE_KHR (0x1 << 1)
----

Enum values for {cl_command_buffer_structure_type_khr_TYPE} allowing the structure
types used for mutating commands between enqueues to be extended by future
extensions built on top of `cl_khr_command_buffer_mutable_dispatch`. Based on
structure pointer chains in Vulkan.
Enum values for {cl_update_config_type_khr_TYPE} allowing the elements of the void
pointer array used for mutating commands between enqueues to be reinterpreted
as the correct struct type.
[source,opencl]
----
CL_STRUCTURE_TYPE_MUTABLE_BASE_CONFIG_KHR 0
CL_STRUCTURE_TYPE_MUTABLE_DISPATCH_CONFIG_KHR 1
CL_STRUCTURE_TYPE_MUTABLE_DISPATCH_CONFIG_KHR 0 // Type of cl_mutable_dispatch_config_khr
----

=== Modifications to section 4.2 of the OpenCL API Specification
Expand Down Expand Up @@ -432,8 +390,7 @@ _mutable-dispatch_ objects, and can be modified through the fields of

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
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
Expand Down Expand Up @@ -485,9 +442,14 @@ individually.

_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.

{clUpdateMutableCommandsKHR} returns {CL_SUCCESS} if all the mutable-command
objects were updated successfully. Otherwise, none of the updates to
Expand All @@ -501,39 +463,34 @@ mutable-command objects are preserved and one of the errors below is returned:
* {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 _config_types_ is `NULL` and _num_configs_ > 0, or
_config_types_ is not `NULL` and _num_configs_ is 0.

* {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 _configs_ is `NULL` and _num_configs_ > 0, or
_configs_ is not `NULL` and _num_configs_ 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 any element of _config_types_ is not a valid
{cl_update_config_type_khr_TYPE} enum.

* {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 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 required by
the OpenCL implementation on the host.

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 values of _local_work_size_ and/or
_global_work_size_ result in an increase to the number of work-groups in the
Expand Down Expand Up @@ -859,8 +816,6 @@ command-buffer submissions.
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 @@ -872,12 +827,15 @@ command-buffer submissions.
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_update_config_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
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_update_config_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 @@ -353,8 +353,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 @@ -367,13 +365,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 @@ -1349,10 +1340,9 @@ server's OpenCL/api-docs repository.
<unused start="5" end="31"/>
</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_update_config_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 @@ -3252,9 +3242,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_update_config_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 @@ -7294,18 +7286,17 @@ server's OpenCL/api-docs repository.
<enum name="CL_QUEUE_JOB_SLOT_ARM"/>
</require>
</extension>
<extension name="cl_khr_command_buffer_mutable_dispatch" requires="cl_khr_command_buffer" comment="version 0.9.0" supported="opencl">
<extension name="cl_khr_command_buffer_mutable_dispatch" requires="cl_khr_command_buffer" comment="version 0.9.1" supported="opencl">
<require>
<type name="CL/cl.h"/>
</require>
<require>
<type name="cl_command_buffer_structure_type_khr"/>
<type name="cl_update_config_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"/>
</require>
<require comment="cl_command_buffer_flags_khr - bitfield">
<enum name="CL_COMMAND_BUFFER_MUTABLE_KHR"/>
Expand Down Expand Up @@ -7338,8 +7329,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_update_config_type_khr">
<enum name="CL_STRUCTURE_TYPE_MUTABLE_DISPATCH_CONFIG_KHR"/>
</require>
<require>
Expand Down

0 comments on commit 8b0b63a

Please sign in to comment.