OpenMP 5.0 metadirective

This patch supports OpenMP 5.0 metadirective features.
It is implemented keeping the OpenMP 5.1 features like dynamic user condition in mind.

A new function, getBestWhenMatchForContext, is defined in llvm/Frontend/OpenMP/OMPContext.h

Currently this function return the index of the when clause with the highest score from the ones applicable in the Context.
But this function is declared with an array which can be used in OpenMP 5.1 implementation to select all the valid when clauses which can be resolved in runtime. Currently this array is set to null by default and its implementation is left for future.

Reviewed By: jdoerfert

Differential Revision: https://reviews.llvm.org/D91944

clang/include/clang-c/Index.h

@@ -2592,7 +2592,11 @@ enum CXCursorKind {
    */
   CXCursor_OMPUnrollDirective = 293,
 
-  CXCursor_LastStmt = CXCursor_OMPUnrollDirective,
+  /** OpenMP metadirective directive.
+   */
+  CXCursor_OMPMetaDirective = 294,
+
+  CXCursor_LastStmt = CXCursor_OMPMetaDirective,
 
   /**
    * Cursor that represents the translation unit itself.

clang/include/clang/AST/RecursiveASTVisitor.h

@@ -2842,6 +2842,9 @@ RecursiveASTVisitor<Derived>::TraverseOMPLoopDirective(OMPLoopDirective *S) {
   return TraverseOMPExecutableDirective(S);
 }
 
+DEF_TRAVERSE_STMT(OMPMetaDirective,
+                  { TRY_TO(TraverseOMPExecutableDirective(S)); })
+
 DEF_TRAVERSE_STMT(OMPParallelDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 

clang/include/clang/AST/StmtOpenMP.h

@@ -5379,6 +5379,44 @@ public:
   }
 };
 
+/// 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()) {}
+
+  void setIfStmt(Stmt *S) { IfStmt = S; }
+
+public:
+  static OMPMetaDirective *Create(const ASTContext &C, SourceLocation StartLoc,
+                                  SourceLocation EndLoc,
+                                  ArrayRef<OMPClause *> Clauses,
+                                  Stmt *AssociatedStmt, Stmt *IfStmt);
+  static OMPMetaDirective *CreateEmpty(const ASTContext &C, unsigned NumClauses,
+                                       EmptyShell);
+  Stmt *getIfStmt() const { return IfStmt; }
+
+  static bool classof(const Stmt *T) {
+    return T->getStmtClass() == OMPMetaDirectiveClass;
+  }
+};
+
 } // end namespace clang
 
 #endif

clang/include/clang/Basic/DiagnosticParseKinds.td

@@ -1436,6 +1436,9 @@ def warn_omp51_compat_attributes : Warning<
   "specifying OpenMP directives with [[]] is incompatible with OpenMP "
   "standards before OpenMP 5.1">,
   InGroup<OpenMPPre51Compat>, DefaultIgnore;
+def err_omp_expected_colon : Error<"missing ':' in %0">;
+def err_omp_expected_context_selector
+    : Error<"expected valid context selector in %0">;
 
 // Pragma loop support.
 def err_pragma_loop_missing_argument : Error<

clang/include/clang/Basic/DiagnosticSemaKinds.td

@@ -10794,6 +10794,8 @@ def err_omp_dispatch_statement_call
 def err_omp_unroll_full_variable_trip_count : Error<
   "loop to be fully unrolled must have a constant trip count">;
 def note_omp_directive_here : Note<"'%0' directive found here">;
+def err_omp_instantiation_not_supported
+    : Error<"instantiation of '%0' not supported yet">;
 } // end of OpenMP category
 
 let CategoryName = "Related Result Type Issue" in {

clang/include/clang/Basic/StmtNodes.td

@@ -219,6 +219,7 @@ def AsTypeExpr : StmtNode<Expr>;
 // OpenMP Directives.
 def OMPCanonicalLoop : StmtNode<Stmt>;
 def OMPExecutableDirective : StmtNode<Stmt, 1>;
+def OMPMetaDirective : StmtNode<OMPExecutableDirective>;
 def OMPLoopBasedDirective : StmtNode<OMPExecutableDirective, 1>;
 def OMPLoopDirective : StmtNode<OMPLoopBasedDirective, 1>;
 def OMPParallelDirective : StmtNode<OMPExecutableDirective>;

clang/include/clang/Sema/Sema.h

@@ -10456,6 +10456,12 @@ public:
   /// \param Init First part of the for loop.
   void ActOnOpenMPLoopInitialization(SourceLocation ForLoc, Stmt *Init);
 
+  /// Called on well-formed '\#pragma omp metadirective' after parsing
+  /// of the  associated statement.
+  StmtResult ActOnOpenMPMetaDirective(ArrayRef<OMPClause *> Clauses,
+                                      Stmt *AStmt, SourceLocation StartLoc,
+                                      SourceLocation EndLoc);
+
   // OpenMP directives and clauses.
   /// Called on correct id-expression from the '#pragma omp
   /// threadprivate'.
@@ -11023,6 +11029,10 @@ public:
                                      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,

clang/include/clang/Serialization/ASTBitCodes.h

@@ -1893,6 +1893,7 @@ enum StmtCode {
   STMT_SEH_TRY,                     // SEHTryStmt
 
   // OpenMP directives
+  STMT_OMP_META_DIRECTIVE,
   STMT_OMP_CANONICAL_LOOP,
   STMT_OMP_PARALLEL_DIRECTIVE,
   STMT_OMP_SIMD_DIRECTIVE,

clang/lib/AST/OpenMPClause.cpp

@@ -160,6 +160,7 @@ const OMPClauseWithPreInit *OMPClauseWithPreInit::get(const OMPClause *C) {
   case OMPC_exclusive:
   case OMPC_uses_allocators:
   case OMPC_affinity:
+  case OMPC_when:
     break;
   default:
     break;
@@ -257,6 +258,7 @@ const OMPClauseWithPostUpdate *OMPClauseWithPostUpdate::get(const OMPClause *C)
   case OMPC_exclusive:
   case OMPC_uses_allocators:
   case OMPC_affinity:
+  case OMPC_when:
     break;
   default:
     break;

clang/lib/AST/StmtOpenMP.cpp

@@ -253,6 +253,25 @@ void OMPLoopDirective::setFinalsConditions(ArrayRef<Expr *> A) {
   llvm::copy(A, getFinalsConditions().begin());
 }
 
+OMPMetaDirective *OMPMetaDirective::Create(const ASTContext &C,
+                                           SourceLocation StartLoc,
+                                           SourceLocation EndLoc,
+                                           ArrayRef<OMPClause *> Clauses,
+                                           Stmt *AssociatedStmt, Stmt *IfStmt) {
+  auto *Dir = createDirective<OMPMetaDirective>(
+      C, Clauses, AssociatedStmt, /*NumChildren=*/1, StartLoc, EndLoc);
+  Dir->setIfStmt(IfStmt);
+  return Dir;
+}
+
+OMPMetaDirective *OMPMetaDirective::CreateEmpty(const ASTContext &C,
+                                                unsigned NumClauses,
+                                                EmptyShell) {
+  return createEmptyDirective<OMPMetaDirective>(C, NumClauses,
+                                                /*HasAssociatedStmt=*/true,
+                                                /*NumChildren=*/1);
+}
+
 OMPParallelDirective *OMPParallelDirective::Create(
     const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
     ArrayRef<OMPClause *> Clauses, Stmt *AssociatedStmt, Expr *TaskRedRef,

clang/lib/AST/StmtPrinter.cpp

@@ -654,6 +654,11 @@ void StmtPrinter::PrintOMPExecutableDirective(OMPExecutableDirective *S,
     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);

clang/lib/AST/StmtProfile.cpp

@@ -903,6 +903,10 @@ void StmtProfiler::VisitOMPLoopDirective(const OMPLoopDirective *S) {
   VisitOMPLoopBasedDirective(S);
 }
 
+void StmtProfiler::VisitOMPMetaDirective(const OMPMetaDirective *S) {
+  VisitOMPExecutableDirective(S);
+}
+
 void StmtProfiler::VisitOMPParallelDirective(const OMPParallelDirective *S) {
   VisitOMPExecutableDirective(S);
 }

clang/lib/Basic/OpenMPKinds.cpp

@@ -185,6 +185,7 @@ unsigned clang::getOpenMPSimpleClauseType(OpenMPClauseKind Kind, StringRef Str,
   case OMPC_exclusive:
   case OMPC_uses_allocators:
   case OMPC_affinity:
+  case OMPC_when:
     break;
   default:
     break;
@@ -428,6 +429,7 @@ const char *clang::getOpenMPSimpleClauseTypeName(OpenMPClauseKind Kind,
   case OMPC_exclusive:
   case OMPC_uses_allocators:
   case OMPC_affinity:
+  case OMPC_when:
     break;
   default:
     break;
@@ -591,6 +593,9 @@ void clang::getOpenMPCaptureRegions(
     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:

clang/lib/CodeGen/CGOpenMPRuntime.cpp

@@ -6740,6 +6740,7 @@ const Expr *CGOpenMPRuntime::getNumTeamsExprForTargetDirective(
   case OMPD_parallel_master_taskloop:
   case OMPD_parallel_master_taskloop_simd:
   case OMPD_requires:
+  case OMPD_metadirective:
   case OMPD_unknown:
     break;
   default:
@@ -7214,6 +7215,7 @@ llvm::Value *CGOpenMPRuntime::emitNumThreadsForTargetDirective(
   case OMPD_parallel_master_taskloop:
   case OMPD_parallel_master_taskloop_simd:
   case OMPD_requires:
+  case OMPD_metadirective:
   case OMPD_unknown:
     break;
   default:
@@ -9851,6 +9853,7 @@ getNestedDistributeDirective(ASTContext &Ctx, const OMPExecutableDirective &D) {
     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.");
@@ -10701,6 +10704,7 @@ void CGOpenMPRuntime::scanForTargetRegionsFunctions(const Stmt *S,
     case OMPD_parallel_master_taskloop:
     case OMPD_parallel_master_taskloop_simd:
     case OMPD_requires:
+    case OMPD_metadirective:
     case OMPD_unknown:
     default:
       llvm_unreachable("Unknown target directive for OpenMP device codegen.");
@@ -11382,6 +11386,7 @@ void CGOpenMPRuntime::emitTargetDataStandAloneCall(
     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.");

clang/lib/CodeGen/CGStmt.cpp

@@ -196,6 +196,9 @@ void CodeGenFunction::EmitStmt(const Stmt *S, ArrayRef<const Attr *> Attrs) {
   case Stmt::SEHTryStmtClass:
     EmitSEHTryStmt(cast<SEHTryStmt>(*S));
     break;
+  case Stmt::OMPMetaDirectiveClass:
+    EmitOMPMetaDirective(cast<OMPMetaDirective>(*S));
+    break;
   case Stmt::OMPCanonicalLoopClass:
     EmitOMPCanonicalLoop(cast<OMPCanonicalLoop>(S));
     break;

clang/lib/CodeGen/CGStmtOpenMP.cpp

@@ -1784,6 +1784,10 @@ void CodeGenFunction::EmitOMPParallelDirective(const OMPParallelDirective &S) {
   checkForLastprivateConditionalUpdate(*this, S);
 }
 
+void CodeGenFunction::EmitOMPMetaDirective(const OMPMetaDirective &S) {
+  EmitStmt(S.getIfStmt());
+}
+
 namespace {
 /// RAII to handle scopes for loop transformation directives.
 class OMPTransformDirectiveScopeRAII {
@@ -5986,6 +5990,7 @@ static void emitOMPAtomicExpr(CodeGenFunction &CGF, OpenMPClauseKind Kind,
   case OMPC_novariants:
   case OMPC_nocontext:
   case OMPC_filter:
+  case OMPC_when:
     llvm_unreachable("Clause is not allowed in 'omp atomic'.");
   }
 }

clang/lib/CodeGen/CodeGenFunction.h

@@ -3464,6 +3464,7 @@ public:
                                        const RegionCodeGenTy &BodyGen,
                                        OMPTargetDataInfo &InputInfo);
 
+  void EmitOMPMetaDirective(const OMPMetaDirective &S);
   void EmitOMPParallelDirective(const OMPParallelDirective &S);
   void EmitOMPSimdDirective(const OMPSimdDirective &S);
   void EmitOMPTileDirective(const OMPTileDirective &S);

clang/lib/Parse/ParseOpenMP.cpp

@@ -2226,6 +2226,7 @@ Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirectiveWithExtDecl(
   case OMPD_target_teams_distribute_simd:
   case OMPD_dispatch:
   case OMPD_masked:
+  case OMPD_metadirective:
     Diag(Tok, diag::err_omp_unexpected_directive)
         << 1 << getOpenMPDirectiveName(DKind);
     break;
@@ -2280,8 +2281,10 @@ Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirectiveWithExtDecl(
 ///
 StmtResult
 Parser::ParseOpenMPDeclarativeOrExecutableDirective(ParsedStmtContext StmtCtx) {
-  assert(Tok.isOneOf(tok::annot_pragma_openmp, tok::annot_attr_openmp) &&
-         "Not an OpenMP directive!");
+  static bool ReadDirectiveWithinMetadirective = false;
+  if (!ReadDirectiveWithinMetadirective)
+    assert(Tok.isOneOf(tok::annot_pragma_openmp, tok::annot_attr_openmp) &&
+           "Not an OpenMP directive!");
   ParsingOpenMPDirectiveRAII DirScope(*this);
   ParenBraceBracketBalancer BalancerRAIIObj(*this);
   SmallVector<OMPClause *, 5> Clauses;
@@ -2290,8 +2293,15 @@ Parser::ParseOpenMPDeclarativeOrExecutableDirective(ParsedStmtContext StmtCtx) {
       FirstClauses(llvm::omp::Clause_enumSize + 1);
   unsigned ScopeFlags = Scope::FnScope | Scope::DeclScope |
                         Scope::CompoundStmtScope | Scope::OpenMPDirectiveScope;
-  SourceLocation Loc = ConsumeAnnotationToken(), EndLoc;
+  SourceLocation Loc = ReadDirectiveWithinMetadirective
+                           ? Tok.getLocation()
+                           : ConsumeAnnotationToken(),
+                 EndLoc;
   OpenMPDirectiveKind DKind = parseOpenMPDirectiveKind(*this);
+  if (ReadDirectiveWithinMetadirective && DKind == OMPD_unknown) {
+    Diag(Tok, diag::err_omp_unknown_directive);
+    return StmtError();
+  }
   OpenMPDirectiveKind CancelRegion = OMPD_unknown;
   // Name of critical directive.
   DeclarationNameInfo DirName;
@@ -2299,6 +2309,141 @@ Parser::ParseOpenMPDeclarativeOrExecutableDirective(ParsedStmtContext StmtCtx) {
   bool HasAssociatedStatement = true;
 
   switch (DKind) {
+  case OMPD_metadirective: {
+    ConsumeToken();
+    SmallVector<VariantMatchInfo, 4> VMIs;
+
+    // First iteration of parsing all clauses of metadirective.
+    // This iteration only parses and collects all context selector ignoring the
+    // associated directives.
+    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 OMPTraitInfo to pass to the When clause
+        parseOMPContextSelectors(Loc, TI);
+        if (TI.Sets.size() == 0) {
+          Diag(Tok, diag::err_omp_expected_context_selector) << "when clause";
+          TPA.Commit();
+          return Directive;
+        }
+
+        // Parse ':'
+        if (Tok.is(tok::colon))
+          ConsumeAnyToken();
+        else {
+          Diag(Tok, diag::err_omp_expected_colon) << "when clause";
+          TPA.Commit();
+          return Directive;
+        }
+      }
+      // Skip Directive for now. We will parse directive in the second iteration
+      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--;
+        if (Tok.is(tok::annot_pragma_openmp_end)) {
+          Diag(Tok, diag::err_omp_expected_punc)
+              << getOpenMPClauseName(CKind) << 0;
+          TPA.Commit();
+          return Directive;
+        }
+        ConsumeAnyToken();
+      }
+      // Parse ')'
+      if (Tok.is(tok::r_paren))
+        T.consumeClose();
+
+      VariantMatchInfo VMI;
+      TI.getAsVariantMatchInfo(ASTContext, VMI);
+
+      VMIs.push_back(VMI);
+    }
+
+    TPA.Revert();
+    // End of the first iteration. Parser is reset to the start of metadirective
+
+    TargetOMPContext OMPCtx(ASTContext, /* DiagUnknownTrait */ nullptr,
+                            /* CurrentFunctionDecl */ nullptr,
+                            ArrayRef<llvm::omp::TraitProperty>());
+
+    // A single match is returned for OpenMP 5.0
+    int BestIdx = getBestVariantMatchForContext(VMIs, OMPCtx);
+
+    int Idx = 0;
+    // In OpenMP 5.0 metadirective is either replaced by another directive or
+    // ignored.
+    // TODO: In OpenMP 5.1 generate multiple directives based upon the matches
+    // found by getBestWhenMatchForContext.
+    while (Tok.isNot(tok::annot_pragma_openmp_end)) {
+      // OpenMP 5.0 implementation - Skip to the best index found.
+      if (Idx++ != BestIdx) {
+        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 the ContextSelectors
+        parseOMPContextSelectors(Loc, TI);
+
+        // Parse ':'
+        ConsumeAnyToken();
+      }
+
+      // If no directive is passed, skip in OpenMP 5.0.
+      // TODO: Generate nothing directive from OpenMP 5.1.
+      if (Tok.is(tok::r_paren)) {
+        SkipUntil(tok::annot_pragma_openmp_end);
+        break;
+      }
+
+      // Parse Directive
+      ReadDirectiveWithinMetadirective = true;
+      Directive = ParseOpenMPDeclarativeOrExecutableDirective(StmtCtx);
+      ReadDirectiveWithinMetadirective = false;
+      break;
+    }
+    break;
+  }
   case OMPD_threadprivate: {
     // FIXME: Should this be permitted in C++?
     if ((StmtCtx & ParsedStmtContext::AllowDeclarationsInC) ==
@@ -2490,6 +2635,13 @@ Parser::ParseOpenMPDeclarativeOrExecutableDirective(ParsedStmtContext StmtCtx) {
     Actions.StartOpenMPDSABlock(DKind, DirName, Actions.getCurScope(), Loc);
 
     while (Tok.isNot(tok::annot_pragma_openmp_end)) {
+      // If we are parsing for a directive within a metadirective, the directive
+      // ends with a ')'.
+      if (ReadDirectiveWithinMetadirective && Tok.is(tok::r_paren)) {
+        while (Tok.isNot(tok::annot_pragma_openmp_end))
+          ConsumeAnyToken();
+        break;
+      }
       bool HasImplicitClause = false;
       if (ImplicitClauseAllowed && Tok.is(tok::l_paren)) {
         HasImplicitClause = true;

clang/lib/Sema/SemaExceptionSpec.cpp

@@ -1496,6 +1496,7 @@ CanThrowResult Sema::canThrow(const Stmt *S) {
   case Stmt::OMPInteropDirectiveClass:
   case Stmt::OMPDispatchDirectiveClass:
   case Stmt::OMPMaskedDirectiveClass:
+  case Stmt::OMPMetaDirectiveClass:
   case Stmt::ReturnStmtClass:
   case Stmt::SEHExceptStmtClass:
   case Stmt::SEHFinallyStmtClass:

clang/lib/Sema/SemaOpenMP.cpp

@@ -4314,6 +4314,7 @@ void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, Scope *CurScope) {
   case OMPD_declare_variant:
   case OMPD_begin_declare_variant:
   case OMPD_end_declare_variant:
+  case OMPD_metadirective:
     llvm_unreachable("OpenMP Directive is not allowed");
   case OMPD_unknown:
   default:
@@ -6363,6 +6364,7 @@ StmtResult Sema::ActOnOpenMPExecutableDirective(
       case OMPC_atomic_default_mem_order:
       case OMPC_device_type:
       case OMPC_match:
+      case OMPC_when:
       default:
         llvm_unreachable("Unexpected clause");
       }
@@ -13300,6 +13302,7 @@ OMPClause *Sema::ActOnOpenMPSingleExprClause(OpenMPClauseKind Kind, Expr *Expr,
   case OMPC_exclusive:
   case OMPC_uses_allocators:
   case OMPC_affinity:
+  case OMPC_when:
   default:
     llvm_unreachable("Clause is not allowed.");
   }
@@ -13456,6 +13459,7 @@ static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
     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:
@@ -13538,6 +13542,7 @@ static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
     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:
@@ -13618,6 +13623,7 @@ static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
     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:
@@ -13698,6 +13704,7 @@ static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
     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:
@@ -13778,6 +13785,7 @@ static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
     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:
@@ -13858,6 +13866,7 @@ static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
     case OMPD_atomic:
     case OMPD_target_teams:
     case OMPD_requires:
+    case OMPD_metadirective:
       llvm_unreachable("Unexpected OpenMP directive with dist_schedule clause");
     case OMPD_unknown:
     default:
@@ -13940,6 +13949,7 @@ static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
     case OMPD_atomic:
     case OMPD_distribute_simd:
     case OMPD_requires:
+    case OMPD_metadirective:
       llvm_unreachable("Unexpected OpenMP directive with device-clause");
     case OMPD_unknown:
     default:
@@ -14022,6 +14032,7 @@ static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
     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:
@@ -14041,6 +14052,15 @@ static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
   case OMPC_filter:
     // Do not capture filter-clause expressions.
     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:
@@ -14568,6 +14588,7 @@ OMPClause *Sema::ActOnOpenMPSimpleClause(
   case OMPC_exclusive:
   case OMPC_uses_allocators:
   case OMPC_affinity:
+  case OMPC_when:
   default:
     llvm_unreachable("Clause is not allowed.");
   }
@@ -14860,6 +14881,7 @@ OMPClause *Sema::ActOnOpenMPSingleExprWithArgClause(
   case OMPC_exclusive:
   case OMPC_uses_allocators:
   case OMPC_affinity:
+  case OMPC_when:
   default:
     llvm_unreachable("Clause is not allowed.");
   }
@@ -15109,6 +15131,7 @@ OMPClause *Sema::ActOnOpenMPClause(OpenMPClauseKind Kind,
   case OMPC_exclusive:
   case OMPC_uses_allocators:
   case OMPC_affinity:
+  case OMPC_when:
   default:
     llvm_unreachable("Clause is not allowed.");
   }
@@ -15651,6 +15674,7 @@ OMPClause *Sema::ActOnOpenMPVarListClause(
   case OMPC_nocontext:
   case OMPC_detach:
   case OMPC_uses_allocators:
+  case OMPC_when:
   default:
     llvm_unreachable("Clause is not allowed.");
   }

clang/lib/Sema/TreeTransform.h

@@ -8504,6 +8504,15 @@ StmtResult TreeTransform<Derived>::TransformOMPExecutableDirective(
       AssociatedStmt.get(), D->getBeginLoc(), D->getEndLoc());
 }
 
+template <typename Derived>
+StmtResult
+TreeTransform<Derived>::TransformOMPMetaDirective(OMPMetaDirective *D) {
+  // TODO: Fix This
+  SemaRef.Diag(D->getBeginLoc(), diag::err_omp_instantiation_not_supported)
+      << getOpenMPDirectiveName(D->getDirectiveKind());
+  return StmtError();
+}
+
 template <typename Derived>
 StmtResult
 TreeTransform<Derived>::TransformOMPParallelDirective(OMPParallelDirective *D) {

clang/lib/Serialization/ASTReaderStmt.cpp

@@ -2307,6 +2307,13 @@ void ASTStmtReader::VisitOMPLoopDirective(OMPLoopDirective *D) {
   VisitOMPLoopBasedDirective(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);
@@ -3183,6 +3190,11 @@ Stmt *ASTReader::ReadStmtFromStream(ModuleFile &F) {
       S = OMPCanonicalLoop::createEmpty(Context);
       break;
 
+    case STMT_OMP_META_DIRECTIVE:
+      S = OMPMetaDirective::CreateEmpty(
+          Context, Record[ASTStmtReader::NumStmtFields], Empty);
+      break;
+
     case STMT_OMP_PARALLEL_DIRECTIVE:
       S =
         OMPParallelDirective::CreateEmpty(Context,

clang/lib/Serialization/ASTWriterStmt.cpp

@@ -2205,6 +2205,13 @@ void ASTStmtWriter::VisitOMPLoopDirective(OMPLoopDirective *D) {
   VisitOMPLoopBasedDirective(D);
 }
 
+void ASTStmtWriter::VisitOMPMetaDirective(OMPMetaDirective *D) {
+  VisitStmt(D);
+  Record.push_back(D->getNumClauses());
+  VisitOMPExecutableDirective(D);
+  Code = serialization::STMT_OMP_META_DIRECTIVE;
+}
+
 void ASTStmtWriter::VisitOMPParallelDirective(OMPParallelDirective *D) {
   VisitStmt(D);
   VisitOMPExecutableDirective(D);

clang/lib/StaticAnalyzer/Core/ExprEngine.cpp

@@ -1298,7 +1298,8 @@ void ExprEngine::Visit(const Stmt *S, ExplodedNode *Pred,
     case Stmt::OMPDispatchDirectiveClass:
     case Stmt::OMPMaskedDirectiveClass:
     case Stmt::CapturedStmtClass:
-    case Stmt::OMPUnrollDirectiveClass: {
+    case Stmt::OMPUnrollDirectiveClass:
+    case Stmt::OMPMetaDirectiveClass: {
       const ExplodedNode *node = Bldr.generateSink(S, Pred, Pred->getState());
       Engine.addAbortedBlock(node, currBldrCtx->getBlock());
       break;

clang/test/OpenMP/metadirective_ast_print.c

@@ -0,0 +1,73 @@
+// RUN: %clang_cc1 -verify -fopenmp -x c -std=c99 -ast-print %s -o - | FileCheck %s
+
+// RUN: %clang_cc1 -verify -fopenmp-simd -x c -std=c99 -ast-print %s -o - | FileCheck %s
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+void bar(void);
+
+#define N 10
+void foo(void) {
+#pragma omp metadirective when(device = {kind(cpu)} \
+                               : parallel) default()
+  bar();
+#pragma omp metadirective when(implementation = {vendor(score(0)  \
+                                                        : llvm)}, \
+                               device = {kind(cpu)}               \
+                               : parallel) default(target teams)
+  bar();
+#pragma omp metadirective when(device = {kind(gpu)}                                 \
+                               : target teams) when(implementation = {vendor(llvm)} \
+                                                    : parallel) default()
+  bar();
+#pragma omp metadirective default(target) when(implementation = {vendor(score(5)  \
+                                                                        : llvm)}, \
+                                               device = {kind(cpu, host)}         \
+                                               : parallel)
+  bar();
+#pragma omp metadirective when(user = {condition(N > 10)}                 \
+                               : target) when(user = {condition(N == 10)} \
+                                              : parallel)
+  bar();
+#pragma omp metadirective when(device = {kind(host)} \
+                               : parallel for)
+  for (int i = 0; i < 100; i++)
+    ;
+#pragma omp metadirective when(implementation = {extension(match_all)} \
+                               : parallel) default(parallel for)
+  for (int i = 0; i < 100; i++)
+    ;
+#pragma omp metadirective when(implementation = {extension(match_any)} \
+                               : parallel) default(parallel for)
+  for (int i = 0; i < 100; i++)
+    ;
+#pragma omp metadirective when(implementation = {extension(match_none)} \
+                               : parallel) default(parallel for)
+  for (int i = 0; i < 100; i++)
+    ;
+}
+
+// CHECK: void bar();
+// CHECK: void foo()
+// CHECK-NEXT: #pragma omp parallel
+// CHECK-NEXT: bar()
+// CHECK-NEXT: #pragma omp parallel
+// CHECK-NEXT: bar()
+// CHECK-NEXT: #pragma omp parallel
+// CHECK-NEXT: bar()
+// CHECK-NEXT: #pragma omp parallel
+// CHECK-NEXT: bar()
+// CHECK-NEXT: #pragma omp parallel
+// CHECK-NEXT: bar()
+// CHECK-NEXT: #pragma omp parallel for
+// CHECK-NEXT: for (int i = 0; i < 100; i++)
+// CHECK: #pragma omp parallel
+// CHECK-NEXT: for (int i = 0; i < 100; i++)
+// CHECK: #pragma omp parallel for
+// CHECK-NEXT: for (int i = 0; i < 100; i++)
+// CHECK: #pragma omp parallel
+// CHECK-NEXT: for (int i = 0; i < 100; i++)
+
+#endif

clang/test/OpenMP/metadirective_device_kind_codegen.c

@@ -0,0 +1,80 @@
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple x86_64-unknown-linux -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple aarch64-unknown-linux -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple ppc64le-unknown-linux -emit-llvm %s -o - | FileCheck %s
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+void bar();
+
+void foo() {
+#pragma omp metadirective when(device = {kind(any)} \
+                               : parallel)
+  bar();
+#pragma omp metadirective when(device = {kind(host, cpu)} \
+                               : parallel for num_threads(4))
+  for (int i = 0; i < 100; i++)
+    ;
+#pragma omp metadirective when(device = {kind(host)} \
+                               : parallel for)
+  for (int i = 0; i < 100; i++)
+    ;
+#pragma omp metadirective when(device = {kind(nohost, gpu)} \
+                               :) when(device = {kind(cpu)} \
+                                       : parallel)
+  bar();
+#pragma omp metadirective when(device = {kind(any, cpu)} \
+                               : parallel)
+  bar();
+#pragma omp metadirective when(device = {kind(any, host)} \
+                               : parallel)
+  bar();
+#pragma omp metadirective when(device = {kind(gpu)} \
+                               : target parallel for) default(parallel for)
+  for (int i = 0; i < 100; i++)
+    ;
+}
+
+// CHECK-LABEL: define {{.+}} void @foo()
+// CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_1:@.+]] to void
+// CHECK-NEXT: @__kmpc_push_num_threads
+// CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_2:@.+]] to void
+// CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_3:@.+]] to void
+// CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_4:@.+]] to void
+// CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_5:@.+]] to void
+// CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_6:@.+]] to void
+// CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_7:@.+]] to void
+// CHECK: ret void
+
+// CHECK: define internal void [[OUTLINED_1]](
+// CHECK: call void {{.+}} @bar
+// CHECK: ret void
+
+// CHECK: define internal void [[OUTLINED_2]](
+// CHECK: call void @__kmpc_for_static_init
+// CHECK: call void @__kmpc_for_static_fini
+// CHECK: ret void
+
+// CHECK: define internal void [[OUTLINED_3]](
+// CHECK: call void @__kmpc_for_static_init
+// CHECK: call void @__kmpc_for_static_fini
+// CHECK: ret void
+
+// CHECK: define internal void [[OUTLINED_4]](
+// CHECK: call void {{.+}} @bar
+// CHECK: ret void
+
+// CHECK: define internal void [[OUTLINED_5]](
+// CHECK: call void {{.+}} @bar
+// CHECK: ret void
+
+// CHECK: define internal void [[OUTLINED_6]](
+// CHECK: call void {{.+}} @bar
+// CHECK: ret void
+
+// CHECK: define internal void [[OUTLINED_7]](
+// CHECK: call void @__kmpc_for_static_init
+// CHECK: call void @__kmpc_for_static_fini
+// CHECK: ret void
+
+#endif

clang/test/OpenMP/metadirective_device_kind_codegen.cpp

@@ -0,0 +1,81 @@
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple aarch64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple ppc64le-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope | FileCheck %s
+
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+void bar();
+
+void foo() {
+#pragma omp metadirective when(device = {kind(any)} \
+                               : parallel)
+  bar();
+#pragma omp metadirective when(device = {kind(host, cpu)} \
+                               : parallel for num_threads(4))
+  for (int i = 0; i < 100; i++)
+    ;
+#pragma omp metadirective when(device = {kind(host)} \
+                               : parallel for)
+  for (int i = 0; i < 100; i++)
+    ;
+#pragma omp metadirective when(device = {kind(nohost, gpu)} \
+                               :) when(device = {kind(cpu)} \
+                                       : parallel)
+  bar();
+#pragma omp metadirective when(device = {kind(any, cpu)} \
+                               : parallel)
+  bar();
+#pragma omp metadirective when(device = {kind(any, host)} \
+                               : parallel)
+  bar();
+#pragma omp metadirective when(device = {kind(gpu)} \
+                               : target parallel for) default(parallel for)
+  for (int i = 0; i < 100; i++)
+    ;
+}
+
+// CHECK-LABEL: define {{.+}} void @_Z3foov()
+// CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_1:@.+]] to void
+// CHECK-NEXT: @__kmpc_push_num_threads
+// CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_2:@.+]] to void
+// CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_3:@.+]] to void
+// CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_4:@.+]] to void
+// CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_5:@.+]] to void
+// CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_6:@.+]] to void
+// CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_7:@.+]] to void
+// CHECK: ret void
+
+// CHECK: define internal void [[OUTLINED_1]](
+// CHECK: void @_Z3barv()
+// CHECK: ret void
+
+// CHECK: define internal void [[OUTLINED_2]](
+// CHECK: call void @__kmpc_for_static_init
+// CHECK: call void @__kmpc_for_static_fini
+// CHECK: ret void
+
+// CHECK: define internal void [[OUTLINED_3]](
+// CHECK: call void @__kmpc_for_static_init
+// CHECK: call void @__kmpc_for_static_fini
+// CHECK: ret void
+
+// CHECK: define internal void [[OUTLINED_4]](
+// CHECK: void @_Z3barv()
+// CHECK: ret void
+
+// CHECK: define internal void [[OUTLINED_5]](
+// CHECK: void @_Z3barv()
+// CHECK: ret void
+
+// CHECK: define internal void [[OUTLINED_6]](
+// CHECK: void @_Z3barv()
+// CHECK: ret void
+
+// CHECK: define internal void [[OUTLINED_7]](
+// CHECK: call void @__kmpc_for_static_init
+// CHECK: call void @__kmpc_for_static_fini
+// CHECK: ret void
+
+#endif

