460 lines · plain
1.. raw:: html2 3 <style type="text/css">4 .none { background-color: #FFCCCC }5 .part { background-color: #FFFF99 }6 .good { background-color: #CCFF99 }7 </style>8 9.. role:: none10.. role:: part11.. role:: good12 13.. contents::14 :local:15 16==================17OpenCL Support18==================19 20Clang has complete support of OpenCL C versions from 1.0 to 3.0.21Support for OpenCL 3.0 is in experimental phase (:ref:`OpenCL 3.0 <opencl_300>`).22 23Clang also supports :ref:`the C++ for OpenCL kernel language <cxx_for_opencl_impl>`.24 25There are also other :ref:`new and experimental features <opencl_experimenal>`26available.27 28Details about usage of clang for OpenCL can be found in :doc:`UsersManual`.29 30Missing features or with limited support31========================================32 33- For general issues and bugs with OpenCL in clang refer to `the GitHub issue34 list35 <https://github.com/llvm/llvm-project/issues?q=is%3Aopen+is%3Aissue+label%3Aopencl>`__.36 37- Command-line flag :option:`-cl-ext` (used to override extensions/38 features supported by a target) is missing support of some functionality i.e. that is39 implemented fully through libraries (see :ref:`library-based features and40 extensions <opencl_ext_libs>`).41 42Internals Manual43================44 45This section acts as internal documentation for OpenCL features design46as well as some important implementation aspects. It is primarily targeted47at the advanced users and the toolchain developers integrating frontend48functionality as a component.49 50OpenCL Metadata51---------------52 53Clang uses metadata to provide additional OpenCL semantics in IR needed for54backends and OpenCL runtime.55 56Each kernel will have function metadata attached to it, specifying the arguments.57Kernel argument metadata is used to provide source level information for querying58at runtime, for example using the `clGetKernelArgInfo59<https://www.khronos.org/registry/OpenCL/specs/opencl-1.2.pdf#167>`_60call.61 62Note that ``-cl-kernel-arg-info`` enables more information about the original63kernel code to be added e.g. kernel parameter names will appear in the OpenCL64metadata along with other information.65 66The IDs used to encode the OpenCL's logical address spaces in the argument info67metadata follows the SPIR address space mapping as defined in the SPIR68specification `section 2.269<https://www.khronos.org/registry/spir/specs/spir_spec-2.0.pdf#18>`_70 71OpenCL Specific Options72-----------------------73 74In addition to the options described in :doc:`UsersManual` there are the75following options specific to the OpenCL frontend.76 77All the options in this section are frontend-only and therefore if used78with regular clang driver they require frontend forwarding, e.g. ``-cc1``79or ``-Xclang``.80 81.. _opencl_finclude_default_header:82 83.. option:: -finclude-default-header84 85Adds most of builtin types and function declarations during compilations. By86default the OpenCL headers are not loaded by the frontend and therefore certain87builtin types and most of builtin functions are not declared. To load them88automatically this flag can be passed to the frontend (see also :ref:`the89section on the OpenCL Header <opencl_header>`):90 91 .. code-block:: console92 93 $ clang -Xclang -finclude-default-header test.cl94 95Alternatively the internal header `opencl-c.h` containing the declarations96can be included manually using ``-include`` or ``-I`` followed by the path97to the header location. The header can be found in the clang source tree or98installation directory.99 100 .. code-block:: console101 102 $ clang -I<path to clang sources>/lib/Headers/opencl-c.h test.cl103 $ clang -I<path to clang installation>/lib/clang/<llvm version>/include/opencl-c.h/opencl-c.h test.cl104 105In this example it is assumed that the kernel code contains106``#include <opencl-c.h>`` just as a regular C include.107 108Because the header is very large and long to parse, PCH (:doc:`PCHInternals`)109and modules (:doc:`Modules`) can be used internally to improve the compilation110speed.111 112To enable modules for OpenCL:113 114 .. code-block:: console115 116 $ 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.cl117 118Another way to circumvent long parsing latency for the OpenCL builtin119declarations is to use mechanism enabled by :ref:`-fdeclare-opencl-builtins120<opencl_fdeclare_opencl_builtins>` flag that is available as an alternative121feature.122 123.. _opencl_fdeclare_opencl_builtins:124 125.. option:: -fdeclare-opencl-builtins126 127In addition to regular header includes with builtin types and functions using128:ref:`-finclude-default-header <opencl_finclude_default_header>`, clang129supports a fast mechanism to declare builtin functions with130``-fdeclare-opencl-builtins``. This does not declare the builtin types and131therefore it has to be used in combination with ``-finclude-default-header``132if full functionality is required.133 134**Example of Use**:135 136 .. code-block:: console137 138 $ clang -Xclang -fdeclare-opencl-builtins test.cl139 140.. _opencl_fake_address_space_map:141 142.. option:: -ffake-address-space-map143 144Overrides the target address space map with a fake map.145This allows adding explicit address space IDs to the bitcode for non-segmented146memory architectures that do not have separate IDs for each of the OpenCL147logical address spaces by default. Passing ``-ffake-address-space-map`` will148add/override address spaces of the target compiled for with the following values:149``1-global``, ``2-constant``, ``3-local``, ``4-generic``. The private address150space is represented by the absence of an address space attribute in the IR (see151also :ref:`the section on the address space attribute <opencl_addrsp>`).152 153 .. code-block:: console154 155 $ clang -cc1 -ffake-address-space-map test.cl156 157.. _opencl_builtins:158 159OpenCL builtins160---------------161 162**Clang builtins**163 164There are some standard OpenCL functions that are implemented as Clang builtins:165 166- All pipe functions from `section 6.13.16.2/6.13.16.3167 <https://www.khronos.org/registry/cl/specs/opencl-2.0-openclc.pdf#160>`_ of168 the OpenCL v2.0 kernel language specification.169 170- Address space qualifier conversion functions ``to_global``/``to_local``/``to_private``171 from `section 6.13.9172 <https://www.khronos.org/registry/cl/specs/opencl-2.0-openclc.pdf#101>`_.173 174- All the ``enqueue_kernel`` functions from `section 6.13.17.1175 <https://www.khronos.org/registry/cl/specs/opencl-2.0-openclc.pdf#164>`_ and176 enqueue query functions from `section 6.13.17.5177 <https://www.khronos.org/registry/cl/specs/opencl-2.0-openclc.pdf#171>`_.178 179**Fast builtin function declarations**180 181The implementation of the fast builtin function declarations (available via the182:ref:`-fdeclare-opencl-builtins option <opencl_fdeclare_opencl_builtins>`) consists183of the following main components:184 185- A TableGen definitions file ``OpenCLBuiltins.td``. This contains a compact186 representation of the supported builtin functions. When adding new builtin187 function declarations, this is normally the only file that needs modifying.188 189- A Clang TableGen emitter defined in ``ClangOpenCLBuiltinEmitter.cpp``. During190 Clang build time, the emitter reads the TableGen definition file and191 generates ``OpenCLBuiltins.inc``. This generated file contains various tables192 and functions that capture the builtin function data from the TableGen193 definitions in a compact manner.194 195- OpenCL specific code in ``SemaLookup.cpp``. When ``Sema::LookupBuiltin``196 encounters a potential builtin function, it will check if the name corresponds197 to a valid OpenCL builtin function. If so, all overloads of the function are198 inserted using ``InsertOCLBuiltinDeclarationsFromTable`` and overload199 resolution takes place.200 201OpenCL Extensions and Features202------------------------------203 204Clang implements various extensions to OpenCL kernel languages.205 206New functionality is accepted as soon as the documentation is detailed to the207level sufficient to be implemented. There should be an evidence that the208extension is designed with implementation feasibility in consideration and209assessment of complexity for C/C++ based compilers. Alternatively, the210documentation can be accepted in a format of a draft that can be further211refined during the implementation.212 213Implementation guidelines214^^^^^^^^^^^^^^^^^^^^^^^^^215 216This section explains how to extend clang with the new functionality.217 218**Parsing functionality**219 220If a new extension is added it needs to be added to the clang frontend source221code. This also means that the associated macro indicating the presence of the222extension should be added to clang.223 224The default flow for adding a new extension into the frontend is to225modify `OpenCLExtensions.def226<https://github.com/llvm/llvm-project/blob/main/clang/include/clang/Basic/OpenCLExtensions.def>`__,227containing the list of all extensions and optional features supported by228the frontend.229 230This will add the macro automatically and also add a field in the target231options ``clang::TargetOptions::OpenCLFeaturesMap`` to control the exposure232of the new extension during the compilation.233 234Note that by default targets like `SPIR-V`, `SPIR` or `X86` expose all the OpenCL235extensions. For all other targets the configuration has to be made explicitly.236 237Note that the target extension support performed by clang can be overridden238with :option:`-cl-ext` command-line flags.239 240.. _opencl_ext_libs:241 242**Library functionality**243 244If an extension adds functionality that does not modify standard language245parsing it may not require modifying anything other than header files and246``OpenCLBuiltins.td`` detailed in :ref:`OpenCL builtins <opencl_builtins>`.247Most commonly such extensions add functionality via libraries (by adding248non-native types or functions) parsed regularly. Similar to other languages this249is the most common way to add new functionality.250 251Clang has standard headers where new types and functions are being added,252for more details refer to253:ref:`the section on the OpenCL Header <opencl_header>`.254**Pragmas**255 256Some extensions alter standard parsing dynamically via pragmas.257 258Clang provides a mechanism to add the standard extension pragma259``OPENCL EXTENSION`` by setting a dedicated flag in the extension list entry of260``OpenCLExtensions.def``. Note that there is no default behavior for the261standard extension pragmas as it is not specified (for the standards up to and262including version 3.0) in a sufficient level of detail and, therefore,263there is no default functionality provided by clang.264 265Pragmas without detailed information of their behavior (e.g. an explanation of266changes it triggers in the parsing) should not be added to clang. Moreover, the267pragmas should provide useful functionality to the user. For example, such268functionality should address a practical use case and not be redundant i.e.269cannot be achieved using existing features.270 271Note that some legacy extensions (published prior to OpenCL 3.0) still272provide some non-conformant functionality for pragmas e.g. add diagnostics on273the use of types or functions. This functionality is not guaranteed to remain in274future releases. However, any future changes should not affect backward275compatibility.276 277.. _opencl_addrsp:278 279Address spaces attribute280------------------------281 282Clang has arbitrary address space support using the ``address_space(N)``283attribute, where ``N`` is an integer number in the range specified in the284Clang source code. This addresses spaces can be used along with the OpenCL285address spaces however when such addresses spaces converted to/from OpenCL286address spaces the behavior is not governed by OpenCL specification.287 288An OpenCL implementation provides a list of standard address spaces using289keywords: ``private``, ``local``, ``global``, and ``generic``. In the AST and290in the IR each of the address spaces will be represented by unique number291provided in the Clang source code. The specific IDs for an address space do not292have to match between the AST and the IR. Typically in the AST address space293numbers represent logical segments while in the IR they represent physical294segments.295Therefore, machines with flat memory segments can map all AST address space296numbers to the same physical segment ID or skip address space attribute297completely while generating the IR. However, if the address space information298is needed by the IR passes e.g. to improve alias analysis, it is recommended299to keep it and only lower to reflect physical memory segments in the late300machine passes. The mapping between logical and target address spaces is301specified in the Clang's source code.302 303.. _cxx_for_opencl_impl:304 305C++ for OpenCL Implementation Status306====================================307 308Clang implements language versions 1.0 and 2021 published in `the official309release of C++ for OpenCL Documentation310<https://github.com/KhronosGroup/OpenCL-Docs/releases/tag/cxxforopencl-docrev2021.12>`_.311 312Limited support of experimental C++ libraries is described in the :ref:`experimental features <opencl_experimenal>`.313 314GitHub issues for this functionality are typically prefixed315with '[C++4OpenCL]' - click `here316<https://github.com/llvm/llvm-project/issues?q=is%3Aissue+is%3Aopen+%5BC%2B%2B4OpenCL%5D>`__317to view the full bug list.318 319 320Missing features or with limited support321----------------------------------------322 323- Support of C++ for OpenCL 2021 is currently in experimental phase. Refer to324 :ref:`OpenCL 3.0 status <opencl_300>` for details of common missing325 functionality from OpenCL 3.0.326 327- IR generation for non-trivial global destructors is incomplete (See:328 `PR48047 <https://llvm.org/PR48047>`_).329 330- Support of `destructors with non-default address spaces331 <https://www.khronos.org/opencl/assets/CXX_for_OpenCL.html#_construction_initialization_and_destruction>`_332 is incomplete (See: `D109609 <https://reviews.llvm.org/D109609>`_).333 334.. _opencl_300:335 336OpenCL C 3.0 Usage337==================338 339OpenCL C 3.0 language standard makes most OpenCL C 2.0 features optional. Optional340functionality in OpenCL C 3.0 is indicated with the presence of feature-test macros341(list of feature-test macros is `here <https://www.khronos.org/registry/OpenCL/specs/3.0-unified/html/OpenCL_C.html#features>`__).342Command-line flag :option:`-cl-ext` can be used to override features supported by a target.343 344For cases when there is an associated extension for a specific feature (fp64 and 3d image writes)345user should specify both (extension and feature) in command-line flag:346 347 .. code-block:: console348 349 $ clang -cl-std=CL3.0 -cl-ext=+cl_khr_fp64,+__opencl_c_fp64 ...350 $ clang -cl-std=CL3.0 -cl-ext=-cl_khr_fp64,-__opencl_c_fp64 ...351 352 353 354OpenCL C 3.0 Implementation Status355----------------------------------356 357The following table provides an overview of features in OpenCL C 3.0 and their358implementation status.359 360+------------------------------+-------------------------+-----------------------------------------+----------------------+--------------------------------------------------------------------------------------------------------------------------------+361| Category | Feature | Status | Reviews |362+==============================+=========================+=========================================+======================+================================================================================================================================+363| Command line interface | New value for ``-cl-std`` flag | :good:`done` | https://reviews.llvm.org/D88300 |364+------------------------------+-------------------------+-----------------------------------------+----------------------+--------------------------------------------------------------------------------------------------------------------------------+365| Predefined macros | New version macro | :good:`done` | https://reviews.llvm.org/D88300 |366+------------------------------+-------------------------+-----------------------------------------+----------------------+--------------------------------------------------------------------------------------------------------------------------------+367| Predefined macros | Feature macros | :good:`done` | https://reviews.llvm.org/D95776 |368+------------------------------+-------------------------+-----------------------------------------+----------------------+--------------------------------------------------------------------------------------------------------------------------------+369| Feature optionality | Generic address space | :good:`done` | https://reviews.llvm.org/D95778 and https://reviews.llvm.org/D103401 |370+------------------------------+-------------------------+-----------------------------------------+----------------------+--------------------------------------------------------------------------------------------------------------------------------+371| Feature optionality | Builtin function overloads with generic address space | :good:`done` | https://reviews.llvm.org/D105526, https://reviews.llvm.org/D107769 |372+------------------------------+-------------------------+-----------------------------------------+----------------------+--------------------------------------------------------------------------------------------------------------------------------+373| Feature optionality | Program scope variables in global memory | :good:`done` | https://reviews.llvm.org/D103191 |374+------------------------------+-------------------------+-----------------------------------------+----------------------+--------------------------------------------------------------------------------------------------------------------------------+375| Feature optionality | 3D image writes including builtin functions | :good:`done` | https://reviews.llvm.org/D106260 (frontend) |376+------------------------------+-------------------------+-----------------------------------------+----------------------+--------------------------------------------------------------------------------------------------------------------------------+377| Feature optionality | read_write images including builtin functions | :good:`done` | https://reviews.llvm.org/D104915 (frontend) and https://reviews.llvm.org/D107539, https://reviews.llvm.org/D117899 (functions) |378+------------------------------+-------------------------+-----------------------------------------+----------------------+--------------------------------------------------------------------------------------------------------------------------------+379| Feature optionality | C11 atomics memory scopes, ordering and builtin function | :good:`done` | https://reviews.llvm.org/D106111, https://reviews.llvm.org/D119420 |380+------------------------------+-------------------------+-----------------------------------------+----------------------+--------------------------------------------------------------------------------------------------------------------------------+381| Feature optionality | Blocks and Device-side kernel enqueue including builtin functions | :good:`done` | https://reviews.llvm.org/D115640, https://reviews.llvm.org/D118605 |382+------------------------------+-------------------------+-----------------------------------------+----------------------+--------------------------------------------------------------------------------------------------------------------------------+383| Feature optionality | Pipes including builtin functions | :good:`done` | https://reviews.llvm.org/D107154 (frontend) and https://reviews.llvm.org/D105858 (functions) |384+------------------------------+-------------------------+-----------------------------------------+----------------------+--------------------------------------------------------------------------------------------------------------------------------+385| Feature optionality | Work group collective builtin functions | :good:`done` | https://reviews.llvm.org/D105858 |386+------------------------------+-------------------------+-----------------------------------------+----------------------+--------------------------------------------------------------------------------------------------------------------------------+387| Feature optionality | Image types and builtin functions | :good:`done` | https://reviews.llvm.org/D103911 (frontend) and https://reviews.llvm.org/D107539 (functions) |388+------------------------------+-------------------------+-----------------------------------------+----------------------+--------------------------------------------------------------------------------------------------------------------------------+389| Feature optionality | Double precision floating point type | :good:`done` | https://reviews.llvm.org/D96524 |390+------------------------------+-------------------------+-----------------------------------------+----------------------+--------------------------------------------------------------------------------------------------------------------------------+391| New functionality | RGBA vector components | :good:`done` | https://reviews.llvm.org/D99969 |392+------------------------------+-------------------------+-----------------------------------------+----------------------+--------------------------------------------------------------------------------------------------------------------------------+393| New functionality | Subgroup functions | :good:`done` | https://reviews.llvm.org/D105858, https://reviews.llvm.org/D118999 |394+------------------------------+-------------------------+-----------------------------------------+----------------------+--------------------------------------------------------------------------------------------------------------------------------+395| New functionality | Atomic mem scopes: subgroup, all devices including functions | :good:`done` | https://reviews.llvm.org/D103241 |396+------------------------------+-------------------------+-----------------------------------------+----------------------+--------------------------------------------------------------------------------------------------------------------------------+397 398.. _opencl_experimenal:399 400Experimental features401=====================402 403Clang provides the following new WIP features for the developers to experiment404and provide early feedback or contribute with further improvements.405Feel free to contact us on `the Discourse forums (Clang Frontend category)406<https://discourse.llvm.org/c/clang/6>`_ or file `a GitHub issue407<https://github.com/llvm/llvm-project/issues/new>`_.408 409.. _opencl_experimental_cxxlibs:410 411C++ libraries for OpenCL412------------------------413 414There is ongoing work to support C++ standard libraries from `LLVM's libcxx415<https://libcxx.llvm.org/>`_ in OpenCL kernel code using C++ for OpenCL mode.416 417It is currently possible to include `type_traits` from C++17 in the kernel418sources when the following clang extensions are enabled419``__cl_clang_function_pointers`` and ``__cl_clang_variadic_functions``,420see :doc:`LanguageExtensions` for more details. The use of non-conformant421features enabled by the extensions does not expose non-conformant behavior422beyond the compilation i.e. does not get generated in IR or binary.423The extension only appear in metaprogramming424mechanism to identify or verify the properties of types. This allows to provide425the full C++ functionality without a loss of portability. To avoid unsafe use426of the extensions it is recommended that the extensions are disabled directly427after the header include.428 429**Example of Use**:430 431The example of kernel code with `type_traits` is illustrated here.432 433.. code-block:: c++434 435 #pragma OPENCL EXTENSION __cl_clang_function_pointers : enable436 #pragma OPENCL EXTENSION __cl_clang_variadic_functions : enable437 #include <type_traits>438 #pragma OPENCL EXTENSION __cl_clang_function_pointers : disable439 #pragma OPENCL EXTENSION __cl_clang_variadic_functions : disable440 441 using sint_type = std::make_signed<unsigned int>::type;442 443 __kernel void foo() {444 static_assert(!std::is_same<sint_type, unsigned int>::value);445 }446 447The possible clang invocation to compile the example is as follows:448 449 .. code-block:: console450 451 $ clang -I<path to libcxx checkout or installation>/include test.clcpp452 453Note that `type_traits` is a header only library and therefore no extra454linking step against the standard libraries is required. See full example455in `Compiler Explorer <https://godbolt.org/z/5WbnTfb65>`_.456 457More OpenCL specific C++ library implementations built on top of libcxx458are available in `libclcxx <https://github.com/KhronosGroup/libclcxx>`_459project.460