1 // RUN: %clang_cc1 -triple spirv64 -x hip -emit-llvm -fcuda-is-device \ 2 // RUN: -o - %s | FileCheck %s 3 // RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -emit-llvm -fcuda-is-device \ 4 // RUN: -o - %s | FileCheck %s 5 6 #define __device__ __attribute__((device)) 7 #define __shared__ __attribute__((shared)) 8 #define __constant__ __attribute__((constant)) 9 10 // CHECK: %struct.foo_t = type { i32, ptr addrspace(4) } 11 12 // CHECK: @d ={{.*}} addrspace(1) externally_initialized global 13 __device__ int d; 14 15 // CHECK: @c ={{.*}} addrspace(1) externally_initialized constant 16 __constant__ int c; 17 18 // CHECK: @s ={{.*}} addrspace(3) global 19 __shared__ int s; 20 21 // CHECK: @foo ={{.*}} addrspace(1) externally_initialized global %struct.foo_t 22 __device__ struct foo_t { 23 int i; 24 int* pi; 25 } foo; 26 27 // Check literals are placed in address space 1 (CrossWorkGroup/__global). 28 // CHECK: @.str ={{.*}} unnamed_addr addrspace(1) constant 29 30 // CHECK: define{{.*}} spir_func noundef ptr addrspace(4) @_Z3barPi(ptr addrspace(4) 31 __device__ int* bar(int *x) { 32 return x; 33 } 34 35 // CHECK: define{{.*}} spir_func noundef ptr addrspace(4) @_Z5baz_dv() 36 __device__ int* baz_d() { 37 // CHECK: ret ptr addrspace(4) addrspacecast (ptr addrspace(1) @d to ptr addrspace(4) 38 return &d; 39 } 40 41 // CHECK: define{{.*}} spir_func noundef ptr addrspace(4) @_Z5baz_cv() 42 __device__ int* baz_c() { 43 // CHECK: ret ptr addrspace(4) addrspacecast (ptr addrspace(1) @c to ptr addrspace(4) 44 return &c; 45 } 46 47 // CHECK: define{{.*}} spir_func noundef ptr addrspace(4) @_Z5baz_sv() 48 __device__ int* baz_s() { 49 // CHECK: ret ptr addrspace(4) addrspacecast (ptr addrspace(3) @s to ptr addrspace(4) 50 return &s; 51 } 52 53 // CHECK: define{{.*}} spir_func noundef ptr addrspace(4) @_Z3quzv() 54 __device__ const char* quz() { 55 return "abc"; 56 } 57