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