1.. raw:: html 2 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:: none 10.. role:: part 11.. role:: good 12 13.. contents:: 14 :local: 15 16================== 17OpenCL Support 18================== 19 20Clang has complete support of OpenCL C versions from 1.0 to 2.0. 21There is an ongoing work to support :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 28For general issues and bugs with OpenCL in clang refer to `the GitHub issue 29list 30<https://github.com/llvm/llvm-project/issues?q=is%3Aopen+is%3Aissue+label%3Aopencl>`__. 31 32Internals Manual 33================ 34 35This section acts as internal documentation for OpenCL features design 36as well as some important implementation aspects. It is primarily targeted 37at the advanced users and the toolchain developers integrating frontend 38functionality as a component. 39 40OpenCL Metadata 41--------------- 42 43Clang uses metadata to provide additional OpenCL semantics in IR needed for 44backends and OpenCL runtime. 45 46Each kernel will have function metadata attached to it, specifying the arguments. 47Kernel argument metadata is used to provide source level information for querying 48at runtime, for example using the `clGetKernelArgInfo 49<https://www.khronos.org/registry/OpenCL/specs/opencl-1.2.pdf#167>`_ 50call. 51 52Note that ``-cl-kernel-arg-info`` enables more information about the original 53kernel code to be added e.g. kernel parameter names will appear in the OpenCL 54metadata along with other information. 55 56The IDs used to encode the OpenCL's logical address spaces in the argument info 57metadata follows the SPIR address space mapping as defined in the SPIR 58specification `section 2.2 59<https://www.khronos.org/registry/spir/specs/spir_spec-2.0.pdf#18>`_ 60 61OpenCL Specific Options 62----------------------- 63 64In addition to the options described in :doc:`UsersManual` there are the 65following options specific to the OpenCL frontend. 66 67All the options in this section are frontend-only and therefore if used 68with regular clang driver they require frontend forwarding, e.g. ``-cc1`` 69or ``-Xclang``. 70 71.. _opencl_cl_ext: 72 73.. option:: -cl-ext 74 75Disables support of OpenCL extensions. All OpenCL targets provide a list 76of extensions that they support. Clang allows to amend this using the ``-cl-ext`` 77flag with a comma-separated list of extensions prefixed with ``'+'`` or ``'-'``. 78The syntax: ``-cl-ext=<(['-'|'+']<extension>[,])+>``, where extensions 79can be either one of `the OpenCL published extensions 80<https://www.khronos.org/registry/OpenCL>`_ 81or any vendor extension. Alternatively, ``'all'`` can be used to enable 82or disable all known extensions. 83 84Example disabling double support for the 64-bit SPIR target: 85 86 .. code-block:: console 87 88 $ clang -cc1 -triple spir64-unknown-unknown -cl-ext=-cl_khr_fp64 test.cl 89 90Enabling all extensions except double support in R600 AMD GPU can be done using: 91 92 .. code-block:: console 93 94 $ clang -cc1 -triple r600-unknown-unknown -cl-ext=-all,+cl_khr_fp16 test.cl 95 96.. _opencl_finclude_default_header: 97 98.. option:: -finclude-default-header 99 100Adds most of builtin types and function declarations during compilations. By 101default the OpenCL headers are not loaded by the frontend and therefore certain 102builtin types and most of builtin functions are not declared. To load them 103automatically this flag can be passed to the frontend (see also :ref:`the 104section on the OpenCL Header <opencl_header>`): 105 106 .. code-block:: console 107 108 $ clang -Xclang -finclude-default-header test.cl 109 110Alternatively the internal header `opencl-c.h` containing the declarations 111can be included manually using ``-include`` or ``-I`` followed by the path 112to the header location. The header can be found in the clang source tree or 113installation directory. 114 115 .. code-block:: console 116 117 $ clang -I<path to clang sources>/lib/Headers/opencl-c.h test.cl 118 $ clang -I<path to clang installation>/lib/clang/<llvm version>/include/opencl-c.h/opencl-c.h test.cl 119 120In this example it is assumed that the kernel code contains 121``#include <opencl-c.h>`` just as a regular C include. 122 123Because the header is very large and long to parse, PCH (:doc:`PCHInternals`) 124and modules (:doc:`Modules`) can be used internally to improve the compilation 125speed. 126 127To enable modules for OpenCL: 128 129 .. code-block:: console 130 131 $ 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 132 133Another way to circumvent long parsing latency for the OpenCL builtin 134declarations is to use mechanism enabled by :ref:`-fdeclare-opencl-builtins 135<opencl_fdeclare_opencl_builtins>` flag that is available as an alternative 136feature. 137 138.. _opencl_fdeclare_opencl_builtins: 139 140.. option:: -fdeclare-opencl-builtins 141 142In addition to regular header includes with builtin types and functions using 143:ref:`-finclude-default-header <opencl_finclude_default_header>`, clang 144supports a fast mechanism to declare builtin functions with 145``-fdeclare-opencl-builtins``. This does not declare the builtin types and 146therefore it has to be used in combination with ``-finclude-default-header`` 147if full functionality is required. 148 149**Example of Use**: 150 151 .. code-block:: console 152 153 $ clang -Xclang -fdeclare-opencl-builtins test.cl 154 155.. _opencl_fake_address_space_map: 156 157.. option:: -ffake-address-space-map 158 159Overrides the target address space map with a fake map. 160This allows adding explicit address space IDs to the bitcode for non-segmented 161memory architectures that do not have separate IDs for each of the OpenCL 162logical address spaces by default. Passing ``-ffake-address-space-map`` will 163add/override address spaces of the target compiled for with the following values: 164``1-global``, ``2-constant``, ``3-local``, ``4-generic``. The private address 165space is represented by the absence of an address space attribute in the IR (see 166also :ref:`the section on the address space attribute <opencl_addrsp>`). 167 168 .. code-block:: console 169 170 $ clang -cc1 -ffake-address-space-map test.cl 171 172.. _opencl_builtins: 173 174OpenCL builtins 175--------------- 176 177**Clang builtins** 178 179There are some standard OpenCL functions that are implemented as Clang builtins: 180 181- All pipe functions from `section 6.13.16.2/6.13.16.3 182 <https://www.khronos.org/registry/cl/specs/opencl-2.0-openclc.pdf#160>`_ of 183 the OpenCL v2.0 kernel language specification. 184 185- Address space qualifier conversion functions ``to_global``/``to_local``/``to_private`` 186 from `section 6.13.9 187 <https://www.khronos.org/registry/cl/specs/opencl-2.0-openclc.pdf#101>`_. 188 189- All the ``enqueue_kernel`` functions from `section 6.13.17.1 190 <https://www.khronos.org/registry/cl/specs/opencl-2.0-openclc.pdf#164>`_ and 191 enqueue query functions from `section 6.13.17.5 192 <https://www.khronos.org/registry/cl/specs/opencl-2.0-openclc.pdf#171>`_. 193 194**Fast builtin function declarations** 195 196The implementation of the fast builtin function declarations (available via the 197:ref:`-fdeclare-opencl-builtins option <opencl_fdeclare_opencl_builtins>`) consists 198of the following main components: 199 200- A TableGen definitions file ``OpenCLBuiltins.td``. This contains a compact 201 representation of the supported builtin functions. When adding new builtin 202 function declarations, this is normally the only file that needs modifying. 203 204- A Clang TableGen emitter defined in ``ClangOpenCLBuiltinEmitter.cpp``. During 205 Clang build time, the emitter reads the TableGen definition file and 206 generates ``OpenCLBuiltins.inc``. This generated file contains various tables 207 and functions that capture the builtin function data from the TableGen 208 definitions in a compact manner. 209 210- OpenCL specific code in ``SemaLookup.cpp``. When ``Sema::LookupBuiltin`` 211 encounters a potential builtin function, it will check if the name corresponds 212 to a valid OpenCL builtin function. If so, all overloads of the function are 213 inserted using ``InsertOCLBuiltinDeclarationsFromTable`` and overload 214 resolution takes place. 215 216OpenCL Extensions and Features 217------------------------------ 218 219Clang implements various extensions to OpenCL kernel languages. 220 221New functionality is accepted as soon as the documentation is detailed to the 222level sufficient to be implemented. There should be an evidence that the 223extension is designed with implementation feasibility in consideration and 224assessment of complexity for C/C++ based compilers. Alternatively, the 225documentation can be accepted in a format of a draft that can be further 226refined during the implementation. 227 228Implementation guidelines 229^^^^^^^^^^^^^^^^^^^^^^^^^ 230 231This section explains how to extend clang with the new functionality. 232 233**Parsing functionality** 234 235If an extension modifies the standard parsing it needs to be added to 236the clang frontend source code. This also means that the associated macro 237indicating the presence of the extension should be added to clang. 238 239The default flow for adding a new extension into the frontend is to 240modify `OpenCLExtensions.def 241<https://github.com/llvm/llvm-project/blob/main/clang/include/clang/Basic/OpenCLExtensions.def>`_ 242 243This will add the macro automatically and also add a field in the target 244options ``clang::TargetOptions::OpenCLFeaturesMap`` to control the exposure 245of the new extension during the compilation. 246 247Note that by default targets like `SPIR` or `X86` expose all the OpenCL 248extensions. For all other targets the configuration has to be made explicitly. 249 250Note that the target extension support performed by clang can be overridden 251with :ref:`-cl-ext <opencl_cl_ext>` command-line flags. 252 253**Library functionality** 254 255If an extension adds functionality that does not modify standard language 256parsing it should not require modifying anything other than header files and 257``OpenCLBuiltins.td`` detailed in :ref:`OpenCL builtins <opencl_builtins>`. 258Most commonly such extensions add functionality via libraries (by adding 259non-native types or functions) parsed regularly. Similar to other languages this 260is the most common way to add new functionality. 261 262Clang has standard headers where new types and functions are being added, 263for more details refer to 264:ref:`the section on the OpenCL Header <opencl_header>`. The macros indicating 265the presence of such extensions can be added in the standard header files 266conditioned on target specific predefined macros or/and language version 267predefined macros. 268 269**Pragmas** 270 271Some extensions alter standard parsing dynamically via pragmas. 272 273Clang provides a mechanism to add the standard extension pragma 274``OPENCL EXTENSION`` by setting a dedicated flag in the extension list entry of 275``OpenCLExtensions.def``. Note that there is no default behavior for the 276standard extension pragmas as it is not specified (for the standards up to and 277including version 3.0) in a sufficient level of detail and, therefore, 278there is no default functionality provided by clang. 279 280Pragmas without detailed information of their behavior (e.g. an explanation of 281changes it triggers in the parsing) should not be added to clang. Moreover, the 282pragmas should provide useful functionality to the user. For example, such 283functionality should address a practical use case and not be redundant i.e. 284cannot be achieved using existing features. 285 286Note that some legacy extensions (published prior to OpenCL 3.0) still 287provide some non-conformant functionality for pragmas e.g. add diagnostics on 288the use of types or functions. This functionality is not guaranteed to remain in 289future releases. However, any future changes should not affect backward 290compatibility. 291 292.. _opencl_addrsp: 293 294Address spaces attribute 295------------------------ 296 297Clang has arbitrary address space support using the ``address_space(N)`` 298attribute, where ``N`` is an integer number in the range specified in the 299Clang source code. This addresses spaces can be used along with the OpenCL 300address spaces however when such addresses spaces converted to/from OpenCL 301address spaces the behavior is not governed by OpenCL specification. 302 303An OpenCL implementation provides a list of standard address spaces using 304keywords: ``private``, ``local``, ``global``, and ``generic``. In the AST and 305in the IR each of the address spaces will be represented by unique number 306provided in the Clang source code. The specific IDs for an address space do not 307have to match between the AST and the IR. Typically in the AST address space 308numbers represent logical segments while in the IR they represent physical 309segments. 310Therefore, machines with flat memory segments can map all AST address space 311numbers to the same physical segment ID or skip address space attribute 312completely while generating the IR. However, if the address space information 313is needed by the IR passes e.g. to improve alias analysis, it is recommended 314to keep it and only lower to reflect physical memory segments in the late 315machine passes. The mapping between logical and target address spaces is 316specified in the Clang's source code. 317 318.. _cxx_for_opencl_impl: 319 320C++ for OpenCL Implementation Status 321==================================== 322 323Clang implements language versions 1.0 and 2021 published in `the official 324release of C++ for OpenCL Documentation 325<https://github.com/KhronosGroup/OpenCL-Docs/releases/tag/cxxforopencl-docrev2021.12>`_. 326 327Limited support of experimental C++ libraries is described in the :ref:`experimental features <opencl_experimenal>`. 328 329GitHub issues for this functionality are typically prefixed 330with '[C++4OpenCL]' - click `here 331<https://github.com/llvm/llvm-project/issues?q=is%3Aissue+is%3Aopen+%5BC%2B%2B4OpenCL%5D>`__ 332to view the full bug list. 333 334 335Missing features or with limited support 336---------------------------------------- 337 338- Support of C++ for OpenCL 2021 is currently in experimental phase. Refer to 339 :ref:`OpenCL 3.0 status <opencl_300>` for details of common missing 340 functionality from OpenCL 3.0. 341 342- IR generation for non-trivial global destructors is incomplete (See: 343 `PR48047 <https://llvm.org/PR48047>`_). 344 345- Support of `destrutors with non-default address spaces 346 <https://www.khronos.org/opencl/assets/CXX_for_OpenCL.html#_construction_initialization_and_destruction>`_ 347 is incomplete (See: `D109609 <https://reviews.llvm.org/D109609>`_). 348 349.. _opencl_300: 350 351OpenCL C 3.0 Usage 352================== 353 354OpenCL C 3.0 language standard makes most OpenCL C 2.0 features optional. Optional 355functionality in OpenCL C 3.0 is indicated with the presence of feature-test macros 356(list of feature-test macros is `here <https://www.khronos.org/registry/OpenCL/specs/3.0-unified/html/OpenCL_C.html#features>`__). 357Command-line flag :ref:`-cl-ext <opencl_cl_ext>` can be used to override features supported by a target. 358 359For cases when there is an associated extension for a specific feature (fp64 and 3d image writes) 360user should specify both (extension and feature) in command-line flag: 361 362 .. code-block:: console 363 364 $ clang -cc1 -cl-std=CL3.0 -cl-ext=+cl_khr_fp64,+__opencl_c_fp64 ... 365 $ clang -cc1 -cl-std=CL3.0 -cl-ext=-cl_khr_fp64,-__opencl_c_fp64 ... 366 367 368OpenCL C 3.0 Implementation Status 369---------------------------------- 370 371The following table provides an overview of features in OpenCL C 3.0 and their 372implementation status. 373 374+------------------------------+-------------------------+-----------------------------------------+----------------------+----------------------------------------------------------------------------------------------+ 375| Category | Feature | Status | Reviews | 376+==============================+=========================+=========================================+======================+==============================================================================================+ 377| Command line interface | New value for ``-cl-std`` flag | :good:`done` | https://reviews.llvm.org/D88300 | 378+------------------------------+-------------------------+-----------------------------------------+----------------------+----------------------------------------------------------------------------------------------+ 379| Predefined macros | New version macro | :good:`done` | https://reviews.llvm.org/D88300 | 380+------------------------------+-------------------------+-----------------------------------------+----------------------+----------------------------------------------------------------------------------------------+ 381| Predefined macros | Feature macros | :good:`done` | https://reviews.llvm.org/D95776 | 382+------------------------------+-------------------------+-----------------------------------------+----------------------+----------------------------------------------------------------------------------------------+ 383| Feature optionality | Generic address space | :good:`done` | https://reviews.llvm.org/D95778 and https://reviews.llvm.org/D103401 | 384+------------------------------+-------------------------+-----------------------------------------+----------------------+----------------------------------------------------------------------------------------------+ 385| Feature optionality | Builtin function overloads with generic address space | :good:`done` | https://reviews.llvm.org/D105526 | 386+------------------------------+-------------------------+-----------------------------------------+----------------------+----------------------------------------------------------------------------------------------+ 387| Feature optionality | Program scope variables in global memory | :good:`done` | https://reviews.llvm.org/D103191 | 388+------------------------------+-------------------------+-----------------------------------------+----------------------+----------------------------------------------------------------------------------------------+ 389| Feature optionality | 3D image writes including builtin functions | :part:`worked on` | https://reviews.llvm.org/D106260 (frontend) | 390+------------------------------+-------------------------+-----------------------------------------+----------------------+----------------------------------------------------------------------------------------------+ 391| Feature optionality | read_write images including builtin functions | :part:`worked on` | https://reviews.llvm.org/D104915 (frontend) and https://reviews.llvm.org/D107539 (functions) | 392+------------------------------+-------------------------+-----------------------------------------+----------------------+----------------------------------------------------------------------------------------------+ 393| Feature optionality | C11 atomics memory scopes, ordering and builtin function | :good:`done` | https://reviews.llvm.org/D106111 | 394+------------------------------+-------------------------+-----------------------------------------+----------------------+----------------------------------------------------------------------------------------------+ 395| Feature optionality | Blocks and Device-side kernel enqueue including builtin functions | :none:`unclaimed` | | 396+------------------------------+-------------------------+-----------------------------------------+----------------------+----------------------------------------------------------------------------------------------+ 397| Feature optionality | Pipes including builtin functions | :good:`done` | https://reviews.llvm.org/D107154 (frontend) and https://reviews.llvm.org/D105858 (functions) | 398+------------------------------+-------------------------+-----------------------------------------+----------------------+----------------------------------------------------------------------------------------------+ 399| Feature optionality | Work group collective builtin functions | :good:`done` | https://reviews.llvm.org/D105858 | 400+------------------------------+-------------------------+-----------------------------------------+----------------------+----------------------------------------------------------------------------------------------+ 401| Feature optionality | Image types and builtin functions | :good:`done` | https://reviews.llvm.org/D103911 (frontend) and https://reviews.llvm.org/D107539 (functions) | 402+------------------------------+-------------------------+-----------------------------------------+----------------------+----------------------------------------------------------------------------------------------+ 403| Feature optionality | Double precision floating point type | :good:`done` | https://reviews.llvm.org/D96524 | 404+------------------------------+-------------------------+-----------------------------------------+----------------------+----------------------------------------------------------------------------------------------+ 405| New functionality | RGBA vector components | :good:`done` | https://reviews.llvm.org/D99969 | 406+------------------------------+-------------------------+-----------------------------------------+----------------------+----------------------------------------------------------------------------------------------+ 407| New functionality | Subgroup functions | :part:`worked on` | https://reviews.llvm.org/D105858 | 408+------------------------------+-------------------------+-----------------------------------------+----------------------+----------------------------------------------------------------------------------------------+ 409| New functionality | Atomic mem scopes: subgroup, all devices including functions | :part:`worked on` | https://reviews.llvm.org/D103241 | 410+------------------------------+-------------------------+-----------------------------------------+----------------------+----------------------------------------------------------------------------------------------+ 411 412.. _opencl_experimenal: 413 414Experimental features 415===================== 416 417Clang provides the following new WIP features for the developers to experiment 418and provide early feedback or contribute with further improvements. 419Feel free to contact us on `cfe-dev 420<https://lists.llvm.org/mailman/listinfo/cfe-dev>`_ or file `a GitHub issue 421<https://github.com/llvm/llvm-project/issues/new>`_. 422 423.. _opencl_experimental_cxxlibs: 424 425C++ libraries for OpenCL 426------------------------ 427 428There is ongoing work to support C++ standard libraries from `LLVM's libcxx 429<https://libcxx.llvm.org/>`_ in OpenCL kernel code using C++ for OpenCL mode. 430 431It is currently possible to include `type_traits` from C++17 in the kernel 432sources when the following clang extensions are enabled 433``__cl_clang_function_pointers`` and ``__cl_clang_variadic_functions``, 434see :doc:`LanguageExtensions` for more details. The use of non-conformant 435features enabled by the extensions does not expose non-conformant behavior 436beyond the compilation i.e. does not get generated in IR or binary. 437The extension only appear in metaprogramming 438mechanism to identify or verify the properties of types. This allows to provide 439the full C++ functionality without a loss of portability. To avoid unsafe use 440of the extensions it is recommended that the extensions are disabled directly 441after the header include. 442 443**Example of Use**: 444 445The example of kernel code with `type_traits` is illustrated here. 446 447.. code-block:: c++ 448 449 #pragma OPENCL EXTENSION __cl_clang_function_pointers : enable 450 #pragma OPENCL EXTENSION __cl_clang_variadic_functions : enable 451 #include <type_traits> 452 #pragma OPENCL EXTENSION __cl_clang_function_pointers : disable 453 #pragma OPENCL EXTENSION __cl_clang_variadic_functions : disable 454 455 using sint_type = std::make_signed<unsigned int>::type; 456 457 __kernel void foo() { 458 static_assert(!std::is_same<sint_type, unsigned int>::value); 459 } 460 461The possible clang invocation to compile the example is as follows: 462 463 .. code-block:: console 464 465 $ clang -I<path to libcxx checkout or installation>/include test.clcpp 466 467Note that `type_traits` is a header only library and therefore no extra 468linking step against the standard libraries is required. See full example 469in `Compiler Explorer <https://godbolt.org/z/5WbnTfb65>`_. 470 471More OpenCL specific C++ library implementations built on top of libcxx 472are available in `libclcxx <https://github.com/KhronosGroup/libclcxx>`_ 473project. 474