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