diff --git a/clang/include/clang-c/Index.h b/clang/include/clang-c/Index.h --- a/clang/include/clang-c/Index.h +++ b/clang/include/clang-c/Index.h @@ -2568,7 +2568,11 @@ */ CXCursor_OMPScanDirective = 287, - CXCursor_LastStmt = CXCursor_OMPScanDirective, + /** OpenMP metadirective directive. + */ + CXCursor_OMPMetaDirective = 288, + + CXCursor_LastStmt = CXCursor_OMPMetaDirective, /** * Cursor that represents the translation unit itself. diff --git a/clang/include/clang/AST/RecursiveASTVisitor.h b/clang/include/clang/AST/RecursiveASTVisitor.h --- a/clang/include/clang/AST/RecursiveASTVisitor.h +++ b/clang/include/clang/AST/RecursiveASTVisitor.h @@ -2780,6 +2780,9 @@ return TraverseOMPExecutableDirective(S); } +DEF_TRAVERSE_STMT(OMPMetaDirective, + { TRY_TO(TraverseOMPExecutableDirective(S)); }) + DEF_TRAVERSE_STMT(OMPParallelDirective, { TRY_TO(TraverseOMPExecutableDirective(S)); }) diff --git a/clang/include/clang/AST/StmtOpenMP.h b/clang/include/clang/AST/StmtOpenMP.h --- a/clang/include/clang/AST/StmtOpenMP.h +++ b/clang/include/clang/AST/StmtOpenMP.h @@ -362,6 +362,44 @@ } }; +/// This represents '#pragma omp metadirective' directive. +/// +/// \code +/// #pragma omp metadirective when(user={condition(N>10)}: parallel for) +/// \endcode +/// In this example directive '#pragma omp metadirective' has clauses 'when' +/// with a dynamic user condition to check if a variable 'N > 10' +/// +class OMPMetaDirective final : public OMPExecutableDirective { + friend class ASTStmtReader; + friend class OMPExecutableDirective; + Stmt *IfStmt; + + OMPMetaDirective(SourceLocation StartLoc, SourceLocation EndLoc) + : OMPExecutableDirective(OMPMetaDirectiveClass, + llvm::omp::OMPD_metadirective, StartLoc, + EndLoc) {} + explicit OMPMetaDirective() + : OMPExecutableDirective(OMPMetaDirectiveClass, + llvm::omp::OMPD_metadirective, SourceLocation(), + SourceLocation()) {} + +public: + static OMPMetaDirective *Create(const ASTContext &C, SourceLocation StartLoc, + SourceLocation EndLoc, + ArrayRef Clauses, + Stmt *AssociatedStmt, Stmt *IfStmt); + static OMPMetaDirective *CreateEmpty(const ASTContext &C, unsigned NumClauses, + EmptyShell); + + void setIfStmt(Stmt *stmt) { IfStmt = stmt; } + Stmt *getIfStmt() const { return IfStmt; } + + static bool classof(const Stmt *T) { + return T->getStmtClass() == OMPMetaDirectiveClass; + } +}; + /// This represents '#pragma omp parallel' directive. /// /// \code diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -10488,6 +10488,9 @@ : Note<"jump bypasses OpenMP structured block">; def note_omp_exits_structured_block : Note<"jump exits scope of OpenMP structured block">; +def err_omp_misplaced_default_clause : Error< + "misplaced default clause! Only one default clause is allowed in " + "metadirective in the end">; } // end of OpenMP category let CategoryName = "Related Result Type Issue" in { diff --git a/clang/include/clang/Basic/StmtNodes.td b/clang/include/clang/Basic/StmtNodes.td --- a/clang/include/clang/Basic/StmtNodes.td +++ b/clang/include/clang/Basic/StmtNodes.td @@ -217,6 +217,7 @@ // OpenMP Directives. def OMPExecutableDirective : StmtNode; +def OMPMetaDirective : StmtNode; def OMPLoopDirective : StmtNode; def OMPParallelDirective : StmtNode; def OMPSimdDirective : StmtNode; diff --git a/clang/include/clang/Sema/Sema.h b/clang/include/clang/Sema/Sema.h --- a/clang/include/clang/Sema/Sema.h +++ b/clang/include/clang/Sema/Sema.h @@ -10300,6 +10300,11 @@ void ActOnOpenMPLoopInitialization(SourceLocation ForLoc, Stmt *Init); // OpenMP directives and clauses. + /// Called on well-formed '\#pragma omp metadirective' after parsing + /// of the associated statement. + StmtResult ActOnOpenMPMetaDirective(ArrayRef Clauses, + Stmt *AStmt, SourceLocation StartLoc, + SourceLocation EndLoc); /// Called on correct id-expression from the '#pragma omp /// threadprivate'. ExprResult ActOnOpenMPIdExpression(Scope *CurScope, CXXScopeSpec &ScopeSpec, @@ -10795,6 +10800,10 @@ SourceLocation StartLoc, SourceLocation LParenLoc, SourceLocation EndLoc); + /// Called on well-formed 'when' clause. + OMPClause *ActOnOpenMPWhenClause(OMPTraitInfo &TI, SourceLocation StartLoc, + SourceLocation LParenLoc, + SourceLocation EndLoc); /// Called on well-formed 'default' clause. OMPClause *ActOnOpenMPDefaultClause(llvm::omp::DefaultKind Kind, SourceLocation KindLoc, diff --git a/clang/include/clang/Serialization/ASTBitCodes.h b/clang/include/clang/Serialization/ASTBitCodes.h --- a/clang/include/clang/Serialization/ASTBitCodes.h +++ b/clang/include/clang/Serialization/ASTBitCodes.h @@ -1830,21 +1830,21 @@ /// A CXXBoolLiteralExpr record. EXPR_CXX_BOOL_LITERAL, - EXPR_CXX_NULL_PTR_LITERAL, // CXXNullPtrLiteralExpr - EXPR_CXX_TYPEID_EXPR, // CXXTypeidExpr (of expr). - EXPR_CXX_TYPEID_TYPE, // CXXTypeidExpr (of type). - EXPR_CXX_THIS, // CXXThisExpr - EXPR_CXX_THROW, // CXXThrowExpr - EXPR_CXX_DEFAULT_ARG, // CXXDefaultArgExpr - EXPR_CXX_DEFAULT_INIT, // CXXDefaultInitExpr - EXPR_CXX_BIND_TEMPORARY, // CXXBindTemporaryExpr + EXPR_CXX_NULL_PTR_LITERAL, // CXXNullPtrLiteralExpr + EXPR_CXX_TYPEID_EXPR, // CXXTypeidExpr (of expr). + EXPR_CXX_TYPEID_TYPE, // CXXTypeidExpr (of type). + EXPR_CXX_THIS, // CXXThisExpr + EXPR_CXX_THROW, // CXXThrowExpr + EXPR_CXX_DEFAULT_ARG, // CXXDefaultArgExpr + EXPR_CXX_DEFAULT_INIT, // CXXDefaultInitExpr + EXPR_CXX_BIND_TEMPORARY, // CXXBindTemporaryExpr EXPR_CXX_SCALAR_VALUE_INIT, // CXXScalarValueInitExpr EXPR_CXX_NEW, // CXXNewExpr EXPR_CXX_DELETE, // CXXDeleteExpr EXPR_CXX_PSEUDO_DESTRUCTOR, // CXXPseudoDestructorExpr - EXPR_EXPR_WITH_CLEANUPS, // ExprWithCleanups + EXPR_EXPR_WITH_CLEANUPS, // ExprWithCleanups EXPR_CXX_DEPENDENT_SCOPE_MEMBER, // CXXDependentScopeMemberExpr EXPR_CXX_DEPENDENT_SCOPE_DECL_REF, // DependentScopeDeclRefExpr @@ -1852,41 +1852,42 @@ EXPR_CXX_UNRESOLVED_MEMBER, // UnresolvedMemberExpr EXPR_CXX_UNRESOLVED_LOOKUP, // UnresolvedLookupExpr - EXPR_CXX_EXPRESSION_TRAIT, // ExpressionTraitExpr - EXPR_CXX_NOEXCEPT, // CXXNoexceptExpr + EXPR_CXX_EXPRESSION_TRAIT, // ExpressionTraitExpr + EXPR_CXX_NOEXCEPT, // CXXNoexceptExpr - EXPR_OPAQUE_VALUE, // OpaqueValueExpr - EXPR_BINARY_CONDITIONAL_OPERATOR, // BinaryConditionalOperator - EXPR_TYPE_TRAIT, // TypeTraitExpr - EXPR_ARRAY_TYPE_TRAIT, // ArrayTypeTraitIntExpr + EXPR_OPAQUE_VALUE, // OpaqueValueExpr + EXPR_BINARY_CONDITIONAL_OPERATOR, // BinaryConditionalOperator + EXPR_TYPE_TRAIT, // TypeTraitExpr + EXPR_ARRAY_TYPE_TRAIT, // ArrayTypeTraitIntExpr - EXPR_PACK_EXPANSION, // PackExpansionExpr - EXPR_SIZEOF_PACK, // SizeOfPackExpr - EXPR_SUBST_NON_TYPE_TEMPLATE_PARM, // SubstNonTypeTemplateParmExpr - EXPR_SUBST_NON_TYPE_TEMPLATE_PARM_PACK,// SubstNonTypeTemplateParmPackExpr - EXPR_FUNCTION_PARM_PACK, // FunctionParmPackExpr - EXPR_MATERIALIZE_TEMPORARY, // MaterializeTemporaryExpr - EXPR_CXX_FOLD, // CXXFoldExpr - EXPR_CONCEPT_SPECIALIZATION,// ConceptSpecializationExpr - EXPR_REQUIRES, // RequiresExpr + EXPR_PACK_EXPANSION, // PackExpansionExpr + EXPR_SIZEOF_PACK, // SizeOfPackExpr + EXPR_SUBST_NON_TYPE_TEMPLATE_PARM, // SubstNonTypeTemplateParmExpr + EXPR_SUBST_NON_TYPE_TEMPLATE_PARM_PACK, // SubstNonTypeTemplateParmPackExpr + EXPR_FUNCTION_PARM_PACK, // FunctionParmPackExpr + EXPR_MATERIALIZE_TEMPORARY, // MaterializeTemporaryExpr + EXPR_CXX_FOLD, // CXXFoldExpr + EXPR_CONCEPT_SPECIALIZATION, // ConceptSpecializationExpr + EXPR_REQUIRES, // RequiresExpr // CUDA - EXPR_CUDA_KERNEL_CALL, // CUDAKernelCallExpr + EXPR_CUDA_KERNEL_CALL, // CUDAKernelCallExpr // OpenCL - EXPR_ASTYPE, // AsTypeExpr + EXPR_ASTYPE, // AsTypeExpr // Microsoft - EXPR_CXX_PROPERTY_REF_EXPR, // MSPropertyRefExpr + EXPR_CXX_PROPERTY_REF_EXPR, // MSPropertyRefExpr EXPR_CXX_PROPERTY_SUBSCRIPT_EXPR, // MSPropertySubscriptExpr - EXPR_CXX_UUIDOF_EXPR, // CXXUuidofExpr (of expr). - EXPR_CXX_UUIDOF_TYPE, // CXXUuidofExpr (of type). - STMT_SEH_LEAVE, // SEHLeaveStmt - STMT_SEH_EXCEPT, // SEHExceptStmt - STMT_SEH_FINALLY, // SEHFinallyStmt - STMT_SEH_TRY, // SEHTryStmt + EXPR_CXX_UUIDOF_EXPR, // CXXUuidofExpr (of expr). + EXPR_CXX_UUIDOF_TYPE, // CXXUuidofExpr (of type). + STMT_SEH_LEAVE, // SEHLeaveStmt + STMT_SEH_EXCEPT, // SEHExceptStmt + STMT_SEH_FINALLY, // SEHFinallyStmt + STMT_SEH_TRY, // SEHTryStmt // OpenMP directives + STMT_OMP_META_DIRECTIVE, STMT_OMP_PARALLEL_DIRECTIVE, STMT_OMP_SIMD_DIRECTIVE, STMT_OMP_FOR_DIRECTIVE, @@ -1946,10 +1947,10 @@ EXPR_OMP_ITERATOR, // ARC - EXPR_OBJC_BRIDGED_CAST, // ObjCBridgedCastExpr + EXPR_OBJC_BRIDGED_CAST, // ObjCBridgedCastExpr - STMT_MS_DEPENDENT_EXISTS, // MSDependentExistsStmt - EXPR_LAMBDA, // LambdaExpr + STMT_MS_DEPENDENT_EXISTS, // MSDependentExistsStmt + EXPR_LAMBDA, // LambdaExpr STMT_COROUTINE_BODY, STMT_CORETURN, EXPR_COAWAIT, diff --git a/clang/lib/AST/OpenMPClause.cpp b/clang/lib/AST/OpenMPClause.cpp --- a/clang/lib/AST/OpenMPClause.cpp +++ b/clang/lib/AST/OpenMPClause.cpp @@ -156,6 +156,7 @@ case OMPC_exclusive: case OMPC_uses_allocators: case OMPC_affinity: + case OMPC_when: break; default: break; @@ -250,6 +251,7 @@ case OMPC_exclusive: case OMPC_uses_allocators: case OMPC_affinity: + case OMPC_when: break; default: break; diff --git a/clang/lib/AST/StmtOpenMP.cpp b/clang/lib/AST/StmtOpenMP.cpp --- a/clang/lib/AST/StmtOpenMP.cpp +++ b/clang/lib/AST/StmtOpenMP.cpp @@ -191,6 +191,25 @@ llvm::copy(A, getFinalsConditions().begin()); } +OMPMetaDirective *OMPMetaDirective::Create(const ASTContext &C, + SourceLocation StartLoc, + SourceLocation EndLoc, + ArrayRef Clauses, + Stmt *AssociatedStmt, Stmt *IfStmt) { + auto *Dir = createDirective( + C, Clauses, AssociatedStmt, /*NumChildren=*/1, StartLoc, EndLoc); + Dir->setIfStmt(IfStmt); + return Dir; +} + +OMPMetaDirective *OMPMetaDirective::CreateEmpty(const ASTContext &C, + unsigned NumClauses, + EmptyShell) { + return createEmptyDirective(C, NumClauses, + /*HasAssociatedStmt=*/true, + /*NumChildren=*/1); +} + OMPParallelDirective *OMPParallelDirective::Create( const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc, ArrayRef Clauses, Stmt *AssociatedStmt, Expr *TaskRedRef, diff --git a/clang/lib/AST/StmtPrinter.cpp b/clang/lib/AST/StmtPrinter.cpp --- a/clang/lib/AST/StmtPrinter.cpp +++ b/clang/lib/AST/StmtPrinter.cpp @@ -650,6 +650,11 @@ PrintStmt(S->getRawStmt()); } +void StmtPrinter::VisitOMPMetaDirective(OMPMetaDirective *Node) { + Indent() << "#pragma omp metadirective"; + PrintOMPExecutableDirective(Node); +} + void StmtPrinter::VisitOMPParallelDirective(OMPParallelDirective *Node) { Indent() << "#pragma omp parallel"; PrintOMPExecutableDirective(Node); diff --git a/clang/lib/AST/StmtProfile.cpp b/clang/lib/AST/StmtProfile.cpp --- a/clang/lib/AST/StmtProfile.cpp +++ b/clang/lib/AST/StmtProfile.cpp @@ -847,6 +847,10 @@ P.Visit(*I); } +void StmtProfiler::VisitOMPMetaDirective(const OMPMetaDirective *S) { + VisitOMPExecutableDirective(S); +} + void StmtProfiler::VisitOMPLoopDirective(const OMPLoopDirective *S) { VisitOMPExecutableDirective(S); } diff --git a/clang/lib/Basic/OpenMPKinds.cpp b/clang/lib/Basic/OpenMPKinds.cpp --- a/clang/lib/Basic/OpenMPKinds.cpp +++ b/clang/lib/Basic/OpenMPKinds.cpp @@ -180,6 +180,7 @@ case OMPC_exclusive: case OMPC_uses_allocators: case OMPC_affinity: + case OMPC_when: break; default: break; @@ -420,6 +421,7 @@ case OMPC_exclusive: case OMPC_uses_allocators: case OMPC_affinity: + case OMPC_when: break; default: break; @@ -578,6 +580,9 @@ OpenMPDirectiveKind DKind) { assert(unsigned(DKind) < llvm::omp::Directive_enumSize); switch (DKind) { + case OMPD_metadirective: + CaptureRegions.push_back(OMPD_metadirective); + break; case OMPD_parallel: case OMPD_parallel_for: case OMPD_parallel_for_simd: diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -6983,6 +6983,7 @@ case OMPD_parallel_master_taskloop_simd: case OMPD_requires: case OMPD_unknown: + case OMPD_metadirective: break; default: break; @@ -9499,6 +9500,7 @@ case OMPD_parallel_master_taskloop: case OMPD_parallel_master_taskloop_simd: case OMPD_requires: + case OMPD_metadirective: case OMPD_unknown: default: llvm_unreachable("Unexpected directive."); @@ -10202,6 +10204,12 @@ if (!S) return; + if (isa(S)) { + const auto &M = *cast(S); + scanForTargetRegionsFunctions(M.getIfStmt(), ParentName); + return; + } + // Codegen OMP target directives that offload compute to the device. bool RequiresDeviceCodegen = isa(S) && @@ -11010,6 +11018,7 @@ case OMPD_target_parallel_for: case OMPD_target_parallel_for_simd: case OMPD_requires: + case OMPD_metadirective: case OMPD_unknown: default: llvm_unreachable("Unexpected standalone target data directive."); diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp --- a/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp @@ -709,6 +709,7 @@ case OMPD_parallel_master_taskloop: case OMPD_parallel_master_taskloop_simd: case OMPD_requires: + case OMPD_metadirective: case OMPD_unknown: default: llvm_unreachable("Unexpected directive."); @@ -790,6 +791,7 @@ case OMPD_parallel_master_taskloop: case OMPD_parallel_master_taskloop_simd: case OMPD_requires: + case OMPD_metadirective: case OMPD_unknown: default: break; @@ -964,6 +966,7 @@ case OMPD_parallel_master_taskloop: case OMPD_parallel_master_taskloop_simd: case OMPD_requires: + case OMPD_metadirective: case OMPD_unknown: default: llvm_unreachable("Unexpected directive."); @@ -1051,6 +1054,7 @@ case OMPD_parallel_master_taskloop: case OMPD_parallel_master_taskloop_simd: case OMPD_requires: + case OMPD_metadirective: case OMPD_unknown: default: break; diff --git a/clang/lib/CodeGen/CGStmt.cpp b/clang/lib/CodeGen/CGStmt.cpp --- a/clang/lib/CodeGen/CGStmt.cpp +++ b/clang/lib/CodeGen/CGStmt.cpp @@ -194,6 +194,9 @@ case Stmt::SEHTryStmtClass: EmitSEHTryStmt(cast(*S)); break; + case Stmt::OMPMetaDirectiveClass: + EmitOMPMetaDirective(cast(*S)); + break; case Stmt::OMPParallelDirectiveClass: EmitOMPParallelDirective(cast(*S)); break; diff --git a/clang/lib/CodeGen/CGStmtOpenMP.cpp b/clang/lib/CodeGen/CGStmtOpenMP.cpp --- a/clang/lib/CodeGen/CGStmtOpenMP.cpp +++ b/clang/lib/CodeGen/CGStmtOpenMP.cpp @@ -5399,6 +5399,7 @@ case OMPC_exclusive: case OMPC_uses_allocators: case OMPC_affinity: + case OMPC_when: default: llvm_unreachable("Clause is not allowed in 'omp atomic'."); } @@ -6587,6 +6588,10 @@ CGM.getOpenMPRuntime().emitMasterRegion(*this, CodeGen, S.getBeginLoc()); } +void CodeGenFunction::EmitOMPMetaDirective(const OMPMetaDirective &S) { + EmitStmt(S.getIfStmt()); +} + void CodeGenFunction::EmitOMPParallelMasterTaskLoopDirective( const OMPParallelMasterTaskLoopDirective &S) { auto &&CodeGen = [this, &S](CodeGenFunction &CGF, PrePostActionTy &Action) { diff --git a/clang/lib/CodeGen/CodeGenFunction.h b/clang/lib/CodeGen/CodeGenFunction.h --- a/clang/lib/CodeGen/CodeGenFunction.h +++ b/clang/lib/CodeGen/CodeGenFunction.h @@ -3387,6 +3387,7 @@ const RegionCodeGenTy &BodyGen, OMPTargetDataInfo &InputInfo); + void EmitOMPMetaDirective(const OMPMetaDirective &S); void EmitOMPParallelDirective(const OMPParallelDirective &S); void EmitOMPSimdDirective(const OMPSimdDirective &S); void EmitOMPForDirective(const OMPForDirective &S); diff --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp --- a/clang/lib/Parse/ParseOpenMP.cpp +++ b/clang/lib/Parse/ParseOpenMP.cpp @@ -2085,6 +2085,7 @@ case OMPD_target_teams_distribute_parallel_for: case OMPD_target_teams_distribute_parallel_for_simd: case OMPD_target_teams_distribute_simd: + case OMPD_metadirective: Diag(Tok, diag::err_omp_unexpected_directive) << 1 << getOpenMPDirectiveName(DKind); break; @@ -2154,9 +2155,177 @@ // Name of critical directive. DeclarationNameInfo DirName; StmtResult Directive = StmtError(); + bool HasAssociatedStatement = true; switch (DKind) { + case OMPD_metadirective: { + ConsumeToken(); + SmallVector VMIs; + OpenMPDirectiveKind DirKind = OMPD_unknown; + StmtResult AssociatedStmt; + TentativeParsingAction TPA(*this); + ASTContext &ASTContext = Actions.getASTContext(); + + BalancedDelimiterTracker T(*this, tok::l_paren, + tok::annot_pragma_openmp_end); + while (Tok.isNot(tok::annot_pragma_openmp_end)) { + OpenMPClauseKind CKind = Tok.isAnnotation() + ? OMPC_unknown + : getOpenMPClauseKind(PP.getSpelling(Tok)); + SourceLocation Loc = ConsumeToken(); + + // Parse '('. + if (T.expectAndConsume(diag::err_expected_lparen_after, + getOpenMPClauseName(CKind).data())) + return Directive; + + OMPTraitInfo &TI = Actions.getASTContext().getNewOMPTraitInfo(); + if (CKind == OMPC_when) { + // parse and get condition expression to pass to the When clause + parseOMPContextSelectors(Loc, TI); + + // Parse ':' + if (Tok.is(tok::colon)) + ConsumeAnyToken(); + else { + Diag(Tok, diag::warn_pragma_expected_colon) << "when clause"; + return Directive; + } + } + // Skip Directive + int paren = 0; + while (Tok.isNot(tok::r_paren) || paren != 0) { + if (Tok.is(tok::l_paren)) + paren++; + if (Tok.is(tok::r_paren)) + paren--; + ConsumeAnyToken(); + } + // Parse ')' + if (Tok.is(tok::r_paren)) + T.consumeClose(); + + VariantMatchInfo VMI; + TI.getAsVariantMatchInfo(ASTContext, VMI); + + VMIs.push_back(VMI); + } + + TPA.Revert(); + TargetOMPContext OMPCtx(ASTContext, nullptr, nullptr); + int BestIdx = getBestWhenMatchForContext(VMIs, OMPCtx); + + int idx = 0; + while (Tok.isNot(tok::annot_pragma_openmp_end)) { + if (idx++ != BestIdx) { // Ignore this clause + ConsumeToken(); // Consume clause name + T.consumeOpen(); // Consume '(' + int paren = 0; + // Skip everything inside the clause + while (Tok.isNot(tok::r_paren) || paren != 0) { + if (Tok.is(tok::l_paren)) + paren++; + if (Tok.is(tok::r_paren)) + paren--; + ConsumeAnyToken(); + } + // Parse ')' + if (Tok.is(tok::r_paren)) + T.consumeClose(); + continue; + } + + OpenMPClauseKind CKind = Tok.isAnnotation() + ? OMPC_unknown + : getOpenMPClauseKind(PP.getSpelling(Tok)); + SourceLocation Loc = ConsumeToken(); + + // Parse '('. + T.consumeOpen(); + + // Skip ContextSelectors for when clause + if (CKind == OMPC_when) { + OMPTraitInfo &TI = Actions.getASTContext().getNewOMPTraitInfo(); + // parse and skip ContextSelectors + parseOMPContextSelectors(Loc, TI); + + // Parse ':' + ConsumeAnyToken(); + } + + // Parse Directive + if (Tok.isNot(tok::r_paren)) { + DirKind = parseOpenMPDirectiveKind(*this); + ConsumeToken(); + if (DirKind != OMPD_unknown) { + ParseScope OMPDirectiveScope(this, ScopeFlags); + Actions.StartOpenMPDSABlock(DirKind, DirName, Actions.getCurScope(), + Loc); + int paren = 0; + while (Tok.isNot(tok::r_paren) || paren != 0) { + if (Tok.is(tok::l_paren)) + paren++; + if (Tok.is(tok::r_paren)) + paren--; + OpenMPClauseKind CKind = + Tok.isAnnotation() ? OMPC_unknown + : getOpenMPClauseKind(PP.getSpelling(Tok)); + Actions.StartOpenMPClause(CKind); + OMPClause *DClause = ParseOpenMPClause( + DirKind, CKind, !FirstClauses[(unsigned)CKind].getInt()); + FirstClauses[(unsigned)CKind].setInt(true); + if (DClause) { + FirstClauses[(unsigned)CKind].setPointer(DClause); + Clauses.push_back(DClause); + } + // Skip ',' if any. + if (Tok.is(tok::comma)) + ConsumeToken(); + Actions.EndOpenMPClause(); + } + while (Tok.isNot(tok::annot_pragma_openmp_end)) + ConsumeAnyToken(); + // End location of the directive. + EndLoc = Tok.getLocation(); + // Consume final annot_pragma_openmp_end. + ConsumeAnnotationToken(); + + Actions.ActOnOpenMPRegionStart(DirKind, getCurScope()); + ParsingOpenMPDirectiveRAII NormalScope(*this, /*Value=*/false); + ParseScope InnerStmtScope(this, Scope::DeclScope, + getLangOpts().C99 || + getLangOpts().CPlusPlus, + Tok.is(tok::l_brace)); + StmtResult AStmt = ParseStatement(); + InnerStmtScope.Exit(); + AssociatedStmt = (Sema::CompoundScopeRAII(Actions), AStmt); + AssociatedStmt = + Actions.ActOnOpenMPRegionEnd(AssociatedStmt, Clauses); + Directive = Actions.ActOnOpenMPExecutableDirective( + DirKind, DirName, CancelRegion, Clauses, AssociatedStmt.get(), + Loc, EndLoc); + // Exit scope. + Actions.EndOpenMPDSABlock(Directive.get()); + OMPDirectiveScope.Exit(); + } else { + // Unknown Directive encountered. Skip directive. + Diag(Tok, diag::err_omp_unknown_directive); + while (Tok.isNot(tok::annot_pragma_openmp_end)) + ConsumeAnyToken(); + ConsumeAnnotationToken(); + } + } else { + // Empty Directive. Skip in OpenMP 5.0. + // Generate nothing directive from OpenMP 5.1. + while (Tok.isNot(tok::annot_pragma_openmp_end)) + ConsumeAnyToken(); + ConsumeAnnotationToken(); + } + break; + } + break; + } case OMPD_threadprivate: { // FIXME: Should this be permitted in C++? if ((StmtCtx & ParsedStmtContext::AllowDeclarationsInC) == @@ -2759,6 +2928,7 @@ case OMPC_threadprivate: case OMPC_uniform: case OMPC_match: + case OMPC_when: if (!WrongDirective) Diag(Tok, diag::err_omp_unexpected_clause) << getOpenMPClauseName(CKind) << getOpenMPDirectiveName(DKind); diff --git a/clang/lib/Sema/SemaExceptionSpec.cpp b/clang/lib/Sema/SemaExceptionSpec.cpp --- a/clang/lib/Sema/SemaExceptionSpec.cpp +++ b/clang/lib/Sema/SemaExceptionSpec.cpp @@ -1484,6 +1484,7 @@ case Stmt::OMPTeamsDistributeParallelForDirectiveClass: case Stmt::OMPTeamsDistributeParallelForSimdDirectiveClass: case Stmt::OMPTeamsDistributeSimdDirectiveClass: + case Stmt::OMPMetaDirectiveClass: case Stmt::ReturnStmtClass: case Stmt::SEHExceptStmtClass: case Stmt::SEHFinallyStmtClass: diff --git a/clang/lib/Sema/SemaOpenMP.cpp b/clang/lib/Sema/SemaOpenMP.cpp --- a/clang/lib/Sema/SemaOpenMP.cpp +++ b/clang/lib/Sema/SemaOpenMP.cpp @@ -3751,6 +3751,7 @@ void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) { switch (DKind) { + case OMPD_metadirective: case OMPD_parallel: case OMPD_parallel_for: case OMPD_parallel_for_simd: @@ -5436,6 +5437,7 @@ if (LangOpts.OpenMP >= 50) AllowedNameModifiers.push_back(OMPD_simd); break; + case OMPD_metadirective: case OMPD_declare_target: case OMPD_end_declare_target: case OMPD_threadprivate: @@ -5554,6 +5556,7 @@ case OMPC_atomic_default_mem_order: case OMPC_device_type: case OMPC_match: + case OMPC_when: default: llvm_unreachable("Unexpected clause"); } @@ -11859,6 +11862,7 @@ case OMPC_exclusive: case OMPC_uses_allocators: case OMPC_affinity: + case OMPC_when: default: llvm_unreachable("Clause is not allowed."); } @@ -12012,6 +12016,7 @@ case OMPD_atomic: case OMPD_teams_distribute: case OMPD_requires: + case OMPD_metadirective: llvm_unreachable("Unexpected OpenMP directive with if-clause"); case OMPD_unknown: default: @@ -12091,6 +12096,7 @@ case OMPD_teams_distribute: case OMPD_teams_distribute_simd: case OMPD_requires: + case OMPD_metadirective: llvm_unreachable("Unexpected OpenMP directive with num_threads-clause"); case OMPD_unknown: default: @@ -12168,6 +12174,7 @@ case OMPD_atomic: case OMPD_distribute_simd: case OMPD_requires: + case OMPD_metadirective: llvm_unreachable("Unexpected OpenMP directive with num_teams-clause"); case OMPD_unknown: default: @@ -12245,6 +12252,7 @@ case OMPD_atomic: case OMPD_distribute_simd: case OMPD_requires: + case OMPD_metadirective: llvm_unreachable("Unexpected OpenMP directive with thread_limit-clause"); case OMPD_unknown: default: @@ -12322,6 +12330,7 @@ case OMPD_distribute_simd: case OMPD_target_teams: case OMPD_requires: + case OMPD_metadirective: llvm_unreachable("Unexpected OpenMP directive with schedule clause"); case OMPD_unknown: default: @@ -12399,6 +12408,7 @@ case OMPD_atomic: case OMPD_target_teams: case OMPD_requires: + case OMPD_metadirective: llvm_unreachable("Unexpected OpenMP directive with schedule clause"); case OMPD_unknown: default: @@ -12476,6 +12486,7 @@ case OMPD_atomic: case OMPD_distribute_simd: case OMPD_requires: + case OMPD_metadirective: llvm_unreachable("Unexpected OpenMP directive with num_teams-clause"); case OMPD_unknown: default: @@ -12555,12 +12566,22 @@ case OMPD_atomic: case OMPD_distribute_simd: case OMPD_requires: + case OMPD_metadirective: llvm_unreachable("Unexpected OpenMP directive with grainsize-clause"); case OMPD_unknown: default: llvm_unreachable("Unknown OpenMP directive"); } break; + case OMPC_when: + if (DKind == OMPD_metadirective) { + CaptureRegion = OMPD_metadirective; + } else if (DKind == OMPD_unknown) { + llvm_unreachable("Unknown OpenMP directive"); + } else { + llvm_unreachable("Unexpected OpenMP directive with when clause"); + } + break; case OMPC_firstprivate: case OMPC_lastprivate: case OMPC_reduction: @@ -13068,6 +13089,7 @@ case OMPC_exclusive: case OMPC_uses_allocators: case OMPC_affinity: + case OMPC_when: default: llvm_unreachable("Clause is not allowed."); } @@ -13307,6 +13329,7 @@ case OMPC_exclusive: case OMPC_uses_allocators: case OMPC_affinity: + case OMPC_when: default: llvm_unreachable("Clause is not allowed."); } @@ -13545,6 +13568,7 @@ case OMPC_exclusive: case OMPC_uses_allocators: case OMPC_affinity: + case OMPC_when: default: llvm_unreachable("Clause is not allowed."); } @@ -13823,6 +13847,7 @@ case OMPC_destroy: case OMPC_detach: case OMPC_uses_allocators: + case OMPC_when: default: llvm_unreachable("Clause is not allowed."); } diff --git a/clang/lib/Sema/TreeTransform.h b/clang/lib/Sema/TreeTransform.h --- a/clang/lib/Sema/TreeTransform.h +++ b/clang/lib/Sema/TreeTransform.h @@ -8379,6 +8379,17 @@ template StmtResult +TreeTransform::TransformOMPMetaDirective(OMPMetaDirective *D) { + DeclarationNameInfo DirName; + getDerived().getSema().StartOpenMPDSABlock(OMPD_metadirective, DirName, + nullptr, D->getBeginLoc()); + StmtResult Res = getDerived().TransformOMPExecutableDirective(D); + getDerived().getSema().EndOpenMPDSABlock(Res.get()); + return Res; +} + +template +StmtResult TreeTransform::TransformOMPParallelDirective(OMPParallelDirective *D) { DeclarationNameInfo DirName; getDerived().getSema().StartOpenMPDSABlock(OMPD_parallel, DirName, nullptr, diff --git a/clang/lib/Serialization/ASTReaderStmt.cpp b/clang/lib/Serialization/ASTReaderStmt.cpp --- a/clang/lib/Serialization/ASTReaderStmt.cpp +++ b/clang/lib/Serialization/ASTReaderStmt.cpp @@ -2286,6 +2286,13 @@ VisitOMPExecutableDirective(D); } +void ASTStmtReader::VisitOMPMetaDirective(OMPMetaDirective *D) { + VisitStmt(D); + // The NumClauses field was read in ReadStmtFromStream. + Record.skipInts(1); + VisitOMPExecutableDirective(D); +} + void ASTStmtReader::VisitOMPParallelDirective(OMPParallelDirective *D) { VisitStmt(D); VisitOMPExecutableDirective(D); @@ -3130,6 +3137,11 @@ nullptr); break; + case STMT_OMP_META_DIRECTIVE: + S = OMPMetaDirective::CreateEmpty( + Context, Record[ASTStmtReader::NumStmtFields], Empty); + break; + case STMT_OMP_PARALLEL_DIRECTIVE: S = OMPParallelDirective::CreateEmpty(Context, diff --git a/clang/lib/Serialization/ASTWriterStmt.cpp b/clang/lib/Serialization/ASTWriterStmt.cpp --- a/clang/lib/Serialization/ASTWriterStmt.cpp +++ b/clang/lib/Serialization/ASTWriterStmt.cpp @@ -2176,6 +2176,13 @@ Record.AddSourceLocation(E->getEndLoc()); } +void ASTStmtWriter::VisitOMPMetaDirective(OMPMetaDirective *D) { + VisitStmt(D); + Record.push_back(D->getNumClauses()); + VisitOMPExecutableDirective(D); + Code = serialization::STMT_OMP_META_DIRECTIVE; +} + void ASTStmtWriter::VisitOMPLoopDirective(OMPLoopDirective *D) { VisitStmt(D); Record.writeUInt32(D->getCollapsedNumber()); diff --git a/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp b/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp --- a/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp +++ b/clang/lib/StaticAnalyzer/Core/ExprEngine.cpp @@ -1292,6 +1292,7 @@ case Stmt::OMPTargetTeamsDistributeParallelForDirectiveClass: case Stmt::OMPTargetTeamsDistributeParallelForSimdDirectiveClass: case Stmt::OMPTargetTeamsDistributeSimdDirectiveClass: + case Stmt::OMPMetaDirectiveClass: case Stmt::CapturedStmtClass: { const ExplodedNode *node = Bldr.generateSink(S, Pred, Pred->getState()); Engine.addAbortedBlock(node, currBldrCtx->getBlock()); diff --git a/clang/test/OpenMP/metadirective_ast_print.cpp b/clang/test/OpenMP/metadirective_ast_print.cpp new file mode 100644 --- /dev/null +++ b/clang/test/OpenMP/metadirective_ast_print.cpp @@ -0,0 +1,26 @@ +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -ast-print %s -o - | FileCheck %s +// expected-no-diagnostics +#define N 10000 +int main() { + int v1[N], v2[N], v3[N]; + for (int i = 0; i < N; i++) { + v1[i] = (i + 1); + v2[i] = -(i + 1); + } +#pragma omp metadirective \ + when(device = {arch("nvptx64")} \ + : target teams distribute parallel for map(v1 [0:N], v2 [0:N], v3 [0:N])) default() + for (int i = 0; i < N; i++) { + v3[i] = v1[i] * v2[i]; + } + return 0; +} +// CHECK: int v1[10000], v2[10000], v3[10000]; +// CHECK: for (int i = 0; i < 10000; i++) { +// CHECK: v1[i] = (i + 1); +// CHECK: v2[i] = -(i + 1); +// CHECK: } +// CHECK: #pragma omp target teams distribute parallel for map(tofrom: v1[0:10000],v2[0:10000],v3[0:10000]) +// CHECK: for (int i = 0; i < 10000; i++) { +// CHECK: v3[i] = v1[i] * v2[i]; +// CHECK: } diff --git a/clang/test/OpenMP/metadirective_codegen.cpp b/clang/test/OpenMP/metadirective_codegen.cpp new file mode 100644 --- /dev/null +++ b/clang/test/OpenMP/metadirective_codegen.cpp @@ -0,0 +1,58 @@ +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s +// expected-no-diagnostics + +#define N 10000 +int func() { + int v1[N], v2[N], v3[N]; + for (int i = 0; i < N; i++) { + v1[i] = (i + 1); + v2[i] = -(i + 1); + } + +#pragma omp metadirective \ + when(device = {arch("nvptx64")} \ + : target teams distribute parallel for map(v1 [0:N], v2 [0:N], v3 [0:N])) default() + for (int i = 0; i < N; i++) { + v3[i] = v1[i] * v2[i]; + } + + return 0; +} +// CHECK: %arrayidx4 = getelementptr inbounds [10000 x i32], [10000 x i32]* %v3, i64 0, i64 0 +// CHECK: %arrayidx5 = getelementptr inbounds [10000 x i32], [10000 x i32]* %v1, i64 0, i64 0 +// CHECK: %arrayidx6 = getelementptr inbounds [10000 x i32], [10000 x i32]* %v2, i64 0, i64 0 +// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_baseptrs, i32 0, i32 0 +// CHECK: %{{[0-9]*}} = bitcast i8** %{{[0-9]*}} to [10000 x i32]** +// CHECK: store [10000 x i32]* %v3, [10000 x i32]** %{{[0-9]*}}, align 8 +// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_ptrs, i32 0, i32 0 +// CHECK: %{{[0-9]*}} = bitcast i8** %{{[0-9]*}} to i32** +// CHECK: store i32* %arrayidx4, i32** %{{[0-9]*}}, align 8 +// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_mappers, i64 0, i64 0 +// CHECK: store i8* null, i8** %{{[0-9]*}}, align 8 +// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_baseptrs, i32 0, i32 1 +// CHECK: %{{[0-9]*}} = bitcast i8** %{{[0-9]*}} to [10000 x i32]** +// CHECK: store [10000 x i32]* %v1, [10000 x i32]** %{{[0-9]*}}, align 8 +// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_ptrs, i32 0, i32 1 +// CHECK: %{{[0-9]*}} = bitcast i8** %{{[0-9]*}} to i32** +// CHECK: store i32* %arrayidx5, i32** %{{[0-9]*}}, align 8 +// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_mappers, i64 0, i64 1 +// CHECK: store i8* null, i8** %{{[0-9]*}}, align 8 +// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_baseptrs, i32 0, i32 2 +// CHECK: %{{[0-9]*}} = bitcast i8** %{{[0-9]*}} to [10000 x i32]** +// CHECK: store [10000 x i32]* %v2, [10000 x i32]** %{{[0-9]*}}, align 8 +// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_ptrs, i32 0, i32 2 +// CHECK: %{{[0-9]*}} = bitcast i8** %{{[0-9]*}} to i32** +// CHECK: store i32* %arrayidx6, i32** %{{[0-9]*}}, align 8 +// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_mappers, i64 0, i64 2 +// CHECK: store i8* null, i8** %{{[0-9]*}}, align 8 +// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_baseptrs, i32 0, i32 0 +// CHECK: %{{[0-9]*}} = getelementptr inbounds [3 x i8*], [3 x i8*]* %.offload_ptrs, i32 0, i32 0 +// CHECK: call void @__kmpc_push_target_tripcount(%struct.ident_t* @3, i64 -1, i64 10000) +// CHECK: %{{[0-9]*}} = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @3, i64 -1, i8* [[OMP_OUTLINED:@.+]].region_id, i32 3, i8** %{{[0-9]*}}, i8** %{{[0-9]*}}, i64* getelementptr inbounds ([3 x i64], [3 x i64]* @.offload_sizes, i32 0, i32 0), i64* getelementptr inbounds ([3 x i64], [3 x i64]* @.offload_maptypes, i32 0, i32 0), i8** null, i8** null, i32 0, i32 0) +// CHECK: %{{[0-9]*}} = icmp ne i32 %{{[0-9]*}}, 0 +// CHECK: br i1 %{{[0-9]*}}, label %omp_offload.failed, label %omp_offload.cont +// CHECK: omp_offload.failed: +// CHECK: call void [[OMP_OUTLINED:@.+]] +// CHECK: br label %omp_offload.cont +// CHECK: omp_offload.cont: +// CHECK: ret i32 0 diff --git a/clang/test/OpenMP/metadirective_construct.cpp b/clang/test/OpenMP/metadirective_construct.cpp new file mode 100644 --- /dev/null +++ b/clang/test/OpenMP/metadirective_construct.cpp @@ -0,0 +1,106 @@ +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s +// expected-no-diagnostics + +#define N 1000 + +#pragma omp declare target +void func1() { + // Test to check the case where metadirective is called within a target construct. + // Here the directive for 'distribute parallel for' will be generated +#pragma omp metadirective \ + when(construct = {"target"} \ + : distribute parallel for) default() + for (int i = 0; i < N; i++) + ; +} +#pragma omp end declare target + +void func() { +#pragma omp target teams + func1(); +} + +// CHECK: define void @_Z5func1v() #0 { +// CHECK: entry: +// CHECK: %.omp.iv = alloca i32, align 4 +// CHECK: %tmp = alloca i32, align 4 +// CHECK: %.omp.comb.lb = alloca i32, align 4 +// CHECK: %.omp.comb.ub = alloca i32, align 4 +// CHECK: %.omp.stride = alloca i32, align 4 +// CHECK: %.omp.is_last = alloca i32, align 4 +// CHECK: %i = alloca i32, align 4 +// CHECK: %{{[0-9]*}} = call i32 @__kmpc_global_thread_num(%struct.ident_t* @2) +// CHECK: store i32 0, i32* %.omp.comb.lb, align 4 +// CHECK: store i32 999, i32* %.omp.comb.ub, align 4 +// CHECK: store i32 1, i32* %.omp.stride, align 4 +// CHECK: store i32 0, i32* %.omp.is_last, align 4 +// CHECK: call void @__kmpc_for_static_init_4(%struct.ident_t* @1, i32 %{{[0-9]*}}, i32 92, i32* %.omp.is_last, i32* %.omp.comb.lb, i32* %.omp.comb.ub, i32* %.omp.stride, i32 1, i32 1) +// CHECK: %{{[0-9]*}} = load i32, i32* %.omp.comb.ub, align 4 +// CHECK: %cmp = icmp sgt i32 %{{[0-9]*}}, 999 +// CHECK: br i1 %cmp, label %cond.true, label %cond.false +// CHECK: cond.true: +// CHECK: br label %cond.end +// CHECK: cond.false: +// CHECK: %{{[0-9]*}} = load i32, i32* %.omp.comb.ub, align 4 +// CHECK: br label %cond.end +// CHECK: cond.end: +// CHECK: %cond = phi i32 [ 999, %cond.true ], [ %{{[0-9]*}}, %cond.false ] +// CHECK: store i32 %cond, i32* %.omp.comb.ub, align 4 +// CHECK: %{{[0-9]*}} = load i32, i32* %.omp.comb.lb, align 4 +// CHECK: store i32 %{{[0-9]*}}, i32* %.omp.iv, align 4 +// CHECK: br label %omp.inner.for.cond +// CHECK: omp.inner.for.cond: +// CHECK: %{{[0-9]*}} = load i32, i32* %.omp.iv, align 4 +// CHECK: %{{[0-9]*}} = load i32, i32* %.omp.comb.ub, align 4 +// CHECK: %cmp1 = icmp sle i32 %{{[0-9]*}}, %{{[0-9]*}} +// CHECK: br i1 %cmp1, label %omp.inner.for.body, label %omp.inner.for.end +// CHECK: omp.inner.for.body: +// CHECK: %{{[0-9]*}} = load i32, i32* %.omp.comb.lb, align 4 +// CHECK: %{{[0-9]*}} = zext i32 %{{[0-9]*}} to i64 +// CHECK: %{{[0-9]*}} = load i32, i32* %.omp.comb.ub, align 4 +// CHECK: %{{[0-9]*}} = zext i32 %{{[0-9]*}} to i64 +// CHECK: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* @2, i32 2, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, i64, i64)* [[OMP_OUTLINED:@.+]] to void (i32*, i32*, ...)*), i64 %{{[0-9]*}}, i64 %{{[0-9]*}}) +// CHECK: br label %omp.inner.for.inc +// CHECK: omp.inner.for.inc: +// CHECK: %{{[0-9]*}} = load i32, i32* %.omp.iv, align 4 +// CHECK: %{{[0-9]*}} = load i32, i32* %.omp.stride, align 4 +// CHECK: %add = add nsw i32 %{{[0-9]*}}, %{{[0-9]*}} +// CHECK: store i32 %add, i32* %.omp.iv, align 4 +// CHECK: br label %omp.inner.for.cond +// CHECK: omp.inner.for.end: +// CHECK: br label %omp.loop.exit +// CHECK: omp.loop.exit: +// CHECK: call void @__kmpc_for_static_fini(%struct.ident_t* @1, i32 %{{[0-9]*}}) +// CHECK: ret void +// CHECK: } + +// Declaration of funtion handling distribute parallel for +// CHECK: define internal void [[OMP_OUTLINED:@.+]](i32* noalias %.global_tid., i32* noalias %.bound_tid., i64 %.previous.lb., i64 %.previous.ub.) + +// CHECK: define void @_Z4funcv() #0 { +// CHECK: entry: +// CHECK: %{{[0-9]*}} = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @2, i64 -1, i8* [[OMP_OFFLOADING:@.+]].region_id +// CHECK: %{{[0-9]*}} = icmp ne i32 %{{[0-9]*}}, 0 +// CHECK: br i1 %{{[0-9]*}}, label %omp_offload.failed, label %omp_offload.cont +// CHECK: omp_offload.failed: +// CHECK: call void [[OMP_OFFLOADING:@.+]]() #1 +// CHECK: br label %omp_offload.cont +// CHECK: omp_offload.cont: +// CHECK: ret void +// CHECK: } + +// CHECK: define internal void [[OMP_OFFLOADING:@.+]]() #2 { +// CHECK: entry: +// CHECK: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @2, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* @.omp_outlined..1 to void (i32*, i32*, ...)*)) +// CHECK: ret void +// CHECK: } + +// CHECK: define internal void @.omp_outlined..1(i32* noalias %.global_tid., i32* noalias %.bound_tid.) #2 { +// CHECK: entry: +// CHECK: %.global_tid..addr = alloca i32*, align 8 +// CHECK: %.bound_tid..addr = alloca i32*, align 8 +// CHECK: store i32* %.global_tid., i32** %.global_tid..addr, align 8 +// CHECK: store i32* %.bound_tid., i32** %.bound_tid..addr, align 8 +// CHECK: call void @_Z5func1v() +// CHECK: ret void +// CHECK: } diff --git a/clang/test/OpenMP/metadirective_empty.cpp b/clang/test/OpenMP/metadirective_empty.cpp new file mode 100644 --- /dev/null +++ b/clang/test/OpenMP/metadirective_empty.cpp @@ -0,0 +1,31 @@ +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s +// expected-no-diagnostics +#define N 1000 +void func() { + // Test where a valid when clause contains empty directive. + // The directive will be ignored and code for a serial for loop will be generated. +#pragma omp metadirective when(implementation = {vendor(llvm)} \ + :) default(parallel for) + for (int i = 0; i < N; i++) + ; +} + +// CHECK: define void @_Z4funcv() #0 { +// CHECK: entry: +// CHECK: %i = alloca i32, align 4 +// CHECK: store i32 0, i32* %i, align 4 +// CHECK: br label %for.cond +// CHECK: for.cond: +// CHECK: %0 = load i32, i32* %i, align 4 +// CHECK: %cmp = icmp slt i32 %0, 1000 +// CHECK: br i1 %cmp, label %for.body, label %for.end +// CHECK: for.body: +// CHECK: br label %for.inc +// CHECK: for.inc: +// CHECK: %1 = load i32, i32* %i, align 4 +// CHECK: %inc = add nsw i32 %1, 1 +// CHECK: store i32 %inc, i32* %i, align 4 +// CHECK: br label %for.cond, !llvm.loop !2 +// CHECK: for.end: +// CHECK: ret void +// CHECK: } diff --git a/clang/test/OpenMP/metadirective_implementation.cpp b/clang/test/OpenMP/metadirective_implementation.cpp new file mode 100644 --- /dev/null +++ b/clang/test/OpenMP/metadirective_implementation.cpp @@ -0,0 +1,36 @@ +// RUN: %clang_cc1 -verify -fopenmp -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -o - | FileCheck %s +// expected-no-diagnostics +void func() { +#pragma omp target + // Test to check implementation and device. + // Since we are building with llvm for nvptx64 device, 512 teams with thread_limit of 32 are generated. +#pragma omp metadirective \ + when(implementation = {vendor(llvm)}, device = {arch(nvptx64)} \ + : teams num_teams(512) thread_limit(32)) \ + when(implementation = {vendor(amd)}, device = {arch(amdgcn)} \ + : teams num_teams(512) thread_limit(64)) default(teams) +#pragma omp distribute parallel for + for (int i = 0; i < 10000; i++) + ; +} + +// CHECK: define void @_Z4funcv() #0 { +// CHECK: entry: +// CHECK: %tmp = alloca i32, align 4 +// CHECK: call void @__kmpc_push_target_tripcount +// CHECK: %0 = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @3, i64 -1, i8* [[OMP_OUTLINED:@.+]].region_id, +// CHECK: %1 = icmp ne i32 %0, 0 +// CHECK: br i1 %1, label %omp_offload.failed, label %omp_offload.cont +// CHECK: omp_offload.failed: ; preds = %entry +// CHECK: call void [[OMP_OUTLINED:@.+]] #2 +// CHECK: br label %omp_offload.cont +// CHECK: omp_offload.cont: ; preds = %omp_offload.failed, %entry +// CHECK: ret void +// CHECK: } +// CHECK: define internal void [[OMP_OUTLINED:@.+]] #1 { +// CHECK: entry: +// CHECK: %0 = call i32 @__kmpc_global_thread_num(%struct.ident_t* @3) +// CHECK: call void @__kmpc_push_num_teams(%struct.ident_t* @3, i32 %0, i32 512, i32 32) +// CHECK: call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_teams(%struct.ident_t* @3, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OMP_OUTLINED:@.+]] to void (i32*, i32*, ...)*)) +// CHECK: ret void +// CHECK: } diff --git a/clang/tools/libclang/CIndex.cpp b/clang/tools/libclang/CIndex.cpp --- a/clang/tools/libclang/CIndex.cpp +++ b/clang/tools/libclang/CIndex.cpp @@ -5525,6 +5525,8 @@ return cxstring::createRef("CXXAccessSpecifier"); case CXCursor_ModuleImportDecl: return cxstring::createRef("ModuleImport"); + case CXCursor_OMPMetaDirective: + return cxstring::createRef("OMPMetaDirective"); case CXCursor_OMPParallelDirective: return cxstring::createRef("OMPParallelDirective"); case CXCursor_OMPSimdDirective: diff --git a/clang/tools/libclang/CXCursor.cpp b/clang/tools/libclang/CXCursor.cpp --- a/clang/tools/libclang/CXCursor.cpp +++ b/clang/tools/libclang/CXCursor.cpp @@ -639,6 +639,9 @@ case Stmt::MSDependentExistsStmtClass: K = CXCursor_UnexposedStmt; break; + case Stmt::OMPMetaDirectiveClass: + K = CXCursor_OMPMetaDirective; + break; case Stmt::OMPParallelDirectiveClass: K = CXCursor_OMPParallelDirective; break; diff --git a/llvm/include/llvm/Frontend/OpenMP/OMP.td b/llvm/include/llvm/Frontend/OpenMP/OMP.td --- a/llvm/include/llvm/Frontend/OpenMP/OMP.td +++ b/llvm/include/llvm/Frontend/OpenMP/OMP.td @@ -59,6 +59,7 @@ let clangClass = "OMPCollapseClause"; let flangClassValue = "ScalarIntConstantExpr"; } +def OMPC_When: Clause<"when"> {} def OMPC_Default : Clause<"default"> { let clangClass = "OMPDefaultClause"; let flangClass = "OmpDefaultClause"; @@ -320,6 +321,14 @@ // Definition of OpenMP directives //===----------------------------------------------------------------------===// +def OMP_Metadirective : Directive<"metadirective"> { + let allowedClauses = [ + VersionedClause + ]; + let allowedOnceClauses = [ + VersionedClause + ]; +} def OMP_ThreadPrivate : Directive<"threadprivate"> {} def OMP_Parallel : Directive<"parallel"> { let allowedClauses = [ diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPContext.h b/llvm/include/llvm/Frontend/OpenMP/OMPContext.h --- a/llvm/include/llvm/Frontend/OpenMP/OMPContext.h +++ b/llvm/include/llvm/Frontend/OpenMP/OMPContext.h @@ -184,6 +184,17 @@ const OMPContext &Ctx, bool DeviceSetOnly = false); +/// Return the index (into \p VMIs) of the When clause with the highest score +/// from the ones applicble in \p Ctx. +/// In OpenMP 5.1 set OrderedMatch to those conditions which need runtime +/// resolution. +int getBestWhenMatchForContext( + const SmallVectorImpl &VMIs, const OMPContext &Ctx, + SmallVectorImpl *OrderedMatch = nullptr); +/// Return the index (into \p VMIs) of the When clause with the highest score +/// from the ones applicble in \p Ctx. +// int getBestWhenMatchForContext(const SmallVectorImpl &VMIs, +// const OMPContext &Ctx); /// Return the index (into \p VMIs) of the variant with the highest score /// from the ones applicble in \p Ctx. See llvm::isVariantApplicableInContext. int getBestVariantMatchForContext(const SmallVectorImpl &VMIs, diff --git a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def --- a/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def +++ b/llvm/include/llvm/Frontend/OpenMP/OMPKinds.def @@ -115,6 +115,7 @@ __OMP_CLAUSE_NO_CLASS(uniform) __OMP_CLAUSE_NO_CLASS(device_type) __OMP_CLAUSE_NO_CLASS(match) +__OMP_CLAUSE_NO_CLASS(when) __OMP_IMPLICIT_CLAUSE_CLASS(depobj, "depobj", OMPDepobjClause) __OMP_IMPLICIT_CLAUSE_CLASS(flush, "flush", OMPFlushClause) diff --git a/llvm/lib/Frontend/OpenMP/OMPContext.cpp b/llvm/lib/Frontend/OpenMP/OMPContext.cpp --- a/llvm/lib/Frontend/OpenMP/OMPContext.cpp +++ b/llvm/lib/Frontend/OpenMP/OMPContext.cpp @@ -332,6 +332,73 @@ return Score; } +int llvm::omp::getBestWhenMatchForContext( + const SmallVectorImpl &VMIs, const OMPContext &Ctx, + SmallVectorImpl *OrderedMatch) { + + APInt BestScore(64, 0); + int BestVMIIdx = -1; + const VariantMatchInfo *BestVMI = nullptr; + + for (unsigned u = 0, e = VMIs.size(); u < e; ++u) { + const VariantMatchInfo &VMI = VMIs[u]; + + SmallVector ConstructMatches; + // Check if its clearly not the best. + APInt Score = getVariantMatchScore(VMI, Ctx, ConstructMatches); + if (Score.ult(BestScore)) + continue; + // Equal score need subset checks. + if (Score.eq(BestScore)) { + // Strict subset are never best. + if (isStrictSubset(VMI, *BestVMI)) + continue; + // Same score and the current best is no strict subset so we keep it. + if (!isStrictSubset(*BestVMI, VMI)) + continue; + } + // New best found. + BestVMI = &VMI; + BestVMIIdx = u; + BestScore = Score; + } + + return BestVMIIdx; +} + +/*int llvm::omp::getBestWhenMatchForContext( + const SmallVectorImpl &VMIs, const OMPContext &Ctx) { + + APInt BestScore(64, 0); + int BestVMIIdx = -1; + const VariantMatchInfo *BestVMI = nullptr; + + for (unsigned u = 0, e = VMIs.size(); u < e; ++u) { + const VariantMatchInfo &VMI = VMIs[u]; + + SmallVector ConstructMatches; + // Check if its clearly not the best. + APInt Score = getVariantMatchScore(VMI, Ctx, ConstructMatches); + if (Score.ult(BestScore)) + continue; + // Equal score need subset checks. + if (Score.eq(BestScore)) { + // Strict subset are never best. + if (isStrictSubset(VMI, *BestVMI)) + continue; + // Same score and the current best is no strict subset so we keep it. + if (!isStrictSubset(*BestVMI, VMI)) + continue; + } + // New best found. + BestVMI = &VMI; + BestVMIIdx = u; + BestScore = Score; + } + + return BestVMIIdx; +}*/ + int llvm::omp::getBestVariantMatchForContext( const SmallVectorImpl &VMIs, const OMPContext &Ctx) {