diff options
Diffstat (limited to 'test/CodeGenCUDA/address-spaces.cu')
| -rw-r--r-- | test/CodeGenCUDA/address-spaces.cu | 31 |
1 files changed, 24 insertions, 7 deletions
diff --git a/test/CodeGenCUDA/address-spaces.cu b/test/CodeGenCUDA/address-spaces.cu index b80820683f3e..b0ef3558e2d0 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*)) +} |