OpenCL Support

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

Clang also supports the C++ for OpenCL kernel language.

There is an ongoing work to support OpenCL 3.0.

There are also other new and experimental features available.

For general issues and bugs with OpenCL in clang refer to Bugzilla.

C++ for OpenCL Implementation Status

Clang implements language version 1.0 published in the official release of C++ for OpenCL Documentation.

Limited support of experimental C++ libraries is described in the experimental features.

Bugzilla bugs for this functionality are typically prefixed with ‘[C++4OpenCL]’ - click here 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).
  • There is no distinct file extension for sources that are to be compiled in C++ for OpenCL mode (See: PR48097)

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

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 or via Bugzilla.

Fast builtin function declarations

In addition to regular header includes with builtin types and functions using -finclude-default-header explained in Clang Compiler User’s Manual, 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:

$ 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 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 Clang Language Extensions 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.

#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:

$ 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.