clean up a few more descriptions of custom devices

api/embedded_profile.asciidoc

@@ -263,19 +263,21 @@ Device Queries>> table.
         The minimum value is 1 KB.
 | {CL_DEVICE_COMPILER_AVAILABLE}
   | {cl_bool_TYPE}
-      | Is {CL_FALSE} if the implementation does not have a compiler available
-        to compile the program source.
+      | Is {CL_TRUE} if a compiler is available to compile programs created from
+        source or IL, or {CL_FALSE} if the implementation does not have a
+        compiler available.
 
-        Is {CL_TRUE} if the compiler is available.
-        This can be {CL_FALSE} for the embedded platform profile only.
+        This can be {CL_FALSE} for devices supporting the embedded profile.
+
+        When an online compiler is not available, OpenCL programs may still be
+        created from binaries or built-in kernels supported by the device.
 | {CL_DEVICE_LINKER_AVAILABLE}
   | {cl_bool_TYPE}
-      | Is {CL_FALSE} if the implementation does not have a linker available.
-        Is {CL_TRUE} if the linker is available.
-
-        This can be {CL_FALSE} for the embedded platform profile only.
+      | Is {CL_TRUE} if a linker is available to link compiled programs, or
+        {CL_FALSE} if the implementation does not have a linker available.
 
-        This must be {CL_TRUE} if {CL_DEVICE_COMPILER_AVAILABLE} is {CL_TRUE}.
+        This can be {CL_FALSE} for devices supporting the embedded profile, but
+        must be {CL_TRUE} if {CL_DEVICE_COMPILER_AVAILABLE} is {CL_TRUE}.
 | {CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE}
   | {cl_uint_TYPE}
       | The max. size of the device queue in bytes.

api/glossary.asciidoc

@@ -83,13 +83,11 @@ Buffer Object ::
 --
 
 Built-in Kernel ::
-    A _built-in kernel_ is a _kernel_ that is executed on an OpenCL _device_
-    or _custom device_ by fixed-function hardware or in firmware.
-    _Applications_ can query the _built-in kernels_ supported by a _device_
-    or _custom device_.
-    A _program object_ can only contain _kernels_ written in OpenCL C or
-    _built-in kernels_ but not both.
-    See also _Kernel_ and _Program_.
+    A _built-in kernel_ is a _kernel_ that is provided by an OpenCL
+    implementation.
+    A _built-in kernel_ is enqueued for execution like other _kernels_, but may
+    execute on specialized hardware that is unavailable to non-built-in kernels.
+    _Applications_ can query the _built-in kernels_ supported by a _device_.
 
 Child kernel ::
     See _Device-side enqueue_.
@@ -171,17 +169,11 @@ Converged control flow ::
     See _Control flow_.
 
 Custom Device ::
-    An OpenCL _device_ that fully implements the OpenCL Runtime but does not
-    support _programs_ written in OpenCL C.
-    A custom device may be specialized non-programmable hardware that is
-    very power efficient and performant for directed tasks or hardware with
-    limited programmable capabilities such as specialized DSPs.
-    Custom devices are not OpenCL conformant.
-    Custom devices may support an online compiler.
-    Programs for custom devices can be created using the OpenCL runtime APIs
-    that allow OpenCL programs to be created from source (if an online
-    compiler is supported) and/or binary, or from _built-in kernels_
-    supported by the _device_.
+    A _custom device_ is a specialized _device_ that supports a subset of the
+    OpenCL runtime APIs for directed tasks but is not OpenCL conformant.
+    A _custom device_ must implement all of the OpenCL runtime APIs, but may not
+    support all of the required minimum device capabilities and may return
+    implementation-defined error codes for unsupported functionality.
     See also _Device_.
 
 Data Parallel Programming Model ::

api/opencl_architecture.asciidoc

@@ -87,11 +87,8 @@ A device may expose special purpose functionality as a _built-in kernel_.
 The platform provides APIs for enumerating and invoking the built-in
 kernels offered by a device, but otherwise does not define their
 construction or semantics.
