Merge branch 'main' into error-code-consistency

Author: bashbaug

.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/cl_khr_unified_svm.asciidoc

@@ -0,0 +1,101 @@
+// Copyright 2025 The Khronos Group Inc.
+// SPDX-License-Identifier: CC-BY-4.0
+
+include::{generated}/meta/{refprefix}cl_khr_unified_svm.txt[]
+
+=== Other Extension Metadata
+
+*Last Modified Date*::
+    2025-08-30
+*IP Status*::
+    No known IP claims.
+*Contributors*::
+  - Brice Videau, Argonne National Laboratory
+  - Kévin Petit, Arm Ltd.
+  - Ewan Crawford, Codeplay Software Ltd.
+  - Paul Fradgley, Imagination Technologies
+  - Ben Ashbaugh, Intel
+  - Pekka Jääskeläinen, Intel
+  - Nikhil Joshi, NVIDIA
+  - Balaji Calidas, Qualcomm Technologies Inc.
+
+=== Description
+
+This extension adds additional types of Shared Virtual Memory (SVM) to OpenCL.
+
+The extension is still under development.
+To learn more about the APIs proposed for this extension, or to provide feedback, please visit:
+
+https://github.com/KhronosGroup/OpenCL-Docs/pull/1282
+
+=== New Commands
+
+  * {clSVMAllocWithPropertiesKHR}[[clSVMAllocWithPropertiesKHR]]
+  * {clSVMFreeWithPropertiesKHR}[[clSVMFreeWithPropertiesKHR]]
+  * {clGetSVMPointerInfoKHR}[[clGetSVMPointerInfoKHR]]
+  * {clGetSVMSuggestedTypeIndexKHR}[[clGetSVMSuggestedTypeIndexKHR]]
+
+=== New Types
+
+  * {cl_svm_capabilities_khr_TYPE}
+  * {cl_svm_alloc_properties_khr_TYPE}
+  * {cl_svm_alloc_access_flags_khr_TYPE}
+  * {cl_svm_free_properties_khr_TYPE}
+  * {cl_svm_free_flags_khr_TYPE}
+  * {cl_svm_pointer_info_khr_TYPE}
+
+=== New Enums
+
+  * {cl_platform_info_TYPE}
+  ** {CL_PLATFORM_SVM_TYPE_CAPABILITIES_KHR_ANCHOR}
+  * {cl_device_info_TYPE}
+  ** {CL_DEVICE_SVM_TYPE_CAPABILITIES_KHR_ANCHOR}
+  * {cl_svm_capabilities_khr_TYPE}
+  ** {CL_SVM_CAPABILITY_SINGLE_ADDRESS_SPACE_KHR_ANCHOR}
+  ** {CL_SVM_CAPABILITY_SYSTEM_ALLOCATED_KHR_ANCHOR}
+  ** {CL_SVM_CAPABILITY_DEVICE_OWNED_KHR_ANCHOR}
+  ** {CL_SVM_CAPABILITY_DEVICE_UNASSOCIATED_KHR_ANCHOR}
+  ** {CL_SVM_CAPABILITY_CONTEXT_ACCESS_KHR_ANCHOR}
+  ** {CL_SVM_CAPABILITY_HOST_OWNED_KHR_ANCHOR}
+  ** {CL_SVM_CAPABILITY_HOST_READ_KHR_ANCHOR}
+  ** {CL_SVM_CAPABILITY_HOST_WRITE_KHR_ANCHOR}
+  ** {CL_SVM_CAPABILITY_HOST_MAP_KHR_ANCHOR}
+  ** {CL_SVM_CAPABILITY_DEVICE_READ_KHR_ANCHOR}
+  ** {CL_SVM_CAPABILITY_DEVICE_WRITE_KHR_ANCHOR}
+  ** {CL_SVM_CAPABILITY_DEVICE_ATOMIC_ACCESS_KHR_ANCHOR}
+  ** {CL_SVM_CAPABILITY_CONCURRENT_ACCESS_KHR_ANCHOR}
+  ** {CL_SVM_CAPABILITY_CONCURRENT_ATOMIC_ACCESS_KHR_ANCHOR}
+  ** {CL_SVM_CAPABILITY_INDIRECT_ACCESS_KHR_ANCHOR}
+  * {cl_svm_alloc_properties_khr_TYPE}
+  ** {CL_SVM_ALLOC_ASSOCIATED_DEVICE_HANDLE_KHR_ANCHOR}
+  ** {CL_SVM_ALLOC_ACCESS_FLAGS_KHR_ANCHOR}
+  ** {CL_SVM_ALLOC_ALIGNMENT_KHR_ANCHOR}
+  * {cl_svm_alloc_access_flags_khr_TYPE}
+  ** {CL_SVM_ALLOC_ACCESS_HOST_NOREAD_KHR_ANCHOR}
+  ** {CL_SVM_ALLOC_ACCESS_HOST_NOWRITE_KHR_ANCHOR}
+  ** {CL_SVM_ALLOC_ACCESS_DEVICE_NOREAD_KHR_ANCHOR}
+  ** {CL_SVM_ALLOC_ACCESS_DEVICE_NOWRITE_KHR_ANCHOR}
+  * {cl_svm_pointer_info_khr_TYPE}
+  ** {CL_SVM_INFO_TYPE_INDEX_KHR_ANCHOR}
+  ** {CL_SVM_INFO_CAPABILITIES_KHR_ANCHOR}
+  ** {CL_SVM_INFO_PROPERTIES_KHR_ANCHOR}
+  ** {CL_SVM_INFO_ACCESS_FLAGS_KHR_ANCHOR}
+  ** {CL_SVM_INFO_BASE_PTR_KHR_ANCHOR}
+  ** {CL_SVM_INFO_SIZE_KHR_ANCHOR}
+  ** {CL_SVM_INFO_ASSOCIATED_DEVICE_HANDLE_KHR_ANCHOR}
+  * {cl_kernel_exec_info_TYPE}
+  ** {CL_KERNEL_EXEC_INFO_SVM_INDIRECT_ACCESS_KHR_ANCHOR}
+
+=== SVM Type Convenience Macros
+
+  * {CL_SVM_TYPE_MACRO_COARSE_GRAIN_BUFFER_KHR_ANCHOR}
+  * {CL_SVM_TYPE_MACRO_FINE_GRAIN_BUFFER_KHR_ANCHOR}
+  * {CL_SVM_TYPE_MACRO_DEVICE_KHR_ANCHOR}
+  * {CL_SVM_TYPE_MACRO_HOST_KHR_ANCHOR}
+  * {CL_SVM_TYPE_MACRO_SINGLE_DEVICE_SHARED_KHR_ANCHOR}
+  * {CL_SVM_TYPE_MACRO_SYSTEM_KHR_ANCHOR}
+
+=== Version History
+
+  * Revision 0.9.0, 2025-08-30
+  ** Initial revision incorporated into the main specification (experimental).

