[PATCH] D64418: [Docs][OpenCL] Documentation of C++ for OpenCL mode

Wed Jul 17 06:52:51 PDT 2019

kpet requested changes to this revision.
kpet added a comment.
This revision now requires changes to proceed.

Very useful to have all of this documented! Thanks!

================
Comment at: docs/LanguageExtensions.rst:1527
+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
----------------
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 [...].

================
Comment at: docs/LanguageExtensions.rst:1527
+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
----------------
kpet wrote:
> 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 [...].
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.

================
Comment at: docs/LanguageExtensions.rst:1535
+
+The following features are not supported:
+
----------------
Probably worth adding exceptions to the list.

================
Comment at: docs/LanguageExtensions.rst:1541
+- Standard C++ libraries. Currently there is no solution for alternative C++
+  libraries provided. Future release will feature library support.
+
----------------
I would remove this last sentence. Only a very small portion of the libray can be supported anyway.

================
Comment at: docs/LanguageExtensions.rst:1553
+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
----------------
I would remove the "yet", unless we can refer to an on-going and visible effort.

================
Comment at: docs/LanguageExtensions.rst:1553
+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
----------------
kpet wrote:
> I would remove the "yet", unless we can refer to an on-going and visible effort.
"Clang extends existing concept" feels like it should be re-worded.

================
Comment at: docs/LanguageExtensions.rst:1555
+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
----------------
"using sets and overlapping" feels incomplete. Pretending I can't guess: sets of what? What is overlapping what? :)

================
Comment at: docs/LanguageExtensions.rst:1557
+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
----------------
Why use `__generic` when `generic` is equally valid and easier on the eyes? The same comment applies to other mentions of address space keywords.

================
Comment at: docs/LanguageExtensions.rst:1558
+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.
----------------
This is ambiguous. One could read that conversions from `__generic` to `__constant` are allowed.

================
Comment at: docs/LanguageExtensions.rst:1564
+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
----------------
Isn't a conversion explicit if a cast operator is used?

================
Comment at: docs/LanguageExtensions.rst:1565
+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.
----------------
At the time of writing `addrspace_cast` doesn't seem to exist in trunk and C-style casts seem to work.

================
Comment at: docs/LanguageExtensions.rst:1565
+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.
----------------
kpet wrote:
> At the time of writing `addrspace_cast` doesn't seem to exist in trunk and C-style casts seem to work.
Didn't you mean "from `__generic` to named address spaces"?

================
Comment at: docs/LanguageExtensions.rst:1577
+- non-pointer/non-reference class members except for static data members that are
+  deduced to ``__global`` address space.
+- non-pointer/non-reference alias declarations.
----------------
Makes me wonder whether const members couldn't be in the `constant` address space.

================
Comment at: docs/LanguageExtensions.rst:1626
+``__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
----------------
the address space?

================
Comment at: docs/LanguageExtensions.rst:1631
+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 methods but to constructors and destructors too.
----------------
I would remove the "for example to take advantage of memory bank accesses" or make the example clear.

================
Comment at: docs/LanguageExtensions.rst:1641
+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
----------------
the overload

================
Comment at: docs/LanguageExtensions.rst:1652-1657
+ __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
+ }
----------------
mantognini wrote:
> `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
> ```
+1, I was about to make the same comment.

================
Comment at: docs/LanguageExtensions.rst:1668
+    // Has the following implicit definition
+    // void () __generic;
+    // void (const __generic C &) __generic;
----------------
Shouldn't this be `C() __generic`? Similar comments can be made for the other members.

================
Comment at: docs/LanguageExtensions.rst:1684
+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.
----------------
template parameterS

================
Comment at: docs/LanguageExtensions.rst:1702
+instantiation. If multiple different address spaces are specified in template definition and
+instantiation compilation of such program will fail with a diagnostic.
+
----------------
"such a program" or "such programs"

================
Comment at: docs/LanguageExtensions.rst:1760
+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
----------------
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.

================
Comment at: docs/LanguageExtensions.rst:1761
+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
----------------
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.

================
Comment at: docs/LanguageExtensions.rst:1762
+API for invoking global constructors. However, an easy workaround would be
+to enqueue constructor initialization kernel that has a name
+``@_GLOBAL__sub_I_<compiled file name>``. This kernel is only present if there
----------------
I would say that "applications are required to run all of these initialization kernels before any others if they exist".

================
Comment at: docs/LanguageExtensions.rst:1779
+
+Global destructors can not be invoked in OpenCL v2.0 drivers. However, all
+memory used for program scope objects is released on ``clReleaseProgram``.
----------------
This could be managed by the application with a scheme similar to the one you're proposing for constructor functions.

================
Comment at: docs/UsersManual.rst:2400

-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:`C++ for OpenCL <opencl_cpp>`).
----------------
Shouldn't "Clang9" be "Clang 9"? You may even want to spell it out "version 9 of Clang".

================
Comment at: docs/UsersManual.rst:2401
+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>`).

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

================
Comment at: docs/UsersManual.rst:2769
+<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.
+
----------------
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.

================
Comment at: docs/UsersManual.rst:2771
+
+There are only a few restrictions on allowed C++ features. For detailed information
+please refer to documentation on Extensions (:doc:`LanguageExtensions`).
----------------
Everything is relative. I would keep this 100% factual: "there are restrictions". This is not a sales pitch :).

================
Comment at: docs/UsersManual.rst:2779
+To enable the new mode pass the following command line option when compiling ``.cl``
+file ``-cl-std=c++`` or ``-std=c++``.
+
----------------
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++`?

CHANGES SINCE LAST ACTION
  https://reviews.llvm.org/D64418/new/

https://reviews.llvm.org/D64418