[clang] e456165 - [OpenCL] Add link to C++ for OpenCL documentation
Anastasia Stulova via cfe-commits
cfe-commits at lists.llvm.org
Fri Jan 3 04:10:48 PST 2020
Author: Anastasia Stulova
Date: 2020-01-03T12:01:03Z
New Revision: e456165f9fec9148566849f21bc4f7dda2fea034
URL: https://github.com/llvm/llvm-project/commit/e456165f9fec9148566849f21bc4f7dda2fea034
DIFF: https://github.com/llvm/llvm-project/commit/e456165f9fec9148566849f21bc4f7dda2fea034.diff
LOG: [OpenCL] Add link to C++ for OpenCL documentation
Remove description of language mode from the language
extensions and add a link to pdf document.
Tags: #clang
Differential Revision: https://reviews.llvm.org/D72076
Added:
Modified:
clang/docs/LanguageExtensions.rst
clang/docs/UsersManual.rst
Removed:
################################################################################
diff --git a/clang/docs/LanguageExtensions.rst b/clang/docs/LanguageExtensions.rst
index d9a4862dbe80..b0f57202e079 100644
--- a/clang/docs/LanguageExtensions.rst
+++ b/clang/docs/LanguageExtensions.rst
@@ -1631,285 +1631,6 @@ parameters of protocol-qualified type.
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 enabling most of
-regular C++ features in OpenCL kernel code. Most functionality from OpenCL C
-is inherited. This section describes minor
diff erences 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
-- Exceptions
-- ``dynamic_cast`` operator
-- Non-placement ``new``/``delete`` operators
-- Standard C++ libraries. Currently there is no solution for alternative C++
- 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, Clang extends the existing concept
-from C and OpenCL. For example conversion rules are extended from qualification
-conversion but the compatibility is determined using notation of sets and
-overlapping of address spaces from Embedded C (ISO/IEC JTC1 SC22 WG14 N1021
-s3.1.3). For OpenCL it means that implicit conversions are allowed from
-a named address space (except for ``__constant``) to ``__generic`` (OpenCL C
-v2.0 6.5.5). Reverse conversion is only allowed explicitly. The ``__constant``
-address space does not overlap with any other and therefore no valid conversion
-between ``__constant`` and other address spaces exists. Most of the rules
-follow this logic.
-
-**Casts**
-
-C-style casts follow OpenCL C v2.0 rules (s6.5.5). All cast operators
-permit conversion to ``__generic`` implicitly. However converting from
-``__generic`` to named address spaces can only be done using ``addrspace_cast``.
-Note that conversions between ``__constant`` and any other address space
-are disallowed.
-
-.. _opencl_cpp_addrsp_deduction:
-
-**Deduction**
-
-Address spaces are not deduced for:
-
-- non-pointer/non-reference template parameters or any dependent types 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`` expressions.
-
-.. 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**
-
-Reference 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
-(see :ref:`Deduction <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 member functions take an implicit object parameter ``this`` that
-is a pointer type. By default this pointer parameter is in the ``__generic``
-address space. All concrete objects passed as an argument to ``this`` parameter
-will be converted to the ``__generic`` address space first if such conversion is
-valid. Therefore programs using objects in the ``__constant`` address space will
-not be compiled unless the address space is explicitly specified using address
-space qualifiers on member functions
-(see :ref:`Member function qualifier <opencl_cpp_addrspace_method_qual>`) as the
-conversion between ``__constant`` and ``__generic`` is disallowed. Member function
-qualifiers can also be used in case conversion to the ``__generic`` address space
-is undesirable (even if it is legal). For example, a method can be implemented to
-exploit memory access coalescing for segments with memory bank. This not only
-applies to regular member functions but to constructors and destructors too.
-
-.. _opencl_cpp_addrspace_method_qual:
-
-**Member function qualifier**
-
-Clang allows specifying an address space qualifier on member functions to signal
-that they are to be used with objects constructed in some specific address space.
-This works just the same as qualifying member functions with ``const`` or any
-other qualifiers. The overloading resolution will select the candidate with the
-most specific address space if multiple candidates are provided. If there is no
-conversion to an address space among candidates, compilation will fail with a
-diagnostic.
-
-.. code-block:: c++
-
- struct C {
- void foo() __local;
- void foo();
- };
-
- __kernel void bar() {
- __local C c1;
- C c2;
- __constant C c3;
- c1.foo(); // will resolve to the first foo
- c2.foo(); // will resolve to the second foo
- c3.foo(); // 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, destructor) will be generated with the ``__generic`` address space.
-
-.. code-block:: c++
-
- class C {
- // Has the following implicit definition
- // void C() __generic;
- // void C(const __generic C &) __generic;
- // void C(__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 (see :ref:`Deduction <opencl_cpp_addrsp_deduction>`).
-The address space of a template parameter is deduced during type deduction if
-it is not explicitly provided in the instantiation.
-
-.. code-block:: c++
-
- 1 template<typename T>
- 2 void foo(T* i){
- 3 T var;
- 4 }
- 5
- 6 __global int g;
- 7 void bar(){
- 8 foo(&g); // error: template instantiation failed as function scope variable
- 9 // appears to be declared in __global address space (see line 3)
- 10 }
-
-It is not legal to specify multiple
diff erent address spaces between template
-definition and instantiation. If multiple
diff erent address spaces are specified in
-template definition and instantiation, compilation of such a 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
- // __global and __private
- }
-
-Once a template has been 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 the ``__private`` address space. If a
-reference with another address space is bound to them, the conversion will be
-generated in case it is 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 (can't convert from __private to __global)
- }
-
-**Initialization of local and constant address space objects**
-
-TODO
-
-Constructing and destroying global objects
-^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
-
-Global objects must be 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 a constructor initialization kernel that has a name
-``@_GLOBAL__sub_I_<compiled 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.
-
-Applications are currently required to run initialization of global objects
-manually before running any kernels in which the objects are used.
-
-.. code-block:: console
-
- clang -cl-std=clc++ test.cl
-
-If there are any global objects to be initialized, the final binary will
-contain the ``@_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 objects is released on ``clReleaseProgram``.
-
Initializer lists for complex numbers in C
==========================================
diff --git a/clang/docs/UsersManual.rst b/clang/docs/UsersManual.rst
index 07201f0d8616..f55ffccc84d0 100644
--- a/clang/docs/UsersManual.rst
+++ b/clang/docs/UsersManual.rst
@@ -2856,6 +2856,10 @@ function to the custom ``my_ext`` extension.
Declaring the same types in
diff erent vendor extensions is disallowed.
+Clang also supports language extensions documented in `The OpenCL C Language
+Extensions Documentation
+<https://github.com/KhronosGroup/Khronosdotorg/blob/master/api/opencl/assets/OpenCL_LangExt.pdf>`_.
+
OpenCL Metadata
---------------
@@ -3031,8 +3035,9 @@ 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.
-For detailed information about restrictions to allowed C++ features please
-refer to :doc:`LanguageExtensions`.
+For detailed information about this language refer to `The C++ for OpenCL
+Programming Language Documentation
+<https://github.com/KhronosGroup/Khronosdotorg/blob/master/api/opencl/assets/CXX_for_OpenCL.pdf>`_.
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
@@ -3060,6 +3065,35 @@ compiling ``.cl`` file ``-cl-std=clc++``, ``-cl-std=CLC++``, ``-std=clc++`` or
clang -cl-std=clc++ test.cl
+Constructing and destroying global objects
+^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^^
+
+Global objects must be 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 a constructor
+initialization kernel that has a name ``@_GLOBAL__sub_I_<compiled 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.
+
+Applications are currently required to run initialization of global objects manually
+before running any kernels in which the objects are used.
+
+ .. code-block:: console
+
+ clang -cl-std=clc++ test.cl
+
+If there are any global objects to be initialized, the final binary will contain
+the ``@_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 objects is released on ``clReleaseProgram``.
+
+
.. _target_features:
Target-Specific Features and Limitations
More information about the cfe-commits
mailing list