forked from OSchip/llvm-project
[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
This commit is contained in:
Gheorghe-Teodor Bercea
parent
8560ea5534
commit
625f59d190
|
|
@ -2552,6 +2552,32 @@ CGOpenMPRuntime::createDispatchNextFunction(unsigned IVSize, bool IVSigned) {
|
|||
return CGM.CreateRuntimeFunction(FnTy, Name);
|
||||
}
|
||||
|
||||
/// Obtain information that uniquely identifies a target entry. This
|
||||
/// consists of the file and device IDs as well as line number associated with
|
||||
/// the relevant entry source location.
|
||||
static void getTargetEntryUniqueInfo(ASTContext &C, SourceLocation Loc,
|
||||
unsigned &DeviceID, unsigned &FileID,
|
||||
unsigned &LineNum) {
|
||||
SourceManager &SM = C.getSourceManager();
|
||||
|
||||
// The loc should be always valid and have a file ID (the user cannot use
|
||||
// #pragma directives in macros)
|
||||
|
||||
assert(Loc.isValid() && "Source location is expected to be always valid.");
|
||||
|
||||
PresumedLoc PLoc = SM.getPresumedLoc(Loc);
|
||||
assert(PLoc.isValid() && "Source location is expected to be always valid.");
|
||||
|
||||
llvm::sys::fs::UniqueID ID;
|
||||
if (auto EC = llvm::sys::fs::getUniqueID(PLoc.getFilename(), ID))
|
||||
SM.getDiagnostics().Report(diag::err_cannot_open_file)
|
||||
<< PLoc.getFilename() << EC.message();
|
||||
|
||||
DeviceID = ID.getDevice();
|
||||
FileID = ID.getFile();
|
||||
LineNum = PLoc.getLine();
|
||||
}
|
||||
|
||||
Address CGOpenMPRuntime::getAddrOfDeclareTargetVar(const VarDecl *VD) {
|
||||
if (CGM.getLangOpts().OpenMPSimd)
|
||||
return Address::invalid();
|
||||
|
|
@ -2563,19 +2589,27 @@ Address CGOpenMPRuntime::getAddrOfDeclareTargetVar(const VarDecl *VD) {
|
|||
SmallString<64> PtrName;
|
||||
{
|
||||
llvm::raw_svector_ostream OS(PtrName);
|
||||
OS << CGM.getMangledName(GlobalDecl(VD)) << "_decl_tgt_ref_ptr";
|
||||
OS << CGM.getMangledName(GlobalDecl(VD));
|
||||
if (!VD->isExternallyVisible()) {
|
||||
unsigned DeviceID, FileID, Line;
|
||||
getTargetEntryUniqueInfo(CGM.getContext(),
|
||||
VD->getCanonicalDecl()->getBeginLoc(),
|
||||
DeviceID, FileID, Line);
|
||||
OS << llvm::format("_%x", FileID);
|
||||
}
|
||||
OS << "_decl_tgt_ref_ptr";
|
||||
}
|
||||
llvm::Value *Ptr = CGM.getModule().getNamedValue(PtrName);
|
||||
if (!Ptr) {
|
||||
QualType PtrTy = CGM.getContext().getPointerType(VD->getType());
|
||||
Ptr = getOrCreateInternalVariable(CGM.getTypes().ConvertTypeForMem(PtrTy),
|
||||
PtrName);
|
||||
if (!CGM.getLangOpts().OpenMPIsDevice) {
|
||||
auto *GV = cast<llvm::GlobalVariable>(Ptr);
|
||||
GV->setLinkage(llvm::GlobalValue::ExternalLinkage);
|
||||
|
||||
auto *GV = cast<llvm::GlobalVariable>(Ptr);
|
||||
GV->setLinkage(llvm::GlobalValue::WeakAnyLinkage);
|
||||
|
||||
if (!CGM.getLangOpts().OpenMPIsDevice)
|
||||
GV->setInitializer(CGM.GetAddrOfGlobal(VD));
|
||||
}
|
||||
CGM.addUsedGlobal(cast<llvm::GlobalValue>(Ptr));
|
||||
registerTargetGlobalVariable(VD, cast<llvm::Constant>(Ptr));
|
||||
}
|
||||
return Address(Ptr, CGM.getContext().getDeclAlign(VD));
|
||||
|
|
@ -2749,32 +2783,6 @@ llvm::Function *CGOpenMPRuntime::emitThreadPrivateVarDefinition(
|
|||
return nullptr;
|
||||
}
|
||||
|
||||
/// Obtain information that uniquely identifies a target entry. This
|
||||
/// consists of the file and device IDs as well as line number associated with
|
||||
/// the relevant entry source location.
|
||||
static void getTargetEntryUniqueInfo(ASTContext &C, SourceLocation Loc,
|
||||
unsigned &DeviceID, unsigned &FileID,
|
||||
unsigned &LineNum) {
|
||||
SourceManager &SM = C.getSourceManager();
|
||||
|
||||
// The loc should be always valid and have a file ID (the user cannot use
|
||||
// #pragma directives in macros)
|
||||
|
||||
assert(Loc.isValid() && "Source location is expected to be always valid.");
|
||||
|
||||
PresumedLoc PLoc = SM.getPresumedLoc(Loc);
|
||||
assert(PLoc.isValid() && "Source location is expected to be always valid.");
|
||||
|
||||
llvm::sys::fs::UniqueID ID;
|
||||
if (auto EC = llvm::sys::fs::getUniqueID(PLoc.getFilename(), ID))
|
||||
SM.getDiagnostics().Report(diag::err_cannot_open_file)
|
||||
<< PLoc.getFilename() << EC.message();
|
||||
|
||||
DeviceID = ID.getDevice();
|
||||
FileID = ID.getFile();
|
||||
LineNum = PLoc.getLine();
|
||||
}
|
||||
|
||||
bool CGOpenMPRuntime::emitDeclareTargetVarDefinition(const VarDecl *VD,
|
||||
llvm::GlobalVariable *Addr,
|
||||
bool PerformInit) {
|
||||
|
|
|
|||
|
|
@ -20,10 +20,10 @@
|
|||
// 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 = common global i32* null
|
||||
// CHECK-DAG: @ggg_decl_tgt_ref_ptr = common global i32* null
|
||||
// CHECK-DAG: @fff_decl_tgt_ref_ptr = common global i32* null
|
||||
// CHECK-DAG: @eee_decl_tgt_ref_ptr = common global i32* null
|
||||
// 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,
|
||||
|
|
@ -32,7 +32,7 @@
|
|||
// 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 [6 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.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}}{{.*}}()
|
||||
|
|
|
|||
|
|
@ -18,24 +18,31 @@
|
|||
#define HEADER
|
||||
|
||||
// HOST-DAG: @c = external global i32,
|
||||
// HOST-DAG: @c_decl_tgt_ref_ptr = global i32* @c
|
||||
// HOST-DAG: @c_decl_tgt_ref_ptr = weak global i32* @c
|
||||
// HOST-DAG: @[[D:.+]] = internal global i32 2
|
||||
// HOST-DAG: @[[D_PTR:.+]] = weak global i32* @[[D]]
|
||||
// DEVICE-NOT: @c =
|
||||
// DEVICE: @c_decl_tgt_ref_ptr = common global i32* null
|
||||
// HOST: [[SIZES:@.+]] = private unnamed_addr constant [2 x i64] [i64 4, i64 4]
|
||||
// HOST: [[MAPTYPES:@.+]] = private unnamed_addr constant [2 x i64] [i64 35, i64 531]
|
||||
// DEVICE: @c_decl_tgt_ref_ptr = weak global i32* null
|
||||
// HOST: [[SIZES:@.+]] = private unnamed_addr constant [3 x i64] [i64 4, i64 4, i64 4]
|
||||
// HOST: [[MAPTYPES:@.+]] = private unnamed_addr constant [3 x i64] [i64 35, i64 531, i64 531]
|
||||
// HOST: @.omp_offloading.entry_name{{.*}} = internal unnamed_addr constant [{{[0-9]+}} x i8] c"c_decl_tgt_ref_ptr\00"
|
||||
// HOST: @.omp_offloading.entry.c_decl_tgt_ref_ptr = weak constant %struct.__tgt_offload_entry { i8* bitcast (i32** @c_decl_tgt_ref_ptr to i8*), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @.omp_offloading.entry_name, i32 0, i32 0), i64 8, i32 1, i32 0 }, section ".omp_offloading.entries", align 1
|
||||
// DEVICE-NOT: internal unnamed_addr constant [{{[0-9]+}} x i8] c"c_decl_tgt_ref_ptr\00"
|
||||
// CHECK: @llvm.used = appending global [1 x i8*] [i8* bitcast (i32** @c_decl_tgt_ref_ptr to i8*)]
|
||||
// DEVICE-NOT: internal unnamed_addr constant [{{[0-9]+}} x i8] c"c_{{.*}}_decl_tgt_ref_ptr\00"
|
||||
// HOST: @.omp_offloading.entry_name{{.*}} = internal unnamed_addr constant [{{[0-9]+}} x i8] c"_{{.*}}d_{{.*}}_decl_tgt_ref_ptr\00"
|
||||
// HOST: @.omp_offloading.entry.[[D_PTR]] = weak constant %struct.__tgt_offload_entry { i8* bitcast (i32** @[[D_PTR]] to i8*), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @.omp_offloading.entry_name{{.*}}, i32 0, i32 0
|
||||
|
||||
extern int c;
|
||||
#pragma omp declare target link(c)
|
||||
|
||||
static int d = 2;
|
||||
#pragma omp declare target link(d)
|
||||
|
||||
int maini1() {
|
||||
int a;
|
||||
#pragma omp target map(tofrom : a)
|
||||
{
|
||||
a = c;
|
||||
d++;
|
||||
}
|
||||
#pragma omp target
|
||||
#pragma omp teams
|
||||
|
|
@ -43,29 +50,38 @@ int maini1() {
|
|||
return 0;
|
||||
}
|
||||
|
||||
// DEVICE: define weak void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-10]](i32* dereferenceable{{[^,]*}}
|
||||
// DEVICE: define weak void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l42(i32* dereferenceable{{[^,]*}}
|
||||
// DEVICE: [[C_REF:%.+]] = load i32*, i32** @c_decl_tgt_ref_ptr,
|
||||
// DEVICE: [[C:%.+]] = load i32, i32* [[C_REF]],
|
||||
// DEVICE: store i32 [[C]], i32* %
|
||||
|
||||
// HOST: define {{.*}}i32 @{{.*}}maini1{{.*}}()
|
||||
// HOST: [[BASEPTRS:%.+]] = alloca [2 x i8*],
|
||||
// HOST: [[PTRS:%.+]] = alloca [2 x i8*],
|
||||
// HOST: getelementptr inbounds [2 x i8*], [2 x i8*]* [[BASEPTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
|
||||
// HOST: getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
|
||||
// HOST: [[BP1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BASEPTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
|
||||
// HOST: [[BASEPTRS:%.+]] = alloca [3 x i8*],
|
||||
// HOST: [[PTRS:%.+]] = alloca [3 x i8*],
|
||||
// HOST: getelementptr inbounds [3 x i8*], [3 x i8*]* [[BASEPTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
|
||||
// HOST: getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
|
||||
|
||||
// HOST: [[BP1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BASEPTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
|
||||
// HOST: [[BP1_CAST:%.+]] = bitcast i8** [[BP1]] to i32***
|
||||
// HOST: store i32** @c_decl_tgt_ref_ptr, i32*** [[BP1_CAST]],
|
||||
// HOST: [[P1:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
|
||||
// HOST: [[P1:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 1
|
||||
// HOST: [[P1_CAST:%.+]] = bitcast i8** [[P1]] to i32**
|
||||
// HOST: store i32* @c, i32** [[P1_CAST]],
|
||||
// HOST: [[BP0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[BASEPTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
|
||||
// HOST: [[P0:%.+]] = getelementptr inbounds [2 x i8*], [2 x i8*]* [[PTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
|
||||
// HOST: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 2, i8** [[BP0]], i8** [[P0]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[SIZES]], i{{[0-9]+}} 0, i{{[0-9]+}} 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[MAPTYPES]], i{{[0-9]+}} 0, i{{[0-9]+}} 0))
|
||||
// HOST: call void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-29]](i32* %{{[^,]+}})
|
||||
// HOST: call i32 @__tgt_target_teams(i64 -1, i8* @.__omp_offloading_{{.+}}_l40.region_id, i32 2, {{.+}})
|
||||
|
||||
// HOST: define internal void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-32]](i32* dereferenceable{{.*}})
|
||||
// HOST: [[BP2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BASEPTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
|
||||
// HOST: [[BP2_CAST:%.+]] = bitcast i8** [[BP2]] to i32***
|
||||
// HOST: store i32** @[[D_PTR]], i32*** [[BP2_CAST]],
|
||||
// HOST: [[P2:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 2
|
||||
// HOST: [[P2_CAST:%.+]] = bitcast i8** [[P2]] to i32**
|
||||
// HOST: store i32* @[[D]], i32** [[P2_CAST]],
|
||||
|
||||
// HOST: [[BP0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[BASEPTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
|
||||
// HOST: [[P0:%.+]] = getelementptr inbounds [3 x i8*], [3 x i8*]* [[PTRS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
|
||||
// HOST: call i32 @__tgt_target(i64 -1, i8* @{{[^,]+}}, i32 3, i8** [[BP0]], i8** [[P0]], i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[SIZES]], i{{[0-9]+}} 0, i{{[0-9]+}} 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* [[MAPTYPES]], i{{[0-9]+}} 0, i{{[0-9]+}} 0))
|
||||
// HOST: call void @__omp_offloading_{{.*}}_{{.*}}_{{.*}}maini1{{.*}}_l42(i32* %{{[^,]+}})
|
||||
// HOST: call i32 @__tgt_target_teams(i64 -1, i8* @.__omp_offloading_{{.+}}_l47.region_id, i32 2, {{.+}})
|
||||
|
||||
// HOST: define internal void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l42(i32* dereferenceable{{.*}})
|
||||
// HOST: [[C:%.*]] = load i32, i32* @c,
|
||||
// HOST: store i32 [[C]], i32* %
|
||||
|
||||
|
|
|
|||
|
|
@ -31,10 +31,10 @@ int bar(int n){
|
|||
}
|
||||
|
||||
// CHECK-HOST: [[VAR:@.+]] = global double 1.000000e+01
|
||||
// CHECK-HOST: [[VAR_DECL_TGT_LINK_PTR:@.+]] = global double* [[VAR]]
|
||||
// CHECK-HOST: [[VAR_DECL_TGT_LINK_PTR:@.+]] = weak global double* [[VAR]]
|
||||
|
||||
// CHECK-HOST: [[TO_VAR:@.+]] = global double 2.000000e+01
|
||||
// CHECK-HOST: [[VAR_DECL_TGT_TO_PTR:@.+]] = global double* [[TO_VAR]]
|
||||
// CHECK-HOST: [[VAR_DECL_TGT_TO_PTR:@.+]] = weak global double* [[TO_VAR]]
|
||||
|
||||
// CHECK-HOST: [[OFFLOAD_SIZES:@.+]] = private unnamed_addr constant [2 x i64] [i64 4, i64 8]
|
||||
// CHECK-HOST: [[OFFLOAD_MAPTYPES:@.+]] = private unnamed_addr constant [2 x i64] [i64 800, i64 800]
|
||||
|
|
@ -45,8 +45,6 @@ int bar(int n){
|
|||
// CHECK-HOST: [[OMP_OFFLOAD_ENTRY_TO_VAR_PTR_NAME:@.+]] = internal unnamed_addr constant [24 x i8]
|
||||
// CHECK-HOST: [[OMP_OFFLOAD_ENTRY_TO_VAR_PTR:@.+]] = weak constant %struct.__tgt_offload_entry { i8* bitcast (double** [[VAR_DECL_TGT_TO_PTR]] to i8*), i8* getelementptr inbounds ([24 x i8], [24 x i8]* [[OMP_OFFLOAD_ENTRY_TO_VAR_PTR_NAME]], i32 0, i32 0), i64 8, i32 0, i32 0 }, section ".omp_offloading.entries"
|
||||
|
||||
// CHECK-HOST: @llvm.used = appending global [2 x i8*] [i8* bitcast (double** [[VAR_DECL_TGT_LINK_PTR]] to i8*), i8* bitcast (double** [[VAR_DECL_TGT_TO_PTR]] to i8*)], section "llvm.metadata"
|
||||
|
||||
// CHECK-HOST: [[N_CASTED:%.+]] = alloca i64
|
||||
// CHECK-HOST: [[SUM_CASTED:%.+]] = alloca i64
|
||||
|
||||
|
|
@ -75,10 +73,8 @@ int bar(int n){
|
|||
|
||||
// CHECK-HOST: call i32 @__tgt_target(i64 -1, i8* @{{.*}}.region_id, i32 2, i8** [[BPTR7]], i8** [[BPTR8]], i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[OFFLOAD_SIZES]], i32 0, i32 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* [[OFFLOAD_MAPTYPES]], i32 0, i32 0))
|
||||
|
||||
// CHECK-DEVICE: [[VAR_LINK:@.+]] = common global double* null
|
||||
// CHECK-DEVICE: [[VAR_TO:@.+]] = common global double* null
|
||||
|
||||
// CHECK-DEVICE: @llvm.used = appending global [2 x i8*] [i8* bitcast (double** [[VAR_LINK]] to i8*), i8* bitcast (double** [[VAR_TO]] to i8*)], section "llvm.metadata"
|
||||
// CHECK-DEVICE: [[VAR_LINK:@.+]] = weak global double* null
|
||||
// CHECK-DEVICE: [[VAR_TO:@.+]] = weak global double* null
|
||||
|
||||
// CHECK-DEVICE: [[VAR_TO_PTR:%.+]] = load double*, double** [[VAR_TO]]
|
||||
// CHECK-DEVICE: load double, double* [[VAR_TO_PTR]]
|
||||
|
|
|
|||
Loading…
Reference in New Issue