clang/test/OpenMP/metadirective_empty.cpp

@@ -0,0 +1,39 @@
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple aarch64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple ppc64le-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck %s
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+#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-LABEL: void @_Z4funcv()
+// CHECK: entry:
+// CHECK:   [[I:%.+]] = alloca i32,
+// CHECK:   store i32 0, i32* [[I]],
+// CHECK:   br label %[[FOR_COND:.+]]
+// CHECK: [[FOR_COND]]:
+// CHECK:   [[ZERO:%.+]] = load i32, i32* [[I]],
+// CHECK:   [[CMP:%.+]] = icmp slt i32 [[ZERO]], 1000
+// CHECK:   br i1 [[CMP]], label %[[FOR_BODY:.+]], label %[[FOR_END:.+]]
+// CHECK: [[FOR_BODY]]:
+// CHECK:   br label %[[FOR_INC:.+]]
+// CHECK: [[FOR_INC]]:
+// CHECK:   [[ONE:%.+]] = load i32, i32* [[I]],
+// CHECK:   [[INC:%.+]] = add nsw i32 [[ONE]], 1
+// CHECK:   store i32 [[INC]], i32* [[I]],
+// CHECK:   br label %[[FOR_COND]],
+// CHECK: [[FOR_END]]:
+// CHECK:   ret void
+// CHECK: }
+
+#endif

