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