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">