clang/test/OpenMP/metadirective_implementation_codegen.c

@@ -0,0 +1,71 @@
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple x86_64-unknown-linux -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple aarch64-unknown-linux -emit-llvm %s -o - | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c -triple ppc64le-unknown-linux -emit-llvm %s -o - | FileCheck %s
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+void bar();
+
+void foo() {
+#pragma omp metadirective when(implementation = {vendor(score(0)  \
+                                                        : llvm)}, \
+                               device = {kind(cpu)}               \
+                               : parallel) default(target teams)
+  bar();
+#pragma omp metadirective when(device = {kind(gpu)}                                 \
+                               : target teams) when(implementation = {vendor(llvm)} \
+                                                    : parallel) default()
+  bar();
+#pragma omp metadirective default(target) when(implementation = {vendor(score(5)  \
+                                                                        : llvm)}, \
+                                               device = {kind(cpu, host)}         \
+                                               : parallel)
+  bar();
+#pragma omp metadirective when(implementation = {extension(match_all)} \
+                               : parallel) default(parallel for)
+  for (int i = 0; i < 100; i++)
+    ;
+#pragma omp metadirective when(implementation = {extension(match_any)} \
+                               : parallel) default(parallel for)
+  for (int i = 0; i < 100; i++)
+    ;
+#pragma omp metadirective when(implementation = {extension(match_none)} \
+                               : parallel) default(parallel for)
+  for (int i = 0; i < 100; i++)
+    ;
+}
+
+// CHECK: void @foo()
+// CHECK-COUNT-6: ...) @__kmpc_fork_call(
+// CHECK: ret void
+
+// CHECK: define internal void @.omp_outlined.(
+// CHECK: @bar
+// CHECK: ret void
+
+// CHECK: define internal void @.omp_outlined..1(
+// CHECK: @bar
+// CHECK: ret void
+
+// CHECK: define internal void @.omp_outlined..2(
+// CHECK: @bar
+// CHECK: ret void
+
+// CHECK: define internal void @.omp_outlined..3(
+// NO-CHECK: call void @__kmpc_for_static_init
+// NO-CHECK: call void @__kmpc_for_static_fini
+// CHECK: ret void
+
+// CHECK: define internal void @.omp_outlined..4(
+// CHECK: call void @__kmpc_for_static_init
+// CHECK: call void @__kmpc_for_static_fini
+// CHECK: ret void
+
+// CHECK: define internal void @.omp_outlined..5(
+// NO-CHECK: call void @__kmpc_for_static_init
+// NO-CHECK: call void @__kmpc_for_static_fini
+// CHECK: ret void
+
+#endif

