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