Skip to content

Commit

Permalink
[C++ for OpenCL] Minor editorial fixes (#742)
Browse files Browse the repository at this point in the history 
This commit contains the following small fixups:
- Formatting and styling improvements;
- Corrections to cross-references;
- Wording improvements and clarifications;
- Improvements to code examples.

Aside from these, it also contains clarification
about releasing version 2021 as provisional.

Co-authored-by: Sven van Haastregt <[email protected]>
  • Loading branch information
@AnastasiaStulova @svenvh
AnastasiaStulova and svenvh authored Dec 20, 2021
1 parent 1525c7b commit 211d598
Show file tree
Hide file tree
Showing 6 changed files with 129 additions and 89 deletions.
147 changes: 88 additions & 59 deletions cxx4opencl/address_spaces.txt
View file
Original file line number Diff line number Diff line change
Expand Up @@ -28,6 +28,7 @@ therefore, no valid conversion between `+__constant+` and any other address spac
exists. This is aligned with rules from `OpenCL C 3.0 s6.7.9` and this logic
regulates semantics described in this section.

[[address_space_casts]]
==== Casts

C-style casts follow rules of `OpenCL C 3.0 s6.7.9`. Conversions of
Expand Down Expand Up @@ -86,7 +87,7 @@ __private int & ref = ...; // references int in __private address space.
By default references refer to generic address space objects if generic
address space is supported or private address space otherwise, except
for dependent types that are not template specializations (see
<<addrspace-deduction, _Deduction_>>).
<<addrspace-deduction, _Address space inference_>>).

[source,cpp]
------------
Expand All @@ -101,11 +102,13 @@ conversion (`OpenCL C 3.0 s6.7.9`).
[source,cpp]
------------
void f(float &ref, __global float &globref) {
const int& tmp = ref; // legal - reference to generic/__private address space object can
// bind to a temporary object created in __private address space.
const int& tmp = ref; // legal - reference to generic/__private address space object
// can bind to a temporary object created in __private
// address space.

__global const int& globtmp = globref; // error reference to global address space object cannot bind
// to a temporary object created in __private address space.
__global const int& globtmp = globref; // error: reference to global address space
// object cannot bind to a temporary object
// created in __private address space.
}
------------

Expand All @@ -132,9 +135,10 @@ that is a pointer type. By default the `this` pointer parameter is in the
generic address space if it is supported and in the private address space
otherwise. All concrete objects passed as an argument to the
implicit `this` parameter will be converted to this default (generic or private)
address space first if such conversion is valid. Therefore, programs using objects
in disjoint address spaces will not be compiled unless the address spaces for the
object parameter `this` is explicitly specified using address space qualifiers
address space first if such conversion is valid. Therefore, when member functions
are called with objects created in disjoint address spaces from the default one,
the compilation must fail. To prevent the failure the address space on implicit
object parameter `this` must be specified using address space qualifiers
on member functions (see <<addrspace-member-function-qualifiers,
_Member function qualifier_>>). For example, use of member functions with
objects in `+__constant+` address space will always require a `+__constant+`
Expand All @@ -148,7 +152,7 @@ be implemented to exploit memory access coalescing for segments with memory bank
Address spaces are not deduced for:

* non-pointer/non-reference template parameters except for template
specializations or non-type type based template parameters.
specializations or non-type based template parameters.
* non-pointer/non-reference class members except for static data members
that are deduced to the `+__global+` address space for {cpp} for OpenCL 1.0
or {cpp} for OpenCL 2021 with the
Expand Down Expand Up @@ -316,37 +320,37 @@ class C {

// C(const C & par); /* 'this'/'par' is a pointer/reference to */
/* object in generic address space */
/* if supported */
/* if supported, or */
/* in private address space otherwise. */

// C(C && par); /* 'this'/'par' is a pointer/r-val reference to */
/* object in generic address space if supported */
/* object in generic address space if supported, or */
/* in private address space otherwise. */

// C & operator=(const C & par); /* 'this'/'par'/return value is */
/* a pointer/reference/reference to */
/* object in generic address space */
/* if supported */
/* if supported, or */
/* in private address space otherwise. */

// C & operator=(C && par)'; /* 'this'/'par'/return value is */
/* a pointer/r-val reference/reference to */
/* object in generic address space, */
/* if it is supported or */
/* if supported, or */
/* in private address space otherwise. */
};
------------

==== Builtin operators

All builtin operators are available in the specific named address spaces, thus
All builtin operators are available with the specific address spaces, thus
no address space conversions (i.e. to generic address space) are performed.

