From: Alexey Bataev Date: Wed, 28 Mar 2018 14:28:54 +0000 (+0000) Subject: [OPENMP] Codegen for ctor|dtor of declare target variables. X-Git-Url: https://granicus.if.org/sourcecode?a=commitdiff_plain;h=c368db314cba827ced921527739fc9c53c04a121;p=clang [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. git-svn-id: https://llvm.org/svn/llvm-project/cfe/trunk@328705 91177308-0d34-0410-b5e6-96231b3b80d8 --- diff --git a/include/clang/Sema/Sema.h b/include/clang/Sema/Sema.h index 2e907562fa..3f8ff6bb92 100644 --- a/include/clang/Sema/Sema.h +++ b/include/clang/Sema/Sema.h @@ -546,7 +546,7 @@ public: /// FieldCollector - Collects CXXFieldDecls during parsing of C++ classes. std::unique_ptr FieldCollector; - typedef llvm::SmallSetVector NamedDeclSetType; + typedef llvm::SmallSetVector NamedDeclSetType; /// \brief Set containing all declared private fields that are not used. NamedDeclSetType UnusedPrivateFields; diff --git a/lib/CodeGen/CGDeclCXX.cpp b/lib/CodeGen/CGDeclCXX.cpp index 4050499d98..d6e97dd577 100644 --- a/lib/CodeGen/CGDeclCXX.cpp +++ b/lib/CodeGen/CGDeclCXX.cpp @@ -379,6 +379,10 @@ CodeGenModule::EmitCXXGlobalVarDeclInitFunc(const VarDecl *D, D->hasAttr())) 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) diff --git a/lib/CodeGen/CGOpenMPRuntime.cpp b/lib/CodeGen/CGOpenMPRuntime.cpp index 891176b009..4b1fac6827 100644 --- a/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/lib/CodeGen/CGOpenMPRuntime.cpp @@ -2522,6 +2522,139 @@ llvm::Function *CGOpenMPRuntime::emitThreadPrivateVarDefinition( 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 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::OffloadEntriesInfoManagerTy:: "code generation."); OffloadEntriesTargetRegion[DeviceID][FileID][ParentName][LineNum] = OffloadEntryInfoTargetRegion(Order, /*Addr=*/nullptr, /*ID=*/nullptr, - /*Flags=*/0); + OMPTargetRegionEntryTargetRegion); ++OffloadingEntriesNum; } @@ -3383,7 +3516,7 @@ void CGOpenMPRuntime::OffloadEntriesInfoManagerTy:: 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::createOffloadEntriesAndInfoMetadata() { 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::createOffloadEntriesAndInfoMetadata() { // - 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::createOffloadEntriesAndInfoMetadata() { 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::loadOffloadInfoMetadata() { return; for (llvm::MDNode *MN : MD->operands()) { - auto getMDInt = [&](unsigned Idx) { + auto GetMdInt = [MN](unsigned Idx) { llvm::ConstantAsMetadata *V = cast(MN->getOperand(Idx)); return cast(V->getValue())->getZExtValue(); }; - auto getMDString = [&](unsigned Idx) { + auto GetMdString = [MN](unsigned Idx) { llvm::MDString *V = cast(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(CodeGenFunction &CGF, SourceLocation Loc, } } -/// \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::emitTargetOutlinedFunctionHelper( // 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::DisableAutoDeclareTargetRAII::~DisableAutoDeclareTargetRAII() { 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( + 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(FD)) if (!Method->isStatic()) diff --git a/lib/CodeGen/CGOpenMPRuntime.h b/lib/CodeGen/CGOpenMPRuntime.h index e3eb65aa8b..b78bab62fd 100644 --- a/lib/CodeGen/CGOpenMPRuntime.h +++ b/lib/CodeGen/CGOpenMPRuntime.h @@ -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 ThreadPrivateWithDefinition; + /// Set of declare target variables with the generated initializer. + llvm::SmallPtrSet 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. diff --git a/lib/CodeGen/CGOpenMPRuntimeNVPTX.h b/lib/CodeGen/CGOpenMPRuntimeNVPTX.h index feb55c5698..20c3419e33 100644 --- a/lib/CodeGen/CGOpenMPRuntimeNVPTX.h +++ b/lib/CodeGen/CGOpenMPRuntimeNVPTX.h @@ -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. diff --git a/lib/Parse/ParseOpenMP.cpp b/lib/Parse/ParseOpenMP.cpp index 77c5d23b14..7528e435f0 100644 --- a/lib/Parse/ParseOpenMP.cpp +++ b/lib/Parse/ParseOpenMP.cpp @@ -719,7 +719,7 @@ Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirectiveWithExtDecl( SourceLocation DTLoc = ConsumeAnyToken(); if (Tok.isNot(tok::annot_pragma_openmp_end)) { // OpenMP 4.5 syntax with list of entities. - llvm::SmallSetVector 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::ParseOpenMPDeclarativeDirectiveWithExtDecl( 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::ParseOpenMPDeclarativeDirectiveWithExtDecl( } SkipUntil(tok::annot_pragma_openmp_end, StopBeforeMatch); ConsumeAnyToken(); - return DeclGroupPtrTy(); + SmallVector 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::ParseOpenMPDeclarativeDirectiveWithExtDecl( 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); diff --git a/test/OpenMP/nvptx_declare_target_var_ctor_dtor_codegen.cpp b/test/OpenMP/nvptx_declare_target_var_ctor_dtor_codegen.cpp new file mode 100644 index 0000000000..3683b6dc5f --- /dev/null +++ b/test/OpenMP/nvptx_declare_target_var_ctor_dtor_codegen.cpp @@ -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 +