clang/test/CodeGenHIP/hipspv-addr-spaces.cpp

   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, ptr 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 literals are placed in address space 1 (CrossWorkGroup/__global).
  26 // CHECK: @.str ={{.*}} unnamed_addr addrspace(1) constant
  27
  28 // CHECK: define{{.*}} spir_func noundef ptr addrspace(4) @_Z3barPi(ptr addrspace(4)
  29 __device__ int* bar(int *x) {
  30   return x;
  31 }
  32
  33 // CHECK: define{{.*}} spir_func noundef ptr addrspace(4) @_Z5baz_dv()
  34 __device__ int* baz_d() {
  35   // CHECK: ret ptr addrspace(4) addrspacecast (ptr addrspace(1) @d to ptr addrspace(4)
  36   return &d;
  37 }
  38
  39 // CHECK: define{{.*}} spir_func noundef ptr addrspace(4) @_Z5baz_cv()
  40 __device__ int* baz_c() {
  41   // CHECK: ret ptr addrspace(4) addrspacecast (ptr addrspace(1) @c to ptr addrspace(4)
  42   return &c;
  43 }
  44
  45 // CHECK: define{{.*}} spir_func noundef ptr addrspace(4) @_Z5baz_sv()
  46 __device__ int* baz_s() {
  47   // CHECK: ret ptr addrspace(4) addrspacecast (ptr addrspace(3) @s to ptr addrspace(4)
  48   return &s;
  49 }
  50
  51 // CHECK: define{{.*}} spir_func noundef ptr addrspace(4) @_Z3quzv()
  52 __device__ const char* quz() {
  53   return "abc";
  54 }