[Docs][OpenCL] Documentation of C++ for OpenCL mode

[Anastasia]

Summary

Added documentation of C++ for OpenCL mode into Clang User Manual and Language Extensions document.

[neil.hickey]

docs/LanguageExtensions.rst

1538

This seems weird, you say the following features are not supported then say Only placement new/delete. I think you mean to say new/delete is not supported except for placement new and delete.

1549

I'd replace ". Therefore" with "; "

1552

need a , between yet and Clang I think, to remove the ambiguity

[mantognini]

Here are a few comments from me but keep in mind that English is not my primary language

docs/LanguageExtensions.rst

1562–1565

I would suggest using __generic or generic address space. The later is a bit more verbose, but makes the sentence more natural I think.

1582

s/knowns/known/

1611–1615

The ref is broken. There's also one too many (.

1613–1614

When references are bound to values address space compatibility check will be performed.

This might just be me so this might actually be okay, but I feel the sentence is simpler when rewritten as "Address space compatibility checks are performed when references are bound to values." What do you think?

The logic will largely follow the rules from address space pointer conversion

Present tense should probably be better. Are there actually any differences? If no, I would drop "largely" from the sentence. Otherwise it might be good to actually list the differences.

1617

AS might be better fully spelled out.

1619

take an implicit

1624

ref is broken and there's one too manu (.

1627

For example if we need to take advantage of memory bank accesses.

Maybe "For example: to take advantage [...]"?

1662

