Expand on command-buffer interactions (#803)

* Expand on command-buffer interactions

Clarifications to the cl_khr_command_buffer extension specification to address
feedback from
[Issue 799](#799).
Does not contain changes to the API or implementation behaviour.

* Elaborates on the role of the extension as a base which other layered
  extensions are intended to build upon to expand functionality. Removing
  outdated references to particular extension names. When a new layered
  extensions is released, then it can be linked back in here.

* The interaction of kernel commands with `clSetKernelArg` is documented as
  part of the *kernel* parameter description.

* Remove "mutable" from `cl_ndrange_kernel_command_properties_khr` type
  description, as other properties could end up being added here in layered
  extensions, not just those related to mutability.

* Update ext/cl_khr_command_buffer.asciidoc

Co-authored-by: jansol <jhs@psonet.com>

Co-authored-by: jansol <jhs@psonet.com>

Author: EwanC

ext/cl_khr_command_buffer.asciidoc

@@ -131,17 +131,59 @@ the capability is optional to enable optimizations on command-buffer recording.
 
 === Interactions with Other Extensions
 
-All command recording entry-points return a mutable-command handle for
-modifying the command between enqueues of the command-buffer. However,
-functionality for mutating commands is not supported by this
-extension. These are provided for other extensions layered on top to use,
-such as `cl_khr_mutable_dispatch`.
-
-Only a single command-queue specified on creation of a command-buffer can
-be used for recording commands to in this `cl_khr_command_buffer` extension.
-The `cl_khr_heterogeneous_replay` extension allows commands to be recorded across
-multiple queues in the same command-buffer, providing replay of heterogeneous
-task graphs.
+The introduction of the command-buffer abstraction enables functionality
+beyond what the `cl_khr_command_buffer` extension currently provides, i.e.
+the recording of immutable commands to a single queue which can then be
+executed without commands synchronizing outside the command-buffer. It is
+intended that extra functionality expanding on this will be provided as layered
+extensions on top of `cl_khr_command_buffer`.
+
+Having `cl_khr_command_buffer` as a minimal base specification means that the
+API defines mechanisms for functionality that is not enabled by this extension,
+these are described in the following sub-sections. `cl_khr_command_buffer` will
+retain its provisional extension status until other layered extensions are
+released, as these may reveal modifications needed to the base specification to
+support their intended use cases.
+
+==== NDRange Kernel Command Properties
+
+The `clCommandNDRangeKernelKHR` entry-point defines a `properties` parameter of
+new type `cl_ndrange_kernel_command_properties_khr`. No properties are defined
+in `cl_khr_command_buffer`, but the parameter is intended to enable future
+functionality that would change the characteristics of the kernel command.
+
+==== Command Handles
+
+All command recording entry-points define a `cl_mutable_command_khr` output
+parameter which provides a handle to the specific command being recorded. Use of
+these output handles is not enabled by the `cl_khr_command_buffer` extension,
+but the handles will allow individual commands in a command-buffer to be
+referenced by the user. In particular, the capability for an application to use
+these handles to modify commands between enqueues of a command-buffer is
+envisaged.
+
+==== List of Queues
+
+Only a single command-queue can be associated with a command-buffer in the
+`cl_khr_command_buffer` extension, but the API is designed with the intention
+that a future extension will allow commands to be recorded across multiple
+queues in the same command-buffer, providing replay of heterogeneous task
+graphs.
+
+Using multiple queue functionality will result in an error without any layered
+extensions to relax usage of the following API features:
+
+* When a command-buffer is created the API enables passing a list of queues
+  that the command-buffer will record commands to. Only a single queue is
+  permitted in `cl_khr_command_buffer`.
+
+* Individual command recording entry-points define a `cl_command_queue`
+  parameter for which of the queues set on command-buffer creation that command
+  should be record to. This must be passed as NULL in `cl_khr_command_buffer`.
+
+* `clEnqueueCommandBufferKHR` takes a list of queues for command-buffer execution,
+  correspond to those set on creation. Only a single queue is permitted in
+  `cl_khr_command_buffer`.
 
 === New Types
 
@@ -169,7 +211,7 @@ typedef cl_uint cl_sync_point_khr;
 // Handle returned on command recording
 typedef struct _cl_mutable_command_khr* cl_mutable_command_khr;
 
-// Mutable properties of a clCommandNDRangeKernelKHR command
+// Properties of a clCommandNDRangeKernelKHR command
 typedef cl_properties cl_ndrange_kernel_command_properties_khr;
 
 // Properties for command-buffer creation
@@ -1530,7 +1572,11 @@ corresponding desired value. The list is terminated with 0. If no properties are
 required, _properties_ may be `NULL`. This extension does not define any
 properties.
 
-_kernel_ A valid kernel object.
+_kernel_ A valid kernel object which **must** have its arguments set. Any
+changes to _kernel_ after calling *clCommandNDRangeKernelKHR*, such as with
+*clSetKernelArg* or *clSetKernelExecInfo*, have no effect on the recorded
+command. If _kernel_ is recorded to a following *clCommandNDRangeKernelKHR*
+command however, then that command will capture the updated state of _kernel_.
 
 _work_dim_, _global_work_offset_, _global_work_size_, _local_work_size_ Refer
 to *clEnqueueNDRangeKernel*.