57 lines · cpp
1// RUN: %clang_cc1 -triple spirv64 -x hip -emit-llvm -fcuda-is-device \2// RUN: -o - %s | FileCheck %s3// RUN: %clang_cc1 -triple spirv64-amd-amdhsa -x hip -emit-llvm -fcuda-is-device \4// RUN: -o - %s | FileCheck %s5 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 global13__device__ int d;14 15// CHECK: @c ={{.*}} addrspace(1) externally_initialized constant16__constant__ int c;17 18// CHECK: @s ={{.*}} addrspace(3) global19__shared__ int s;20 21// CHECK: @foo ={{.*}} addrspace(1) externally_initialized global %struct.foo_t22__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) constant29 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