1 // REQUIRES: amdgpu-registered-target
2 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
3 // RUN: -std=c++17 -O3 -mllvm -amdgpu-internalize-symbols -emit-llvm -o - \
4 // RUN: -cuid=123 | FileCheck -check-prefix=DEV %s
5 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip %s \
6 // RUN: -std=c++17 -O3 -emit-llvm -o - -cuid=123 | FileCheck -check-prefix=HOST %s
7
8 // Negative tests.
9
10 // RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -fcuda-is-device -x hip %s \
11 // RUN: -std=c++17 -O3 -mllvm -amdgpu-internalize-symbols -emit-llvm -o - \
12 // RUN: | FileCheck -check-prefix=DEV-NEG %s
13 // RUN: %clang_cc1 -triple x86_64-unknown-linux-gnu -x hip %s \
14 // RUN: -std=c++17 -O3 -emit-llvm -o - | FileCheck -check-prefix=HOST-NEG %s
15
16 #include "Inputs/cuda.h"
17
18 // DEV-DAG: @v1
19 __device__ int v1;
20
21 // DEV-DAG: @v2
22 __constant__ int v2;
23
24 // Check device variables used by neither host nor device functioins are not kept.
25
26 // DEV-NEG-NOT: @_ZL2v3
27 static __device__ int v3;
28
29 // Check device variables used by host functions are kept.
30
31 // DEV-DAG: @u1
32 __device__ int u1;
33
34 // DEV-DAG: @u2
35 __constant__ int u2;
36
37 // Check host-used static device var is in llvm.compiler.used.
38 // DEV-DAG: @_ZL2u3
39 static __device__ int u3;
40
41 // Check device-used static device var is emitted but is not in llvm.compiler.used.
42 // DEV-DAG: @_ZL2u4
43 static __device__ int u4;
44
45 // Check device variables with used attribute are always kept.
46 // DEV-DAG: @u5
47 __device__ __attribute__((used)) int u5;
48
49 // Test external device variable ODR-used by host code is not emitted or registered.
50 // DEV-NEG-NOT: @ext_var
51 extern __device__ int ext_var;
52
53 // DEV-DAG: @inline_var = linkonce_odr addrspace(1) externally_initialized global i32 0
54 __device__ inline int inline_var;
55
56 template<typename T>
57 using func_t = T (*) (T, T);
58
59 template <typename T>
add_func(T x,T y)60 __device__ T add_func (T x, T y)
61 {
62 return x + y;
63 }
64
65 // DEV-DAG: @_Z10p_add_funcIiE = linkonce_odr addrspace(1) externally_initialized global ptr @_Z8add_funcIiET_S0_S0_
66 template <typename T>
67 __device__ func_t<T> p_add_func = add_func<T>;
68
69 // Check non-constant constexpr variables ODR-used by host code only is not emitted.
70 // DEV-NEG-NOT: constexpr_var1a
71 // DEV-NEG-NOT: constexpr_var1b
72 constexpr int constexpr_var1a = 1;
73 inline constexpr int constexpr_var1b = 1;
74
75 // Check constant constexpr variables ODR-used by host code only.
76 // Device-side constexpr variables accessed by host code should be externalized and kept.
77 // DEV-DAG: @_ZL15constexpr_var2a = addrspace(4) externally_initialized constant i32 2
78 // DEV-DAG: @constexpr_var2b = linkonce_odr addrspace(4) externally_initialized constant i32 2
79 __constant__ constexpr int constexpr_var2a = 2;
80 inline __constant__ constexpr int constexpr_var2b = 2;
81
82 void use(func_t<int> p);
83 __host__ __device__ void use(const int *p);
84
85 // Check static device variable in host function.
86 // DEV-DAG: @_ZZ4fun1vE11static_var1 = addrspace(1) externally_initialized global i32 3
fun1()87 void fun1() {
88 static __device__ int static_var1 = 3;
89 use(&u1);
90 use(&u2);
91 use(&u3);
92 use(&ext_var);
93 use(&inline_var);
94 use(p_add_func<int>);
95 use(&constexpr_var1a);
96 use(&constexpr_var1b);
97 use(&constexpr_var2a);
98 use(&constexpr_var2b);
99 use(&static_var1);
100 }
101
102 // Check static variable in host device function.
103 // DEV-DAG: @_ZZ4fun2vE11static_var2 = internal addrspace(1) global i32 4
104 // DEV-DAG: @_ZZ4fun2vE11static_var3 = addrspace(1) global i32 4
fun2()105 __host__ __device__ void fun2() {
106 static int static_var2 = 4;
107 static __device__ int static_var3 = 4;
108 use(&static_var2);
109 use(&static_var3);
110 }
111
kern1(int ** x)112 __global__ void kern1(int **x) {
113 *x = &u4;
114 fun2();
115 }
116
117 // Check static variables of lambda functions.
118
119 // Lambda functions are implicit host device functions.
120 // Default static variables in lambda functions should be treated
121 // as host variables on host side, therefore should not be forced
122 // to be emitted on device.
123
124 // DEV-DAG: @_ZZZN21TestStaticVarInLambda3funEvENKUlPcE_clES0_E4var2 = addrspace(1) externally_initialized global i32 5
125 // DEV-NEG-NOT: @_ZZZN21TestStaticVarInLambda3funEvENKUlPcE_clES0_E4var1
126 namespace TestStaticVarInLambda {
127 class A {
128 public:
129 A(char *);
130 };
fun()131 void fun() {
132 (void) [](char *c) {
133 static A var1(c);
134 static __device__ int var2 = 5;
135 (void) var1;
136 (void) var2;
137 };
138 }
139 }
140
141 // Check implicit constant variable ODR-used by host code is not emitted.
142
143 // AST contains instantiation of al<ar>, which triggers AST instantiation
144 // of x::al<ar>::am, which triggers AST instatiation of x::ap<ar>,
145 // which triggers AST instantiation of aw<ar>::c, which has type
146 // ar. ar has base class x which has member ah. x::ah is initialized
147 // with function pointer pointing to ar:as, which returns an object
148 // of type ou. The constexpr aw<ar>::c is an implicit constant variable
149 // which is ODR-used by host function x::ap<ar>. An incorrect implementation
150 // will force aw<ar>::c to be emitted on device side, which will trigger
151 // emit of x::as and further more ctor of ou and variable o.
152 // The ODR-use of aw<ar>::c in x::ap<ar> should be treated as a host variable
153 // instead of device variable.
154
155 // DEV-NEG-NOT: _ZN16TestConstexprVar1oE
156 namespace TestConstexprVar {
157 char o;
158 class ou {
159 public:
ou(char)160 ou(char) { __builtin_strlen(&o); }
161 };
162 template < typename ao > struct aw { static constexpr ao c; };
163 class x {
164 protected:
165 typedef ou (*y)(const x *);
x(y ag)166 constexpr x(y ag) : ah(ag) {}
167 template < bool * > struct ak;
168 template < typename > struct al {
169 static bool am;
170 static ak< &am > an;
171 };
ap()172 template < typename ao > static x ap() { (void)aw< ao >::c; return x(nullptr); }
173 y ah;
174 };
175 template < typename ao > bool x::al< ao >::am(&ap< ao >);
176 class ar : x {
177 public:
ar()178 constexpr ar() : x(as) {}
as(const x *)179 static ou as(const x *) { return 0; }
180 al< ar > av;
181 };
182 }
183
184 // Check the exact list of variables to ensure @_ZL2u4 is not among them.
185 // DEV: @llvm.compiler.used = {{[^@]*}} @_Z10p_add_funcIiE
186 // DEV-SAME: {{^[^@]*}} @_ZL15constexpr_var2a
187 // DEV-SAME: {{^[^@]*}} @_ZL2u3
188 // DEV-SAME: {{^[^@]*}} @_ZZ4fun1vE11static_var1
189 // DEV-SAME: {{^[^@]*}} @_ZZZN21TestStaticVarInLambda3funEvENKUlPcE_clES0_E4var2
190 // DEV-SAME: {{^[^@]*}} @__hip_cuid_{{[0-9a-f]+}}
191 // DEV-SAME: {{^[^@]*}} @constexpr_var2b
192 // DEV-SAME: {{^[^@]*}} @inline_var
193 // DEV-SAME: {{^[^@]*}} @u1
194 // DEV-SAME: {{^[^@]*}} @u2
195 // DEV-SAME: {{^[^@]*}} @u5
196 // DEV-SAME: {{^[^@]*$}}
197
198 // HOST-DAG: hipRegisterVar{{.*}}@u1
199 // HOST-DAG: hipRegisterVar{{.*}}@u2
200 // HOST-DAG: hipRegisterVar{{.*}}@_ZL2u3
201 // HOST-DAG: hipRegisterVar{{.*}}@constexpr_var2b
202 // HOST-DAG: hipRegisterVar{{.*}}@u5
203 // HOST-DAG: hipRegisterVar{{.*}}@inline_var
204 // HOST-DAG: hipRegisterVar{{.*}}@_Z10p_add_funcIiE
205 // HOST-NEG-NOT: hipRegisterVar{{.*}}@_ZZ4fun1vE11static_var1
206 // HOST-NEG-NOT: hipRegisterVar{{.*}}@_ZZ4fun2vE11static_var2
207 // HOST-NEG-NOT: hipRegisterVar{{.*}}@_ZZ4fun2vE11static_var3
208 // HOST-NEG-NOT: hipRegisterVar{{.*}}@_ZZZN21TestStaticVarInLambda3funEvENKUlPcE_clES0_E4var2
209 // HOST-NEG-NOT: hipRegisterVar{{.*}}@_ZZZN21TestStaticVarInLambda3funEvENKUlPcE_clES0_E4var1
210 // HOST-NEG-NOT: hipRegisterVar{{.*}}@ext_var
211 // HOST-NEG-NOT: hipRegisterVar{{.*}}@_ZL2u4
212 // HOST-NEG-NOT: hipRegisterVar{{.*}}@constexpr_var1a
213 // HOST-NEG-NOT: hipRegisterVar{{.*}}@constexpr_var1b
214 // HOST-NEG-NOT: hipRegisterVar{{.*}}@constexpr_var2a
215