Anastasia updated this revision to Diff 210064. Anastasia marked 5 inline comments as done. Anastasia added a comment.
Addressed some review comments. Refs are still to be fixed! CHANGES SINCE LAST ACTION https://reviews.llvm.org/D64418/new/ https://reviews.llvm.org/D64418 Files: docs/LanguageExtensions.rst docs/UsersManual.rst
Index: docs/UsersManual.rst =================================================================== --- docs/UsersManual.rst +++ docs/UsersManual.rst @@ -2397,7 +2397,8 @@ This will produce a generic test.bc file that can be used in vendor toolchains to perform machine code generation. -Clang currently supports OpenCL C language standards up to v2.0. +Clang currently supports OpenCL C language standards up to v2.0. Starting from Clang9 +C++ mode is available for OpenCL (see :ref:`<opencl_cpp>`). OpenCL Specific Options ----------------------- @@ -2756,6 +2757,45 @@ enqueue query functions from `section 6.13.17.5 <https://www.khronos.org/registry/cl/specs/opencl-2.0-openclc.pdf#171>`_. +.. _opencl_cpp: + +C++ for OpenCL +-------------- + +Starting from Clang9 kernel code can contain C++17 features: classes, templates, +function overloading, type deduction, etc. Please note that this is not an +implementation of `OpenCL C++ +<https://www.khronos.org/registry/OpenCL/specs/2.2/pdf/OpenCL_Cxx.pdf>`_ and +there is no plan to support it in clang in any new releases in the near future. + +There are only a few restrictions on allowed C++ features. For detailed information +please refer to documentation on Extensions (:doc:`LanguageExtensions`). + +Since C++ features are to be used on top of OpenCL C functionality, all existing +restrictions from OpenCL C v2.0 will inherently apply. All OpenCL C builtin types +and function libraries are supported and can be used in the new mode. + +To enable the new mode pass the following command line option when compiling ``.cl`` +file ``-cl-std=c++`` or ``-std=c++``. + + .. code-block:: c++ + + template<class T> T add( T x, T y ) + { + return x + y; + } + + __kernel void test( __global float* a, __global float* b) + { + auto index = get_global_id(0); + a[index] = add(b[index], b[index+1]); + } + + + .. code-block:: console + + clang -cl-std=c++ test.cl + .. _target_features: Target-Specific Features and Limitations Index: docs/LanguageExtensions.rst =================================================================== --- docs/LanguageExtensions.rst +++ docs/LanguageExtensions.rst @@ -1516,6 +1516,268 @@ Query the presence of this new mangling with ``__has_feature(objc_protocol_qualifier_mangling)``. + +OpenCL Features +=============== + +C++ for OpenCL +-------------- + +This functionality is built on top of OpenCL C v2.0 and C++17. Regular C++ +features can be used in OpenCL kernel code. All functionality from OpenCL C +is inherited. This section describes minor differences to OpenCL C and any +limitations related to C++ support as well as interactions between OpenCL and +C++ features that are not documented elsewhere. + +Restrictions to C++17 +^^^^^^^^^^^^^^^^^^^^^ + +The following features are not supported: + +- Virtual functions +- ``dynamic_cast`` operator +- Non-placement ``new``/``delete`` operators +- Standard C++ libraries. Currently there is no solution for alternative + libraries provided. Future release will feature library support. + + +Interplay of OpenCL and C++ features +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ + +Address space behavior +"""""""""""""""""""""" + +Address spaces are part of the type qualifiers; many rules are just inherited +from the qualifier behavior documented in OpenCL C v2.0 s6.5 and Embedded C +extension ISO/IEC JTC1 SC22 WG14 N1021 s3.1. Note that since the address space +behavior in C++ is not documented formally yet, Clang extends existing concept +from C and OpenCL. For example conversion rules are extended from qualification +conversions but the compatibility is determined using sets and overlapping from +Embedded C (ISO/IEC JTC1 SC22 WG14 N1021 s3.1.3). For OpenCL it means that +implicit conversions are allowed from named to ``__generic`` but not vice versa +(OpenCL C v2.0 s6.5.5) except for constant address space. Most of the rules +are built on top of this behavior. + +**Casts** + +C style cast will follow OpenCL C v2.0 rules (s6.5.5). All cast operators will +permit implicit conversion to ``__generic``. However converting from named +address spaces to ``__generic`` can only be done using ``addrspace_cast``. Note +that conversions between ``__constant`` and any other is still disallowed. + +.. _opencl_cpp_addrsp_deduction: + +**Deduction** + +Address spaces are not deduced for: + +- non-pointer/non-reference template parameters or any dependent types are except + for template specializations. +- non-pointer/non-reference class members except for static data members that are + deduced to __global address space. +- non-pointer/non-reference alias declarations. +- ``decltype`` expression. + +.. code-block:: c++ + + template <typename T> + void foo() { + T m; // address space of m will be known at template instantiation time. + T * ptr; // ptr points to __generic address space object. + T & ref = ...; // ref references an object in __generic address space. + }; + + template <int N> + struct S { + int i; // i has no address space + static int ii; // ii is in global address space + int * ptr; // ptr points to __generic address space int. + int & ref = ...; // ref references int in __generic address space. + }; + + template <int N> + void bar() + { + S<N> s; // s is in __private address space + } + +TODO: Add example for type alias and decltype! + +**References** + +References types can be qualified with an address space. + +.. code-block:: c++ + + __private int & ref = ...; // references int in __private address space + +By default references will refer to ``__generic`` address space objects (except for +dependent types that are not template specializations +(:ref:`<opencl_cpp_addrsp_deduction>`). Address space compatibility checks are +performed when references are bound to values. The logic follows the rules from +address space pointer conversion (OpenCL v2.0 s6.5.5). + +**Default address space** + +All non-static methods take an implicit object parameter ``this`` that is a pointer +type. By default this pointer parameter is in ``__generic`` address space. +All concrete objects passed as an argument to ``this`` parameter will be +converted to ``__generic`` address space first if the conversion is valid. +Therefore programs using objects in ``__constant`` address space won't be compiled +unless address space is explicitly specified using address space method qualifiers +(:ref:`<opencl_cpp_addrspace_method_qual>`) as the conversion between +``__constant`` and ``__generic`` is disallowed. Method qualifiers can also be +in case conversion to ``__generic`` address space is undesirable (even if it is +legal), for example to take advantage of memory bank accesses. Note this not only +applies to regular methods but to constructors and destructors too. + +.. _opencl_cpp_addrspace_method_qual: + +**Method qualifier** + +Clang allows specifying address space qualifier on methods to signal that the method +is to be used with objects constructed in some specific address space. This works just +the same as qualifying methods with const or any other qualifiers. The overloading +resolution will select overload with most specific address space if multiple candidates +are provided. If there is no conversion to an address space among existing overloads +compilation will fail with a diagnostic. + +.. code-block:: c++ + + struct C { + void foo() __local; + void foo(); + }; + + __kernel void bar() { + __local C c1; // will resolve to the first foo + C c2; // will resolve to the second foo + __constant C c3; // error due to mismatching address spaces - can't convert to + // __local or __generic + } + +**Implicit special members** + +All implicit special members (default, copy, or move constructor, copy or move assignment) +will be generated with ``__generic`` address space. + +.. code-block:: c++ + + class C { + // Has the following implicit definition + // void () __generic; + // void (const __generic C &) __generic; + // void (__generic C &&) __generic; + // operator= '__generic C &(__generic C &&)' + // operator= '__generic C &(const __generic C &) __generic + } + +**Builtin operators** + +All builtin operators are available in the specific address spaces, thus no conversion +to ``__generic`` is performed. + +**Templates** + +There is no deduction of address spaces in non-pointer/non-reference template parameters and +dependent types (:ref:`<opencl_cpp_addrsp_deduction>`). The address space of template +parameter is deduced during the type deduction if it's not explicitly provided in +instantiation. + +.. code-block:: c++ + + 1 template<typename T> + 2 void foo(T* i){ + 3 T ii; + 4 } + 5 + 6 __global int g; + 7 void bar(){ + 8 foo(&g); // error: template instantiation failed as function scope variable appears to + 9 // be declared in global address space (see line 3) + 10 } + +It is not legal to specify multiple different address spaces between template definition and +instantiation. If multiple different address spaces are specified in template definition and +instantiation compilation of such program will fail with a diagnostic. + +.. code-block:: c++ + + template <typename T> + void foo() { + __private T var; + } + + void bar() { + foo<__global int>(); // error: conflicting address space qualifiers are provided + } + +Once template is instantiated regular restrictions for address spaces will apply. + +.. code-block:: c++ + + template<typename T> + void foo(){ + T var; + } + + void bar(){ + foo<__global int>(); // error: function scope variable cannot be declared in global + // address space + } + +**Temporary materialization** + +All temporaries are materialized in private address space. If a reference with some other +address space is bound to them, the conversion will be generated in case it's valid +otherwise compilation will fail with a diagnostic. + +.. code-block:: c++ + + int bar(const unsigned int &i); + + void foo() { + bar(1); // temporary is created in __private address space but converted + // to __generic address space of parameter reference + } + + __global const int& f(__global float &ref) { + return ref; // error: address space mismatch between temporary object + // created to hold value converted float->int and return + // value type + } + +**Initialization of local and constant address space objects** + +TODO + +Constructing and destroying global objects +^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^ + +Global objects are constructed before the first kernel using the global +objects is executed and destroyed just after the last kernel using the +program objects is executed. In OpenCL v2.0 drivers there is no specific +API for invoking global constructors. However, an easy workaround would be +to enqueue constructor initialization kernel that has a name +``@_GLOBAL__sub_I_<compled file name>``. This kernel is only present if there +are any global objects to be initialized in the compiled binary. One way to +check this is by passing ``CL_PROGRAM_KERNEL_NAMES`` to ``clGetProgramInfo`` +(OpenCL v2.0 s5.8.7). + +Note that if multiple files are compiled and linked into libraries multiple +kernels that initialize global objects for multiple modules would have to be +invoked. + +.. code-block:: console + + clang -cl-std=c++ test.cl + +If there are any global objects to be initialized the final binary will +contain ``@_GLOBAL__sub_I_test.cl`` kernel to be enqueued. + +Global destructors can not be invoked in OpenCL v2.0 drivers. However, all +memory used for program scope object is released on ``clReleaseProgram``. + Initializer lists for complex numbers in C ==========================================
_______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
