The C++ for OpenCL 1.0 and 2021 Programming Language Documentation

// Example assumes generic address space support.
int * genptr; // points to generic address space.

// generic -> named address space conversions.
__private float * ptrfloat = reinterpret_cast<__private float*>(genptr); // illegal.
__private float * ptrfloat = addrspace_cast<__private float*>(genptr); // illegal.
__private int * ptr = addrspace_cast<__private int*>(genptr); // legal.

// named -> generic address space conversion.
float * genptrfloat = reinterpret_cast<float*>(ptr); // legal.

// disjoint address space conversion.
__constant int * constptr = addrspace_cast<__constant int*>(genptr); // illegal.

// Example without generic address space support.
int * privptr; // points to private address space.

// The same address space conversions.
__private float * ptrfloat = reinterpret_cast<__private float*>(privptr); // legal.
__private float * ptrfloat = addrspace_cast<__private float*>(privptr); // illegal.
__private int * ptr = addrspace_cast<__private int*>(privptr); // legal, no op.
float * privptrfloat = reinterpret_cast<float*>(ptr); // legal.

// disjoint address space conversion.
__constant int * constptr = addrspace_cast<__constant int*>(privptr); // illegal.

__private int & ref = ...; // references int in __private address space.

int & ref = ...; // references int in generic address space if it is
                 // supported otherwise in __private address space.

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.

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

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 when it is
           // supported otherwise to __private address space.
  T & ref = ...; // 'ref' references an object in generic address space when
                 // it is supported otherwise in __private address space.
};

template <int N>
struct S {
  int i; // 'i' has no address space.
  static int ii; // 'ii' is in global address space if program scope variables
                 // are supported; otherwise this statement is not legal.
  int * ptr; // 'ptr' points to int in generic address space if it is supported;
             // otherwise to __private address space.
  int & ref = ...; // 'ref' references int in generic address space if it is
                   // supported; otherwise in __private address space.
};

template <int N>
void bar()
{
  S<N> s; // 's' is in __private address space.
}

struct c1 {};
using alias_c1 = c1; // 'alias_c1' is 'c1'.
using alias_c1_ptr = c1 *; // 'alias_c1_ptr' is a generic address space pointer to
                           // 'c1' when generic address space is supported; otherwise
                           // it points to 'c1' located in __private address space.

__kernel void foo()
{
  __local int i;
  decltype(i)* ii; // type of 'ii' is '__local int *__private'.
}

// This example assumes that generic address space is supported.
__kernel void foo()
{
  __local int i;
  constexpr int c = 1;

  __constant auto cai = c; // type of 'cai' is '__constant int' (no deduction).

  auto aii = cai; // type of 'aii' is '__private int' (regular deduction).

  auto *ptr = &i; // type of 'ptr' is '__local int * __private'
                  // (addr space of a pointer is deduced regularly,
                  // addr space of its pointee is taken from 'i').

  auto *&refptr = ptr; // type of 'refptr' is '__local int * generic & __private'
                       // (addr space of a reference and type of referencing object
                       // is deduced regularly,
                       // addr space of a pointee is taken from the pointee of 'ptr').
}

struct C {
  C() __local {};
  C() __private {};
  constexpr C() __constant {};

  void foo() __local;
  void foo(); // This is implicitly qualified by generic address space
              // if it is supported otherwise by '__private'.
};

__kernel void bar() {
  __local C c1;      // will resolve to the first constructor overload.
  __private C c2;    // will resolve to the second constructor overload.
  __constant C c3{}; // will resolve to the third constructor overload.
  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/'__private' address spaces.
}

__kernel void foo() {
  auto priv1 = []() __private {};
  priv1();
  auto priv2 = []() __global {};
  priv2(); // error: lambda object and its expression have mismatching address space.
  __constant auto const3 = []() __constant{};
  const3();

  [&] () __global {} (); // error: lambda temporary is in __private address space.

  [&] () mutable __private {} ();
  [&] () __private mutable {} (); // error: mutable specifier should precede address
                                  // space.
}

class C {
  // Has the following implicitly defined member functions.

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

  // C(const C & par); /* 'this'/'par' is a pointer/reference to */
                       /*  object in generic address space */
                       /*  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, 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, 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 supported, or */
                               /* in private address space otherwise. */
};

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 'var'
9             // appears to be declared in __global address space (see line 3).
10 }

template <typename T>
void foo() {
  __private T var;
}

void bar() {
  foo<__global int>(); // error: conflicting address space qualifiers are provided
                       // for 'var', '__global' and '__private'.
}

template<typename T>
void foo(){
  T var;
}

void bar(){
  foo<__global int>(); // error: function scope variable 'var' cannot be declared
                       // in '__global' address space.
}

int bar(const unsigned int &i); // references generic address space object
                                // if generic address space is supported
                                // otherwise private address space object.

void foo() {
  bar(1); // temporary is created in __private address space but (if generic
          // address space is supported) converted to generic address space
          // of parameter reference.
}

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).
}

class C {
  int m;
// If generic address space is not supported or for performance optimization
// purposes the following members might be required.
public:
  __local C & operator=(const C & par) __local;
  ~C() __local;
};
__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.
  locobj.~C(); // local address space object destructors are not invoked
               // automatically.
}

struct C1 {
  int m;
};

struct C2 {
  int m;
  constexpr C2(int init) __constant : m(init) {};
};

__constant C1 c1obj1 = {1};
__constant C1 c1obj2 = C1();
__constant C2 c2obj1(1);

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  }

// 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.
}

local int * * locdefptr;
constant int * * cnstdefptr;
int * * defdefptr;
defdefptr = const_cast<int * *>(locdefptr); // illegal.
defdefptr = static_cast<int * *>(cnstdefptr); // illegal.
defdefptr = addrspace_cast<int * *>(cnstdefptr); // illegal.
defdefptr = reinterpret_cast<int * *>(locdefptr); // legal.
defdefptr = reinterpret_cast<int * *>(cnstdefptr); // legal.

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

void bar() {
  __global int* ptr;
  foo(ptr);
}

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.

queue_t q;
reserve_id_t id = reinterpret_cast<reserve_id_t>(q); // illegal.
int i = reinterpret_cast<int>(id); // illegal.