clang/test/OpenMP/metadirective_implementation_codegen.cpp

@@ -0,0 +1,76 @@
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple x86_64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple aarch64-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple ppc64le-unknown-linux -emit-llvm %s -fexceptions -fcxx-exceptions -o - -fsanitize-address-use-after-scope | FileCheck %s
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+void bar();
+
+void foo() {
+#pragma omp metadirective when(implementation = {vendor(score(0)  \
+                                                        : llvm)}, \
+                               device = {kind(cpu)}               \
+                               : parallel) default(target teams)
+  bar();
+#pragma omp metadirective when(device = {kind(gpu)}                                 \
+                               : target teams) when(implementation = {vendor(llvm)} \
+                                                    : parallel) default()
+  bar();
+#pragma omp metadirective default(target) when(implementation = {vendor(score(5)  \
+                                                                        : llvm)}, \
+                                               device = {kind(cpu, host)}         \
+                                               : parallel)
+  bar();
+#pragma omp metadirective when(implementation = {extension(match_all)} \
+                               : parallel) default(parallel for)
+  for (int i = 0; i < 100; i++)
+    ;
+#pragma omp metadirective when(implementation = {extension(match_any)} \
+                               : parallel) default(parallel for)
+  for (int i = 0; i < 100; i++)
+    ;
+#pragma omp metadirective when(implementation = {extension(match_none)} \
+                               : parallel) default(parallel for)
+  for (int i = 0; i < 100; i++)
+    ;
+}
+
+// CHECK-LABEL: void @_Z3foov()
+// CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_2:@.+]] to void
+// CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_3:@.+]] to void
+// CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_4:@.+]] to void
+// CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_5:@.+]] to void
+// CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_6:@.+]] to void
+// CHECK: @__kmpc_fork_call(%struct.ident_t* {{.+}}, i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* [[OUTLINED_7:@.+]] to void
+// CHECK: ret void
+
+// CHECK: define internal void [[OUTLINED_2]](
+// CHECK: @_Z3barv
+// CHECK: ret void
+
+// CHECK: define internal void [[OUTLINED_3]](
+// CHECK: @_Z3barv
+// CHECK: ret void
+
+// CHECK: define internal void [[OUTLINED_4]](
+// CHECK: @_Z3barv
+// CHECK: ret void
+
+// CHECK: define internal void [[OUTLINED_5]](
+// NO-CHECK: call void @__kmpc_for_static_init
+// NO-CHECK: call void @__kmpc_for_static_fini
+// CHECK: ret void
+
+// CHECK: define internal void [[OUTLINED_6]](
+// CHECK: call void @__kmpc_for_static_init
+// CHECK: call void @__kmpc_for_static_fini
+// CHECK: ret void
+
+// CHECK: define internal void [[OUTLINED_7]](
+// NO-CHECK: call void @__kmpc_for_static_init
+// NO-CHECK: call void @__kmpc_for_static_fini
+// CHECK: ret void
+
+#endif

