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