[OpenCL] Add generic AS to 'this' pointer

Address spaces are cast into generic before invoking the constructor.

Added support for a trailing Qualifiers object in FunctionProtoType.

Note: This recommits the previously reverted patch, 
      but now it is commited together with a fix for lldb.

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

llvm-svn: 349019

3 files changed, 166 insertions, 7 deletions

diff --git a/clang/test/CodeGenOpenCLCXX/addrspace-of-this.cl b/clang/test/CodeGenOpenCLCXX/addrspace-of-this.cl
new file mode 100644
index 00000000000..af0e3b1a68c
--- /dev/null
+++ b/clang/test/CodeGenOpenCLCXX/addrspace-of-this.cl
@@ -0,0 +1,154 @@
+// 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; }
+  // FIXME: Does not work yet.
+  // 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;
+  //}
+
+  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;
+  // FIXME: Does not work yet.
+  // 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
+
+// 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) #4
+// 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: %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)*))
+
+// 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
+
+// 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)
+
+#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 method.
+// 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.
+// 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 a constructor.
+// CHECK: %3 = addrspacecast %class.C* %c2 to %class.C addrspace(4)*
+// CHECK: call void @_ZNU3AS41CC1Ev(%class.C addrspace(4)* %3)
+
+// 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)
+
+TEST(__private)
+
+// CHECK-LABEL: @_Z13test__privatev
+
+// 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)
+
+// 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)
+
+// 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)
+
+// 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)
+
+TEST()
+
+// CHECK-LABEL: @_Z4testv()
+// 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
+
+// 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)
+
+// 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)
+
+// 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)
diff --git a/clang/test/CodeGenOpenCLCXX/template-address-spaces.cl b/clang/test/CodeGenOpenCLCXX/template-address-spaces.cl
index eb274ea4560..7c722537fca 100644
--- a/clang/test/CodeGenOpenCLCXX/template-address-spaces.cl
+++ b/clang/test/CodeGenOpenCLCXX/template-address-spaces.cl
@@ -9,13 +9,16 @@ struct S{
 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: %struct.S = type { i32 }
+// CHECK: %struct.S.0 = type { i32 addrspace(4)* }
+// CHECK: %struct.S.1 = type { i32 addrspace(1)* }
 
-//CHECK: i32 @_ZN1SIiE3fooEv(%struct.S* %this)
-//CHECK: i32 addrspace(4)* @_ZN1SIPU3AS4iE3fooEv(%struct.S.0* %this)
-//CHECK: i32 addrspace(1)* @_ZN1SIPU3AS1iE3fooEv(%struct.S.1* %this)
+// 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;
diff --git a/clang/test/SemaOpenCLCXX/address-space-templates.cl b/clang/test/SemaOpenCLCXX/address-space-templates.cl
index 48fbdc7642d..80762fce8a9 100644
--- a/clang/test/SemaOpenCLCXX/address-space-templates.cl
+++ b/clang/test/SemaOpenCLCXX/address-space-templates.cl
@@ -4,7 +4,9 @@ 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}}
-  void f2(T); // expected-error{{parameter 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}}
+
 };
 
 template <typename T>