4 files changed, 285 insertions, 0 deletions
diff --git a/test/CodeGenOpenCLCXX/address-space-deduction.cl b/test/CodeGenOpenCLCXX/address-space-deduction.cl
new file mode 100644
index 000000000000..ecc0066d45ae
--- /dev/null
+++ b/test/CodeGenOpenCLCXX/address-space-deduction.cl
@@ -0,0 +1,46 @@
+// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -O0 -emit-llvm -o - | FileCheck %s -check-prefixes=COMMON,PTR
+// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -O0 -emit-llvm -o - -DREF | FileCheck %s -check-prefixes=COMMON,REF
+
+#ifdef REF
+#define PTR &
+#define ADR(x) x
+#else
+#define PTR *
+#define ADR(x) &x
+#endif
+
+//COMMON: @glob = addrspace(1) global i32
+int glob;
+//PTR: @glob_p = addrspace(1) global i32 addrspace(4)* addrspacecast (i32 addrspace(1)* @glob to i32 addrspace(4)*)
+//REF: @glob_p = addrspace(1) global i32 addrspace(4)* null
+int PTR glob_p = ADR(glob);
+
+//COMMON: @_ZZ3fooi{{P|R}}U3AS4iE6loc_st = internal addrspace(1) global i32
+//PTR: @_ZZ3fooiPU3AS4iE8loc_st_p = internal addrspace(1) global i32 addrspace(4)* addrspacecast (i32 addrspace(1)* @_ZZ3fooiPU3AS4iE6loc_st to i32 addrspace(4)*)
+//REF: @_ZZ3fooiRU3AS4iE8loc_st_p = internal addrspace(1) global i32 addrspace(4)* null
+//COMMON: @loc_ext_p = external addrspace(1) {{global|constant}} i32 addrspace(4)*
+//COMMON: @loc_ext = external addrspace(1) global i32
+
+//REF: store i32 addrspace(4)* addrspacecast (i32 addrspace(1)* @glob to i32 addrspace(4)*), i32 addrspace(4)* addrspace(1)* @glob_p
+
+//COMMON: define spir_func i32 @_Z3fooi{{P|R}}U3AS4i(i32 %par, i32 addrspace(4)*{{.*}} %par_p)
+int foo(int par, int PTR par_p){
+  //COMMON: %loc = alloca i32
+  int loc;
+  //COMMON: %loc_p = alloca i32 addrspace(4)*
+  //COMMON: %loc_p_const = alloca i32*
+  //COMMON: [[GAS:%[0-9]+]] = addrspacecast i32* %loc to i32 addrspace(4)*
+  //COMMON: store i32 addrspace(4)* [[GAS]], i32 addrspace(4)** %loc_p
+  int PTR loc_p = ADR(loc);
+  //COMMON: store i32* %loc, i32** %loc_p_const
+  const __private int PTR loc_p_const = ADR(loc);
+
+  // CHECK directives for the following code are located above.
+  static int loc_st;
+  //REF: store i32 addrspace(4)* addrspacecast (i32 addrspace(1)* @_ZZ3fooiRU3AS4iE6loc_st to i32 addrspace(4)*), i32 addrspace(4)* addrspace(1)* @_ZZ3fooiRU3AS4iE8loc_st_p
+  static int PTR loc_st_p = ADR(loc_st);
+  extern int loc_ext;
+  extern int PTR loc_ext_p;
+  (void)loc_ext_p;
+  return loc_ext;
+}
diff --git a/test/CodeGenOpenCLCXX/address-space-deduction2.cl b/test/CodeGenOpenCLCXX/address-space-deduction2.cl
new file mode 100644
index 000000000000..6772cdcb11d0
--- /dev/null
+++ b/test/CodeGenOpenCLCXX/address-space-deduction2.cl
@@ -0,0 +1,20 @@
+// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -O0 -emit-llvm -o - | FileCheck %s
+
+class P {
+public:
+  P(const P &Rhs) = default;
+
+  long A;
+  long B;
+};
+
+void foo(__global P *GPtr) {
+// CHECK: call void @llvm.memcpy{{.*}}, {{.*}}, i32 16
+  P Val = GPtr[0];
+}
+
+struct __attribute__((packed)) A { int X; };
+int test(__global A *GPtr) {
+// CHECK: {{.*}} = load i32, {{.*}}, align 1
+  return static_cast<__generic A &>(*GPtr).X;
+}
diff --git a/test/CodeGenOpenCLCXX/addrspace-of-this.cl b/test/CodeGenOpenCLCXX/addrspace-of-this.cl
new file mode 100644
index 000000000000..83af3fbcb3b6
--- /dev/null
+++ b/test/CodeGenOpenCLCXX/addrspace-of-this.cl
@@ -0,0 +1,188 @@
+// RUN: %clang_cc1 %s -triple spir-unknown-unknown -cl-std=c++ -emit-llvm -pedantic -verify -O0 -o - | FileCheck %s
+// expected-no-diagnostics
+
+// Test that the 'this' pointer is in the __generic address space.
+
+// FIXME: Add support for __constant address space.
+
+class C {
+public:
+  int v;
+  C() { v = 2; }
+  C(C &&c) { v = c.v; }
+  C(const C &c) { v = c.v; }
+  C &operator=(const C &c) {
+    v = c.v;
+    return *this;
+  }
+  C &operator=(C &&c) & {
+    v = c.v;
+    return *this;
+  }
+
+  C operator+(const C& c) {
+    v += c.v;
+    return *this;
+  }
+
+  int get() { return v; }
+
+  int outside();
+};
+
+int C::outside() {
+  return v;
+}
+
+extern C&& foo();
+
+__global C c;
+
+__kernel void test__global() {
+  int i = c.get();
+  int i2 = c.outside();
+  C c1(c);
+  C c2;
+  c2 = c1;
+  C c3 = c1 + c2;
+  C c4(foo());
+  C c5 = foo();
+}
+
+// CHECK-LABEL: @__cxx_global_var_init()
+// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
+
+// Test that the address space is __generic for the constructor
+// CHECK-LABEL: @_ZNU3AS41CC1Ev(%class.C addrspace(4)* %this)
+// CHECK: entry:
+// CHECK:   %this.addr = alloca %class.C addrspace(4)*, align 4
+// CHECK:   store %class.C addrspace(4)* %this, %class.C addrspace(4)** %this.addr, align 4
+// CHECK:   %this1 = load %class.C addrspace(4)*, %class.C addrspace(4)** %this.addr, align 4
+// CHECK:   call void @_ZNU3AS41CC2Ev(%class.C addrspace(4)* %this1)
+// CHECK:   ret void
+
+// CHECK-LABEL: @_Z12test__globalv()
+
+// Test the address space of 'this' when invoking a method.
+// CHECK: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
+
+// Test the address space of 'this' when invoking a method that is declared in the file contex.
+// CHECK: %call1 = call i32 @_ZNU3AS41C7outsideEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
+
+// Test the address space of 'this' when invoking copy-constructor.
+// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
+
+// Test the address space of 'this' when invoking a constructor.
+// CHECK:   [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK:   call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]])
+
+// Test the address space of 'this' when invoking assignment operator.
+// CHECK:  [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK:  [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK:  %call2 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]])
+
+// Test the address space of 'this' when invoking the operator+
+// CHECK: [[C3GEN:%[0-9]+]] = addrspacecast %class.C* %c3 to %class.C addrspace(4)*
+// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK: call void @_ZNU3AS41CplERU3AS4KS_(%class.C* sret %ref.tmp, %class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) [[C2GEN]])
+// CHECK: [[REFGEN:%[0-9]+]] = addrspacecast %class.C* %ref.tmp to %class.C addrspace(4)*
+// CHECK: call void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* [[C3GEN]], %class.C addrspace(4)* dereferenceable(4) [[REFGEN]])
+
+// Test the address space of 'this' when invoking the move constructor
+// CHECK: [[C4GEN:%[0-9]+]] = addrspacecast %class.C* %c4 to %class.C addrspace(4)*
+// CHECK: %call3 = call spir_func dereferenceable(4) %class.C addrspace(4)* @_Z3foov()
+// CHECK: call void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* [[C4GEN]], %class.C addrspace(4)* dereferenceable(4) %call3)
+
+// Test the address space of 'this' when invoking the move assignment
+// CHECK: [[C5GEN:%[0-9]+]] = addrspacecast %class.C* %c5 to %class.C addrspace(4)*
+// CHECK: %call4 = call spir_func dereferenceable(4) %class.C addrspace(4)* @_Z3foov() #5
+// CHECK: call void @_ZNU3AS41CC1EOU3AS4S_(%class.C addrspace(4)* [[C5GEN:%[0-9]+]], %class.C addrspace(4)* dereferenceable(4) %call4)
+
+
+#define TEST(AS)             \
+  __kernel void test##AS() { \
+    AS C c;                  \
+    int i = c.get();         \
+    C c1(c);                 \
+    C c2;                    \
+    c2 = c1;                 \
+  }
+
+TEST(__local)
+
+// CHECK-LABEL: _Z11test__localv
+// CHECK: @__cxa_guard_acquire
+
+// Test the address space of 'this' when invoking a constructor for an object in non-default address space
+// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*))
+
+// Test the address space of 'this' when invoking a method.
+// CHECK: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*))
+
+
+// Test the address space of 'this' when invoking copy-constructor.
+// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) addrspacecast (%class.C addrspace(3)* @_ZZ11test__localvE1c to %class.C addrspace(4)*))
+
+// Test the address space of 'this' when invoking a constructor.
+// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]])
+
+// Test the address space of 'this' when invoking assignment operator.
+// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]])
+
+TEST(__private)
+
+// CHECK-LABEL: @_Z13test__privatev
+
+// Test the address space of 'this' when invoking a constructor for an object in non-default address space
+// CHECK: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[CGEN]])
+
+// Test the address space of 'this' when invoking a method.
+// CHECK: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// CHECK: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* [[CGEN]])
+
+// Test the address space of 'this' when invoking a copy-constructor.
+// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// CHECK: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) [[CGEN]])
+
+// Test the address space of 'this' when invoking a constructor.
+// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]])
+
+// Test the address space of 'this' when invoking a copy-assignment.
+// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]])
+
+TEST()
+
+// CHECK-LABEL: @_Z4testv()
+
+// Test the address space of 'this' when invoking a constructor for an object in non-default address space
+// CHECK: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[CGEN]])
+
+// Test the address space of 'this' when invoking a method.
+// CHECK: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// CHECK: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* [[CGEN]])
+
+// Test the address space of 'this' when invoking a copy-constructor.
+// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK: [[CGEN:%[0-9]+]] = addrspacecast %class.C* %c to %class.C addrspace(4)*
+// CHECK: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* [[C1GEN]], %class.C addrspace(4)* dereferenceable(4) [[CGEN]])
+
+// Test the address space of 'this' when invoking a constructor.
+// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* [[C2GEN]])
+
+// Test the address space of 'this' when invoking a copy-assignment.
+// CHECK: [[C1GEN:%[0-9]+]] = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
+// CHECK: [[C2GEN:%[0-9]+]] = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* [[C2GEN]], %class.C addrspace(4)* dereferenceable(4) [[C1GEN]])
diff --git a/test/CodeGenOpenCLCXX/template-address-spaces.cl b/test/CodeGenOpenCLCXX/template-address-spaces.cl
new file mode 100644
index 000000000000..7c722537fca3
--- /dev/null
+++ b/test/CodeGenOpenCLCXX/template-address-spaces.cl
@@ -0,0 +1,31 @@
+// RUN: %clang_cc1 -cl-std=c++ %s -emit-llvm -o - -O0 -triple spir-unknown-unknown | FileCheck %s
+
+template <typename T>
+struct S{
+  T a;
+  T foo();
+};
+
+template<typename T>
+T S<T>::foo() { return a;}
+
+// CHECK: %struct.S = type { i32 }
+// CHECK: %struct.S.0 = type { i32 addrspace(4)* }
+// CHECK: %struct.S.1 = type { i32 addrspace(1)* }
+
+// CHECK:  %0 = addrspacecast %struct.S* %sint to %struct.S addrspace(4)*
+// CHECK:  %call = call i32 @_ZNU3AS41SIiE3fooEv(%struct.S addrspace(4)* %0) #1
+// CHECK:  %1 = addrspacecast %struct.S.0* %sintptr to %struct.S.0 addrspace(4)*
+// CHECK:  %call1 = call i32 addrspace(4)* @_ZNU3AS41SIPU3AS4iE3fooEv(%struct.S.0 addrspace(4)* %1) #1
+// CHECK:  %2 = addrspacecast %struct.S.1* %sintptrgl to %struct.S.1 addrspace(4)*
+// CHECK:  %call2 = call i32 addrspace(1)* @_ZNU3AS41SIPU3AS1iE3fooEv(%struct.S.1 addrspace(4)* %2) #1
+
+void bar(){
+  S<int> sint;
+  S<int*> sintptr;
+  S<__global int*> sintptrgl;
+
+  sint.foo();
+  sintptr.foo();
+  sintptrgl.foo();
+}