==== Templates

There is no deduction of address spaces in non-pointer/non-reference
template parameters and dependent types (see <<addrspace-deduction,
_Deduction_>>). The address space of a template parameter is deduced
_Address space inference_>>). The address space of a template parameter is deduced
during type deduction if it is not explicitly provided in the
instantiation.

Expand Down Expand Up @@ -385,7 +389,7 @@ void bar() {
------------

Once a template has been instantiated, regular restrictions for
address spaces will apply.
address spaces will apply as described in `OpenCL C 3.0 s6.7`.

[source,cpp]
------------
Expand Down Expand Up @@ -418,10 +422,11 @@ void foo() {
// 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).
void f(__global float &ref) {
__global const int& newref = ref; // error: address space mismatch between
// temporary object created to hold value
// converted float->int and local variable
// (can't convert from __private to __global).
}
------------

Expand Down Expand Up @@ -458,8 +463,8 @@ public:
};
__kernel void foo() {
__local C locobj{}; // error: local address space objects can't be initialized
__local C locobj; // uninitialised object
locobj = {}; // calling copy assignment operator is allowed
__local C locobj; // uninitialised object.
locobj = {}; // calling copy assignment operator is allowed.
locobj.~C(); // local address space object destructors are not invoked
// automatically.
}
Expand All @@ -474,7 +479,7 @@ Objects in `+__constant+` address space can be initialized using:
* Literal expressions;
* Uniform initialization syntax `{}`;
* Using implicit constructors.
* Using constexpr constructors.
* Using `constexpr` constructors.

[source,cpp]
------------
Expand All @@ -487,44 +492,67 @@ struct C2 {
constexpr C2(int init) __constant : m(init) {};
};

__constant C1 cobj1 = {1};
__constant C1 cobj2 = C1();
__constant C2 cobj3(1);
__constant C1 c1obj1 = {1};
__constant C1 c1obj2 = C1();
__constant C2 c2obj1(1);
------------

Non-trivial destructors for objects in non-default address spaces (i.e. all
other than generic address space) are not required to be supported by
implementations. The macro `+__opencl_cpp_destructor_with_address_spaces+`,
other than generic address space when it is supported or `+__private+`
otherwise) are not required to be supported by implementations.
The macro `+__opencl_cpp_destructor_with_address_spaces+`,
which is defined if and only if such destructors are supported by an
implementation, can be used to check whether this functionality can be used
in kernel sources. Additionally destructors with global objects might not be
supported even if address spaces are supported with destructors in general.
Such functionality is indicated by the presence of the
`+__opencl_cpp_global_destructor+` macro. If the macro
`+__opencl_cpp_global_destructor+` is defined then
`+__opencl_cpp_destructor_with_address_spaces+` must also be defined. Note that
the destruction of objects in named address spaces global, local, or private
can be performed using destructors with default address space (i.e. generic)
by utilising address space conversions.
`+__opencl_cpp_destructor_with_address_spaces+` must also be defined.

Note that the destruction of objects in named address spaces `+__global+`,
`+__local+`, or `+__private+` can be performed using destructors with
default address space (i.e. generic) by utilising address space conversions.

[source,cpp]
------------
1 class C {
2 public:
3 #ifdef __opencl_cpp_destructor_with_address_spaces
4 ~C() __local;
5 #else
6 ~C();
7 #endif
8 };
9
10 kernel void foo() {
11 __local C locobj;
12 locobj.~C(); // uses destructor in local address space (on line 4)
13 // if such destructors are supported,
14 // otherwise uses generic address space destructor (on line 6)
15 // converting to generic address prior to call into destructor.
16 }
1 // Example assumes generic address space support.
2 class C {
3 public:
4 #ifdef __opencl_cpp_destructor_with_address_spaces
5 ~C() __local;
6 #else
7 ~C();
8 #endif
9 };
10
11 kernel void foo() {
12 __local C locobj;
13 locobj.~C(); // uses destructor in local address space (on line 5)
14 // if such destructors are supported,
15 // otherwise uses generic address space destructor (on line 7)
16 // converting to generic address prior to call into destructor.
17 }
------------

However, when generic address space feature is unsupported, absence of
destructor support with address spaces results in a compilation failure
when such destructors overloaded with non-default address spaces are
encountered in the kernel code.

[source,cpp]
------------
// Example assumes generic address space is not supported.
class C {
public:
~C();
};

