[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 ↗

(On Diff #208677)

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 ↗

(On Diff #208677)

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

1552 ↗

(On Diff #208677)

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 ↗

(On Diff #208677)

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 ↗

(On Diff #208677)

s/knowns/known/

1611–1615 ↗

(On Diff #208677)

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

1613–1614 ↗

(On Diff #208677)

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 ↗

(On Diff #208677)

AS might be better fully spelled out.

1619 ↗

(On Diff #208677)

take an implicit

1624 ↗

(On Diff #208677)

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

1627 ↗

(On Diff #208677)

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

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

1662 ↗

(On Diff #208677)

nitpicking: missing space before {.

1671 ↗

(On Diff #208677)

Typo: Builtin

1673 ↗

(On Diff #208677)

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

1679 ↗

(On Diff #208677)

The reference seems broken as well.

1708 ↗

(On Diff #208677)

s/foo3/foo/

1717 ↗

(On Diff #208677)

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

1753–1761 ↗

(On Diff #208677)

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

1768 ↗

(On Diff #208677)

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 ↗

(On Diff #208677)

s/end/final/

docs/UsersManual.rst

2401 ↗

(On Diff #208677)

The reference is broken.

2771–2772 ↗

(On Diff #208677)

I would make to sentences instead of using a ,.

2797 ↗

(On Diff #208677)

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 ↗

(On Diff #208677)

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

docs/UsersManual.rst

2774 ↗

(On Diff #208677)

functionality,

2778 ↗

(On Diff #208677)

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

[Anastasia]

docs/LanguageExtensions.rst

1772 ↗

(On Diff #208677)

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 ↗

(On Diff #208677)

Ok, specific here is as opposite to generic.

1753–1761 ↗

(On Diff #208677)

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

1758 ↗

(On Diff #208677)

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 ↗

(On Diff #208677)

This is the example I am explaining just below.

docs/UsersManual.rst

2797 ↗

(On Diff #208677)

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 ↗

(On Diff #210101)

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

1652–1657 ↗

(On Diff #210101)

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 ↗

(On Diff #210101)

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 ↗

(On Diff #210101)

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 ↗

(On Diff #210101)

Probably worth adding exceptions to the list.

1541 ↗

(On Diff #210101)

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

1553 ↗

(On Diff #210101)

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

1553 ↗

(On Diff #210101)

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

1555 ↗

(On Diff #210101)

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

1557 ↗

(On Diff #210101)

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 ↗

(On Diff #210101)

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

1564 ↗

(On Diff #210101)

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

1565 ↗

(On Diff #210101)

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

1565 ↗

(On Diff #210101)

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

1577 ↗

(On Diff #210101)

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

1626 ↗

(On Diff #210101)

the address space?

1631 ↗

(On Diff #210101)

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

1641 ↗

(On Diff #210101)

the overload

1652–1657 ↗

(On Diff #210101)

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

1668 ↗

(On Diff #210101)

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

1684 ↗

(On Diff #210101)

template parameterS

1702 ↗

(On Diff #210101)

"such a program" or "such programs"

1760 ↗

(On Diff #210101)

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 ↗

(On Diff #210101)

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 ↗

(On Diff #210101)

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

1779 ↗

(On Diff #210101)

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

docs/UsersManual.rst

2400 ↗

(On Diff #210101)

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

2401 ↗

(On Diff #210101)

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

2769 ↗

(On Diff #210101)

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 ↗

(On Diff #210101)

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

2779 ↗

(On Diff #210101)

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 ↗

(On Diff #215106)

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

1541 ↗

(On Diff #210101)

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

1553 ↗

(On Diff #210101)

I need to understand what needs re-wording

1557 ↗

(On Diff #210101)

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 ↗

(On Diff #210101)

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 ↗

(On Diff #210101)

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 ↗

(On Diff #210101)

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 ↗

(On Diff #210101)

Not sure whether you imply any action here.

1779 ↗

(On Diff #210101)

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 ↗

(On Diff #210101)

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

[svenvh]

docs/LanguageExtensions.rst

1561 ↗

(On Diff #215106)

"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 ↗

(On Diff #215106)

+The+ `__constant` address space

1570 ↗

(On Diff #215106)

"C-style casts follow"

1630 ↗

(On Diff #215106)

+the+ `__generic` address space

(also in the next few lines)

1647 ↗

(On Diff #215106)

+an+ address space qualifier

1650 ↗

(On Diff #215106)

ithe -> the

1651 ↗

(On Diff #215106)

Duplicate "to" (to to)

1652 ↗

(On Diff #215106)

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

1715 ↗

(On Diff #215106)

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

1730 ↗

(On Diff #215106)

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

1746 ↗

(On Diff #215106)

the `__private` address space

1747 ↗

(On Diff #215106)

some other -> another

1748 ↗

(On Diff #215106)

it's -> it is

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

1776 ↗

(On Diff #215106)

enqueue *a* constructor

1782 ↗

(On Diff #215106)

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

1793 ↗

(On Diff #215106)

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

1794 ↗

(On Diff #215106)

contain +the+

docs/UsersManual.rst

2765 ↗

(On Diff #215106)

Clang9 -> Clang 9 (with space).

2776 ↗

(On Diff #215106)

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

2778 ↗

(On Diff #215106)

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:

2769 ↗

(On Diff #210101)

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

Agreed.

2771 ↗

(On Diff #210101)

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

[Anastasia]

• Addressed comments from Sven.

[Anastasia]

docs/UsersManual.rst

2771 ↗

(On Diff #210101)

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!

cfe/trunk/docs/LanguageExtensions.rst

 
 
 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
+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 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
+- 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 yet, Clang extends existing concept
+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 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.
+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 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.
+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`` expression.
+- ``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.
.. code-block:: c++
   {
     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.
+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 ``__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 qualifiers on
-member functions
+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 ``__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 member functions but to
-constructors and destructors too.
+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 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 overload with most specific
-address space if multiple candidates are provided. If there is no conversion to
-to an address space among existing overloads compilation will fail with a
+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 ``__generic`` address space.
+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.
+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 template parameter is deduced during the type deduction if it's not explicitly
-provided in instantiation.
+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 appears to
-  9            // be declared in __global address space (see line 3)
+  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 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.
+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 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
+    foo<__global int>(); // error: conflicting address space qualifiers are provided
+                         // __global and __private
   }
 
-Once template is instantiated regular restrictions for address spaces will apply.
+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
+    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.
+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 are constructed before the first kernel using the global
+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 constructor initialization kernel that has a name
+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
+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 ``@_GLOBAL__sub_I_test.cl`` kernel to be enqueued.
+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
 ==========================================
 
 clang supports an extension which allows the following in C:

cfe/trunk/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. Starting from Clang9
-C++ mode is available for OpenCL (see :ref:`C++ for OpenCL <opencl_cpp>`).
+Clang currently supports OpenCL C language standards up to v2.0. Starting from
+clang 9 a C++ mode is available for OpenCL (see :ref:`C++ for OpenCL <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:
- All the ``enqueue_kernel`` functions from `section 6.13.17.1
   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,
+Starting from clang 9 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`).
+For detailed information about restrictions to allowed C++ features please
+refer to :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.
+and function libraries are supported and can be used in this mode.
 
-To enable the new mode pass the following command line option when compiling ``.cl``
-file ``-cl-std=clc++``, ``-cl-std=CLC++``, ``-std=clc++`` or ``-std=CLC++``.
+To enable the C++ for OpenCL mode, pass one of following command line options when
+compiling ``.cl`` file ``-cl-std=clc++``, ``-cl-std=CLC++``, ``-std=clc++`` or
+``-std=CLC++``.
 
    .. code-block:: c++
 
      template<class T> T add( T x, T y )
      {
        return x + y;
      }