Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

Use array for clUpdateMutableCommandsKHR. #1045

Merged
merged 1 commit into from
Jul 16, 2024
Merged

Use array for clUpdateMutableCommandsKHR. #1045

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
60 changes: 23 additions & 37 deletions api/cl_khr_command_buffer_mutable_dispatch.asciidoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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-06-19
*IP Status*::
No known IP claims.
*Contributors*::
Expand Down Expand Up @@ -43,32 +43,15 @@ 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_EXT}.
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_EXT} 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_EXT} 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 Commands

Expand All @@ -79,8 +62,7 @@ also be updated between enqueues of the command-buffer.

* {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}
Expand Down Expand Up @@ -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
** {CL_INVALID_MUTABLE_COMMAND_KHR}
Expand Down Expand Up @@ -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 */,
Expand All @@ -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);
}

Expand Down Expand Up @@ -374,3 +357,6 @@ may be a introduced as a stand alone extension.
* 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-06-19
** Change {clUpdateMutableCommandsKHR} API to pass configs as an array rather
than linked list (provisional).
93 changes: 45 additions & 48 deletions api/opencl_runtime_layer.asciidoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -15620,7 +15620,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
Expand All @@ -15631,11 +15631,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
Expand Down Expand Up @@ -15676,8 +15675,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

Expand All @@ -15691,16 +15695,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
Expand All @@ -15726,19 +15727,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
Comment on lines +15730 to +15734
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I think we should try to improve how this is described so we make the list of errors that may be returned a bit clearer. This is not something that this PR changes so I'm prefectly fine with capturing an issue that we'd look at separately.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Creating an issue makes sense, is there any particular way you imagine the improvement looking? e.g. Being more explicit about the errors rather than referencing core API entry-points, or just making the layout of the existing information more readable.

Copy link
Contributor

@kpet kpet Jul 10, 2024 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The first idea that came to mind was to categorise errors in the description of those commands such that we could refer to the categories here but I have not though enough about this to really have good input or opinions.

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Created issue #1209 to track this

{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
Expand Down Expand Up @@ -15766,24 +15765,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 encapsulates all aspects of
mutation 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_anchor}
* _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']
--
Expand All @@ -15793,9 +15793,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_anchor}.
* _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.
Expand Down
34 changes: 12 additions & 22 deletions xml/cl.xml
View file
Open in desktop
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_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>
Expand Down Expand Up @@ -370,8 +370,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 @@ -384,13 +382,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 @@ -1370,10 +1361,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">
Expand Down Expand Up @@ -3280,9 +3270,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>
Expand Down Expand Up @@ -7319,18 +7311,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">
Expand Down Expand Up @@ -7364,8 +7355,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">
Expand Down