-A _custom device_ supports only built-in kernels, and cannot be programmed
-via a kernel language.
 
-NOTE: Built-in kernels and custom devices are <<unified-spec, missing before>>
-version 1.2.
+NOTE: Built-in kernels are <<unified-spec, missing before>> version 1.2.
 
 All device types support the OpenCL execution model, the OpenCL memory
 model, and the APIs used in OpenCL to manage devices.
@@ -685,10 +682,9 @@ The OpenCL execution model supports three types of kernels:
   * *Built-in kernels* are tied to particular device and are not built at
     runtime from source code in a program object.
     The common use of built in kernels is to expose fixed-function hardware
-    or firmware associated with a particular OpenCL device or custom device.
-    The semantics of a built-in kernel may be defined outside of OpenCL and
+    or firmware associated with a particular OpenCL device.
+    The semantics of a built-in kernel are defined outside of OpenCL and
     hence are implementation-defined.
-    Note: Built-in kernels are <<unified-spec, missing before>> version 1.2.
 
 
 All three types of kernels are manipulated through the OpenCL command-queues

api/opencl_platform_layer.asciidoc

@@ -446,8 +446,11 @@ include::{generated}/api/version-notes/CL_DEVICE_TYPE_ACCELERATOR.asciidoc[]
 | {CL_DEVICE_TYPE_CUSTOM_anchor}
 
 include::{generated}/api/version-notes/CL_DEVICE_TYPE_CUSTOM.asciidoc[]
-  | Specialized devices that implement some of the OpenCL runtime APIs but
-    do not support all of the required OpenCL functionality.
+  | Specialized devices that support a subset of the OpenCL runtime APIs for
+    directed tasks but are not OpenCL conformant.
+    A custom device must implement all of the OpenCL runtime APIs, but may not
+    support all of the required minimum device capabilities and may return
+    implementation-defined error codes for unsupported functionality.
 
 | {CL_DEVICE_TYPE_DEFAULT_anchor}
 
@@ -1138,7 +1141,7 @@ include::{generated}/api/version-notes/CL_DEVICE_PROFILING_TIMER_RESOLUTION.asci
 include::{generated}/api/version-notes/CL_DEVICE_ENDIAN_LITTLE.asciidoc[]
   | {cl_bool_TYPE}
       | Is {CL_TRUE} if the OpenCL device is a little endian device and
-        {CL_FALSE} otherwise
+        {CL_FALSE} otherwise.
 | {CL_DEVICE_AVAILABLE_anchor}
 
 include::{generated}/api/version-notes/CL_DEVICE_AVAILABLE.asciidoc[]
@@ -1150,21 +1153,23 @@ include::{generated}/api/version-notes/CL_DEVICE_AVAILABLE.asciidoc[]
 
 include::{generated}/api/version-notes/CL_DEVICE_COMPILER_AVAILABLE.asciidoc[]
   | {cl_bool_TYPE}
-      | Is {CL_FALSE} if the implementation does not have a compiler available
-        to compile the program source.
+      | Is {CL_TRUE} if a compiler is available to compile programs created from
+        source or IL, or {CL_FALSE} if the implementation does not have a
+        compiler available.
 
-        Is {CL_TRUE} if the compiler is available.
-        This can be {CL_FALSE} for the embedded platform profile only.
+        This must be {CL_TRUE} for devices supporting the full profile.
+
+        When an online compiler is not available, OpenCL programs may still be
+        created from binaries or built-in kernels supported by the device.
 | {CL_DEVICE_LINKER_AVAILABLE_anchor}
 
 include::{generated}/api/version-notes/CL_DEVICE_LINKER_AVAILABLE.asciidoc[]
   | {cl_bool_TYPE}
-      | Is {CL_FALSE} if the implementation does not have a linker available.
-        Is {CL_TRUE} if the linker is available.
-
-        This can be {CL_FALSE} for the embedded platform profile only.
+      | Is {CL_TRUE} if a linker is available to link compiled programs, or
+        {CL_FALSE} if the implementation does not have a linker available.
 
