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