Author: abataev Date: Fri Aug 23 09:11:14 2019 New Revision: 369775 URL: http://llvm.org/viewvc/llvm-project?rev=369775&view=rev Log: [OPENMP5.0]Add support for device_type clause in declare target construct.
OpenMP 5.0 introduced new clause for declare target directive, device_type clause, which may accept values host, nohost, and any. Host means that the function must be emitted only for the host, nohost - only for the device, and any - for both, device and the host. Modified: cfe/trunk/include/clang/AST/RecursiveASTVisitor.h cfe/trunk/include/clang/Basic/Attr.td cfe/trunk/include/clang/Basic/AttrDocs.td cfe/trunk/include/clang/Basic/DiagnosticParseKinds.td cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td cfe/trunk/include/clang/Basic/OpenMPKinds.def cfe/trunk/include/clang/Basic/OpenMPKinds.h cfe/trunk/include/clang/Sema/Sema.h cfe/trunk/lib/AST/ASTTypeTraits.cpp cfe/trunk/lib/AST/OpenMPClause.cpp cfe/trunk/lib/Basic/OpenMPKinds.cpp cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp cfe/trunk/lib/CodeGen/CodeGenModule.cpp cfe/trunk/lib/Parse/ParseOpenMP.cpp cfe/trunk/lib/Sema/Sema.cpp cfe/trunk/lib/Sema/SemaExpr.cpp cfe/trunk/lib/Sema/SemaOpenMP.cpp cfe/trunk/lib/Serialization/ASTReaderDecl.cpp cfe/trunk/test/OpenMP/declare_target_ast_print.cpp cfe/trunk/test/OpenMP/declare_target_codegen.cpp cfe/trunk/test/OpenMP/declare_target_messages.cpp cfe/trunk/test/OpenMP/nvptx_allocate_messages.cpp cfe/trunk/test/OpenMP/nvptx_asm_delayed_diags.c cfe/trunk/test/OpenMP/target_vla_messages.cpp Modified: cfe/trunk/include/clang/AST/RecursiveASTVisitor.h URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/AST/RecursiveASTVisitor.h?rev=369775&r1=369774&r2=369775&view=diff ============================================================================== --- cfe/trunk/include/clang/AST/RecursiveASTVisitor.h (original) +++ cfe/trunk/include/clang/AST/RecursiveASTVisitor.h Fri Aug 23 09:11:14 2019 @@ -2843,6 +2843,7 @@ bool RecursiveASTVisitor<Derived>::Trave #include "clang/Basic/OpenMPKinds.def" case OMPC_threadprivate: case OMPC_uniform: + case OMPC_device_type: case OMPC_unknown: break; } Modified: cfe/trunk/include/clang/Basic/Attr.td URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/Attr.td?rev=369775&r1=369774&r2=369775&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/Attr.td (original) +++ cfe/trunk/include/clang/Basic/Attr.td Fri Aug 23 09:11:14 2019 @@ -3207,11 +3207,16 @@ def OMPDeclareTargetDecl : InheritableAt let Args = [ EnumArgument<"MapType", "MapTypeTy", [ "to", "link" ], - [ "MT_To", "MT_Link" ]> + [ "MT_To", "MT_Link" ]>, + EnumArgument<"DevType", "DevTypeTy", + [ "host", "nohost", "any" ], + [ "DT_Host", "DT_NoHost", "DT_Any" ]> ]; let AdditionalMembers = [{ void printPrettyPragma(raw_ostream &OS, const PrintingPolicy &Policy) const { // Use fake syntax because it is for testing and debugging purpose only. + if (getDevType() != DT_Any) + OS << " device_type(" << ConvertDevTypeTyToStr(getDevType()) << ")"; if (getMapType() != MT_To) OS << ' ' << ConvertMapTypeTyToStr(getMapType()); } @@ -3224,6 +3229,14 @@ def OMPDeclareTargetDecl : InheritableAt return llvm::None; } + static llvm::Optional<DevTypeTy> getDeviceType(const ValueDecl *VD) { + if (!VD->hasAttrs()) + return llvm::None; + if (const auto *Attr = VD->getAttr<OMPDeclareTargetDeclAttr>()) + return Attr->getDevType(); + + return llvm::None; + } }]; } Modified: cfe/trunk/include/clang/Basic/AttrDocs.td URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/AttrDocs.td?rev=369775&r1=369774&r2=369775&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/AttrDocs.td (original) +++ cfe/trunk/include/clang/Basic/AttrDocs.td Fri Aug 23 09:11:14 2019 @@ -3176,6 +3176,27 @@ The syntax of the declare target directi #pragma omp declare target new-line declarations-definition-seq #pragma omp end declare target new-line + +or + + .. code-block:: c + + #pragma omp declare target (extended-list) new-line + +or + + .. code-block:: c + + #pragma omp declare target clause[ [,] clause ... ] new-line + +where clause is one of the following: + + + .. code-block:: c + + to(extended-list) + link(list) + device_type(host | nohost | any) }]; } Modified: cfe/trunk/include/clang/Basic/DiagnosticParseKinds.td URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticParseKinds.td?rev=369775&r1=369774&r2=369775&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/DiagnosticParseKinds.td (original) +++ cfe/trunk/include/clang/Basic/DiagnosticParseKinds.td Fri Aug 23 09:11:14 2019 @@ -1195,13 +1195,16 @@ def err_omp_declare_simd_inbranch_notinb def err_expected_end_declare_target : Error< "expected '#pragma omp end declare target'">; def err_omp_declare_target_unexpected_clause: Error< - "unexpected '%0' clause, only 'to' or 'link' clauses expected">; + "unexpected '%0' clause, only %select{'to' or 'link'|'to', 'link' or 'device_type'}1 clauses expected">; def err_omp_expected_clause: Error< "expected at least one clause on '#pragma omp %0' directive">; def err_omp_mapper_illegal_identifier : Error< "illegal OpenMP user-defined mapper identifier">; def err_omp_mapper_expected_declarator : Error< "expected declarator on 'omp declare mapper' directive">; +def warn_omp_more_one_device_type_clause : Warning< + "more than one 'device_type' clause is specified">, + InGroup<OpenMPClauses>; // Pragma loop support. def err_pragma_loop_missing_argument : Error< Modified: cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td?rev=369775&r1=369774&r2=369775&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td (original) +++ cfe/trunk/include/clang/Basic/DiagnosticSemaKinds.td Fri Aug 23 09:11:14 2019 @@ -9331,6 +9331,14 @@ def err_omp_unsupported_type : Error < "host requires %0 bit size %1 type support, but device '%2' does not support it">; def err_omp_lambda_capture_in_declare_target_not_to : Error< "variable captured in declare target region must appear in a to clause">; +def err_omp_device_type_mismatch : Error< + "'device_type(%0)' does not match previously specified 'device_type(%1)' for the same declaration">; +def err_omp_wrong_device_function_call : Error< + "function with 'device_type(%0)' is not available on %select{device|host}1">; +def note_omp_marked_device_type_here : Note<"marked as 'device_type(%0)' here">; +def warn_omp_declare_target_after_first_use : Warning< + "declaration marked as declare target after first use, it may lead to incorrect results">, + InGroup<OpenMPTarget>; } // end of OpenMP category let CategoryName = "Related Result Type Issue" in { Modified: cfe/trunk/include/clang/Basic/OpenMPKinds.def URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/OpenMPKinds.def?rev=369775&r1=369774&r2=369775&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/OpenMPKinds.def (original) +++ cfe/trunk/include/clang/Basic/OpenMPKinds.def Fri Aug 23 09:11:14 2019 @@ -191,6 +191,9 @@ #ifndef OPENMP_ALLOCATE_CLAUSE # define OPENMP_ALLOCATE_CLAUSE(Name) #endif +#ifndef OPENMP_DEVICE_TYPE_KIND +#define OPENMP_DEVICE_TYPE_KIND(Name) +#endif // OpenMP directives. OPENMP_DIRECTIVE(threadprivate) @@ -950,6 +953,12 @@ OPENMP_TASKGROUP_CLAUSE(allocate) // Clauses allowed for OpenMP directive 'declare mapper'. OPENMP_DECLARE_MAPPER_CLAUSE(map) +// Device types for 'device_type' clause. +OPENMP_DEVICE_TYPE_KIND(host) +OPENMP_DEVICE_TYPE_KIND(nohost) +OPENMP_DEVICE_TYPE_KIND(any) + +#undef OPENMP_DEVICE_TYPE_KIND #undef OPENMP_ALLOCATE_CLAUSE #undef OPENMP_DECLARE_MAPPER_CLAUSE #undef OPENMP_TASKGROUP_CLAUSE Modified: cfe/trunk/include/clang/Basic/OpenMPKinds.h URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Basic/OpenMPKinds.h?rev=369775&r1=369774&r2=369775&view=diff ============================================================================== --- cfe/trunk/include/clang/Basic/OpenMPKinds.h (original) +++ cfe/trunk/include/clang/Basic/OpenMPKinds.h Fri Aug 23 09:11:14 2019 @@ -35,6 +35,7 @@ enum OpenMPClauseKind { #include "clang/Basic/OpenMPKinds.def" OMPC_threadprivate, OMPC_uniform, + OMPC_device_type, OMPC_unknown }; @@ -152,6 +153,14 @@ enum OpenMPAtomicDefaultMemOrderClauseKi OMPC_ATOMIC_DEFAULT_MEM_ORDER_unknown }; +/// OpenMP device type for 'device_type' clause. +enum OpenMPDeviceType { +#define OPENMP_DEVICE_TYPE_KIND(Name) \ + OMPC_DEVICE_TYPE_##Name, +#include "clang/Basic/OpenMPKinds.def" + OMPC_DEVICE_TYPE_unknown +}; + /// Scheduling data for loop-based OpenMP directives. struct OpenMPScheduleTy final { OpenMPScheduleClauseKind Schedule = OMPC_SCHEDULE_unknown; Modified: cfe/trunk/include/clang/Sema/Sema.h URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/include/clang/Sema/Sema.h?rev=369775&r1=369774&r2=369775&view=diff ============================================================================== --- cfe/trunk/include/clang/Sema/Sema.h (original) +++ cfe/trunk/include/clang/Sema/Sema.h Fri Aug 23 09:11:14 2019 @@ -9001,10 +9001,18 @@ private: 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. @@ -9151,11 +9159,16 @@ public: bool ActOnStartOpenMPDeclareTargetDirective(SourceLocation Loc); /// Called at the end of target region i.e. '#pragme omp end declare target'. void ActOnFinishOpenMPDeclareTargetDirective(); + /// Searches for the provided declaration name for OpenMP declare target + /// directive. + NamedDecl * + lookupOpenMPDeclareTargetName(Scope *CurScope, CXXScopeSpec &ScopeSpec, + const DeclarationNameInfo &Id, + NamedDeclSetType &SameDirectiveDecls); /// Called on correct id-expression from the '#pragma omp declare target'. - void ActOnOpenMPDeclareTargetName(Scope *CurScope, CXXScopeSpec &ScopeSpec, - const DeclarationNameInfo &Id, + void ActOnOpenMPDeclareTargetName(NamedDecl *ND, SourceLocation Loc, OMPDeclareTargetDeclAttr::MapTypeTy MT, - NamedDeclSetType &SameDirectiveDecls); + OMPDeclareTargetDeclAttr::DevTypeTy DT); /// Check declaration inside target region. void checkDeclIsAllowedInOpenMPTarget(Expr *E, Decl *D, @@ -10539,6 +10552,21 @@ public: /// // Otherwise, continue parsing as normal. DeviceDiagBuilder diagIfOpenMPDeviceCode(SourceLocation Loc, unsigned DiagID); + /// Creates a DeviceDiagBuilder that emits the diagnostic if the current + /// context is "used as host code". + /// + /// - If CurContext is a `declare target` function or it is known that the + /// function is emitted for the host, emits the diagnostics immediately. + /// - If CurContext is a non-host function, just ignore it. + /// + /// Example usage: + /// + /// // Variable-length arrays are not allowed in NVPTX device code. + /// if (diagIfOpenMPHostode(Loc, diag::err_vla_unsupported)) + /// return ExprError(); + /// // Otherwise, continue parsing as normal. + DeviceDiagBuilder diagIfOpenMPHostCode(SourceLocation Loc, unsigned DiagID); + DeviceDiagBuilder targetDiag(SourceLocation Loc, unsigned DiagID); enum CUDAFunctionTarget { Modified: cfe/trunk/lib/AST/ASTTypeTraits.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/ASTTypeTraits.cpp?rev=369775&r1=369774&r2=369775&view=diff ============================================================================== --- cfe/trunk/lib/AST/ASTTypeTraits.cpp (original) +++ cfe/trunk/lib/AST/ASTTypeTraits.cpp Fri Aug 23 09:11:14 2019 @@ -116,6 +116,7 @@ ASTNodeKind ASTNodeKind::getFromNode(con #include "clang/Basic/OpenMPKinds.def" case OMPC_threadprivate: case OMPC_uniform: + case OMPC_device_type: case OMPC_unknown: llvm_unreachable("unexpected OpenMP clause kind"); } Modified: cfe/trunk/lib/AST/OpenMPClause.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/AST/OpenMPClause.cpp?rev=369775&r1=369774&r2=369775&view=diff ============================================================================== --- cfe/trunk/lib/AST/OpenMPClause.cpp (original) +++ cfe/trunk/lib/AST/OpenMPClause.cpp Fri Aug 23 09:11:14 2019 @@ -43,6 +43,7 @@ OMPClause::child_range OMPClause::used_c #include "clang/Basic/OpenMPKinds.def" case OMPC_threadprivate: case OMPC_uniform: + case OMPC_device_type: case OMPC_unknown: break; } @@ -127,6 +128,7 @@ const OMPClauseWithPreInit *OMPClauseWit case OMPC_reverse_offload: case OMPC_dynamic_allocators: case OMPC_atomic_default_mem_order: + case OMPC_device_type: break; } @@ -203,6 +205,7 @@ const OMPClauseWithPostUpdate *OMPClause case OMPC_reverse_offload: case OMPC_dynamic_allocators: case OMPC_atomic_default_mem_order: + case OMPC_device_type: break; } Modified: cfe/trunk/lib/Basic/OpenMPKinds.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Basic/OpenMPKinds.cpp?rev=369775&r1=369774&r2=369775&view=diff ============================================================================== --- cfe/trunk/lib/Basic/OpenMPKinds.cpp (original) +++ cfe/trunk/lib/Basic/OpenMPKinds.cpp Fri Aug 23 09:11:14 2019 @@ -55,6 +55,7 @@ OpenMPClauseKind clang::getOpenMPClauseK #define OPENMP_CLAUSE(Name, Class) .Case(#Name, OMPC_##Name) #include "clang/Basic/OpenMPKinds.def" .Case("uniform", OMPC_uniform) + .Case("device_type", OMPC_device_type) .Default(OMPC_unknown); } @@ -71,6 +72,8 @@ const char *clang::getOpenMPClauseName(O return "uniform"; case OMPC_threadprivate: return "threadprivate or thread local"; + case OMPC_device_type: + return "device_type"; } llvm_unreachable("Invalid OpenMP clause kind"); } @@ -145,6 +148,11 @@ unsigned clang::getOpenMPSimpleClauseTyp .Case(#Name, OMPC_ATOMIC_DEFAULT_MEM_ORDER_##Name) #include "clang/Basic/OpenMPKinds.def" .Default(OMPC_ATOMIC_DEFAULT_MEM_ORDER_unknown); + case OMPC_device_type: + return llvm::StringSwitch<OpenMPDeviceType>(Str) +#define OPENMP_DEVICE_TYPE_KIND(Name) .Case(#Name, OMPC_DEVICE_TYPE_##Name) +#include "clang/Basic/OpenMPKinds.def" + .Default(OMPC_DEVICE_TYPE_unknown); case OMPC_unknown: case OMPC_threadprivate: case OMPC_if: @@ -328,6 +336,16 @@ const char *clang::getOpenMPSimpleClause #include "clang/Basic/OpenMPKinds.def" } llvm_unreachable("Invalid OpenMP 'atomic_default_mem_order' clause type"); + case OMPC_device_type: + switch (Type) { + case OMPC_DEVICE_TYPE_unknown: + return "unknown"; +#define OPENMP_DEVICE_TYPE_KIND(Name) \ + case OMPC_DEVICE_TYPE_##Name: \ + return #Name; +#include "clang/Basic/OpenMPKinds.def" + } + llvm_unreachable("Invalid OpenMP 'device_type' clause type"); case OMPC_unknown: case OMPC_threadprivate: case OMPC_if: Modified: cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp?rev=369775&r1=369774&r2=369775&view=diff ============================================================================== --- cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp (original) +++ cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp Fri Aug 23 09:11:14 2019 @@ -9604,14 +9604,28 @@ void CGOpenMPRuntime::scanForTargetRegio bool CGOpenMPRuntime::emitTargetFunctions(GlobalDecl GD) { // If emitting code for the host, we do not process FD here. Instead we do // the normal code generation. - if (!CGM.getLangOpts().OpenMPIsDevice) + if (!CGM.getLangOpts().OpenMPIsDevice) { + if (const auto *FD = dyn_cast<FunctionDecl>(GD.getDecl())) { + Optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy = + OMPDeclareTargetDeclAttr::getDeviceType(FD); + // Do not emit device_type(nohost) functions for the host. + if (DevTy && *DevTy == OMPDeclareTargetDeclAttr::DT_NoHost) + return true; + } return false; + } const ValueDecl *VD = cast<ValueDecl>(GD.getDecl()); StringRef Name = CGM.getMangledName(GD); // Try to detect target regions in the function. - if (const auto *FD = dyn_cast<FunctionDecl>(VD)) + if (const auto *FD = dyn_cast<FunctionDecl>(VD)) { scanForTargetRegionsFunctions(FD->getBody(), Name); + Optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy = + OMPDeclareTargetDeclAttr::getDeviceType(FD); + // Do not emit device_type(nohost) functions for the host. + if (DevTy && *DevTy == OMPDeclareTargetDeclAttr::DT_Host) + return true; + } // Do not to emit function if it is not marked as declare target. return !OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD) && Modified: cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp?rev=369775&r1=369774&r2=369775&view=diff ============================================================================== --- cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp (original) +++ cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp Fri Aug 23 09:11:14 2019 @@ -4021,6 +4021,7 @@ static void emitOMPAtomicExpr(CodeGenFun case OMPC_reverse_offload: case OMPC_dynamic_allocators: case OMPC_atomic_default_mem_order: + case OMPC_device_type: llvm_unreachable("Clause is not allowed in 'omp atomic'."); } } Modified: cfe/trunk/lib/CodeGen/CodeGenModule.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/CodeGen/CodeGenModule.cpp?rev=369775&r1=369774&r2=369775&view=diff ============================================================================== --- cfe/trunk/lib/CodeGen/CodeGenModule.cpp (original) +++ cfe/trunk/lib/CodeGen/CodeGenModule.cpp Fri Aug 23 09:11:14 2019 @@ -2122,6 +2122,10 @@ void CodeGenModule::EmitDeferred() { if (!GV->isDeclaration()) continue; + // If this is OpenMP, check if it is legal to emit this global normally. + if (LangOpts.OpenMP && OpenMPRuntime && OpenMPRuntime->emitTargetGlobal(D)) + continue; + // Otherwise, emit the definition and move on to the next one. EmitGlobalDefinition(D, GV); @@ -2318,11 +2322,20 @@ bool CodeGenModule::MustBeEmitted(const } bool CodeGenModule::MayBeEmittedEagerly(const ValueDecl *Global) { - if (const auto *FD = dyn_cast<FunctionDecl>(Global)) + if (const auto *FD = dyn_cast<FunctionDecl>(Global)) { if (FD->getTemplateSpecializationKind() == TSK_ImplicitInstantiation) // Implicit template instantiations may change linkage if they are later // explicitly instantiated, so they should not be emitted eagerly. return false; + // In OpenMP 5.0 function may be marked as device_type(nohost) and we should + // not emit them eagerly unless we sure that the function must be emitted on + // the host. + if (LangOpts.OpenMP >= 50 && !LangOpts.OpenMPSimd && + !LangOpts.OpenMPIsDevice && + !OMPDeclareTargetDeclAttr::getDeviceType(FD) && + !FD->isUsed(/*CheckUsedAttr=*/false) && !FD->isReferenced()) + return false; + } if (const auto *VD = dyn_cast<VarDecl>(Global)) if (Context.getInlineVariableDefinitionKind(VD) == ASTContext::InlineVariableDefinitionKind::WeakUnknown) @@ -2445,8 +2458,7 @@ void CodeGenModule::EmitGlobal(GlobalDec } if (LangOpts.OpenMP) { - // If this is OpenMP device, check if it is legal to emit this global - // normally. + // If this is OpenMP, check if it is legal to emit this global normally. if (OpenMPRuntime && OpenMPRuntime->emitTargetGlobal(GD)) return; if (auto *DRD = dyn_cast<OMPDeclareReductionDecl>(Global)) { Modified: cfe/trunk/lib/Parse/ParseOpenMP.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Parse/ParseOpenMP.cpp?rev=369775&r1=369774&r2=369775&view=diff ============================================================================== --- cfe/trunk/lib/Parse/ParseOpenMP.cpp (original) +++ cfe/trunk/lib/Parse/ParseOpenMP.cpp Fri Aug 23 09:11:14 2019 @@ -782,25 +782,114 @@ Parser::ParseOMPDeclareSimdClauses(Parse LinModifiers, Steps, SourceRange(Loc, EndLoc)); } +/// Parsing of simple OpenMP clauses like 'default' or 'proc_bind'. +/// +/// default-clause: +/// 'default' '(' 'none' | 'shared' ') +/// +/// proc_bind-clause: +/// 'proc_bind' '(' 'master' | 'close' | 'spread' ') +/// +/// device_type-clause: +/// 'device_type' '(' 'host' | 'nohost' | 'any' )' +namespace { + struct SimpleClauseData { + unsigned Type; + SourceLocation Loc; + SourceLocation LOpen; + SourceLocation TypeLoc; + SourceLocation RLoc; + SimpleClauseData(unsigned Type, SourceLocation Loc, SourceLocation LOpen, + SourceLocation TypeLoc, SourceLocation RLoc) + : Type(Type), Loc(Loc), LOpen(LOpen), TypeLoc(TypeLoc), RLoc(RLoc) {} + }; +} // anonymous namespace + +static Optional<SimpleClauseData> +parseOpenMPSimpleClause(Parser &P, OpenMPClauseKind Kind) { + const Token &Tok = P.getCurToken(); + SourceLocation Loc = Tok.getLocation(); + SourceLocation LOpen = P.ConsumeToken(); + // Parse '('. + BalancedDelimiterTracker T(P, tok::l_paren, tok::annot_pragma_openmp_end); + if (T.expectAndConsume(diag::err_expected_lparen_after, + getOpenMPClauseName(Kind))) + return llvm::None; + + unsigned Type = getOpenMPSimpleClauseType( + Kind, Tok.isAnnotation() ? "" : P.getPreprocessor().getSpelling(Tok)); + SourceLocation TypeLoc = Tok.getLocation(); + if (Tok.isNot(tok::r_paren) && Tok.isNot(tok::comma) && + Tok.isNot(tok::annot_pragma_openmp_end)) + P.ConsumeAnyToken(); + + // Parse ')'. + SourceLocation RLoc = Tok.getLocation(); + if (!T.consumeClose()) + RLoc = T.getCloseLocation(); + + return SimpleClauseData(Type, Loc, LOpen, TypeLoc, RLoc); +} + Parser::DeclGroupPtrTy Parser::ParseOMPDeclareTargetClauses() { // OpenMP 4.5 syntax with list of entities. Sema::NamedDeclSetType SameDirectiveDecls; + SmallVector<std::tuple<OMPDeclareTargetDeclAttr::MapTypeTy, SourceLocation, + NamedDecl *>, + 4> + DeclareTargetDecls; + OMPDeclareTargetDeclAttr::DevTypeTy DT = OMPDeclareTargetDeclAttr::DT_Any; + SourceLocation DeviceTypeLoc; while (Tok.isNot(tok::annot_pragma_openmp_end)) { OMPDeclareTargetDeclAttr::MapTypeTy MT = OMPDeclareTargetDeclAttr::MT_To; if (Tok.is(tok::identifier)) { IdentifierInfo *II = Tok.getIdentifierInfo(); StringRef ClauseName = II->getName(); - // Parse 'to|link' clauses. - if (!OMPDeclareTargetDeclAttr::ConvertStrToMapTypeTy(ClauseName, MT)) { - Diag(Tok, diag::err_omp_declare_target_unexpected_clause) << ClauseName; + bool IsDeviceTypeClause = + getLangOpts().OpenMP >= 50 && + getOpenMPClauseKind(ClauseName) == OMPC_device_type; + // Parse 'to|link|device_type' clauses. + if (!OMPDeclareTargetDeclAttr::ConvertStrToMapTypeTy(ClauseName, MT) && + !IsDeviceTypeClause) { + Diag(Tok, diag::err_omp_declare_target_unexpected_clause) + << ClauseName << (getLangOpts().OpenMP >= 50 ? 1 : 0); break; } + // Parse 'device_type' clause and go to next clause if any. + if (IsDeviceTypeClause) { + Optional<SimpleClauseData> DevTypeData = + parseOpenMPSimpleClause(*this, OMPC_device_type); + if (DevTypeData.hasValue()) { + if (DeviceTypeLoc.isValid()) { + // We already saw another device_type clause, diagnose it. + Diag(DevTypeData.getValue().Loc, + diag::warn_omp_more_one_device_type_clause); + } + switch(static_cast<OpenMPDeviceType>(DevTypeData.getValue().Type)) { + case OMPC_DEVICE_TYPE_any: + DT = OMPDeclareTargetDeclAttr::DT_Any; + break; + case OMPC_DEVICE_TYPE_host: + DT = OMPDeclareTargetDeclAttr::DT_Host; + break; + case OMPC_DEVICE_TYPE_nohost: + DT = OMPDeclareTargetDeclAttr::DT_NoHost; + break; + case OMPC_DEVICE_TYPE_unknown: + llvm_unreachable("Unexpected device_type"); + } + DeviceTypeLoc = DevTypeData.getValue().Loc; + } + continue; + } ConsumeToken(); } - auto &&Callback = [this, MT, &SameDirectiveDecls]( - CXXScopeSpec &SS, DeclarationNameInfo NameInfo) { - Actions.ActOnOpenMPDeclareTargetName(getCurScope(), SS, NameInfo, MT, - SameDirectiveDecls); + auto &&Callback = [this, MT, &DeclareTargetDecls, &SameDirectiveDecls]( + CXXScopeSpec &SS, DeclarationNameInfo NameInfo) { + NamedDecl *ND = Actions.lookupOpenMPDeclareTargetName( + getCurScope(), SS, NameInfo, SameDirectiveDecls); + if (ND) + DeclareTargetDecls.emplace_back(MT, NameInfo.getLoc(), ND); }; if (ParseOpenMPSimpleVarList(OMPD_declare_target, Callback, /*AllowScopeSpecifier=*/true)) @@ -812,6 +901,15 @@ Parser::DeclGroupPtrTy Parser::ParseOMPD } SkipUntil(tok::annot_pragma_openmp_end, StopBeforeMatch); ConsumeAnyToken(); + for (auto &MTLocDecl : DeclareTargetDecls) { + OMPDeclareTargetDeclAttr::MapTypeTy MT; + SourceLocation Loc; + NamedDecl *ND; + std::tie(MT, Loc, ND) = MTLocDecl; + // device_type clause is applied only to functions. + Actions.ActOnOpenMPDeclareTargetName( + ND, Loc, MT, isa<VarDecl>(ND) ? OMPDeclareTargetDeclAttr::DT_Any : DT); + } SmallVector<Decl *, 4> Decls(SameDirectiveDecls.begin(), SameDirectiveDecls.end()); if (Decls.empty()) @@ -1712,6 +1810,7 @@ OMPClause *Parser::ParseOpenMPClause(Ope case OMPC_allocate: Clause = ParseOpenMPVarListClause(DKind, CKind, WrongDirective); break; + case OMPC_device_type: case OMPC_unknown: Diag(Tok, diag::warn_omp_extra_tokens_at_eol) << getOpenMPDirectiveName(DKind); @@ -1811,29 +1910,12 @@ OMPClause *Parser::ParseOpenMPSingleExpr /// OMPClause *Parser::ParseOpenMPSimpleClause(OpenMPClauseKind Kind, bool ParseOnly) { - SourceLocation Loc = Tok.getLocation(); - SourceLocation LOpen = ConsumeToken(); - // Parse '('. - BalancedDelimiterTracker T(*this, tok::l_paren, tok::annot_pragma_openmp_end); - if (T.expectAndConsume(diag::err_expected_lparen_after, - getOpenMPClauseName(Kind))) - return nullptr; - - unsigned Type = getOpenMPSimpleClauseType( - Kind, Tok.isAnnotation() ? "" : PP.getSpelling(Tok)); - SourceLocation TypeLoc = Tok.getLocation(); - if (Tok.isNot(tok::r_paren) && Tok.isNot(tok::comma) && - Tok.isNot(tok::annot_pragma_openmp_end)) - ConsumeAnyToken(); - - // Parse ')'. - SourceLocation RLoc = Tok.getLocation(); - if (!T.consumeClose()) - RLoc = T.getCloseLocation(); - - if (ParseOnly) + llvm::Optional<SimpleClauseData> Val = parseOpenMPSimpleClause(*this, Kind); + if (!Val || ParseOnly) return nullptr; - return Actions.ActOnOpenMPSimpleClause(Kind, Type, TypeLoc, LOpen, Loc, RLoc); + return Actions.ActOnOpenMPSimpleClause( + Kind, Val.getValue().Type, Val.getValue().TypeLoc, Val.getValue().LOpen, + Val.getValue().Loc, Val.getValue().RLoc); } /// Parsing of OpenMP clauses like 'ordered'. Modified: cfe/trunk/lib/Sema/Sema.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/Sema.cpp?rev=369775&r1=369774&r2=369775&view=diff ============================================================================== --- cfe/trunk/lib/Sema/Sema.cpp (original) +++ cfe/trunk/lib/Sema/Sema.cpp Fri Aug 23 09:11:14 2019 @@ -913,6 +913,10 @@ void Sema::ActOnEndOfTranslationUnitFrag PerformPendingInstantiations(); } + // Finalize analysis of OpenMP-specific constructs. + if (LangOpts.OpenMP) + finalizeOpenMPDelayedAnalysis(); + assert(LateParsedInstantiations.empty() && "end of TU template instantiation should not create more " "late-parsed templates"); @@ -1542,8 +1546,9 @@ void Sema::markKnownEmitted( } Sema::DeviceDiagBuilder Sema::targetDiag(SourceLocation Loc, unsigned DiagID) { - if (LangOpts.OpenMP && LangOpts.OpenMPIsDevice) - return diagIfOpenMPDeviceCode(Loc, DiagID); + if (LangOpts.OpenMP) + return LangOpts.OpenMPIsDevice ? diagIfOpenMPDeviceCode(Loc, DiagID) + : diagIfOpenMPHostCode(Loc, DiagID); if (getLangOpts().CUDA) return getLangOpts().CUDAIsDevice ? CUDADiagIfDeviceCode(Loc, DiagID) : CUDADiagIfHostCode(Loc, DiagID); Modified: cfe/trunk/lib/Sema/SemaExpr.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaExpr.cpp?rev=369775&r1=369774&r2=369775&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaExpr.cpp (original) +++ cfe/trunk/lib/Sema/SemaExpr.cpp Fri Aug 23 09:11:14 2019 @@ -15350,9 +15350,13 @@ void Sema::MarkFunctionReferenced(Source CheckCompleteParameterTypesForMangler(*this, Func, Loc); Func->markUsed(Context); + } - if (LangOpts.OpenMP && LangOpts.OpenMPIsDevice) + if (LangOpts.OpenMP) { + if (LangOpts.OpenMPIsDevice) checkOpenMPDeviceFunction(Loc, Func); + else + checkOpenMPHostFunction(Loc, Func); } } @@ -17745,4 +17749,4 @@ ExprResult Sema::ActOnObjCAvailabilityCh return new (Context) ObjCAvailabilityCheckExpr(Version, AtLoc, RParen, Context.BoolTy); -} \ No newline at end of file +} Modified: cfe/trunk/lib/Sema/SemaOpenMP.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Sema/SemaOpenMP.cpp?rev=369775&r1=369774&r2=369775&view=diff ============================================================================== --- cfe/trunk/lib/Sema/SemaOpenMP.cpp (original) +++ cfe/trunk/lib/Sema/SemaOpenMP.cpp Fri Aug 23 09:11:14 2019 @@ -1556,32 +1556,102 @@ static bool isOpenMPDeviceDelayedContext !S.isInOpenMPDeclareTargetContext(); } +namespace { +/// Status of the function emission on the host/device. +enum class FunctionEmissionStatus { + Emitted, + Discarded, + Unknown, +}; +} // anonymous namespace + /// Do we know that we will eventually codegen the given function? -static bool isKnownEmitted(Sema &S, FunctionDecl *FD) { +static FunctionEmissionStatus isKnownDeviceEmitted(Sema &S, FunctionDecl *FD) { assert(S.LangOpts.OpenMP && S.LangOpts.OpenMPIsDevice && "Expected OpenMP device compilation."); // Templates are emitted when they're instantiated. if (FD->isDependentContext()) - return false; + return FunctionEmissionStatus::Discarded; - if (OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration( - FD->getCanonicalDecl())) - return true; + Optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy = + OMPDeclareTargetDeclAttr::getDeviceType(FD->getCanonicalDecl()); + if (DevTy.hasValue()) + return (*DevTy == OMPDeclareTargetDeclAttr::DT_Host) + ? FunctionEmissionStatus::Discarded + : FunctionEmissionStatus::Emitted; // Otherwise, the function is known-emitted if it's in our set of // known-emitted functions. - return S.DeviceKnownEmittedFns.count(FD) > 0; + return (S.DeviceKnownEmittedFns.count(FD) > 0) + ? FunctionEmissionStatus::Emitted + : FunctionEmissionStatus::Unknown; } Sema::DeviceDiagBuilder Sema::diagIfOpenMPDeviceCode(SourceLocation Loc, unsigned DiagID) { assert(LangOpts.OpenMP && LangOpts.OpenMPIsDevice && "Expected OpenMP device compilation."); - return DeviceDiagBuilder((isOpenMPDeviceDelayedContext(*this) && - !isKnownEmitted(*this, getCurFunctionDecl())) - ? DeviceDiagBuilder::K_Deferred - : DeviceDiagBuilder::K_Immediate, - Loc, DiagID, getCurFunctionDecl(), *this); + FunctionEmissionStatus FES = + isKnownDeviceEmitted(*this, getCurFunctionDecl()); + DeviceDiagBuilder::Kind Kind = DeviceDiagBuilder::K_Nop; + switch (FES) { + case FunctionEmissionStatus::Emitted: + Kind = DeviceDiagBuilder::K_Immediate; + break; + case FunctionEmissionStatus::Unknown: + Kind = isOpenMPDeviceDelayedContext(*this) ? DeviceDiagBuilder::K_Deferred + : DeviceDiagBuilder::K_Immediate; + break; + case FunctionEmissionStatus::Discarded: + Kind = DeviceDiagBuilder::K_Nop; + break; + } + + return DeviceDiagBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this); +} + +/// Do we know that we will eventually codegen the given function? +static FunctionEmissionStatus isKnownHostEmitted(Sema &S, FunctionDecl *FD) { + assert(S.LangOpts.OpenMP && !S.LangOpts.OpenMPIsDevice && + "Expected OpenMP host compilation."); + // In OpenMP 4.5 all the functions are host functions. + if (S.LangOpts.OpenMP <= 45) + return FunctionEmissionStatus::Emitted; + + Optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy = + OMPDeclareTargetDeclAttr::getDeviceType(FD->getCanonicalDecl()); + if (DevTy.hasValue()) + return (*DevTy == OMPDeclareTargetDeclAttr::DT_NoHost) + ? FunctionEmissionStatus::Discarded + : FunctionEmissionStatus::Emitted; + + // Otherwise, the function is known-emitted if it's in our set of + // known-emitted functions. + return (S.DeviceKnownEmittedFns.count(FD) > 0) + ? FunctionEmissionStatus::Emitted + : FunctionEmissionStatus::Unknown; +} + +Sema::DeviceDiagBuilder Sema::diagIfOpenMPHostCode(SourceLocation Loc, + unsigned DiagID) { + assert(LangOpts.OpenMP && !LangOpts.OpenMPIsDevice && + "Expected OpenMP host compilation."); + FunctionEmissionStatus FES = + isKnownHostEmitted(*this, getCurFunctionDecl()); + DeviceDiagBuilder::Kind Kind = DeviceDiagBuilder::K_Nop; + switch (FES) { + case FunctionEmissionStatus::Emitted: + Kind = DeviceDiagBuilder::K_Immediate; + break; + case FunctionEmissionStatus::Unknown: + Kind = DeviceDiagBuilder::K_Deferred; + break; + case FunctionEmissionStatus::Discarded: + Kind = DeviceDiagBuilder::K_Nop; + break; + } + + return DeviceDiagBuilder(Kind, Loc, DiagID, getCurFunctionDecl(), *this); } void Sema::checkOpenMPDeviceFunction(SourceLocation Loc, FunctionDecl *Callee, @@ -1589,21 +1659,75 @@ void Sema::checkOpenMPDeviceFunction(Sou 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 && + (isKnownDeviceEmitted(*this, Caller) == FunctionEmissionStatus::Emitted || + (!isOpenMPDeviceDelayedContext(*this) && + isKnownDeviceEmitted(*this, Caller) == + FunctionEmissionStatus::Unknown)) && + isKnownDeviceEmitted(*this, Callee) == + FunctionEmissionStatus::Discarded) { + StringRef HostDevTy = + getOpenMPSimpleClauseTypeName(OMPC_device_type, OMPC_DEVICE_TYPE_host); + Diag(Loc, diag::err_omp_wrong_device_function_call) << HostDevTy << 0; + Diag(Callee->getAttr<OMPDeclareTargetDeclAttr>()->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 && isKnownEmitted(*this, Caller))) + (Caller && + isKnownDeviceEmitted(*this, Caller) == FunctionEmissionStatus::Emitted)) markKnownEmitted(*this, Caller, Callee, Loc, [CheckForDelayedContext](Sema &S, FunctionDecl *FD) { - return CheckForDelayedContext && isKnownEmitted(S, FD); + return CheckForDelayedContext && + isKnownDeviceEmitted(S, 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 && + isKnownHostEmitted(*this, Caller) == FunctionEmissionStatus::Emitted && + isKnownHostEmitted(*this, Callee) == FunctionEmissionStatus::Discarded) { + StringRef NoHostDevTy = getOpenMPSimpleClauseTypeName( + OMPC_device_type, OMPC_DEVICE_TYPE_nohost); + Diag(Loc, diag::err_omp_wrong_device_function_call) << NoHostDevTy << 1; + Diag(Callee->getAttr<OMPDeclareTargetDeclAttr>()->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 ((!CheckCaller && !Caller) || + (Caller && + isKnownHostEmitted(*this, Caller) == FunctionEmissionStatus::Emitted)) + markKnownEmitted( + *this, Caller, Callee, Loc, [CheckCaller](Sema &S, FunctionDecl *FD) { + return CheckCaller && + isKnownHostEmitted(S, 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."); @@ -1970,6 +2094,54 @@ bool Sema::isOpenMPTargetCapturedDecl(co void Sema::DestroyDataSharingAttributesStack() { delete DSAStack; } +void Sema::finalizeOpenMPDelayedAnalysis() { + assert(LangOpts.OpenMP && "Expected OpenMP compilation mode."); + // Diagnose implicit declare target functions and their callees. + for (const auto &CallerCallees : DeviceCallGraph) { + Optional<OMPDeclareTargetDeclAttr::DevTypeTy> 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<CanonicalDeclPtr<FunctionDecl>, SourceLocation> + &Callee : CallerCallees.getSecond()) { + const FunctionDecl *FD = Callee.first->getMostRecentDecl(); + Optional<OMPDeclareTargetDeclAttr::DevTypeTy> 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<OMPDeclareTargetDeclAttr>()->getLocation(), + diag::note_omp_marked_device_type_here) + << HostDevTy; + continue; + } + 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(FD->getAttr<OMPDeclareTargetDeclAttr>()->getLocation(), + diag::note_omp_marked_device_type_here) + << NoHostDevTy; + continue; + } + } + } +} + void Sema::StartOpenMPDSABlock(OpenMPDirectiveKind DKind, const DeclarationNameInfo &DirName, Scope *CurScope, SourceLocation Loc) { @@ -4415,6 +4587,7 @@ StmtResult Sema::ActOnOpenMPExecutableDi case OMPC_reverse_offload: case OMPC_dynamic_allocators: case OMPC_atomic_default_mem_order: + case OMPC_device_type: llvm_unreachable("Unexpected clause"); } for (Stmt *CC : C->children()) { @@ -9642,6 +9815,7 @@ OMPClause *Sema::ActOnOpenMPSingleExprCl case OMPC_reverse_offload: case OMPC_dynamic_allocators: case OMPC_atomic_default_mem_order: + case OMPC_device_type: llvm_unreachable("Clause is not allowed."); } return Res; @@ -10184,6 +10358,7 @@ static OpenMPDirectiveKind getOpenMPCapt case OMPC_reverse_offload: case OMPC_dynamic_allocators: case OMPC_atomic_default_mem_order: + case OMPC_device_type: llvm_unreachable("Unexpected OpenMP clause."); } return CaptureRegion; @@ -10577,6 +10752,7 @@ OMPClause *Sema::ActOnOpenMPSimpleClause case OMPC_unified_shared_memory: case OMPC_reverse_offload: case OMPC_dynamic_allocators: + case OMPC_device_type: llvm_unreachable("Clause is not allowed."); } return Res; @@ -10755,6 +10931,7 @@ OMPClause *Sema::ActOnOpenMPSingleExprWi case OMPC_reverse_offload: case OMPC_dynamic_allocators: case OMPC_atomic_default_mem_order: + case OMPC_device_type: llvm_unreachable("Clause is not allowed."); } return Res; @@ -10964,6 +11141,7 @@ OMPClause *Sema::ActOnOpenMPClause(OpenM case OMPC_use_device_ptr: case OMPC_is_device_ptr: case OMPC_atomic_default_mem_order: + case OMPC_device_type: llvm_unreachable("Clause is not allowed."); } return Res; @@ -11170,6 +11348,7 @@ OMPClause *Sema::ActOnOpenMPVarListClaus case OMPC_reverse_offload: case OMPC_dynamic_allocators: case OMPC_atomic_default_mem_order: + case OMPC_device_type: llvm_unreachable("Clause is not allowed."); } return Res; @@ -15333,16 +15512,15 @@ void Sema::ActOnFinishOpenMPDeclareTarge --DeclareTargetNestingLevel; } -void Sema::ActOnOpenMPDeclareTargetName(Scope *CurScope, - CXXScopeSpec &ScopeSpec, - const DeclarationNameInfo &Id, - OMPDeclareTargetDeclAttr::MapTypeTy MT, - NamedDeclSetType &SameDirectiveDecls) { +NamedDecl * +Sema::lookupOpenMPDeclareTargetName(Scope *CurScope, CXXScopeSpec &ScopeSpec, + const DeclarationNameInfo &Id, + NamedDeclSetType &SameDirectiveDecls) { LookupResult Lookup(*this, Id, LookupOrdinaryName); LookupParsedName(Lookup, CurScope, &ScopeSpec, true); if (Lookup.isAmbiguous()) - return; + return nullptr; Lookup.suppressDiagnostics(); if (!Lookup.isSingleResult()) { @@ -15353,33 +15531,56 @@ void Sema::ActOnOpenMPDeclareTargetName( diagnoseTypo(Corrected, PDiag(diag::err_undeclared_var_use_suggest) << Id.getName()); checkDeclIsAllowedInOpenMPTarget(nullptr, Corrected.getCorrectionDecl()); - return; + return nullptr; } Diag(Id.getLoc(), diag::err_undeclared_var_use) << Id.getName(); - return; + return nullptr; } NamedDecl *ND = Lookup.getAsSingle<NamedDecl>(); - if (isa<VarDecl>(ND) || isa<FunctionDecl>(ND) || - isa<FunctionTemplateDecl>(ND)) { - if (!SameDirectiveDecls.insert(cast<NamedDecl>(ND->getCanonicalDecl()))) - Diag(Id.getLoc(), diag::err_omp_declare_target_multiple) << Id.getName(); - llvm::Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res = - OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration( - cast<ValueDecl>(ND)); - if (!Res) { - auto *A = OMPDeclareTargetDeclAttr::CreateImplicit(Context, MT); - ND->addAttr(A); - if (ASTMutationListener *ML = Context.getASTMutationListener()) - ML->DeclarationMarkedOpenMPDeclareTarget(ND, A); - checkDeclIsAllowedInOpenMPTarget(nullptr, ND, Id.getLoc()); - } else if (*Res != MT) { - Diag(Id.getLoc(), diag::err_omp_declare_target_to_and_link) - << Id.getName(); - } - } else { + if (!isa<VarDecl>(ND) && !isa<FunctionDecl>(ND) && + !isa<FunctionTemplateDecl>(ND)) { Diag(Id.getLoc(), diag::err_omp_invalid_target_decl) << Id.getName(); + return nullptr; + } + if (!SameDirectiveDecls.insert(cast<NamedDecl>(ND->getCanonicalDecl()))) + Diag(Id.getLoc(), diag::err_omp_declare_target_multiple) << Id.getName(); + return ND; +} + +void Sema::ActOnOpenMPDeclareTargetName( + NamedDecl *ND, SourceLocation Loc, OMPDeclareTargetDeclAttr::MapTypeTy MT, + OMPDeclareTargetDeclAttr::DevTypeTy DT) { + assert((isa<VarDecl>(ND) || isa<FunctionDecl>(ND) || + isa<FunctionTemplateDecl>(ND)) && + "Expected variable, function or function template."); + + // Diagnose marking after use as it may lead to incorrect diagnosis and + // codegen. + if (LangOpts.OpenMP >= 50 && + (ND->isUsed(/*CheckUsedAttr=*/false) || ND->isReferenced())) + Diag(Loc, diag::warn_omp_declare_target_after_first_use); + + Optional<OMPDeclareTargetDeclAttr::DevTypeTy> DevTy = + OMPDeclareTargetDeclAttr::getDeviceType(cast<ValueDecl>(ND)); + if (DevTy.hasValue() && *DevTy != DT) { + Diag(Loc, diag::err_omp_device_type_mismatch) + << OMPDeclareTargetDeclAttr::ConvertDevTypeTyToStr(DT) + << OMPDeclareTargetDeclAttr::ConvertDevTypeTyToStr(*DevTy); + return; + } + Optional<OMPDeclareTargetDeclAttr::MapTypeTy> Res = + OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(cast<ValueDecl>(ND)); + if (!Res) { + auto *A = OMPDeclareTargetDeclAttr::CreateImplicit(Context, MT, DT, + SourceRange(Loc, Loc)); + ND->addAttr(A); + if (ASTMutationListener *ML = Context.getASTMutationListener()) + ML->DeclarationMarkedOpenMPDeclareTarget(ND, A); + checkDeclIsAllowedInOpenMPTarget(nullptr, ND, Loc); + } else if (*Res != MT) { + Diag(Loc, diag::err_omp_declare_target_to_and_link) << ND; } } @@ -15453,8 +15654,14 @@ void Sema::checkDeclIsAllowedInOpenMPTar return; } // Mark the function as must be emitted for the device. - if (LangOpts.OpenMPIsDevice && Res.hasValue() && IdLoc.isValid()) + Optional<OMPDeclareTargetDeclAttr::DevTypeTy> 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<ValueDecl>(D)) { // Problem if any with var declared with incomplete type will be reported @@ -15467,7 +15674,8 @@ void Sema::checkDeclIsAllowedInOpenMPTar if (isa<VarDecl>(D) || isa<FunctionDecl>(D) || isa<FunctionTemplateDecl>(D)) { auto *A = OMPDeclareTargetDeclAttr::CreateImplicit( - Context, OMPDeclareTargetDeclAttr::MT_To); + Context, OMPDeclareTargetDeclAttr::MT_To, + OMPDeclareTargetDeclAttr::DT_Any, SourceRange(IdLoc, IdLoc)); D->addAttr(A); if (ASTMutationListener *ML = Context.getASTMutationListener()) ML->DeclarationMarkedOpenMPDeclareTarget(D, A); Modified: cfe/trunk/lib/Serialization/ASTReaderDecl.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/lib/Serialization/ASTReaderDecl.cpp?rev=369775&r1=369774&r2=369775&view=diff ============================================================================== --- cfe/trunk/lib/Serialization/ASTReaderDecl.cpp (original) +++ cfe/trunk/lib/Serialization/ASTReaderDecl.cpp Fri Aug 23 09:11:14 2019 @@ -4571,12 +4571,15 @@ void ASTDeclReader::UpdateDecl(Decl *D, break; } - case UPD_DECL_MARKED_OPENMP_DECLARETARGET: + case UPD_DECL_MARKED_OPENMP_DECLARETARGET: { + OMPDeclareTargetDeclAttr::MapTypeTy MapType = + static_cast<OMPDeclareTargetDeclAttr::MapTypeTy>(Record.readInt()); + OMPDeclareTargetDeclAttr::DevTypeTy DevType = + static_cast<OMPDeclareTargetDeclAttr::DevTypeTy>(Record.readInt()); D->addAttr(OMPDeclareTargetDeclAttr::CreateImplicit( - Reader.getContext(), - static_cast<OMPDeclareTargetDeclAttr::MapTypeTy>(Record.readInt()), - ReadSourceRange())); + Reader.getContext(), MapType, DevType, ReadSourceRange())); break; + } case UPD_ADDED_ATTR_TO_RECORD: AttrVec Attrs; Modified: cfe/trunk/test/OpenMP/declare_target_ast_print.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/declare_target_ast_print.cpp?rev=369775&r1=369774&r2=369775&view=diff ============================================================================== --- cfe/trunk/test/OpenMP/declare_target_ast_print.cpp (original) +++ cfe/trunk/test/OpenMP/declare_target_ast_print.cpp Fri Aug 23 09:11:14 2019 @@ -2,6 +2,10 @@ // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -I %S/Inputs -emit-pch -o %t %s // RUN: %clang_cc1 -fopenmp -std=c++11 -include-pch %t -fsyntax-only -I %S/Inputs -verify %s -ast-print | FileCheck %s +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=50 -I %S/Inputs -ast-print %s | FileCheck %s --check-prefix=CHECK --check-prefix=OMP50 +// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -x c++ -std=c++11 -I %S/Inputs -emit-pch -o %t %s +// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -std=c++11 -include-pch %t -fsyntax-only -I %S/Inputs -verify %s -ast-print | FileCheck %s --check-prefix=CHECK --check-prefix=OMP50 + // RUN: %clang_cc1 -verify -fopenmp-simd -I %S/Inputs -ast-print %s | FileCheck %s // RUN: %clang_cc1 -fopenmp-simd -x c++ -std=c++11 -I %S/Inputs -emit-pch -o %t %s // RUN: %clang_cc1 -fopenmp-simd -std=c++11 -include-pch %t -fsyntax-only -I %S/Inputs -verify %s -ast-print | FileCheck %s @@ -10,7 +14,29 @@ #ifndef HEADER #define HEADER +#if _OPENMP == 201811 +void bar(); +#pragma omp declare target to(bar) device_type(any) +// OMP50: #pragma omp declare target{{$}} +// OMP50: void bar(); +// OMP50: #pragma omp end declare target{{$}} +void baz(); +#pragma omp declare target to(baz) device_type(nohost) +// OMP50: #pragma omp declare target device_type(nohost){{$}} +// OMP50: void baz(); +// OMP50: #pragma omp end declare target{{$}} +void bazz(); +#pragma omp declare target to(bazz) device_type(host) +// OMP50: #pragma omp declare target device_type(host){{$}} +// OMP50: void bazz(); +// OMP50: #pragma omp end declare target{{$}} +#endif // _OPENMP + int out_decl_target = 0; +#if _OPENMP == 201811 +#pragma omp declare target (out_decl_target) +#endif // _OPENMP + // CHECK: #pragma omp declare target{{$}} // CHECK: int out_decl_target = 0; // CHECK: #pragma omp end declare target{{$}} Modified: cfe/trunk/test/OpenMP/declare_target_codegen.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/declare_target_codegen.cpp?rev=369775&r1=369774&r2=369775&view=diff ============================================================================== --- cfe/trunk/test/OpenMP/declare_target_codegen.cpp (original) +++ cfe/trunk/test/OpenMP/declare_target_codegen.cpp Fri Aug 23 09:11:14 2019 @@ -3,6 +3,14 @@ // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -include-pch %t -o - | FileCheck %s +// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - -fopenmp-version=50 -DOMP5 | FileCheck %s --check-prefix HOST5 +// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc -fopenmp-version=50 -DOMP5 +// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -fopenmp-version=50 -DOMP5 | FileCheck %s --check-prefix DEV5 + +// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - -fopenmp-version=50 -DOMP5 | FileCheck %s --check-prefix SIMD-ONLY +// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc -fopenmp-version=50 -DOMP5 +// RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - -fopenmp-version=50 -DOMP5 | FileCheck %s --check-prefix SIMD-ONLY + // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm-bc %s -o %t-ppc-host.bc // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o -| FileCheck %s --check-prefix SIMD-ONLY // RUN: %clang_cc1 -verify -fopenmp-simd -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -emit-pch -o %t @@ -32,13 +40,13 @@ // CHECK-DAG: [[STAT:@.+stat]] = internal global %struct.S zeroinitializer, // CHECK-DAG: [[STAT_REF:@.+]] = internal constant %struct.S* [[STAT]] // CHECK-DAG: @out_decl_target ={{ dso_local | }}global i32 0, -// CHECK-DAG: @llvm.used = appending global [2 x i8*] [i8* bitcast (void ()* @__omp_offloading__{{.+}}_globals_l[[@LINE+80]]_ctor to i8*), i8* bitcast (void ()* @__omp_offloading__{{.+}}_stat_l[[@LINE+81]]_ctor to i8*)], +// CHECK-DAG: @llvm.used = appending global [2 x i8*] [i8* bitcast (void ()* @__omp_offloading__{{.+}}_globals_l[[@LINE+84]]_ctor to i8*), i8* bitcast (void ()* @__omp_offloading__{{.+}}_stat_l[[@LINE+85]]_ctor to i8*)], // CHECK-DAG: @llvm.compiler.used = appending global [1 x i8*] [i8* bitcast (%struct.S** [[STAT_REF]] to i8*)], // CHECK-DAG: define {{.*}}i32 @{{.*}}{{foo|bar|baz2|baz3|FA|f_method}}{{.*}}() // CHECK-DAG: define {{.*}}void @{{.*}}TemplateClass{{.*}}(%class.TemplateClass* %{{.*}}) // CHECK-DAG: define {{.*}}i32 @{{.*}}TemplateClass{{.*}}f_method{{.*}}(%class.TemplateClass* %{{.*}}) -// CHECK-DAG: define {{.*}}void @__omp_offloading__{{.*}}_globals_l[[@LINE+74]]_ctor() +// CHECK-DAG: define {{.*}}void @__omp_offloading__{{.*}}_globals_l[[@LINE+78]]_ctor() #ifndef HEADER #define HEADER @@ -68,6 +76,10 @@ int hhh = 0; #pragma omp declare target link(eee, fff, ggg, hhh) int out_decl_target = 0; +#ifdef OMP5 +#pragma omp declare target(out_decl_target) +#endif + #pragma omp declare target void lambda () { #ifdef __cpp_lambdas @@ -224,4 +236,18 @@ int main() { // CHECK-DAG: !{i32 1, !"ccc", i32 0, i32 {{[0-9]+}}} // CHECK-DAG: !{{{.+}}virtual_foo +#ifdef OMP5 +void host_fun() {} +#pragma omp declare target to(host_fun) device_type(host) +void device_fun() {} +#pragma omp declare target to(device_fun) device_type(nohost) +// HOST5-NOT: define {{.*}}void {{.*}}device_fun{{.*}} +// HOST5: define {{.*}}void {{.*}}host_fun{{.*}} +// HOST5-NOT: define {{.*}}void {{.*}}device_fun{{.*}} + +// DEV5-NOT: define {{.*}}void {{.*}}host_fun{{.*}} +// DEV5: define {{.*}}void {{.*}}device_fun{{.*}} +// DEV5-NOT: define {{.*}}void {{.*}}host_fun{{.*}} +#endif // OMP5 + #endif // HEADER Modified: cfe/trunk/test/OpenMP/declare_target_messages.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/declare_target_messages.cpp?rev=369775&r1=369774&r2=369775&view=diff ============================================================================== --- cfe/trunk/test/OpenMP/declare_target_messages.cpp (original) +++ cfe/trunk/test/OpenMP/declare_target_messages.cpp Fri Aug 23 09:11:14 2019 @@ -1,7 +1,10 @@ -// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp -fnoopenmp-use-tls -ferror-limit 100 -o - %s -// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp5 -fopenmp -fopenmp-version=50 -fnoopenmp-use-tls -ferror-limit 100 -o - %s - -// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify -fopenmp-simd -fnoopenmp-use-tls -ferror-limit 100 -o - %s +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp45 -fopenmp -fnoopenmp-use-tls -ferror-limit 100 -o - %s +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp5,host5 -fopenmp -fopenmp-version=50 -fnoopenmp-use-tls -ferror-limit 100 -o - %s +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp5,dev5 -fopenmp -fopenmp-is-device -fopenmp-targets=x86_64-apple-macos10.7.0 -aux-triple x86_64-apple-macos10.7.0 -fopenmp-version=50 -fnoopenmp-use-tls -ferror-limit 100 -o - %s + +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp5,host5 -fopenmp-simd -fopenmp-version=50 -fnoopenmp-use-tls -ferror-limit 100 -o - %s +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp5,host5 -fopenmp-simd -fopenmp-is-device -fopenmp-version=50 -fnoopenmp-use-tls -ferror-limit 100 -o - %s +// RUN: %clang_cc1 -triple x86_64-apple-macos10.7.0 -verify=expected,omp45 -fopenmp-simd -fnoopenmp-use-tls -ferror-limit 100 -o - %s #pragma omp end declare target // expected-error {{unexpected OpenMP directive '#pragma omp end declare target'}} @@ -14,17 +17,23 @@ __thread int t; // expected-note {{defin void f(); #pragma omp end declare target shared(a) // expected-warning {{extra tokens at the end of '#pragma omp end declare target' are ignored}} -#pragma omp declare target map(a) // expected-error {{unexpected 'map' clause, only 'to' or 'link' clauses expected}} +#pragma omp declare target map(a) // omp45-error {{unexpected 'map' clause, only 'to' or 'link' clauses expected}} omp5-error {{unexpected 'map' clause, only 'to', 'link' or 'device_type' clauses expected}} #pragma omp declare target to(foo1) // expected-error {{use of undeclared identifier 'foo1'}} #pragma omp declare target link(foo2) // expected-error {{use of undeclared identifier 'foo2'}} +#pragma omp declare target to(f) device_type(any) device_type(any) device_type(host) // omp45-error {{unexpected 'device_type' clause, only 'to' or 'link' clauses expected}} omp5-warning 2 {{more than one 'device_type' clause is specified}} omp5-error {{'device_type(host)' does not match previously specified 'device_type(any)' for the same declaration}} + void c(); void func() {} // expected-note {{'func' defined here}} -#pragma omp declare target link(func) allocate(a) // expected-error {{function name is not allowed in 'link' clause}} expected-error {{unexpected 'allocate' clause, only 'to' or 'link' clauses expected}} +#pragma omp declare target link(func) allocate(a) // expected-error {{function name is not allowed in 'link' clause}} omp45-error {{unexpected 'allocate' clause, only 'to' or 'link' clauses expected}} omp5-error {{unexpected 'allocate' clause, only 'to', 'link' or 'device_type' clauses expected}} + +void bar(); +void baz() {bar();} +#pragma omp declare target(bar) // omp5-warning {{declaration marked as declare target after first use, it may lead to incorrect results}} extern int b; @@ -152,4 +161,30 @@ namespace { #pragma omp declare target to(x) to(x) // expected-error {{'x' appears multiple times in clauses on the same declare target directive}} #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}} +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();} +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();} +#pragma omp declare target to(host2) +void device() {host1();} +#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();} +#pragma omp declare target to(host3) + +#pragma omp declare target +void any1() {any();} +void any2() {host1();} // dev5-error {{function with 'device_type(host)' is not available on device}} +void any3() {device();} // host5-error {{function with 'device_type(nohost)' is not available on host}} +void any4() {any2();} +#pragma omp end declare target + +void any5() {any();} +void any6() {host1();} // dev5-error {{function with 'device_type(host)' is not available on device}} +void any7() {device();} // host5-error {{function with 'device_type(nohost)' is not available on host}} +void any8() {any2();} + #pragma omp declare target // expected-error {{expected '#pragma omp end declare target'}} expected-note {{to match this '#pragma omp declare target'}} Modified: cfe/trunk/test/OpenMP/nvptx_allocate_messages.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/nvptx_allocate_messages.cpp?rev=369775&r1=369774&r2=369775&view=diff ============================================================================== --- cfe/trunk/test/OpenMP/nvptx_allocate_messages.cpp (original) +++ cfe/trunk/test/OpenMP/nvptx_allocate_messages.cpp Fri Aug 23 09:11:14 2019 @@ -58,7 +58,7 @@ template <class T> T foo() { #pragma omp allocate(v) allocator(omp_cgroup_mem_alloc) v = ST<T>::m; #if defined(DEVICE) && !defined(REQUIRES) -// expected-error@+2 2 {{expected an allocator expression inside of the target region; provide an allocator expression or use 'requires' directive with the 'dynamic_allocators' clause}} +// expected-error@+2 {{expected an allocator expression inside of the target region; provide an allocator expression or use 'requires' directive with the 'dynamic_allocators' clause}} #endif // DEVICE && !REQUIRES #pragma omp parallel private(v) allocate(v) v = 0; Modified: cfe/trunk/test/OpenMP/nvptx_asm_delayed_diags.c URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/nvptx_asm_delayed_diags.c?rev=369775&r1=369774&r2=369775&view=diff ============================================================================== --- cfe/trunk/test/OpenMP/nvptx_asm_delayed_diags.c (original) +++ cfe/trunk/test/OpenMP/nvptx_asm_delayed_diags.c Fri Aug 23 09:11:14 2019 @@ -2,6 +2,10 @@ // RUN: %clang_cc1 -verify -fopenmp -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized // RUN: %clang_cc1 -verify -DDIAGS -DIMMEDIATE -fopenmp -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized // RUN: %clang_cc1 -verify -DDIAGS -DDELAYED -fopenmp -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized +// RUN: %clang_cc1 -fopenmp -x c -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -fopenmp-version=50 -emit-llvm-bc %s -o %t-x86-host.bc +// RUN: %clang_cc1 -verify=expected,omp5 -fopenmp -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -fopenmp-version=50 %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized +// RUN: %clang_cc1 -verify=expected,omp5 -DDIAGS -DOMP5 -DIMMEDIATE -fopenmp -fopenmp-version=50 -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized +// RUN: %clang_cc1 -verify=expected,omp5 -DDIAGS -DOMP5 -DDELAYED -fopenmp -fopenmp-version=50 -x c -triple nvptx-unknown-unknown -aux-triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -fsyntax-only -Wuninitialized // REQUIRES: x86-registered-target // REQUIRES: nvptx-registered-target @@ -9,6 +13,22 @@ // expected-no-diagnostics #endif // DIAGS +#ifdef OMP5 +void bar(int r) { +#ifdef IMMEDIATE +// omp5-error@+4 {{invalid input constraint 'mx' in asm}} +#endif // IMMEDIATE + __asm__("PR3908 %[lf] %[xx] %[li] %[r]" + : [ r ] "+r"(r) + : [ lf ] "mx"(0), [ li ] "mr"(0), [ xx ] "x"((double)(0))); +} +#ifdef IMMEDIATE +#pragma omp declare target to(bar) device_type(nohost) +#else +#pragma omp declare target to(bar) device_type(host) +#endif // IMMEDIATE +#endif // OMP5 + void foo(int r) { #ifdef IMMEDIATE // expected-error@+4 {{invalid input constraint 'mx' in asm}} Modified: cfe/trunk/test/OpenMP/target_vla_messages.cpp URL: http://llvm.org/viewvc/llvm-project/cfe/trunk/test/OpenMP/target_vla_messages.cpp?rev=369775&r1=369774&r2=369775&view=diff ============================================================================== --- cfe/trunk/test/OpenMP/target_vla_messages.cpp (original) +++ cfe/trunk/test/OpenMP/target_vla_messages.cpp Fri Aug 23 09:11:14 2019 @@ -47,7 +47,7 @@ void target_template(int arg) { #pragma omp target { #ifdef NO_VLA - // expected-error@+2 2 {{variable length arrays are not supported for the current target}} + // expected-error@+2 {{variable length arrays are not supported for the current target}} #endif T vla[arg]; } _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits