clang/test/OpenMP/declare_target_codegen.cpp

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