1 // REQUIRES: x86-registered-target
2 // REQUIRES: amdgpu-registered-target
3
4 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -std=c++11 \
5 // RUN: -emit-llvm -o - -x hip %s | FileCheck \
6 // RUN: -check-prefixes=DEV %s
7
8 // RUN: %clang_cc1 -triple x86_64-gnu-linux -std=c++11 \
9 // RUN: -emit-llvm -o - -x hip %s | FileCheck \
10 // RUN: -check-prefixes=HOST %s
11
12 #include "Inputs/cuda.h"
13
14 // Test function scope static device variable, which should not be externalized.
15 // DEV-DAG: @_ZZ6kernelPiPPKiE1w = internal addrspace(4) constant i32 1
16 // DEV-DAG: @_ZZ6kernelPiPPKiE21local_static_constant = internal addrspace(4) constant i32 42
17 // DEV-DAG: @_ZZ6kernelPiPPKiE19local_static_device = internal addrspace(1) constant i32 43
18
19 // Check a static device variable referenced by host function is externalized.
20 // DEV-DAG: @_ZL1x = addrspace(1) externally_initialized global i32 0
21 // HOST-DAG: @_ZL1x = internal global i32 undef
22 // HOST-DAG: @[[DEVNAMEX:[0-9]+]] = {{.*}}c"_ZL1x\00"
23
24 static __device__ int x;
25
26 // Check a static device variables referenced only by device functions and kernels
27 // is not externalized.
28 // DEV-DAG: @_ZL2x2 = internal addrspace(1) global i32 0
29 static __device__ int x2;
30
31 // Check a static device variable referenced by host device function is externalized.
32 // DEV-DAG: @_ZL2x3 = addrspace(1) externally_initialized global i32 0
33 static __device__ int x3;
34
35 // Check a static device variable referenced in file scope is externalized.
36 // DEV-DAG: @_ZL2x4 = addrspace(1) externally_initialized global i32 0
37 static __device__ int x4;
38 int& x4_ref = x4;
39
40 // Check a static device variable in anonymous namespace.
41 // DEV-DAG: @_ZN12_GLOBAL__N_12x5E = addrspace(1) externally_initialized global i32 0
42 namespace {
43 static __device__ int x5;
44 }
45
46 // Check a static constant variable referenced by host is externalized.
47 // DEV-DAG: @_ZL1y = addrspace(4) externally_initialized global i32 0
48 // HOST-DAG: @_ZL1y = internal global i32 undef
49 // HOST-DAG: @[[DEVNAMEY:[0-9]+]] = {{.*}}c"_ZL1y\00"
50
51 static __constant__ int y;
52
53 // Test static host variable, which should not be externalized nor registered.
54 // HOST-DAG: @_ZL1z = internal global i32 0
55 // DEV-NOT: @_ZL1z
56 static int z;
57
58 // Test implicit static constant variable, which should not be externalized.
59 // HOST-DAG: @_ZL2z2 = internal constant i32 456
60 // DEV-DAG: @_ZL2z2 = internal addrspace(4) constant i32 456
61
62 static constexpr int z2 = 456;
63
64 // Test static device variable in inline function, which should not be
65 // externalized nor registered.
66 // DEV-DAG: @_ZZ6devfunPPKiE1p = linkonce_odr addrspace(4) constant i32 2, comdat
67
68 // Check a static device variable referenced by host function only is externalized.
69 // DEV-DAG: @_ZL1w = addrspace(1) externally_initialized global i32 0
70 // HOST-DAG: @_ZL1w = internal global i32 undef
71 // HOST-DAG: @[[DEVNAMEW:[0-9]+]] = {{.*}}c"_ZL1w\00"
72
73 static __device__ int w;
74
devfun(const int ** b)75 inline __device__ void devfun(const int ** b) {
76 const static int p = 2;
77 b[0] = &p;
78 b[1] = &x2;
79 }
80
kernel(int * a,const int ** b)81 __global__ void kernel(int *a, const int **b) {
82 const static int w = 1;
83 const static __constant__ int local_static_constant = 42;
84 const static __device__ int local_static_device = 43;
85 a[0] = x;
86 a[1] = y;
87 a[2] = x2;
88 a[3] = x3;
89 a[4] = x4;
90 a[5] = x5;
91 b[0] = &w;
92 b[1] = &z2;
93 b[2] = &local_static_constant;
94 b[3] = &local_static_device;
95 devfun(b);
96 }
97
hdf(int * a)98 __host__ __device__ void hdf(int *a) {
99 a[0] = x3;
100 }
101
102 int* getDeviceSymbol(int *x);
103
foo(const int ** a)104 void foo(const int **a) {
105 getDeviceSymbol(&x);
106 getDeviceSymbol(&x5);
107 getDeviceSymbol(&y);
108 getDeviceSymbol(&w);
109 z = 123;
110 a[0] = &z2;
111 }
112
113 // HOST: __hipRegisterVar({{.*}}@_ZL1x {{.*}}@[[DEVNAMEX]]
114 // HOST: __hipRegisterVar({{.*}}@_ZL1y {{.*}}@[[DEVNAMEY]]
115 // HOST: __hipRegisterVar({{.*}}@_ZL1w {{.*}}@[[DEVNAMEW]]
116 // HOST-NOT: __hipRegisterVar({{.*}}@_ZZ6kernelPiPPKiE1w
117 // HOST-NOT: __hipRegisterVar({{.*}}@_ZZ6devfunPPKiE1p
118