clang/test/OpenMP/metadirective_messages.cpp

@@ -0,0 +1,20 @@
+// RUN: %clang_cc1 -triple=x86_64-pc-linux-gnu -verify -fopenmp -x c++ -std=c++14 -fexceptions -fcxx-exceptions %s
+
+// RUN: %clang_cc1 -triple=x86_64-pc-linux-gnu -verify -fopenmp-simd -x c++ -std=c++14 -fexceptions -fcxx-exceptions %s
+
+void foo() {
+#pragma omp metadirective // expected-error {{expected expression}}
+  ;
+#pragma omp metadirective when() // expected-error {{expected valid context selector in when clause}} expected-error {{expected expression}} expected-warning {{expected identifier or string literal describing a context set; set skipped}} expected-note {{context set options are: 'construct' 'device' 'implementation' 'user'}} expected-note {{the ignored set spans until here}}
+  ;
+#pragma omp metadirective when(device{}) // expected-warning {{expected '=' after the context set name "device"; '=' assumed}} expected-warning {{expected identifier or string literal describing a context selector; selector skipped}} expected-note {{context selector options are: 'kind' 'arch' 'isa'}} expected-note {{the ignored selector spans until here}} expected-error {{expected valid context selector in when clause}} expected-error {{expected expression}}
+  ;
+#pragma omp metadirective when(device{arch(nvptx)}) // expected-error {{missing ':' in when clause}} expected-error {{expected expression}} expected-warning {{expected '=' after the context set name "device"; '=' assumed}}
+  ;
+#pragma omp metadirective when(device{arch(nvptx)}: ) default() // expected-warning {{expected '=' after the context set name "device"; '=' assumed}}
+  ;
+#pragma omp metadirective when(device = {arch(nvptx)} : ) default(xyz) // expected-error {{expected an OpenMP directive}} expected-error {{use of undeclared identifier 'xyz'}}
+  ;
+#pragma omp metadirective when(device = {arch(nvptx)} : parallel default() // expected-error {{expected ',' or ')' in 'when' clause}} expected-error {{expected expression}}
+  ;
+}

