1 files changed, 24 insertions, 7 deletions

diff --git a/test/CodeGenCUDA/address-spaces.cu b/test/CodeGenCUDA/address-spaces.cu
index b808206..b0ef355 100644
--- a/test/CodeGenCUDA/address-spaces.cu
+++ b/test/CodeGenCUDA/address-spaces.cu
@@ -27,25 +27,25 @@ struct MyStruct {
 // CHECK: @b = addrspace(3) global float 0.000000e+00
 
 __device__ void foo() {
-  // CHECK: load i32* addrspacecast (i32 addrspace(1)* @i to i32*)
+  // CHECK: load i32, i32* addrspacecast (i32 addrspace(1)* @i to i32*)
   i++;
 
-  // CHECK: load i32* addrspacecast (i32 addrspace(4)* @j to i32*)
+  // CHECK: load i32, i32* addrspacecast (i32 addrspace(4)* @j to i32*)
   j++;
 
-  // CHECK: load i32* addrspacecast (i32 addrspace(3)* @k to i32*)
+  // CHECK: load i32, i32* addrspacecast (i32 addrspace(3)* @k to i32*)
   k++;
 
   static int li;
-  // CHECK: load i32* addrspacecast (i32 addrspace(1)* @_ZZ3foovE2li to i32*)
+  // CHECK: load i32, i32* addrspacecast (i32 addrspace(1)* @_ZZ3foovE2li to i32*)
   li++;
 
   __constant__ int lj;
-  // CHECK: load i32* addrspacecast (i32 addrspace(4)* @_ZZ3foovE2lj to i32*)
+  // CHECK: load i32, i32* addrspacecast (i32 addrspace(4)* @_ZZ3foovE2lj to i32*)
   lj++;
 
   __shared__ int lk;
-  // CHECK: load i32* addrspacecast (i32 addrspace(3)* @_ZZ3foovE2lk to i32*)
+  // CHECK: load i32, i32* addrspacecast (i32 addrspace(3)* @_ZZ3foovE2lk to i32*)
   lk++;
 }
 
@@ -75,7 +75,7 @@ __device__ void func2() {
   *ap = 1.0f;
 }
 // CHECK: define void @_Z5func2v()
-// CHECK: store float* getelementptr inbounds ([256 x float]* addrspacecast ([256 x float] addrspace(3)* @_ZZ5func2vE1a to [256 x float]*), i32 0, i32 128), float** %ap
+// CHECK: store float* getelementptr inbounds ([256 x float], [256 x float]* addrspacecast ([256 x float] addrspace(3)* @_ZZ5func2vE1a to [256 x float]*), i32 0, i32 128), float** %ap
 
 __device__ void func3() {
   __shared__ float a;
@@ -100,3 +100,20 @@ __device__ float *func5() {
 }
 // CHECK: define float* @_Z5func5v()
 // CHECK: ret float* addrspacecast (float addrspace(3)* @b to float*)
+
+struct StructWithCtor {
+  __device__ StructWithCtor(): data(1) {}
+  __device__ StructWithCtor(const StructWithCtor &second): data(second.data) {}
+  __device__ int getData() { return data; }
+  int data;
+};
+
+__device__ int construct_shared_struct() {
+// CHECK-LABEL: define i32 @_Z23construct_shared_structv()
+  __shared__ StructWithCtor s;
+// CHECK: call void @_ZN14StructWithCtorC1Ev(%struct.StructWithCtor* addrspacecast (%struct.StructWithCtor addrspace(3)* @_ZZ23construct_shared_structvE1s to %struct.StructWithCtor*))
+  __shared__ StructWithCtor t(s);
+// CHECK: call void @_ZN14StructWithCtorC1ERKS_(%struct.StructWithCtor* addrspacecast (%struct.StructWithCtor addrspace(3)* @_ZZ23construct_shared_structvE1t to %struct.StructWithCtor*), %struct.StructWithCtor* dereferenceable(4) addrspacecast (%struct.StructWithCtor addrspace(3)* @_ZZ23construct_shared_structvE1s to %struct.StructWithCtor*))
+  return t.getData();
+// CHECK: call i32 @_ZN14StructWithCtor7getDataEv(%struct.StructWithCtor* addrspacecast (%struct.StructWithCtor addrspace(3)* @_ZZ23construct_shared_structvE1t to %struct.StructWithCtor*))
+}