[OpenCL] Set generic addr space of 'this' in special class members.

Set address spaces of 'this' param correctly for implicit special
class members.

This also changes initialization conversion sequence to separate
address space conversion from other qualifiers in case of binding
reference to a temporary. In this case address space conversion
should happen after the binding (unlike for other quals). This is
needed to materialize it correctly in the alloca address space.

Initial patch by Mikael Nilssoni!

Differential Revision: https://reviews.llvm.org/D56066

2 files changed, 83 insertions, 51 deletions

diff --git a/clang/test/CodeGenOpenCLCXX/addrspace-of-this.cl b/clang/test/CodeGenOpenCLCXX/addrspace-of-this.cl
index af0e3b1a68c..83af3fbcb3b 100644
--- a/clang/test/CodeGenOpenCLCXX/addrspace-of-this.cl
+++ b/clang/test/CodeGenOpenCLCXX/addrspace-of-this.cl
@@ -9,18 +9,21 @@ class C {
 public:
   int v;
   C() { v = 2; }
-  // FIXME: Does not work yet.
-  // C(C &&c) { v = c.v; }
+  C(C &&c) { v = c.v; }
   C(const C &c) { v = c.v; }
   C &operator=(const C &c) {
     v = c.v;
     return *this;
   }
-  // FIXME: Does not work yet.
-  //C &operator=(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; }
 
@@ -41,15 +44,13 @@ __kernel void test__global() {
   C c1(c);
   C c2;
   c2 = c1;
-  // FIXME: Does not work yet.
-  // C c3 = c1 + c2;
-  // C c4(foo());
-  // C c5 = foo();
-
+  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)*)) #4
+// 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)
@@ -57,7 +58,7 @@ __kernel void test__global() {
 // 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) #4
+// CHECK:   call void @_ZNU3AS41CC2Ev(%class.C addrspace(4)* %this1)
 // CHECK:   ret void
 
 // CHECK-LABEL: @_Z12test__globalv()
@@ -69,17 +70,36 @@ __kernel void test__global() {
 // 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: %0 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
-// CHECK: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* %0, %class.C addrspace(4)* dereferenceable(4) addrspacecast (%class.C addrspace(1)* @c to %class.C addrspace(4)*))
+// 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:   %1 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// CHECK:   call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* %1) #4
+// 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:   %2 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
-// CHECK:   %3 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// CHECK:   %call2 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* %3, %class.C addrspace(4)* dereferenceable(4) %2)
+// 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() { \
@@ -95,60 +115,74 @@ TEST(__local)
 // CHECK-LABEL: _Z11test__localv
 // CHECK: @__cxa_guard_acquire
 
-// Test the address space of 'this' when invoking a method.
+// 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 copy-constructor.
+// 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: %3 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* %3)
+// 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:  %4 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
-// CHECK:  %5 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// CHECK:  %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* %5, %class.C addrspace(4)* dereferenceable(4) %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: %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:   %1 = addrspacecast %class.C* %c to %class.C addrspace(4)*
-// CHECK:   %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* %1)
+// 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: %2 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
-// CHECK: %3 = addrspacecast %class.C* %c to %class.C addrspace(4)*
-// CHECK: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* %2, %class.C addrspace(4)* dereferenceable(4) %3)
+// 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: %4 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// CHECK:   call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* %4)
+// 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:   %5 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
-// CHECK:   %6 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// CHECK:   %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* %6, %class.C addrspace(4)* dereferenceable(4) %5)
+// 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: %1 = addrspacecast %class.C* %c to %class.C addrspace(4)*
-// CHECK: %call = call i32 @_ZNU3AS41C3getEv(%class.C addrspace(4)* %1) #4
+// 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: %2 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
-// CHECK: %3 = addrspacecast %class.C* %c to %class.C addrspace(4)*
-// CHECK: call void @_ZNU3AS41CC1ERU3AS4KS_(%class.C addrspace(4)* %2, %class.C addrspace(4)* dereferenceable(4) %3)
+// 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: %4 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* %4)
+// 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: %5 = addrspacecast %class.C* %c1 to %class.C addrspace(4)*
-// CHECK: %6 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
-// CHECK: %call1 = call dereferenceable(4) %class.C addrspace(4)* @_ZNU3AS41CaSERU3AS4KS_(%class.C addrspace(4)* %6, %class.C addrspace(4)* dereferenceable(4) %5)
+// 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/clang/test/SemaOpenCLCXX/address-space-templates.cl b/clang/test/SemaOpenCLCXX/address-space-templates.cl
index 80762fce8a9..48fbdc7642d 100644
--- a/clang/test/SemaOpenCLCXX/address-space-templates.cl
+++ b/clang/test/SemaOpenCLCXX/address-space-templates.cl
@@ -4,9 +4,7 @@ template <typename T>
 struct S {
   T a;        // expected-error{{field may not be qualified with an address space}}
   T f1();     // expected-error{{function type may not be qualified with an address space}}
-  // FIXME: Should only get the error message once.
-  void f2(T); // expected-error{{parameter may not be qualified with an address space}} expected-error{{parameter may not be qualified with an address space}}
-
+  void f2(T); // expected-error{{parameter may not be qualified with an address space}}
 };
 
 template <typename T>