1 // RUN: %clang_cc1 -triple spirv64 -x hip -emit-llvm -fcuda-is-device \
2 // RUN:   -o - %s | FileCheck %s
3 
4 #define __device__ __attribute__((device))
5 #define __shared__ __attribute__((shared))
6 #define __constant__ __attribute__((constant))
7 
8 // CHECK: %struct.foo_t = type { i32, i32 addrspace(4)* }
9 
10 // CHECK: @d ={{.*}} addrspace(1) externally_initialized global
11 __device__ int d;
12 
13 // CHECK: @c ={{.*}} addrspace(1) externally_initialized global
14 __constant__ int c;
15 
16 // CHECK: @s ={{.*}} addrspace(3) global
17 __shared__ int s;
18 
19 // CHECK: @foo ={{.*}} addrspace(1) externally_initialized global %struct.foo_t
20 __device__ struct foo_t {
21   int i;
22   int* pi;
23 } foo;
24 
25 // CHECK: define{{.*}} spir_func noundef i32 addrspace(4)* @_Z3barPi(i32 addrspace(4)*
26 __device__ int* bar(int *x) {
27   return x;
28 }
29 
30 // CHECK: define{{.*}} spir_func noundef i32 addrspace(4)* @_Z5baz_dv()
31 __device__ int* baz_d() {
32   // CHECK: ret i32 addrspace(4)* addrspacecast (i32 addrspace(1)* @d to i32 addrspace(4)*
33   return &d;
34 }
35 
36 // CHECK: define{{.*}} spir_func noundef i32 addrspace(4)* @_Z5baz_cv()
37 __device__ int* baz_c() {
38   // CHECK: ret i32 addrspace(4)* addrspacecast (i32 addrspace(1)* @c to i32 addrspace(4)*
39   return &c;
40 }
41 
42 // CHECK: define{{.*}} spir_func noundef i32 addrspace(4)* @_Z5baz_sv()
43 __device__ int* baz_s() {
44   // CHECK: ret i32 addrspace(4)* addrspacecast (i32 addrspace(3)* @s to i32 addrspace(4)*
45   return &s;
46 }
47