setting kernel arguments before recording it in command buffer

.gitignore

@@ -2,9 +2,10 @@
 *~
 out/
 *.pyc
+gen/
 
 # Files generated from cl.xml
-generated/api
+generated/
 
 # Files generated by extraction from spec source
 man/*.txt

api/cl_khr_command_buffer_mutable_dispatch.asciidoc

@@ -111,6 +111,8 @@ void pointer using {cl_command_buffer_update_type_khr_TYPE}.
   ** {CL_COMMAND_BUFFER_MUTABLE_DISPATCH_ASSERTS_KHR}
   * {cl_command_buffer_update_type_khr_TYPE}
   ** {CL_STRUCTURE_TYPE_MUTABLE_DISPATCH_CONFIG_KHR}
+  * {cl_command_buffer_state_khr_TYPE}
+  ** {CL_COMMAND_BUFFER_STATE_FINALIZED_KHR}
   * New Error Codes
   ** {CL_INVALID_MUTABLE_COMMAND_KHR}
 
@@ -379,3 +381,6 @@ may be a introduced as a stand alone extension.
     **  Move `CL_COMMAND_BUFFER_SIMULTANEOUS_USE_KHR` and
         `CL_COMMAND_BUFFER_CAPABILITY_SIMULTANEOUS_USE_KHR` in this
         extension from the base extension (experimental).
+  * Revision 0.9.5 2025-08-28
+  ** Permitting recording ND-range kernel commands without having set all of
+     their arguments (experimental).

api/opencl_runtime_layer.asciidoc

@@ -14768,9 +14768,16 @@ the recording state.
 Recording:: Initial state of a command-buffer on creation, where commands can be
 recorded to the command-buffer.
 
+[[finalized]]
+Finalized:: State after command recording has finished with
+{clFinalizeCommandBufferKHR}, but there is at least one command for which not
+all arguments or parameters have been set (this is a valid state only for
+mutable command buffers).
+
 [[executable]]
 Executable:: State after command recording has finished with
-{clFinalizeCommandBufferKHR} and the command-buffer may be enqueued.
+{clFinalizeCommandBufferKHR} and all the arguments and parameters of all
+commands have been set. In this state the command-buffer may be enqueued.
 
 // Image generated from the following mermaid diagram description using https://mermaid.live
 // Ideally we'd use the asciidoctor-diagram extension to generate the rendered diagram, but
@@ -14780,7 +14787,9 @@ Executable:: State after command recording has finished with
 // ....
 // stateDiagram-v2
 //     [*] --> Recording: Create
-//     Recording -->Executable: Finalize
+//     Recording -->Finalized: Finalize
+//     Finalized --> Executable: All commands arguments/parameters set
+//     Executable
 // ....
 
 image::images/commandbuffer_lifecycle.png[align="center", title="Lifecycle of a command-buffer."]
@@ -15026,9 +15035,10 @@ include::{generated}/api/version-notes/clFinalizeCommandBufferKHR.asciidoc[]
 
 [NOTE]
 ====
-{clFinalizeCommandBufferKHR} places the command-buffer in the
-<<executable, Executable>> state where commands can no longer be recorded, at
-this point the command-buffer is ready to be enqueued.
+{clFinalizeCommandBufferKHR} places the command-buffer either in the
+<<finalized, Finalized>> or <<executable, Executable>> states where commands can
+no longer be recorded. If the new state is <<executable, Executable>>, the
+command-buffer is ready to be enqueued.
 ====
 
 // refError
@@ -16035,7 +16045,14 @@ ifdef::cl_khr_command_buffer_mutable_dispatch[]
     in the <<ndrange-kernel-properties-table, List of supported properties
     by clCommandNDRangeKernelKHR>> table.
 endif::cl_khr_command_buffer_mutable_dispatch[]
+ifndef::cl_khr_command_buffer_mutable_dispatch[]
   * _kernel_ is a valid kernel object which **must** have its arguments set.
+endif::cl_khr_command_buffer_mutable_dispatch[]
+ifdef::cl_khr_command_buffer_mutable_dispatch[]
+  * _kernel_ is a valid kernel object, which **must** have its arguments set
+    unless the command has the {CL_MUTABLE_DISPATCH_ARGUMENTS_KHR} property set
+    (see description of argument _mutable_handle_).
+endif::cl_khr_command_buffer_mutable_dispatch[]
     Any changes to _kernel_ after calling {clCommandNDRangeKernelKHR}, such
     as with {clSetKernelArg} or {clSetKernelExecInfo}, have no effect on the
     recorded command.
@@ -16213,6 +16230,14 @@ The starting local ID is always (0, 0, ... 0).
 successfully.
 Otherwise, it returns the errors defined by {clEnqueueNDRangeKernel} except:
 
+ifdef::cl_khr_command_buffer_mutable_dispatch[]
+{CL_INVALID_KERNEL_ARGS} is replaced with:
+
+  * {CL_INVALID_KERNEL_ARGS} if the kernel argument values have not been
+    specified and the {CL_MUTABLE_DISPATCH_ARGUMENTS_KHR} flag was not set in
+    the _properties_ parameter.
+endif::cl_khr_command_buffer_mutable_dispatch[]
+
 {CL_INVALID_COMMAND_QUEUE} is replaced with:
 
   * {CL_INVALID_COMMAND_QUEUE} if the
@@ -16653,6 +16678,13 @@ device occupancy high by avoiding blocking in host code.
 
 [open,refpage='clUpdateMutableCommandsKHR',desc='Modify configuration of mutable-command handles to update behavior for future enqueues',type='protos']
 --
+
+When recording an ND-range kernel command, the kernel's arguments do not have to
+be set, and setting them may be postponed to after the command buffer's
+finalization (in which case, the command buffer is in Finalized state).
+All the arguments must be set before the command buffer can be enqueued
+for execution (Executable state).
+
 To modify the configuration of mutable-command handles returned during
 _command_buffer_ recording, updating the behavior of those commands in
 future enqueues of _command_buffer_, call the function
@@ -16909,8 +16941,15 @@ include::{generated}/api/version-notes/CL_COMMAND_BUFFER_STATE_KHR.asciidoc[]
 
 include::{generated}/api/version-notes/CL_COMMAND_BUFFER_STATE_RECORDING_KHR.asciidoc[]
 
+        {CL_COMMAND_BUFFER_STATE_FINALIZED_KHR_anchor} is returned when an
+        instance of _command_buffer_ has been finalized, but there is at least
+        one command for which not all arguments or parameters have been set.
+
+include::{generated}/api/version-notes/CL_COMMAND_BUFFER_STATE_FINALIZED_KHR.asciidoc[]
+
         {CL_COMMAND_BUFFER_STATE_EXECUTABLE_KHR_anchor} is returned when
-        _command_buffer_ has been finalized.
+        _command_buffer_ has been finalized, all the arguments and parameters of
+        all commands have been set.
 
 include::{generated}/api/version-notes/CL_COMMAND_BUFFER_STATE_EXECUTABLE_KHR.asciidoc[]
 

xml/cl.xml

@@ -1371,6 +1371,7 @@ server's OpenCL/api-docs repository.
     <enums name="cl_command_buffer_state_khr" vendor="Khronos">
         <enum value="0"            name="CL_COMMAND_BUFFER_STATE_RECORDING_KHR"/>
         <enum value="1"            name="CL_COMMAND_BUFFER_STATE_EXECUTABLE_KHR"/>
+        <enum value="2"            name="CL_COMMAND_BUFFER_STATE_FINALIZED_KHR"/>
     </enums>
     <enums name="cl_mutable_dispatch_fields_khr" vendor="Khronos" type="bitmask">
         <enum bitpos="0"            name="CL_MUTABLE_DISPATCH_GLOBAL_OFFSET_KHR"/>
@@ -7507,7 +7508,7 @@ 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.4" supported="opencl" depends="cl_khr_command_buffer" ratified="opencl" experimental="true" comment="requires cl_khr_command_buffer 0.9.8 or later">
+        <extension name="cl_khr_command_buffer_mutable_dispatch" revision="0.9.5" supported="opencl" depends="cl_khr_command_buffer" ratified="opencl" experimental="true" comment="requires cl_khr_command_buffer 0.9.8 or later">
             <require>
                 <type name="CL/cl.h"/>
             </require>
@@ -7573,6 +7574,9 @@ server's OpenCL/api-docs repository.
                 <command name="clUpdateMutableCommandsKHR"/>
                 <command name="clGetMutableCommandInfoKHR"/>
             </require>
+            <require comment="cl_command_buffer_state_khr">
+                <enum name="CL_COMMAND_BUFFER_STATE_FINALIZED_KHR"/>
+            </require>
         </extension>
         <extension name="cl_ext_image_from_buffer" revision="1.0.0" depends="CL_VERSION_3_0+cl_ext_image_requirements_info" condition="defined(CL_VERSION_3_0)" supported="opencl">
             <require comment="cl_image_requirements_info_ext">