summaryrefslogtreecommitdiff
path: root/test/OpenMP/declare_target_codegen.cpp
diff options
context:
space:
mode:
authorDimitry Andric <dim@FreeBSD.org>2019-01-19 10:04:05 +0000
committerDimitry Andric <dim@FreeBSD.org>2019-01-19 10:04:05 +0000
commit676fbe8105eeb6ff4bb2ed261cb212fcfdbe7b63 (patch)
tree02a1ac369cb734d0abfa5000dd86e5b7797e6a74 /test/OpenMP/declare_target_codegen.cpp
parentc7e70c433efc6953dc3888b9fbf9f3512d7da2b0 (diff)
Diffstat (limited to 'test/OpenMP/declare_target_codegen.cpp')
-rw-r--r--test/OpenMP/declare_target_codegen.cpp132
1 files changed, 126 insertions, 6 deletions
diff --git a/test/OpenMP/declare_target_codegen.cpp b/test/OpenMP/declare_target_codegen.cpp
index 96bdc874ece26..cc7525a44b20e 100644
--- a/test/OpenMP/declare_target_codegen.cpp
+++ b/test/OpenMP/declare_target_codegen.cpp
@@ -12,7 +12,17 @@
// SIMD-ONLY-NOT: {{__kmpc|__tgt}}
-// CHECK-NOT: define {{.*}}{{baz1|baz4|maini1}}
+// CHECK-NOT: define {{.*}}{{baz1|baz4|maini1|Base|virtual_}}
+// CHECK-DAG: Bake
+// CHECK-NOT: @{{hhh|ggg|fff|eee}} =
+// CHECK-DAG: @aaa = external global i32,
+// CHECK-DAG: @bbb = global i32 0,
+// CHECK-DAG: @ccc = external global i32,
+// CHECK-DAG: @ddd = global i32 0,
+// CHECK-DAG: @hhh_decl_tgt_link_ptr = common global i32* null
+// CHECK-DAG: @ggg_decl_tgt_link_ptr = common global i32* null
+// CHECK-DAG: @fff_decl_tgt_link_ptr = common global i32* null
+// CHECK-DAG: @eee_decl_tgt_link_ptr = common global i32* null
// CHECK-DAG: @{{.*}}maini1{{.*}}aaa = internal global i64 23,
// CHECK-DAG: @b = global i32 15,
// CHECK-DAG: @d = global i32 0,
@@ -20,17 +30,45 @@
// CHECK-DAG: @globals = global %struct.S zeroinitializer,
// CHECK-DAG: [[STAT:@.+stat]] = internal global %struct.S zeroinitializer,
// CHECK-DAG: [[STAT_REF:@.+]] = internal constant %struct.S* [[STAT]]
-// CHECK-DAG: @llvm.used = appending global [2 x i8*] [i8* bitcast (void ()* @__omp_offloading__{{.+}}_globals_l[[@LINE+42]]_ctor to i8*), i8* bitcast (void ()* @__omp_offloading__{{.+}}_stat_l[[@LINE+43]]_ctor to i8*)],
+// CHECK-DAG: @out_decl_target = global i32 0,
+// CHECK-DAG: @llvm.used = appending global [6 x i8*] [i8* bitcast (void ()* @__omp_offloading__{{.+}}_globals_l[[@LINE+69]]_ctor to i8*), i8* bitcast (void ()* @__omp_offloading__{{.+}}_stat_l[[@LINE+70]]_ctor to i8*),
// CHECK-DAG: @llvm.compiler.used = appending global [1 x i8*] [i8* bitcast (%struct.S** [[STAT_REF]] to i8*)],
// CHECK-DAG: define {{.*}}i32 @{{.*}}{{foo|bar|baz2|baz3|FA|f_method}}{{.*}}()
// CHECK-DAG: define {{.*}}void @{{.*}}TemplateClass{{.*}}(%class.TemplateClass* %{{.*}})
// CHECK-DAG: define {{.*}}i32 @{{.*}}TemplateClass{{.*}}f_method{{.*}}(%class.TemplateClass* %{{.*}})
-// CHECK-DAG: define {{.*}}void @__omp_offloading__{{.*}}_globals_l[[@LINE+36]]_ctor()
+// CHECK-DAG: define {{.*}}void @__omp_offloading__{{.*}}_globals_l[[@LINE+63]]_ctor()
#ifndef HEADER
#define HEADER
+#pragma omp declare target
+extern int aaa;
+int bbb = 0;
+extern int ccc;
+int ddd = 0;
+#pragma omp end declare target
+
+extern int eee;
+int fff = 0;
+extern int ggg;
+int hhh = 0;
+#pragma omp declare target link(eee, fff, ggg, hhh)
+
+int out_decl_target = 0;
+#pragma omp declare target
+void lambda () {
+#ifdef __cpp_lambdas
+ (void)[&] { (void)out_decl_target; };
+#else
+#pragma clang __debug captured
+ {
+ (void)out_decl_target;
+ }
+#endif
+};
+#pragma omp end declare target
+
template <typename T>
class TemplateClass {
T a;
@@ -41,7 +79,7 @@ public:
int foo();
-int baz1();
+static int baz1() { return 0; }
int baz2();
@@ -71,7 +109,7 @@ int bar() { return 1 + foo() + bar() + baz1() + baz2(); }
int maini1() {
int a;
- static long aa = 32;
+ static long aa = 32 + bbb + ccc + fff + ggg;
// CHECK-DAG: define weak void @__omp_offloading_{{.*}}maini1{{.*}}_l[[@LINE+1]](i32* dereferenceable{{.*}}, i64 {{.*}}, i64 {{.*}})
#pragma omp target map(tofrom \
: a, b)
@@ -91,5 +129,87 @@ int baz2() {
return 2 + baz3();
}
-// CHECK-NOT: define {{.*}}{{baz1|baz4|maini1}}
+extern int create() throw();
+
+static __typeof(create) __t_create __attribute__((__weakref__("__create")));
+
+int baz5() {
+ bool a;
+// CHECK-DAG: define weak void @__omp_offloading_{{.*}}baz5{{.*}}_l[[@LINE+1]](i64 {{.*}})
+#pragma omp target
+ a = __extension__(void *) & __t_create != 0;
+ return a;
+}
+
+template <typename T>
+struct Base {
+ virtual ~Base() {}
+};
+
+template class Base<int>;
+
+template <typename T>
+struct Bake {
+ virtual ~Bake() {}
+};
+
+#pragma omp declare target
+template class Bake<int>;
+#pragma omp end declare target
+
+struct BaseNonT {
+ virtual ~BaseNonT() {}
+};
+
+#pragma omp declare target
+struct BakeNonT {
+ virtual ~BakeNonT() {}
+};
+#pragma omp end declare target
+
+template <typename T>
+struct B {
+ virtual void virtual_foo();
+};
+
+void new_bar() { new B<int>(); }
+
+template <typename T>
+void B<T>::virtual_foo() {
+#pragma omp target
+ {}
+}
+
+struct A {
+ virtual void emitted() {}
+};
+
+template <typename T>
+struct C : public A {
+ virtual void emitted();
+};
+
+template <typename T>
+void C<T>::emitted() {
+#pragma omp target
+ {}
+}
+
+int main() {
+ A *X = new C<int>();
+ X->emitted();
+ return 0;
+}
+
+// CHECK-DAG: define {{.*}}void @__omp_offloading_{{.*}}virtual_foo{{.*}}_l[[@LINE-25]]()
+// CHECK-DAG: define {{.*}}void @__omp_offloading_{{.*}}emitted{{.*}}_l[[@LINE-11]]()
+
+// CHECK-DAG: declare extern_weak signext i32 @__create()
+
+// CHECK-NOT: define {{.*}}{{baz1|baz4|maini1|Base|virtual_}}
+
+// CHECK-DAG: !{i32 1, !"aaa", i32 0, i32 {{[0-9]+}}}
+// CHECK-DAG: !{i32 1, !"ccc", i32 0, i32 {{[0-9]+}}}
+// CHECK-DAG: !{{{.+}}virtual_foo
+
#endif // HEADER