Anastasia created this revision. Anastasia added reviewers: svenvh, mantognini, stuart, neil.hickey. Herald added subscribers: ebevhan, yaxunl.
Added documentation of C++ for OpenCL mode into Clang User Manual and Language Extensions document. 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,261 @@ 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 +- Only 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. Therefore 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 knowns 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>`). When references are bound to values address +space compatibility check will be performed. The logic will largely follow the rules +from address space pointer conversion (OpenCL v2.0 s6.5.5). + +**Default AS** + +All non-static methods take implicit object parameter 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 useful in case conversion to generic address space is undesirable (even if +it's legal). For example if we need to take advantage of memory bank accesses. +Please 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 + } + +**Built in operators** + +All builtin operators will be 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 ii; + } + + void bar() { + foo3<__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 end binary will contain +``@_GLOBAL__sub_I_test.cl`` kernel to be enqueued. + + 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
