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