kernel void foo() {
__local C locobj;
locobj.~C(); // error due to illegal conversion of 'this' from __local
// to __private address space pointer.
}
------------

==== Nested pointers
Expand Down Expand Up @@ -553,29 +581,30 @@ defdefptr = reinterpret_cast<int * *>(cnstdefptr); // legal.
------------

[[remove-addrspace]]
==== Address space removal
==== Address space removal type trait

`+template<class T> struct __remove_address_space;+`

{cpp} for OpenCL 2021 supports utility `+__remove_address_space+` which allows to
derive types that have address space qualifiers removed. Its effect is analogous
to `+__remove_const+` and other similar traits in {cpp}17 `[meta.trans]`. The
utility only removes an address space qualifier from a given type, therefore, other
type qualifiers such as `+const+` or `+volatile+` remain unchanged.
{cpp} for OpenCL 2021 supports the type trait `+__remove_address_space+`
that provides the member `typedef type` which is the same as `T`, except
that its topmost address space qualifier is removed. Its effect is analogous
to `+remove_const+` and other similar type traits in {cpp}17 `[meta.trans]`.
The trait only removes an address space qualifier from a given type, therefore,
all other type qualifiers such as `+const+` or `+volatile+` remain unchanged.

[source,cpp]
----------
------------
template<typename T>
void foo(T *par) {
T var1; // error: function scope variable cannot be declared in global
// address space
// address space.
__private T var2; // error: conflicting address space qualifiers are provided
// between types '__private T' and '__global int'
__private __remove_address_space<T>::type var3; // type of var3 is __private int
// between types '__private T' and '__global int'.
__private __remove_address_space<T>::type var3; // type of var3 is __private int.
}

void bar() {
__global int* ptr;
foo(ptr);
}
----------
------------
30 changes: 16 additions & 14 deletions cxx4opencl/cxxcasts.txt
View file
Original file line number Diff line number Diff line change
Expand Up @@ -5,9 +5,10 @@
[[cxxcasts]]
=== {cpp} casts

{cpp} has 3 casts in addition to C-style casts. `static_cast` and `const_cast`
function the same way as in {cpp}, but `reinterpret_cast` has some additional
functionality:
{cpp} has three cast operators in addition to C-style casts. Additional
logic specific to address spaces are applied to all casts as detailed in
<<address_space_casts, _conversions with address spaces_>>.
`reinterpret_cast` has some additional functionality:

* Conversion between vectors and scalars are allowed.
* Conversion between OpenCL types are disallowed.
Expand All @@ -20,26 +21,27 @@ pointers. In {cpp} for openCL this also includes vector types, and so using
size of the vectors are the same.

[source,cpp]
----------
------------
int i;
short2 s2 = reinterpret_cast<short2>(i); // legal.
int2 i2 = reinterpret_cast<int2>(i); // illegal.

short8 s8;
int4 i4 = reinterpret_cast<int4>(s8); // legal.
long l4 = reinterpret_cast<long>(s8); // illegal.
----------
------------

==== OpenCL types

Some of the OpenCL types are the same size as integers, or can be implemented as
integers, but since they are not conceptually integral, they can not be used
with `reinterpret_cast`. Therefore these are all illegal conversions:
Some of the OpenCL-specific types, defined as "Other Built-in Data Types" in
`OpenCL C 3.0 s6.3.3`, are convertible to integer literals, but since they
are not conceptually integral, they can not be used with `reinterpret_cast`.
Therefore conversions of an OpenCL-specific type to any distinct type are
illegal.

[source,cpp]
----------
reserve_id_t id;
reserve_id_t id2 = reinterpret_cast<reserve_id_t>(id); // illegal.
int i = reinterpret_cast<int>(r); // illegal.
long l = reinterpret_cast<long>(r); // illegal.
----------
------------
queue_t q;
reserve_id_t id = reinterpret_cast<reserve_id_t>(q); // illegal.
int i = reinterpret_cast<int>(id); // illegal.
------------
2 changes: 1 addition & 1 deletion cxx4opencl/diff2cxx.txt
View file
Original file line number Diff line number Diff line change
Expand Up @@ -35,4 +35,4 @@ macro{fn-feature-macro}.
The list above only contains extra restrictions that are not detailed in OpenCL
C specification. As OpenCL restricts a number of C features, the same restrictions
are inherited by C++ for OpenCL. The detailed list of C feature restrictions
is provided in `OpenCL C 3.0 section 6.11`.
is provided in `OpenCL C 3.0 s6.11`.
Loading

0 comments on commit 211d598

Please sign in to comment.