nitpicking: missing space before {.

1671

Typo: Builtin

1673

"specific" seems to imply to the user will have to explicitly write down the AS. Maybe "requested" or "desired" would fit better?

1679

The reference seems broken as well.

1708

s/foo3/foo/

1717

nitpicking: indentation. Might as well rename it to ii to match the other example.

1753–1761

Is it expected of the user to call this special kernel, or is it expected that the driver takes care of that?

1768

As it stands, I'm not sure why this command is mentioned in this paragraph. Maybe it's a leftover of things that are now in the User's Manual?

1770

s/end/final/

docs/UsersManual.rst

2401

The reference is broken.

2771–2772

I would make to sentences instead of using a ,.

2797

I would remove this code sample; it is already explained above that -cl-std=c++ should be used.

[keryell]

Good improvement!

docs/LanguageExtensions.rst

1758

"<compiled"?
Is there an issue if you link a program from several files with the same file names?

docs/UsersManual.rst

2774

functionality,

2778

To enable the new mode, pass the following command line option when compiling a `.cl`

[Anastasia]

docs/LanguageExtensions.rst

1772

I think I should add something about dtors too.

[Anastasia]

Addressed some review comments. Refs are still to be fixed!

[Anastasia]

docs/LanguageExtensions.rst

1673

Ok, specific here is as opposite to generic.

1753–1761

In v2.0 drivers this is a manual step. In v2.2 drivers it is expected to be done automatically.

1758

I think IR linker makes them unique i.e. adds a suffix. Although I think the files (including the path) linked are supposed to be unique. Anyway, good point to look into.

No idea about obj linker. We currently don't support this feature.

1768

This is the example I am explaining just below.

docs/UsersManual.rst

2797

Manual is good for examples. Some new developers might find it useful to have the full line even though it's quite simple.

[mantognini]

Beside my two comments, I think this looks good.

docs/LanguageExtensions.rst

1614

( -> , (there's no matching closing parenthesis.)

1652–1657

foo() isn't actually called here. You probably meant to write this:

__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

[kpet]

Very useful to have all of this documented! Thanks!

docs/LanguageExtensions.rst

1527

I'm tempted to suggest you merge the first two sentences and tone down the claim in the second. "This functionality [...] and C++17, enabling _some_ regular C++ features [...].

1527

s/All/Most/? Stating that all functionality from OpenCL C is available right before stating that the differences will be described feels like a contradiction.

1535

Probably worth adding exceptions to the list.

1541

I would remove this last sentence. Only a very small portion of the libray can be supported anyway.

1553

I would remove the "yet", unless we can refer to an on-going and visible effort.

1553

"Clang extends existing concept" feels like it should be re-worded.

1555

"using sets and overlapping" feels incomplete. Pretending I can't guess: sets of what? What is overlapping what? :)

1557

Why use __generic when generic is equally valid and easier on the eyes? The same comment applies to other mentions of address space keywords.

1558

This is ambiguous. One could read that conversions from __generic to __constant are allowed.

1564

Isn't a conversion explicit if a cast operator is used?

1565

At the time of writing addrspace_cast doesn't seem to exist in trunk and C-style casts seem to work.

1565

Didn't you mean "from __generic to named address spaces"?

1577

Makes me wonder whether const members couldn't be in the constant address space.

1626

the address space?

1631

I would remove the "for example to take advantage of memory bank accesses" or make the example clear.

1641

the overload

1652–1657

+1, I was about to make the same comment.

1668

Shouldn't this be C() __generic? Similar comments can be made for the other members.

1684

template parameterS

1702

"such a program" or "such programs"

1760

I would say "must be constructed before the first kernel begins execution" and "may be destroyed once the last kernel has completed its execution". These are requirements for a runtime environment, not functionality provided by Clang. The documentation should only describe the interface with the runtime and the requirements on the runtime.

1761

OpenCL doesn't have an API for this per-se, regardless of the version. The only mechanism currently specified is the SPIR-V Initializer and Finalizer execution modes for entrypoints, either of which currently require an OpenCL 2.2 implementation.

1762

I would say that "applications are required to run all of these initialization kernels before any others if they exist".

1779

This could be managed by the application with a scheme similar to the one you're proposing for constructor functions.

docs/UsersManual.rst

2400–2401

Shouldn't "Clang9" be "Clang 9"? You may even want to spell it out "version 9 of Clang".

2401

"A C++ mode" or "A non-standard C++ language extension" to be more precise?

2769

No plans that we're aware of ;). For all we know, there could be a company/individual working on this. I think conjectures on people's plans don't belong in the documentation.

2771

Everything is relative. I would keep this 100% factual: "there are restrictions". This is not a sales pitch :).

2779

I find -std=c++ confusing/misleading. Maybe we should change the way -cl-std works so we can have -std=clc++ and -cl-std=c++?

[Anastasia]

In D64418#1589401, @kpet wrote:

Very useful to have all of this documented! Thanks!

Would it be ok if I fix those in a separate commit? I would really like to commit the core part before the release branch is taken.

[kpet]

Would it be ok if I fix those in a separate commit? I would really like to commit the core part before the release branch is taken.

I'm fine with this as long as we can continue the discussion for open comments in this review. Accepting the revision as I don't seem to be able to go back on my request for another patch.

[Anastasia]

Reopening to continue reviews.

[Anastasia]

Addressed various comments from Kevin.

[Anastasia]

docs/LanguageExtensions.rst

327

FYI, this is not part of my change. It was just hard to produce combined diff.

1541

I think it's valuable especially for metaprogramming. I was being asked about it.

1553

I need to understand what needs re-wording

1557

I prefer to use __ for two reasons (1) it's more official in OpenCL docs and tutorials (2) I sometimes omit the noun and just use it as a short term in the text.

1564

It's complicated to explain but basically some cast operators perform implicit conversions of what they are not expected to do because implicit conversions are allowed everywhere even if the operator isn't intended to do the conversion explicitly.

Similar example in C++, most of operators accept converting to const objects.

1565

At the time of writing addrspace_cast doesn't seem to exist in trunk and C-style casts seem to work.

I think it's ok to document full extension even if it's WIP.

1577

No constant specify memory region but const indicates that object is not changed during lifetime. You can cast away const but you can't convert constant object to any other address space.

1761

Not sure whether you imply any action here.

1779

The Itanium C++ ABIs for destructors are different and we haven't got it working yet for OpenCL. I would leave the description as is for now.

docs/UsersManual.rst

2771

I think it's fair to say few restrictions because the majority of C++ just works in OpenCL.

[svenvh]

docs/LanguageExtensions.rst

1561

"For OpenCL it means that implicit conversions are allowed from a ​named address space except for conversion from the `__constant to the __generic` address space."

I would omit "that is a superset of all others except for `__constant`" as that is already documented in the OpenCL spec and it might confuse readers here. If you want to keep it, I suggest rephrasing it as a separate sentence.

1563

+The+ `__constant` address space

1570

"C-style casts follow"

1630

+the+ `__generic` address space

(also in the next few lines)

1647

+an+ address space qualifier

1650

ithe -> the

1651

Duplicate "to" (to to)

1652

Add comma:
... overloads, compilation ...

1715

Add comma:
... instantiation, compilation ...

1730

Once a template has been instantiated, regular restrictions for address spaces will apply.

1746

the `__private` address space

1747

some other -> another

1748

it's -> it is

add comma:
... valid, otherwise ...

1776

enqueue *a* constructor

1782

Add comma:
... libraries, multiple ...

1793

Add comma:
... initialized, the ...

1794

contain +the+

docs/UsersManual.rst

2765

Clang9 -> Clang 9 (with space).

2769

I think conjectures on people's plans don't belong in the documentation.

Agreed.

2771

I agree with @kpet and also prefer "There are restrictions" instead of an unquantified statement.

2776

I wouldn't say "new" mode as this is part of a persistent user manual.

2778

Again I would avoid "new mode".

"To enable the C++ for OpenCL mode,"

Also: pass *one of* the following command line options when compiling a `.cl`

So in total:
To enable the C++ for OpenCL mode, pass one of the following command line options when compiling a `.cl` file:

[Anastasia]

• Addressed comments from Sven.

[Anastasia]

docs/UsersManual.rst

2771

Ok, I don't understand what value does it bring to make information less specific but potentially it's not important enough to continue the argument. The idea of this sentence was indeed to point out that the mode is just very close to C++ with very few restrictions. This isn't a program so it can't have a formal definition for every statement. :P

[Anastasia]

Added small corrections in various parts.

[svenvh]

LGTM!

docs/LanguageExtensions.rst

 option for a warning and returns true if that is a valid warning option.
 
 .. code-block:: c++
 
   #if __has_warning("-Wformat")
   ...
   #endif
 
 Builtin Macros
 ==============
 
 ``__BASE_FILE__``
   Defined to a string that contains the name of the main input file passed to
   Clang.
   
 ``__FILE_NAME__``
   Clang-specific extension that functions similar to ``__FILE__`` but only

 
 This was a non-backward compatible mangling change to the ABI.  This change
 allows proper overloading, and also prevents mangling conflicts with template
 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. 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
 ==========================================
 
 clang supports an extension which allows the following in C:
 
 .. code-block:: c++
 
   #include <math.h>
   #include <complex.h>
   complex float x = { 1.0f, INFINITY }; // Init to (1, Inf)
 
 This construct is useful because there is no way to separately initialize the
 real and imaginary parts of a complex variable in standard C, given that clang
 does not support ``_Imaginary``.  (Clang also supports the ``__real__`` and
 ``__imag__`` extensions from gcc, which help in some cases, but are not usable
 in static initializers.)
 
 Note that this extension does not allow eliding the braces; the meaning of the
 following two lines is different:
 
 .. code-block:: c++
 
   complex float x[] = { { 1.0f, 1.0f } }; // [0] = (1, 1)
   complex float x[] = { 1.0f, 1.0f }; // [0] = (1, 0), [1] = (1, 0)
 
 This extension also works in C++ mode, as far as that goes, but does not apply
 to the C++ ``std::complex``.  (In C++11, list initialization allows the same
 syntax to be used with ``std::complex`` with the same meaning.)
 

docs/UsersManual.rst

Compiling to bitcode can be done as follows:
 
    .. code-block:: console
 
      $ clang -c -emit-llvm test.cl
 
 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
 -----------------------
 
 Most of the OpenCL build options from `the specification v2.0 section 5.8.4
 <https://www.khronos.org/registry/cl/specs/opencl-2.0.pdf#200>`_ are available.
 
 Examples:
- Address space qualifier conversion functions ``to_global``/``to_local``/``to_private``
   from `section 6.13.9
   <https://www.khronos.org/registry/cl/specs/opencl-2.0-openclc.pdf#101>`_.
 
 - All the ``enqueue_kernel`` functions from `section 6.13.17.1
   <https://www.khronos.org/registry/cl/specs/opencl-2.0-openclc.pdf#164>`_ and
   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
 ========================================
 
 CPU Architectures Features and Limitations
 ------------------------------------------