181 lines
11 KiB
ReStructuredText
181 lines
11 KiB
ReStructuredText
|
.. 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.
|