Index: clang/include/clang/Sema/ExternalSemaSource.h =================================================================== --- clang/include/clang/Sema/ExternalSemaSource.h +++ clang/include/clang/Sema/ExternalSemaSource.h @@ -193,6 +193,15 @@ llvm::MapVector> &LPTMap) {} + /// Read the set of decls to be checked for deferred diags. + /// + /// The external source should append its own potentially emitted function + /// and variable decls which may cause deferred diags. Note that this routine + /// may be invoked multiple times; the external source should take care not to + /// introduce the same declarations repeatedly. + virtual void ReadDeclsToCheckForDeferredDiags( + llvm::SmallVector &Decls) {} + /// \copydoc Sema::CorrectTypo /// \note LookupKind must correspond to a valid Sema::LookupNameKind /// Index: clang/include/clang/Sema/MultiplexExternalSemaSource.h =================================================================== --- clang/include/clang/Sema/MultiplexExternalSemaSource.h +++ clang/include/clang/Sema/MultiplexExternalSemaSource.h @@ -332,6 +332,15 @@ llvm::MapVector> &LPTMap) override; + /// Read the set of decls to be checked for deferred diags. + /// + /// The external source should append its own potentially emitted function + /// and variable decls which may cause deferred diags. Note that this routine + /// may be invoked multiple times; the external source should take care not to + /// introduce the same declarations repeatedly. + void ReadDeclsToCheckForDeferredDiags( + llvm::SmallVector &Decls) override; + /// \copydoc ExternalSemaSource::CorrectTypo /// \note Returns the first nonempty correction. TypoCorrection CorrectTypo(const DeclarationNameInfo &Typo, Index: clang/include/clang/Sema/Sema.h =================================================================== --- clang/include/clang/Sema/Sema.h +++ clang/include/clang/Sema/Sema.h @@ -1492,6 +1492,18 @@ void emitAndClearUnusedLocalTypedefWarnings(); + private: + /// Function or variable declarations to be checked for whether the deferred + /// diagnostics should be emitted. + SmallVector DeclsToCheckForDeferredDiags; + + public: + // Emit all deferred diagnostics. + void emitDeferredDiags(); + // Emit any deferred diagnostics for FD and erase them from the map in which + // they're stored. + void emitDeferredDiags(FunctionDecl *FD, bool ShowCallStack); + enum TUFragmentKind { /// The global module fragment, between 'module;' and a module-declaration. Global, @@ -3767,7 +3779,8 @@ TemplateDiscarded, // Discarded due to uninstantiated templates Unknown, }; - FunctionEmissionStatus getEmissionStatus(FunctionDecl *Decl); + FunctionEmissionStatus getEmissionStatus(FunctionDecl *Decl, + bool Final = false); // Whether the callee should be ignored in CUDA/HIP/OpenMP host/device check. bool shouldIgnoreInHostDeviceCheck(FunctionDecl *Callee); @@ -9767,22 +9780,10 @@ /// Pop OpenMP function region for non-capturing function. void popOpenMPFunctionRegion(const sema::FunctionScopeInfo *OldFSI); - /// Check whether we're allowed to call Callee from the current function. - void checkOpenMPDeviceFunction(SourceLocation Loc, FunctionDecl *Callee, - bool CheckForDelayedContext = true); - - /// Check whether we're allowed to call Callee from the current function. - void checkOpenMPHostFunction(SourceLocation Loc, FunctionDecl *Callee, - bool CheckCaller = true); - /// Check if the expression is allowed to be used in expressions for the /// OpenMP devices. void checkOpenMPDeviceExpr(const Expr *E); - /// Finishes analysis of the deferred functions calls that may be declared as - /// host/nohost during device/host compilation. - void finalizeOpenMPDelayedAnalysis(); - /// Checks if a type or a declaration is disabled due to the owning extension /// being disabled, and emits diagnostic messages if it is disabled. /// \param D type or declaration to be checked. @@ -9973,6 +9974,11 @@ void checkDeclIsAllowedInOpenMPTarget(Expr *E, Decl *D, SourceLocation IdLoc = SourceLocation()); + /// Finishes analysis of the deferred functions calls that may be declared as + /// host/nohost during device/host compilation. + void finalizeOpenMPDelayedAnalysis(const FunctionDecl *Caller, + const FunctionDecl *Callee, + SourceLocation Loc); /// Return true inside OpenMP declare target region. bool isInOpenMPDeclareTargetContext() const { return DeclareTargetNestingLevel > 0; @@ -11348,18 +11354,6 @@ /* Caller = */ FunctionDeclAndLoc> DeviceKnownEmittedFns; - /// A partial call graph maintained during CUDA/OpenMP device code compilation - /// to support deferred diagnostics. - /// - /// Functions are only added here if, at the time they're considered, they are - /// not known-emitted. As soon as we discover that a function is - /// known-emitted, we remove it and everything it transitively calls from this - /// set and add those functions to DeviceKnownEmittedFns. - llvm::DenseMap, - /* Callees = */ llvm::MapVector, - SourceLocation>> - DeviceCallGraph; - /// Diagnostic builder for CUDA/OpenMP devices errors which may or may not be /// deferred. /// @@ -11434,14 +11428,6 @@ llvm::Optional PartialDiagId; }; - /// Indicate that this function (and thus everything it transtively calls) - /// will be codegen'ed, and emit any deferred diagnostics on this function and - /// its (transitive) callees. - void markKnownEmitted( - Sema &S, FunctionDecl *OrigCaller, FunctionDecl *OrigCallee, - SourceLocation OrigLoc, - const llvm::function_ref IsKnownEmitted); - /// Creates a DeviceDiagBuilder that emits the diagnostic if the current context /// is "used as device code". /// Index: clang/include/clang/Serialization/ASTBitCodes.h =================================================================== --- clang/include/clang/Serialization/ASTBitCodes.h +++ clang/include/clang/Serialization/ASTBitCodes.h @@ -650,7 +650,10 @@ PP_CONDITIONAL_STACK = 62, /// A table of skipped ranges within the preprocessing record. - PPD_SKIPPED_RANGES = 63 + PPD_SKIPPED_RANGES = 63, + + /// Record code for the Decls to be checked for deferred diags. + DECLS_TO_CHECK_FOR_DEFERRED_DIAGS = 64, }; /// Record types used within a source manager block. Index: clang/include/clang/Serialization/ASTReader.h =================================================================== --- clang/include/clang/Serialization/ASTReader.h +++ clang/include/clang/Serialization/ASTReader.h @@ -890,6 +890,12 @@ // A list of late parsed template function data. SmallVector LateParsedTemplates; + /// The IDs of all decls to be checked for deferred diags. + /// + /// Sema tracks these to emit deferred diags. + SmallVector DeclsToCheckForDeferredDiags; + + public: struct ImportedSubmodule { serialization::SubmoduleID ID; @@ -1983,6 +1989,9 @@ void ReadUnusedLocalTypedefNameCandidates( llvm::SmallSetVector &Decls) override; + void ReadDeclsToCheckForDeferredDiags( + llvm::SmallVector &Decls) override; + void ReadReferencedSelectors( SmallVectorImpl> &Sels) override; Index: clang/lib/Sema/MultiplexExternalSemaSource.cpp =================================================================== --- clang/lib/Sema/MultiplexExternalSemaSource.cpp +++ clang/lib/Sema/MultiplexExternalSemaSource.cpp @@ -275,6 +275,12 @@ Sources[i]->ReadExtVectorDecls(Decls); } +void MultiplexExternalSemaSource::ReadDeclsToCheckForDeferredDiags( + llvm::SmallVector &Decls) { + for(size_t i = 0; i < Sources.size(); ++i) + Sources[i]->ReadDeclsToCheckForDeferredDiags(Decls); +} + void MultiplexExternalSemaSource::ReadUnusedLocalTypedefNameCandidates( llvm::SmallSetVector &Decls) { for(size_t i = 0; i < Sources.size(); ++i) Index: clang/lib/Sema/Sema.cpp =================================================================== --- clang/lib/Sema/Sema.cpp +++ clang/lib/Sema/Sema.cpp @@ -11,6 +11,7 @@ // //===----------------------------------------------------------------------===// +#include "UsedDeclVisitor.h" #include "clang/AST/ASTContext.h" #include "clang/AST/ASTDiagnostic.h" #include "clang/AST/DeclCXX.h" @@ -955,9 +956,7 @@ PerformPendingInstantiations(); } - // Finalize analysis of OpenMP-specific constructs. - if (LangOpts.OpenMP) - finalizeOpenMPDelayedAnalysis(); + emitDeferredDiags(); assert(LateParsedInstantiations.empty() && "end of TU template instantiation should not create more " @@ -1452,27 +1451,108 @@ // Emit any deferred diagnostics for FD and erase them from the map in which // they're stored. -static void emitDeferredDiags(Sema &S, FunctionDecl *FD, bool ShowCallStack) { - auto It = S.DeviceDeferredDiags.find(FD); - if (It == S.DeviceDeferredDiags.end()) +void Sema::emitDeferredDiags(FunctionDecl *FD, bool ShowCallStack) { + auto It = DeviceDeferredDiags.find(FD); + if (It == DeviceDeferredDiags.end()) return; bool HasWarningOrError = false; + bool FirstDiag = true; for (PartialDiagnosticAt &PDAt : It->second) { const SourceLocation &Loc = PDAt.first; const PartialDiagnostic &PD = PDAt.second; - HasWarningOrError |= S.getDiagnostics().getDiagnosticLevel( + HasWarningOrError |= getDiagnostics().getDiagnosticLevel( PD.getDiagID(), Loc) >= DiagnosticsEngine::Warning; - DiagnosticBuilder Builder(S.Diags.Report(Loc, PD.getDiagID())); - Builder.setForceEmit(); - PD.Emit(Builder); + { + DiagnosticBuilder Builder(Diags.Report(Loc, PD.getDiagID())); + Builder.setForceEmit(); + PD.Emit(Builder); + } + + // Emit the note on the first diagnostic in case too many diagnostics cause + // the note not emitted. + if (FirstDiag && HasWarningOrError && ShowCallStack) { + emitCallStackNotes(*this, FD); + FirstDiag = false; + } } - S.DeviceDeferredDiags.erase(It); - // FIXME: Should this be called after every warning/error emitted in the loop - // above, instead of just once per function? That would be consistent with - // how we handle immediate errors, but it also seems like a bit much. - if (HasWarningOrError && ShowCallStack) - emitCallStackNotes(S, FD); +} + +namespace { +/// Helper class that emits deferred diagnostic messages if an entity directly +/// or indirectly using the function that causes the deferred diagnostic +/// messages is known to be emitted. +class DeferredDiagnosticsEmitter + : public UsedDeclVisitor { +public: + typedef UsedDeclVisitor Inherited; + llvm::SmallSet, 4> Visited; + llvm::SmallVector, 4> UseStack; + bool ShouldEmit; + unsigned InOMPDeviceContext; + + DeferredDiagnosticsEmitter(Sema &S) + : Inherited(S), ShouldEmit(false), InOMPDeviceContext(0) {} + + void VisitOMPTargetDirective(OMPTargetDirective *Node) { + ++InOMPDeviceContext; + Inherited::VisitOMPTargetDirective(Node); + --InOMPDeviceContext; + } + + void visitUsedDecl(SourceLocation Loc, Decl *D) { + if (auto *FD = dyn_cast(D)) { + FunctionDecl *Caller = UseStack.empty() ? nullptr : UseStack.back(); + auto IsKnownEmitted = S.getEmissionStatus(FD, /*Final=*/true) == + Sema::FunctionEmissionStatus::Emitted; + if (!Caller) + ShouldEmit = IsKnownEmitted; + if ((!ShouldEmit && !S.getLangOpts().OpenMP && !Caller) || + S.shouldIgnoreInHostDeviceCheck(FD) || Visited.count(D)) + return; + // Finalize analysis of OpenMP-specific constructs. + if (Caller && S.LangOpts.OpenMP && UseStack.size() == 1) + S.finalizeOpenMPDelayedAnalysis(Caller, FD, Loc); + if (Caller) + S.DeviceKnownEmittedFns[FD] = {Caller, Loc}; + if (ShouldEmit || InOMPDeviceContext) + S.emitDeferredDiags(FD, Caller); + Visited.insert(D); + UseStack.push_back(FD); + if (auto *S = FD->getBody()) { + this->Visit(S); + } + UseStack.pop_back(); + Visited.erase(D); + } else if (auto *VD = dyn_cast(D)) { + if (auto *Init = VD->getInit()) { + auto DevTy = OMPDeclareTargetDeclAttr::getDeviceType(VD); + bool IsDev = DevTy && (*DevTy == OMPDeclareTargetDeclAttr::DT_NoHost || + *DevTy == OMPDeclareTargetDeclAttr::DT_Any); + if (IsDev) + ++InOMPDeviceContext; + this->Visit(Init); + if (IsDev) + --InOMPDeviceContext; + } + } else + Inherited::visitUsedDecl(Loc, D); + } +}; +} // namespace + +void Sema::emitDeferredDiags() { + if (ExternalSource) + ExternalSource->ReadDeclsToCheckForDeferredDiags( + DeclsToCheckForDeferredDiags); + + if ((DeviceDeferredDiags.empty() && !LangOpts.OpenMP) || + DeclsToCheckForDeferredDiags.empty()) + return; + + DeferredDiagnosticsEmitter DDE(*this); + for (auto D : DeclsToCheckForDeferredDiags) + DDE.visitUsedDecl(SourceLocation(), D); } // In CUDA, there are some constructs which may appear in semantically-valid @@ -1545,71 +1625,6 @@ } } -// Indicate that this function (and thus everything it transtively calls) will -// be codegen'ed, and emit any deferred diagnostics on this function and its -// (transitive) callees. -void Sema::markKnownEmitted( - Sema &S, FunctionDecl *OrigCaller, FunctionDecl *OrigCallee, - SourceLocation OrigLoc, - const llvm::function_ref IsKnownEmitted) { - // Nothing to do if we already know that FD is emitted. - if (IsKnownEmitted(S, OrigCallee)) { - assert(!S.DeviceCallGraph.count(OrigCallee)); - return; - } - - // We've just discovered that OrigCallee is known-emitted. Walk our call - // graph to see what else we can now discover also must be emitted. - - struct CallInfo { - FunctionDecl *Caller; - FunctionDecl *Callee; - SourceLocation Loc; - }; - llvm::SmallVector Worklist = {{OrigCaller, OrigCallee, OrigLoc}}; - llvm::SmallSet, 4> Seen; - Seen.insert(OrigCallee); - while (!Worklist.empty()) { - CallInfo C = Worklist.pop_back_val(); - assert(!IsKnownEmitted(S, C.Callee) && - "Worklist should not contain known-emitted functions."); - S.DeviceKnownEmittedFns[C.Callee] = {C.Caller, C.Loc}; - emitDeferredDiags(S, C.Callee, C.Caller); - - // If this is a template instantiation, explore its callgraph as well: - // Non-dependent calls are part of the template's callgraph, while dependent - // calls are part of to the instantiation's call graph. - if (auto *Templ = C.Callee->getPrimaryTemplate()) { - FunctionDecl *TemplFD = Templ->getAsFunction(); - if (!Seen.count(TemplFD) && !S.DeviceKnownEmittedFns.count(TemplFD)) { - Seen.insert(TemplFD); - Worklist.push_back( - {/* Caller = */ C.Caller, /* Callee = */ TemplFD, C.Loc}); - } - } - - // Add all functions called by Callee to our worklist. - auto CGIt = S.DeviceCallGraph.find(C.Callee); - if (CGIt == S.DeviceCallGraph.end()) - continue; - - for (std::pair, SourceLocation> FDLoc : - CGIt->second) { - FunctionDecl *NewCallee = FDLoc.first; - SourceLocation CallLoc = FDLoc.second; - if (Seen.count(NewCallee) || IsKnownEmitted(S, NewCallee)) - continue; - Seen.insert(NewCallee); - Worklist.push_back( - {/* Caller = */ C.Callee, /* Callee = */ NewCallee, CallLoc}); - } - - // C.Callee is now known-emitted, so we no longer need to maintain its list - // of callees in DeviceCallGraph. - S.DeviceCallGraph.erase(CGIt); - } -} - Sema::DeviceDiagBuilder Sema::targetDiag(SourceLocation Loc, unsigned DiagID) { if (LangOpts.OpenMP) return LangOpts.OpenMPIsDevice ? diagIfOpenMPDeviceCode(Loc, DiagID) Index: clang/lib/Sema/SemaCUDA.cpp =================================================================== --- clang/lib/Sema/SemaCUDA.cpp +++ clang/lib/Sema/SemaCUDA.cpp @@ -675,25 +675,6 @@ // Otherwise, mark the call in our call graph so we can traverse it later. bool CallerKnownEmitted = getEmissionStatus(Caller) == FunctionEmissionStatus::Emitted; - if (CallerKnownEmitted) { - // Host-side references to a __global__ function refer to the stub, so the - // function itself is never emitted and therefore should not be marked. - if (!shouldIgnoreInHostDeviceCheck(Callee)) - markKnownEmitted( - *this, Caller, Callee, Loc, [](Sema &S, FunctionDecl *FD) { - return S.getEmissionStatus(FD) == FunctionEmissionStatus::Emitted; - }); - } else { - // If we have - // host fn calls kernel fn calls host+device, - // the HD function does not get instantiated on the host. We model this by - // omitting at the call to the kernel from the callgraph. This ensures - // that, when compiling for host, only HD functions actually called from the - // host get marked as known-emitted. - if (!shouldIgnoreInHostDeviceCheck(Callee)) - DeviceCallGraph[Caller].insert({Callee, Loc}); - } - DeviceDiagBuilder::Kind DiagKind = [this, Caller, Callee, CallerKnownEmitted] { switch (IdentifyCUDAPreference(Caller, Callee)) { Index: clang/lib/Sema/SemaDecl.cpp =================================================================== --- clang/lib/Sema/SemaDecl.cpp +++ clang/lib/Sema/SemaDecl.cpp @@ -12244,6 +12244,8 @@ VDecl->setInitStyle(VarDecl::ListInit); } + if (LangOpts.OpenMP && VDecl->hasGlobalStorage()) + DeclsToCheckForDeferredDiags.push_back(VDecl); CheckCompleteVariableDeclaration(VDecl); } @@ -14360,6 +14362,13 @@ DiscardCleanupsInEvaluationContext(); } + if (LangOpts.OpenMP || LangOpts.CUDA) { + auto ES = getEmissionStatus(FD); + if (ES == Sema::FunctionEmissionStatus::Emitted || + ES == Sema::FunctionEmissionStatus::Unknown) + DeclsToCheckForDeferredDiags.push_back(FD); + } + return dcl; } @@ -18023,7 +18032,8 @@ return (dyn_cast_or_null(CurContext)); } -Sema::FunctionEmissionStatus Sema::getEmissionStatus(FunctionDecl *FD) { +Sema::FunctionEmissionStatus Sema::getEmissionStatus(FunctionDecl *FD, + bool Final) { // Templates are emitted when they're instantiated. if (FD->isDependentContext()) return FunctionEmissionStatus::TemplateDiscarded; @@ -18035,8 +18045,10 @@ if (DevTy.hasValue()) { if (*DevTy == OMPDeclareTargetDeclAttr::DT_Host) OMPES = FunctionEmissionStatus::OMPDiscarded; - else if (DeviceKnownEmittedFns.count(FD) > 0) + else if (*DevTy == OMPDeclareTargetDeclAttr::DT_NoHost || + *DevTy == OMPDeclareTargetDeclAttr::DT_Any) { OMPES = FunctionEmissionStatus::Emitted; + } } } else if (LangOpts.OpenMP) { // In OpenMP 4.5 all the functions are host functions. @@ -18052,10 +18064,11 @@ if (DevTy.hasValue()) { if (*DevTy == OMPDeclareTargetDeclAttr::DT_NoHost) { OMPES = FunctionEmissionStatus::OMPDiscarded; - } else if (DeviceKnownEmittedFns.count(FD) > 0) { + } else if (*DevTy == OMPDeclareTargetDeclAttr::DT_Host || + *DevTy == OMPDeclareTargetDeclAttr::DT_Any) OMPES = FunctionEmissionStatus::Emitted; - } - } + } else if (Final) + OMPES = FunctionEmissionStatus::Emitted; } } if (OMPES == FunctionEmissionStatus::OMPDiscarded || @@ -18090,9 +18103,7 @@ // Otherwise, the function is known-emitted if it's in our set of // known-emitted functions. - return (DeviceKnownEmittedFns.count(FD) > 0) - ? FunctionEmissionStatus::Emitted - : FunctionEmissionStatus::Unknown; + return FunctionEmissionStatus::Unknown; } bool Sema::shouldIgnoreInHostDeviceCheck(FunctionDecl *Callee) { Index: clang/lib/Sema/SemaExpr.cpp =================================================================== --- clang/lib/Sema/SemaExpr.cpp +++ clang/lib/Sema/SemaExpr.cpp @@ -15990,13 +15990,8 @@ Func->markUsed(Context); } - if (LangOpts.OpenMP) { + if (LangOpts.OpenMP) markOpenMPDeclareVariantFuncsReferenced(Loc, Func, MightBeOdrUse); - if (LangOpts.OpenMPIsDevice) - checkOpenMPDeviceFunction(Loc, Func); - else - checkOpenMPHostFunction(Loc, Func); - } } /// Directly mark a variable odr-used. Given a choice, prefer to use Index: clang/lib/Sema/SemaOpenMP.cpp =================================================================== --- clang/lib/Sema/SemaOpenMP.cpp +++ clang/lib/Sema/SemaOpenMP.cpp @@ -1768,92 +1768,6 @@ return DeviceDiagBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this); } -void Sema::checkOpenMPDeviceFunction(SourceLocation Loc, FunctionDecl *Callee, - bool CheckForDelayedContext) { - assert(LangOpts.OpenMP && LangOpts.OpenMPIsDevice && - "Expected OpenMP device compilation."); - assert(Callee && "Callee may not be null."); - Callee = Callee->getMostRecentDecl(); - FunctionDecl *Caller = getCurFunctionDecl(); - - // host only function are not available on the device. - if (Caller) { - FunctionEmissionStatus CallerS = getEmissionStatus(Caller); - FunctionEmissionStatus CalleeS = getEmissionStatus(Callee); - assert(CallerS != FunctionEmissionStatus::CUDADiscarded && - CalleeS != FunctionEmissionStatus::CUDADiscarded && - "CUDADiscarded unexpected in OpenMP device function check"); - if ((CallerS == FunctionEmissionStatus::Emitted || - (!isOpenMPDeviceDelayedContext(*this) && - CallerS == FunctionEmissionStatus::Unknown)) && - CalleeS == FunctionEmissionStatus::OMPDiscarded) { - StringRef HostDevTy = getOpenMPSimpleClauseTypeName( - OMPC_device_type, OMPC_DEVICE_TYPE_host); - Diag(Loc, diag::err_omp_wrong_device_function_call) << HostDevTy << 0; - Diag(Callee->getAttr()->getLocation(), - diag::note_omp_marked_device_type_here) - << HostDevTy; - return; - } - } - // If the caller is known-emitted, mark the callee as known-emitted. - // Otherwise, mark the call in our call graph so we can traverse it later. - if ((CheckForDelayedContext && !isOpenMPDeviceDelayedContext(*this)) || - (!Caller && !CheckForDelayedContext) || - (Caller && getEmissionStatus(Caller) == FunctionEmissionStatus::Emitted)) - markKnownEmitted(*this, Caller, Callee, Loc, - [CheckForDelayedContext](Sema &S, FunctionDecl *FD) { - return CheckForDelayedContext && - S.getEmissionStatus(FD) == - FunctionEmissionStatus::Emitted; - }); - else if (Caller) - DeviceCallGraph[Caller].insert({Callee, Loc}); -} - -void Sema::checkOpenMPHostFunction(SourceLocation Loc, FunctionDecl *Callee, - bool CheckCaller) { - assert(LangOpts.OpenMP && !LangOpts.OpenMPIsDevice && - "Expected OpenMP host compilation."); - assert(Callee && "Callee may not be null."); - Callee = Callee->getMostRecentDecl(); - FunctionDecl *Caller = getCurFunctionDecl(); - - // device only function are not available on the host. - if (Caller) { - FunctionEmissionStatus CallerS = getEmissionStatus(Caller); - FunctionEmissionStatus CalleeS = getEmissionStatus(Callee); - assert( - (LangOpts.CUDA || (CallerS != FunctionEmissionStatus::CUDADiscarded && - CalleeS != FunctionEmissionStatus::CUDADiscarded)) && - "CUDADiscarded unexpected in OpenMP host function check"); - if (CallerS == FunctionEmissionStatus::Emitted && - CalleeS == FunctionEmissionStatus::OMPDiscarded) { - StringRef NoHostDevTy = getOpenMPSimpleClauseTypeName( - OMPC_device_type, OMPC_DEVICE_TYPE_nohost); - Diag(Loc, diag::err_omp_wrong_device_function_call) << NoHostDevTy << 1; - Diag(Callee->getAttr()->getLocation(), - diag::note_omp_marked_device_type_here) - << NoHostDevTy; - return; - } - } - // If the caller is known-emitted, mark the callee as known-emitted. - // Otherwise, mark the call in our call graph so we can traverse it later. - if (!shouldIgnoreInHostDeviceCheck(Callee)) { - if ((!CheckCaller && !Caller) || - (Caller && - getEmissionStatus(Caller) == FunctionEmissionStatus::Emitted)) - markKnownEmitted( - *this, Caller, Callee, Loc, [CheckCaller](Sema &S, FunctionDecl *FD) { - return CheckCaller && - S.getEmissionStatus(FD) == FunctionEmissionStatus::Emitted; - }); - else if (Caller) - DeviceCallGraph[Caller].insert({Callee, Loc}); - } -} - void Sema::checkOpenMPDeviceExpr(const Expr *E) { assert(getLangOpts().OpenMP && getLangOpts().OpenMPIsDevice && "OpenMP device compilation mode is expected."); @@ -2313,52 +2227,43 @@ void Sema::DestroyDataSharingAttributesStack() { delete DSAStack; } -void Sema::finalizeOpenMPDelayedAnalysis() { +void Sema::finalizeOpenMPDelayedAnalysis(const FunctionDecl *Caller, + const FunctionDecl *Callee, + SourceLocation Loc) { assert(LangOpts.OpenMP && "Expected OpenMP compilation mode."); - // Diagnose implicit declare target functions and their callees. - for (const auto &CallerCallees : DeviceCallGraph) { - Optional DevTy = - OMPDeclareTargetDeclAttr::getDeviceType( - CallerCallees.getFirst()->getMostRecentDecl()); - // Ignore host functions during device analyzis. - if (LangOpts.OpenMPIsDevice && DevTy && - *DevTy == OMPDeclareTargetDeclAttr::DT_Host) - continue; - // Ignore nohost functions during host analyzis. - if (!LangOpts.OpenMPIsDevice && DevTy && - *DevTy == OMPDeclareTargetDeclAttr::DT_NoHost) - continue; - for (const std::pair, SourceLocation> - &Callee : CallerCallees.getSecond()) { - const FunctionDecl *FD = Callee.first->getMostRecentDecl(); - Optional DevTy = - OMPDeclareTargetDeclAttr::getDeviceType(FD); - if (LangOpts.OpenMPIsDevice && DevTy && - *DevTy == OMPDeclareTargetDeclAttr::DT_Host) { - // Diagnose host function called during device codegen. - StringRef HostDevTy = getOpenMPSimpleClauseTypeName( - OMPC_device_type, OMPC_DEVICE_TYPE_host); - Diag(Callee.second, diag::err_omp_wrong_device_function_call) - << HostDevTy << 0; - Diag(FD->getAttr()->getLocation(), - diag::note_omp_marked_device_type_here) - << HostDevTy; - continue; - } + Optional DevTy = + OMPDeclareTargetDeclAttr::getDeviceType(Caller->getMostRecentDecl()); + // Ignore host functions during device analyzis. + if (LangOpts.OpenMPIsDevice && DevTy && + *DevTy == OMPDeclareTargetDeclAttr::DT_Host) + return; + // Ignore nohost functions during host analyzis. + if (!LangOpts.OpenMPIsDevice && DevTy && + *DevTy == OMPDeclareTargetDeclAttr::DT_NoHost) + return; + const FunctionDecl *FD = Callee->getMostRecentDecl(); + DevTy = OMPDeclareTargetDeclAttr::getDeviceType(FD); + if (LangOpts.OpenMPIsDevice && DevTy && + *DevTy == OMPDeclareTargetDeclAttr::DT_Host) { + // Diagnose host function called during device codegen. + StringRef HostDevTy = + getOpenMPSimpleClauseTypeName(OMPC_device_type, OMPC_DEVICE_TYPE_host); + Diag(Loc, diag::err_omp_wrong_device_function_call) << HostDevTy << 0; + Diag(FD->getAttr()->getLocation(), + diag::note_omp_marked_device_type_here) + << HostDevTy; + return; + } if (!LangOpts.OpenMPIsDevice && DevTy && *DevTy == OMPDeclareTargetDeclAttr::DT_NoHost) { // Diagnose nohost function called during host codegen. StringRef NoHostDevTy = getOpenMPSimpleClauseTypeName( OMPC_device_type, OMPC_DEVICE_TYPE_nohost); - Diag(Callee.second, diag::err_omp_wrong_device_function_call) - << NoHostDevTy << 1; + Diag(Loc, diag::err_omp_wrong_device_function_call) << NoHostDevTy << 1; Diag(FD->getAttr()->getLocation(), diag::note_omp_marked_device_type_here) << NoHostDevTy; - continue; } - } - } } void Sema::StartOpenMPDSABlock(OpenMPDirectiveKind DKind, @@ -17654,15 +17559,6 @@ Diag(FD->getLocation(), diag::note_defined_here) << FD; return; } - // Mark the function as must be emitted for the device. - Optional DevTy = - OMPDeclareTargetDeclAttr::getDeviceType(FD); - if (LangOpts.OpenMPIsDevice && Res.hasValue() && IdLoc.isValid() && - *DevTy != OMPDeclareTargetDeclAttr::DT_Host) - checkOpenMPDeviceFunction(IdLoc, FD, /*CheckForDelayedContext=*/false); - if (!LangOpts.OpenMPIsDevice && Res.hasValue() && IdLoc.isValid() && - *DevTy != OMPDeclareTargetDeclAttr::DT_NoHost) - checkOpenMPHostFunction(IdLoc, FD, /*CheckCaller=*/false); } if (auto *VD = dyn_cast(D)) { // Problem if any with var declared with incomplete type will be reported Index: clang/lib/Sema/UsedDeclVisitor.h =================================================================== --- clang/lib/Sema/UsedDeclVisitor.h +++ clang/lib/Sema/UsedDeclVisitor.h @@ -84,6 +84,18 @@ void VisitCXXDefaultArgExpr(CXXDefaultArgExpr *E) { asImpl().Visit(E->getExpr()); } + + void visitUsedDecl(SourceLocation Loc, Decl *D) { + if (auto *CD = dyn_cast(D)) { + if (auto *S = CD->getBody()) { + asImpl().Visit(S); + } + } else if (auto *CD = dyn_cast(D)) { + if (auto *S = CD->getBody()) { + asImpl().Visit(S); + } + } + } }; } // end namespace clang Index: clang/lib/Serialization/ASTReader.cpp =================================================================== --- clang/lib/Serialization/ASTReader.cpp +++ clang/lib/Serialization/ASTReader.cpp @@ -3773,6 +3773,11 @@ } break; } + + case DECLS_TO_CHECK_FOR_DEFERRED_DIAGS: + for (unsigned I = 0, N = Record.size(); I != N; ++I) + DeclsToCheckForDeferredDiags.push_back(getGlobalDeclID(F, Record[I])); + break; } } } @@ -8180,6 +8185,19 @@ UnusedLocalTypedefNameCandidates.clear(); } +void ASTReader::ReadDeclsToCheckForDeferredDiags( + llvm::SmallVector &Decls) { + for (unsigned I = 0, N = DeclsToCheckForDeferredDiags.size(); I != N; + ++I) { + auto *D = dyn_cast_or_null( + GetDecl(DeclsToCheckForDeferredDiags[I])); + if (D) + Decls.push_back(D); + } + DeclsToCheckForDeferredDiags.clear(); +} + + void ASTReader::ReadReferencedSelectors( SmallVectorImpl> &Sels) { if (ReferencedSelectorsData.empty()) Index: clang/lib/Serialization/ASTWriter.cpp =================================================================== --- clang/lib/Serialization/ASTWriter.cpp +++ clang/lib/Serialization/ASTWriter.cpp @@ -756,6 +756,7 @@ RECORD(DELETE_EXPRS_TO_ANALYZE); RECORD(CUDA_PRAGMA_FORCE_HOST_DEVICE_DEPTH); RECORD(PP_CONDITIONAL_STACK); + RECORD(DECLS_TO_CHECK_FOR_DEFERRED_DIAGS); // SourceManager Block. BLOCK(SOURCE_MANAGER_BLOCK); @@ -4671,6 +4672,11 @@ Buffer.data(), Buffer.size()); } + // Build a record containing all of the DeclsToCheckForDeferredDiags. + RecordData DeclsToCheckForDeferredDiags; + for (auto *D : SemaRef.DeclsToCheckForDeferredDiags) + AddDeclRef(D, DeclsToCheckForDeferredDiags); + RecordData DeclUpdatesOffsetsRecord; // Keep writing types, declarations, and declaration update records @@ -4762,6 +4768,11 @@ if (!SemaDeclRefs.empty()) Stream.EmitRecord(SEMA_DECL_REFS, SemaDeclRefs); + // Write the record containing decls to be checked for deferred diags. + if (!DeclsToCheckForDeferredDiags.empty()) + Stream.EmitRecord(DECLS_TO_CHECK_FOR_DEFERRED_DIAGS, + DeclsToCheckForDeferredDiags); + // Write the record containing CUDA-specific declaration references. if (!CUDASpecialDeclRefs.empty()) Stream.EmitRecord(CUDA_SPECIAL_DECL_REFS, CUDASpecialDeclRefs); Index: clang/test/OpenMP/declare_target_messages.cpp =================================================================== --- clang/test/OpenMP/declare_target_messages.cpp +++ clang/test/OpenMP/declare_target_messages.cpp @@ -162,17 +162,17 @@ #pragma omp declare target link(x) // expected-error {{'x' must not appear in both clauses 'to' and 'link'}} void bazz() {} -#pragma omp declare target to(bazz) device_type(nohost) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}} host5-note {{marked as 'device_type(nohost)' here}} +#pragma omp declare target to(bazz) device_type(nohost) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}} host5-note 3{{marked as 'device_type(nohost)' here}} void bazzz() {bazz();} #pragma omp declare target to(bazzz) device_type(nohost) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}} void any() {bazz();} // host5-error {{function with 'device_type(nohost)' is not available on host}} -void host1() {bazz();} -#pragma omp declare target to(host1) device_type(host) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}} dev5-note 2 {{marked as 'device_type(host)' here}} -void host2() {bazz();} +void host1() {bazz();} // host5-error {{function with 'device_type(nohost)' is not available on host}} +#pragma omp declare target to(host1) device_type(host) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}} dev5-note 4 {{marked as 'device_type(host)' here}} +void host2() {bazz();} //host5-error {{function with 'device_type(nohost)' is not available on host}} #pragma omp declare target to(host2) -void device() {host1();} +void device() {host1();} // dev5-error {{function with 'device_type(host)' is not available on device}} #pragma omp declare target to(device) device_type(nohost) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}} host5-note 2 {{marked as 'device_type(nohost)' here}} -void host3() {host1();} +void host3() {host1();} // dev5-error {{function with 'device_type(host)' is not available on device}} #pragma omp declare target to(host3) #pragma omp declare target Index: clang/test/OpenMP/nvptx_target_exceptions_messages.cpp =================================================================== --- clang/test/OpenMP/nvptx_target_exceptions_messages.cpp +++ clang/test/OpenMP/nvptx_target_exceptions_messages.cpp @@ -38,7 +38,7 @@ #pragma omp end declare target int c; -int bar() { return 1 + foo() + bar() + baz1() + baz2(); } +int bar() { return 1 + foo() + bar() + baz1() + baz2(); } // expected-note {{called by 'bar'}} int maini1() { int a; @@ -49,7 +49,7 @@ { S s(a); static long aaa = 23; - a = foo() + bar() + b + c + d + aa + aaa + FA(); + a = foo() + bar() + b + c + d + aa + aaa + FA(); // expected-note{{called by 'maini1'}} if (!a) throw "Error"; // expected-error {{cannot use 'throw' with exceptions disabled}} } Index: clang/test/SemaCUDA/bad-calls-on-same-line.cu =================================================================== --- clang/test/SemaCUDA/bad-calls-on-same-line.cu +++ clang/test/SemaCUDA/bad-calls-on-same-line.cu @@ -33,8 +33,8 @@ void host_fn() { hd(); - hd(); // expected-note {{function template specialization 'hd'}} + hd(); // expected-note@-1 {{called by 'host_fn'}} - hd(); // expected-note {{function template specialization 'hd'}} + hd(); // expected-note@-1 {{called by 'host_fn'}} } Index: clang/test/SemaCUDA/call-device-fn-from-host.cu =================================================================== --- clang/test/SemaCUDA/call-device-fn-from-host.cu +++ clang/test/SemaCUDA/call-device-fn-from-host.cu @@ -1,7 +1,7 @@ // RUN: %clang_cc1 %s --std=c++11 -triple x86_64-unknown-linux -emit-llvm -o - \ // RUN: -verify -verify-ignore-unexpected=note // RUN: %clang_cc1 %s --std=c++11 -triple x86_64-unknown-linux -emit-llvm -o - \ -// RUN: -verify -verify-ignore-unexpected=note -fopenmp +// RUN: -verify=expected,omp -verify-ignore-unexpected=note -fopenmp // Note: This test won't work with -fsyntax-only, because some of these errors // are emitted during codegen. @@ -39,7 +39,7 @@ } template __host__ __device__ void hd2() { device_fn(); } -// expected-error@-1 2 {{reference to __device__ function 'device_fn' in __host__ __device__ function}} +// expected-error@-1 {{reference to __device__ function 'device_fn' in __host__ __device__ function}} void host_fn() { hd2(); } __host__ __device__ void hd() { device_fn(); } Index: clang/test/SemaCUDA/call-host-fn-from-device.cu =================================================================== --- clang/test/SemaCUDA/call-host-fn-from-device.cu +++ clang/test/SemaCUDA/call-host-fn-from-device.cu @@ -56,14 +56,14 @@ } template __host__ __device__ void hd2() { host_fn(); } -// expected-error@-1 2 {{reference to __host__ function 'host_fn' in __host__ __device__ function}} +// expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}} __global__ void kernel() { hd2(); } __host__ __device__ void hd() { host_fn(); } // expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}} template __host__ __device__ void hd3() { host_fn(); } -// expected-error@-1 2 {{reference to __host__ function 'host_fn' in __host__ __device__ function}} +// expected-error@-1 {{reference to __host__ function 'host_fn' in __host__ __device__ function}} __device__ void device_fn() { hd3(); } // No error because this is never instantiated. Index: clang/test/SemaCUDA/openmp-target.cu =================================================================== --- clang/test/SemaCUDA/openmp-target.cu +++ clang/test/SemaCUDA/openmp-target.cu @@ -16,9 +16,9 @@ void bazzz() {bazz();} #pragma omp declare target to(bazzz) device_type(nohost) void any() {bazz();} // expected-error {{function with 'device_type(nohost)' is not available on host}} -void host1() {bazz();} +void host1() {bazz();} // expected-error {{function with 'device_type(nohost)' is not available on host}} #pragma omp declare target to(host1) device_type(host) -void host2() {bazz();} +void host2() {bazz();} // expected-error {{function with 'device_type(nohost)' is not available on host}} #pragma omp declare target to(host2) void device() {host1();} #pragma omp declare target to(device) device_type(nohost) Index: clang/test/SemaCUDA/trace-through-global.cu =================================================================== --- clang/test/SemaCUDA/trace-through-global.cu +++ clang/test/SemaCUDA/trace-through-global.cu @@ -38,7 +38,7 @@ // Notice that these two diagnostics are different: Because the call to hd1 // is not dependent on T, the call to hd1 comes from 'launch_kernel', while // the call to hd3, being dependent, comes from 'launch_kernel'. - hd1(); // expected-note {{called by 'launch_kernel'}} + hd1(); // expected-note {{called by 'launch_kernel'}} hd3(T()); // expected-note {{called by 'launch_kernel'}} }