304 lines · cpp
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 -DLOAD2// 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 %s3// 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 %t4// 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 %s5 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 HOST57// 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 -DOMP58// 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 DEV59 10// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - -DOMP5 | FileCheck %s --check-prefix KMPC-ONLY11 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-ONLY13// 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 -DOMP514// 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-ONLY15 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=4517// 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-ONLY18// 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=4519// 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-ONLY20 21// expected-no-diagnostics22 23// SIMD-ONLY-NOT: {{__kmpc|__tgt}}24// KMPC-ONLY-NOT: __tgt25 26// CHECK-NOT: define {{.*}}{{baz1|baz4|maini1|Base|virtual_}}27// CHECK-DAG: Bake28// 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 null37// CHECK-DAG: @ggg_decl_tgt_ref_ptr = weak global ptr null38// CHECK-DAG: @fff_decl_tgt_ref_ptr = weak global ptr null39// CHECK-DAG: @eee_decl_tgt_ref_ptr = weak global ptr null40// CHECK-DAG: @{{.*}}maini1{{.*}}aaa = internal global i64 23,41// CHECK-DAG: @pair = {{.*}}addrspace(3) global %struct.PAIR undef42// 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 HEADER57#define HEADER58 59int dx = 0;60extern int dx;61#pragma omp declare target to(dx)62 63int dy = 0;64#pragma omp begin declare target65 66extern int dy;67#pragma omp end declare target68 69#pragma omp declare target70bool flag [[clang::loader_uninitialized]];71extern int bbb;72#pragma omp end declare target73#pragma omp declare target74extern int bbb;75#pragma omp end declare target76 77#pragma omp declare target78extern int aaa;79int bbb = 0;80extern int ccc;81int ddd = ccc;82#pragma omp end declare target83 84#pragma omp declare target85extern int bbb;86#pragma omp end declare target87 88extern int eee;89int fff = 0;90extern int ggg;91int hhh = 0;92#pragma omp declare target link(eee, fff, ggg, hhh)93 94int out_decl_target = 0;95#ifdef OMP596#pragma omp declare target(out_decl_target)97#endif98 99#pragma omp declare target100void lambda () {101#ifdef __cpp_lambdas102 (void)[&] { (void)out_decl_target; };103#else104#pragma clang __debug captured105 {106 (void)out_decl_target;107 }108#endif109};110#pragma omp end declare target111 112template <typename T>113class TemplateClass {114 T a;115public:116 TemplateClass() {}117 T f_method() const { return a; }118};119 120int foo();121 122static int baz1() { return 0; }123 124int baz2();125 126int baz4() { return 5; }127 128template <typename T>129T FA() {130 TemplateClass<T> s;131 return s.f_method();132}133 134#pragma omp declare target135struct S {136 int a;137 S(int a) : a(a) {}138};139 140int foo() { return 0; }141int b = 15;142int d;143S globals(d);144static S stat(d);145#pragma omp end declare target146int c;147 148int bar() { return 1 + foo() + bar() + baz1() + baz2(); }149 150int 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 164int baz3() { return 2 + baz2(); }165int baz2() {166// CHECK-DAG: define weak_odr protected void @__omp_offloading_{{.*}}baz2{{.*}}_l[[@LINE+1]](ptr {{.*}}, i64 {{.*}})167#pragma omp target parallel168 ++c;169 return 2 + baz3();170}171 172extern int create() throw();173 174static __typeof(create) __t_create __attribute__((__weakref__("__create")));175 176int baz5() {177 bool a;178// CHECK-DAG: define weak_odr protected void @__omp_offloading_{{.*}}baz5{{.*}}_l[[@LINE+1]](ptr {{.*}}, i64 {{.*}})179#pragma omp target180 a = __extension__(void *) & __t_create != 0;181 return a;182}183 184template <typename T>185struct Base {186 virtual ~Base() {}187};188 189template class Base<int>;190 191template <typename T>192struct Bake {193 virtual ~Bake() {}194};195 196#pragma omp declare target197template class Bake<int>;198#pragma omp end declare target199 200struct BaseNonT {201 virtual ~BaseNonT() {}202};203 204#pragma omp declare target205struct BakeNonT {206 virtual ~BakeNonT() {}207};208#pragma omp end declare target209 210template <typename T>211struct B {212 virtual void virtual_foo();213};214 215void new_bar() { new B<int>(); }216 217template <typename T>218void B<T>::virtual_foo() {219#pragma omp target220 {}221}222 223struct A {224 virtual void emitted() {}225};226 227template <typename T>228struct C : public A {229 virtual void emitted();230};231 232template <typename T>233void C<T>::emitted() {234#pragma omp target235 {}236}237 238int 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 247template <typename T>248struct TTT {249 virtual void emitted() {250#pragma omp target251 {}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_foo262 263#ifdef OMP5264void host_fun() {}265#pragma omp declare target to(host_fun) device_type(host)266void 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 // OMP5276 277struct PAIR {278 int a;279 int b;280};281 282#pragma omp declare target283PAIR pair __attribute__((address_space(3), loader_uninitialized));284#pragma omp end declare target285 286#endif // HEADER287 288#ifdef LOAD289#pragma omp declare target290void new_bar1() {291 TTT<int> *X = new TTT<int>();292 X->emitted();293}294#pragma omp end declare target295 296struct SS {297#pragma omp declare target298 static int SSS;299#pragma omp end declare target300};301int SS::SSS = 1;302 303#endif304