1 files changed, 70 insertions, 4 deletions

diff --git a/test/CodeGenCUDA/address-spaces.cu b/test/CodeGenCUDA/address-spaces.cu
index 04344526f40e..b80820683f3e 100644
--- a/test/CodeGenCUDA/address-spaces.cu
+++ b/test/CodeGenCUDA/address-spaces.cu
@@ -1,6 +1,9 @@
 // RUN: %clang_cc1 -emit-llvm %s -o - -fcuda-is-device -triple nvptx-unknown-unknown | FileCheck %s
 
-#include "../SemaCUDA/cuda.h"
+// Verifies Clang emits correct address spaces and addrspacecast instructions
+// for CUDA code.
+
+#include "Inputs/cuda.h"
 
 // CHECK: @i = addrspace(1) global
 __device__ int i;
@@ -11,6 +14,18 @@ __constant__ int j;
 // CHECK: @k = addrspace(3) global
 __shared__ int k;
 
+struct MyStruct {
+  int data1;
+  int data2;
+};
+
+// CHECK: @_ZZ5func0vE1a = internal addrspace(3) global %struct.MyStruct zeroinitializer
+// CHECK: @_ZZ5func1vE1a = internal addrspace(3) global float 0.000000e+00
+// CHECK: @_ZZ5func2vE1a = internal addrspace(3) global [256 x float] zeroinitializer
+// CHECK: @_ZZ5func3vE1a = internal addrspace(3) global float 0.000000e+00
+// CHECK: @_ZZ5func4vE1a = internal addrspace(3) global float 0.000000e+00
+// CHECK: @b = addrspace(3) global float 0.000000e+00
+
 __device__ void foo() {
   // CHECK: load i32* addrspacecast (i32 addrspace(1)* @i to i32*)
   i++;
@@ -22,15 +37,66 @@ __device__ void foo() {
   k++;
 
   static int li;
-  // CHECK: load i32 addrspace(1)* @_ZZ3foovE2li
+  // CHECK: load i32* addrspacecast (i32 addrspace(1)* @_ZZ3foovE2li to i32*)
   li++;
 
   __constant__ int lj;
-  // CHECK: load i32 addrspace(4)* @_ZZ3foovE2lj
+  // CHECK: load i32* addrspacecast (i32 addrspace(4)* @_ZZ3foovE2lj to i32*)
   lj++;
 
   __shared__ int lk;
-  // CHECK: load i32 addrspace(3)* @_ZZ3foovE2lk
+  // CHECK: load i32* addrspacecast (i32 addrspace(3)* @_ZZ3foovE2lk to i32*)
   lk++;
 }
 
+__device__ void func0() {
+  __shared__ MyStruct a;
+  MyStruct *ap = &a; // composite type
+  ap->data1 = 1;
+  ap->data2 = 2;
+}
+// CHECK: define void @_Z5func0v()
+// CHECK: store %struct.MyStruct* addrspacecast (%struct.MyStruct addrspace(3)* @_ZZ5func0vE1a to %struct.MyStruct*), %struct.MyStruct** %ap
+
+__device__ void callee(float *ap) {
+  *ap = 1.0f;
+}
+
+__device__ void func1() {
+  __shared__ float a;
+  callee(&a); // implicit cast from parameters
+}
+// CHECK: define void @_Z5func1v()
+// CHECK: call void @_Z6calleePf(float* addrspacecast (float addrspace(3)* @_ZZ5func1vE1a to float*))
+
+__device__ void func2() {
+  __shared__ float a[256];
+  float *ap = &a[128]; // implicit cast from a decayed array
+  *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
+
+__device__ void func3() {
+  __shared__ float a;
+  float *ap = reinterpret_cast<float *>(&a); // explicit cast
+  *ap = 1.0f;
+}
+// CHECK: define void @_Z5func3v()
+// CHECK: store float* addrspacecast (float addrspace(3)* @_ZZ5func3vE1a to float*), float** %ap
+
+__device__ void func4() {
+  __shared__ float a;
+  float *ap = (float *)&a; // explicit c-style cast
+  *ap = 1.0f;
+}
+// CHECK: define void @_Z5func4v()
+// CHECK: store float* addrspacecast (float addrspace(3)* @_ZZ5func4vE1a to float*), float** %ap
+
+__shared__ float b;
+
+__device__ float *func5() {
+  return &b; // implicit cast from a return value
+}
+// CHECK: define float* @_Z5func5v()
+// CHECK: ret float* addrspacecast (float addrspace(3)* @b to float*)