api/opencl_runtime_layer.asciidoc

@@ -1247,8 +1247,8 @@ Otherwise, they return one of the following errors:
     ** if _buffer_origin_, _host_origin_, or _region_ is `NULL`
     ** if the region being read or written specified by _buffer_origin_, _region_, _buffer_row_pitch_, and _buffer_slice_pitch_ is out of bounds
     ** if any _region_ array element is zero
-    ** if _buffer_row_pitch_ is not equal to zero and is less than _region_[0].
-    ** if _buffer_slice_pitch_ is not equal to zero and is less than the_region_[1] times _buffer_row_pitch_
+    ** if _buffer_row_pitch_ is not equal to zero and is less than _region_[0]
+    ** if _buffer_slice_pitch_ is not equal to zero and is less than _region_[1] times _buffer_row_pitch_
     ** if _buffer_slice_pitch_ is not equal to zero and is not a multiple of _buffer_row_pitch_
     ** if _host_row_pitch_ is not equal to zero and is less than _region_[0]
     ** if _host_slice_pitch_ is not equal to zero and is less than _region_[1] times _host_row_pitch_
@@ -1415,7 +1415,7 @@ Otherwise, it returns one of the following errors:
     ** if _event_wait_list_ is not `NULL` and _num_events_in_wait_list_ is zero
     ** if event objects in _event_wait_list_ are not valid events
   * {CL_MEM_COPY_OVERLAP}
