spec source for OpenCL 3.1 (#1554)

Author: bashbaug

.gitlab-ci.yml

@@ -0,0 +1,46 @@
+# Copyright 2018-2025 The Khronos Group Inc.
+
+# Gitlab CI file for OpenCL specs
+
+# All stages use the same Docker image, so there are no prerequisites
+# Refer to the container by its SHA instead of the name, to prevent
+# caching problems when updating the image.
+# image: khronosgroup/docker-images:asciidoctor-spec.202506
+# There are no 'before_script' tags for most stages, because all
+# prerequisites are in the default image.
+image: khronosgroup/docker-images@sha256:0f91e60e1af2bdd889783af3907f63279c08f573f2eccbc31094e348d1a32a4f
+
+# Specify which gitlab runner to use
+default:
+  tags:
+    - khrmedium
+
+# Build the OpenCL specification
+spec-generate:
+  stage: build
+  script:
+    - make -C xml validate
+    - python3 makeSpec -clean -spec core    OUTDIR=out.core     -j -O api c env
+    - python3 makeSpec -clean -spec khr+ext OUTDIR=out.ext      -j -O api c env
+    - python3 makeSpec -clean -spec khr+ext OUTDIR=out.refpages -j -O manhtmlpages
+  artifacts:
+    when: always
+    paths:
+      - out.core/
+      - out.ext/
+      - out.refpages/
+    expire_in: 1 week
+
+spec-deploy-pages:
+  stage: deploy
+  script:
+    - mkdir -p public
+    - cp -r out.ext/* public/
+    - cp -r out.refpages/* public/
+  artifacts:
+    paths:
+      - public
+  pages: true
+  rules:
+    - if: '$CI_COMMIT_BRANCH == "github"'
+    - if: '$CI_COMMIT_TAG =~ /^v3\.1/'

OpenCL_C.txt

@@ -38,6 +38,8 @@ MonPing Wang, Apple +
 Tanya Lattner, Apple +
 Mikael Bourges-Sevenier, Aptina +
 Brice Videau, Argonne National Laboratory +
+Thomas Applencourt, Argonne National Laboratory +
+Nevin Liber, Argonne National Laboratory +
 Anton Lokhmotov, ARM +
 Dave Shreiner, ARM +
 Einar Hov, ARM +
@@ -58,11 +60,15 @@ Maria Rovatsou, Codeplay +
 Alistair Donaldson, Codeplay +
 Alastair Murray, Codeplay +
 Ewan Crawford, Codeplay +
+Ewa Gałamon, Cognizant +
+Marcin Hajder, Cognizant +
+Faith Ekstrand, Collabora +
 Stephen Frye, Electronic Arts +
 Eric Schenk, Electronic Arts +
 Daniel Laroche, Freescale +
 David Neto, Google +
 James Price, Google +
+Romaric Jodin, Google +
 Robin Grosman, Huawei +
 Craig Davies, Huawei +
 Brian Horton, IBM +
@@ -73,6 +79,7 @@ Joaquin Madruga, IBM +
 Mark Nutter, IBM +
 Mike Perks, IBM +
 Sean Wagner, IBM +
+Ahmed Amrani Akdi, Imagination Technologies +
 Jeremy Kemp, Imagination Technologies +
 Jon Parr, Imagination Technologies +
 Paul Fradgley, Imagination Technologies +
@@ -98,6 +105,7 @@ Hong Jiang, Intel +
 Jayanth Rao, Intel +
 Josh Fryman, Intel +
 Kevin Stevens, Intel +
+Konrad Trifunovic, Intel +
 Larry Seiler, Intel +
 Michael Kinsner, Intel +
 Michal Mrozek, Intel +
@@ -114,6 +122,7 @@ Roy Ju, Mediatek +
 Bor-Sung Liang, Mediatek +
 Rahul Agarwal, Mediatek +
 Michal Witaszek, Mobica +
+Aharon Abramson, Mobileye +
 JenqKuen Lee, NTHU +
 Amit Rao, NVIDIA +
 Ashish Srivastava, NVIDIA +
@@ -142,9 +151,11 @@ Yuan Lin, NVIDIA +
 Mayuresh Pise, NVIDIA +
 Allan Tzeng, QUALCOMM +
 Alex Bourd, QUALCOMM +
+Alexander Galazin, QUALCOMM +
 Andrew Gruber, QUALCOMM +
 Andrzej Mamona, QUALCOMM +
 Anirudh Acharya, QUALCOMM +
+Arvind Sudarsanam, QUALCOMM +
 Balaji Calidas, QUALCOMM +
 Benedict Gaster, QUALCOMM +
 Bill Torzewski, QUALCOMM +
@@ -157,6 +168,7 @@ David Ligon, QUALCOMM +
 Hongqiang Wang, QUALCOMM +
 Jay Yun, QUALCOMM +
 Jian Liu, QUALCOMM +
+Jose Lopez, QUALCOMM +
 Joshua Kelly, QUALCOMM +
 Lee Howes, QUALCOMM +
 Lihan Bin, QUALCOMM +
@@ -170,7 +182,11 @@ Vineet Goel, QUALCOMM +
 Vlad Shimanskiy, QUALCOMM +
 Yu-Chi Huang, QUALCOMM +
 Yuehai Du, QUALCOMM +
+Karol Herbst, Red Hat +
+Austin Annestrand, Samsung +
 Raun Krisch, Samsung +
+Gowtham Tammana, Samsung +
+Pavan Lanka, Samsung +
 Tasneem Brutch, Samsung +
 Yoonseo Choi, Samsung +
 Dennis Adams, Sony +
@@ -181,6 +197,7 @@ Anton Gorenko, StreamHPC +
 Jakub Szuppe, StreamHPC +
 Máté Ferenc Nagy-Egri, StreamHPC +
 Vincent Hindriksen, StreamHPC +
+Pekka Jääskeläinen, Tampere University +
 Ajay Jayaraj, Texas Instruments +
 Alan Ward, Texas Instruments +
 Yuan Zhao, Texas Instruments +

api/acknowledgements.asciidoc

@@ -165,7 +165,7 @@ For example:
 
 Vector data type components may also be accessed using the `.rgba` field
 naming convention, similar to how they are used within the OpenCL C 3.0
-language.
+or newer language.
 Use of the `.rgba` field naming convention only allows accessing of the
 first 4 component fields.
 Support of these notations is identified by the `CL_HAS_NAMED_RGBA_VECTOR_FIELDS`

api/appendix_c.asciidoc

@@ -200,7 +200,7 @@ runtime (_sections 4 and 5_):
 
   * Shared virtual memory.  The associated API additions are:
   ** {clSetKernelArgSVMPointer} to control which shared virtual memory (SVM)
-     pointer to associate with a kernel instance.
+     pointer to associate with a kernel-instance.
   ** {clSVMAlloc}, {clSVMFree} and {clEnqueueSVMFree} to allocate and free
      memory for use with SVM.
   ** {clEnqueueSVMMap} and {clEnqueueSVMUnmap} to map and unmap to update
@@ -326,6 +326,7 @@ The OpenCL 2.0 kernel language will still be consumed by OpenCL 2.1
 runtimes.
 
 The SPIR-V and OpenCL SPIR-V Environment specifications have been added.
+OpenCL 2.1 requires support for the SPIR-V 1.0 intermediate language.
 
 == Summary of Changes from OpenCL 2.1 to OpenCL 2.2
 
@@ -353,6 +354,9 @@ runtime (section 4 and 5):
 Added definition of Deprecation and Specialization constants to the
 glossary.
 
+OpenCL 2.2 requires support for the SPIR-V 1.0, SPIR-V 1.1, and SPIR-V 1.2
+intermediate languages.
+
 == Summary of Changes from OpenCL 2.2 to OpenCL 3.0
 
 OpenCL 3.0 is a major revision that breaks backwards compatibility with
@@ -457,8 +461,6 @@ conformance process:
 
   * {CL_DEVICE_LATEST_CONFORMANCE_VERSION_PASSED}
 
-== Summary of Changes to OpenCL 3.0
-
 The first non-experimental version of the OpenCL 3.0 specifications was *v3.0.5*.
 
 Changes from *v3.0.5* to *v3.0.6*:
@@ -597,7 +599,7 @@ Changes from *v3.0.14* to *v3.0.15*:
       ** Added an semaphore-specific device handle list enum, see {khronos-opencl-pr}/956[#956].
       ** Restricted semaphores to a single associated device, see {khronos-opencl-pr}/996[#996].
   * {cl_khr_subgroup_rotate_EXT}:
-      ** Clarified that only rotating within a subgroup is supported, see {khronos-opencl-pr}/967[#967].
+      ** Clarified that only rotating within a sub-group is supported, see {khronos-opencl-pr}/967[#967].
 
 Changes from *v3.0.15* to *v3.0.16*:
 
@@ -707,5 +709,54 @@ Changes from *v3.0.18* to *v3.0.19*:
   * The following extension has been finalized and is no longer experimental:
       ** {cl_khr_kernel_clock_EXT}
   * Added new extensions:
-      ** {cl_khr_external_memory_android_hardware_buffer_EXT}
-      ** {cl_khr_spirv_queries_EXT}
+      ** {cl_khr_external_memory_android_hardware_buffer_EXT} (experimental)
+      ** {cl_khr_spirv_queries_EXT}
+
+== Summary of Changes from OpenCL 3.0 to OpenCL 3.1
+
+OpenCL 3.1 adds the OpenCL 3.1 C kernel language.
+Please refer to the OpenCL C specification for details.
+
+OpenCL 3.1 requires support for the SPIR-V 1.0, SPIR-V 1.1, SPIR-V 1.2, SPIR-V
+1.3, and SPIR-V 1.4 intermediate languages.
+Please refer to the OpenCL SPIR-V Environment specification for details.
+
+OpenCL 3.1 removes the OpenCL Extension Specification.
+The OpenCL Extension Specification is no longer needed now that EXT and KHR extensions are documented in the main OpenCL specifications.
+See {khronos-opencl-pr}/1516[#1516].
+
+Other changes in OpenCL 3.1:
+
+  * Added required support for sub-groups, see {CL_DEVICE_MAX_NUM_SUB_GROUPS}.
+  * Relaxed the definition of inclusive scopes, see internal issue 367.
+  * Clarified and un-deprecated the {CL_DEVICE_HOST_UNIFIED_MEMORY} query, see internal issue 370.
+  * Updated the memory model so observing that an event is {CL_COMPLETE} is a synchronization point, see internal issue 373.
+  * Allowed {clSetKernelArg} to set a local memory kernel argument to zero, see internal issue 374.
+  * Deprecated the confusingly named {CL_DEVICE_MAX_WORK_ITEM_SIZES} query, use {CL_DEVICE_MAX_WORK_GROUP_SIZES} instead, see internal issue 375.
+  * Cleaned up inconsistencies in the error condition descriptions for {clSetKernelExecInfo}, see {khronos-opencl-pr}/1419[#1419].
+  * Improved error code documentation consistency for many OpenCL APIs, see {khronos-opencl-pr}/1439[#1439] and others.
+  * Added missing error conditions for {clCompileProgram} and {clLinkProgram}, see {khronos-opencl-pr}/1453[#1453].
+  * Clarified the description of sub-group functions, see {khronos-opencl-pr}/1483[#1483].
+  * Refactored and clarified the description for {clSetKernelArg}, see {khronos-opencl-pr}/1493[#1493].
+  * Clarified the description of atomic operations, see {khronos-opencl-pr}/1500[#1500].
+  * Cleaned up a few more descriptions of custom devices, see {khronos-opencl-pr}/1540[#1540].
+  * Added missing error condition for {clEnqueueNDRangeKernel} when the local work-group size is zero, see {khronos-opencl-pr}/1542[#1542].
+  * {cl_khr_command_buffer_mutable_dispatch_EXT} (experimental):
+      ** Relaxed the requirement to set all kernel arguments before recording a command buffer, see {khronos-opencl-pr}/1382[#1382].
+      ** Redefined and clarified command-buffer simultaneous use, see {khronos-opencl-pr}/1411[#1411].
+  * {cl_khr_external_memory_android_hardware_buffer_EXT} (experimental)
+      ** Clarified that images cannot be created if the format is `AHARDWAREBUFFER_FORMAT_BLOB`, see {khronos-opencl-pr}/1477[#1477].
+  * {cl_ext_buffer_device_address}
+      ** Added a missing error condition for {clSetKernelArgDevicePointerEXT}, see {khronos-opencl-pr}/1493[#1492].
+  * Added new extensions:
+      ** {cl_khr_unified_svm_EXT} (experimental)
+  * Promoted the following extensions to the core API:
+      ** {cl_khr_device_uuid_EXT}
+      ** {cl_khr_extended_bit_ops_EXT}
+      ** {cl_khr_integer_dot_product_EXT}
+      ** {cl_khr_spirv_queries_EXT}
+      ** {cl_khr_subgroup_extended_types_EXT}
+      ** {cl_khr_subgroup_rotate_EXT}
+      ** {cl_khr_subgroup_shuffle_EXT}
+      ** {cl_khr_subgroup_shuffle_relative_EXT}
+      ** {cl_khr_suggested_local_work_size_EXT}

api/appendix_e.asciidoc

@@ -30,7 +30,7 @@ include::{generated}/meta/interfaces/cl_ext_immutable_memory_objects.txt[]
 1) Can {CL_MEM_READ_ONLY} be used instead of {CL_MEM_IMMUTABLE_EXT}?
 --
 *RESOLVED*: No. Memory objects created with {CL_MEM_READ_ONLY} can be modified
-by copy or fill commands and this behaviour cannot be changed without breaking
+by copy or fill commands and this behavior cannot be changed without breaking
 backwards compatibility.
 --
 

api/cl_ext_immutable_memory_objects.asciidoc

@@ -103,7 +103,7 @@ Here's a recommended policy:
   Adding {CL_DEVICE_BUILT_IN_KERNELS_WITH_VERSION_KHR}.
 --
 
-. What is the behaviour of the queries that return an array of structures when
+. What is the behavior of the queries that return an array of structures when
 there are no elements to return?
 +
 --

api/cl_khr_extended_versioning.asciidoc

@@ -323,5 +323,5 @@ profile devices is:
 |====
 
 For embedded profiles devices that support reading from and writing to the same
-image object from the same kernel instance (see {CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS})
+image object from the same kernel-instance (see {CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS})
 there is no required minimum list of supported image formats.

api/embedded_profile.asciidoc

@@ -48,6 +48,12 @@ The value of {CL_COMPLETE} and {CL_SUCCESS} are the same. \
 {clGetDeviceIDs} may return all or a subset of the actual physical devices present in the platform and that match _device_type_. \
 ]
 
+:fn-host-unified-memory: pass:n[ \
+When the query for {CL_DEVICE_HOST_UNIFIED_MEMORY} is {CL_TRUE}, allocating OpenCL memory will likely reduce the amount of host memory available to the system. \
+Likewise, allocating host memory will likely reduce the amount of memory available to OpenCL. \
+If the memory described by {CL_DEVICE_GLOBAL_MEM_SIZE} is primarily host memory, such as for CPUs, integrated GPUs, and other devices with a relatively small amount of dedicated device memory, then the device should return {CL_TRUE} for {CL_DEVICE_HOST_UNIFIED_MEMORY}. \
+]
+
 :fn-image-array-performance: pass:n[ \
 Note that reading and writing 2D image arrays from a kernel with `image_array_size` equal to one may perform worse than 2D images. \
 ]
@@ -68,6 +74,10 @@ This value for *memory_scope* can only be used with *atomic_work_item_fence* wit
 Note that the performance of 64-bit integer arithmetic can vary significantly between embedded devices. \
 ]
 
+:fn-local-arg-size-zero: pass:n[ \
+When the size of a `local` argument is set to zero, the value of the pointer within the kernel is implementation-defined. \
+]
+
 :fn-map-count-usage: pass:n[ \
 The map count returned should be considered immediately stale. \
 It is unsuitable for general use in applications. \

api/footnotes.asciidoc

@@ -364,7 +364,8 @@ In-order Execution ::
 
 Intermediate Language ::
     A lower-level language that may be used to create programs.
-    SPIR-V is a required intermediate language (IL) for OpenCL 2.1 and 2.2 devices.
+    SPIR-V is a required intermediate language (IL) for OpenCL 2.1, OpenCL 2.2,
+    and OpenCL 3.1 or newer devices.
     Other OpenCL devices may optionally support SPIR-V or other ILs.
 
 Kernel ::
@@ -376,7 +377,7 @@ Kernel ::
 Kernel-instance ::
     The work carried out by an OpenCL program occurs through the execution
     of kernel-instances on devices.
-    The kernel instance is the _kernel object_, the values associated with
+    The kernel-instance is the _kernel object_, the values associated with
     the arguments to the kernel, and the parameters that define the
     _ND-range_ index space.
 
@@ -493,9 +494,9 @@ Pipe ::
     items.
     A pipe has two endpoints: a write endpoint into which data items are
     inserted, and a read endpoint from which data items are removed.
-    At any one time, only one kernel instance may write into a pipe, and
-    only one kernel instance may read from a pipe.
-    To support the producer consumer design pattern, one kernel instance
+    At any one time, only one kernel-instance may write into a pipe, and
+    only one kernel-instance may read from a pipe.
+    To support the producer consumer design pattern, one kernel-instance
     connects to the write endpoint (the producer) while another kernel
     instance connects to the reading endpoint (the consumer).
 
@@ -636,16 +637,12 @@ Sampler ::
     image coordinate is a normalized or unnormalized value.
 
 Scope inclusion ::
-    Two actions *A* and *B* are defined to have an inclusive scope if they
-    have the same scope *P* such that: (1) if *P* is
-    *memory_scope_sub_group*, and *A* and *B* are executed by work-items
-    within the same sub-group, or (2) if *P* is *memory_scope_work_group*,
-    and *A* and *B* are executed by work-items within the same work-group,
-    or (3) if *P* is *memory_scope_device*, and *A* and *B* are executed by
-    work-items on the same device, or (4) if *P* is
-    *memory_scope_all_svm_devices* or *memory_scope_all_devices*, if *A* and *B*
-    are executed by host threads or by work-items on one or more devices that
-    can share SVM memory with each other and the host process.
+    Two actions *A* and *B* have an inclusive scope if the work-item or host
+    thread performing *A* is in the memory scope of *B* and the work-item or
+    host thread performing *B* is in the memory scope of *A*.
+    Prior to OpenCL 3.1, the memory scope of *A* and *B* additionally must be
+    the same for the actions to have an inclusive scope.
+    For OpenCL 3.1 and newer, the memory scope of *A* and *B* may be different.
 
 Sequenced before ::
     A relation between evaluations executed by a single unit of execution.