forked from OSchip/llvm-project
[OPENMP] Codegen for declare target with link clause.
If the link clause is used on the declare target directive, the object should be linked on target or target data directives, not during the codegen. Patch adds support for this clause. llvm-svn: 328544
This commit is contained in:
Alexey Bataev
parent
3f72a6b7a1
commit
92327c50d3
|
|
@ -9494,10 +9494,13 @@ bool ASTContext::DeclMustBeEmitted(const Decl *D) {
|
|||
return true;
|
||||
|
||||
// If the decl is marked as `declare target`, it should be emitted.
|
||||
for (const auto *Decl = D->getMostRecentDecl(); Decl;
|
||||
Decl = Decl->getPreviousDecl())
|
||||
if (Decl->hasAttr<OMPDeclareTargetDeclAttr>())
|
||||
return true;
|
||||
for (const auto *Decl : D->redecls()) {
|
||||
if (!Decl->hasAttrs())
|
||||
continue;
|
||||
if (const auto *Attr = Decl->getAttr<OMPDeclareTargetDeclAttr>())
|
||||
if (Attr->getMapType() != OMPDeclareTargetDeclAttr::MT_Link)
|
||||
return true;
|
||||
}
|
||||
|
||||
return false;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -2210,6 +2210,22 @@ static LValue EmitThreadPrivateVarDeclLValue(
|
|||
return CGF.MakeAddrLValue(Addr, T, AlignmentSource::Decl);
|
||||
}
|
||||
|
||||
static Address emitDeclTargetLinkVarDeclLValue(CodeGenFunction &CGF,
|
||||
const VarDecl *VD, QualType T) {
|
||||
for (const auto *D : VD->redecls()) {
|
||||
if (!VD->hasAttrs())
|
||||
continue;
|
||||
if (const auto *Attr = D->getAttr<OMPDeclareTargetDeclAttr>())
|
||||
if (Attr->getMapType() == OMPDeclareTargetDeclAttr::MT_Link) {
|
||||
QualType PtrTy = CGF.getContext().getPointerType(VD->getType());
|
||||
Address Addr =
|
||||
CGF.CGM.getOpenMPRuntime().getAddrOfDeclareTargetLink(CGF, VD);
|
||||
return CGF.EmitLoadOfPointer(Addr, PtrTy->castAs<PointerType>());
|
||||
}
|
||||
}
|
||||
return Address::invalid();
|
||||
}
|
||||
|
||||
Address
|
||||
CodeGenFunction::EmitLoadOfReference(LValue RefLVal,
|
||||
LValueBaseInfo *PointeeBaseInfo,
|
||||
|
|
@ -2259,6 +2275,13 @@ static LValue EmitGlobalVarDeclLValue(CodeGenFunction &CGF,
|
|||
if (VD->getTLSKind() == VarDecl::TLS_Dynamic &&
|
||||
CGF.CGM.getCXXABI().usesThreadWrapperFunction())
|
||||
return CGF.CGM.getCXXABI().EmitThreadLocalVarDeclLValue(CGF, VD, T);
|
||||
// Check if the variable is marked as declare target with link clause in
|
||||
// device codegen.
|
||||
if (CGF.getLangOpts().OpenMPIsDevice) {
|
||||
Address Addr = emitDeclTargetLinkVarDeclLValue(CGF, VD, T);
|
||||
if (Addr.isValid())
|
||||
return CGF.MakeAddrLValue(Addr, T, AlignmentSource::Decl);
|
||||
}
|
||||
|
||||
llvm::Value *V = CGF.CGM.GetAddrOfGlobalVar(VD);
|
||||
llvm::Type *RealVarTy = CGF.getTypes().ConvertTypeForMem(VD->getType());
|
||||
|
|
|
|||
|
|
@ -893,6 +893,17 @@ static void EmitOMPAggregateInit(CodeGenFunction &CGF, Address DestAddr,
|
|||
CGF.EmitBlock(DoneBB, /*IsFinished=*/true);
|
||||
}
|
||||
|
||||
static llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy>
|
||||
isDeclareTargetDeclaration(const ValueDecl *VD) {
|
||||
for (const auto *D : VD->redecls()) {
|
||||
if (!D->hasAttrs())
|
||||
continue;
|
||||
if (const auto *Attr = D->getAttr<OMPDeclareTargetDeclAttr>())
|
||||
return Attr->getMapType();
|
||||
}
|
||||
return llvm::None;
|
||||
}
|
||||
|
||||
LValue ReductionCodeGen::emitSharedLValue(CodeGenFunction &CGF, const Expr *E) {
|
||||
return CGF.EmitOMPSharedLValue(E);
|
||||
}
|
||||
|
|
@ -2326,6 +2337,28 @@ llvm::Constant *CGOpenMPRuntime::createDispatchNextFunction(unsigned IVSize,
|
|||
return CGM.CreateRuntimeFunction(FnTy, Name);
|
||||
}
|
||||
|
||||
Address CGOpenMPRuntime::getAddrOfDeclareTargetLink(CodeGenFunction &CGF,
|
||||
const VarDecl *VD) {
|
||||
llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
|
||||
isDeclareTargetDeclaration(VD);
|
||||
if (Res && *Res == OMPDeclareTargetDeclAttr::MT_Link) {
|
||||
SmallString<64> PtrName;
|
||||
{
|
||||
llvm::raw_svector_ostream OS(PtrName);
|
||||
OS << CGM.getMangledName(GlobalDecl(VD)) << "_decl_tgt_link_ptr";
|
||||
}
|
||||
llvm::Value *Ptr = CGM.getModule().getNamedValue(PtrName);
|
||||
if (!Ptr) {
|
||||
QualType PtrTy = CGM.getContext().getPointerType(VD->getType());
|
||||
Ptr = getOrCreateInternalVariable(CGM.getTypes().ConvertTypeForMem(PtrTy),
|
||||
PtrName);
|
||||
CGF.CGM.addUsedGlobal(cast<llvm::GlobalValue>(Ptr));
|
||||
}
|
||||
return Address(Ptr, CGM.getContext().getDeclAlign(VD));
|
||||
}
|
||||
return Address::invalid();
|
||||
}
|
||||
|
||||
llvm::Constant *
|
||||
CGOpenMPRuntime::getOrCreateThreadPrivateCache(const VarDecl *VD) {
|
||||
assert(!CGM.getLangOpts().OpenMPUseTLS ||
|
||||
|
|
@ -6320,6 +6353,50 @@ private:
|
|||
return ConstLength.getSExtValue() != 1;
|
||||
}
|
||||
|
||||
/// \brief Return the adjusted map modifiers if the declaration a capture
|
||||
/// refers to appears in a first-private clause. This is expected to be used
|
||||
/// only with directives that start with 'target'.
|
||||
unsigned adjustMapModifiersForPrivateClauses(const CapturedStmt::Capture &Cap,
|
||||
unsigned CurrentModifiers) {
|
||||
assert(Cap.capturesVariable() && "Expected capture by reference only!");
|
||||
|
||||
// A first private variable captured by reference will use only the
|
||||
// 'private ptr' and 'map to' flag. Return the right flags if the captured
|
||||
// declaration is known as first-private in this handler.
|
||||
if (FirstPrivateDecls.count(Cap.getCapturedVar()))
|
||||
return MappableExprsHandler::OMP_MAP_PRIVATE |
|
||||
MappableExprsHandler::OMP_MAP_TO;
|
||||
// Reduction variable will use only the 'private ptr' and 'map to_from'
|
||||
// flag.
|
||||
if (ReductionDecls.count(Cap.getCapturedVar())) {
|
||||
return MappableExprsHandler::OMP_MAP_TO |
|
||||
MappableExprsHandler::OMP_MAP_FROM;
|
||||
}
|
||||
|
||||
// We didn't modify anything.
|
||||
return CurrentModifiers;
|
||||
}
|
||||
|
||||
public:
|
||||
MappableExprsHandler(const OMPExecutableDirective &Dir, CodeGenFunction &CGF)
|
||||
: CurDir(Dir), CGF(CGF) {
|
||||
// Extract firstprivate clause information.
|
||||
for (const auto *C : Dir.getClausesOfKind<OMPFirstprivateClause>())
|
||||
for (const auto *D : C->varlists())
|
||||
FirstPrivateDecls.insert(
|
||||
cast<VarDecl>(cast<DeclRefExpr>(D)->getDecl())->getCanonicalDecl());
|
||||
for (const auto *C : Dir.getClausesOfKind<OMPReductionClause>()) {
|
||||
for (const auto *D : C->varlists()) {
|
||||
ReductionDecls.insert(
|
||||
cast<VarDecl>(cast<DeclRefExpr>(D)->getDecl())->getCanonicalDecl());
|
||||
}
|
||||
}
|
||||
// Extract device pointer clause information.
|
||||
for (const auto *C : Dir.getClausesOfKind<OMPIsDevicePtrClause>())
|
||||
for (auto L : C->component_lists())
|
||||
DevPointersMap[L.first].push_back(L.second);
|
||||
}
|
||||
|
||||
/// \brief Generate the base pointers, section pointers, sizes and map type
|
||||
/// bits for the provided map type, map modifier, and expression components.
|
||||
/// \a IsFirstComponent should be set to true if the provided set of
|
||||
|
|
@ -6445,6 +6522,7 @@ private:
|
|||
|
||||
// Track if the map information being generated is the first for a capture.
|
||||
bool IsCaptureFirstInfo = IsFirstComponentList;
|
||||
bool IsLink = false; // Is this variable a "declare target link"?
|
||||
|
||||
// Scan the components from the base to the complete expression.
|
||||
auto CI = Components.rbegin();
|
||||
|
|
@ -6464,6 +6542,20 @@ private:
|
|||
// The base is the reference to the variable.
|
||||
// BP = &Var.
|
||||
BP = CGF.EmitOMPSharedLValue(I->getAssociatedExpression()).getPointer();
|
||||
if (const auto *VD =
|
||||
dyn_cast_or_null<VarDecl>(I->getAssociatedDeclaration())) {
|
||||
if (llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
|
||||
isDeclareTargetDeclaration(VD)) {
|
||||
assert(*Res == OMPDeclareTargetDeclAttr::MT_Link &&
|
||||
"Declare target link is expected.");
|
||||
// Avoid warning in release build.
|
||||
(void)*Res;
|
||||
IsLink = true;
|
||||
BP = CGF.CGM.getOpenMPRuntime()
|
||||
.getAddrOfDeclareTargetLink(CGF, VD)
|
||||
.getPointer();
|
||||
}
|
||||
}
|
||||
|
||||
// If the variable is a pointer and is being dereferenced (i.e. is not
|
||||
// the last component), the base has to be the pointer itself, not its
|
||||
|
|
@ -6552,9 +6644,10 @@ private:
|
|||
// same expression except for the first one. We also need to signal
|
||||
// this map is the first one that relates with the current capture
|
||||
// (there is a set of entries for each capture).
|
||||
Types.push_back(DefaultFlags | getMapTypeBits(MapType, MapTypeModifier,
|
||||
!IsExpressionFirstInfo,
|
||||
IsCaptureFirstInfo));
|
||||
Types.push_back(DefaultFlags |
|
||||
getMapTypeBits(MapType, MapTypeModifier,
|
||||
!IsExpressionFirstInfo || IsLink,
|
||||
IsCaptureFirstInfo && !IsLink));
|
||||
|
||||
// If we have a final array section, we are done with this expression.
|
||||
if (IsFinalArraySection)
|
||||
|
|
@ -6570,50 +6663,6 @@ private:
|
|||
}
|
||||
}
|
||||
|
||||
/// \brief Return the adjusted map modifiers if the declaration a capture
|
||||
/// refers to appears in a first-private clause. This is expected to be used
|
||||
/// only with directives that start with 'target'.
|
||||
unsigned adjustMapModifiersForPrivateClauses(const CapturedStmt::Capture &Cap,
|
||||
unsigned CurrentModifiers) {
|
||||
assert(Cap.capturesVariable() && "Expected capture by reference only!");
|
||||
|
||||
// A first private variable captured by reference will use only the
|
||||
// 'private ptr' and 'map to' flag. Return the right flags if the captured
|
||||
// declaration is known as first-private in this handler.
|
||||
if (FirstPrivateDecls.count(Cap.getCapturedVar()))
|
||||
return MappableExprsHandler::OMP_MAP_PRIVATE |
|
||||
MappableExprsHandler::OMP_MAP_TO;
|
||||
// Reduction variable will use only the 'private ptr' and 'map to_from'
|
||||
// flag.
|
||||
if (ReductionDecls.count(Cap.getCapturedVar())) {
|
||||
return MappableExprsHandler::OMP_MAP_TO |
|
||||
MappableExprsHandler::OMP_MAP_FROM;
|
||||
}
|
||||
|
||||
// We didn't modify anything.
|
||||
return CurrentModifiers;
|
||||
}
|
||||
|
||||
public:
|
||||
MappableExprsHandler(const OMPExecutableDirective &Dir, CodeGenFunction &CGF)
|
||||
: CurDir(Dir), CGF(CGF) {
|
||||
// Extract firstprivate clause information.
|
||||
for (const auto *C : Dir.getClausesOfKind<OMPFirstprivateClause>())
|
||||
for (const auto *D : C->varlists())
|
||||
FirstPrivateDecls.insert(
|
||||
cast<VarDecl>(cast<DeclRefExpr>(D)->getDecl())->getCanonicalDecl());
|
||||
for (const auto *C : Dir.getClausesOfKind<OMPReductionClause>()) {
|
||||
for (const auto *D : C->varlists()) {
|
||||
ReductionDecls.insert(
|
||||
cast<VarDecl>(cast<DeclRefExpr>(D)->getDecl())->getCanonicalDecl());
|
||||
}
|
||||
}
|
||||
// Extract device pointer clause information.
|
||||
for (const auto *C : Dir.getClausesOfKind<OMPIsDevicePtrClause>())
|
||||
for (auto L : C->component_lists())
|
||||
DevPointersMap[L.first].push_back(L.second);
|
||||
}
|
||||
|
||||
/// \brief Generate all the base pointers, section pointers, sizes and map
|
||||
/// types for the extracted mappable expressions. Also, for each item that
|
||||
/// relates with a device pointer, a pair of the relevant declaration and
|
||||
|
|
@ -7254,6 +7303,25 @@ void CGOpenMPRuntime::emitTargetCall(CodeGenFunction &CGF,
|
|||
Sizes.append(CurSizes.begin(), CurSizes.end());
|
||||
MapTypes.append(CurMapTypes.begin(), CurMapTypes.end());
|
||||
}
|
||||
// Map other list items in the map clause which are not captured variables
|
||||
// but "declare target link" global variables.
|
||||
for (const auto *C : D.getClausesOfKind<OMPMapClause>()) {
|
||||
for (auto L : C->component_lists()) {
|
||||
if (!L.first)
|
||||
continue;
|
||||
const auto *VD = dyn_cast<VarDecl>(L.first);
|
||||
if (!VD)
|
||||
continue;
|
||||
llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
|
||||
isDeclareTargetDeclaration(VD);
|
||||
if (!Res || *Res != OMPDeclareTargetDeclAttr::MT_Link)
|
||||
continue;
|
||||
MEHandler.generateInfoForComponentList(
|
||||
C->getMapType(), C->getMapTypeModifier(), L.second, BasePointers,
|
||||
Pointers, Sizes, MapTypes, /*IsFirstComponentList=*/true,
|
||||
C->isImplicit());
|
||||
}
|
||||
}
|
||||
|
||||
TargetDataInfo Info;
|
||||
// Fill up the arrays and create the arguments.
|
||||
|
|
@ -7406,14 +7474,7 @@ bool CGOpenMPRuntime::emitTargetFunctions(GlobalDecl GD) {
|
|||
scanForTargetRegionsFunctions(FD.getBody(), CGM.getMangledName(GD));
|
||||
|
||||
// Do not to emit function if it is not marked as declare target.
|
||||
if (!GD.getDecl()->hasAttrs())
|
||||
return true;
|
||||
|
||||
for (const auto *D = FD.getMostRecentDecl(); D; D = D->getPreviousDecl())
|
||||
if (D->hasAttr<OMPDeclareTargetDeclAttr>())
|
||||
return false;
|
||||
|
||||
return true;
|
||||
return !isDeclareTargetDeclaration(&FD);
|
||||
}
|
||||
|
||||
bool CGOpenMPRuntime::emitTargetGlobalVariable(GlobalDecl GD) {
|
||||
|
|
@ -7439,15 +7500,9 @@ bool CGOpenMPRuntime::emitTargetGlobalVariable(GlobalDecl GD) {
|
|||
}
|
||||
|
||||
// Do not to emit variable if it is not marked as declare target.
|
||||
if (!GD.getDecl()->hasAttrs())
|
||||
return true;
|
||||
|
||||
for (const Decl *D = GD.getDecl()->getMostRecentDecl(); D;
|
||||
D = D->getPreviousDecl())
|
||||
if (D->hasAttr<OMPDeclareTargetDeclAttr>())
|
||||
return false;
|
||||
|
||||
return true;
|
||||
llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
|
||||
isDeclareTargetDeclaration(cast<ValueDecl>(GD.getDecl()));
|
||||
return !Res || *Res == OMPDeclareTargetDeclAttr::MT_Link;
|
||||
}
|
||||
|
||||
bool CGOpenMPRuntime::emitTargetGlobal(GlobalDecl GD) {
|
||||
|
|
@ -7477,9 +7532,8 @@ bool CGOpenMPRuntime::markAsGlobalTarget(const FunctionDecl *D) {
|
|||
return true;
|
||||
// Do not to emit function if it is marked as declare target as it was already
|
||||
// emitted.
|
||||
for (const auto *FD = D->getMostRecentDecl(); FD; FD = FD->getPreviousDecl())
|
||||
if (FD->hasAttr<OMPDeclareTargetDeclAttr>())
|
||||
return true;
|
||||
if (isDeclareTargetDeclaration(D))
|
||||
return true;
|
||||
|
||||
const FunctionDecl *FD = D->getCanonicalDecl();
|
||||
// Do not mark member functions except for static.
|
||||
|
|
|
|||
|
|
@ -952,6 +952,11 @@ public:
|
|||
Address VDAddr,
|
||||
SourceLocation Loc);
|
||||
|
||||
/// Returns the address of the variable marked as declare target with link
|
||||
/// clause.
|
||||
virtual Address getAddrOfDeclareTargetLink(CodeGenFunction &CGF,
|
||||
const VarDecl *VD);
|
||||
|
||||
/// \brief Emit a code for initialization of threadprivate variable. It emits
|
||||
/// a call to runtime library which adds initial value to the newly created
|
||||
/// threadprivate variable (if it is not constant) and registers destructor
|
||||
|
|
|
|||
|
|
@ -1215,6 +1215,17 @@ void Sema::popOpenMPFunctionRegion(const FunctionScopeInfo *OldFSI) {
|
|||
DSAStack->popFunction(OldFSI);
|
||||
}
|
||||
|
||||
static llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy>
|
||||
isDeclareTargetDeclaration(const ValueDecl *VD) {
|
||||
for (const auto *D : VD->redecls()) {
|
||||
if (!D->hasAttrs())
|
||||
continue;
|
||||
if (const auto *Attr = D->getAttr<OMPDeclareTargetDeclAttr>())
|
||||
return Attr->getMapType();
|
||||
}
|
||||
return llvm::None;
|
||||
}
|
||||
|
||||
bool Sema::IsOpenMPCapturedByRef(ValueDecl *D, unsigned Level) {
|
||||
assert(LangOpts.OpenMP && "OpenMP is not allowed");
|
||||
|
||||
|
|
@ -1392,10 +1403,8 @@ VarDecl *Sema::IsOpenMPCapturedDecl(ValueDecl *D) {
|
|||
// If the declaration is enclosed in a 'declare target' directive,
|
||||
// then it should not be captured.
|
||||
//
|
||||
for (const auto *Var = VD->getMostRecentDecl(); Var;
|
||||
Var = Var->getPreviousDecl())
|
||||
if (Var->hasAttr<OMPDeclareTargetDeclAttr>())
|
||||
return nullptr;
|
||||
if (isDeclareTargetDeclaration(VD))
|
||||
return nullptr;
|
||||
return VD;
|
||||
}
|
||||
|
||||
|
|
@ -1929,7 +1938,10 @@ public:
|
|||
return;
|
||||
|
||||
// Skip internally declared static variables.
|
||||
if (VD->hasGlobalStorage() && !CS->capturesVariable(VD))
|
||||
llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
|
||||
isDeclareTargetDeclaration(VD);
|
||||
if (VD->hasGlobalStorage() && !CS->capturesVariable(VD) &&
|
||||
(!Res || *Res != OMPDeclareTargetDeclAttr::MT_Link))
|
||||
return;
|
||||
|
||||
auto ELoc = E->getExprLoc();
|
||||
|
|
@ -1976,7 +1988,7 @@ public:
|
|||
IsFirstprivate =
|
||||
IsFirstprivate ||
|
||||
(VD->getType().getNonReferenceType()->isScalarType() &&
|
||||
Stack->getDefaultDMA() != DMA_tofrom_scalar);
|
||||
Stack->getDefaultDMA() != DMA_tofrom_scalar && !Res);
|
||||
if (IsFirstprivate)
|
||||
ImplicitFirstprivate.emplace_back(E);
|
||||
else
|
||||
|
|
|
|||
|
|
@ -0,0 +1,64 @@
|
|||
// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix HOST --check-prefix CHECK
|
||||
// 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 --check-prefix DEVICE --check-prefix CHECK
|
||||
// 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 --check-prefix DEVICE --check-prefix CHECK
|
||||
|
||||
// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -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-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}}
|
||||
|
||||
#ifndef HEADER
|
||||
#define HEADER
|
||||
|
||||
// HOST: @c = external global i32,
|
||||
// DEVICE-NOT: @c =
|
||||
// CHECK: @c_decl_tgt_link_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]
|
||||
// CHECK: @llvm.used = appending global [1 x i8*] [i8* bitcast (i32** @c_decl_tgt_link_ptr to i8*)]
|
||||
|
||||
extern int c;
|
||||
#pragma omp declare target link(c)
|
||||
|
||||
int maini1() {
|
||||
int a;
|
||||
#pragma omp target map(tofrom : a)
|
||||
{
|
||||
a = c;
|
||||
}
|
||||
return 0;
|
||||
}
|
||||
|
||||
// DEVICE: define void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-7]](i32* dereferenceable{{[^,]*}}
|
||||
// DEVICE: [[C_REF:%.+]] = load i32*, i32** @c_decl_tgt_link_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: [[BP1_CAST:%.+]] = bitcast i8** [[BP1]] to i32***
|
||||
// HOST: store i32** @c_decl_tgt_link_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_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-26]](i32* %{{[^,]+}})
|
||||
|
||||
// HOST: define internal void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-28]](i32* dereferenceable{{.*}})
|
||||
// HOST: [[C:%.*]] = load i32, i32* @c,
|
||||
// HOST: store i32 [[C]], i32* %
|
||||
|
||||
#endif // HEADER
|
||||
Loading…
Reference in New Issue