-    ** if _src_buffer_ and _dst_buffer_ are the same buffer or sub-buffer object and the source and destination regions overlap.
+    ** if _src_buffer_ and _dst_buffer_ are the same buffer or sub-buffer object and the source and destination regions overlap
     ** if _src_buffer_ and _dst_buffer_ are different sub-buffers of the same associated buffer object and they overlap
     ** Refer to <<check-copy-overlap, Checking for Memory Copy Overlap>> for
     details on how to determine if source and destination regions overlap.
@@ -1563,7 +1563,7 @@ Otherwise, it returns one of the following errors:
     ** if _event_wait_list_ is not `NULL` and _num_events_in_wait_list_ is zero
     ** if event objects in _event_wait_list_ are not valid events
   * {CL_MEM_COPY_OVERLAP}
-    ** if _src_buffer_ and _dst_buffer_ are the same buffer or sub-buffer object and the source and destination regions overlap.
+    ** if _src_buffer_ and _dst_buffer_ are the same buffer or sub-buffer object and the source and destination regions overlap
     ** if _src_buffer_ and _dst_buffer_ are different sub-buffers of the same associated buffer object and they overlap
     ** Refer to <<check-copy-overlap, Checking for Memory Copy Overlap>> for
     details on how to determine if source and destination regions overlap.
@@ -1657,7 +1657,7 @@ Otherwise, it returns one of the following errors:
   * {CL_INVALID_MEM_OBJECT}
     ** if _buffer_ is not a valid buffer object
   * {CL_INVALID_VALUE}
-    ** if the region being read or written specified by _offset_ and _size_ is out of bounds
+    ** if the region being filled specified by _offset_ and _size_ is out of bounds
     ** if _pattern_ is `NULL`
     ** if _pattern_size_ is zero
     ** if _pattern_size_ is not a power of two
@@ -9806,6 +9806,8 @@ successfully.
 Otherwise, it returns one of the following errors:
 
   * {CL_INVALID_PROGRAM} if _program_ is not a valid program object.
+  * {CL_INVALID_CONTEXT} if the context associated with _program_ and
+    programs in _input_headers_ are not the same.
   * {CL_INVALID_VALUE} if _device_list_ is `NULL` and _num_devices_ is greater
     than zero, or if _device_list_ is not `NULL` and _num_devices_ is zero.
   * {CL_INVALID_VALUE} if _num_input_headers_ is zero and
@@ -9945,6 +9947,8 @@ check if the link was successful or not.
 The list of errors that can be returned are:
 
   * {CL_INVALID_CONTEXT} if _context_ is not a valid context.
+  * {CL_INVALID_CONTEXT} if the context associated with programs in
+    _input_programs_ is not the same as _context_.
   * {CL_INVALID_VALUE} if _device_list_ is `NULL` and _num_devices_ is greater
     than zero, or if _device_list_ is not `NULL` and _num_devices_ is zero.
   * {CL_INVALID_VALUE} if _num_input_programs_ is zero and _input_programs_ is
@@ -14843,9 +14847,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
@@ -14855,7 +14866,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."]
@@ -15101,9 +15114,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
@@ -16110,7 +16124,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.
@@ -16288,6 +16309,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
@@ -16728,6 +16757,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
@@ -16984,8 +17020,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[]
 

env/validation_rules.asciidoc

@@ -32,12 +32,9 @@ For all *OpTypeImage* type-declaration instructions:
   * _Image Format_ must be *Unknown*, indicating that the image does not
      have a specified format.
   * The optional image _Access Qualifier_ must be present.
-  
-The image write instruction *OpImageWrite* must not include any optional
-_Image Operands_.
 
-The image read instructions *OpImageRead* and *OpImageSampleExplicitLod* 
-must not include the optional _Image Operand_ *ConstOffset*.
+For all image read and write instructions, if an optional _Image Operand_ is
+present, then it must not include *ConstOffset*.
 
 For all *Atomic Instructions*: