1 // 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 2 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s 3 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t 4 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - | FileCheck %s 5 6 // 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 7 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o -| FileCheck %s --check-prefix SIMD-ONLY 8 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t 9 // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify -o - | FileCheck %s --check-prefix SIMD-ONLY 10 11 // expected-no-diagnostics 12 13 // SIMD-ONLY-NOT: {{__kmpc|__tgt}} 14 15 // CHECK-NOT: define {{.*}}{{baz1|baz4|maini1}} 16 // CHECK-NOT: @{{hhh|ggg|fff|eee}} = 17 // CHECK-DAG: @aaa = external global i32, 18 // CHECK-DAG: @bbb = global i32 0, 19 // CHECK-DAG: @ccc = external global i32, 20 // CHECK-DAG: @ddd = global i32 0, 21 // CHECK-DAG: @hhh_decl_tgt_link_ptr = common global i32* null 22 // CHECK-DAG: @ggg_decl_tgt_link_ptr = common global i32* null 23 // CHECK-DAG: @fff_decl_tgt_link_ptr = common global i32* null 24 // CHECK-DAG: @eee_decl_tgt_link_ptr = common global i32* null 25 // CHECK-DAG: @{{.*}}maini1{{.*}}aaa = internal global i64 23, 26 // CHECK-DAG: @b = global i32 15, 27 // CHECK-DAG: @d = global i32 0, 28 // CHECK-DAG: @c = external global i32, 29 // CHECK-DAG: @globals = global %struct.S zeroinitializer, 30 // CHECK-DAG: [[STAT:@.+stat]] = internal global %struct.S zeroinitializer, 31 // CHECK-DAG: [[STAT_REF:@.+]] = internal constant %struct.S* [[STAT]] 32 // CHECK-DAG: @out_decl_target = global i32 0, 33 // CHECK-DAG: @llvm.used = appending global [6 x i8*] [i8* bitcast (void ()* @__omp_offloading__{{.+}}_globals_l[[@LINE+69]]_ctor to i8*), i8* bitcast (void ()* @__omp_offloading__{{.+}}_stat_l[[@LINE+70]]_ctor to i8*), 34 // CHECK-DAG: @llvm.compiler.used = appending global [1 x i8*] [i8* bitcast (%struct.S** [[STAT_REF]] to i8*)], 35 36 // CHECK-DAG: define {{.*}}i32 @{{.*}}{{foo|bar|baz2|baz3|FA|f_method}}{{.*}}() 37 // CHECK-DAG: define {{.*}}void @{{.*}}TemplateClass{{.*}}(%class.TemplateClass* %{{.*}}) 38 // CHECK-DAG: define {{.*}}i32 @{{.*}}TemplateClass{{.*}}f_method{{.*}}(%class.TemplateClass* %{{.*}}) 39 // CHECK-DAG: define {{.*}}void @__omp_offloading__{{.*}}_globals_l[[@LINE+63]]_ctor() 40 41 #ifndef HEADER 42 #define HEADER 43 44 #pragma omp declare target 45 extern int aaa; 46 int bbb = 0; 47 extern int ccc; 48 int ddd = 0; 49 #pragma omp end declare target 50 51 extern int eee; 52 int fff = 0; 53 extern int ggg; 54 int hhh = 0; 55 #pragma omp declare target link(eee, fff, ggg, hhh) 56 57 int out_decl_target = 0; 58 #pragma omp declare target 59 void lambda () { 60 #ifdef __cpp_lambdas 61 (void)[&] { (void)out_decl_target; }; 62 #else 63 #pragma clang __debug captured 64 { 65 (void)out_decl_target; 66 } 67 #endif 68 }; 69 #pragma omp end declare target 70 71 template <typename T> 72 class TemplateClass { 73 T a; 74 public: 75 TemplateClass() {} 76 T f_method() const { return a; } 77 }; 78 79 int foo(); 80 81 static int baz1() { return 0; } 82 83 int baz2(); 84 85 int baz4() { return 5; } 86 87 template <typename T> 88 T FA() { 89 TemplateClass<T> s; 90 return s.f_method(); 91 } 92 93 #pragma omp declare target 94 struct S { 95 int a; 96 S(int a) : a(a) {} 97 }; 98 99 int foo() { return 0; } 100 int b = 15; 101 int d; 102 S globals(d); 103 static S stat(d); 104 #pragma omp end declare target 105 int c; 106 107 int bar() { return 1 + foo() + bar() + baz1() + baz2(); } 108 109 int maini1() { 110 int a; 111 static long aa = 32 + bbb + ccc + fff + ggg; 112 // CHECK-DAG: define weak void @__omp_offloading_{{.*}}maini1{{.*}}_l[[@LINE+1]](i32* dereferenceable{{.*}}, i64 {{.*}}, i64 {{.*}}) 113 #pragma omp target map(tofrom \ 114 : a, b) 115 { 116 S s(a); 117 static long aaa = 23; 118 a = foo() + bar() + b + c + d + aa + aaa + FA<int>(); 119 } 120 return baz4(); 121 } 122 123 int baz3() { return 2 + baz2(); } 124 int baz2() { 125 // CHECK-DAG: define weak void @__omp_offloading_{{.*}}baz2{{.*}}_l[[@LINE+1]](i64 {{.*}}) 126 #pragma omp target parallel 127 ++c; 128 return 2 + baz3(); 129 } 130 131 extern int create() throw(); 132 133 static __typeof(create) __t_create __attribute__((__weakref__("__create"))); 134 135 int baz5() { 136 bool a; 137 // CHECK-DAG: define weak void @__omp_offloading_{{.*}}baz5{{.*}}_l[[@LINE+1]](i64 {{.*}}) 138 #pragma omp target 139 a = __extension__(void *) & __t_create != 0; 140 return a; 141 } 142 143 // CHECK-DAG: declare extern_weak signext i32 @__create() 144 145 // CHECK-NOT: define {{.*}}{{baz1|baz4|maini1}} 146 147 // CHECK-DAG: !{i32 1, !"aaa", i32 0, i32 {{[0-9]+}}} 148 // CHECK-DAG: !{i32 1, !"ccc", i32 0, i32 {{[0-9]+}}} 149 150 #endif // HEADER 151