.. raw:: html

  <style type="text/css">
    .none { background-color: #FFCCCC }
    .partial { background-color: #FFFF99 }
    .good { background-color: #CCFF99 }
  </style>

.. role:: none
.. role:: partial
.. role:: good

.. contents::
   :local:

==================
OpenCL Support
==================

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

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

There is an ongoing work to support :ref:`OpenCL 3.0 <opencl_300>`.

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

For general issues and bugs with OpenCL in clang refer to `Bugzilla
<https://bugs.llvm.org/buglist.cgi?component=OpenCL&list_id=172679&product=clang&resolution=--->`__.

.. _cxx_for_opencl_impl:

C++ for OpenCL Implementation Status
====================================

Clang implements language version 1.0 published in `the official
release of C++ for OpenCL Documentation
<https://github.com/KhronosGroup/OpenCL-Docs/releases/tag/cxxforopencl-v1.0-r1>`_.

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

Bugzilla bugs for this functionality are typically prefixed
with '[C++4OpenCL]' - click `here
<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>`_
to view the full bug list.


Missing features or with limited support
----------------------------------------

- Use of ObjC blocks is disabled and therefore the ``enqueue_kernel`` builtin
  function is not supported currently. It is expected that if support for this
  feature is added in the future, it will utilize C++ lambdas instead of ObjC
  blocks.

- IR generation for global destructors is incomplete (See:
  `PR48047 <https://llvm.org/PR48047>`_).

- There is no distinct file extension for sources that are to be compiled
  in C++ for OpenCL mode (See: `PR48097 <https://llvm.org/PR48097>`_)

.. _opencl_300:

OpenCL 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                               | :good:`done`         | https://reviews.llvm.org/D88300                                           |
+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+
| Predefined macros            | New version macro                                            | :good:`done`         | https://reviews.llvm.org/D88300                                           |
+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+
| Predefined macros            | Feature macros                                               | :part:`worked on`    | https://reviews.llvm.org/D89869                                           |
+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+
| Feature optionality          | Generic address space                                        | :none:`unclaimed`    |                                                                           |
+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+
| Feature optionality          | Builtin function overloads with generic address space        | :part:`worked on`    | https://reviews.llvm.org/D92004                                           |
+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+
| Feature optionality          | Program scope variables in global memory                     | :none:`unclaimed`    |                                                                           |
+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+
| Feature optionality          | 3D image writes including builtin functions                  | :none:`unclaimed`    |                                                                           |
+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+
| Feature optionality          | read_write images including builtin functions                | :none:`unclaimed`    |                                                                           |
+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+
| Feature optionality          | C11 atomics memory scopes, ordering and builtin function     | :part:`worked on`    | https://reviews.llvm.org/D92004 (functions only)                          |
+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+
| Feature optionality          | Device-side kernel enqueue including builtin functions       | :none:`unclaimed`    |                                                                           |
+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+
| Feature optionality          | Pipes including builtin functions                            | :part:`worked on`    | https://reviews.llvm.org/D92004 (functions only)                          |
+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+
| Feature optionality          | Work group collective functions                              | :part:`worked on`    | https://reviews.llvm.org/D92004                                           |
+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+
| New functionality            | RGBA vector components                                       | :none:`unclaimed`    |                                                                           |
+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+
| New functionality            | Subgroup functions                                           | :part:`worked on`    | https://reviews.llvm.org/D92004                                           |
+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+
| New functionality            | Atomic mem scopes: subgroup, all devices including functions | :part:`worked on`    | https://reviews.llvm.org/D92004 (functions only)                          |
+------------------------------+--------------------------------------------------------------+----------------------+---------------------------------------------------------------------------+

.. _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 `cfe-dev
<https://lists.llvm.org/mailman/listinfo/cfe-dev>`_ or via `Bugzilla
<https://bugs.llvm.org/>`__.

Fast builtin function declarations
----------------------------------

In addition to regular header includes with builtin types and functions using
``-finclude-default-header`` explained in :doc:`UsersManual`, 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**:

    .. code-block:: console
 
      $ clang -Xclang -finclude-default-header test.cl

Note that this is a frontend-only flag and therefore it requires the use of
flags that forward options to the frontend, e.g. ``-cc1`` or ``-Xclang``.

As this feature is still in experimental phase some changes might still occur
on the command line interface side.

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.

.. code-block:: 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:

   .. code-block:: console

     $ clang -cl-std=clc++  -I<path to libcxx checkout or installation>/include test.cl

Note that `type_traits` is a header only library and therefore no extra
linking step against the standard libraries is required.