llvm/clang/test/SemaOpenCLCXX/address-space-deduction.clcpp

//RUN: %clang_cc1 %s -pedantic -ast-dump -verify | FileCheck %s

//expected-no-diagnostics

//CHECK: |-VarDecl {{.*}} foo 'const __global int'
constexpr int foo = 0;

//CHECK: |-VarDecl {{.*}} foo1 'T' cinit
//CHECK: `-VarTemplateSpecializationDecl {{.*}} used foo1 '__global long' implicit_instantiation cinit
template <typename T>
T foo1 = 0;

class c {
public:
  //CHECK: `-VarDecl {{.*}} foo2 'const __global int'
  static constexpr int foo2 = 0;
};

struct c1 {};

// We only deduce addr space in type alias in pointer types.
//CHECK: TypeAliasDecl {{.*}} alias_c1 'c1'
using alias_c1 = c1;
//CHECK: TypeAliasDecl {{.*}} alias_c1_ptr '__generic c1 *'
using alias_c1_ptr = c1 *;

struct c2 {
  alias_c1 y;
  alias_c1_ptr ptr = &y;
};


// Addr spaces for pointee of dependent types are not deduced
// during parsing but during template instantiation instead.

template <class T>
struct x1 {
//CHECK: -CXXMethodDecl {{.*}} operator= 'x1<T> &(const x1<T> &){{( __attribute__.*)?}} __generic'
//CHECK: -CXXMethodDecl {{.*}} operator= '__generic x1<int> &(const __generic x1<int> &__private){{( __attribute__.*)?}} __generic'
  x1<T>& operator=(const x1<T>& xx) {
    y = xx.y;
    return *this;
  }
  int y;
};

template <class T>
struct x2 {
//CHECK: -CXXMethodDecl {{.*}} foo 'void (x1<T> *){{( __attribute__.*)?}} __generic'
//CHECK: -CXXMethodDecl {{.*}} foo 'void (__generic x1<int> *__private){{( __attribute__.*)?}} __generic'
  void foo(x1<T>* xx) {
    m[0] = *xx;
  }
//CHECK: -FieldDecl {{.*}}  m 'x1<int>[2]'
  x1<T> m[2];
};

void bar(__global x1<int> *xx, __global x2<int> *bar) {
  bar->foo(xx);
}

template <typename T>
class x3 : public T {
public:
  //CHECK: -CXXConstructorDecl {{.*}} x3<T> 'void (const x3<T> &){{( __attribute__.*)?}} __generic'
  x3(const x3 &t);
};
//CHECK: -CXXConstructorDecl {{.*}} x3<T> 'void (const x3<T> &){{( __attribute__.*)?}} __generic'
template <typename T>
x3<T>::x3(const x3<T> &t) {}

template <class T>
T xxx(T *in1, T in2) {
  // This pointer can't be deduced to generic because addr space
  // will be taken from the template argument.
  //CHECK: `-VarDecl {{.*}} 'T *' cinit
  //CHECK: `-VarDecl {{.*}} i '__private int *__private' cinit
  T *i = in1;
  T ii;
  __private T *ptr = &ii;
  ptr = &in2;
  return *i;
}

__kernel void test() {
  int foo[10];
  xxx<__private int>(&foo[0], foo[0]);
  // FIXME: Template param deduction fails here because
  // temporaries are not in the __private address space.
  // It is probably reasonable to put them in __private
  // considering that stack and function params are
  // implicitly in __private.
  // However, if temporaries are left in default addr
  // space we should at least pretty print the __private
  // addr space. Otherwise diagnostic apprears to be
  // confusing.
  //xxx(&foo[0], foo[0]);
}

// Addr space for pointer/reference to an array
//CHECK: FunctionDecl {{.*}} t1 'void (const float (__generic &__private)[2])'
void t1(const float (&fYZ)[2]);
//CHECK: FunctionDecl {{.*}} t2 'void (const float (__generic *__private)[2])'
void t2(const float (*fYZ)[2]);
//CHECK: FunctionDecl {{.*}} t3 'void (float (((__generic *__private)))[2])'
void t3(float(((*fYZ)))[2]);
//CHECK: FunctionDecl {{.*}} t4 'void (float (((__generic *__generic *__private)))[2])'
void t4(float(((**fYZ)))[2]);
//CHECK: FunctionDecl {{.*}} t5 'void (float (__generic *(__generic *__private))[2])'
void t5(float (*(*fYZ))[2]);

__kernel void k() {
  __local float x[2];
  float(*p)[2];
  t1(x);
  t2(&x);
  t3(&x);
  t4(&p);
  t5(&p);
  long f1 = foo1<long>;
}