Use array for clUpdateMutableCommandsKHR.

Proposal to pass the update configs to `clUpdateMutableCommandsKHR` as
an array, rather than pointer changed linked list.

See #1041 for
motivation.

Author: Ewan Crawford

api/cl_khr_command_buffer_mutable_dispatch.asciidoc

@@ -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-03-22
 *IP Status*::
     No known IP claims.
 *Contributors*::
@@ -43,39 +43,21 @@ 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.
-
-[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.
+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 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 @@ 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, Error Codes>>
   ** {CL_INVALID_MUTABLE_COMMAND_KHR}
@@ -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 */,
@@ -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);
     }
 
@@ -376,3 +359,6 @@ include::provisional_notice.asciidoc[]
   * 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).

api/opencl_runtime_layer.asciidoc

@@ -15649,7 +15649,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
@@ -15660,11 +15660,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
@@ -15705,8 +15704,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
 
@@ -15720,16 +15724,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
@@ -15755,19 +15756,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
@@ -15795,24 +15794,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 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 @@ 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}.
-  * _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.

xml/cl.xml

@@ -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>
@@ -354,8 +354,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>
@@ -368,13 +366,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. -->
@@ -1354,10 +1345,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">
@@ -3259,9 +3249,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>
@@ -7308,18 +7300,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">
@@ -7353,8 +7344,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">