Use array for clUpdateMutableCommandsKHR.
Proposal to pass the update configs to `clUpdateMutableCommandsKHR` as
an array, rather than pointer changed linked list.
See https://github.com/KhronosGroup/OpenCL-Docs/issues/1041 for
motivation.
diff --git a/api/cl_khr_command_buffer_mutable_dispatch.asciidoc b/api/cl_khr_command_buffer_mutable_dispatch.asciidoc
index beda74a..0dca75d 100644
--- a/api/cl_khr_command_buffer_mutable_dispatch.asciidoc
+++ b/api/cl_khr_command_buffer_mutable_dispatch.asciidoc
@@ -6,7 +6,7 @@
=== Other Extension Metadata
*Last Modified Date*::
- 2022-08-31
+ 2024-03-22
*IP Status*::
No known IP claims.
*Contributors*::
@@ -43,39 +43,21 @@
=== 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_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 Types
* {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}
@@ -115,8 +97,7 @@
** {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, Error Codes>>
** {CL_INVALID_MUTABLE_COMMAND_KHR}
@@ -274,8 +255,6 @@
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 */,
@@ -287,12 +266,16 @@
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);
}
@@ -376,3 +359,6 @@
* 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-03-22
+ ** Change {clUpdateMutableCommandsKHR} API to pass configs as an array rather
+ than linked list (provisional).
diff --git a/api/opencl_runtime_layer.asciidoc b/api/opencl_runtime_layer.asciidoc
index cb79089..75d4bc4 100644
--- a/api/opencl_runtime_layer.asciidoc
+++ b/api/opencl_runtime_layer.asciidoc
@@ -15649,7 +15649,7 @@
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
@@ -15660,11 +15660,10 @@
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
@@ -15705,8 +15704,13 @@
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
@@ -15720,16 +15724,13 @@
* {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
@@ -15755,19 +15756,17 @@
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
@@ -15795,24 +15794,25 @@
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 is TODO Add fuller
-description here 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}
- * _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']
--
@@ -15822,9 +15822,6 @@
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}.
- * _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.
diff --git a/xml/cl.xml b/xml/cl.xml
index fa6b299..da5bc02 100644
--- a/xml/cl.xml
+++ b/xml/cl.xml
@@ -249,7 +249,7 @@
<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>
@@ -354,8 +354,6 @@
<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>
@@ -368,13 +366,6 @@
<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. -->
@@ -1354,10 +1345,9 @@
<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">
@@ -3259,9 +3249,11 @@
<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>
@@ -7308,18 +7300,17 @@
<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">
@@ -7353,8 +7344,7 @@
<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">
