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|Base|virtual_}} 16 // CHECK-DAG: Bake 17 // CHECK-NOT: @{{hhh|ggg|fff|eee}} = 18 // CHECK-DAG: @aaa = external global i32, 19 // CHECK-DAG: @bbb = global i32 0, 20 // CHECK-DAG: @ccc = external global i32, 21 // CHECK-DAG: @ddd = global i32 0, 22 // CHECK-DAG: @hhh_decl_tgt_link_ptr = common global i32* null 23 // CHECK-DAG: @ggg_decl_tgt_link_ptr = common global i32* null 24 // CHECK-DAG: @fff_decl_tgt_link_ptr = common global i32* null 25 // CHECK-DAG: @eee_decl_tgt_link_ptr = common global i32* null 26 // CHECK-DAG: @{{.*}}maini1{{.*}}aaa = internal global i64 23, 27 // CHECK-DAG: @b = global i32 15, 28 // CHECK-DAG: @d = global i32 0, 29 // CHECK-DAG: @c = external global i32, 30 // CHECK-DAG: @globals = global %struct.S zeroinitializer, 31 // CHECK-DAG: [[STAT:@.+stat]] = internal global %struct.S zeroinitializer, 32 // CHECK-DAG: [[STAT_REF:@.+]] = internal constant %struct.S* [[STAT]] 33 // CHECK-DAG: @out_decl_target = global i32 0, 34 // 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*), 35 // CHECK-DAG: @llvm.compiler.used = appending global [1 x i8*] [i8* bitcast (%struct.S** [[STAT_REF]] to i8*)], 36 37 // CHECK-DAG: define {{.*}}i32 @{{.*}}{{foo|bar|baz2|baz3|FA|f_method}}{{.*}}() 38 // CHECK-DAG: define {{.*}}void @{{.*}}TemplateClass{{.*}}(%class.TemplateClass* %{{.*}}) 39 // CHECK-DAG: define {{.*}}i32 @{{.*}}TemplateClass{{.*}}f_method{{.*}}(%class.TemplateClass* %{{.*}}) 40 // CHECK-DAG: define {{.*}}void @__omp_offloading__{{.*}}_globals_l[[@LINE+63]]_ctor() 41 42 #ifndef HEADER 43 #define HEADER 44 45 #pragma omp declare target 46 extern int aaa; 47 int bbb = 0; 48 extern int ccc; 49 int ddd = 0; 50 #pragma omp end declare target 51 52 extern int eee; 53 int fff = 0; 54 extern int ggg; 55 int hhh = 0; 56 #pragma omp declare target link(eee, fff, ggg, hhh) 57 58 int out_decl_target = 0; 59 #pragma omp declare target 60 void lambda () { 61 #ifdef __cpp_lambdas 62 (void)[&] { (void)out_decl_target; }; 63 #else 64 #pragma clang __debug captured 65 { 66 (void)out_decl_target; 67 } 68 #endif 69 }; 70 #pragma omp end declare target 71 72 template <typename T> 73 class TemplateClass { 74 T a; 75 public: 76 TemplateClass() {} 77 T f_method() const { return a; } 78 }; 79 80 int foo(); 81 82 static int baz1() { return 0; } 83 84 int baz2(); 85 86 int baz4() { return 5; } 87 88 template <typename T> 89 T FA() { 90 TemplateClass<T> s; 91 return s.f_method(); 92 } 93 94 #pragma omp declare target 95 struct S { 96 int a; 97 S(int a) : a(a) {} 98 }; 99 100 int foo() { return 0; } 101 int b = 15; 102 int d; 103 S globals(d); 104 static S stat(d); 105 #pragma omp end declare target 106 int c; 107 108 int bar() { return 1 + foo() + bar() + baz1() + baz2(); } 109 110 int maini1() { 111 int a; 112 static long aa = 32 + bbb + ccc + fff + ggg; 113 // CHECK-DAG: define weak void @__omp_offloading_{{.*}}maini1{{.*}}_l[[@LINE+1]](i32* dereferenceable{{.*}}, i64 {{.*}}, i64 {{.*}}) 114 #pragma omp target map(tofrom \ 115 : a, b) 116 { 117 S s(a); 118 static long aaa = 23; 119 a = foo() + bar() + b + c + d + aa + aaa + FA<int>(); 120 } 121 return baz4(); 122 } 123 124 int baz3() { return 2 + baz2(); } 125 int baz2() { 126 // CHECK-DAG: define weak void @__omp_offloading_{{.*}}baz2{{.*}}_l[[@LINE+1]](i64 {{.*}}) 127 #pragma omp target parallel 128 ++c; 129 return 2 + baz3(); 130 } 131 132 extern int create() throw(); 133 134 static __typeof(create) __t_create __attribute__((__weakref__("__create"))); 135 136 int baz5() { 137 bool a; 138 // CHECK-DAG: define weak void @__omp_offloading_{{.*}}baz5{{.*}}_l[[@LINE+1]](i64 {{.*}}) 139 #pragma omp target 140 a = __extension__(void *) & __t_create != 0; 141 return a; 142 } 143 144 template <typename T> 145 struct Base { 146 virtual ~Base() {} 147 }; 148 149 template class Base<int>; 150 151 template <typename T> 152 struct Bake { 153 virtual ~Bake() {} 154 }; 155 156 #pragma omp declare target 157 template class Bake<int>; 158 #pragma omp end declare target 159 160 struct BaseNonT { 161 virtual ~BaseNonT() {} 162 }; 163 164 #pragma omp declare target 165 struct BakeNonT { 166 virtual ~BakeNonT() {} 167 }; 168 #pragma omp end declare target 169 170 template <typename T> 171 struct B { 172 virtual void virtual_foo(); 173 }; 174 175 void new_bar() { new B<int>(); } 176 177 template <typename T> 178 void B<T>::virtual_foo() { 179 #pragma omp target 180 {} 181 } 182 183 struct A { 184 virtual void emitted() {} 185 }; 186 187 template <typename T> 188 struct C : public A { 189 virtual void emitted(); 190 }; 191 192 template <typename T> 193 void C<T>::emitted() { 194 #pragma omp target 195 {} 196 } 197 198 int main() { 199 A *X = new C<int>(); 200 X->emitted(); 201 return 0; 202 } 203 204 // CHECK-DAG: define {{.*}}void @__omp_offloading_{{.*}}virtual_foo{{.*}}_l[[@LINE-25]]() 205 // CHECK-DAG: define {{.*}}void @__omp_offloading_{{.*}}emitted{{.*}}_l[[@LINE-11]]() 206 207 // CHECK-DAG: declare extern_weak signext i32 @__create() 208 209 // CHECK-NOT: define {{.*}}{{baz1|baz4|maini1|Base|virtual_}} 210 211 // CHECK-DAG: !{i32 1, !"aaa", i32 0, i32 {{[0-9]+}}} 212 // CHECK-DAG: !{i32 1, !"ccc", i32 0, i32 {{[0-9]+}}} 213 // CHECK-DAG: !{{{.+}}virtual_foo 214 215 #endif // HEADER 216