From 6ec48c4400ed69e6789f9834596c3e6fd19d99d3 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Pekka=20J=C3=A4=C3=A4skel=C3=A4inen?=
 <pekka.jaaskelainen@intel.com>
Date: Mon, 9 Sep 2024 16:54:49 +0300
Subject: [PATCH] cl_ext_buffer_device_address updates

Changed the CL_MEM_DEVICE_ADDRESS_EXT wording for multi-device
cases "all", not "any", covering a case where not all devices
can ensure the same address across the context. In that case
CL_INVALID_VALUE can be returned.  Defined sub-buffer address
computation to be 'base_addr + origin'. Added error conditions
for clSetKernelExecInfo when the device doesn't support
device pointers.
---
 .../cl_ext_buffer_device_address.asciidoc     | 38 ++++++++++++++++---
 1 file changed, 32 insertions(+), 6 deletions(-)

diff --git a/extensions/cl_ext_buffer_device_address.asciidoc b/extensions/cl_ext_buffer_device_address.asciidoc
index 72fe5529c..dabf5d1cf 100644
--- a/extensions/cl_ext_buffer_device_address.asciidoc
+++ b/extensions/cl_ext_buffer_device_address.asciidoc
@@ -46,7 +46,7 @@ Draft.
 == Version
 
 Built On: {docdate} +
-Revision: 0.1.0
+Revision: 0.2.0
 
 == Dependencies
 
@@ -147,7 +147,7 @@ Add new allocation flags <<clCreateBuffer, List of supported memory flag values
 include::{generated}/api/version-notes/CL_MEM_DEVICE_ADDRESS_EXT.asciidoc[]
   | This flag specifies that the buffer must have a single fixed address
    for its lifetime and the address should be unique at least across the devices
-   of the context, but not necessarily withing the host (virtual) memory.
+   of the context, but not necessarily within the host (virtual) memory.
 
    The flag might imply that the buffer will be "pinned" permanently to
    a device's memory, but might not be necessarily so, as long as the address
@@ -158,18 +158,28 @@ include::{generated}/api/version-notes/CL_MEM_DEVICE_ADDRESS_EXT.asciidoc[]
 
    The device-specific buffer content updates are still performed by
    implicit or explicit buffer migrations performed by the runtime or the
-   client code. If any of the devices in the context does not support
+   client code. If all of the devices in the context do not support
    this type of allocations, an error (CL_INVALID_VALUE) is returned.
+
+   The device addresses of sub-buffers derived from CL_MEM_DEVICE_ADDRESS_EXT
+   allocated buffers can be computed by adding the sub-buffer origin to the
+   start address.
+
 | {CL_MEM_DEVICE_PRIVATE_EXT_anchor}
 
 include::{generated}/api/version-notes/CL_MEM_DEVICE_PRIVATE_EXT.asciidoc[]
   | If this flag is combined with CL_MEM_DEVICE_ADDRESS_EXT, each device in
-   the context can have their own (fixed) device-side address and copy of
+   the context can have their own (fixed) device-side address and a copy of
    the created buffer which are synchronized implicitly by the runtime.
    The main difference to a default cl_mem allocation in that case is then
    that the addresses are queriable with CL_MEM_DEVICE_PTRS_EXT and the
    per-device address is guaranteed to be the same for the entire lifetime
    of the cl_mem.
+
+   The device addresses of sub-buffers derived from CL_MEM_DEVICE_PRIVATE_EXT
+   allocated buffers can be computed by adding the sub-buffer origin to the
+   device-specific start address.
+
 |====
 
 // refError
@@ -218,7 +228,7 @@ include::{generated}/api/version-notes/clSetKernelArgDevicePointerEXT.asciidoc[]
     argument specified by _arg_index_.
     The device pointer specified is the value used by all API calls that enqueue
     _kernel_ ({clEnqueueNDRangeKernel} and {clEnqueueTask}) until the argument
-    value is changed by a call to {clSetKernelArgSVMPointer} for _kernel_.
+    value is changed by a call to {clSetKernelArgDevicePointer} for _kernel_.
     The device pointer can only be used for arguments that are declared to be a
     pointer to `global` memory allocated with clCreateBuffer() with the
     CL_MEM_DEVICE_ADDRESS_EXT flag. The pointer value specified as the argument value
@@ -260,6 +270,14 @@ include::{generated}/api/version-notes/CL_KERNEL_EXEC_INFO_DEVICE_PTRS_EXT.ascii
 
 // refError
 
+Change the descriptions for when returning CL_INVALID_OPERATION from {clSetKernelExecInfo}
+as follows:
+
+  * {CL_INVALID_OPERATION} if passing {CL_KERNEL_EXEC_INFO_SVM_PTRS} or
+  {CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM} with _param_value_ set to CL_TRUE
+  and no device in the context associated with _kernel_ support SVM.
+  * {CL_INVALID_OPERATION} if passing {CL_KERNEL_EXEC_INFO_DEVICE_PTRS_EXT} and no
+  device in the context associated with _kernel_ support device pointers.
 
 == Interactions with Other Extensions
 
@@ -280,7 +298,15 @@ None.
 [grid="rows"]
 [options="header"]
 |====
-| *Version* | *Date*       | *Author*             | *Changes*
+| *Version* | *Date*       | *Author*                         | *Changes*
+| 0.2.0     | 2024-09-09   | Pekka Jääskeläinen, Karol Herbst |
+                           Changed the CL_MEM_DEVICE_ADDRESS_EXT wording for multi-device
+                           cases "all", not "any", covering a case where not all devices
+                           can ensure the same address across the context. In that case
+                           CL_INVALID_VALUE can be returned.  Defined sub-buffer address
+                           computation to be 'base_addr + origin'. Added error conditions
+                           for clSetKernelExecInfo when the device doesn't support
+                           device pointers.
 | 0.1.0     | 2024-05-07   | Pekka Jääskeläinen   | First draft text for feedback.
                            This version describes the first API version that was prototyped
                            in PoCL and RustiCL using temporary placeholder flag/enum values.