clarify the conditions where the global work size may be null or zero (#1381)

Author: bashbaug

api/opencl_runtime_layer.asciidoc

@@ -11909,12 +11909,6 @@ include::{generated}/api/version-notes/clEnqueueNDRangeKernel.asciidoc[]
     and work-items in the work-group.
     _work_dim_ must be greater than zero and less than or equal to
     {CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS}.
-    If _global_work_size_ is `NULL`, or the value in any passed dimension is 0
-    then the kernel command will trivially succeed after its event dependencies
-    are satisfied and subsequently update its completion event.
-    The behavior in this situation is similar to that of an enqueued marker,
-    except that unlike a marker, an enqueued kernel with no events passed to
-    _event_wait_list_ may run at any time.
   * _global_work_offset_ can be used to specify an array of _work_dim_ unsigned
     values that describe the offset used to calculate the global ID of a
     work-item.
@@ -11925,6 +11919,14 @@ include::{generated}/api/version-notes/clEnqueueNDRangeKernel.asciidoc[]
     execute the kernel function.
     The total number of global work-items is computed as _global_work_size_[0]
     {times} ... {times} _global_work_size_[_work_dim_ - 1].
+    If the device associated with _command_queue_ is an OpenCL 2.1 or newer
+    device, and _global_work_size_ is `NULL` or the value in any passed
+    dimension is zero,
+    then the kernel command will trivially succeed after its event dependencies
+    are satisfied and subsequently update its completion event.
+    The behavior in this situation is similar to that of an enqueued marker,
+    except that unlike a marker, an enqueued kernel with no events passed to
+    _event_wait_list_ may run at any time.
   * _local_work_size_ points to an array of _work_dim_ unsigned values that
     describe the number of work-items that make up a work-group (also referred
     to as the size of the work-group) that will execute the kernel specified by
@@ -12016,7 +12018,7 @@ successfully queued.
 Otherwise, it returns one of the following errors:
 
   * {CL_INVALID_PROGRAM_EXECUTABLE} if there is no successfully built program
-    executable available for device associated with _command_queue_.
+    executable available for the device associated with _command_queue_.
   * {CL_INVALID_COMMAND_QUEUE} if _command_queue_ is not a valid host
     command-queue.
   * {CL_INVALID_KERNEL} if _kernel_ is not a valid kernel object.
@@ -12029,9 +12031,9 @@ Otherwise, it returns one of the following errors:
     value between 1 and {CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS}).
   * {CL_INVALID_GLOBAL_WORK_SIZE} if _global_work_size_ is NULL or if any of
     the values specified in _global_work_size_[0], ...
-    _global_work_size_[_work_dim_ - 1] are 0.
-    Returning this error code under these circumstances is  <<unified-spec,
-    deprecated by>> version 2.1.
+    _global_work_size_[_work_dim_ - 1] are zero.
+    This error condition does not apply when the device associated with
+    _command_queue_ supports OpenCL 2.1 or newer.
   * {CL_INVALID_GLOBAL_WORK_SIZE} if any of the values specified in
     _global_work_size_[0], ... _global_work_size_[_work_dim_ - 1] exceed the
     maximum value representable by {size_t_TYPE} on the device on which the