-        This must be {CL_TRUE} if {CL_DEVICE_COMPILER_AVAILABLE} is {CL_TRUE}.
+        This must be {CL_TRUE} for devices supporting the full profile or if
+        {CL_DEVICE_COMPILER_AVAILABLE} is {CL_TRUE}.
 | {CL_DEVICE_EXECUTION_CAPABILITIES_anchor}
 
 include::{generated}/api/version-notes/CL_DEVICE_EXECUTION_CAPABILITIES.asciidoc[]

api/opencl_runtime_layer.asciidoc

@@ -9145,20 +9145,12 @@ include::{generated}/api/version-notes/clCreateProgramWithSource.asciidoc[]
   * _errcode_ret_ will return an appropriate error code.
     If _errcode_ret_ is `NULL`, no error code is returned.
 
-The source code specified by _strings_ will be loaded into the program
-object.
-
 The devices associated with the program object are the devices associated
 with _context_.
-The source code specified by _strings_ is either an OpenCL C program source,
-header or implementation-defined source for custom devices that support an
-online compiler.
-OpenCL {cpp} is not supported as an online-compiled kernel language through
-this interface.
-ifdef::cl_ext_cxx_for_opencl[]
-If the {cl_ext_cxx_for_opencl_EXT} extension is supported, the source code
-specified by _strings_ may also be a C++ for OpenCL program source or header.
-endif::cl_ext_cxx_for_opencl[]
+
+The source code specified by _strings_ are loaded into the program object.
+The source languages supported for _strings_ are defined by device queries such
+as {CL_DEVICE_OPENCL_C_ALL_VERSIONS}.
 
 // refError
 
@@ -11501,22 +11493,19 @@ include::{generated}/api/version-notes/clGetKernelWorkGroupInfo.asciidoc[]
 
 include::{generated}/api/version-notes/CL_KERNEL_GLOBAL_WORK_SIZE.asciidoc[]
   | {size_t_TYPE}[3]
-      | This provides a mechanism for the application to query the maximum
-        global size that can be used to execute a kernel (i.e. the
-        _global_work_size_ argument to {clEnqueueNDRangeKernel}) on a custom
-        device given by _device_ or a built-in kernel on an OpenCL device
-        given by _device_.
-
-        If _device_ is not a custom device and _kernel_ is not a built-in
-        kernel, {clGetKernelWorkGroupInfo} returns the error
-        {CL_INVALID_VALUE}.
+      | Provides a mechanism for the application to query the maximum
+        global size that can be used to execute a built-in kernel on a
+        specific device given by _device_.
+
+        If _kernel_ is not a built-in kernel, querying
+        {CL_KERNEL_GLOBAL_WORK_SIZE) returns the error {CL_INVALID_VALUE}.
 | {CL_KERNEL_WORK_GROUP_SIZE_anchor}
 
 include::{generated}/api/version-notes/CL_KERNEL_WORK_GROUP_SIZE.asciidoc[]
   | {size_t_TYPE}
-      | This provides a mechanism for the application to query the maximum
+      | Provides a mechanism for the application to query the maximum
         work-group size that can be used to execute the kernel on a specific
-        device given by device.
+        device given by _device_.
         The OpenCL implementation uses the resource requirements of the
         kernel (register usage etc.) to determine what this work-group size
         should be.
@@ -11587,7 +11576,7 @@ Otherwise, it returns one of the following errors:
     <<kernel-workgroup-info-table, Kernel Object Device Queries>> table
     and _param_value_ is not `NULL`.
   * {CL_INVALID_VALUE} if _param_name_ is {CL_KERNEL_GLOBAL_WORK_SIZE} and
-    _device_ is not a custom device and _kernel_ is not a built-in kernel.
+    _kernel_ is not a built-in kernel.
   * {CL_OUT_OF_RESOURCES} if there is a failure to allocate resources required
     by the OpenCL implementation on the device.
   * {CL_OUT_OF_HOST_MEMORY} if there is a failure to allocate resources