clang/tools/libclang/CIndex.cpp

@@ -5582,6 +5582,8 @@ CXString clang_getCursorKindSpelling(enum CXCursorKind Kind) {
     return cxstring::createRef("ModuleImport");
   case CXCursor_OMPCanonicalLoop:
     return cxstring::createRef("OMPCanonicalLoop");
+  case CXCursor_OMPMetaDirective:
+    return cxstring::createRef("OMPMetaDirective");
   case CXCursor_OMPParallelDirective:
     return cxstring::createRef("OMPParallelDirective");
   case CXCursor_OMPSimdDirective:

clang/tools/libclang/CXCursor.cpp

@@ -643,6 +643,9 @@ CXCursor cxcursor::MakeCXCursor(const Stmt *S, const Decl *Parent,
   case Stmt::OMPCanonicalLoopClass:
     K = CXCursor_OMPCanonicalLoop;
     break;
+  case Stmt::OMPMetaDirectiveClass:
+    K = CXCursor_OMPMetaDirective;
+    break;
   case Stmt::OMPParallelDirectiveClass:
     K = CXCursor_OMPParallelDirective;
     break;

flang/lib/Semantics/check-omp-structure.cpp

@@ -1391,6 +1391,7 @@ CHECK_SIMPLE_CLAUSE(Use, OMPC_use)
 CHECK_SIMPLE_CLAUSE(Novariants, OMPC_novariants)
 CHECK_SIMPLE_CLAUSE(Nocontext, OMPC_nocontext)
 CHECK_SIMPLE_CLAUSE(Filter, OMPC_filter)
+CHECK_SIMPLE_CLAUSE(When, OMPC_when)
 
 CHECK_REQ_SCALAR_INT_CLAUSE(Grainsize, OMPC_grainsize)
 CHECK_REQ_SCALAR_INT_CLAUSE(NumTasks, OMPC_num_tasks)

llvm/include/llvm/Frontend/OpenMP/OMP.td

@@ -339,6 +339,7 @@ def OMPC_Filter : Clause<"filter"> {
   let clangClass = "OMPFilterClause";
   let flangClass = "ScalarIntExpr";
 }
+def OMPC_When: Clause<"when"> {}
 
 //===----------------------------------------------------------------------===//
 // Definition of OpenMP directives
@@ -1703,6 +1704,10 @@ def OMP_masked : Directive<"masked"> {
     VersionedClause<OMPC_Filter>
   ];
 }
+def OMP_Metadirective : Directive<"metadirective"> {
+  let allowedClauses = [VersionedClause<OMPC_When>];
+  let allowedOnceClauses = [VersionedClause<OMPC_Default>];
+}
 def OMP_Unknown : Directive<"unknown"> {
   let isDefault = true;
 }