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 // 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 }