r328705 - [OPENMP] Codegen for ctor|dtor of declare target variables.
Alexey Bataev via cfe-commits
cfe-commits at lists.llvm.org
Wed Mar 28 07:28:54 PDT 2018
Author: abataev
Date: Wed Mar 28 07:28:54 2018
New Revision: 328705
URL: http://llvm.org/viewvc/llvm-project?rev=328705&view=rev
Log:
[OPENMP] Codegen for ctor|dtor of declare target variables.
When the declare target variables are emitted for the device,
constructors|destructors for these variables must emitted and registered
by the runtime in the offloading sections.
Added:
cfe/trunk/test/OpenMP/nvptx_declare_target_var_ctor_dtor_codegen.cpp
Modified:
cfe/trunk/include/clang/Sema/Sema.h
cfe/trunk/lib/CodeGen/CGDeclCXX.cpp
cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp
cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h
cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.h
cfe/trunk/lib/Parse/ParseOpenMP.cpp
Modified: cfe/trunk/include/clang/Sema/Sema.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=328705&r1=328704&r2=328705&view=diff
==============================================================================
--- cfe/trunk/include/clang/Sema/Sema.h (original)
+++ cfe/trunk/include/clang/Sema/Sema.h Wed Mar 28 07:28:54 2018
@@ -546,7 +546,7 @@ public:
/// FieldCollector - Collects CXXFieldDecls during parsing of C++ classes.
std::unique_ptr<CXXFieldCollector> FieldCollector;
- typedef llvm::SmallSetVector<const NamedDecl*, 16> NamedDeclSetType;
+ typedef llvm::SmallSetVector<NamedDecl *, 16> NamedDeclSetType;
/// \brief Set containing all declared private fields that are not used.
NamedDeclSetType UnusedPrivateFields;
Modified: cfe/trunk/lib/CodeGen/CGDeclCXX.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGDeclCXX.cpp?rev=328705&r1=328704&r2=328705&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGDeclCXX.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGDeclCXX.cpp Wed Mar 28 07:28:54 2018
@@ -379,6 +379,10 @@ CodeGenModule::EmitCXXGlobalVarDeclInitF
D->hasAttr<CUDASharedAttr>()))
return;
+ if (getLangOpts().OpenMP &&
+ getOpenMPRuntime().emitDeclareTargetVarDefinition(D, Addr, PerformInit))
+ return;
+
// Check if we've already initialized this decl.
auto I = DelayedCXXInitPosition.find(D);
if (I != DelayedCXXInitPosition.end() && I->second == ~0U)
Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp?rev=328705&r1=328704&r2=328705&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp Wed Mar 28 07:28:54 2018
@@ -2522,6 +2522,139 @@ llvm::Function *CGOpenMPRuntime::emitThr
return nullptr;
}
+/// \brief 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) {
+
+ auto &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.");
+ assert(Loc.isFileID() && "Source location is expected to refer to a file.");
+
+ PresumedLoc PLoc = SM.getPresumedLoc(Loc);
+ assert(PLoc.isValid() && "Source location is expected to be always valid.");
+
+ llvm::sys::fs::UniqueID ID;
+ if (llvm::sys::fs::getUniqueID(PLoc.getFilename(), ID))
+ llvm_unreachable("Source file with target region no longer exists!");
+
+ DeviceID = ID.getDevice();
+ FileID = ID.getFile();
+ LineNum = PLoc.getLine();
+}
+
+bool CGOpenMPRuntime::emitDeclareTargetVarDefinition(const VarDecl *VD,
+ llvm::GlobalVariable *Addr,
+ bool PerformInit) {
+ Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res =
+ isDeclareTargetDeclaration(VD);
+ if (!Res || *Res == OMPDeclareTargetDeclAttr::MT_Link)
+ return false;
+ VD = VD->getDefinition(CGM.getContext());
+ if (VD && !DeclareTargetWithDefinition.insert(VD).second)
+ return CGM.getLangOpts().OpenMPIsDevice;
+
+ QualType ASTTy = VD->getType();
+
+ SourceLocation Loc = VD->getCanonicalDecl()->getLocStart();
+ // Produce the unique prefix to identify the new target regions. We use
+ // the source location of the variable declaration which we know to not
+ // conflict with any target region.
+ unsigned DeviceID;
+ unsigned FileID;
+ unsigned Line;
+ getTargetEntryUniqueInfo(CGM.getContext(), Loc, DeviceID, FileID, Line);
+ SmallString<128> Buffer, Out;
+ {
+ llvm::raw_svector_ostream OS(Buffer);
+ OS << "__omp_offloading_" << llvm::format("_%x", DeviceID)
+ << llvm::format("_%x_", FileID) << VD->getName() << "_l" << Line;
+ }
+
+ const Expr *Init = VD->getAnyInitializer();
+ if (CGM.getLangOpts().CPlusPlus && PerformInit) {
+ llvm::Constant *Ctor;
+ llvm::Constant *ID;
+ if (CGM.getLangOpts().OpenMPIsDevice) {
+ // Generate function that re-emits the declaration's initializer into
+ // the threadprivate copy of the variable VD
+ CodeGenFunction CtorCGF(CGM);
+
+ const CGFunctionInfo &FI = CGM.getTypes().arrangeNullaryFunction();
+ llvm::FunctionType *FTy = CGM.getTypes().GetFunctionType(FI);
+ llvm::Function *Fn = CGM.CreateGlobalInitOrDestructFunction(
+ FTy, Twine(Buffer, "_ctor"), FI, Loc);
+ auto NL = ApplyDebugLocation::CreateEmpty(CtorCGF);
+ CtorCGF.StartFunction(GlobalDecl(), CGM.getContext().VoidTy, Fn, FI,
+ FunctionArgList(), Loc, Loc);
+ auto AL = ApplyDebugLocation::CreateArtificial(CtorCGF);
+ CtorCGF.EmitAnyExprToMem(Init,
+ Address(Addr, CGM.getContext().getDeclAlign(VD)),
+ Init->getType().getQualifiers(),
+ /*IsInitializer=*/true);
+ CtorCGF.FinishFunction();
+ Ctor = Fn;
+ ID = llvm::ConstantExpr::getBitCast(Fn, CGM.Int8PtrTy);
+ } else {
+ Ctor = new llvm::GlobalVariable(
+ CGM.getModule(), CGM.Int8Ty, /*isConstant=*/true,
+ llvm::GlobalValue::PrivateLinkage,
+ llvm::Constant::getNullValue(CGM.Int8Ty), Twine(Buffer, "_ctor"));
+ ID = Ctor;
+ }
+
+ // Register the information for the entry associated with the constructor.
+ Out.clear();
+ OffloadEntriesInfoManager.registerTargetRegionEntryInfo(
+ DeviceID, FileID, Twine(Buffer, "_ctor").toStringRef(Out), Line, Ctor,
+ ID, OMPTargetRegionEntryCtor);
+ }
+ if (VD->getType().isDestructedType() != QualType::DK_none) {
+ llvm::Constant *Dtor;
+ llvm::Constant *ID;
+ if (CGM.getLangOpts().OpenMPIsDevice) {
+ // Generate function that emits destructor call for the threadprivate
+ // copy of the variable VD
+ CodeGenFunction DtorCGF(CGM);
+
+ const CGFunctionInfo &FI = CGM.getTypes().arrangeNullaryFunction();
+ llvm::FunctionType *FTy = CGM.getTypes().GetFunctionType(FI);
+ llvm::Function *Fn = CGM.CreateGlobalInitOrDestructFunction(
+ FTy, Twine(Buffer, "_dtor"), FI, Loc);
+ auto NL = ApplyDebugLocation::CreateEmpty(DtorCGF);
+ DtorCGF.StartFunction(GlobalDecl(), CGM.getContext().VoidTy, Fn, FI,
+ FunctionArgList(), Loc, Loc);
+ // Create a scope with an artificial location for the body of this
+ // function.
+ auto AL = ApplyDebugLocation::CreateArtificial(DtorCGF);
+ DtorCGF.emitDestroy(Address(Addr, CGM.getContext().getDeclAlign(VD)),
+ ASTTy, DtorCGF.getDestroyer(ASTTy.isDestructedType()),
+ DtorCGF.needsEHCleanup(ASTTy.isDestructedType()));
+ DtorCGF.FinishFunction();
+ Dtor = Fn;
+ ID = llvm::ConstantExpr::getBitCast(Fn, CGM.Int8PtrTy);
+ } else {
+ Dtor = new llvm::GlobalVariable(
+ CGM.getModule(), CGM.Int8Ty, /*isConstant=*/true,
+ llvm::GlobalValue::PrivateLinkage,
+ llvm::Constant::getNullValue(CGM.Int8Ty), Twine(Buffer, "_dtor"));
+ ID = Dtor;
+ }
+ // Register the information for the entry associated with the destructor.
+ Out.clear();
+ OffloadEntriesInfoManager.registerTargetRegionEntryInfo(
+ DeviceID, FileID, Twine(Buffer, "_dtor").toStringRef(Out), Line, Dtor,
+ ID, OMPTargetRegionEntryDtor);
+ }
+ return CGM.getLangOpts().OpenMPIsDevice;
+}
+
Address CGOpenMPRuntime::getAddrOfArtificialThreadPrivate(CodeGenFunction &CGF,
QualType VarType,
StringRef Name) {
@@ -3375,7 +3508,7 @@ void CGOpenMPRuntime::OffloadEntriesInfo
"code generation.");
OffloadEntriesTargetRegion[DeviceID][FileID][ParentName][LineNum] =
OffloadEntryInfoTargetRegion(Order, /*Addr=*/nullptr, /*ID=*/nullptr,
- /*Flags=*/0);
+ OMPTargetRegionEntryTargetRegion);
++OffloadingEntriesNum;
}
@@ -3383,7 +3516,7 @@ void CGOpenMPRuntime::OffloadEntriesInfo
registerTargetRegionEntryInfo(unsigned DeviceID, unsigned FileID,
StringRef ParentName, unsigned LineNum,
llvm::Constant *Addr, llvm::Constant *ID,
- int32_t Flags) {
+ OMPTargetRegionEntryKind Flags) {
// If we are emitting code for a target, the entry is already initialized,
// only has to be registered.
if (CGM.getLangOpts().OpenMPIsDevice) {
@@ -3641,12 +3774,12 @@ void CGOpenMPRuntime::createOffloadEntri
llvm::NamedMDNode *MD = M.getOrInsertNamedMetadata("omp_offload.info");
// Auxiliary methods to create metadata values and strings.
- auto getMDInt = [&](unsigned v) {
+ auto GetMdInt = [&C](unsigned V) {
return llvm::ConstantAsMetadata::get(
- llvm::ConstantInt::get(llvm::Type::getInt32Ty(C), v));
+ llvm::ConstantInt::get(llvm::Type::getInt32Ty(C), V));
};
- auto getMDString = [&](StringRef v) { return llvm::MDString::get(C, v); };
+ auto GetMdString = [&C](StringRef V) { return llvm::MDString::get(C, V); };
// Create function that emits metadata for each target region entry;
auto &&TargetRegionMetadataEmitter = [&](
@@ -3662,12 +3795,12 @@ void CGOpenMPRuntime::createOffloadEntri
// - Entry 4 -> Line in the file where the entry was identified.
// - Entry 5 -> Order the entry was created.
// The first element of the metadata node is the kind.
- Ops.push_back(getMDInt(E.getKind()));
- Ops.push_back(getMDInt(DeviceID));
- Ops.push_back(getMDInt(FileID));
- Ops.push_back(getMDString(ParentName));
- Ops.push_back(getMDInt(Line));
- Ops.push_back(getMDInt(E.getOrder()));
+ Ops.push_back(GetMdInt(E.getKind()));
+ Ops.push_back(GetMdInt(DeviceID));
+ Ops.push_back(GetMdInt(FileID));
+ Ops.push_back(GetMdString(ParentName));
+ Ops.push_back(GetMdInt(Line));
+ Ops.push_back(GetMdInt(E.getOrder()));
// Save this entry in the right position of the ordered entries array.
OrderedEntries[E.getOrder()] = &E;
@@ -3686,7 +3819,8 @@ void CGOpenMPRuntime::createOffloadEntri
E)) {
assert(CE->getID() && CE->getAddress() &&
"Entry ID and Addr are invalid!");
- createOffloadEntry(CE->getID(), CE->getAddress(), /*Size=*/0);
+ createOffloadEntry(CE->getID(), CE->getAddress(), /*Size=*/0,
+ CE->getFlags());
} else
llvm_unreachable("Unsupported entry kind.");
}
@@ -3720,27 +3854,27 @@ void CGOpenMPRuntime::loadOffloadInfoMet
return;
for (llvm::MDNode *MN : MD->operands()) {
- auto getMDInt = [&](unsigned Idx) {
+ auto GetMdInt = [MN](unsigned Idx) {
llvm::ConstantAsMetadata *V =
cast<llvm::ConstantAsMetadata>(MN->getOperand(Idx));
return cast<llvm::ConstantInt>(V->getValue())->getZExtValue();
};
- auto getMDString = [&](unsigned Idx) {
+ auto GetMdString = [MN](unsigned Idx) {
llvm::MDString *V = cast<llvm::MDString>(MN->getOperand(Idx));
return V->getString();
};
- switch (getMDInt(0)) {
+ switch (GetMdInt(0)) {
default:
llvm_unreachable("Unexpected metadata!");
break;
case OffloadEntriesInfoManagerTy::OffloadEntryInfo::
- OFFLOAD_ENTRY_INFO_TARGET_REGION:
+ OffloadingEntryInfoTargetRegion:
OffloadEntriesInfoManager.initializeTargetRegionEntryInfo(
- /*DeviceID=*/getMDInt(1), /*FileID=*/getMDInt(2),
- /*ParentName=*/getMDString(3), /*Line=*/getMDInt(4),
- /*Order=*/getMDInt(5));
+ /*DeviceID=*/GetMdInt(1), /*FileID=*/GetMdInt(2),
+ /*ParentName=*/GetMdString(3), /*Line=*/GetMdInt(4),
+ /*Order=*/GetMdInt(5));
break;
}
}
@@ -5871,33 +6005,6 @@ void CGOpenMPRuntime::emitCancelCall(Cod
}
}
-/// \brief 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) {
-
- auto &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.");
- assert(Loc.isFileID() && "Source location is expected to refer to a file.");
-
- PresumedLoc PLoc = SM.getPresumedLoc(Loc);
- assert(PLoc.isValid() && "Source location is expected to be always valid.");
-
- llvm::sys::fs::UniqueID ID;
- if (llvm::sys::fs::getUniqueID(PLoc.getFilename(), ID))
- llvm_unreachable("Source file with target region no longer exists!");
-
- DeviceID = ID.getDevice();
- FileID = ID.getFile();
- LineNum = PLoc.getLine();
-}
-
void CGOpenMPRuntime::emitTargetOutlinedFunction(
const OMPExecutableDirective &D, StringRef ParentName,
llvm::Function *&OutlinedFn, llvm::Constant *&OutlinedFnID,
@@ -5970,7 +6077,7 @@ void CGOpenMPRuntime::emitTargetOutlined
// Register the information for the entry associated with this target region.
OffloadEntriesInfoManager.registerTargetRegionEntryInfo(
DeviceID, FileID, ParentName, Line, OutlinedFn, OutlinedFnID,
- /*Flags=*/0);
+ OMPTargetRegionEntryTargetRegion);
}
/// discard all CompoundStmts intervening between two constructs
@@ -7530,12 +7637,20 @@ CGOpenMPRuntime::DisableAutoDeclareTarge
bool CGOpenMPRuntime::markAsGlobalTarget(const FunctionDecl *D) {
if (!CGM.getLangOpts().OpenMPIsDevice || !ShouldMarkAsGlobal)
return true;
+
+ const FunctionDecl *FD = D->getCanonicalDecl();
// Do not to emit function if it is marked as declare target as it was already
// emitted.
- if (isDeclareTargetDeclaration(D))
+ if (isDeclareTargetDeclaration(D)) {
+ if (D->hasBody() && AlreadyEmittedTargetFunctions.count(FD) == 0) {
+ if (auto *F = dyn_cast_or_null<llvm::Function>(
+ CGM.GetGlobalValue(CGM.getMangledName(D))))
+ return !F->isDeclaration();
+ return false;
+ }
return true;
+ }
- const FunctionDecl *FD = D->getCanonicalDecl();
// Do not mark member functions except for static.
if (const auto *Method = dyn_cast<CXXMethodDecl>(FD))
if (!Method->isStatic())
Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h?rev=328705&r1=328704&r2=328705&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h Wed Mar 28 07:28:54 2018
@@ -217,7 +217,7 @@ protected:
/// \brief Creates offloading entry for the provided entry ID \a ID,
/// address \a Addr, size \a Size, and flags \a Flags.
virtual void createOffloadEntry(llvm::Constant *ID, llvm::Constant *Addr,
- uint64_t Size, int32_t Flags = 0);
+ uint64_t Size, int32_t Flags);
/// \brief Helper to emit outlined function for 'target' directive.
/// \param D Directive to emit.
@@ -382,6 +382,15 @@ private:
/// // entries (non inclusive).
/// };
QualType TgtBinaryDescriptorQTy;
+ /// Kind of the target registry entry.
+ enum OMPTargetRegionEntryKind {
+ /// Mark the entry as target region.
+ OMPTargetRegionEntryTargetRegion = 0x0,
+ /// Mark the entry as a global constructor.
+ OMPTargetRegionEntryCtor = 0x02,
+ /// Mark the entry as a global destructor.
+ OMPTargetRegionEntryDtor = 0x04,
+ };
/// \brief Entity that registers the offloading constants that were emitted so
/// far.
class OffloadEntriesInfoManagerTy {
@@ -394,31 +403,31 @@ private:
/// Base class of the entries info.
class OffloadEntryInfo {
public:
- /// Kind of a given entry. Currently, only target regions are
- /// supported.
+ /// Kind of a given entry.
enum OffloadingEntryInfoKinds : unsigned {
- // Entry is a target region.
- OFFLOAD_ENTRY_INFO_TARGET_REGION = 0,
- // Invalid entry info.
- OFFLOAD_ENTRY_INFO_INVALID = ~0u
+ /// Entry is a target region.
+ OffloadingEntryInfoTargetRegion = 0,
+ /// Invalid entry info.
+ OffloadingEntryInfoInvalid = ~0u
};
OffloadEntryInfo()
- : Flags(0), Order(~0u), Kind(OFFLOAD_ENTRY_INFO_INVALID) {}
+ : Flags(OMPTargetRegionEntryTargetRegion), Order(~0u),
+ Kind(OffloadingEntryInfoInvalid) {}
explicit OffloadEntryInfo(OffloadingEntryInfoKinds Kind, unsigned Order,
- int32_t Flags)
+ OMPTargetRegionEntryKind Flags)
: Flags(Flags), Order(Order), Kind(Kind) {}
bool isValid() const { return Order != ~0u; }
unsigned getOrder() const { return Order; }
OffloadingEntryInfoKinds getKind() const { return Kind; }
int32_t getFlags() const { return Flags; }
- void setFlags(int32_t NewFlags) { Flags = NewFlags; }
+ void setFlags(OMPTargetRegionEntryKind NewFlags) { Flags = NewFlags; }
static bool classof(const OffloadEntryInfo *Info) { return true; }
private:
/// Flags associated with the device global.
- int32_t Flags;
+ OMPTargetRegionEntryKind Flags;
/// Order this entry was emitted.
unsigned Order;
@@ -445,27 +454,28 @@ private:
public:
OffloadEntryInfoTargetRegion()
- : OffloadEntryInfo(OFFLOAD_ENTRY_INFO_TARGET_REGION, ~0u,
- /*Flags=*/0),
+ : OffloadEntryInfo(OffloadingEntryInfoTargetRegion, ~0u,
+ OMPTargetRegionEntryTargetRegion),
Addr(nullptr), ID(nullptr) {}
explicit OffloadEntryInfoTargetRegion(unsigned Order,
llvm::Constant *Addr,
- llvm::Constant *ID, int32_t Flags)
- : OffloadEntryInfo(OFFLOAD_ENTRY_INFO_TARGET_REGION, Order, Flags),
+ llvm::Constant *ID,
+ OMPTargetRegionEntryKind Flags)
+ : OffloadEntryInfo(OffloadingEntryInfoTargetRegion, Order, Flags),
Addr(Addr), ID(ID) {}
llvm::Constant *getAddress() const { return Addr; }
llvm::Constant *getID() const { return ID; }
void setAddress(llvm::Constant *V) {
- assert(!Addr && "Address as been set before!");
+ assert(!Addr && "Address has been set before!");
Addr = V;
}
void setID(llvm::Constant *V) {
- assert(!ID && "ID as been set before!");
+ assert(!ID && "ID has been set before!");
ID = V;
}
static bool classof(const OffloadEntryInfo *Info) {
- return Info->getKind() == OFFLOAD_ENTRY_INFO_TARGET_REGION;
+ return Info->getKind() == OffloadingEntryInfoTargetRegion;
}
};
/// \brief Initialize target region entry.
@@ -476,7 +486,7 @@ private:
void registerTargetRegionEntryInfo(unsigned DeviceID, unsigned FileID,
StringRef ParentName, unsigned LineNum,
llvm::Constant *Addr, llvm::Constant *ID,
- int32_t Flags);
+ OMPTargetRegionEntryKind Flags);
/// \brief Return true if a target region entry with the provided
/// information exists.
bool hasTargetRegionEntryInfo(unsigned DeviceID, unsigned FileID,
@@ -580,6 +590,9 @@ private:
/// \brief Set of threadprivate variables with the generated initializer.
llvm::SmallPtrSet<const VarDecl *, 4> ThreadPrivateWithDefinition;
+ /// Set of declare target variables with the generated initializer.
+ llvm::SmallPtrSet<const VarDecl *, 4> DeclareTargetWithDefinition;
+
/// \brief Emits initialization code for the threadprivate variables.
/// \param VDAddr Address of the global variable \a VD.
/// \param Ctor Pointer to a global init function for \a VD.
@@ -970,6 +983,14 @@ public:
SourceLocation Loc, bool PerformInit,
CodeGenFunction *CGF = nullptr);
+ /// \brief Emit a code for initialization of declare target variable.
+ /// \param VD Declare target variable.
+ /// \param Addr Address of the global variable \a VD.
+ /// \param PerformInit true if initialization expression is not constant.
+ virtual bool emitDeclareTargetVarDefinition(const VarDecl *VD,
+ llvm::GlobalVariable *Addr,
+ bool PerformInit);
+
/// Creates artificial threadprivate variable with name \p Name and type \p
/// VarType.
/// \param VarType Type of the artificial threadprivate variable.
Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.h
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.h?rev=328705&r1=328704&r2=328705&view=diff
==============================================================================
--- cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.h (original)
+++ cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.h Wed Mar 28 07:28:54 2018
@@ -81,7 +81,7 @@ private:
/// \brief Creates offloading entry for the provided entry ID \a ID,
/// address \a Addr, size \a Size, and flags \a Flags.
void createOffloadEntry(llvm::Constant *ID, llvm::Constant *Addr,
- uint64_t Size, int32_t Flags = 0) override;
+ uint64_t Size, int32_t Flags) override;
/// \brief Emit outlined function specialized for the Fork-Join
/// programming model for applicable target directives on the NVPTX device.
Modified: cfe/trunk/lib/Parse/ParseOpenMP.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Parse/ParseOpenMP.cpp?rev=328705&r1=328704&r2=328705&view=diff
==============================================================================
--- cfe/trunk/lib/Parse/ParseOpenMP.cpp (original)
+++ cfe/trunk/lib/Parse/ParseOpenMP.cpp Wed Mar 28 07:28:54 2018
@@ -719,7 +719,7 @@ Parser::DeclGroupPtrTy Parser::ParseOpen
SourceLocation DTLoc = ConsumeAnyToken();
if (Tok.isNot(tok::annot_pragma_openmp_end)) {
// OpenMP 4.5 syntax with list of entities.
- llvm::SmallSetVector<const NamedDecl*, 16> SameDirectiveDecls;
+ Sema::NamedDeclSetType SameDirectiveDecls;
while (Tok.isNot(tok::annot_pragma_openmp_end)) {
OMPDeclareTargetDeclAttr::MapTypeTy MT =
OMPDeclareTargetDeclAttr::MT_To;
@@ -736,11 +736,12 @@ Parser::DeclGroupPtrTy Parser::ParseOpen
ConsumeToken();
}
auto Callback = [this, MT, &SameDirectiveDecls](
- CXXScopeSpec &SS, DeclarationNameInfo NameInfo) {
+ CXXScopeSpec &SS, DeclarationNameInfo NameInfo) {
Actions.ActOnOpenMPDeclareTargetName(getCurScope(), SS, NameInfo, MT,
SameDirectiveDecls);
};
- if (ParseOpenMPSimpleVarList(OMPD_declare_target, Callback, true))
+ if (ParseOpenMPSimpleVarList(OMPD_declare_target, Callback,
+ /*AllowScopeSpecifier=*/true))
break;
// Consume optional ','.
@@ -749,7 +750,13 @@ Parser::DeclGroupPtrTy Parser::ParseOpen
}
SkipUntil(tok::annot_pragma_openmp_end, StopBeforeMatch);
ConsumeAnyToken();
- return DeclGroupPtrTy();
+ SmallVector<Decl *, 4> Decls;
+ Decls.reserve(SameDirectiveDecls.size());
+ for (Decl *D : SameDirectiveDecls)
+ Decls.emplace_back(D);
+ if (Decls.empty())
+ return DeclGroupPtrTy();
+ return Actions.BuildDeclaratorGroup(Decls);
}
// Skip the last annot_pragma_openmp_end.
@@ -802,8 +809,7 @@ Parser::DeclGroupPtrTy Parser::ParseOpen
Diag(DTLoc, diag::note_matching) << "'#pragma omp declare target'";
}
Actions.ActOnFinishOpenMPDeclareTargetDirective();
- return DeclGroupPtrTy::make(DeclGroupRef::Create(
- Actions.getASTContext(), Decls.begin(), Decls.size()));
+ return Actions.BuildDeclaratorGroup(Decls);
}
case OMPD_unknown:
Diag(Tok, diag::err_omp_unknown_directive);
Added: cfe/trunk/test/OpenMP/nvptx_declare_target_var_ctor_dtor_codegen.cpp
URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/nvptx_declare_target_var_ctor_dtor_codegen.cpp?rev=328705&view=auto
==============================================================================
--- cfe/trunk/test/OpenMP/nvptx_declare_target_var_ctor_dtor_codegen.cpp (added)
+++ cfe/trunk/test/OpenMP/nvptx_declare_target_var_ctor_dtor_codegen.cpp Wed Mar 28 07:28:54 2018
@@ -0,0 +1,104 @@
+// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s --check-prefix HOST --check-prefix CHECK
+// RUN: %clang_cc1 -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -fopenmp -x c++ -triple nvptx64-nvidia-cuda -fopenmp-targets=nvptx64-nvidia-cuda -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 -fopenmp -x c++ -triple nvptx64-nvidia-cuda -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t
+// RUN: %clang_cc1 -fopenmp -x c++ -triple nvptx64-nvidia-cuda -fopenmp-targets=nvptx64-nvidia-cuda -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 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o - | FileCheck %s --check-prefix SIMD-ONLY
+// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -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 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t
+// RUN: %clang_cc1 -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - | FileCheck %s --check-prefix SIMD-ONLY
+
+#ifndef HEADER
+#define HEADER
+
+// SIMD-ONLY-NOT: {{__kmpc|__tgt}}
+
+// DEVICE-DAG: [[C_ADDR:@.+]] = internal global i32 0,
+// DEVICE-DAG: [[CD_ADDR:@.+]] = global %struct.S zeroinitializer,
+// HOST-DAG: [[C_ADDR:@.+]] = internal global i32 0,
+// HOST-DAG: [[CD_ADDR:@.+]] = global %struct.S zeroinitializer,
+
+#pragma omp declare target
+int foo() { return 0; }
+#pragma omp end declare target
+int bar() { return 0; }
+#pragma omp declare target (bar)
+int baz() { return 0; }
+
+#pragma omp declare target
+int doo() { return 0; }
+#pragma omp end declare target
+int car() { return 0; }
+#pragma omp declare target (bar)
+int caz() { return 0; }
+
+// DEVICE-DAG: define i32 [[FOO:@.*foo.*]]()
+// DEVICE-DAG: define i32 [[BAR:@.*bar.*]]()
+// DEVICE-DAG: define i32 [[BAZ:@.*baz.*]]()
+// DEVICE-DAG: define i32 [[DOO:@.*doo.*]]()
+// DEVICE-DAG: define i32 [[CAR:@.*car.*]]()
+// DEVICE-DAG: define i32 [[CAZ:@.*caz.*]]()
+
+static int c = foo() + bar() + baz();
+#pragma omp declare target (c)
+// HOST-DAG: @[[C_CTOR:__omp_offloading__.+_c_l44_ctor]] = private constant i8 0
+// DEVICE-DAG: define internal void [[C_CTOR:@__omp_offloading__.+_c_l44_ctor]]()
+// DEVICE-DAG: call i32 [[FOO]]()
+// DEVICE-DAG: call i32 [[BAR]]()
+// DEVICE-DAG: call i32 [[BAZ]]()
+// DEVICE-DAG: ret void
+
+struct S {
+ int a;
+ S() = default;
+ S(int a) : a(a) {}
+ ~S() { a = 0; }
+};
+
+#pragma omp declare target
+S cd = doo() + car() + caz() + baz();
+#pragma omp end declare target
+// HOST-DAG: @[[CD_CTOR:__omp_offloading__.+_cd_l61_ctor]] = private constant i8 0
+// DEVICE-DAG: define internal void [[CD_CTOR:@__omp_offloading__.+_cd_l61_ctor]]()
+// DEVICE-DAG: call i32 [[DOO]]()
+// DEVICE-DAG: call i32 [[CAR]]()
+// DEVICE-DAG: call i32 [[CAZ]]()
+// DEVICE-DAG: ret void
+
+// HOST-DAG: @[[CD_DTOR:__omp_offloading__.+_cd_l61_dtor]] = private constant i8 0
+// DEVICE-DAG: define internal void [[CD_DTOR:@__omp_offloading__.+_cd_l61_dtor]]()
+// DEVICE-DAG: call void
+// DEVICE-DAG: ret void
+
+// HOST: @.omp_offloading.entry_name{{.*}} = internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[C_CTOR]]\00"
+// HOST: @.omp_offloading.entry.[[C_CTOR]] = constant %struct.__tgt_offload_entry { i8* @[[C_CTOR]], i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @.omp_offloading.entry_name{{.*}}, i32 0, i32 0), i64 0, i32 2, i32 0 }, section ".omp_offloading.entries", align 1
+// HOST: @.omp_offloading.entry_name{{.*}}= internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[CD_CTOR]]\00"
+// HOST: @.omp_offloading.entry.[[CD_CTOR]] = constant %struct.__tgt_offload_entry { i8* @[[CD_CTOR]], i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @.omp_offloading.entry_name{{.*}}, i32 0, i32 0), i64 0, i32 2, i32 0 }, section ".omp_offloading.entries", align 1
+// HOST: @.omp_offloading.entry_name{{.*}}= internal unnamed_addr constant [{{[0-9]+}} x i8] c"[[CD_DTOR]]\00"
+// HOST: @.omp_offloading.entry.[[CD_DTOR]] = constant %struct.__tgt_offload_entry { i8* @[[CD_DTOR]], i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @.omp_offloading.entry_name{{.*}}, i32 0, i32 0), i64 0, i32 4, i32 0 }, section ".omp_offloading.entries", align 1
+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:%.+]] = load i32, i32* [[C_ADDR]],
+// DEVICE: store i32 [[C]], i32* %
+
+// HOST: define internal void @__omp_offloading_{{.*}}_{{.*}}maini1{{.*}}_l[[@LINE-11]](i32* dereferenceable{{.*}})
+// HOST: [[C:%.*]] = load i32, i32* [[C_ADDR]],
+// HOST: store i32 [[C]], i32* %
+
+// DEVICE: !nvvm.annotations
+// DEVICE-DAG: !{void ()* [[C_CTOR]], !"kernel", i32 1}
+// DEVICE-DAG: !{void ()* [[CD_CTOR]], !"kernel", i32 1}
+// DEVICE-DAG: !{void ()* [[CD_DTOR]], !"kernel", i32 1}
+
+#endif // HEADER
+
More information about the cfe-commits
mailing list