clang/docs/OpenCLSupport.md - llvm-project - Git at Google


            
            ```{raw} html

            
            <style type="text/css">

            
              .none { background-color: #FFCCCC }

            
              .part { background-color: #FFFF99 }

            
              .good { background-color: #CCFF99 }

            
            </style>

            
            ```


            # OpenCL Support


            Clang has complete support of OpenCL C versions from 1.0 to 3.0.

            
            Support for OpenCL 3.0 is in experimental phase ({ref}`OpenCL 3.0 <opencl_300>`).


            Clang also supports {ref}`the C++ for OpenCL kernel language <cxx_for_opencl_impl>`.


            There are also other {ref}`new and experimental features <opencl_experimenal>`

            
            available.


            Details about usage of clang for OpenCL can be found in {doc}`UsersManual`.


            ## Missing features or with limited support


            - For general issues and bugs with OpenCL in clang refer to [the GitHub issue

            
              list](https://github.com/llvm/llvm-project/issues?q=is%3Aopen+is%3Aissue+label%3Aopencl).

            
            - Command-line flag {option}`-cl-ext` (used to override extensions/

            
              features supported by a target) is missing support of some functionality i.e. that is

            
              implemented fully through libraries (see {ref}`library-based features and

            
              extensions <opencl_ext_libs>`).


            ## Internals Manual


            This section acts as internal documentation for OpenCL features design

            
            as well as some important implementation aspects. It is primarily targeted

            
            at the advanced users and the toolchain developers integrating frontend

            
            functionality as a component.


            ### OpenCL Metadata


            Clang uses metadata to provide additional OpenCL semantics in IR needed for

            
            backends and OpenCL runtime.


            Each kernel will have function metadata attached to it, specifying the arguments.

            
            Kernel argument metadata is used to provide source level information for querying

            
            at runtime, for example using the [clGetKernelArgInfo](https://www.khronos.org/registry/OpenCL/specs/opencl-1.2.pdf#167)

            
            call.


            Note that `-cl-kernel-arg-info` enables more information about the original

            
            kernel code to be added e.g. kernel parameter names will appear in the OpenCL

            
            metadata along with other information.


            The IDs used to encode the OpenCL's logical address spaces in the argument info

            
            metadata follows the SPIR address space mapping as defined in the SPIR

            
            specification [section 2.2](https://www.khronos.org/registry/spir/specs/spir_spec-2.0.pdf#18)


            ### OpenCL Specific Options


            In addition to the options described in {doc}`UsersManual` there are the

            
            following options specific to the OpenCL frontend.


            All the options in this section are frontend-only and therefore if used

            
            with regular clang driver they require frontend forwarding, e.g. `-cc1`

            
            or `-Xclang`.


            (opencl_finclude_default_header)=

            
            (opencl-finclude-default-header)=


            ```{eval-rst}

            
            .. option:: -finclude-default-header

            
            ```


            Adds most of builtin types and function declarations during compilations. By

            
            default the OpenCL headers are not loaded by the frontend and therefore certain

            
            builtin types and most of builtin functions are not declared. To load them

            
            automatically this flag can be passed to the frontend (see also {ref}`the

            
            section on the OpenCL Header <opencl_header>`):


            > ```console

            
            > $ clang -Xclang -finclude-default-header test.cl

            
            > ```


            Alternatively the internal header `opencl-c.h` containing the declarations

            
            can be included manually using `-include` or `-I` followed by the path

            
            to the header location. The header can be found in the clang source tree or

            
            installation directory.


            > ```console

            
            > $ clang -I<path to clang sources>/lib/Headers/opencl-c.h test.cl

            
            > $ clang -I<path to clang installation>/lib/clang/<llvm version>/include/opencl-c.h/opencl-c.h test.cl

            
            > ```


            In this example it is assumed that the kernel code contains

            
            `#include <opencl-c.h>` just as a regular C include.


            Because the header is very large and long to parse, PCH ({doc}`PCHInternals`)

            
            and modules ({doc}`Modules`) can be used internally to improve the compilation

            
            speed.


            To enable modules for OpenCL:


            > ```console

            
            > $ clang --target=spir-unknown-unknown -c -emit-llvm -Xclang -finclude-default-header -fmodules -fimplicit-module-maps -fmodules-cache-path=<path to the generated module> test.cl

            
            > ```


            Another way to circumvent long parsing latency for the OpenCL builtin

            
            declarations is to use mechanism enabled by {ref}`-fdeclare-opencl-builtins

            
            <opencl_fdeclare_opencl_builtins>` flag that is available as an alternative

            
            feature.


            (opencl_fdeclare_opencl_builtins)=

            
            (opencl-fdeclare-opencl-builtins)=


            ```{eval-rst}

            
            .. option:: -fdeclare-opencl-builtins

            
            ```


            In addition to regular header includes with builtin types and functions using

            
            {ref}`-finclude-default-header <opencl_finclude_default_header>`, clang

            
            supports a fast mechanism to declare builtin functions with

            
            `-fdeclare-opencl-builtins`. This does not declare the builtin types and

            
            therefore it has to be used in combination with `-finclude-default-header`

            
            if full functionality is required.


            **Example of Use**:


            > ```console

            
            > $ clang -Xclang -fdeclare-opencl-builtins test.cl

            
            > ```


            (opencl_fake_address_space_map)=

            
            (opencl-fake-address-space-map)=


            ```{eval-rst}

            
            .. option:: -ffake-address-space-map

            
            ```


            Overrides the target address space map with a fake map.

            
            This allows adding explicit address space IDs to the bitcode for non-segmented

            
            memory architectures that do not have separate IDs for each of the OpenCL

            
            logical address spaces by default. Passing `-ffake-address-space-map` will

            
            add/override address spaces of the target compiled for with the following values:

            
            `1-global`, `2-constant`, `3-local`, `4-generic`. The private address

            
            space is represented by the absence of an address space attribute in the IR (see

            
            also {ref}`the section on the address space attribute <opencl_addrsp>`).


            > ```console

            
            > $ clang -cc1 -ffake-address-space-map test.cl

            
            > ```


            (opencl_builtins)=

            
            (opencl-builtins)=


            ### OpenCL builtins


            **Clang builtins**


            There are some standard OpenCL functions that are implemented as Clang builtins:


            - All pipe functions from [section 6.13.16.2/6.13.16.3](https://www.khronos.org/registry/cl/specs/opencl-2.0-openclc.pdf#160) of

            
              the OpenCL v2.0 kernel language specification.

            
            - Address space qualifier conversion functions `to_global`/`to_local`/`to_private`

            
              from [section 6.13.9](https://www.khronos.org/registry/cl/specs/opencl-2.0-openclc.pdf#101).

            
            - All the `enqueue_kernel` functions from [section 6.13.17.1](https://www.khronos.org/registry/cl/specs/opencl-2.0-openclc.pdf#164) and

            
              enqueue query functions from [section 6.13.17.5](https://www.khronos.org/registry/cl/specs/opencl-2.0-openclc.pdf#171).


            **Fast builtin function declarations**


            The implementation of the fast builtin function declarations (available via the

            
            {ref}`-fdeclare-opencl-builtins option <opencl_fdeclare_opencl_builtins>`) consists

            
            of the following main components:


            - A TableGen definitions file `OpenCLBuiltins.td`. This contains a compact

            
              representation of the supported builtin functions. When adding new builtin

            
              function declarations, this is normally the only file that needs modifying.

            
            - A Clang TableGen emitter defined in `ClangOpenCLBuiltinEmitter.cpp`. During

            
              Clang build time, the emitter reads the TableGen definition file and

            
              generates `OpenCLBuiltins.inc`. This generated file contains various tables

            
              and functions that capture the builtin function data from the TableGen

            
              definitions in a compact manner.

            
            - OpenCL specific code in `SemaLookup.cpp`. When `Sema::LookupBuiltin`

            
              encounters a potential builtin function, it will check if the name corresponds

            
              to a valid OpenCL builtin function. If so, all overloads of the function are

            
              inserted using `InsertOCLBuiltinDeclarationsFromTable` and overload

            
              resolution takes place.


            ### OpenCL Extensions and Features


            Clang implements various extensions to OpenCL kernel languages.


            New functionality is accepted as soon as the documentation is detailed to the

            
            level sufficient to be implemented. There should be an evidence that the

            
            extension is designed with implementation feasibility in consideration and

            
            assessment of complexity for C/C++ based compilers. Alternatively, the

            
            documentation can be accepted in a format of a draft that can be further

            
            refined during the implementation.


            #### Implementation guidelines


            This section explains how to extend clang with the new functionality.


            **Parsing functionality**


            If a new extension is added it needs to be added to the clang frontend source

            
            code. This also means that the associated macro indicating the presence of the

            
            extension should be added to clang.


            The default flow for adding a new extension into the frontend is to

            
            modify [OpenCLExtensions.def](https://github.com/llvm/llvm-project/blob/main/clang/include/clang/Basic/OpenCLExtensions.def),

            
            containing the list of all extensions and optional features supported by

            
            the frontend.


            This will add the macro automatically and also add a field in the target

            
            options `clang::TargetOptions::OpenCLFeaturesMap` to control the exposure

            
            of the new extension during the compilation.


            Note that by default targets like `SPIR-V`, `SPIR` or `X86` expose all the OpenCL

            
            extensions. For all other targets the configuration has to be made explicitly.


            Note that the target extension support performed by clang can be overridden

            
            with {option}`-cl-ext` command-line flags.


            (opencl_ext_libs)=

            
            (opencl-ext-libs)=


            **Library functionality**


            If an extension adds functionality that does not modify standard language

            
            parsing it may not require modifying anything other than header files and

            
            `OpenCLBuiltins.td` detailed in {ref}`OpenCL builtins <opencl_builtins>`.

            
            Most commonly such extensions add functionality via libraries (by adding

            
            non-native types or functions) parsed regularly. Similar to other languages this

            
            is the most common way to add new functionality.


            Clang has standard headers where new types and functions are being added,

            
            for more details refer to

            
            {ref}`the section on the OpenCL Header <opencl_header>`.

            
            **Pragmas**


            Some extensions alter standard parsing dynamically via pragmas.


            Clang provides a mechanism to add the standard extension pragma

            
            `OPENCL EXTENSION` by setting a dedicated flag in the extension list entry of

            
            `OpenCLExtensions.def`. Note that there is no default behavior for the

            
            standard extension pragmas as it is not specified (for the standards up to and

            
            including version 3.0) in a sufficient level of detail and, therefore,

            
            there is no default functionality provided by clang.


            Pragmas without detailed information of their behavior (e.g. an explanation of

            
            changes it triggers in the parsing) should not be added to clang. Moreover, the

            
            pragmas should provide useful functionality to the user. For example, such

            
            functionality should address a practical use case and not be redundant i.e.

            
            cannot be achieved using existing features.


            Note that some legacy extensions (published prior to OpenCL 3.0) still

            
            provide some non-conformant functionality for pragmas e.g. add diagnostics on

            
            the use of types or functions. This functionality is not guaranteed to remain in

            
            future releases. However, any future changes should not affect backward

            
            compatibility.


            (opencl_addrsp)=

            
            (opencl-addrsp)=


            ### Address spaces attribute


            Clang has arbitrary address space support using the `address_space(N)`

            
            attribute, where `N` is an integer number in the range specified in the

            
            Clang source code. This addresses spaces can be used along with the OpenCL

            
            address spaces however when such addresses spaces converted to/from OpenCL

            
            address spaces the behavior is not governed by OpenCL specification.


            An OpenCL implementation provides a list of standard address spaces using

            
            keywords: `private`, `local`, `global`, and `generic`. In the AST and

            
            in the IR each of the address spaces will be represented by unique number

            
            provided in the Clang source code. The specific IDs for an address space do not

            
            have to match between the AST and the IR. Typically in the AST address space

            
            numbers represent logical segments while in the IR they represent physical

            
            segments.

            
            Therefore, machines with flat memory segments can map all AST address space

            
            numbers to the same physical segment ID or skip address space attribute

            
            completely while generating the IR. However, if the address space information

            
            is needed by the IR passes e.g. to improve alias analysis, it is recommended

            
            to keep it and only lower to reflect physical memory segments in the late

            
            machine passes. The mapping between logical and target address spaces is

            
            specified in the Clang's source code.


            (cxx_for_opencl_impl)=

            
            (cxx-for-opencl-impl)=


            ## C++ for OpenCL Implementation Status


            Clang implements language versions 1.0 and 2021 published in [the official

            
            release of C++ for OpenCL Documentation](https://github.com/KhronosGroup/OpenCL-Docs/releases/tag/cxxforopencl-docrev2021.12).


            Limited support of experimental C++ libraries is described in the {ref}`experimental features <opencl_experimenal>`.


            GitHub issues for this functionality are typically prefixed

            
            with '[C++4OpenCL]' - click [here](https://github.com/llvm/llvm-project/issues?q=is%3Aissue+is%3Aopen+%5BC%2B%2B4OpenCL%5D)

            
            to view the full bug list.


            ### Missing features or with limited support


            - Support of C++ for OpenCL 2021 is currently in experimental phase. Refer to

            
              {ref}`OpenCL 3.0 status <opencl_300>` for details of common missing

            
              functionality from OpenCL 3.0.

            
            - IR generation for non-trivial global destructors is incomplete (See:

            
              [PR48047](https://llvm.org/PR48047)).

            
            - Support of [destructors with non-default address spaces](https://www.khronos.org/opencl/assets/CXX_for_OpenCL.html#_construction_initialization_and_destruction)

            
              is incomplete (See: [D109609](https://reviews.llvm.org/D109609)).


            (opencl_300)=

            
            (opencl-300)=


            ## OpenCL C 3.0 Usage


            OpenCL C 3.0 language standard makes most OpenCL C 2.0 features optional. Optional

            
            functionality in OpenCL C 3.0 is indicated with the presence of feature-test macros

            
            (list of feature-test macros is [here](https://www.khronos.org/registry/OpenCL/specs/3.0-unified/html/OpenCL_C.html#features)).

            
            Command-line flag {option}`-cl-ext` can be used to override features supported by a target.


            For cases when there is an associated extension for a specific feature (fp64 and 3d image writes)

            
            user should specify both (extension and feature) in command-line flag:


            > ```console

            
            > $ clang -cl-std=CL3.0 -cl-ext=+cl_khr_fp64,+__opencl_c_fp64 ...

            
            > $ clang -cl-std=CL3.0 -cl-ext=-cl_khr_fp64,-__opencl_c_fp64 ...

            
            > ```


            ### OpenCL C 3.0 Implementation Status


            The following table provides an overview of features in OpenCL C 3.0 and their

            
            implementation status.


            | Category               | Feature                                                           | Status                         | Reviews                                                                              |

            
            | ---------------------- | ----------------------------------------------------------------- | ------------------------------ | ------------------------------------------------------------------------------------ |

            
            | Command line interface | New value for `-cl-std` flag                                      | <span class="good">done</span> | [D88300][D88300]                                                                     |

            
            | Predefined macros      | New version macro                                                 | <span class="good">done</span> | [D88300][D88300]                                                                     |

            
            | Predefined macros      | Feature macros                                                    | <span class="good">done</span> | [D95776][D95776]                                                                     |

            
            | Feature optionality    | Generic address space                                             | <span class="good">done</span> | [D95778][D95778] and [D103401][D103401]                                              |

            
            | Feature optionality    | Builtin function overloads with generic address space             | <span class="good">done</span> | [D105526][D105526], [D107769][D107769]                                               |

            
            | Feature optionality    | Program scope variables in global memory                          | <span class="good">done</span> | [D103191][D103191]                                                                   |

            
            | Feature optionality    | 3D image writes including builtin functions                       | <span class="good">done</span> | [D106260][D106260] (frontend)                                                        |

            
            | Feature optionality    | read_write images including builtin functions                     | <span class="good">done</span> | [D104915][D104915] (frontend) and [D107539][D107539], [D117899][D117899] (functions) |

            
            | Feature optionality    | C11 atomics memory scopes, ordering and builtin function          | <span class="good">done</span> | [D106111][D106111], [D119420][D119420]                                               |

            
            | Feature optionality    | Blocks and Device-side kernel enqueue including builtin functions | <span class="good">done</span> | [D115640][D115640], [D118605][D118605]                                               |

            
            | Feature optionality    | Pipes including builtin functions                                 | <span class="good">done</span> | [D107154][D107154] (frontend) and [D105858][D105858] (functions)                     |

            
            | Feature optionality    | Work group collective builtin functions                           | <span class="good">done</span> | [D105858][D105858]                                                                   |

            
            | Feature optionality    | Image types and builtin functions                                 | <span class="good">done</span> | [D103911][D103911] (frontend) and [D107539][D107539] (functions)                     |

            
            | Feature optionality    | Double precision floating point type                              | <span class="good">done</span> | [D96524][D96524]                                                                     |

            
            | New functionality      | RGBA vector components                                            | <span class="good">done</span> | [D99969][D99969]                                                                     |

            
            | New functionality      | Subgroup functions                                                | <span class="good">done</span> | [D105858][D105858], [D118999][D118999]                                               |

            
            | New functionality      | Atomic mem scopes: subgroup, all devices including functions      | <span class="good">done</span> | [D103241][D103241]                                                                   |


            [D88300]: https://reviews.llvm.org/D88300

            
            [D95776]: https://reviews.llvm.org/D95776

            
            [D95778]: https://reviews.llvm.org/D95778

            
            [D103401]: https://reviews.llvm.org/D103401

            
            [D105526]: https://reviews.llvm.org/D105526

            
            [D107769]: https://reviews.llvm.org/D107769

            
            [D103191]: https://reviews.llvm.org/D103191

            
            [D106260]: https://reviews.llvm.org/D106260

            
            [D104915]: https://reviews.llvm.org/D104915

            
            [D107539]: https://reviews.llvm.org/D107539

            
            [D117899]: https://reviews.llvm.org/D117899

            
            [D106111]: https://reviews.llvm.org/D106111

            
            [D119420]: https://reviews.llvm.org/D119420

            
            [D115640]: https://reviews.llvm.org/D115640

            
            [D118605]: https://reviews.llvm.org/D118605

            
            [D107154]: https://reviews.llvm.org/D107154

            
            [D105858]: https://reviews.llvm.org/D105858

            
            [D103911]: https://reviews.llvm.org/D103911

            
            [D96524]: https://reviews.llvm.org/D96524

            
            [D99969]: https://reviews.llvm.org/D99969

            
            [D118999]: https://reviews.llvm.org/D118999

            
            [D103241]: https://reviews.llvm.org/D103241


            (opencl_experimenal)=

            
            (opencl-experimenal)=


            ## Experimental features


            Clang provides the following new WIP features for the developers to experiment

            
            and provide early feedback or contribute with further improvements.

            
            Feel free to contact us on [the Discourse forums (Clang Frontend category)](https://discourse.llvm.org/c/clang/6) or file [a GitHub issue](https://github.com/llvm/llvm-project/issues/new).


            (opencl-experimental-cxxlibs)=


            ### C++ libraries for OpenCL


            There is ongoing work to support C++ standard libraries from [LLVM's libcxx](https://libcxx.llvm.org/) in OpenCL kernel code using C++ for OpenCL mode.


            It is currently possible to include `type_traits` from C++17 in the kernel

            
            sources when the following clang extensions are enabled

            
            `__cl_clang_function_pointers` and `__cl_clang_variadic_functions`,

            
            see {doc}`LanguageExtensions` for more details. The use of non-conformant

            
            features enabled by the extensions does not expose non-conformant behavior

            
            beyond the compilation i.e. does not get generated in IR or binary.

            
            The extension only appear in metaprogramming

            
            mechanism to identify or verify the properties of types. This allows to provide

            
            the full C++ functionality without a loss of portability. To avoid unsafe use

            
            of the extensions it is recommended that the extensions are disabled directly

            
            after the header include.


            **Example of Use**:


            The example of kernel code with `type_traits` is illustrated here.


            ```c++

            
            #pragma OPENCL EXTENSION __cl_clang_function_pointers : enable

            
            #pragma OPENCL EXTENSION __cl_clang_variadic_functions : enable

            
            #include <type_traits>

            
            #pragma OPENCL EXTENSION __cl_clang_function_pointers : disable

            
            #pragma OPENCL EXTENSION __cl_clang_variadic_functions : disable


            using sint_type = std::make_signed<unsigned int>::type;


            __kernel void foo() {

            
              static_assert(!std::is_same<sint_type, unsigned int>::value);

            
            }

            
            ```


            The possible clang invocation to compile the example is as follows:


            > ```console

            
            > $ clang -I<path to libcxx checkout or installation>/include test.clcpp

            
            > ```


            Note that `type_traits` is a header only library and therefore no extra

            
            linking step against the standard libraries is required. See full example

            
            in [Compiler Explorer](https://godbolt.org/z/5WbnTfb65).


            More OpenCL specific C++ library implementations built on top of libcxx

            
            are available in [libclcxx](https://github.com/KhronosGroup/libclcxx)

            
            project.
	```{raw} html
	<style type="text/css">
	.none { background-color: #FFCCCC }
	.part { background-color: #FFFF99 }
	.good { background-color: #CCFF99 }
	</style>
	```


	# OpenCL Support

	Clang has complete support of OpenCL C versions from 1.0 to 3.0.
	Support for OpenCL 3.0 is in experimental phase ({ref}`OpenCL 3.0 <opencl_300>`).

	Clang also supports {ref}`the C++ for OpenCL kernel language <cxx_for_opencl_impl>`.

	There are also other {ref}`new and experimental features <opencl_experimenal>`
	available.

	Details about usage of clang for OpenCL can be found in {doc}`UsersManual`.

	## Missing features or with limited support

	- For general issues and bugs with OpenCL in clang refer to [the GitHub issue
	list](https://github.com/llvm/llvm-project/issues?q=is%3Aopen+is%3Aissue+label%3Aopencl).
	- Command-line flag {option}`-cl-ext` (used to override extensions/
	features supported by a target) is missing support of some functionality i.e. that is
	implemented fully through libraries (see {ref}`library-based features and
	extensions <opencl_ext_libs>`).

	## Internals Manual

	This section acts as internal documentation for OpenCL features design
	as well as some important implementation aspects. It is primarily targeted
	at the advanced users and the toolchain developers integrating frontend
	functionality as a component.

	### OpenCL Metadata

	Clang uses metadata to provide additional OpenCL semantics in IR needed for
	backends and OpenCL runtime.

	Each kernel will have function metadata attached to it, specifying the arguments.
	Kernel argument metadata is used to provide source level information for querying
	at runtime, for example using the [clGetKernelArgInfo](https://www.khronos.org/registry/OpenCL/specs/opencl-1.2.pdf#167)
	call.

	Note that `-cl-kernel-arg-info` enables more information about the original
	kernel code to be added e.g. kernel parameter names will appear in the OpenCL
	metadata along with other information.

	The IDs used to encode the OpenCL's logical address spaces in the argument info
	metadata follows the SPIR address space mapping as defined in the SPIR
	specification [section 2.2](https://www.khronos.org/registry/spir/specs/spir_spec-2.0.pdf#18)

	### OpenCL Specific Options

	In addition to the options described in {doc}`UsersManual` there are the
	following options specific to the OpenCL frontend.

	All the options in this section are frontend-only and therefore if used
	with regular clang driver they require frontend forwarding, e.g. `-cc1`
	or `-Xclang`.

	(opencl_finclude_default_header)=
	(opencl-finclude-default-header)=

	```{eval-rst}
	.. option:: -finclude-default-header
	```

	Adds most of builtin types and function declarations during compilations. By
	default the OpenCL headers are not loaded by the frontend and therefore certain
	builtin types and most of builtin functions are not declared. To load them
	automatically this flag can be passed to the frontend (see also {ref}`the
	section on the OpenCL Header <opencl_header>`):

	> ```console
	> $ clang -Xclang -finclude-default-header test.cl
	> ```

	Alternatively the internal header `opencl-c.h` containing the declarations
	can be included manually using `-include` or `-I` followed by the path
	to the header location. The header can be found in the clang source tree or
	installation directory.

	> ```console
	> $ clang -I<path to clang sources>/lib/Headers/opencl-c.h test.cl
	> $ clang -I<path to clang installation>/lib/clang/<llvm version>/include/opencl-c.h/opencl-c.h test.cl
	> ```

	In this example it is assumed that the kernel code contains
	`#include <opencl-c.h>` just as a regular C include.

	Because the header is very large and long to parse, PCH ({doc}`PCHInternals`)
	and modules ({doc}`Modules`) can be used internally to improve the compilation
	speed.

	To enable modules for OpenCL:

	> ```console
	> $ clang --target=spir-unknown-unknown -c -emit-llvm -Xclang -finclude-default-header -fmodules -fimplicit-module-maps -fmodules-cache-path=<path to the generated module> test.cl
	> ```

	Another way to circumvent long parsing latency for the OpenCL builtin
	declarations is to use mechanism enabled by {ref}`-fdeclare-opencl-builtins
	<opencl_fdeclare_opencl_builtins>` flag that is available as an alternative
	feature.

	(opencl_fdeclare_opencl_builtins)=
	(opencl-fdeclare-opencl-builtins)=

	```{eval-rst}
	.. option:: -fdeclare-opencl-builtins
	```

	In addition to regular header includes with builtin types and functions using
	{ref}`-finclude-default-header <opencl_finclude_default_header>`, clang
	supports a fast mechanism to declare builtin functions with
	`-fdeclare-opencl-builtins`. This does not declare the builtin types and
	therefore it has to be used in combination with `-finclude-default-header`
	if full functionality is required.

	Example of Use:

	> ```console
	> $ clang -Xclang -fdeclare-opencl-builtins test.cl
	> ```

	(opencl_fake_address_space_map)=
	(opencl-fake-address-space-map)=

	```{eval-rst}
	.. option:: -ffake-address-space-map
	```

	Overrides the target address space map with a fake map.
	This allows adding explicit address space IDs to the bitcode for non-segmented
	memory architectures that do not have separate IDs for each of the OpenCL
	logical address spaces by default. Passing `-ffake-address-space-map` will
	add/override address spaces of the target compiled for with the following values:
	`1-global`, `2-constant`, `3-local`, `4-generic`. The private address
	space is represented by the absence of an address space attribute in the IR (see
	also {ref}`the section on the address space attribute <opencl_addrsp>`).

	> ```console
	> $ clang -cc1 -ffake-address-space-map test.cl
	> ```

	(opencl_builtins)=
	(opencl-builtins)=

	### OpenCL builtins

	Clang builtins

	There are some standard OpenCL functions that are implemented as Clang builtins:

	- All pipe functions from [section 6.13.16.2/6.13.16.3](https://www.khronos.org/registry/cl/specs/opencl-2.0-openclc.pdf#160) of
	the OpenCL v2.0 kernel language specification.
	- Address space qualifier conversion functions `to_global`/`to_local`/`to_private`
	from [section 6.13.9](https://www.khronos.org/registry/cl/specs/opencl-2.0-openclc.pdf#101).
	- All the `enqueue_kernel` functions from [section 6.13.17.1](https://www.khronos.org/registry/cl/specs/opencl-2.0-openclc.pdf#164) and
	enqueue query functions from [section 6.13.17.5](https://www.khronos.org/registry/cl/specs/opencl-2.0-openclc.pdf#171).

	Fast builtin function declarations

	The implementation of the fast builtin function declarations (available via the
	{ref}`-fdeclare-opencl-builtins option <opencl_fdeclare_opencl_builtins>`) consists
	of the following main components:

	- A TableGen definitions file `OpenCLBuiltins.td`. This contains a compact
	representation of the supported builtin functions. When adding new builtin
	function declarations, this is normally the only file that needs modifying.
	- A Clang TableGen emitter defined in `ClangOpenCLBuiltinEmitter.cpp`. During
	Clang build time, the emitter reads the TableGen definition file and
	generates `OpenCLBuiltins.inc`. This generated file contains various tables
	and functions that capture the builtin function data from the TableGen
	definitions in a compact manner.
	- OpenCL specific code in `SemaLookup.cpp`. When `Sema::LookupBuiltin`
	encounters a potential builtin function, it will check if the name corresponds
	to a valid OpenCL builtin function. If so, all overloads of the function are
	inserted using `InsertOCLBuiltinDeclarationsFromTable` and overload
	resolution takes place.

	### OpenCL Extensions and Features

	Clang implements various extensions to OpenCL kernel languages.

	New functionality is accepted as soon as the documentation is detailed to the
	level sufficient to be implemented. There should be an evidence that the
	extension is designed with implementation feasibility in consideration and
	assessment of complexity for C/C++ based compilers. Alternatively, the
	documentation can be accepted in a format of a draft that can be further
	refined during the implementation.

	#### Implementation guidelines

	This section explains how to extend clang with the new functionality.

	Parsing functionality

	If a new extension is added it needs to be added to the clang frontend source
	code. This also means that the associated macro indicating the presence of the
	extension should be added to clang.

	The default flow for adding a new extension into the frontend is to
	modify [OpenCLExtensions.def](https://github.com/llvm/llvm-project/blob/main/clang/include/clang/Basic/OpenCLExtensions.def),
	containing the list of all extensions and optional features supported by
	the frontend.

	This will add the macro automatically and also add a field in the target
	options `clang::TargetOptions::OpenCLFeaturesMap` to control the exposure
	of the new extension during the compilation.

	Note that by default targets like `SPIR-V`, `SPIR` or `X86` expose all the OpenCL
	extensions. For all other targets the configuration has to be made explicitly.

	Note that the target extension support performed by clang can be overridden
	with {option}`-cl-ext` command-line flags.

	(opencl_ext_libs)=
	(opencl-ext-libs)=

	Library functionality

	If an extension adds functionality that does not modify standard language
	parsing it may not require modifying anything other than header files and
	`OpenCLBuiltins.td` detailed in {ref}`OpenCL builtins <opencl_builtins>`.
	Most commonly such extensions add functionality via libraries (by adding
	non-native types or functions) parsed regularly. Similar to other languages this
	is the most common way to add new functionality.

	Clang has standard headers where new types and functions are being added,
	for more details refer to
	{ref}`the section on the OpenCL Header <opencl_header>`.
	Pragmas

	Some extensions alter standard parsing dynamically via pragmas.

	Clang provides a mechanism to add the standard extension pragma
	`OPENCL EXTENSION` by setting a dedicated flag in the extension list entry of
	`OpenCLExtensions.def`. Note that there is no default behavior for the
	standard extension pragmas as it is not specified (for the standards up to and
	including version 3.0) in a sufficient level of detail and, therefore,
	there is no default functionality provided by clang.

	Pragmas without detailed information of their behavior (e.g. an explanation of
	changes it triggers in the parsing) should not be added to clang. Moreover, the
	pragmas should provide useful functionality to the user. For example, such
	functionality should address a practical use case and not be redundant i.e.
	cannot be achieved using existing features.

	Note that some legacy extensions (published prior to OpenCL 3.0) still
	provide some non-conformant functionality for pragmas e.g. add diagnostics on
	the use of types or functions. This functionality is not guaranteed to remain in
	future releases. However, any future changes should not affect backward
	compatibility.

	(opencl_addrsp)=
	(opencl-addrsp)=

	### Address spaces attribute

	Clang has arbitrary address space support using the `address_space(N)`
	attribute, where `N` is an integer number in the range specified in the
	Clang source code. This addresses spaces can be used along with the OpenCL
	address spaces however when such addresses spaces converted to/from OpenCL
	address spaces the behavior is not governed by OpenCL specification.

	An OpenCL implementation provides a list of standard address spaces using
	keywords: `private`, `local`, `global`, and `generic`. In the AST and
	in the IR each of the address spaces will be represented by unique number
	provided in the Clang source code. The specific IDs for an address space do not
	have to match between the AST and the IR. Typically in the AST address space
	numbers represent logical segments while in the IR they represent physical
	segments.
	Therefore, machines with flat memory segments can map all AST address space
	numbers to the same physical segment ID or skip address space attribute
	completely while generating the IR. However, if the address space information
	is needed by the IR passes e.g. to improve alias analysis, it is recommended
	to keep it and only lower to reflect physical memory segments in the late
	machine passes. The mapping between logical and target address spaces is
	specified in the Clang's source code.

	(cxx_for_opencl_impl)=
	(cxx-for-opencl-impl)=

	## C++ for OpenCL Implementation Status

	Clang implements language versions 1.0 and 2021 published in [the official
	release of C++ for OpenCL Documentation](https://github.com/KhronosGroup/OpenCL-Docs/releases/tag/cxxforopencl-docrev2021.12).

	Limited support of experimental C++ libraries is described in the {ref}`experimental features <opencl_experimenal>`.

	GitHub issues for this functionality are typically prefixed
	with '[C++4OpenCL]' - click [here](https://github.com/llvm/llvm-project/issues?q=is%3Aissue+is%3Aopen+%5BC%2B%2B4OpenCL%5D)
	to view the full bug list.

	### Missing features or with limited support

	- Support of C++ for OpenCL 2021 is currently in experimental phase. Refer to
	{ref}`OpenCL 3.0 status <opencl_300>` for details of common missing
	functionality from OpenCL 3.0.
	- IR generation for non-trivial global destructors is incomplete (See:
	[PR48047](https://llvm.org/PR48047)).
	- Support of [destructors with non-default address spaces](https://www.khronos.org/opencl/assets/CXX_for_OpenCL.html#_construction_initialization_and_destruction)
	is incomplete (See: [D109609](https://reviews.llvm.org/D109609)).

	(opencl_300)=
	(opencl-300)=

	## OpenCL C 3.0 Usage

	OpenCL C 3.0 language standard makes most OpenCL C 2.0 features optional. Optional
	functionality in OpenCL C 3.0 is indicated with the presence of feature-test macros
	(list of feature-test macros is [here](https://www.khronos.org/registry/OpenCL/specs/3.0-unified/html/OpenCL_C.html#features)).
	Command-line flag {option}`-cl-ext` can be used to override features supported by a target.

	For cases when there is an associated extension for a specific feature (fp64 and 3d image writes)
	user should specify both (extension and feature) in command-line flag:

	> ```console
	> $ clang -cl-std=CL3.0 -cl-ext=+cl_khr_fp64,+__opencl_c_fp64 ...
	> $ clang -cl-std=CL3.0 -cl-ext=-cl_khr_fp64,-__opencl_c_fp64 ...
	> ```

	### OpenCL C 3.0 Implementation Status

	The following table provides an overview of features in OpenCL C 3.0 and their
	implementation status.

	\| Category \| Feature \| Status \| Reviews \|
	\| ---------------------- \| ----------------------------------------------------------------- \| ------------------------------ \| ------------------------------------------------------------------------------------ \|
	\| Command line interface \| New value for `-cl-std` flag \| <span class="good">done</span> \| [D88300][D88300] \|
	\| Predefined macros \| New version macro \| <span class="good">done</span> \| [D88300][D88300] \|
	\| Predefined macros \| Feature macros \| <span class="good">done</span> \| [D95776][D95776] \|
	\| Feature optionality \| Generic address space \| <span class="good">done</span> \| [D95778][D95778] and [D103401][D103401] \|
	\| Feature optionality \| Builtin function overloads with generic address space \| <span class="good">done</span> \| [D105526][D105526], [D107769][D107769] \|
	\| Feature optionality \| Program scope variables in global memory \| <span class="good">done</span> \| [D103191][D103191] \|
	\| Feature optionality \| 3D image writes including builtin functions \| <span class="good">done</span> \| [D106260][D106260] (frontend) \|
	\| Feature optionality \| read_write images including builtin functions \| <span class="good">done</span> \| [D104915][D104915] (frontend) and [D107539][D107539], [D117899][D117899] (functions) \|
	\| Feature optionality \| C11 atomics memory scopes, ordering and builtin function \| <span class="good">done</span> \| [D106111][D106111], [D119420][D119420] \|
	\| Feature optionality \| Blocks and Device-side kernel enqueue including builtin functions \| <span class="good">done</span> \| [D115640][D115640], [D118605][D118605] \|
	\| Feature optionality \| Pipes including builtin functions \| <span class="good">done</span> \| [D107154][D107154] (frontend) and [D105858][D105858] (functions) \|
	\| Feature optionality \| Work group collective builtin functions \| <span class="good">done</span> \| [D105858][D105858] \|
	\| Feature optionality \| Image types and builtin functions \| <span class="good">done</span> \| [D103911][D103911] (frontend) and [D107539][D107539] (functions) \|
	\| Feature optionality \| Double precision floating point type \| <span class="good">done</span> \| [D96524][D96524] \|
	\| New functionality \| RGBA vector components \| <span class="good">done</span> \| [D99969][D99969] \|
	\| New functionality \| Subgroup functions \| <span class="good">done</span> \| [D105858][D105858], [D118999][D118999] \|
	\| New functionality \| Atomic mem scopes: subgroup, all devices including functions \| <span class="good">done</span> \| [D103241][D103241] \|

	[D88300]: https://reviews.llvm.org/D88300
	[D95776]: https://reviews.llvm.org/D95776
	[D95778]: https://reviews.llvm.org/D95778
	[D103401]: https://reviews.llvm.org/D103401
	[D105526]: https://reviews.llvm.org/D105526
	[D107769]: https://reviews.llvm.org/D107769
	[D103191]: https://reviews.llvm.org/D103191
	[D106260]: https://reviews.llvm.org/D106260
	[D104915]: https://reviews.llvm.org/D104915
	[D107539]: https://reviews.llvm.org/D107539
	[D117899]: https://reviews.llvm.org/D117899
	[D106111]: https://reviews.llvm.org/D106111
	[D119420]: https://reviews.llvm.org/D119420
	[D115640]: https://reviews.llvm.org/D115640
	[D118605]: https://reviews.llvm.org/D118605
	[D107154]: https://reviews.llvm.org/D107154
	[D105858]: https://reviews.llvm.org/D105858
	[D103911]: https://reviews.llvm.org/D103911
	[D96524]: https://reviews.llvm.org/D96524
	[D99969]: https://reviews.llvm.org/D99969
	[D118999]: https://reviews.llvm.org/D118999
	[D103241]: https://reviews.llvm.org/D103241

	(opencl_experimenal)=
	(opencl-experimenal)=

	## Experimental features

	Clang provides the following new WIP features for the developers to experiment
	and provide early feedback or contribute with further improvements.
	Feel free to contact us on [the Discourse forums (Clang Frontend category)](https://discourse.llvm.org/c/clang/6) or file [a GitHub issue](https://github.com/llvm/llvm-project/issues/new).

	(opencl-experimental-cxxlibs)=

	### C++ libraries for OpenCL

	There is ongoing work to support C++ standard libraries from [LLVM's libcxx](https://libcxx.llvm.org/) in OpenCL kernel code using C++ for OpenCL mode.

	It is currently possible to include `type_traits` from C++17 in the kernel
	sources when the following clang extensions are enabled
	`__cl_clang_function_pointers` and `__cl_clang_variadic_functions`,
	see {doc}`LanguageExtensions` for more details. The use of non-conformant
	features enabled by the extensions does not expose non-conformant behavior
	beyond the compilation i.e. does not get generated in IR or binary.
	The extension only appear in metaprogramming
	mechanism to identify or verify the properties of types. This allows to provide
	the full C++ functionality without a loss of portability. To avoid unsafe use
	of the extensions it is recommended that the extensions are disabled directly
	after the header include.

	Example of Use:

	The example of kernel code with `type_traits` is illustrated here.

	```c++
	#pragma OPENCL EXTENSION __cl_clang_function_pointers : enable
	#pragma OPENCL EXTENSION __cl_clang_variadic_functions : enable
	#include <type_traits>
	#pragma OPENCL EXTENSION __cl_clang_function_pointers : disable
	#pragma OPENCL EXTENSION __cl_clang_variadic_functions : disable

	using sint_type = std::make_signed<unsigned int>::type;

	__kernel void foo() {
	static_assert(!std::is_same<sint_type, unsigned int>::value);
	}
	```

	The possible clang invocation to compile the example is as follows:

	> ```console
	> $ clang -I<path to libcxx checkout or installation>/include test.clcpp
	> ```

	Note that `type_traits` is a header only library and therefore no extra
	linking step against the standard libraries is required. See full example
	in [Compiler Explorer](https://godbolt.org/z/5WbnTfb65).

	More OpenCL specific C++ library implementations built on top of libcxx
	are available in [libclcxx](https://github.com/KhronosGroup/libclcxx)
	project.