1 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc -DLOAD 2 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-target-device -fvisibility=protected -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -DLOAD | FileCheck %s 3 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown %s -fopenmp-is-target-device -fvisibility=protected -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t 4 // RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-target-device -fvisibility=protected -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - -DLOAD | FileCheck %s 5 6 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - -DOMP5 | FileCheck %s --check-prefix HOST5 7 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc -DOMP5 8 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-target-device -fvisibility=protected -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -DOMP5 | FileCheck %s --check-prefix DEV5 9 10 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - -DOMP5 | FileCheck %s --check-prefix KMPC-ONLY 11 12 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - -DOMP5 | FileCheck %s --check-prefix SIMD-ONLY 13 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc -DOMP5 14 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-target-device -fvisibility=protected -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -DOMP5 | FileCheck %s --check-prefix SIMD-ONLY 15 16 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc -fopenmp-version=45 17 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-target-device -fvisibility=protected -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -fopenmp-version=45 | FileCheck %s --check-prefix SIMD-ONLY 18 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown %s -fopenmp-is-target-device -fvisibility=protected -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t -fopenmp-version=45 19 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -fopenmp-is-target-device -fvisibility=protected -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify -o - -fopenmp-version=45 | FileCheck %s --check-prefix SIMD-ONLY 20 21 // expected-no-diagnostics 22 23 // SIMD-ONLY-NOT: {{__kmpc|__tgt}} 24 // KMPC-ONLY-NOT: __tgt 25 26 // CHECK-NOT: define {{.*}}{{baz1|baz4|maini1|Base|virtual_}} 27 // CHECK-DAG: Bake 28 // CHECK-NOT: @{{hhh|ggg|fff|eee}} = 29 // CHECK-DAG: @flag = protected global i8 undef, 30 // CHECK-DAG: @dx = {{protected | }}global i32 0, 31 // CHECK-DAG: @dy = {{protected | }}global i32 0, 32 // CHECK-DAG: @bbb = {{protected | }}global i32 0, 33 // CHECK-DAG: weak constant %struct.__tgt_offload_entry { 34 // CHECK-DAG: @ccc = external global i32, 35 // CHECK-DAG: @ddd = {{protected | }}global i32 0, 36 // CHECK-DAG: @hhh_decl_tgt_ref_ptr = weak global ptr null 37 // CHECK-DAG: @ggg_decl_tgt_ref_ptr = weak global ptr null 38 // CHECK-DAG: @fff_decl_tgt_ref_ptr = weak global ptr null 39 // CHECK-DAG: @eee_decl_tgt_ref_ptr = weak global ptr null 40 // CHECK-DAG: @{{.*}}maini1{{.*}}aaa = internal global i64 23, 41 // CHECK-DAG: @pair = {{.*}}addrspace(3) global %struct.PAIR undef 42 // CHECK-DAG: @_ZN2SS3SSSE ={{ protected | }}global i32 1, 43 // CHECK-DAG: @b ={{ protected | }}global i32 15, 44 // CHECK-DAG: @d ={{ protected | }}global i32 0, 45 // CHECK-DAG: @c = external global i32, 46 // CHECK-DAG: @globals ={{ protected | }}global %struct.S zeroinitializer, 47 // CHECK-DAG: [[STAT:@.+stat]] = internal global %struct.S zeroinitializer, 48 // CHECK-DAG: [[STAT_REF:@.+]] = internal constant ptr [[STAT]] 49 // CHECK-DAG: @out_decl_target ={{ protected | }}global i32 0, 50 // CHECK-DAG: @llvm.compiler.used = appending global [1 x ptr] [ptr [[STAT_REF]]], 51 52 // CHECK-DAG: define {{.*}}i32 @{{.*}}{{foo|bar|baz2|baz3|FA|f_method}}{{.*}}() 53 // CHECK-DAG: define {{.*}}void @{{.*}}TemplateClass{{.*}}(ptr {{[^,]*}} %{{.*}}) 54 // CHECK-DAG: define {{.*}}i32 @{{.*}}TemplateClass{{.*}}f_method{{.*}}(ptr {{[^,]*}} %{{.*}}) 55 56 #ifndef HEADER 57 #define HEADER 58 59 int dx = 0; 60 extern int dx; 61 #pragma omp declare target to(dx) 62 63 int dy = 0; 64 #pragma omp begin declare target 65 66 extern int dy; 67 #pragma omp end declare target 68 69 #pragma omp declare target 70 bool flag [[clang::loader_uninitialized]]; 71 extern int bbb; 72 #pragma omp end declare target 73 #pragma omp declare target 74 extern int bbb; 75 #pragma omp end declare target 76 77 #pragma omp declare target 78 extern int aaa; 79 int bbb = 0; 80 extern int ccc; 81 int ddd = ccc; 82 #pragma omp end declare target 83 84 #pragma omp declare target 85 extern int bbb; 86 #pragma omp end declare target 87 88 extern int eee; 89 int fff = 0; 90 extern int ggg; 91 int hhh = 0; 92 #pragma omp declare target link(eee, fff, ggg, hhh) 93 94 int out_decl_target = 0; 95 #ifdef OMP5 96 #pragma omp declare target(out_decl_target) 97 #endif 98 99 #pragma omp declare target 100 void lambda () { 101 #ifdef __cpp_lambdas 102 (void)[&] { (void)out_decl_target; }; 103 #else 104 #pragma clang __debug captured 105 { 106 (void)out_decl_target; 107 } 108 #endif 109 }; 110 #pragma omp end declare target 111 112 template <typename T> 113 class TemplateClass { 114 T a; 115 public: 116 TemplateClass() {} 117 T f_method() const { return a; } 118 }; 119 120 int foo(); 121 122 static int baz1() { return 0; } 123 124 int baz2(); 125 126 int baz4() { return 5; } 127 128 template <typename T> 129 T FA() { 130 TemplateClass<T> s; 131 return s.f_method(); 132 } 133 134 #pragma omp declare target 135 struct S { 136 int a; 137 S(int a) : a(a) {} 138 }; 139 140 int foo() { return 0; } 141 int b = 15; 142 int d; 143 S globals(d); 144 static S stat(d); 145 #pragma omp end declare target 146 int c; 147 148 int bar() { return 1 + foo() + bar() + baz1() + baz2(); } 149 150 int maini1() { 151 int a; 152 static long aa = 32 + bbb + ccc + fff + ggg; 153 // CHECK-DAG: define weak_odr protected void @__omp_offloading_{{.*}}maini1{{.*}}_l[[@LINE+1]](ptr {{.*}}, ptr noundef nonnull align {{[0-9]+}} dereferenceable({{[0-9]+}}) %{{.*}}, i64 {{.*}}, i64 {{.*}}) 154 #pragma omp target map(tofrom \ 155 : a, b) 156 { 157 S s(a); 158 static long aaa = 23; 159 a = foo() + bar() + b + c + d + aa + aaa + FA<int>(); 160 } 161 return baz4(); 162 } 163 164 int baz3() { return 2 + baz2(); } 165 int baz2() { 166 // CHECK-DAG: define weak_odr protected void @__omp_offloading_{{.*}}baz2{{.*}}_l[[@LINE+1]](ptr {{.*}}, i64 {{.*}}) 167 #pragma omp target parallel 168 ++c; 169 return 2 + baz3(); 170 } 171 172 extern int create() throw(); 173 174 static __typeof(create) __t_create __attribute__((__weakref__("__create"))); 175 176 int baz5() { 177 bool a; 178 // CHECK-DAG: define weak_odr protected void @__omp_offloading_{{.*}}baz5{{.*}}_l[[@LINE+1]](ptr {{.*}}, i64 {{.*}}) 179 #pragma omp target 180 a = __extension__(void *) & __t_create != 0; 181 return a; 182 } 183 184 template <typename T> 185 struct Base { 186 virtual ~Base() {} 187 }; 188 189 template class Base<int>; 190 191 template <typename T> 192 struct Bake { 193 virtual ~Bake() {} 194 }; 195 196 #pragma omp declare target 197 template class Bake<int>; 198 #pragma omp end declare target 199 200 struct BaseNonT { 201 virtual ~BaseNonT() {} 202 }; 203 204 #pragma omp declare target 205 struct BakeNonT { 206 virtual ~BakeNonT() {} 207 }; 208 #pragma omp end declare target 209 210 template <typename T> 211 struct B { 212 virtual void virtual_foo(); 213 }; 214 215 void new_bar() { new B<int>(); } 216 217 template <typename T> 218 void B<T>::virtual_foo() { 219 #pragma omp target 220 {} 221 } 222 223 struct A { 224 virtual void emitted() {} 225 }; 226 227 template <typename T> 228 struct C : public A { 229 virtual void emitted(); 230 }; 231 232 template <typename T> 233 void C<T>::emitted() { 234 #pragma omp target 235 {} 236 } 237 238 int main() { 239 A *X = new C<int>(); 240 X->emitted(); 241 return 0; 242 } 243 244 // CHECK-DAG: define {{.*}}void @__omp_offloading_{{.*}}virtual_foo{{.*}}_l[[@LINE-25]](ptr {{.*}}) 245 // CHECK-DAG: define {{.*}}void @__omp_offloading_{{.*}}emitted{{.*}}_l[[@LINE-11]](ptr {{.*}}) 246 247 template <typename T> 248 struct TTT { 249 virtual void emitted() { 250 #pragma omp target 251 {} 252 } 253 }; 254 255 // CHECK-DAG: define {{.*}}void @__omp_offloading_{{.*}}emitted{{.*}}_l[[@LINE-5]](ptr {{.*}}) 256 257 // CHECK-DAG: declare extern_weak noundef signext i32 @__create() 258 259 // CHECK-NOT: define {{.*}}{{baz1|baz4|maini1|Base|virtual_}} 260 261 // CHECK-DAG: !{{{.+}}virtual_foo 262 263 #ifdef OMP5 264 void host_fun() {} 265 #pragma omp declare target to(host_fun) device_type(host) 266 void device_fun() {} 267 #pragma omp declare target to(device_fun) device_type(nohost) 268 // HOST5-NOT: define {{.*}}void {{.*}}device_fun{{.*}} 269 // HOST5: define {{.*}}void {{.*}}host_fun{{.*}} 270 // HOST5-NOT: define {{.*}}void {{.*}}device_fun{{.*}} 271 272 // DEV5-NOT: define {{.*}}void {{.*}}host_fun{{.*}} 273 // DEV5: define {{.*}}void {{.*}}device_fun{{.*}} 274 // DEV5-NOT: define {{.*}}void {{.*}}host_fun{{.*}} 275 #endif // OMP5 276 277 struct PAIR { 278 int a; 279 int b; 280 }; 281 282 #pragma omp declare target 283 PAIR pair __attribute__((address_space(3), loader_uninitialized)); 284 #pragma omp end declare target 285 286 #endif // HEADER 287 288 #ifdef LOAD 289 #pragma omp declare target 290 void new_bar1() { 291 TTT<int> *X = new TTT<int>(); 292 X->emitted(); 293 } 294 #pragma omp end declare target 295 296 struct SS { 297 #pragma omp declare target 298 static int SSS; 299 #pragma omp end declare target 300 }; 301 int SS::SSS = 1; 302 303 #endif 304