shylie/llvm-project
1
0
Fork 0
Code Issues Pull Requests Actions 6 Packages Projects Releases Wiki Activity
llvm-project/clang/test/OpenMP/declare_target_codegen.cpp
Gheorghe-Teodor Bercea 625f59d190 [OpenMP] Fix declare target link implementation 
Summary:
This patch fixes the case where variables in different compilation units or the same compilation unit are under the declare target link clause AND have the same name.
This also fixes the name clash error that occurs when unified memory is activated.
The changes in this patch include:
- Pointers to internal variables are given unique names.
- Externally visible variables are given the same name as before.
- All pointer variables (external or internal) are weakly linked.

Reviewers: ABataev, jdoerfert, caomhin

Reviewed By: ABataev

Subscribers: lebedev.ri, guansong, cfe-commits

Tags: #clang

Differential Revision: https://reviews.llvm.org/D64592

llvm-svn: 367613
2019-08-01 21:15:58 +00:00

228 lines
6.9 KiB
C++
Raw Blame History

// 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
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - | FileCheck %s
// 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
// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o -| FileCheck %s --check-prefix SIMD-ONLY
// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t
// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -verify -o - | FileCheck %s --check-prefix SIMD-ONLY
// expected-no-diagnostics
// SIMD-ONLY-NOT: {{__kmpc|__tgt}}
// 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 ={{ dso_local | }}global i32 0,
// CHECK-DAG: weak constant %struct.__tgt_offload_entry { i8* bitcast (i32* @bbb to i8*),
// CHECK-DAG: @ccc = external global i32,
// CHECK-DAG: @ddd ={{ dso_local | }}global i32 0,
// CHECK-DAG: @hhh_decl_tgt_ref_ptr = weak global i32* null
// CHECK-DAG: @ggg_decl_tgt_ref_ptr = weak global i32* null
// CHECK-DAG: @fff_decl_tgt_ref_ptr = weak global i32* null
// CHECK-DAG: @eee_decl_tgt_ref_ptr = weak global i32* null
// CHECK-DAG: @{{.*}}maini1{{.*}}aaa = internal global i64 23,
// CHECK-DAG: @b ={{ dso_local | }}global i32 15,
// CHECK-DAG: @d ={{ dso_local | }}global i32 0,
// CHECK-DAG: @c = external global i32,
// CHECK-DAG: @globals ={{ dso_local | }}global %struct.S zeroinitializer,
// CHECK-DAG: [[STAT:@.+stat]] = internal global %struct.S zeroinitializer,
// CHECK-DAG: [[STAT_REF:@.+]] = internal constant %struct.S* [[STAT]]
// CHECK-DAG: @out_decl_target ={{ dso_local | }}global i32 0,
// CHECK-DAG: @llvm.used = appending global [2 x i8*] [i8* bitcast (void ()* @__omp_offloading__{{.+}}_globals_l[[@LINE+80]]_ctor to i8*), i8* bitcast (void ()* @__omp_offloading__{{.+}}_stat_l[[@LINE+81]]_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+74]]_ctor()
#ifndef HEADER
#define HEADER
#pragma omp declare target
extern int bbb;
#pragma omp end declare target
#pragma omp declare target
extern int bbb;
#pragma omp end declare target
#pragma omp declare target
extern int aaa;
int bbb = 0;
extern int ccc;
int ddd = 0;
#pragma omp end declare target
#pragma omp declare target
extern int bbb;
#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;
public:
TemplateClass() {}
T f_method() const { return a; }
};
int foo();
static int baz1() { return 0; }
int baz2();
int baz4() { return 5; }
template <typename T>
T FA() {
TemplateClass<T> s;
return s.f_method();
}
#pragma omp declare target
struct S {
int a;
S(int a) : a(a) {}
};
int foo() { return 0; }
int b = 15;
int d;
S globals(d);
static S stat(d);
#pragma omp end declare target
int c;
int bar() { return 1 + foo() + bar() + baz1() + baz2(); }
int maini1() {
int a;
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)
{
S s(a);
static long aaa = 23;
a = foo() + bar() + b + c + d + aa + aaa + FA<int>();
}
return baz4();
}
int baz3() { return 2 + baz2(); }
int baz2() {
// CHECK-DAG: define weak void @__omp_offloading_{{.*}}baz2{{.*}}_l[[@LINE+1]](i64 {{.*}})
#pragma omp target parallel
++c;
return 2 + baz3();
}
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
Reference in New Issue View Git Blame Copy Permalink