1.. raw:: html 2 3 <style type="text/css"> 4 .none { background-color: #FFCCCC } 5 .partial { background-color: #FFFF99 } 6 .good { background-color: #CCFF99 } 7 </style> 8 9.. role:: none 10.. role:: partial 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. 21 22Clang also supports :ref:`the C++ for OpenCL kernel language <cxx_for_opencl_impl>`. 23 24There is an ongoing work to support :ref:`OpenCL 3.0 <opencl_300>`. 25 26There are also other :ref:`new and experimental features <opencl_experimenal>` available. 27 28For general issues and bugs with OpenCL in clang refer to `Bugzilla 29<https://bugs.llvm.org/buglist.cgi?component=OpenCL&list_id=172679&product=clang&resolution=--->`__. 30 31.. _cxx_for_opencl_impl: 32 33C++ for OpenCL Implementation Status 34==================================== 35 36Clang implements language version 1.0 published in `the official 37release of C++ for OpenCL Documentation 38<https://github.com/KhronosGroup/OpenCL-Docs/releases/tag/cxxforopencl-v1.0-r1>`_. 39 40Limited support of experimental C++ libraries is described in the :ref:`experimental features <opencl_experimenal>`. 41 42Bugzilla bugs for this functionality are typically prefixed 43with '[C++4OpenCL]' - click `here 44<https://bugs.llvm.org/buglist.cgi?component=OpenCL&list_id=204139&product=clang&query_format=advanced&resolution=---&sh ort_desc=%5BC%2B%2B4OpenCL%5D&short_desc_type=allwordssubstr>`_ 45to view the full bug list. 46 47 48Missing features or with limited support 49---------------------------------------- 50 51- Use of ObjC blocks is disabled and therefore the ``enqueue_kernel`` builtin 52 function is not supported currently. It is expected that if support for this 53 feature is added in the future, it will utilize C++ lambdas instead of ObjC 54 blocks. 55 56- IR generation for global destructors is incomplete (See: 57 `PR48047 <https://llvm.org/PR48047>`_). 58 59- There is no distinct file extension for sources that are to be compiled 60 in C++ for OpenCL mode (See: `PR48097 <https://llvm.org/PR48097>`_) 61 62.. _opencl_300: 63 64OpenCL 3.0 Implementation Status 65================================ 66 67The following table provides an overview of features in OpenCL C 3.0 and their 68implementation status. 69 70+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+ 71| Category | Feature | Status | Reviews | 72+==============================+==============================================================+======================+===========================================================================+ 73| Command line interface | New value for ``-cl-std`` flag | :good:`done` | https://reviews.llvm.org/D88300 | 74+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+ 75| Predefined macros | New version macro | :good:`done` | https://reviews.llvm.org/D88300 | 76+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+ 77| Predefined macros | Feature macros | :part:`worked on` | https://reviews.llvm.org/D89869 | 78+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+ 79| Feature optionality | Generic address space | :none:`unclaimed` | | 80+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+ 81| Feature optionality | Builtin function overloads with generic address space | :part:`worked on` | https://reviews.llvm.org/D92004 | 82+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+ 83| Feature optionality | Program scope variables in global memory | :none:`unclaimed` | | 84+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+ 85| Feature optionality | 3D image writes including builtin functions | :none:`unclaimed` | | 86+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+ 87| Feature optionality | read_write images including builtin functions | :none:`unclaimed` | | 88+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+ 89| Feature optionality | C11 atomics memory scopes, ordering and builtin function | :part:`worked on` | https://reviews.llvm.org/D92004 (functions only) | 90+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+ 91| Feature optionality | Device-side kernel enqueue including builtin functions | :none:`unclaimed` | | 92+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+ 93| Feature optionality | Pipes including builtin functions | :part:`worked on` | https://reviews.llvm.org/D92004 (functions only) | 94+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+ 95| Feature optionality | Work group collective functions | :part:`worked on` | https://reviews.llvm.org/D92004 | 96+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+ 97| New functionality | RGBA vector components | :none:`unclaimed` | | 98+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+ 99| New functionality | Subgroup functions | :part:`worked on` | https://reviews.llvm.org/D92004 | 100+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+ 101| New functionality | Atomic mem scopes: subgroup, all devices including functions | :part:`worked on` | https://reviews.llvm.org/D92004 (functions only) | 102+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+ 103 104.. _opencl_experimenal: 105 106Experimental features 107===================== 108 109Clang provides the following new WIP features for the developers to experiment 110and provide early feedback or contribute with further improvements. 111Feel free to contact us on `cfe-dev 112<https://lists.llvm.org/mailman/listinfo/cfe-dev>`_ or via `Bugzilla 113<https://bugs.llvm.org/>`__. 114 115Fast builtin function declarations 116---------------------------------- 117 118In addition to regular header includes with builtin types and functions using 119``-finclude-default-header`` explained in :doc:`UsersManual`, clang 120supports a fast mechanism to declare builtin functions with 121``-fdeclare-opencl-builtins``. This does not declare the builtin types and 122therefore it has to be used in combination with ``-finclude-default-header`` 123if full functionality is required. 124 125**Example of Use**: 126 127 .. code-block:: console 128 129 $ clang -Xclang -finclude-default-header test.cl 130 131Note that this is a frontend-only flag and therefore it requires the use of 132flags that forward options to the frontend, e.g. ``-cc1`` or ``-Xclang``. 133 134As this feature is still in experimental phase some changes might still occur 135on the command line interface side. 136 137C++ libraries for OpenCL 138------------------------ 139 140There is ongoing work to support C++ standard libraries from `LLVM's libcxx 141<https://libcxx.llvm.org/>`_ in OpenCL kernel code using C++ for OpenCL mode. 142 143It is currently possible to include `type_traits` from C++17 in the kernel 144sources when the following clang extensions are enabled 145``__cl_clang_function_pointers`` and ``__cl_clang_variadic_functions``, 146see :doc:`LanguageExtensions` for more details. The use of non-conformant 147features enabled by the extensions does not expose non-conformant behavior 148beyond the compilation i.e. does not get generated in IR or binary. 149The extension only appear in metaprogramming 150mechanism to identify or verify the properties of types. This allows to provide 151the full C++ functionality without a loss of portability. To avoid unsafe use 152of the extensions it is recommended that the extensions are disabled directly 153after the header include. 154 155**Example of Use**: 156 157The example of kernel code with `type_traits` is illustrated here. 158 159.. code-block:: c++ 160 161 #pragma OPENCL EXTENSION __cl_clang_function_pointers : enable 162 #pragma OPENCL EXTENSION __cl_clang_variadic_functions : enable 163 #include <type_traits> 164 #pragma OPENCL EXTENSION __cl_clang_function_pointers : disable 165 #pragma OPENCL EXTENSION __cl_clang_variadic_functions : disable 166 167 using sint_type = std::make_signed<unsigned int>::type; 168 169 __kernel void foo() { 170 static_assert(!std::is_same<sint_type, unsigned int>::value); 171 } 172 173The possible clang invocation to compile the example is as follows: 174 175 .. code-block:: console 176 177 $ clang -cl-std=clc++ -I<path to libcxx checkout or installation>/include test.cl 178 179Note that `type_traits` is a header only library and therefore no extra 180linking step against the standard libraries is required. 181