[clang] de59f56 - [OpenMP] Support OpenMP 5.1 attributes

[Aaron Ballman via cfe-commits]

Author: Aaron Ballman
Date: 2021-07-12T06:51:19-04:00
New Revision: de59f564400de1b0fe30ae07f3c800562a025e27

URL: https://github.com/llvm/llvm-project/commit/de59f564400de1b0fe30ae07f3c800562a025e27
DIFF: https://github.com/llvm/llvm-project/commit/de59f564400de1b0fe30ae07f3c800562a025e27.diff

LOG: [OpenMP] Support OpenMP 5.1 attributes

OpenMP 5.1 added support for writing OpenMP directives using [[]]
syntax in addition to using #pragma and this introduces support for the
new syntax.

In OpenMP, the attributes take one of two forms:
[[omp::directive(...)]] or [[omp::sequence(...)]]. A directive
attribute contains an OpenMP directive clause that is identical to the
analogous #pragma syntax. A sequence attribute can contain either
sequence or directive arguments and is used to ensure that the
attributes are processed sequentially for situations where the order of
the attributes matter (remember:
https://eel.is/c++draft/dcl.attr.grammar#4.sentence-4).

The approach taken here is somewhat novel and deserves mention. We
could refactor much of the OpenMP parsing logic to work for either
pragma annotation tokens or for attribute clauses. It would be a fair
amount of effort to share the logic for both, but it's certainly
doable. However, the semantic attribute system is not designed to
handle the arbitrarily complex arguments that OpenMP directives
contain. Adding support to thread the novel parsed information until we
can produce a semantic attribute would be considerably more effort.
What's more, existing OpenMP constructs are not (often) represented as
semantic attributes. So doing this through Attr.td would be a massive
undertaking that would likely only benefit OpenMP and comes with
additional risks. Rather than walk down that path, I am taking
advantage of the fact that the syntax of the directives within the
directive clause is identical to that of the #pragma form. Once the
parser recognizes that we're processing an OpenMP attribute, it caches
all of the directive argument tokens and then replays them as though
the user wrote a pragma. This reuses the same OpenMP parsing and
semantic logic directly, but does come with a risk if the OpenMP
committee decides to purposefully diverge their pragma and attribute
syntaxes. So, despite this being a novel approach that does token
replay, I think it's actually a better approach than trying to do this
through the declarative syntax in Attr.td.

Added: 
    clang/test/OpenMP/allocate_codegen_attr.cpp
    clang/test/OpenMP/assumes_messages_attr.c
    clang/test/OpenMP/critical_codegen_attr.cpp
    clang/test/OpenMP/masked_messages_attr.cpp
    clang/test/OpenMP/openmp_attribute.cpp
    clang/test/OpenMP/openmp_attribute_compat.cpp
    clang/test/OpenMP/openmp_attribute_parsing.cpp
    clang/test/OpenMP/target_map_names_attr.cpp
    clang/test/OpenMP/taskloop_reduction_messages_attr.cpp
    clang/test/OpenMP/teams_distribute_parallel_for_simd_num_teams_messages_attr.cpp
    clang/test/OpenMP/unroll_codegen_unroll_for_attr.cpp

Modified: 
    clang/docs/OpenMPSupport.rst
    clang/include/clang/Basic/DiagnosticGroups.td
    clang/include/clang/Basic/DiagnosticParseKinds.td
    clang/include/clang/Basic/TokenKinds.def
    clang/include/clang/Parse/Parser.h
    clang/lib/Basic/Attributes.cpp
    clang/lib/Parse/ParseCXXInlineMethods.cpp
    clang/lib/Parse/ParseDecl.cpp
    clang/lib/Parse/ParseDeclCXX.cpp
    clang/lib/Parse/ParseOpenMP.cpp
    clang/lib/Parse/ParseStmt.cpp
    clang/lib/Parse/Parser.cpp

Removed: 
    


################################################################################
diff  --git a/clang/docs/OpenMPSupport.rst b/clang/docs/OpenMPSupport.rst
index 3bd1a0c211486..88d3107c21e7a 100644
--- a/clang/docs/OpenMPSupport.rst
+++ b/clang/docs/OpenMPSupport.rst
@@ -268,7 +268,7 @@ want to help with the implementation.
 +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
 | atomic extension             | 'fail' clause on atomic construct                            | :none:`unclaimed`        |                                                                       |
 +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
-| base language                | C++ attribute specifier syntax                               | :part:`worked on`        |                                                                       |
+| base language                | C++ attribute specifier syntax                               | :good:`done`             | D105648                                                               |
 +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+
 | device extension             | 'present' map type modifier                                  | :good:`done`             | D83061, D83062, D84422                                                |
 +------------------------------+--------------------------------------------------------------+--------------------------+-----------------------------------------------------------------------+

diff  --git a/clang/include/clang/Basic/DiagnosticGroups.td b/clang/include/clang/Basic/DiagnosticGroups.td
index f35c105964a38..18097cfcfffe4 100644
--- a/clang/include/clang/Basic/DiagnosticGroups.td
+++ b/clang/include/clang/Basic/DiagnosticGroups.td
@@ -1186,8 +1186,11 @@ def OpenMPClauses : DiagGroup<"openmp-clauses">;
 def OpenMPLoopForm : DiagGroup<"openmp-loop-form">;
 def OpenMPMapping : DiagGroup<"openmp-mapping">;
 def OpenMPTarget : DiagGroup<"openmp-target", [OpenMPMapping]>;
+def OpenMPPre51Compat : DiagGroup<"pre-openmp-51-compat">;
+def OpenMP51Ext : DiagGroup<"openmp-51-extensions">;
 def OpenMP : DiagGroup<"openmp", [
-    SourceUsesOpenMP, OpenMPClauses, OpenMPLoopForm, OpenMPTarget, OpenMPMapping
+    SourceUsesOpenMP, OpenMPClauses, OpenMPLoopForm, OpenMPTarget,
+    OpenMPMapping, OpenMP51Ext
   ]>;
 
 // Backend warnings.

diff  --git a/clang/include/clang/Basic/DiagnosticParseKinds.td b/clang/include/clang/Basic/DiagnosticParseKinds.td
index aecc5859c718e..7e4b0841e06b1 100644
--- a/clang/include/clang/Basic/DiagnosticParseKinds.td
+++ b/clang/include/clang/Basic/DiagnosticParseKinds.td
@@ -1427,6 +1427,15 @@ def err_omp_expected_interop_type : Error<
 def warn_omp_more_one_interop_type
   : Warning<"interop type '%0' cannot be specified more than once">,
     InGroup<OpenMPClauses>;
+def err_expected_sequence_or_directive : Error<
+  "expected an OpenMP 'directive' or 'sequence' attribute argument">;
+def ext_omp_attributes : ExtWarn<
+  "specifying OpenMP directives with [[]] is an OpenMP 5.1 extension">,
+  InGroup<OpenMP51Ext>;
+def warn_omp51_compat_attributes : Warning<
+  "specifying OpenMP directives with [[]] is incompatible with OpenMP "
+  "standards before OpenMP 5.1">,
+  InGroup<OpenMPPre51Compat>, DefaultIgnore;
 
 // Pragma loop support.
 def err_pragma_loop_missing_argument : Error<

diff  --git a/clang/include/clang/Basic/TokenKinds.def b/clang/include/clang/Basic/TokenKinds.def
index fd9909628b3b7..48a664e3494e2 100644
--- a/clang/include/clang/Basic/TokenKinds.def
+++ b/clang/include/clang/Basic/TokenKinds.def
@@ -863,6 +863,13 @@ PRAGMA_ANNOTATION(pragma_ms_pragma)
 PRAGMA_ANNOTATION(pragma_opencl_extension)
 
 // Annotations for OpenMP pragma directives - #pragma omp ...
+// The parser produces this annotation token when it parses an [[omp::*]]
+// attribute. The tokens from the attribute argument list are replayed to the
+// token stream with this leading token (and a trailing pragma_openmp_end) so
+// that the parser can reuse the OpenMP parsing logic but still be able to
+// distinguish between a real pragma and a converted pragma. It is not marked
+// as a PRAGMA_ANNOTATION because it doesn't get generated from a #pragma.
+ANNOTATION(attr_openmp)
 // The lexer produces these so that they only take effect when the parser
 // handles #pragma omp ... directives.
 PRAGMA_ANNOTATION(pragma_openmp)

diff  --git a/clang/include/clang/Parse/Parser.h b/clang/include/clang/Parse/Parser.h
index 605292f08df47..8eb3f9029d9d8 100644
--- a/clang/include/clang/Parse/Parser.h
+++ b/clang/include/clang/Parse/Parser.h
@@ -2772,6 +2772,16 @@ class Parser : public CodeCompletionHandler {
                           IdentifierInfo *ScopeName, SourceLocation ScopeLoc,
                           ParsedAttr::Syntax Syntax);
 
+  void ReplayOpenMPAttributeTokens(CachedTokens &OpenMPTokens) {
+    // If parsing the attributes found an OpenMP directive, emit those tokens
+    // to the parse stream now.
+    if (!OpenMPTokens.empty()) {
+      PP.EnterToken(Tok, /*IsReinject*/ true);
+      PP.EnterTokenStream(OpenMPTokens, /*DisableMacroExpansion*/ true,
+                          /*IsReinject*/ true);
+      ConsumeAnyToken(/*ConsumeCodeCompletionTok*/ true);
+    }
+  }
   void MaybeParseCXX11Attributes(Declarator &D) {
     if (standardAttributesAllowed() && isCXX11AttributeSpecifier()) {
       ParsedAttributesWithRange attrs(AttrFactory);
@@ -2801,8 +2811,18 @@ class Parser : public CodeCompletionHandler {
     return false;
   }
 
-  void ParseCXX11AttributeSpecifier(ParsedAttributes &attrs,
-                                    SourceLocation *EndLoc = nullptr);
+  void ParseOpenMPAttributeArgs(IdentifierInfo *AttrName,
+                                CachedTokens &OpenMPTokens);
+
+  void ParseCXX11AttributeSpecifierInternal(ParsedAttributes &Attrs,
+                                            CachedTokens &OpenMPTokens,
+                                            SourceLocation *EndLoc = nullptr);
+  void ParseCXX11AttributeSpecifier(ParsedAttributes &Attrs,
+                                    SourceLocation *EndLoc = nullptr) {
+    CachedTokens OpenMPTokens;
+    ParseCXX11AttributeSpecifierInternal(Attrs, OpenMPTokens, EndLoc);
+    ReplayOpenMPAttributeTokens(OpenMPTokens);
+  }
   void ParseCXX11Attributes(ParsedAttributesWithRange &attrs,
                             SourceLocation *EndLoc = nullptr);
   /// Parses a C++11 (or C2x)-style attribute argument list. Returns true
@@ -2811,7 +2831,8 @@ class Parser : public CodeCompletionHandler {
                                SourceLocation AttrNameLoc,
                                ParsedAttributes &Attrs, SourceLocation *EndLoc,
                                IdentifierInfo *ScopeName,
-                               SourceLocation ScopeLoc);
+                               SourceLocation ScopeLoc,
+                               CachedTokens &OpenMPTokens);
 
   IdentifierInfo *TryParseCXX11AttributeIdentifier(SourceLocation &Loc);
 

diff  --git a/clang/lib/Basic/Attributes.cpp b/clang/lib/Basic/Attributes.cpp
index ff6dbf870fcf4..62eea9c590825 100644
--- a/clang/lib/Basic/Attributes.cpp
+++ b/clang/lib/Basic/Attributes.cpp
@@ -20,6 +20,14 @@ int clang::hasAttribute(AttrSyntax Syntax, const IdentifierInfo *Scope,
   else if (ScopeName == "_Clang")
     ScopeName = "clang";
 
+  // As a special case, look for the omp::sequence and omp::directive
+  // attributes. We support those, but not through the typical attribute
+  // machinery that goes through TableGen. We support this in all OpenMP modes
+  // so long as double square brackets are enabled.
+  if (LangOpts.OpenMP && LangOpts.DoubleSquareBracketAttributes &&
+      ScopeName == "omp")
+    return (Name == "directive" || Name == "sequence") ? 1 : 0;
+
 #include "clang/Basic/AttrHasAttributeImpl.inc"
 
   return 0;

diff  --git a/clang/lib/Parse/ParseCXXInlineMethods.cpp b/clang/lib/Parse/ParseCXXInlineMethods.cpp
index b0335905b6f88..116724a0d50ba 100644
--- a/clang/lib/Parse/ParseCXXInlineMethods.cpp
+++ b/clang/lib/Parse/ParseCXXInlineMethods.cpp
@@ -778,6 +778,7 @@ void Parser::ParseLexedPragma(LateParsedPragma &LP) {
   ConsumeAnyToken(/*ConsumeCodeCompletionTok=*/true);
   assert(Tok.isAnnotation() && "Expected annotation token.");
   switch (Tok.getKind()) {
+  case tok::annot_attr_openmp:
   case tok::annot_pragma_openmp: {
     AccessSpecifier AS = LP.getAccessSpecifier();
     ParsedAttributesWithRange Attrs(AttrFactory);

diff  --git a/clang/lib/Parse/ParseDecl.cpp b/clang/lib/Parse/ParseDecl.cpp
index 76d031b62e9db..c1f20b2e42a22 100644
--- a/clang/lib/Parse/ParseDecl.cpp
+++ b/clang/lib/Parse/ParseDecl.cpp
@@ -4301,7 +4301,7 @@ void Parser::ParseStructUnionBody(SourceLocation RecordLoc,
       continue;
     }
 
-    if (Tok.is(tok::annot_pragma_openmp)) {
+    if (Tok.isOneOf(tok::annot_pragma_openmp, tok::annot_attr_openmp)) {
       // Result can be ignored, because it must be always empty.
       AccessSpecifier AS = AS_none;
       ParsedAttributesWithRange Attrs(AttrFactory);

diff  --git a/clang/lib/Parse/ParseDeclCXX.cpp b/clang/lib/Parse/ParseDeclCXX.cpp
index 37a5ae1f13770..ca5c013a51fe0 100644
--- a/clang/lib/Parse/ParseDeclCXX.cpp
+++ b/clang/lib/Parse/ParseDeclCXX.cpp
@@ -2667,6 +2667,13 @@ Parser::ParseCXXClassMemberDeclaration(AccessSpecifier AS,
   ParsedAttributesViewWithRange FnAttrs;
   // Optional C++11 attribute-specifier
   MaybeParseCXX11Attributes(attrs);
+
+  // The next token may be an OpenMP pragma annotation token. That would
+  // normally be handled from ParseCXXClassMemberDeclarationWithPragmas, but in
+  // this case, it came from an *attribute* rather than a pragma. Handle it now.
+  if (Tok.is(tok::annot_attr_openmp))
+    return ParseOpenMPDeclarativeDirectiveWithExtDecl(AS, attrs);
+
   // We need to keep these attributes for future diagnostic
   // before they are taken over by declaration specifier.
   FnAttrs.addAll(attrs.begin(), attrs.end());
@@ -3261,6 +3268,7 @@ Parser::DeclGroupPtrTy Parser::ParseCXXClassMemberDeclarationWithPragmas(
     return nullptr;
   }
 
+  case tok::annot_attr_openmp:
   case tok::annot_pragma_openmp:
     return ParseOpenMPDeclarativeDirectiveWithExtDecl(
         AS, AccessAttrs, /*Delayed=*/true, TagType, TagDecl);
@@ -4135,6 +4143,70 @@ IdentifierInfo *Parser::TryParseCXX11AttributeIdentifier(SourceLocation &Loc) {
   }
 }
 
+void Parser::ParseOpenMPAttributeArgs(IdentifierInfo *AttrName,
+                                      CachedTokens &OpenMPTokens) {
+  // Both 'sequence' and 'directive' attributes require arguments, so parse the
+  // open paren for the argument list.
+  BalancedDelimiterTracker T(*this, tok::l_paren);
+  if (T.consumeOpen()) {
+    Diag(Tok, diag::err_expected) << tok::l_paren;
+    return;
+  }
+
+  if (AttrName->isStr("directive")) {
+    // If the attribute is named `directive`, we can consume its argument list
+    // and push the tokens from it into the cached token stream for a new OpenMP
+    // pragma directive.
+    Token OMPBeginTok;
+    OMPBeginTok.startToken();
+    OMPBeginTok.setKind(tok::annot_attr_openmp);
+    OMPBeginTok.setLocation(Tok.getLocation());
+    OpenMPTokens.push_back(OMPBeginTok);
+
+    ConsumeAndStoreUntil(tok::r_paren, OpenMPTokens, /*StopAtSemi=*/false,
+                         /*ConsumeFinalToken*/ false);
+    Token OMPEndTok;
+    OMPEndTok.startToken();
+    OMPEndTok.setKind(tok::annot_pragma_openmp_end);
+    OMPEndTok.setLocation(Tok.getLocation());
+    OpenMPTokens.push_back(OMPEndTok);
+  } else {
+    assert(AttrName->isStr("sequence") &&
+           "Expected either 'directive' or 'sequence'");
+    // If the attribute is named 'sequence', its argument is a list of one or
+    // more OpenMP attributes (either 'omp::directive' or 'omp::sequence',
+    // where the 'omp::' is optional).
+    do {
+      // We expect to see one of the following:
+      //  * An identifier (omp) for the attribute namespace followed by ::
+      //  * An identifier (directive) or an identifier (sequence).
+      SourceLocation IdentLoc;
+      IdentifierInfo *Ident = TryParseCXX11AttributeIdentifier(IdentLoc);
+
+      // If there is an identifier and it is 'omp', a double colon is required
+      // followed by the actual identifier we're after.
+      if (Ident && Ident->isStr("omp") && !ExpectAndConsume(tok::coloncolon))
+        Ident = TryParseCXX11AttributeIdentifier(IdentLoc);
+
+      // If we failed to find an identifier (scoped or otherwise), or we found
+      // an unexpected identifier, diagnose.
+      if (!Ident || (!Ident->isStr("directive") && !Ident->isStr("sequence"))) {
+        Diag(Tok.getLocation(), diag::err_expected_sequence_or_directive);
+        SkipUntil(tok::r_paren, StopBeforeMatch);
+        continue;
+      }
+      // We read an identifier. If the identifier is one of the ones we
+      // expected, we can recurse to parse the args.
+      ParseOpenMPAttributeArgs(Ident, OpenMPTokens);
+
+      // There may be a comma to signal that we expect another directive in the
+      // sequence.
+    } while (TryConsumeToken(tok::comma));
+  }
+  // Parse the closing paren for the argument list.
+  T.consumeClose();
+}
+
 static bool IsBuiltInOrStandardCXX11Attribute(IdentifierInfo *AttrName,
                                               IdentifierInfo *ScopeName) {
   switch (
@@ -4175,7 +4247,8 @@ bool Parser::ParseCXX11AttributeArgs(IdentifierInfo *AttrName,
                                      ParsedAttributes &Attrs,
                                      SourceLocation *EndLoc,
                                      IdentifierInfo *ScopeName,
-                                     SourceLocation ScopeLoc) {
+                                     SourceLocation ScopeLoc,
+                                     CachedTokens &OpenMPTokens) {
   assert(Tok.is(tok::l_paren) && "Not a C++11 attribute argument list");
   SourceLocation LParenLoc = Tok.getLocation();
   const LangOptions &LO = getLangOpts();
@@ -4200,6 +4273,18 @@ bool Parser::ParseCXX11AttributeArgs(IdentifierInfo *AttrName,
     return true;
   }
 
+  if (ScopeName && ScopeName->isStr("omp")) {
+    Diag(AttrNameLoc, getLangOpts().OpenMP >= 51
+                          ? diag::warn_omp51_compat_attributes
+                                    : diag::ext_omp_attributes);
+
+    ParseOpenMPAttributeArgs(AttrName, OpenMPTokens);
+
+    // We claim that an attribute was parsed and added so that one is not
+    // created for us by the caller.
+    return true;
+  }
+
   unsigned NumArgs;
   // Some Clang-scoped attributes have some special parsing behavior.
   if (ScopeName && (ScopeName->isStr("clang") || ScopeName->isStr("_Clang")))
@@ -4259,11 +4344,12 @@ bool Parser::ParseCXX11AttributeArgs(IdentifierInfo *AttrName,
 ///
 /// [C++11] attribute-namespace:
 ///         identifier
-void Parser::ParseCXX11AttributeSpecifier(ParsedAttributes &attrs,
-                                          SourceLocation *endLoc) {
+void Parser::ParseCXX11AttributeSpecifierInternal(ParsedAttributes &Attrs,
+                                                  CachedTokens &OpenMPTokens,
+                                                  SourceLocation *EndLoc) {
   if (Tok.is(tok::kw_alignas)) {
     Diag(Tok.getLocation(), diag::warn_cxx98_compat_alignas);
-    ParseAlignmentSpecifier(attrs, endLoc);
+    ParseAlignmentSpecifier(Attrs, EndLoc);
     return;
   }
 
@@ -4345,11 +4431,11 @@ void Parser::ParseCXX11AttributeSpecifier(ParsedAttributes &attrs,
 
     // Parse attribute arguments
     if (Tok.is(tok::l_paren))
-      AttrParsed = ParseCXX11AttributeArgs(AttrName, AttrLoc, attrs, endLoc,
-                                           ScopeName, ScopeLoc);
+      AttrParsed = ParseCXX11AttributeArgs(AttrName, AttrLoc, Attrs, EndLoc,
+                                           ScopeName, ScopeLoc, OpenMPTokens);
 
     if (!AttrParsed) {
-      attrs.addNew(
+      Attrs.addNew(
           AttrName,
           SourceRange(ScopeLoc.isValid() ? ScopeLoc : AttrLoc, AttrLoc),
           ScopeName, ScopeLoc, nullptr, 0,
@@ -4374,8 +4460,8 @@ void Parser::ParseCXX11AttributeSpecifier(ParsedAttributes &attrs,
     SkipUntil(tok::r_square);
   else if (Tok.is(tok::r_square))
     checkCompoundToken(CloseLoc, tok::r_square, CompoundToken::AttrEnd);
-  if (endLoc)
-    *endLoc = Tok.getLocation();
+  if (EndLoc)
+    *EndLoc = Tok.getLocation();
   if (ExpectAndConsume(tok::r_square))
     SkipUntil(tok::r_square);
 }

diff  --git a/clang/lib/Parse/ParseOpenMP.cpp b/clang/lib/Parse/ParseOpenMP.cpp
index d2152fd117938..18e43c3734acb 100644
--- a/clang/lib/Parse/ParseOpenMP.cpp
+++ b/clang/lib/Parse/ParseOpenMP.cpp
@@ -1857,7 +1857,8 @@ void Parser::ParseOMPEndDeclareTargetDirective(OpenMPDirectiveKind BeginDKind,
 Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirectiveWithExtDecl(
     AccessSpecifier &AS, ParsedAttributesWithRange &Attrs, bool Delayed,
     DeclSpec::TST TagType, Decl *Tag) {
-  assert(Tok.is(tok::annot_pragma_openmp) && "Not an OpenMP directive!");
+  assert(Tok.isOneOf(tok::annot_pragma_openmp, tok::annot_attr_openmp) &&
+         "Not an OpenMP directive!");
   ParsingOpenMPDirectiveRAII DirScope(*this);
   ParenBraceBracketBalancer BalancerRAIIObj(*this);
 
@@ -1875,7 +1876,7 @@ Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirectiveWithExtDecl(
       Toks.push_back(Tok);
       while (Cnt && Tok.isNot(tok::eof)) {
         (void)ConsumeAnyToken();
-        if (Tok.is(tok::annot_pragma_openmp))
+        if (Tok.isOneOf(tok::annot_pragma_openmp, tok::annot_attr_openmp))
           ++Cnt;
         else if (Tok.is(tok::annot_pragma_openmp_end))
           --Cnt;
@@ -2098,7 +2099,7 @@ Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirectiveWithExtDecl(
     ConsumeAnyToken();
 
     DeclGroupPtrTy Ptr;
-    if (Tok.is(tok::annot_pragma_openmp)) {
+    if (Tok.isOneOf(tok::annot_pragma_openmp, tok::annot_attr_openmp)) {
       Ptr = ParseOpenMPDeclarativeDirectiveWithExtDecl(AS, Attrs, Delayed,
                                                        TagType, Tag);
     } else if (Tok.isNot(tok::r_brace) && !isEofOrEom()) {
@@ -2275,7 +2276,8 @@ Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirectiveWithExtDecl(
 ///
 StmtResult
 Parser::ParseOpenMPDeclarativeOrExecutableDirective(ParsedStmtContext StmtCtx) {
-  assert(Tok.is(tok::annot_pragma_openmp) && "Not an OpenMP directive!");
+  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;

diff  --git a/clang/lib/Parse/ParseStmt.cpp b/clang/lib/Parse/ParseStmt.cpp
index 7dd8dd1ff3b71..ebfe048513b1f 100644
--- a/clang/lib/Parse/ParseStmt.cpp
+++ b/clang/lib/Parse/ParseStmt.cpp
@@ -401,7 +401,12 @@ StmtResult Parser::ParseStatementOrDeclarationAfterAttributes(
     return HandlePragmaCaptured();
 
   case tok::annot_pragma_openmp:
+    // Prohibit attributes that are not OpenMP attributes, but only before
+    // processing a #pragma omp clause.
     ProhibitAttributes(Attrs);
+    LLVM_FALLTHROUGH;
+  case tok::annot_attr_openmp:
+    // Do not prohibit attributes if they were OpenMP attributes.
     return ParseOpenMPDeclarativeOrExecutableDirective(StmtCtx);
 
   case tok::annot_pragma_ms_pointers_to_members:

diff  --git a/clang/lib/Parse/Parser.cpp b/clang/lib/Parse/Parser.cpp
index 55b25d20db51c..c81dd03ffaaae 100644
--- a/clang/lib/Parse/Parser.cpp
+++ b/clang/lib/Parse/Parser.cpp
@@ -309,6 +309,7 @@ bool Parser::SkipUntil(ArrayRef<tok::TokenKind> Toks, SkipUntilFlags Flags) {
       return false;
 
     case tok::annot_pragma_openmp:
+    case tok::annot_attr_openmp:
     case tok::annot_pragma_openmp_end:
       // Stop before an OpenMP pragma boundary.
       if (OpenMPDirectiveParsing)
@@ -798,6 +799,7 @@ Parser::ParseExternalDeclaration(ParsedAttributesWithRange &attrs,
   case tok::annot_pragma_opencl_extension:
     HandlePragmaOpenCLExtension();
     return nullptr;
+  case tok::annot_attr_openmp:
   case tok::annot_pragma_openmp: {
     AccessSpecifier AS = AS_none;
     return ParseOpenMPDeclarativeDirectiveWithExtDecl(AS, attrs);

diff  --git a/clang/test/OpenMP/allocate_codegen_attr.cpp b/clang/test/OpenMP/allocate_codegen_attr.cpp
new file mode 100644
index 0000000000000..0ef5c77a2738f
--- /dev/null
+++ b/clang/test/OpenMP/allocate_codegen_attr.cpp
@@ -0,0 +1,112 @@
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=51 -triple x86_64-apple-darwin10.6.0 -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -triple x86_64-apple-darwin10.6.0 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -triple x86_64-apple-darwin10.6.0 -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=51 -triple x86_64-unknown-linux-gnu -emit-llvm -o - %s | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fnoopenmp-use-tls -triple x86_64-unknown-linux-gnu -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fnoopenmp-use-tls -triple x86_64-unknown-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s
+
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=51 -triple x86_64-apple-darwin10.6.0 -emit-llvm -o - %s | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -triple x86_64-apple-darwin10.6.0 -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -triple x86_64-apple-darwin10.6.0 -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=51 -triple x86_64-unknown-linux-gnu -emit-llvm -o - %s | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -fnoopenmp-use-tls -triple x86_64-unknown-linux-gnu -x c++ -std=c++11 -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -fnoopenmp-use-tls -triple x86_64-unknown-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+enum omp_allocator_handle_t {
+  omp_null_allocator = 0,
+  omp_default_mem_alloc = 1,
+  omp_large_cap_mem_alloc = 2,
+  omp_const_mem_alloc = 3,
+  omp_high_bw_mem_alloc = 4,
+  omp_low_lat_mem_alloc = 5,
+  omp_cgroup_mem_alloc = 6,
+  omp_pteam_mem_alloc = 7,
+  omp_thread_mem_alloc = 8,
+  KMP_ALLOCATOR_MAX_HANDLE = __UINTPTR_MAX__
+};
+
+struct St{
+ int a;
+};
+
+struct St1{
+ int a;
+ static int b;
+ [[omp::directive(allocate(b) allocator(omp_default_mem_alloc))]];
+} d;
+
+int a, b, c;
+[[omp::directive(allocate(a) allocator(omp_large_cap_mem_alloc)),
+       directive(allocate(b) allocator(omp_const_mem_alloc)),
+       directive(allocate(d, c) allocator(omp_high_bw_mem_alloc))]];
+
+template <class T>
+struct ST {
+  static T m;
+  [[omp::directive(allocate(m) allocator(omp_low_lat_mem_alloc))]];
+};
+
+template <class T> T foo() {
+  T v;
+  [[omp::directive(allocate(v) allocator(omp_cgroup_mem_alloc))]];
+  v = ST<T>::m;
+  return v;
+}
+
+namespace ns{
+  int a;
+}
+[[omp::directive(allocate(ns::a) allocator(omp_pteam_mem_alloc))]];
+
+// CHECK-NOT:  call {{.+}} {{__kmpc_alloc|__kmpc_free}}
+
+// CHECK-LABEL: @main
+int main () {
+  static int a;
+  [[omp::directive(allocate(a) allocator(omp_thread_mem_alloc))]];
+  a=2;
+  // CHECK-NOT:  {{__kmpc_alloc|__kmpc_free}}
+  // CHECK:      alloca double,
+  // CHECK-NOT:  {{__kmpc_alloc|__kmpc_free}}
+  double b = 3;
+  [[omp::directive(allocate(b))]];
+  return (foo<int>());
+}
+
+// CHECK: define {{.*}}i32 @{{.+}}foo{{.+}}()
+// CHECK:      [[GTID:%.+]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @{{.+}})
+// CHECK-NEXT: [[V_VOID_ADDR:%.+]] = call i8* @__kmpc_alloc(i32 [[GTID]], i64 4, i8* inttoptr (i64 6 to i8*))
+// CHECK-NEXT: [[V_ADDR:%.+]] = bitcast i8* [[V_VOID_ADDR]] to i32*
+// CHECK-NOT:  {{__kmpc_alloc|__kmpc_free}}
+// CHECK:      store i32 %{{.+}}, i32* [[V_ADDR]],
+// CHECK-NEXT: [[V_VAL:%.+]] = load i32, i32* [[V_ADDR]],
+// CHECK-NEXT: [[V_VOID_ADDR:%.+]] = bitcast i32* [[V_ADDR]] to i8*
+// CHECK-NEXT: call void @__kmpc_free(i32 [[GTID]], i8* [[V_VOID_ADDR]], i8* inttoptr (i64 6 to i8*))
+// CHECK-NOT:  {{__kmpc_alloc|__kmpc_free}}
+// CHECK:      ret i32 [[V_VAL]]
+
+// CHECK-NOT:  call {{.+}} {{__kmpc_alloc|__kmpc_free}}
+extern template int ST<int>::m;
+
+// CHECK: define{{.*}} void @{{.+}}bar{{.+}}(i32 %{{.+}}, float* {{.+}})
+void bar(int a, float &z) {
+// CHECK: [[A_VOID_PTR:%.+]] = call i8* @__kmpc_alloc(i32 [[GTID:%.+]], i64 4, i8* inttoptr (i64 1 to i8*))
+// CHECK: [[A_ADDR:%.+]] = bitcast i8* [[A_VOID_PTR]] to i32*
+// CHECK: store i32 %{{.+}}, i32* [[A_ADDR]],
+// CHECK: [[Z_VOID_PTR:%.+]] = call i8* @__kmpc_alloc(i32 [[GTID]], i64 8, i8* inttoptr (i64 1 to i8*))
+// CHECK: [[Z_ADDR:%.+]] = bitcast i8* [[Z_VOID_PTR]] to float**
+// CHECK: store float* %{{.+}}, float** [[Z_ADDR]],
+[[omp::directive(allocate(a,z) allocator(omp_default_mem_alloc))]];
+// CHECK-NEXT: [[Z_VOID_PTR:%.+]] = bitcast float** [[Z_ADDR]] to i8*
+// CHECK: call void @__kmpc_free(i32 [[GTID]], i8* [[Z_VOID_PTR]], i8* inttoptr (i64 1 to i8*))
+// CHECK-NEXT: [[A_VOID_PTR:%.+]] = bitcast i32* [[A_ADDR]] to i8*
+// CHECK: call void @__kmpc_free(i32 [[GTID]], i8* [[A_VOID_PTR]], i8* inttoptr (i64 1 to i8*))
+// CHECK: ret void
+}
+#endif
+

diff  --git a/clang/test/OpenMP/assumes_messages_attr.c b/clang/test/OpenMP/assumes_messages_attr.c
new file mode 100644
index 0000000000000..bd06e85cbf3b5
--- /dev/null
+++ b/clang/test/OpenMP/assumes_messages_attr.c
@@ -0,0 +1,57 @@
+// RUN: %clang_cc1 -triple=x86_64-pc-win32 -verify -fopenmp -fopenmp-version=51 -std=c99 -fms-extensions -fdouble-square-bracket-attributes -Wno-pragma-pack %s
+// RUN: %clang_cc1 -triple=x86_64-pc-win32 -verify -fopenmp-simd -fopenmp-version=51 -std=c99 -fms-extensions -fdouble-square-bracket-attributes -Wno-pragma-pack %s
+
+[[omp::directive(assumes)]]; // expected-error {{expected at least one 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism' clause for '#pragma omp assumes'}}
+[[omp::directive(begin)]]; // expected-error {{expected an OpenMP directive}}
+[[omp::directive(begin assumes)]]; // expected-error {{expected at least one 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism' clause for '#pragma omp begin assumes'}}
+[[omp::directive(end assumes)]];
+
+[[omp::directive(assumes foobar)]]; // expected-warning {{valid assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}}
+[[omp::directive(begin assumes foobar)]]; // expected-warning {{valid begin assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}}
+[[omp::directive(end assumes)]];
+
+[[omp::directive(begin assumes foobar(foo 2 baz))]]; // expected-warning {{valid begin assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; tokens will be ignored}} expected-note {{the ignored tokens spans until here}}
+[[omp::directive(assumes foobar(foo 2 baz))]]; // expected-warning {{valid assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; tokens will be ignored}} expected-note {{the ignored tokens spans until here}}
+[[omp::directive(end assumes)]];
+
+[[omp::directive(assumes no_openmp(1))]]; // expected-warning {{'no_openmp' clause should not be followed by arguments; tokens will be ignored}} expected-note {{the ignored tokens spans until here}}
+[[omp::directive(begin assumes no_openmp(1 2 3))]]; // expected-warning {{'no_openmp' clause should not be followed by arguments; tokens will be ignored}} expected-note {{the ignored tokens spans until here}}
+[[omp::directive(end assumes no_openmp(1))]];
+
+[[omp::directive(assumes foobar no_openmp bazbaz)]]; // expected-warning {{valid assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}} expected-warning {{valid assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}}
+[[omp::directive(begin assumes foobar no_openmp bazbaz)]]; // expected-warning {{valid begin assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}} expected-warning {{valid begin assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}}
+[[omp::directive(end assumes)]];
+
+[[omp::directive(begin assumes foobar(foo 2 baz) no_openmp bazbaz(foo 2 baz))]]; // expected-warning {{valid begin assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; tokens will be ignored}} expected-warning {{valid begin assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; tokens will be ignored}} expected-note {{the ignored tokens spans until here}} expected-note {{the ignored tokens spans until here}}
+[[omp::directive(assumes foobar(foo 2 baz) no_openmp bazbaz(foo 2 baz))]]; // expected-warning {{valid assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; tokens will be ignored}} expected-warning {{valid assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; tokens will be ignored}} expected-note {{the ignored tokens spans until here}} expected-note {{the ignored tokens spans until here}}
+[[omp::directive(end assumes)]];
+
+[[omp::directive(assumes no_openmp foobar no_openmp)]]; // expected-warning {{valid assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}}
+[[omp::directive(begin assumes no_openmp foobar no_openmp)]]; // expected-warning {{valid begin assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}}
+[[omp::directive(end assumes)]];
+
+[[omp::directive(assumes holds(1, 2 3))]];
+[[omp::directive(begin assumes holds(1, 2 3))]];
+[[omp::directive(end assumes)]];
+
+[[omp::directive(assumes absent(1, 2 3))]];
+[[omp::directive(begin assumes absent(1, 2 3))]];
+[[omp::directive(end assumes)]];
+
+[[omp::directive(assumes contains(1, 2 3))]];
+[[omp::directive(begin assumes contains(1, 2 3))]];
+[[omp::directive(end assumes)]];
+
+[[omp::directive(assumes ext)]]; // expected-warning {{valid assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}}
+[[omp::directive(begin assumes ext)]]; // expected-warning {{valid begin assumes clauses start with 'ext_', 'absent', 'contains', 'holds', 'no_openmp', 'no_openmp_routines', 'no_parallelism'; token will be ignored}}
+[[omp::directive(end assumes)]];
+
+[[omp::directive(assumes ext_123(not allowed))]]; // expected-warning {{'ext_123' clause should not be followed by arguments; tokens will be ignored}} expected-note {{the ignored tokens spans until here}}
+[[omp::directive(begin assumes ext_123(not allowed))]]; // expected-warning {{'ext_123' clause should not be followed by arguments; tokens will be ignored}} expected-note {{the ignored tokens spans until here}}
+[[omp::directive(end assumes)]];
+
+[[omp::directive(end assumes)]]; // expected-error {{'#pragma omp end assumes' with no matching '#pragma omp begin assumes'}}
+
+// TODO: we should emit a warning at least.
+[[omp::directive(begin assumes ext_abc)]];
+

diff  --git a/clang/test/OpenMP/critical_codegen_attr.cpp b/clang/test/OpenMP/critical_codegen_attr.cpp
new file mode 100644
index 0000000000000..e3890ff6058cc
--- /dev/null
+++ b/clang/test/OpenMP/critical_codegen_attr.cpp
@@ -0,0 +1,130 @@
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=51 -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck %s --check-prefixes=ALL,NORMAL
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -std=c++11 -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefixes=ALL,NORMAL
+// RUN: %clang_cc1 -verify -triple x86_64-apple-darwin10 -fopenmp -fopenmp-version=51 -fexceptions -fcxx-exceptions -debug-info-kind=line-tables-only -emit-llvm %s -o - | FileCheck %s --check-prefix=TERM_DEBUG
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=51 -fopenmp-enable-irbuilder -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck %s --check-prefixes=ALL,IRBUILDER
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-enable-irbuilder -std=c++11 -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fopenmp-enable-irbuilder -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefixes=ALL,IRBUILDER
+
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=51 -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -std=c++11 -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp-simd -fopenmp-version=51 -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// RUN: %clang_cc1 -verify -triple x86_64-apple-darwin10 -fopenmp-simd -fopenmp-version=51 -fexceptions -fcxx-exceptions -debug-info-kind=line-tables-only -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s
+// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+// ALL:       [[IDENT_T_TY:%.+]] = type { i32, i32, i32, i32, i8* }
+// ALL:       [[UNNAMED_LOCK:@.+]] = common global [8 x i32] zeroinitializer
+// ALL:       [[THE_NAME_LOCK:@.+]] = common global [8 x i32] zeroinitializer
+// ALL:       [[THE_NAME_LOCK1:@.+]] = common global [8 x i32] zeroinitializer
+
+// ALL:       define {{.*}}void [[FOO:@.+]]()
+
+void foo() { extern void mayThrow(); mayThrow(); }
+
+// ALL-LABEL: @main
+// TERM_DEBUG-LABEL: @main
+int main() {
+  // ALL:       [[A_ADDR:%.+]] = alloca i8
+  char a;
+
+// ALL:       			[[GTID:%.+]] = call {{.*}}i32 @__kmpc_global_thread_num([[IDENT_T_TY]]* [[DEFAULT_LOC:@.+]])
+// ALL:       			call {{.*}}void @__kmpc_critical([[IDENT_T_TY]]* [[DEFAULT_LOC]], i32 [[GTID]], [8 x i32]* [[UNNAMED_LOCK]])
+// ALL-NEXT:  			store i8 2, i8* [[A_ADDR]]
+// ALL-NEXT:  			call {{.*}}void @__kmpc_end_critical([[IDENT_T_TY]]* [[DEFAULT_LOC]], i32 [[GTID]], [8 x i32]* [[UNNAMED_LOCK]])
+  [[omp::directive(critical)]]
+  a = 2;
+// IRBUILDER:       [[GTID:%.+]] = call {{.*}}i32 @__kmpc_global_thread_num([[IDENT_T_TY]]* [[DEFAULT_LOC:@.+]])
+// ALL:       			call {{.*}}void @__kmpc_critical([[IDENT_T_TY]]* [[DEFAULT_LOC]], i32 [[GTID]], [8 x i32]* [[THE_NAME_LOCK]])
+// IRBUILDER-NEXT:	call {{.*}}void [[FOO]]()
+// NORMAL-NEXT:  		invoke {{.*}}void [[FOO]]()
+// ALL:      				call {{.*}}void @__kmpc_end_critical([[IDENT_T_TY]]* [[DEFAULT_LOC]], i32 [[GTID]], [8 x i32]* [[THE_NAME_LOCK]])
+  [[omp::directive(critical(the_name))]]
+  foo();
+// IRBUILDER:   		[[GTID:%.+]] = call {{.*}}i32 @__kmpc_global_thread_num([[IDENT_T_TY]]* [[DEFAULT_LOC:@.+]])
+// ALL: 	      		call {{.*}}void @__kmpc_critical_with_hint([[IDENT_T_TY]]* [[DEFAULT_LOC]], i32 [[GTID]], [8 x i32]* [[THE_NAME_LOCK1]], i{{64|32}} 23)
+// IRBUILDER-NEXT:	call {{.*}}void [[FOO]]()
+// NORMAL-NEXT:		  invoke {{.*}}void [[FOO]]()
+// ALL:		       		call {{.*}}void @__kmpc_end_critical([[IDENT_T_TY]]* [[DEFAULT_LOC]], i32 [[GTID]], [8 x i32]* [[THE_NAME_LOCK1]])
+  [[omp::directive(critical(the_name1) hint(23))]]
+  foo();
+  // IRBUILDER:   		[[GTID:%.+]] = call {{.*}}i32 @__kmpc_global_thread_num([[IDENT_T_TY]]* [[DEFAULT_LOC:@.+]])
+  // ALL:       call {{.*}}void @__kmpc_critical([[IDENT_T_TY]]* [[DEFAULT_LOC]], i32 [[GTID]], [8 x i32]* [[THE_NAME_LOCK]])
+  // ALL:       br label
+  // ALL-NOT:   call {{.*}}void @__kmpc_end_critical(
+  // ALL:       br label
+  // ALL-NOT:   call {{.*}}void @__kmpc_end_critical(
+  // NORMAL:       br label
+  if (a)
+    [[omp::directive(critical(the_name))]]
+    while (1)
+      ;
+  // ALL:  call {{.*}}void [[FOO]]()
+  foo();
+  // ALL-NOT:   call void @__kmpc_critical
+  // ALL-NOT:   call void @__kmpc_end_critical
+  return a;
+}
+
+// ALL-LABEL:        lambda_critical
+// TERM_DEBUG-LABEL: lambda_critical
+void lambda_critical(int a, int b) {
+  auto l = [=]() {
+    [[omp::directive(critical)]]
+    {
+      // ALL: call void @__kmpc_critical(
+      int c = a + b;
+    }
+  };
+
+  l();
+
+  auto l1 = [=]() {
+    [[omp::sequence(directive(parallel), directive(critical))]]
+    {
+      // ALL: call void @__kmpc_critical(
+      int c = a + b;
+    }
+  };
+
+  l1();
+}
+
+struct S {
+  int a;
+};
+// ALL-LABEL: critical_ref
+void critical_ref(S &s) {
+  // ALL: [[S_ADDR:%.+]] = alloca %struct.S*,
+  // ALL: [[S_REF:%.+]] = load %struct.S*, %struct.S** [[S_ADDR]],
+  // ALL: [[S_A_REF:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[S_REF]], i32 0, i32 0
+  ++s.a;
+  // ALL: call void @__kmpc_critical(
+  [[omp::directive(critical)]]
+  // ALL: [[S_REF:%.+]] = load %struct.S*, %struct.S** [[S_ADDR]],
+  // ALL: [[S_A_REF:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[S_REF]], i32 0, i32 0
+  ++s.a;
+  // ALL: call void @__kmpc_end_critical(
+}
+
+// ALL-LABEL:      parallel_critical
+// TERM_DEBUG-LABEL: parallel_critical
+void parallel_critical() {
+  [[omp::sequence(directive(parallel), directive(critical))]]
+  // TERM_DEBUG-NOT: __kmpc_global_thread_num
+  // TERM_DEBUG:     call void @__kmpc_critical({{.+}}), !dbg [[DBG_LOC_START:![0-9]+]]
+  // TERM_DEBUG:     invoke void {{.*}}foo{{.*}}()
+  // TERM_DEBUG:     unwind label %[[TERM_LPAD:.+]],
+  // TERM_DEBUG-NOT: __kmpc_global_thread_num
+  // TERM_DEBUG:     call void @__kmpc_end_critical({{.+}}), !dbg [[DBG_LOC_END:![0-9]+]]
+  // TERM_DEBUG:     [[TERM_LPAD]]
+  // TERM_DEBUG:     call void @__clang_call_terminate
+  // TERM_DEBUG:     unreachable
+  foo();
+}
+// TERM_DEBUG-DAG: [[DBG_LOC_START]] = !DILocation(line: [[@LINE-12]],
+// TERM_DEBUG-DAG: [[DBG_LOC_END]] = !DILocation(line: [[@LINE-3]],
+#endif
+

diff  --git a/clang/test/OpenMP/masked_messages_attr.cpp b/clang/test/OpenMP/masked_messages_attr.cpp
new file mode 100644
index 0000000000000..37252930b1f0f
--- /dev/null
+++ b/clang/test/OpenMP/masked_messages_attr.cpp
@@ -0,0 +1,86 @@
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=51 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=51 %s -Wuninitialized
+
+void xxx(int argc) {
+  int x; // expected-note {{initialize the variable 'x' to silence this warning}}
+  [[omp::directive(masked)]]
+  argc = x; // expected-warning {{variable 'x' is uninitialized when used here}}
+}
+
+void yyy(int argc) {
+  int x; // expected-note {{initialize the variable 'x' to silence this warning}}
+  [[omp::directive(masked filter(1))]]
+  argc = x; // expected-warning {{variable 'x' is uninitialized when used here}}
+}
+
+int foo();
+
+int main() {
+  [[omp::directive(masked)]]
+  ;
+  [[omp::directive(masked filter(1) filter(2))]] // expected-error {{directive '#pragma omp masked' cannot contain more than one 'filter' clause}}
+  ;
+  int x,y,z;
+  [[omp::directive(masked filter(x) filter(y) filter(z))]] // expected-error 2 {{directive '#pragma omp masked' cannot contain more than one 'filter' clause}}
+  ;
+  [[omp::directive(masked nowait)]] // expected-error {{unexpected OpenMP clause 'nowait' in directive '#pragma omp masked'}}
+  [[omp::directive(masked unknown)]] // expected-warning {{extra tokens at the end of '#pragma omp masked' are ignored}}
+  foo();
+  {
+	[[omp::directive(masked)]]
+  } // expected-error {{expected statement}}
+  {
+	[[omp::directive(masked filter(2))]]
+  } // expected-error {{expected statement}}
+  [[omp::directive(for)]]
+  for (int i = 0; i < 10; ++i) {
+    foo();
+    [[omp::directive(masked filter(1))]] // expected-error {{region cannot be closely nested inside 'for' region}}
+    foo();
+  }
+  [[omp::directive(sections)]]
+  {
+    foo();
+    [[omp::directive(masked)]] // expected-error {{region cannot be closely nested inside 'sections' region}}
+    foo();
+  }
+  [[omp::directive(single)]]
+  for (int i = 0; i < 10; ++i) {
+    foo();
+    [[omp::directive(masked allocate(i))]] // expected-error {{region cannot be closely nested inside 'single' region}} expected-error {{unexpected OpenMP clause 'allocate' in directive '#pragma omp masked'}}
+    foo();
+  }
+  [[omp::directive(masked)]]
+  for (int i = 0; i < 10; ++i) {
+    foo();
+    [[omp::directive(masked)]]
+    foo();
+  }
+  [[omp::directive(for ordered)]]
+  for (int i = 0; i < 10; ++i)
+  [[omp::directive(masked)]] // expected-error {{region cannot be closely nested inside 'for' region}}
+  {
+    foo();
+  }
+
+  return 0;
+}
+
+int foo() {
+  L1: // expected-note {{jump exits scope of OpenMP structured block}}
+    foo();
+  [[omp::directive(masked filter(0))]]
+  {
+    foo();
+    goto L1; // expected-error {{cannot jump from this goto statement to its label}}
+  }
+  goto L2; // expected-error {{cannot jump from this goto statement to its label}}
+  [[omp::directive(masked filter(-2))]]
+  { // expected-note {{jump bypasses OpenMP structured block}}
+    L2:
+    foo();
+  }
+
+  return 0;
+}
+

diff  --git a/clang/test/OpenMP/openmp_attribute.cpp b/clang/test/OpenMP/openmp_attribute.cpp
new file mode 100644
index 0000000000000..5b6f64e6a2073
--- /dev/null
+++ b/clang/test/OpenMP/openmp_attribute.cpp
@@ -0,0 +1,41 @@
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fsyntax-only -verify -DSUPPORTED=1 %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fsyntax-only -verify -DSUPPORTED=1 -x c -std=c2x %s
+// RUN: %clang_cc1 -fsyntax-only -verify -DSUPPORTED=0 %s
+// RUN: %clang_cc1 -fsyntax-only -verify -DSUPPORTED=0 -x c -std=c2x %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fsyntax-only -verify -DSUPPORTED=1 %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fsyntax-only -verify -DSUPPORTED=1 -x c -std=c2x %s
+// expected-no-diagnostics
+
+#ifndef SUPPORTED
+#error "Someone messed up a RUN line"
+#endif
+
+#ifdef __cplusplus
+#if __has_cpp_attribute(omp::sequence) != SUPPORTED
+#error "No idea what you're talking about"
+#endif
+
+#if __has_cpp_attribute(omp::directive) != SUPPORTED
+#error "No idea what you're talking about"
+#endif
+
+#if __has_cpp_attribute(omp::totally_bogus)
+#error "No idea what you're talking about"
+#endif
+
+#else // __cplusplus
+
+#if __has_c_attribute(omp::sequence) != SUPPORTED
+#error "No idea what you're talking about"
+#endif
+
+#if __has_c_attribute(omp::directive) != SUPPORTED
+#error "No idea what you're talking about"
+#endif
+
+#if __has_c_attribute(omp::totally_bogus)
+#error "No idea what you're talking about"
+#endif
+
+#endif
+

diff  --git a/clang/test/OpenMP/openmp_attribute_compat.cpp b/clang/test/OpenMP/openmp_attribute_compat.cpp
new file mode 100644
index 0000000000000..5f51e30102de1
--- /dev/null
+++ b/clang/test/OpenMP/openmp_attribute_compat.cpp
@@ -0,0 +1,19 @@
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fsyntax-only -verify=pre -Wpre-openmp-51-compat %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fsyntax-only -verify=off %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fsyntax-only -verify=ext -Wopenmp %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fsyntax-only -verify=ext -Wopenmp-51-extensions %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fsyntax-only -verify=ext -Wall %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fsyntax-only -verify=ext %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fsyntax-only -verify=off -Wno-openmp %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fsyntax-only -verify=off -Wno-openmp-51-extensions %s
+
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fsyntax-only -verify=pre -Wpre-openmp-51-compat -x c -fdouble-square-bracket-attributes %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -fsyntax-only -verify=off -x c -std=c2x %s
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=50 -fsyntax-only -verify=ext -Wopenmp -x c -std=c2x %s
+
+// off-no-diagnostics
+
+int x;
+[[omp::directive(threadprivate(x))]]; // pre-warning {{specifying OpenMP directives with [[]] is incompatible with OpenMP standards before OpenMP 5.1}} \
+                                      // ext-warning {{specifying OpenMP directives with [[]] is an OpenMP 5.1 extension}}
+

diff  --git a/clang/test/OpenMP/openmp_attribute_parsing.cpp b/clang/test/OpenMP/openmp_attribute_parsing.cpp
new file mode 100644
index 0000000000000..4d2c84cd323a5
--- /dev/null
+++ b/clang/test/OpenMP/openmp_attribute_parsing.cpp
@@ -0,0 +1,77 @@
+// RUN: %clang_cc1 -std=c++17 -fopenmp -fopenmp-version=51 -fsyntax-only -verify %s
+
+// This file tests the custom parsing logic for the OpenMP 5.1 attribute
+// syntax. It does not test actual OpenMP directive syntax, just the attribute
+// parsing bits.
+
+// FIXME: the diagnostic here is a bit unsatisfying. We handle the custom omp
+// attribute parsing logic when parsing the attribute argument list, and we
+// only process an attribute argument list when we see an open paren after the
+// attribute name. So this means we never hit the omp-specific parsing and
+// instead handle this through the usual Sema attribute handling in
+// SemaDeclAttr.cpp, which diagnoses this as an unknown attribute.
+[[omp::directive]]; // expected-warning {{unknown attribute 'directive' ignored}}
+[[omp::sequence]]; // expected-warning {{unknown attribute 'sequence' ignored}}
+[[omp::unknown]]; // expected-warning {{unknown attribute 'unknown' ignored}}
+
+[[omp::directive()]]; // expected-error {{expected an OpenMP directive}}
+[[omp::sequence()]]; // expected-error {{expected an OpenMP 'directive' or 'sequence' attribute argument}}
+
+// Both sequence and directive require an argument list, test that we diagnose
+// when the inner directive or sequence is missing its argument list.
+[[omp::sequence(directive)]]; // expected-error {{expected '('}}
+[[omp::sequence(sequence)]]; // expected-error {{expected '('}}
+[[omp::sequence(omp::directive)]]; // expected-error {{expected '('}}
+[[omp::sequence(omp::sequence)]]; // expected-error {{expected '('}}
+
+// All of the diagnostics here come from the inner sequence and directive not
+// being given an argument, but this tests that we can parse either with or
+// without the 'omp::'.
+[[omp::sequence(directive(), sequence())]]; // expected-error {{expected an OpenMP directive}} expected-error {{expected an OpenMP 'directive' or 'sequence' attribute argument}}
+[[omp::sequence(omp::directive(), sequence())]]; // expected-error {{expected an OpenMP directive}} expected-error {{expected an OpenMP 'directive' or 'sequence' attribute argument}}
+[[omp::sequence(directive(), omp::sequence())]]; // expected-error {{expected an OpenMP directive}} expected-error {{expected an OpenMP 'directive' or 'sequence' attribute argument}}
+[[omp::sequence(omp::directive(), omp::sequence())]]; // expected-error {{expected an OpenMP directive}} expected-error {{expected an OpenMP 'directive' or 'sequence' attribute argument}}
+
+// Test that we properly diagnose missing parens within the inner arguments of
+// a sequence attribute.
+[[omp::sequence( // expected-note {{to match this '('}}
+  directive(
+)]]; // expected-error {{expected ')'}} expected-error {{expected an OpenMP directive}}
+[[omp::sequence( // expected-note {{to match this '('}}
+  sequence(
+)]]; // expected-error {{expected ')'}} expected-error {{expected an OpenMP 'directive' or 'sequence' attribute argument}}
+
+// Test that we properly handle the using attribute syntax.
+[[using omp: directive()]]; // expected-error {{expected an OpenMP directive}}
+[[using omp: sequence()]]; // expected-error {{expected an OpenMP 'directive' or 'sequence' attribute argument}}
+[[using omp: sequence(omp::directive())]]; // expected-error {{expected an OpenMP directive}}
+[[using omp: sequence(directive())]]; // expected-error {{expected an OpenMP directive}}
+
+// Test that we give a sensible error on an unknown attribute in the omp
+// namespace that has an argument list.
+[[omp::unknown()]]; // expected-warning {{unknown attribute 'unknown' ignored}}
+[[using omp: unknown()]]; // expected-warning {{unknown attribute 'unknown' ignored}}
+
+// Test that unknown arguments to the omp::sequence are rejected, regardless of
+// what level they're at.
+[[omp::sequence(unknown)]]; // expected-error {{expected an OpenMP 'directive' or 'sequence' attribute argument}}
+[[omp::sequence(sequence(unknown))]]; // expected-error {{expected an OpenMP 'directive' or 'sequence' attribute argument}}
+[[omp::sequence(omp::unknown)]]; // expected-error {{expected an OpenMP 'directive' or 'sequence' attribute argument}}
+[[omp::sequence(sequence(omp::unknown))]]; // expected-error {{expected an OpenMP 'directive' or 'sequence' attribute argument}}
+
+// FIXME: combining non-openmp attributes with openmp attributes has surprising
+// results due to the replay of tokens. We properly parse the non-openmp
+// attributes, but we also replay the OpenMP tokens. The attributes then get
+// passed to the OpenMP parsing functions and it does not attach the attribute
+// to the declaration statement AST node as you might expect. This means that
+// the expected diagnostics are not issued. Thankfully, due to the positioning
+// of OpenMP attributes and what they appertain to, this should not be a
+// frequent issue (hopefully).
+int x;
+[[deprecated, omp::directive(threadprivate(x))]] int y; // FIXME-expected-note {{'y' has been explicitly marked deprecated here}}
+[[omp::directive(threadprivate(x)), deprecated]] int z; // FIXME-expected-note {{'z' has been explicitly marked deprecated here}}
+void test() {
+  x = 1;
+  y = 1; // FIXME-expected-warning {{warning: 'y' is deprecated}}
+  z = 1; // FIXME-expected-warning {{warning: 'z' is deprecated}}
+}

diff  --git a/clang/test/OpenMP/target_map_names_attr.cpp b/clang/test/OpenMP/target_map_names_attr.cpp
new file mode 100644
index 0000000000000..c45e7c8b996bb
--- /dev/null
+++ b/clang/test/OpenMP/target_map_names_attr.cpp
@@ -0,0 +1,215 @@
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -debug-info-kind=limited -emit-llvm %s -o - | FileCheck %s --check-prefix DEBUG
+// RUN: %clang_cc1 -fopenmp -fopenmp-version=51 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK
+#ifndef HEADER
+#define HEADER
+
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";d;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";i;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";i[1:23];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";p;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";p[1:24];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.i;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.s.f;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.p[:22];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.ps->s.i;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.ps->ps;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.ps->ps->ps;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.ps->ps->s.f[:22];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->i;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->s.f;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->p[:22];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->ps->s.i;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->ps->ps;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->ps->ps->ps;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->ps->ps->s.f[:22];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.f[:22];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.p[:33];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";ps->p[:33];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.s;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+
+struct S1 {
+    int i;
+    float f[50];
+};
+
+struct S2 {
+    int i;
+    float f[50];
+    S1 s;
+    double *p;
+    struct S2 *ps;
+};
+
+void foo() {
+  double d;
+  int i[100];
+  float *p;
+
+  S2 s;
+  S2 *ps;
+
+  [[omp::directive(target map(d))]]
+  { }
+  [[omp::directive(target map(i))]]
+  { }
+  [[omp::directive(target map(i[1:23]))]]
+  { }
+  [[omp::directive(target map(p))]]
+  { }
+  [[omp::directive(target map(p[1:24]))]]
+  { }
+  [[omp::directive(target map(s))]]
+  { }
+  [[omp::directive(target map(s.i))]]
+  { }
+  [[omp::directive(target map(s.s.f))]]
+  { }
+  [[omp::directive(target map(s.p))]]
+  { }
+  [[omp::directive(target map(to: s.p[:22]))]]
+  { }
+  [[omp::directive(target map(s.ps))]]
+  { }
+  [[omp::directive(target map(from: s.ps->s.i))]]
+  { }
+  [[omp::directive(target map(to: s.ps->ps))]]
+  { }
+  [[omp::directive(target map(s.ps->ps->ps))]]
+  { }
+  [[omp::directive(target map(to: s.ps->ps->s.f[:22]))]]
+  { }
+  [[omp::directive(target map(ps))]]
+  { }
+  [[omp::directive(target map(ps->i))]]
+  { }
+  [[omp::directive(target map(ps->s.f))]]
+  { }
+  [[omp::directive(target map(from: ps->p))]]
+  { }
+  [[omp::directive(target map(to: ps->p[:22]))]]
+  { }
+  [[omp::directive(target map(ps->ps))]]
+  { }
+  [[omp::directive(target map(from: ps->ps->s.i))]]
+  { }
+  [[omp::directive(target map(from: ps->ps->ps))]]
+  { }
+  [[omp::directive(target map(ps->ps->ps->ps))]]
+  { }
+  [[omp::directive(target map(to: ps->ps->ps->s.f[:22]))]]
+  { }
+  [[omp::directive(target map(to: s.f[:22]) map(from: s.p[:33]))]]
+  { }
+  [[omp::directive(target map(from: s.f[:22]) map(to: ps->p[:33]))]]
+  { }
+  [[omp::directive(target map(from: s.f[:22], s.s) map(to: ps->p[:33]))]]
+  { }
+}
+
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";B;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";unknown;unknown;0;0;;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";A;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";x;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";fn;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+// DEBUG: @{{.+}} = private constant [7 x i8*] [i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @{{[0-9]+}}, i32 0, i32 0), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @{{[0-9]+}}, i32 0, i32 0), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @{{[0-9]+}}, i32 0, i32 0), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @{{[0-9]+}}, i32 0, i32 0), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @{{[0-9]+}}, i32 0, i32 0), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @{{[0-9]+}}, i32 0, i32 0), i8* getelementptr inbounds ([{{[0-9]+}} x i8], [{{[0-9]+}} x i8]* @{{[0-9]+}}, i32 0, i32 0)]
+
+void bar(int N) {
+  double B[10];
+  double A[N];
+  double x;
+  S1 s;
+  auto fn = [&x]() { return x; };
+  [[omp::directive(target)]]
+  {
+    (void)B;
+    (void)A;
+    (void)fn();
+    (void)s.f;
+  }
+}
+
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";t;{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+
+[[omp::directive(declare target)]];
+double t;
+[[omp::directive(end declare target)]];
+
+void baz() {
+  [[omp::directive(target map(to:t))]]
+  { }
+  [[omp::directive(target map(to:t) nowait)]]
+  { }
+  [[omp::directive(target teams map(to:t))]]
+  { }
+  [[omp::directive(target teams map(to:t) nowait)]]
+  { }
+  [[omp::directive(target data map(to:t))]]
+  { }
+  [[omp::sequence(directive(target enter data map(to:t)),
+                  directive(target enter data map(to:t) nowait),
+                  directive(target exit data map(from:t)),
+                  directive(target exit data map(from:t) nowait),
+                  directive(target update from(t)),
+                  directive(target update to(t)),
+                  directive(target update from(t) nowait),
+                  directive(target update to(t) nowait))]];
+}
+
+struct S3 {
+  double Z[64];
+};
+
+[[omp::directive(declare mapper(id: S3 s) map(s.Z[0:64]))]]
+void qux() {
+  S3 s;
+  [[omp::directive(target map(mapper(id), to:s))]]
+  { }
+}
+
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";s.Z[0:64];{{.*}}.cpp;{{[0-9]+}};{{[0-9]+}};;\00"
+
+// Clang used to mistakenly generate the map name "x" for both x and y on this
+// directive.  Conditions to reproduce the bug: a single map clause has two
+// variables, and at least the second is used in the associated statement.
+//
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";x;{{.*}}.cpp;[[@LINE+3]];7;;\00"
+// DEBUG: @{{[0-9]+}} = private unnamed_addr constant [{{[0-9]+}} x i8] c";y;{{.*}}.cpp;[[@LINE+2]];10;;\00"
+void secondMapNameInClause() {
+  int x, y;
+  [[omp::directive(target map(to: x, y))]];
+  x = y = 1;
+}
+
+// DEBUG: %{{.+}} = call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{.+}}, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* {{.+}}, i64* {{.+}}, i8** getelementptr inbounds ([{{[0-9]+}} x i8*], [{{[0-9]+}} x i8*]* @.offload_mapnames{{.*}}, i32 0, i32 0), i8** {{.+}})
+// DEBUG: %{{.+}} = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{.+}}, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* {{.+}}, i64* {{.+}}, i8** getelementptr inbounds ([{{[0-9]+}} x i8*], [{{[0-9]+}} x i8*]* @.offload_mapnames{{.*}}, i32 0, i32 0), i8** {{.+}}, i32 {{.+}}, i32 {{.+}})
+// DEBUG: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* {{.+}}, i64* {{.+}}, i8** getelementptr inbounds ([{{[0-9]+}} x i8*], [{{[0-9]+}} x i8*]* @.offload_mapnames{{.*}}, i32 0, i32 0), i8** {{.+}})
+// DEBUG: call void @__tgt_target_data_end_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* {{.+}}, i64* {{.+}}, i8** getelementptr inbounds ([{{[0-9]+}} x i8*], [{{[0-9]+}} x i8*]* @.offload_mapnames{{.*}}, i32 0, i32 0), i8** {{.+}})
+// DEBUG: call void @__tgt_target_data_update_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* {{.+}}, i64* {{.+}}, i8** getelementptr inbounds ([{{[0-9]+}} x i8*], [{{[0-9]+}} x i8*]* @.offload_mapnames{{.*}}, i32 0, i32 0), i8** {{.+}})
+// DEBUG: %{{.+}} = call i32 @__tgt_target_nowait_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{.+}}, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* %{{.+}}, i64* {{.+}}, i8** getelementptr inbounds ([{{[0-9]+}} x i8*], [{{[0-9]+}} x i8*]* @.offload_mapnames{{.*}}, i32 0, i32 0), i8** {{.+}})
+// DEBUG: %{{.+}} = call i32 @__tgt_target_teams_nowait_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{.+}}, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* {{.+}}, i64* {{.+}}, i8** getelementptr inbounds ([{{[0-9]+}} x i8*], [{{[0-9]+}} x i8*]* @.offload_mapnames{{.*}}, i32 0, i32 0), i8** {{.+}}, i32 {{.+}}, i32 {{.+}})
+// DEBUG: call void @__tgt_target_data_begin_nowait_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* {{.+}}, i64* {{.+}}, i8** getelementptr inbounds ([{{[0-9]+}} x i8*], [{{[0-9]+}} x i8*]* @.offload_mapnames{{.*}}, i32 0, i32 0), i8** {{.+}})
+// DEBUG: call void @__tgt_target_data_end_nowait_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* {{.+}}, i64* {{.+}}, i8** getelementptr inbounds ([{{[0-9]+}} x i8*], [{{[0-9]+}} x i8*]* @.offload_mapnames{{.*}}, i32 0, i32 0), i8** {{.+}})
+// DEBUG: call void @__tgt_target_data_update_nowait_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* {{.+}}, i64* {{.+}}, i8** getelementptr inbounds ([{{[0-9]+}} x i8*], [{{[0-9]+}} x i8*]* @.offload_mapnames{{.*}}, i32 0, i32 0), i8** {{.+}})
+
+// CHECK: %{{.+}} = call i32 @__tgt_target_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{.+}}, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* {{.+}}, i64* {{.+}}, i8** null, i8** {{.+}})
+// CHECK: %{{.+}} = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{.+}}, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* {{.+}}, i64* {{.+}}, i8** null, i8** {{.+}}, i32 {{.+}}, i32 {{.+}})
+// CHECK: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* {{.+}}, i64* {{.+}}, i8** null, i8** {{.+}})
+// CHECK: call void @__tgt_target_data_end_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* {{.+}}, i64* {{.+}}, i8** null, i8** {{.+}})
+// CHECK: call void @__tgt_target_data_update_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* {{.+}}, i64* {{.+}}, i8** null, i8** {{.+}})
+// CHECK: %{{.+}} = call i32 @__tgt_target_nowait_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{.+}}, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* %{{.+}}, i64* {{.+}}, i8** null, i8** {{.+}})
+// CHECK: %{{.+}} = call i32 @__tgt_target_teams_nowait_mapper(%struct.ident_t* @{{.+}}, i64 -1, i8* @{{.+}}, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* {{.+}}, i64* {{.+}}, i8** null, i8** {{.+}}, i32 {{.+}}, i32 {{.+}})
+// CHECK: call void @__tgt_target_data_begin_nowait_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* {{.+}}, i64* {{.+}}, i8** null, i8** {{.+}})
+// CHECK: call void @__tgt_target_data_end_nowait_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* {{.+}}, i64* {{.+}}, i8** null, i8** {{.+}})
+// CHECK: call void @__tgt_target_data_update_nowait_mapper(%struct.ident_t* @{{.+}}, i64 -1, i32 1, i8** %{{.+}}, i8** %{{.+}}, i64* {{.+}}, i64* {{.+}}, i8** null, i8** {{.+}})
+
+
+// DEBUG: void @.omp_mapper._ZTS2S3.id(i8* {{.*}}, i8* {{.*}}, i8* {{.*}}, i64 {{.*}}, i64 {{.*}}, i8* [[NAME_ARG:%.+]])
+// DEBUG: store i8* [[NAME_ARG]], i8** [[NAME_STACK:%.+]]
+// DEBUG: [[MAPPER_NAME:%.+]] = load i8*, i8** [[NAME_STACK]]
+// DEBUG: call void @__tgt_push_mapper_component(i8* %{{.*}}, i8* %{{.*}}, i8* %{{.*}}, i64 %{{.*}}, i64 %{{.*}}, i8* [[MAPPER_NAME]])
+
+#endif
+

diff  --git a/clang/test/OpenMP/taskloop_reduction_messages_attr.cpp b/clang/test/OpenMP/taskloop_reduction_messages_attr.cpp
new file mode 100644
index 0000000000000..3dc2258620004
--- /dev/null
+++ b/clang/test/OpenMP/taskloop_reduction_messages_attr.cpp
@@ -0,0 +1,311 @@
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=51 -fsyntax-only %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=51 -std=c++11 -fsyntax-only %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=51 -fsyntax-only %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fopenmp-simd -fopenmp-version=51 -std=c++11 -fsyntax-only %s -Wuninitialized
+
+typedef void **omp_allocator_handle_t;
+extern const omp_allocator_handle_t omp_null_allocator;
+extern const omp_allocator_handle_t omp_default_mem_alloc;
+extern const omp_allocator_handle_t omp_large_cap_mem_alloc;
+extern const omp_allocator_handle_t omp_const_mem_alloc;
+extern const omp_allocator_handle_t omp_high_bw_mem_alloc;
+extern const omp_allocator_handle_t omp_low_lat_mem_alloc;
+extern const omp_allocator_handle_t omp_cgroup_mem_alloc;
+extern const omp_allocator_handle_t omp_pteam_mem_alloc;
+extern const omp_allocator_handle_t omp_thread_mem_alloc;
+
+void xxx(int argc) {
+  int fp; // expected-note {{initialize the variable 'fp' to silence this warning}}
+  [[omp::directive(taskloop reduction(+:fp))]] // expected-warning {{variable 'fp' is uninitialized when used here}}
+  for (int i = 0; i < 10; ++i)
+    ;
+}
+
+void foo() {
+}
+
+bool foobool(int argc) {
+  return argc;
+}
+
+void foobar(int &ref) {
+  [[omp::directive(taskloop reduction(+:ref))]]
+  for (int i = 0; i < 10; ++i)
+    foo();
+}
+
+struct S1; // expected-note {{declared here}} expected-note 4 {{forward declaration of 'S1'}}
+extern S1 a;
+class S2 {
+  mutable int a;
+  S2 &operator+(const S2 &arg) { return (*this); } // expected-note 3 {{implicitly declared private here}}
+
+public:
+  S2() : a(0) {}
+  S2(S2 &s2) : a(s2.a) {}
+  static float S2s; // expected-note 2 {{static data member is predetermined as shared}}
+  static const float S2sc; // expected-note 2 {{'S2sc' declared here}}
+};
+const float S2::S2sc = 0;
+S2 b;                     // expected-note 3 {{'b' defined here}}
+const S2 ba[5];           // expected-note 2 {{'ba' defined here}}
+class S3 {
+  int a;
+
+public:
+  int b;
+  S3() : a(0) {}
+  S3(const S3 &s3) : a(s3.a) {}
+  S3 operator+(const S3 &arg1) { return arg1; }
+};
+int operator+(const S3 &arg1, const S3 &arg2) { return 5; }
+S3 c;               // expected-note 3 {{'c' defined here}}
+const S3 ca[5];     // expected-note 2 {{'ca' defined here}}
+extern const int f; // expected-note 4 {{'f' declared here}}
+class S4 {
+  int a;
+  S4(); // expected-note {{implicitly declared private here}}
+  S4(const S4 &s4);
+  S4 &operator+(const S4 &arg) { return (*this); }
+
+public:
+  S4(int v) : a(v) {}
+};
+S4 &operator&=(S4 &arg1, S4 &arg2) { return arg1; }
+class S5 {
+  int a:32;
+  S5() : a(0) {} // expected-note {{implicitly declared private here}}
+  S5(const S5 &s5) : a(s5.a) {}
+  S5 &operator+(const S5 &arg);
+
+public:
+  S5(int v) : a(v) {}
+};
+class S6 { // expected-note 3 {{candidate function (the implicit copy assignment operator) not viable: no known conversion from 'int' to 'const S6' for 1st argument}}
+#if __cplusplus >= 201103L // C++11 or later
+// expected-note at -2 3 {{candidate function (the implicit move assignment operator) not viable}}
+#endif
+  int a;
+
+public:
+  S6() : a(6) {}
+  operator int() { return 6; }
+} o;
+
+struct S7 {
+  int a: 32;
+  S7() {
+    [[omp::directive(taskloop reduction(+:a))]] // expected-error {{expected addressable reduction item for the task-based directives}}
+    for (int i = 0; i < 10; ++i)
+      ++a;
+  }
+};
+
+S3 h, k;
+[[omp::directive(threadprivate(h))]]; // expected-note 2 {{defined as threadprivate or thread local}}
+
+template <class T>       // expected-note {{declared here}}
+T tmain(T argc) {
+  const T d = T();       // expected-note 4 {{'d' defined here}}
+  const T da[5] = {T()}; // expected-note 2 {{'da' defined here}}
+  T qa[5] = {T()};
+  T i, z;
+  T &j = i;                        // expected-note 4 {{'j' defined here}}
+  S3 &p = k;                       // expected-note 2 {{'p' defined here}}
+  const T &r = da[(int)i];         // expected-note 2 {{'r' defined here}}
+  T &q = qa[(int)i];               // expected-note 2 {{'q' defined here}}
+  T fl;
+  [[omp::directive(taskloop reduction)]] // expected-error {{expected '(' after 'reduction'}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction +)]] // expected-error {{expected '(' after 'reduction'}} expected-warning {{extra tokens at the end of '#pragma omp taskloop' are ignored}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction())]] // expected-error {{expected unqualified-id}} expected-warning {{missing ':' after reduction identifier - ignoring}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(*))]] // expected-warning {{missing ':' after reduction identifier - ignoring}} expected-error {{expected expression}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(\))]] // expected-error {{expected unqualified-id}} expected-warning {{missing ':' after reduction identifier - ignoring}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(foo : argc))]] //expected-error {{incorrect reduction identifier, expected one of '+', '-', '*', '&', '|', '^', '&&', '||', 'min' or 'max' or declare reduction for type 'float'}} expected-error {{incorrect reduction identifier, expected one of '+', '-', '*', '&', '|', '^', '&&', '||', 'min' or 'max' or declare reduction for type 'int'}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(^ : T))]] // expected-error {{'T' does not refer to a value}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(+ : z, a, b, c, d, f))]] // expected-error {{a reduction list item with incomplete type 'S1'}} expected-error 3 {{const-qualified variable cannot be reduction}} expected-error 2 {{'operator+' is a private member of 'S2'}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(min : a, b, c, d, f))]] // expected-error {{a reduction list item with incomplete type 'S1'}} expected-error 4 {{arguments of OpenMP clause 'reduction' for 'min' or 'max' must be of arithmetic type}} expected-error 3 {{const-qualified variable cannot be reduction}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(max : h.b))]] // expected-error {{expected variable name, array element or array section}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(+ : ba))]] // expected-error {{const-qualified variable cannot be reduction}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(* : ca))]] // expected-error {{const-qualified variable cannot be reduction}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(- : da))]] // expected-error {{const-qualified variable cannot be reduction}} expected-error {{const-qualified variable cannot be reduction}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(^ : fl))]] // expected-error {{invalid operands to binary expression ('float' and 'float')}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(&& : S2::S2s))]] // expected-error {{shared variable cannot be reduction}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(&& : S2::S2sc))]] // expected-error {{const-qualified variable cannot be reduction}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(+ : h, k))]] // expected-error {{threadprivate or thread local variable cannot be reduction}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(+ : o))]] // expected-error 2 {{no viable overloaded '='}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop private(i), reduction(+ : j), reduction(+ : q))]] // expected-error 4 {{argument of OpenMP clause 'reduction' must reference the same object in all threads}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::sequence(directive(parallel private(k)), directive(taskloop reduction(+ : p), reduction(+ : p)))]] // expected-error 2 {{argument of OpenMP clause 'reduction' must reference the same object in all threads}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(+ : p), reduction(+ : p))]] // expected-error 2 {{variable can appear only once in OpenMP 'reduction' clause}} expected-note 2 {{previously referenced here}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(+ : r))]] // expected-error 2 {{const-qualified variable cannot be reduction}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::sequence(directive(parallel shared(i)), directive(parallel reduction(min : i)), directive(taskloop reduction(max : j)))]] // expected-error 2 {{argument of OpenMP clause 'reduction' must reference the same object in all threads}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::sequence(directive(parallel private(fl)), directive(taskloop reduction(+ : fl) allocate(omp_thread_mem_alloc: fl)))]] // expected-warning 2 {{allocator with the 'thread' trait access has unspecified behavior on 'taskloop' directive}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::sequence(directive(parallel reduction(* : fl)), directive(taskloop reduction(+ : fl)))]]
+  for (int i = 0; i < 10; ++i)
+    foo();
+
+  return T();
+}
+
+namespace A {
+double x;
+[[omp::directive(threadprivate(x))]]; // expected-note {{defined as threadprivate or thread local}}
+}
+namespace B {
+using A::x;
+}
+
+int main(int argc, char **argv) {
+  const int d = 5;       // expected-note 2 {{'d' defined here}}
+  const int da[5] = {0}; // expected-note {{'da' defined here}}
+  int qa[5] = {0};
+  S4 e(4);
+  S5 g(5);
+  int i, z;
+  int &j = i;                      // expected-note 2 {{'j' defined here}}
+  S3 &p = k;                       // expected-note 2 {{'p' defined here}}
+  const int &r = da[i];            // expected-note {{'r' defined here}}
+  int &q = qa[i];                  // expected-note {{'q' defined here}}
+  float fl;
+  [[omp::directive(taskloop reduction)]] // expected-error {{expected '(' after 'reduction'}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction +)]] // expected-error {{expected '(' after 'reduction'}} expected-warning {{extra tokens at the end of '#pragma omp taskloop' are ignored}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction())]] // expected-error {{expected unqualified-id}} expected-warning {{missing ':' after reduction identifier - ignoring}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(*))]] // expected-warning {{missing ':' after reduction identifier - ignoring}} expected-error {{expected expression}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(\))]] // expected-error {{expected unqualified-id}} expected-warning {{missing ':' after reduction identifier - ignoring}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(~ : argc))]] // expected-error {{expected unqualified-id}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(&& : argc, z))]]
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(^ : S1))]] // expected-error {{'S1' does not refer to a value}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(+ : a, b, c, d, f))]] // expected-error {{a reduction list item with incomplete type 'S1'}} expected-error 2 {{const-qualified variable cannot be reduction}} expected-error {{'operator+' is a private member of 'S2'}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(min : a, b, c, d, f))]] // expected-error {{a reduction list item with incomplete type 'S1'}} expected-error 2 {{arguments of OpenMP clause 'reduction' for 'min' or 'max' must be of arithmetic type}} expected-error 2 {{const-qualified variable cannot be reduction}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(max : h.b))]] // expected-error {{expected variable name, array element or array section}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(+ : ba))]] // expected-error {{const-qualified variable cannot be reduction}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(* : ca))]] // expected-error {{const-qualified variable cannot be reduction}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(- : da))]] // expected-error {{const-qualified variable cannot be reduction}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(^ : fl))]] // expected-error {{invalid operands to binary expression ('float' and 'float')}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(&& : S2::S2s))]] // expected-error {{shared variable cannot be reduction}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(&& : S2::S2sc))]] // expected-error {{const-qualified variable cannot be reduction}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(& : e, g))]] // expected-error {{calling a private constructor of class 'S4'}} expected-error {{calling a private constructor of class 'S5'}} expected-error {{invalid operands to binary expression ('S5' and 'S5')}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(+ : h, k, B::x))]] // expected-error 2 {{threadprivate or thread local variable cannot be reduction}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(+ : o))]] // expected-error {{no viable overloaded '='}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop private(i), reduction(+ : j), reduction(+ : q))]] // expected-error 2 {{argument of OpenMP clause 'reduction' must reference the same object in all threads}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::sequence(directive(parallel private(k)), directive(taskloop reduction(+ : p), reduction(+ : p)))]] // expected-error 2 {{argument of OpenMP clause 'reduction' must reference the same object in all threads}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(+ : p), reduction(+ : p))]] // expected-error {{variable can appear only once in OpenMP 'reduction' clause}} expected-note {{previously referenced here}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::directive(taskloop reduction(+ : r))]] // expected-error {{const-qualified variable cannot be reduction}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::sequence(directive(parallel shared(i)), directive(parallel reduction(min : i)), directive(taskloop reduction(max : j)))]] // expected-error {{argument of OpenMP clause 'reduction' must reference the same object in all threads}}
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::sequence(directive(parallel private(fl)), directive(taskloop reduction(+ : fl)))]]
+  for (int i = 0; i < 10; ++i)
+    foo();
+  [[omp::sequence(directive(parallel reduction(* : fl)), directive(taskloop reduction(+ : fl)))]]
+  for (int i = 0; i < 10; ++i)
+    foo();
+  static int m;
+  [[omp::directive(taskloop reduction(+ : m))]] // OK
+  for (int i = 0; i < 10; ++i)
+    m++;
+  [[omp::directive(taskloop reduction(task, + : m))]] // expected-error {{'reduction' clause with 'task' modifier allowed only on non-simd parallel or worksharing constructs}}
+  for (int i = 0; i < 10; ++i)
+    m++;
+  [[omp::directive(taskloop nogroup reduction(+ : m))]] // expected-error {{'reduction' clause cannot be used with 'nogroup' clause}}
+  for (int i = 0; i < 10; ++i)
+    m++;
+
+  return tmain(argc) + tmain(fl); // expected-note {{in instantiation of function template specialization 'tmain<int>' requested here}} expected-note {{in instantiation of function template specialization 'tmain<float>' requested here}}
+}
+

diff  --git a/clang/test/OpenMP/teams_distribute_parallel_for_simd_num_teams_messages_attr.cpp b/clang/test/OpenMP/teams_distribute_parallel_for_simd_num_teams_messages_attr.cpp
new file mode 100644
index 0000000000000..9c4bde791d3ed
--- /dev/null
+++ b/clang/test/OpenMP/teams_distribute_parallel_for_simd_num_teams_messages_attr.cpp
@@ -0,0 +1,73 @@
+// RUN: %clang_cc1 -verify -fsyntax-only -fopenmp -fopenmp-version=51 -std=c++11 %s -Wuninitialized
+// RUN: %clang_cc1 -verify -fsyntax-only -fopenmp-simd -fopenmp-version=51 -std=c++11 %s -Wuninitialized
+
+void foo() {
+}
+
+bool foobool(int argc) {
+  return argc;
+}
+
+struct S1; // expected-note 2 {{declared here}}
+
+template <typename T, int C> // expected-note {{declared here}}
+T tmain(T argc) {
+  char **a;
+  T k;
+  [[omp::sequence(directive(target), directive(teams distribute parallel for simd num_teams(C)))]]
+  for (int i=0; i<100; i++) foo();
+  [[omp::sequence(directive(target), directive(teams distribute parallel for simd num_teams(T)))]] // expected-error {{'T' does not refer to a value}}
+  for (int i=0; i<100; i++) foo();
+  [[omp::sequence(directive(target), directive(teams distribute parallel for simd num_teams))]] // expected-error {{expected '(' after 'num_teams'}}
+  for (int i=0; i<100; i++) foo();
+  [[omp::sequence(directive(target), directive(teams distribute parallel for simd num_teams()))]] // expected-error {{expected expression}}
+  for (int i=0; i<100; i++) foo();
+  [[omp::sequence(directive(target), directive(teams distribute parallel for simd num_teams(argc > 0 ? a[1] : a[2])))]] // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}}
+  for (int i=0; i<100; i++) foo();
+  [[omp::sequence(directive(target), directive(teams distribute parallel for simd num_teams(argc + k)))]]
+  for (int i=0; i<100; i++) foo();
+  [[omp::sequence(directive(target), directive(teams distribute parallel for simd num_teams(argc), num_teams (argc+1)))]] // expected-error {{directive '#pragma omp teams distribute parallel for simd' cannot contain more than one 'num_teams' clause}}
+  for (int i=0; i<100; i++) foo();
+  [[omp::sequence(directive(target), directive(teams distribute parallel for simd num_teams(S1)))]] // expected-error {{'S1' does not refer to a value}}
+  for (int i=0; i<100; i++) foo();
+  [[omp::sequence(directive(target), directive(teams distribute parallel for simd num_teams(-2)))]] // expected-error {{argument to 'num_teams' clause must be a strictly positive integer value}}
+  for (int i=0; i<100; i++) foo();
+  [[omp::sequence(directive(target), directive(teams distribute parallel for simd num_teams(-10u)))]]
+  for (int i=0; i<100; i++) foo();
+  [[omp::sequence(directive(target), directive(teams distribute parallel for simd num_teams(3.14)))]] // expected-error 2 {{expression must have integral or unscoped enumeration type, not 'double'}}
+  for (int i=0; i<100; i++) foo();
+
+  return 0;
+}
+
+int main(int argc, char **argv) {
+  int k;
+  [[omp::sequence(directive(target), directive(teams distribute parallel for simd num_teams))]] // expected-error {{expected '(' after 'num_teams'}}
+  for (int i=0; i<100; i++) foo();
+
+  [[omp::sequence(directive(target), directive(teams distribute parallel for simd num_teams ()))]] // expected-error {{expected expression}}
+  for (int i=0; i<100; i++) foo();
+
+  [[omp::sequence(directive(target), directive(teams distribute parallel for simd num_teams (argc > 0 ? argv[1] : argv[2])))]] // expected-error {{expression must have integral or unscoped enumeration type, not 'char *'}}
+  for (int i=0; i<100; i++) foo();
+
+  [[omp::sequence(directive(target), directive(teams distribute parallel for simd num_teams (argc + k)))]]
+  for (int i=0; i<100; i++) foo();
+
+  [[omp::sequence(directive(target), directive(teams distribute parallel for simd num_teams (argc), num_teams (argc+1)))]] // expected-error {{directive '#pragma omp teams distribute parallel for simd' cannot contain more than one 'num_teams' clause}}
+  for (int i=0; i<100; i++) foo();
+
+  [[omp::sequence(directive(target), directive(teams distribute parallel for simd num_teams (S1)))]] // expected-error {{'S1' does not refer to a value}}
+  for (int i=0; i<100; i++) foo();
+
+  [[omp::sequence(directive(target), directive(teams distribute parallel for simd num_teams (-2)))]] // expected-error {{argument to 'num_teams' clause must be a strictly positive integer value}}
+  for (int i=0; i<100; i++) foo();
+
+  [[omp::sequence(directive(target), directive(teams distribute parallel for simd num_teams (-10u)))]]
+  for (int i=0; i<100; i++) foo();
+
+  [[omp::sequence(directive(target), directive(teams distribute parallel for simd num_teams (3.14)))]] // expected-error {{expression must have integral or unscoped enumeration type, not 'double'}}
+  for (int i=0; i<100; i++) foo();
+
+  return tmain<int, 10>(argc); // expected-note {{in instantiation of function template specialization 'tmain<int, 10>' requested here}}
+}

diff  --git a/clang/test/OpenMP/unroll_codegen_unroll_for_attr.cpp b/clang/test/OpenMP/unroll_codegen_unroll_for_attr.cpp
new file mode 100644
index 0000000000000..f431201bcf476
--- /dev/null
+++ b/clang/test/OpenMP/unroll_codegen_unroll_for_attr.cpp
@@ -0,0 +1,237 @@
+// Check code generation
+// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp -fopenmp-version=51 -emit-llvm %s -o - | FileCheck %s --check-prefix=IR
+
+// Check same results after serialization round-trip
+// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp -fopenmp-version=51 -emit-pch -o %t %s
+// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp -fopenmp-version=51 -include-pch %t -emit-llvm %s -o - | FileCheck %s --check-prefix=IR
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+// placeholder for loop body code.
+extern "C" void body(...) {}
+
+
+// IR-LABEL: @func(
+// IR-NEXT:  [[ENTRY:.*]]:
+// IR-NEXT:    %[[START_ADDR:.+]] = alloca i32, align 4
+// IR-NEXT:    %[[END_ADDR:.+]] = alloca i32, align 4
+// IR-NEXT:    %[[STEP_ADDR:.+]] = alloca i32, align 4
+// IR-NEXT:    %[[DOTOMP_IV:.+]] = alloca i32, align 4
+// IR-NEXT:    %[[TMP:.+]] = alloca i32, align 4
+// IR-NEXT:    %[[I:.+]] = alloca i32, align 4
+// IR-NEXT:    %[[DOTCAPTURE_EXPR_:.+]] = alloca i32, align 4
+// IR-NEXT:    %[[DOTCAPTURE_EXPR_1:.+]] = alloca i32, align 4
+// IR-NEXT:    %[[DOTCAPTURE_EXPR_2:.+]] = alloca i32, align 4
+// IR-NEXT:    %[[DOTCAPTURE_EXPR_3:.+]] = alloca i32, align 4
+// IR-NEXT:    %[[DOTUNROLLED_IV_I:.+]] = alloca i32, align 4
+// IR-NEXT:    %[[DOTCAPTURE_EXPR_6:.+]] = alloca i32, align 4
+// IR-NEXT:    %[[DOTCAPTURE_EXPR_8:.+]] = alloca i32, align 4
+// IR-NEXT:    %[[DOTCAPTURE_EXPR_12:.+]] = alloca i32, align 4
+// IR-NEXT:    %[[DOTCAPTURE_EXPR_14:.+]] = alloca i32, align 4
+// IR-NEXT:    %[[DOTUNROLLED_IV__UNROLLED_IV_I:.+]] = alloca i32, align 4
+// IR-NEXT:    %[[DOTOMP_LB:.+]] = alloca i32, align 4
+// IR-NEXT:    %[[DOTOMP_UB:.+]] = alloca i32, align 4
+// IR-NEXT:    %[[DOTOMP_STRIDE:.+]] = alloca i32, align 4
+// IR-NEXT:    %[[DOTOMP_IS_LAST:.+]] = alloca i32, align 4
+// IR-NEXT:    %[[DOTUNROLLED_IV__UNROLLED_IV_I18:.+]] = alloca i32, align 4
+// IR-NEXT:    %[[DOTUNROLL_INNER_IV__UNROLLED_IV_I:.+]] = alloca i32, align 4
+// IR-NEXT:    %[[DOTUNROLL_INNER_IV_I:.+]] = alloca i32, align 4
+// IR-NEXT:    %[[TMP0:.+]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @2)
+// IR-NEXT:    store i32 %[[START:.+]], i32* %[[START_ADDR]], align 4
+// IR-NEXT:    store i32 %[[END:.+]], i32* %[[END_ADDR]], align 4
+// IR-NEXT:    store i32 %[[STEP:.+]], i32* %[[STEP_ADDR]], align 4
+// IR-NEXT:    %[[TMP1:.+]] = load i32, i32* %[[START_ADDR]], align 4
+// IR-NEXT:    store i32 %[[TMP1]], i32* %[[I]], align 4
+// IR-NEXT:    %[[TMP2:.+]] = load i32, i32* %[[START_ADDR]], align 4
+// IR-NEXT:    store i32 %[[TMP2]], i32* %[[DOTCAPTURE_EXPR_]], align 4
+// IR-NEXT:    %[[TMP3:.+]] = load i32, i32* %[[END_ADDR]], align 4
+// IR-NEXT:    store i32 %[[TMP3]], i32* %[[DOTCAPTURE_EXPR_1]], align 4
+// IR-NEXT:    %[[TMP4:.+]] = load i32, i32* %[[STEP_ADDR]], align 4
+// IR-NEXT:    store i32 %[[TMP4]], i32* %[[DOTCAPTURE_EXPR_2]], align 4
+// IR-NEXT:    %[[TMP5:.+]] = load i32, i32* %[[DOTCAPTURE_EXPR_1]], align 4
+// IR-NEXT:    %[[TMP6:.+]] = load i32, i32* %[[DOTCAPTURE_EXPR_]], align 4
+// IR-NEXT:    %[[SUB:.+]] = sub i32 %[[TMP5]], %[[TMP6]]
+// IR-NEXT:    %[[SUB4:.+]] = sub i32 %[[SUB]], 1
+// IR-NEXT:    %[[TMP7:.+]] = load i32, i32* %[[DOTCAPTURE_EXPR_2]], align 4
+// IR-NEXT:    %[[ADD:.+]] = add i32 %[[SUB4]], %[[TMP7]]
+// IR-NEXT:    %[[TMP8:.+]] = load i32, i32* %[[DOTCAPTURE_EXPR_2]], align 4
+// IR-NEXT:    %[[DIV:.+]] = udiv i32 %[[ADD]], %[[TMP8]]
+// IR-NEXT:    %[[SUB5:.+]] = sub i32 %[[DIV]], 1
+// IR-NEXT:    store i32 %[[SUB5]], i32* %[[DOTCAPTURE_EXPR_3]], align 4
+// IR-NEXT:    store i32 0, i32* %[[DOTUNROLLED_IV_I]], align 4
+// IR-NEXT:    %[[TMP9:.+]] = load i32, i32* %[[DOTCAPTURE_EXPR_3]], align 4
+// IR-NEXT:    %[[ADD7:.+]] = add i32 %[[TMP9]], 1
+// IR-NEXT:    store i32 %[[ADD7]], i32* %[[DOTCAPTURE_EXPR_6]], align 4
+// IR-NEXT:    %[[TMP10:.+]] = load i32, i32* %[[DOTCAPTURE_EXPR_6]], align 4
+// IR-NEXT:    %[[SUB9:.+]] = sub i32 %[[TMP10]], -1
+// IR-NEXT:    %[[DIV10:.+]] = udiv i32 %[[SUB9]], 2
+// IR-NEXT:    %[[SUB11:.+]] = sub i32 %[[DIV10]], 1
+// IR-NEXT:    store i32 %[[SUB11]], i32* %[[DOTCAPTURE_EXPR_8]], align 4
+// IR-NEXT:    %[[TMP11:.+]] = load i32, i32* %[[DOTCAPTURE_EXPR_8]], align 4
+// IR-NEXT:    %[[ADD13:.+]] = add i32 %[[TMP11]], 1
+// IR-NEXT:    store i32 %[[ADD13]], i32* %[[DOTCAPTURE_EXPR_12]], align 4
+// IR-NEXT:    %[[TMP12:.+]] = load i32, i32* %[[DOTCAPTURE_EXPR_12]], align 4
+// IR-NEXT:    %[[SUB15:.+]] = sub i32 %[[TMP12]], -1
+// IR-NEXT:    %[[DIV16:.+]] = udiv i32 %[[SUB15]], 2
+// IR-NEXT:    %[[SUB17:.+]] = sub i32 %[[DIV16]], 1
+// IR-NEXT:    store i32 %[[SUB17]], i32* %[[DOTCAPTURE_EXPR_14]], align 4
+// IR-NEXT:    store i32 0, i32* %[[DOTUNROLLED_IV__UNROLLED_IV_I]], align 4
+// IR-NEXT:    %[[TMP13:.+]] = load i32, i32* %[[DOTCAPTURE_EXPR_12]], align 4
+// IR-NEXT:    %[[CMP:.+]] = icmp ult i32 0, %[[TMP13]]
+// IR-NEXT:    br i1 %[[CMP]], label %[[OMP_PRECOND_THEN:.+]], label %[[OMP_PRECOND_END:.+]]
+// IR-EMPTY:
+// IR-NEXT:  [[OMP_PRECOND_THEN]]:
+// IR-NEXT:    store i32 0, i32* %[[DOTOMP_LB]], align 4
+// IR-NEXT:    %[[TMP14:.+]] = load i32, i32* %[[DOTCAPTURE_EXPR_14]], align 4
+// IR-NEXT:    store i32 %[[TMP14]], i32* %[[DOTOMP_UB]], align 4
+// IR-NEXT:    store i32 1, i32* %[[DOTOMP_STRIDE]], align 4
+// IR-NEXT:    store i32 0, i32* %[[DOTOMP_IS_LAST]], align 4
+// IR-NEXT:    call void @__kmpc_for_static_init_4u(%struct.ident_t* @1, i32 %[[TMP0]], i32 34, i32* %[[DOTOMP_IS_LAST]], i32* %[[DOTOMP_LB]], i32* %[[DOTOMP_UB]], i32* %[[DOTOMP_STRIDE]], i32 1, i32 1)
+// IR-NEXT:    %[[TMP15:.+]] = load i32, i32* %[[DOTOMP_UB]], align 4
+// IR-NEXT:    %[[TMP16:.+]] = load i32, i32* %[[DOTCAPTURE_EXPR_14]], align 4
+// IR-NEXT:    %[[CMP19:.+]] = icmp ugt i32 %[[TMP15]], %[[TMP16]]
+// IR-NEXT:    br i1 %[[CMP19]], label %[[COND_TRUE:.+]], label %[[COND_FALSE:.+]]
+// IR-EMPTY:
+// IR-NEXT:  [[COND_TRUE]]:
+// IR-NEXT:    %[[TMP17:.+]] = load i32, i32* %[[DOTCAPTURE_EXPR_14]], align 4
+// IR-NEXT:    br label %[[COND_END:.+]]
+// IR-EMPTY:
+// IR-NEXT:  [[COND_FALSE]]:
+// IR-NEXT:    %[[TMP18:.+]] = load i32, i32* %[[DOTOMP_UB]], align 4
+// IR-NEXT:    br label %[[COND_END]]
+// IR-EMPTY:
+// IR-NEXT:  [[COND_END]]:
+// IR-NEXT:    %[[COND:.+]] = phi i32 [ %[[TMP17]], %[[COND_TRUE]] ], [ %[[TMP18]], %[[COND_FALSE]] ]
+// IR-NEXT:    store i32 %[[COND]], i32* %[[DOTOMP_UB]], align 4
+// IR-NEXT:    %[[TMP19:.+]] = load i32, i32* %[[DOTOMP_LB]], align 4
+// IR-NEXT:    store i32 %[[TMP19]], i32* %[[DOTOMP_IV]], align 4
+// IR-NEXT:    br label %[[OMP_INNER_FOR_COND:.+]]
+// IR-EMPTY:
+// IR-NEXT:  [[OMP_INNER_FOR_COND]]:
+// IR-NEXT:    %[[TMP20:.+]] = load i32, i32* %[[DOTOMP_IV]], align 4
+// IR-NEXT:    %[[TMP21:.+]] = load i32, i32* %[[DOTOMP_UB]], align 4
+// IR-NEXT:    %[[ADD20:.+]] = add i32 %[[TMP21]], 1
+// IR-NEXT:    %[[CMP21:.+]] = icmp ult i32 %[[TMP20]], %[[ADD20]]
+// IR-NEXT:    br i1 %[[CMP21]], label %[[OMP_INNER_FOR_BODY:.+]], label %[[OMP_INNER_FOR_END:.+]]
+// IR-EMPTY:
+// IR-NEXT:  [[OMP_INNER_FOR_BODY]]:
+// IR-NEXT:    %[[TMP22:.+]] = load i32, i32* %[[DOTOMP_IV]], align 4
+// IR-NEXT:    %[[MUL:.+]] = mul i32 %[[TMP22]], 2
+// IR-NEXT:    %[[ADD22:.+]] = add i32 0, %[[MUL]]
+// IR-NEXT:    store i32 %[[ADD22]], i32* %[[DOTUNROLLED_IV__UNROLLED_IV_I18]], align 4
+// IR-NEXT:    %[[TMP23:.+]] = load i32, i32* %[[DOTUNROLLED_IV__UNROLLED_IV_I18]], align 4
+// IR-NEXT:    store i32 %[[TMP23]], i32* %[[DOTUNROLL_INNER_IV__UNROLLED_IV_I]], align 4
+// IR-NEXT:    br label %[[FOR_COND:.+]]
+// IR-EMPTY:
+// IR-NEXT:  [[FOR_COND]]:
+// IR-NEXT:    %[[TMP24:.+]] = load i32, i32* %[[DOTUNROLL_INNER_IV__UNROLLED_IV_I]], align 4
+// IR-NEXT:    %[[TMP25:.+]] = load i32, i32* %[[DOTUNROLLED_IV__UNROLLED_IV_I18]], align 4
+// IR-NEXT:    %[[ADD23:.+]] = add i32 %[[TMP25]], 2
+// IR-NEXT:    %[[CMP24:.+]] = icmp ule i32 %[[TMP24]], %[[ADD23]]
+// IR-NEXT:    br i1 %[[CMP24]], label %[[LAND_RHS:.+]], label %[[LAND_END:.+]]
+// IR-EMPTY:
+// IR-NEXT:  [[LAND_RHS]]:
+// IR-NEXT:    %[[TMP26:.+]] = load i32, i32* %[[DOTUNROLL_INNER_IV__UNROLLED_IV_I]], align 4
+// IR-NEXT:    %[[TMP27:.+]] = load i32, i32* %[[DOTCAPTURE_EXPR_8]], align 4
+// IR-NEXT:    %[[ADD25:.+]] = add i32 %[[TMP27]], 1
+// IR-NEXT:    %[[CMP26:.+]] = icmp ule i32 %[[TMP26]], %[[ADD25]]
+// IR-NEXT:    br label %[[LAND_END]]
+// IR-EMPTY:
+// IR-NEXT:  [[LAND_END]]:
+// IR-NEXT:    %[[TMP28:.+]] = phi i1 [ false, %[[FOR_COND]] ], [ %[[CMP26]], %[[LAND_RHS]] ]
+// IR-NEXT:    br i1 %[[TMP28]], label %[[FOR_BODY:.+]], label %[[FOR_END41:.+]]
+// IR-EMPTY:
+// IR-NEXT:  [[FOR_BODY]]:
+// IR-NEXT:    %[[TMP29:.+]] = load i32, i32* %[[DOTUNROLL_INNER_IV__UNROLLED_IV_I]], align 4
+// IR-NEXT:    %[[MUL27:.+]] = mul i32 %[[TMP29]], 2
+// IR-NEXT:    %[[ADD28:.+]] = add i32 0, %[[MUL27]]
+// IR-NEXT:    store i32 %[[ADD28]], i32* %[[DOTUNROLLED_IV_I]], align 4
+// IR-NEXT:    %[[TMP30:.+]] = load i32, i32* %[[DOTUNROLLED_IV_I]], align 4
+// IR-NEXT:    store i32 %[[TMP30]], i32* %[[DOTUNROLL_INNER_IV_I]], align 4
+// IR-NEXT:    br label %[[FOR_COND29:.+]]
+// IR-EMPTY:
+// IR-NEXT:  [[FOR_COND29]]:
+// IR-NEXT:    %[[TMP31:.+]] = load i32, i32* %[[DOTUNROLL_INNER_IV_I]], align 4
+// IR-NEXT:    %[[TMP32:.+]] = load i32, i32* %[[DOTUNROLLED_IV_I]], align 4
+// IR-NEXT:    %[[ADD30:.+]] = add i32 %[[TMP32]], 2
+// IR-NEXT:    %[[CMP31:.+]] = icmp ule i32 %[[TMP31]], %[[ADD30]]
+// IR-NEXT:    br i1 %[[CMP31]], label %[[LAND_RHS32:.+]], label %[[LAND_END35:.+]]
+// IR-EMPTY:
+// IR-NEXT:  [[LAND_RHS32]]:
+// IR-NEXT:    %[[TMP33:.+]] = load i32, i32* %[[DOTUNROLL_INNER_IV_I]], align 4
+// IR-NEXT:    %[[TMP34:.+]] = load i32, i32* %[[DOTCAPTURE_EXPR_3]], align 4
+// IR-NEXT:    %[[ADD33:.+]] = add i32 %[[TMP34]], 1
+// IR-NEXT:    %[[CMP34:.+]] = icmp ule i32 %[[TMP33]], %[[ADD33]]
+// IR-NEXT:    br label %[[LAND_END35]]
+// IR-EMPTY:
+// IR-NEXT:  [[LAND_END35]]:
+// IR-NEXT:    %[[TMP35:.+]] = phi i1 [ false, %[[FOR_COND29]] ], [ %[[CMP34]], %[[LAND_RHS32]] ]
+// IR-NEXT:    br i1 %[[TMP35]], label %[[FOR_BODY36:.+]], label %[[FOR_END:.+]]
+// IR-EMPTY:
+// IR-NEXT:  [[FOR_BODY36]]:
+// IR-NEXT:    %[[TMP36:.+]] = load i32, i32* %[[DOTCAPTURE_EXPR_]], align 4
+// IR-NEXT:    %[[TMP37:.+]] = load i32, i32* %[[DOTUNROLL_INNER_IV_I]], align 4
+// IR-NEXT:    %[[TMP38:.+]] = load i32, i32* %[[DOTCAPTURE_EXPR_2]], align 4
+// IR-NEXT:    %[[MUL37:.+]] = mul i32 %[[TMP37]], %[[TMP38]]
+// IR-NEXT:    %[[ADD38:.+]] = add i32 %[[TMP36]], %[[MUL37]]
+// IR-NEXT:    store i32 %[[ADD38]], i32* %[[I]], align 4
+// IR-NEXT:    %[[TMP39:.+]] = load i32, i32* %[[START_ADDR]], align 4
+// IR-NEXT:    %[[TMP40:.+]] = load i32, i32* %[[END_ADDR]], align 4
+// IR-NEXT:    %[[TMP41:.+]] = load i32, i32* %[[STEP_ADDR]], align 4
+// IR-NEXT:    %[[TMP42:.+]] = load i32, i32* %[[I]], align 4
+// IR-NEXT:    call void (...) @body(i32 %[[TMP39]], i32 %[[TMP40]], i32 %[[TMP41]], i32 %[[TMP42]])
+// IR-NEXT:    br label %[[FOR_INC:.+]]
+// IR-EMPTY:
+// IR-NEXT:  [[FOR_INC]]:
+// IR-NEXT:    %[[TMP43:.+]] = load i32, i32* %[[DOTUNROLL_INNER_IV_I]], align 4
+// IR-NEXT:    %[[INC:.+]] = add i32 %[[TMP43]], 1
+// IR-NEXT:    store i32 %[[INC]], i32* %[[DOTUNROLL_INNER_IV_I]], align 4
+// IR-NEXT:    br label %[[FOR_COND29]], !llvm.loop ![[LOOP2:[0-9]+]]
+// IR-EMPTY:
+// IR-NEXT:  [[FOR_END]]:
+// IR-NEXT:    br label %[[FOR_INC39:.+]]
+// IR-EMPTY:
+// IR-NEXT:  [[FOR_INC39]]:
+// IR-NEXT:    %[[TMP44:.+]] = load i32, i32* %[[DOTUNROLL_INNER_IV__UNROLLED_IV_I]], align 4
+// IR-NEXT:    %[[INC40:.+]] = add i32 %[[TMP44]], 1
+// IR-NEXT:    store i32 %[[INC40]], i32* %[[DOTUNROLL_INNER_IV__UNROLLED_IV_I]], align 4
+// IR-NEXT:    br label %[[FOR_COND]], !llvm.loop ![[LOOP5:[0-9]+]]
+// IR-EMPTY:
+// IR-NEXT:  [[FOR_END41]]:
+// IR-NEXT:    br label %[[OMP_BODY_CONTINUE:.+]]
+// IR-EMPTY:
+// IR-NEXT:  [[OMP_BODY_CONTINUE]]:
+// IR-NEXT:    br label %[[OMP_INNER_FOR_INC:.+]]
+// IR-EMPTY:
+// IR-NEXT:  [[OMP_INNER_FOR_INC]]:
+// IR-NEXT:    %[[TMP45:.+]] = load i32, i32* %[[DOTOMP_IV]], align 4
+// IR-NEXT:    %[[ADD42:.+]] = add i32 %[[TMP45]], 1
+// IR-NEXT:    store i32 %[[ADD42]], i32* %[[DOTOMP_IV]], align 4
+// IR-NEXT:    br label %[[OMP_INNER_FOR_COND]]
+// IR-EMPTY:
+// IR-NEXT:  [[OMP_INNER_FOR_END]]:
+// IR-NEXT:    br label %[[OMP_LOOP_EXIT:.+]]
+// IR-EMPTY:
+// IR-NEXT:  [[OMP_LOOP_EXIT]]:
+// IR-NEXT:    call void @__kmpc_for_static_fini(%struct.ident_t* @1, i32 %[[TMP0]])
+// IR-NEXT:    br label %[[OMP_PRECOND_END]]
+// IR-EMPTY:
+// IR-NEXT:  [[OMP_PRECOND_END]]:
+// IR-NEXT:    call void @__kmpc_barrier(%struct.ident_t* @3, i32 %[[TMP0]])
+// IR-NEXT:    ret void
+// IR-NEXT:  }
+extern "C" void func(int start, int end, int step) {
+  [[omp::sequence(directive(for), directive(unroll partial), directive(unroll partial))]]
+  for (int i = start; i < end; i+=step)
+    body(start, end, step, i);
+}
+
+#endif /* HEADER */
+
+
+// IR: ![[LOOP2]] = distinct !{![[LOOP2]], ![[LOOPPROP3:[0-9]+]], ![[LOOPPROP4:[0-9]+]]}
+// IR: ![[LOOPPROP3]] = !{!"llvm.loop.mustprogress"}
+// IR: ![[LOOPPROP4]] = !{!"llvm.loop.unroll.count", i32 2}
+// IR: ![[LOOP5]] = distinct !{![[LOOP5]], ![[LOOPPROP3]], ![[LOOPPROP4]]}