[OpenMP] Implement '#pragma omp tile'

[Meinersbur]

Summary

The tile directive is in OpenMP's Technical Report 8 and foreseeably will be part of the upcoming OpenMP 5.1 standard.

This implementation is based on an AST transformation providing a de-sugared loop nest. This makes it simple to forward the de-sugared transformation to loop associated directives taking the tiled loops. In contrast to other loop associated directives, the OMPTileDirective does not use CapturedStmts. Letting loop associated directives consume loops from different capture context would be difficult.

A significant amount of code generation logic is taking place in the Sema class. Eventually, I would prefer if these would move into the CodeGen component such that we could make use of the OpenMPIRBuilder, together with flang. Only expressions converting between the language's iteration variable and the logical iteration space need to take place in the semantic analyzer: Getting the of iterations (e.g. the overload resolution of std::distance) and converting the logical iteration number to the iteration variable (e.g. overload resolution of iteration + .omp.iv). In clang, only CXXForRangeStmt is also represented by its de-sugared components. However, OpenMP loop are not defined as syntatic sugar. Starting with an AST-based approach allows us to gradually move generated AST statements into CodeGen, instead all at once.

I would also like to refactor checkOpenMPLoop into its functionalities in a follow-up. In this patch it is used twice. Once for checking proper nesting and emitting diagnostics, and additionally for deriving the logical iteration space per-loop (instead of for the loop nest).

[Meinersbur]

ping

[Meinersbur]

• Rebase
• Remove unrelated change

[Meinersbur]

• Rebase after clang-format of cindex

[JonChesterfield]

edit: actually you've already done the clang-format on trunk as I hoped, phabricator mislead me. Apologies for the noise

[Meinersbur]

In D76342#1941591, @JonChesterfield wrote:

edit: actually you've already done the clang-format on trunk as I hoped, phabricator mislead me. Apologies for the noise

The previous diff was against a branch with only the whitespace changes to get rid of the noise in the diff. It seems that the pre-merge check instead compared formatting to the master (probably Arcanist/Phabricator looked for the common ancestor of origin/master and my git branch). This let to weird formatting warnings that weren't even changes highlighted in this diff. They are gone now with clang-formatted origin/master. Fotunately.

[Meinersbur]

ping

[Meinersbur]

ping

[Meinersbur]

pong

[ABataev]

clang/include/clang/AST/StmtOpenMP.h

4781–4784

Not sure that this is a good idea to treat this directive as the executable directive. To me, it looks like kind of AttributedStmt. Maybe better to introduce some kind of a new base node for this and similar constructs, which does not own the loop but is its kind of attribute-like entity?
Also, can we have something like:

#pragma omp simd
#pragma omp tile ...
for(...) ;

Thoughts?

[Meinersbur]

clang/include/clang/AST/StmtOpenMP.h

4781–4784

While not executed at runtime, syntactically it is parsed like a executable (loop-associated) directive. IMHO it does 'own' the loop, but produces another one for to be owned(/associated) by a different directive, as in your tile/simd example, which should already work. Allowing this was the motivation to do the transformation on the AST-level for now.

[Meinersbur]

ping

[Meinersbur]

ping

[Meinersbur]

ping

[ABataev]

clang/include/clang/AST/StmtOpenMP.h

4781–4784

I'm not saying that we should separate parsing of this directive from others, it is just better to treat this directive as a little bit different node. Currently, it introduces too many changes in the base classes. Better to create a new base class, that does not relies on CapturedStmt as the base, and derive OMPExecutableDirective and this directive and other similar (+ maybe, OMPSimdDirective) from this new base class.

[Meinersbur]

clang/include/clang/AST/StmtOpenMP.h

4781–4784

Unless you tell me otherwise, OMPLoopDirective represents a loop-associated directive. #pragma omp tile is a loop-associated directive. OMPLoopDirective contains all the functionality to parse associated loops, and unfortunately if derived from OMPExecutableDirective.

You seem to ask me to create a new class "OMPDirectiveAssociatedWithLoopButNotExecutable" that duplicates the parsing part of "OMPLoopDirective"? This will either be a lot of duplicated code or result in even more changes to the base classes due to the refactoring.

By the OpenMP specification, simd and tile are executable directives, so structurally I think the class hierarchy as-is makes sense. From the glossary of the upcoming OpenMP 5.1:

An OpenMP directive that appears in an executable context and results in implementation code and/or prescribes the manner in which associated user code must execute.

Avoiding a CapturedStmt when not needed would a modification of clang::getOpenMPCaptureRegions which currently adds a capture of type OMPD_unknown for such directives. This is unrelated to loop-associated directives.

[ABataev]

clang/include/clang/AST/StmtOpenMP.h

4781–4784

No, this is not what I'm asking for. I asked you to think about adding a new level of abstraction somewhere between the OMPLoop... and OMPExecutableDirective classes to minimize the functional changes and make the classes more robust. Anyway, give me a week or so to try to find possible better abstractions and if there are no better ones, we'll proceed with your original solution.

[Meinersbur]

clang/include/clang/AST/StmtOpenMP.h

4781–4784

derive OMPExecutableDirective and this directive and other similar (+ maybe, OMPSimdDirective) from this new base class

sounded like the new class hierarchy should be

Stmt -> NewClass -> OMPExecutableDirective -> OMPLoopDirective -> OMPForDirective
                \
                 -> OMPTileDirective
                \
                 -> OMPSimdDirective

Since OMPSimdDirective seems to be affected as well, this seems more like a refactoring of the existing structure than something specific to #pragma omp tile.

Looking forward to your idea.

[ABataev]

clang/include/clang/AST/StmtOpenMP.h

4781–4784

I have a little bit different scheme in my mind, but you got the main idea. Let me check, if it works, or I can find a better design.

[Meinersbur]

ping

@ABataev Have you checked the alternative scheme?

[ABataev]

In D76342#2128739, @Meinersbur wrote:

ping

@ABataev Have you checked the alternative scheme?

Working on it, need a little bit more time.

[Meinersbur]

• Rebase, especially omp_gen
• Some cleanup

[Meinersbur]

Rebase after D83261

[ABataev]

In D76342#2262414, @Meinersbur wrote:

Rebase after D83261

I will try to adapt it to the new interface (it won't work for simd, still need to capture the variables) and will send the patch back to you.

[Meinersbur]

Thank you

[ABataev]

Rework, bug fixes.

[Meinersbur]

Thank you for the update.

From comparing the diffs, here is the list of main changes I found:

1. Create OMPLoopBasedDirective as a new base class for OMPLoopDirective (and OMPTileDirective) for CollapsedNum and some utility methods for analyzing the CapturedStmt/ForStmt structure (replacing getTopmostAssociatedStructuredBlock, collectAssociatedLoops)
2. Add PreInits child to OMPTileDirective (instead of wrapping them in a CompoundStmt of the transformed body/collectAssociatedLoops)
3. Use meaningful SourceLocations instead of my placeholder {}
4. New OMPTransformDirectiveScopeRAII
   1. to collect nested PreInits (instead by collectAssociatedLoops)
   2. to assign values to CapturedDecls (instead of adding a Capturing argument to various function)
5. no call to checkOpenMPLoop for the entire lop nest (but still once per loop)

My remarks to these changes:

1. Saves us to store all the loop-specific subexpressions in the AST node. However, they are still computed in ActOnOpenMPTileDirective. With the OpenMPIRBuilder (D94973) these will also not be necessary for other loop-associated statements.
2. looks more consistent with other directives
3. This was just laziness on my size. Thank you.
4. is what I just did not know. 4.b. still feels like a hack because there are captures variables outside of any CapturedStmt and therefore complicated the AST. Comparable directive such as master don't do this.
5. The additional call on the entire loop nest felt necessary to check constraints such as invariant loop bounds and disallow nesting. However, both properties are still checked now.

clang/include/clang/AST/RecursiveASTVisitor.h

2794–2795

[no change request] OMPLoopBasedDirective is not represented in the visitor.

Well, nor does anything call TraverseOMPLoopDirective, so the possibility to only override the base class traverse-function is unused.

clang/include/clang/AST/StmtOpenMP.h

452

Please add a documentation what this class represents.

457

Inaccurate for the tile directive.

465

Inaccurate for the tile directive.

660–686

Please add doxygen comments.

IMHO, using callbacks makes the callers significantly more complicated. Why not fill a SmallVectorImpl<Stmt*> with the result?

668

I do not see why making this a forwarding reference.

689–718

Use isOpenMPLoopDirective() instead?

[ABataev]

In D76342#2535321, @Meinersbur wrote:

Thank you for the update.

From comparing the diffs, here is the list of main changes I found:

1. Create OMPLoopBasedDirective as a new base class for OMPLoopDirective (and OMPTileDirective) for CollapsedNum and some utility methods for analyzing the CapturedStmt/ForStmt structure (replacing getTopmostAssociatedStructuredBlock, collectAssociatedLoops)
2. Add PreInits child to OMPTileDirective (instead of wrapping them in a CompoundStmt of the transformed body/collectAssociatedLoops)
3. Use meaningful SourceLocations instead of my placeholder {}
4. New OMPTransformDirectiveScopeRAII
   1. to collect nested PreInits (instead by collectAssociatedLoops)
   2. to assign values to CapturedDecls (instead of adding a Capturing argument to various function)
5. no call to checkOpenMPLoop for the entire lop nest (but still once per loop)

My remarks to these changes:

1. Saves us to store all the loop-specific subexpressions in the AST node. However, they are still computed in ActOnOpenMPTileDirective. With the OpenMPIRBuilder (D94973) these will also not be necessary for other loop-associated statements.
2. looks more consistent with other directives
3. This was just laziness on my size. Thank you.
4. is what I just did not know. 4.b. still feels like a hack because there are captures variables outside of any CapturedStmt and therefore complicated the AST. Comparable directive such as master don't do this.
5. The additional call on the entire loop nest felt necessary to check constraints such as invariant loop bounds and disallow nesting. However, both properties are still checked now.

Just one thing. Preinits does not directly relate to CapturedStmt, you can consider it as a temp way to optimize the codegen. Also, it allows to support the use of data members as loop counters in classes/structures.

clang/include/clang/AST/StmtOpenMP.h

457

We can rename it to something like NumAssociatedLoops

660–686

It just hides all internal representation in the class and the user/caller should not worry about proper implementation of the loops processing. Consider it as a way to encapsulate representation details.

668

There is a type mismatch in callback types, Stmt * and const Stmt *

689–718

I'll check if it can be used instead, though I just want to be consistent with what we have for other statements/expressions/directives.

[Meinersbur]

clang/include/clang/AST/StmtOpenMP.h

457

Sounds good.

660–686

That could also be done by returning and iterator/filling an array with the requested data. The latter is one of the primary uses of SmallVectorImpl.

coroutines should make returning an iterator much simpler, eventually when we are allowed to use C++20.

668

I understand why the additional lambda is needed, but wondered about why the declaration is not

auto NewCallback = [Callback](unsigned Cnt, Stmt *CurStmt) {

689–718

It would reduce the maintenance cost of not forgetting to update this list as well.

[ABataev]

Just a note:

4.b. still feels like a hack because there are captures variables outside of any CapturedStmt and therefore complicated the AST. Comparable directive such as master don't do this.

It is not a hack, this was an optimization to minimize the number of recalculations of some expressions for the loops. For master or other directives we just did not need to do it at all.

clang/include/clang/AST/StmtOpenMP.h

660–686

I just prefer not to expose internals if they could be processed within the class instance.

668

I believe it just saves us from the extra constructor call.

[ABataev]

Update+Rebase

[Meinersbur]

Who is going to commit?

clang/include/clang/AST/StmtOpenMP.h

660–686

Using SmallVectorImpl for this is still the dominant style in clang. For instance:

• Sema::getUndefinedButUsed
• Sema::FindHiddenVirtualMethods
• Sema::CollectMultipleMethodsInGlobalPool
• Sema::SelectBestMethod
• Sema::CollectIvarsToConstructOrDestruct
• Sema::collectUnexpandedParameterPacks
• Sema::FindProtocolDeclaration
• Sema::GatherArgumentsForCall
• Sema::GatherGlobalCodeCompletions
• Sema::getUndefinedButUsed
• ASTContext::getOverriddenMethods
• ASTContext::DeepCollectObjCIvars
• ASTContext::getInjectedTemplateArgs
  
  ...

668

It does not: https://godbolt.org/z/81ac7o

[ABataev]

In D76342#2562097, @Meinersbur wrote:

Who is going to commit?

I can commit it as soon as you accepted it if you don't mind of course.

clang/include/clang/AST/StmtOpenMP.h

660–686

There is also another motivation behind this solution: to reuse the code as much as possible and avoid code duplication (loops to process inner loops/bodies) and focus just on loops/bodies processing.

668

Yes, currently it is not, it was before. But even now there is a difference in AST:

`-CompoundStmt <col:34, line:10:1>
  |-DeclStmt <line:4:5, line:6:6>
  | `-VarDecl <line:4:5, line:6:5> line:4:17 used Lambda '(lambda at line:4:26)':'(lambda at line:4:26)' cinit
  |   `-ExprWithCleanups <col:26, line:6:5> '(lambda at line:4:26)':'(lambda at line:4:26)'
  |     `-CXXConstructExpr <line:4:26, line:6:5> '(lambda at line:4:26)':'(lambda at line:4:26)' 'void ((lambda at line:4:26) &&) noexcept' elidable
  |       `-MaterializeTemporaryExpr <line:4:26, line:6:5> '(lambda at line:4:26)' xvalue

With &&:

`-CompoundStmt <col:34, line:10:1>
  |-DeclStmt <line:4:5, line:6:6>
  | `-VarDecl <line:4:5, line:6:5> line:4:17 used Lambda '(lambda at line:4:26) &&' cinit
  |   `-ExprWithCleanups <col:26, line:6:5> '(lambda at line:4:26)' xvalue
  |     `-MaterializeTemporaryExpr <line:4:26, line:6:5> '(lambda at line:4:26)' xvalue extended by Var 0x562226678ad8 'Lambda' '(lambda at line:4:26) &&'

Just the constructor call is marked as elidable and is not emitted anymore, but it still saves some memory/time because we don't need to emit an extra AST node.

[Meinersbur]

In D76342#2565686, @ABataev wrote:

I can commit it as soon as you accepted it if you don't mind of course.

I do not mind

[ABataev]

In D76342#2565945, @Meinersbur wrote:

In D76342#2565686, @ABataev wrote:

I can commit it as soon as you accepted it if you don't mind of course.

I do not mind

Ok, will commit it in a few minutes, thanks!

[clementval]

@Meinersbur / @ABataev This patch is making Flang buildbots failing.

Since you have added a clause in OMP.td you have to add lines here as well -> https://github.com/llvm/llvm-project/blob/adfd3c7083f9808d145239153c10f72eece485d8/flang/lib/Semantics/check-omp-structure.cpp#L669

Adding this line

CHECK_SIMPLE_CLAUSE(Sizes, OMPC_sizes)

After

CHECK_SIMPLE_CLAUSE(Simd, OMPC_simd)

Fixes the build problem

[clementval]

A fix is ready here https://reviews.llvm.org/D96808

Fix has landed so no need to revert.

clang/include/clang-c/Index.h

enum CXCursorKind {
   /** OpenMP depobj directive.
    */
   CXCursor_OMPDepobjDirective = 286,
 
   /** OpenMP scan directive.
    */
   CXCursor_OMPScanDirective = 287,
 
-  CXCursor_LastStmt = CXCursor_OMPScanDirective,
+  /** OpenMP tile directive.
+   */
+  CXCursor_OMPTileDirective = 288,
+
+  CXCursor_LastStmt = CXCursor_OMPTileDirective,
 
   /**
    * Cursor that represents the translation unit itself.
    *
    * The translation unit cursor exists primarily to act as the root
    * cursor for traversing the contents of a translation unit.
    */
   CXCursor_TranslationUnit = 300,

clang/include/clang/AST/OpenMPClause.h

const_child_range used_children() const {
     return const_child_range(const_child_iterator(), const_child_iterator());
   }
 
   static bool classof(const OMPClause *T) {
     return T->getClauseKind() == OMPC_exclusive;
   }
 };
 
+/// This represents the 'sizes' clause in the '#pragma omp tile' directive.
+///
+/// \code
+/// #pragma omp tile sizes(5,5)
+/// for (int i = 0; i < 64; ++i)
+///   for (int j = 0; j < 64; ++j)
+/// \endcode
+class OMPSizesClause final
+    : public OMPClause,
+      private llvm::TrailingObjects<OMPSizesClause, Expr *> {
+  friend class OMPClauseReader;
+  friend class llvm::TrailingObjects<OMPSizesClause, Expr *>;
+
+  /// Location of '('.
+  SourceLocation LParenLoc;
+
+  /// Number of tile sizes in the clause.
+  unsigned NumSizes;
+
+  /// Build an empty clause.
+  explicit OMPSizesClause(int NumSizes)
+      : OMPClause(OMPC_sizes, SourceLocation(), SourceLocation()),
+        NumSizes(NumSizes) {}
+
+public:
+  /// Build a 'sizes' AST node.
+  ///
+  /// \param C         Context of the AST.
+  /// \param StartLoc  Location of the 'sizes' identifier.
+  /// \param LParenLoc Location of '('.
+  /// \param EndLoc    Location of ')'.
+  /// \param Sizes     Content of the clause.
+  static OMPSizesClause *Create(const ASTContext &C, SourceLocation StartLoc,
+                                SourceLocation LParenLoc, SourceLocation EndLoc,
+                                ArrayRef<Expr *> Sizes);
+
+  /// Build an empty 'sizes' AST node for deserialization.
+  ///
+  /// \param C     Context of the AST.
+  /// \param Sizes Number of items in the clause.
+  static OMPSizesClause *CreateEmpty(const ASTContext &C, unsigned NumSizes);
+
+  /// Sets the location of '('.
+  void setLParenLoc(SourceLocation Loc) { LParenLoc = Loc; }
+
+  /// Returns the location of '('.
+  SourceLocation getLParenLoc() const { return LParenLoc; }
+
+  /// Returns the number of list items.
+  unsigned getNumSizes() const { return NumSizes; }
+
+  /// Returns the tile size expressions.
+  MutableArrayRef<Expr *> getSizesRefs() {
+    return MutableArrayRef<Expr *>(static_cast<OMPSizesClause *>(this)
+                                       ->template getTrailingObjects<Expr *>(),
+                                   NumSizes);
+  }
+  ArrayRef<Expr *> getSizesRefs() const {
+    return ArrayRef<Expr *>(static_cast<const OMPSizesClause *>(this)
+                                ->template getTrailingObjects<Expr *>(),
+                            NumSizes);
+  }
+
+  /// Sets the tile size expressions.
+  void setSizesRefs(ArrayRef<Expr *> VL) {
+    assert(VL.size() == NumSizes);
+    std::copy(VL.begin(), VL.end(),
+              static_cast<OMPSizesClause *>(this)
+                  ->template getTrailingObjects<Expr *>());
+  }
+
+  child_range children() {
+    MutableArrayRef<Expr *> Sizes = getSizesRefs();
+    return child_range(reinterpret_cast<Stmt **>(Sizes.begin()),
+                       reinterpret_cast<Stmt **>(Sizes.end()));
+  }
+  const_child_range children() const {
+    ArrayRef<Expr *> Sizes = getSizesRefs();
+    return const_child_range(reinterpret_cast<Stmt *const *>(Sizes.begin()),
+                             reinterpret_cast<Stmt *const *>(Sizes.end()));
+  }
+
+  child_range used_children() {
+    return child_range(child_iterator(), child_iterator());
+  }
+  const_child_range used_children() const {
+    return const_child_range(const_child_iterator(), const_child_iterator());
+  }
+
+  static bool classof(const OMPClause *T) {
+    return T->getClauseKind() == OMPC_sizes;
+  }
+};
+
 /// This class implements a simple visitor for OMPClause
 /// subclasses.
 template<class ImplClass, template <typename> class Ptr, typename RetTy>
 class OMPClauseVisitorBase {
 public:
 #define PTR(CLASS) Ptr<CLASS>
 #define DISPATCH(CLASS) \
   return static_cast<ImplClass*>(this)->Visit##CLASS(static_cast<PTR(CLASS)>(S))

clang/include/clang/AST/RecursiveASTVisitor.h

 //  traverses the entire AST.
 //
 //===----------------------------------------------------------------------===//
 #ifndef LLVM_CLANG_AST_RECURSIVEASTVISITOR_H
 #define LLVM_CLANG_AST_RECURSIVEASTVISITOR_H
 
 #include "clang/AST/Attr.h"
 #include "clang/AST/Decl.h"
-#include "clang/AST/DeclarationName.h"
 #include "clang/AST/DeclBase.h"
 #include "clang/AST/DeclCXX.h"
 #include "clang/AST/DeclFriend.h"
 #include "clang/AST/DeclObjC.h"
 #include "clang/AST/DeclOpenMP.h"
 #include "clang/AST/DeclTemplate.h"
+#include "clang/AST/DeclarationName.h"
 #include "clang/AST/Expr.h"
-#include "clang/AST/ExprConcepts.h"
 #include "clang/AST/ExprCXX.h"
+#include "clang/AST/ExprConcepts.h"
 #include "clang/AST/ExprObjC.h"
 #include "clang/AST/ExprOpenMP.h"
 #include "clang/AST/LambdaCapture.h"
 #include "clang/AST/NestedNameSpecifier.h"
 #include "clang/AST/OpenMPClause.h"
 #include "clang/AST/Stmt.h"
 #include "clang/AST/StmtCXX.h"
 #include "clang/AST/StmtObjC.h"

 }
 
 DEF_TRAVERSE_STMT(OMPParallelDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
 DEF_TRAVERSE_STMT(OMPSimdDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
 DEF_TRAVERSE_STMT(OMPForDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
 DEF_TRAVERSE_STMT(OMPForSimdDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
 DEF_TRAVERSE_STMT(OMPSectionsDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
 DEF_TRAVERSE_STMT(OMPSectionDirective,
DEF_TRAVERSE_STMT(OMPTargetTeamsDistributeParallelForDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
 DEF_TRAVERSE_STMT(OMPTargetTeamsDistributeParallelForSimdDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
 DEF_TRAVERSE_STMT(OMPTargetTeamsDistributeSimdDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
+DEF_TRAVERSE_STMT(OMPTileDirective,
+                  { TRY_TO(TraverseOMPExecutableDirective(S)); })
+
 // OpenMP clauses.
 template <typename Derived>
 bool RecursiveASTVisitor<Derived>::TraverseOMPClause(OMPClause *C) {
   if (!C)
     return true;
   switch (C->getClauseKind()) {
 #define OPENMP_CLAUSE(Name, Class)                                             \
   case OMPC_##Name:                                                            \

 }
 
 template <typename Derived>
 bool RecursiveASTVisitor<Derived>::VisitOMPOrderClause(OMPOrderClause *) {
   return true;
 }
 
 template <typename Derived>
+bool RecursiveASTVisitor<Derived>::VisitOMPSizesClause(OMPSizesClause *C) {
+  for (Expr *E : C->getSizesRefs())
+    TRY_TO(TraverseStmt(E));
+  return true;
+}
+
+template <typename Derived>
 bool RecursiveASTVisitor<Derived>::VisitOMPDetachClause(OMPDetachClause *C) {
   TRY_TO(TraverseStmt(C->getEventHandler()));
   return true;
 }
 
 // FIXME: look at the following tricky-seeming exprs to see if we
 // need to recurse on anything.  These are ones that have methods
 // returning decls or qualtypes or nestednamespecifier -- though I'm

clang/include/clang/AST/StmtOpenMP.h

protected:
   ///
   /// /param S Associated statement.
   ///
   void setAssociatedStmt(Stmt *S) {
     assert(hasAssociatedStmt() && "no associated statement.");
     *child_begin() = S;
   }
 
+  /// Returns the number of Stmt*'s trailing the object.
+  ///
+  /// This is not the number of elements returned by children().
+  unsigned getNumStmts() const { return NumChildren; }
+
 public:
   /// Iterates over expressions/statements used in the construct.
   class used_clauses_child_iterator
       : public llvm::iterator_adaptor_base<
             used_clauses_child_iterator, ArrayRef<OMPClause *>::iterator,
             std::forward_iterator_tag, Stmt *, ptrdiff_t, Stmt *, Stmt *> {
     ArrayRef<OMPClause *>::iterator End;
     OMPClause::child_iterator ChildI, ChildEnd;
Stmt *getAssociatedStmt() {
     return *child_begin();
   }
 
   /// Returns the captured statement associated with the
   /// component region within the (combined) directive.
   //
   // \param RegionKind Component region kind.
   const CapturedStmt *getCapturedStmt(OpenMPDirectiveKind RegionKind) const {
+    assert(!isOpenMPLoopTransformationDirective(RegionKind) &&
+           "Loop transformations do not capture");
     SmallVector<OpenMPDirectiveKind, 4> CaptureRegions;
     getOpenMPCaptureRegions(CaptureRegions, getDirectiveKind());
     assert(std::any_of(
                CaptureRegions.begin(), CaptureRegions.end(),
                [=](const OpenMPDirectiveKind K) { return K == RegionKind; }) &&
            "RegionKind not found in OpenMP CaptureRegions.");
     auto *CS = cast<CapturedStmt>(getAssociatedStmt());
     for (auto ThisCaptureRegion : CaptureRegions) {
       if (ThisCaptureRegion == RegionKind)
         return CS;
       CS = cast<CapturedStmt>(CS->getCapturedStmt());
     }
     llvm_unreachable("Incorrect RegionKind specified for directive.");
   }
 
   /// Get innermost captured statement for the construct.
   CapturedStmt *getInnermostCapturedStmt() {
+    assert(
+        !isOpenMPLoopTransformationDirective(getDirectiveKind()) &&
+        "Loop transformations do not capture; consider using ignoreCaptures()");
     assert(hasAssociatedStmt() && getAssociatedStmt() &&
            "Must have associated statement.");
     SmallVector<OpenMPDirectiveKind, 4> CaptureRegions;
     getOpenMPCaptureRegions(CaptureRegions, getDirectiveKind());
     assert(!CaptureRegions.empty() &&
            "At least one captured statement must be provided.");
     auto *CS = cast<CapturedStmt>(getAssociatedStmt());
     for (unsigned Level = CaptureRegions.size(); Level > 1; --Level)
       CS = cast<CapturedStmt>(CS->getCapturedStmt());
     return CS;
   }
 
   const CapturedStmt *getInnermostCapturedStmt() const {
     return const_cast<OMPExecutableDirective *>(this)
         ->getInnermostCapturedStmt();
   }
 
+  /// Return the associated code without captures.
+  ///
+  /// Methods with similar functionality:
+  ///
+  ///  * IgnoreContainers(true) - Not specific to OpenMP; also ignores
+  ///  containers/captures potentially inside the associated code.
+  ///
+  ///  * getCapturedStmt(Kind) - Return the capture of the specified kind;
+  ///  requires the captured decl/stmt to exist.
+  ///
+  ///  * getInnermostCapturedStmt() - Returns the innermost capture containing
+  ///  the associated code; also cannot be used if there is no CapturedStmt,
+  ///  e.g. for loop transformations.
+  ///
+  ///  * getBody() - In addition to ignoring captures, also skips associated
+  ///  loops.
+  ///
+  ///  * getStructuredBlock() - For loop-associated directives, returns the
+  ///  body; for others, returns the captured statement.
+  ///
+  Stmt *ignoreCaptures();
+  const Stmt *ignoreCaptures() const {
+    return const_cast<OMPExecutableDirective *>(this)->ignoreCaptures();
+  }
+
   OpenMPDirectiveKind getDirectiveKind() const { return Kind; }
 
   static bool classof(const Stmt *S) {
     return S->getStmtClass() >= firstOMPExecutableDirectiveConstant &&
            S->getStmtClass() <= lastOMPExecutableDirectiveConstant;
   }
 
   child_range children() {
public:
   /// Return true if current directive has inner cancel directive.
   bool hasCancel() const { return HasCancel; }
 
   static bool classof(const Stmt *T) {
     return T->getStmtClass() == OMPParallelDirectiveClass;
   }
 };
 
 /// This is a common base class for loop directives ('omp simd', 'omp
 /// for', 'omp for simd' etc.). It is responsible for the loop code generation.
 ///
 class OMPLoopDirective : public OMPExecutableDirective {
   friend class ASTStmtReader;
   /// Number of collapsed loops as specified by 'collapse' clause.
   unsigned CollapsedNum;
 
   /// Offsets to the stored exprs.
   /// This enumeration contains offsets to all the pointers to children
   /// expressions stored in OMPLoopDirective.
   /// The first 9 children are necessary for all the loop directives,
   /// the next 8 are specific to the worksharing ones, and the next 11 are
   /// used for combined constructs containing two pragmas associated to loops.
   /// After the fixed children, three arrays of length CollapsedNum are
   /// allocated: loop counters, their updates and final values.
   /// PrevLowerBound and PrevUpperBound are used to communicate blocking
   /// information in composite constructs which require loop blocking
   /// DistInc is used to generate the increment expression for the distribute
   /// loop when combined with a further nested loop
   /// PrevEnsureUpperBound is used as the EnsureUpperBound expression for the
   /// for loop when combined with a previous distribute loop in the same pragma
   /// (e.g. 'distribute parallel for')
   ///
+  /// Loop transformations, such as OMPTileDirective, only use
+  /// AssociatedStmtOffset. Its specific statments/expressions are stored as
+  /// object members.
+  ///
   enum {
     AssociatedStmtOffset = 0,
     IterationVariableOffset = 1,
     LastIterationOffset = 2,
     CalcLastIterationOffset = 3,
     PreConditionOffset = 4,
     CondOffset = 5,
     InitOffset = 6,
enum {
     // Offset to the end (and start of the following
     // counters/updates/finals/dependent_counters/dependent_inits/finals_conditions
     // arrays) for combined distribute loop directives.
     CombinedDistributeEnd = 30,
   };
 
   /// Get the counters storage.
   MutableArrayRef<Expr *> getCounters() {
+    assert(!isOpenMPLoopTransformationDirective(getDirectiveKind()));
     Expr **Storage = reinterpret_cast<Expr **>(
         &(*(std::next(child_begin(), getArraysOffset(getDirectiveKind())))));
     return MutableArrayRef<Expr *>(Storage, CollapsedNum);
   }
 
   /// Get the private counters storage.
   MutableArrayRef<Expr *> getPrivateCounters() {
+    assert(!isOpenMPLoopTransformationDirective(getDirectiveKind()));
     Expr **Storage = reinterpret_cast<Expr **>(&*std::next(
         child_begin(), getArraysOffset(getDirectiveKind()) + CollapsedNum));
     return MutableArrayRef<Expr *>(Storage, CollapsedNum);
   }
 
   /// Get the updates storage.
   MutableArrayRef<Expr *> getInits() {
+    assert(!isOpenMPLoopTransformationDirective(getDirectiveKind()));
     Expr **Storage = reinterpret_cast<Expr **>(
         &*std::next(child_begin(),
                     getArraysOffset(getDirectiveKind()) + 2 * CollapsedNum));
     return MutableArrayRef<Expr *>(Storage, CollapsedNum);
   }
 
   /// Get the updates storage.
   MutableArrayRef<Expr *> getUpdates() {
+    assert(!isOpenMPLoopTransformationDirective(getDirectiveKind()));
     Expr **Storage = reinterpret_cast<Expr **>(
         &*std::next(child_begin(),
                     getArraysOffset(getDirectiveKind()) + 3 * CollapsedNum));
     return MutableArrayRef<Expr *>(Storage, CollapsedNum);
   }
 
   /// Get the final counter updates storage.
   MutableArrayRef<Expr *> getFinals() {
+    assert(!isOpenMPLoopTransformationDirective(getDirectiveKind()));
     Expr **Storage = reinterpret_cast<Expr **>(
         &*std::next(child_begin(),
                     getArraysOffset(getDirectiveKind()) + 4 * CollapsedNum));
     return MutableArrayRef<Expr *>(Storage, CollapsedNum);
   }
 
   /// Get the dependent counters storage.
   MutableArrayRef<Expr *> getDependentCounters() {
+    assert(!isOpenMPLoopTransformationDirective(getDirectiveKind()));
     Expr **Storage = reinterpret_cast<Expr **>(
         &*std::next(child_begin(),
                     getArraysOffset(getDirectiveKind()) + 5 * CollapsedNum));
     return MutableArrayRef<Expr *>(Storage, CollapsedNum);
   }
 
   /// Get the dependent inits storage.
   MutableArrayRef<Expr *> getDependentInits() {
+    assert(!isOpenMPLoopTransformationDirective(getDirectiveKind()));
     Expr **Storage = reinterpret_cast<Expr **>(
         &*std::next(child_begin(),
                     getArraysOffset(getDirectiveKind()) + 6 * CollapsedNum));
     return MutableArrayRef<Expr *>(Storage, CollapsedNum);
   }
 
   /// Get the finals conditions storage.
   MutableArrayRef<Expr *> getFinalsConditions() {
+    assert(!isOpenMPLoopTransformationDirective(getDirectiveKind()));
     Expr **Storage = reinterpret_cast<Expr **>(
         &*std::next(child_begin(),
                     getArraysOffset(getDirectiveKind()) + 7 * CollapsedNum));
     return MutableArrayRef<Expr *>(Storage, CollapsedNum);
   }
 
 protected:
   /// Build instance of loop directive of class \a Kind.
OMPLoopDirective(const T *That, StmtClass SC, OpenMPDirectiveKind Kind,
                    unsigned NumSpecialChildren = 0)
       : OMPExecutableDirective(That, SC, Kind, StartLoc, EndLoc, NumClauses,
                                numLoopChildren(CollapsedNum, Kind) +
                                    NumSpecialChildren),
         CollapsedNum(CollapsedNum) {}
 
   /// Offset to the start of children expression arrays.
   static unsigned getArraysOffset(OpenMPDirectiveKind Kind) {
+    if (isOpenMPLoopTransformationDirective(Kind))
+      return AssociatedStmtOffset + 1;
     if (isOpenMPLoopBoundSharingDirective(Kind))
       return CombinedDistributeEnd;
     if (isOpenMPWorksharingDirective(Kind) || isOpenMPTaskLoopDirective(Kind) ||
         isOpenMPDistributeDirective(Kind))
       return WorksharingEnd;
     return DefaultEnd;
   }
 
   /// Children number.
   static unsigned numLoopChildren(unsigned CollapsedNum,
                                   OpenMPDirectiveKind Kind) {
+    if (isOpenMPLoopTransformationDirective(Kind))
+      return getArraysOffset(Kind);
     return getArraysOffset(Kind) +
            8 * CollapsedNum; // Counters, PrivateCounters, Inits,
                              // Updates, Finals, DependentCounters,
                              // DependentInits, FinalsConditions.
   }
 
   void setIterationVariable(Expr *IV) {
     *std::next(child_begin(), IterationVariableOffset) = IV;
void setCond(Expr *Cond) {
     *std::next(child_begin(), CondOffset) = Cond;
   }
   void setInit(Expr *Init) { *std::next(child_begin(), InitOffset) = Init; }
   void setInc(Expr *Inc) { *std::next(child_begin(), IncOffset) = Inc; }
   void setPreInits(Stmt *PreInits) {
     *std::next(child_begin(), PreInitsOffset) = PreInits;
   }
   void setIsLastIterVariable(Expr *IL) {
     assert((isOpenMPWorksharingDirective(getDirectiveKind()) ||
             isOpenMPTaskLoopDirective(getDirectiveKind()) ||
             isOpenMPDistributeDirective(getDirectiveKind())) &&
            "expected worksharing loop directive");
     *std::next(child_begin(), IsLastIterVariableOffset) = IL;
   }
   void setLowerBoundVariable(Expr *LB) {
     assert((isOpenMPWorksharingDirective(getDirectiveKind()) ||
             isOpenMPTaskLoopDirective(getDirectiveKind()) ||
             isOpenMPDistributeDirective(getDirectiveKind())) &&
            "expected worksharing loop directive");
     *std::next(child_begin(), LowerBoundVariableOffset) = LB;
   }
   void setUpperBoundVariable(Expr *UB) {
     assert((isOpenMPWorksharingDirective(getDirectiveKind()) ||
             isOpenMPTaskLoopDirective(getDirectiveKind()) ||
             isOpenMPDistributeDirective(getDirectiveKind())) &&
            "expected worksharing loop directive");
     *std::next(child_begin(), UpperBoundVariableOffset) = UB;
   }
   void setStrideVariable(Expr *ST) {
     assert((isOpenMPWorksharingDirective(getDirectiveKind()) ||
             isOpenMPTaskLoopDirective(getDirectiveKind()) ||
             isOpenMPDistributeDirective(getDirectiveKind())) &&
            "expected worksharing loop directive");
     *std::next(child_begin(), StrideVariableOffset) = ST;
   }
   void setEnsureUpperBound(Expr *EUB) {
     assert((isOpenMPWorksharingDirective(getDirectiveKind()) ||
             isOpenMPTaskLoopDirective(getDirectiveKind()) ||
             isOpenMPDistributeDirective(getDirectiveKind())) &&
            "expected worksharing loop directive");
     *std::next(child_begin(), EnsureUpperBoundOffset) = EUB;
   }
   void setNextLowerBound(Expr *NLB) {
     assert((isOpenMPWorksharingDirective(getDirectiveKind()) ||
             isOpenMPTaskLoopDirective(getDirectiveKind()) ||
             isOpenMPDistributeDirective(getDirectiveKind())) &&
            "expected worksharing loop directive");
     *std::next(child_begin(), NextLowerBoundOffset) = NLB;
   }
   void setNextUpperBound(Expr *NUB) {
     assert((isOpenMPWorksharingDirective(getDirectiveKind()) ||
             isOpenMPTaskLoopDirective(getDirectiveKind()) ||
             isOpenMPDistributeDirective(getDirectiveKind())) &&
            "expected worksharing loop directive");
     *std::next(child_begin(), NextUpperBoundOffset) = NUB;
   }
   void setNumIterations(Expr *NI) {
     assert((isOpenMPWorksharingDirective(getDirectiveKind()) ||
             isOpenMPTaskLoopDirective(getDirectiveKind()) ||
             isOpenMPDistributeDirective(getDirectiveKind())) &&
            "expected worksharing loop directive");
     *std::next(child_begin(), NumIterationsOffset) = NI;
   }
   void setPrevLowerBoundVariable(Expr *PrevLB) {
     assert(isOpenMPLoopBoundSharingDirective(getDirectiveKind()) &&
            "expected loop bound sharing directive");
     *std::next(child_begin(), PrevLowerBoundVariableOffset) = PrevLB;
   }
   void setPrevUpperBoundVariable(Expr *PrevUB) {
     assert(isOpenMPLoopBoundSharingDirective(getDirectiveKind()) &&
            "expected loop bound sharing directive");
     *std::next(child_begin(), PrevUpperBoundVariableOffset) = PrevUB;
   }
   void setDistInc(Expr *DistInc) {
     assert(isOpenMPLoopBoundSharingDirective(getDirectiveKind()) &&
struct HelperExprs {
     /// expression IV = IV + ST
     Expr *DistInc;
     /// PrevEUB - expression similar to EUB but to be used when loop
     /// scheduling uses PrevLB and PrevUB (e.g.  in 'distribute parallel for'
     /// when ensuring that the UB is either the calculated UB by the runtime or
     /// the end of the assigned distribute chunk)
     /// expression UB = min (UB, PrevUB)
     Expr *PrevEUB;
+    /// The associated loops.
+    SmallVector<Stmt *, 4> Loops;
+    /// The body of each associated loop.
+    SmallVector<Stmt *, 4> Bodys;
     /// Counters Loop counters.
     SmallVector<Expr *, 4> Counters;
     /// PrivateCounters Loop counters.
     SmallVector<Expr *, 4> PrivateCounters;
     /// Expressions for loop counters inits for CodeGen.
     SmallVector<Expr *, 4> Inits;
     /// Expressions for loop counters update for CodeGen.
     SmallVector<Expr *, 4> Updates;
void clear(unsigned Size) {
       Counters.resize(Size);
       PrivateCounters.resize(Size);
       Inits.resize(Size);
       Updates.resize(Size);
       Finals.resize(Size);
       DependentCounters.resize(Size);
       DependentInits.resize(Size);
       FinalsConditions.resize(Size);
+      Loops.resize(Size);
+      Bodys.resize(Size);
       for (unsigned i = 0; i < Size; ++i) {
         Counters[i] = nullptr;
         PrivateCounters[i] = nullptr;
         Inits[i] = nullptr;
         Updates[i] = nullptr;
         Finals[i] = nullptr;
         DependentCounters[i] = nullptr;
         DependentInits[i] = nullptr;
         FinalsConditions[i] = nullptr;
+        Loops[i] = nullptr;
+        Bodys[i] = nullptr;
       }
       PreInits = nullptr;
       DistCombinedFields.LB = nullptr;
       DistCombinedFields.UB = nullptr;
       DistCombinedFields.EUB = nullptr;
       DistCombinedFields.Init = nullptr;
       DistCombinedFields.Cond = nullptr;
       DistCombinedFields.NLB = nullptr;
       DistCombinedFields.NUB = nullptr;
       DistCombinedFields.DistCond = nullptr;
       DistCombinedFields.ParForInDistCond = nullptr;
     }
   };
 
   /// Get number of collapsed loops.
   unsigned getCollapsedNumber() const { return CollapsedNum; }
 
   Expr *getIterationVariable() const {
+    assert(!isOpenMPLoopTransformationDirective(getDirectiveKind()));
     return const_cast<Expr *>(reinterpret_cast<const Expr *>(
         *std::next(child_begin(), IterationVariableOffset)));
   }
   Expr *getLastIteration() const {
+    assert(!isOpenMPLoopTransformationDirective(getDirectiveKind()));
     return const_cast<Expr *>(reinterpret_cast<const Expr *>(
         *std::next(child_begin(), LastIterationOffset)));
   }
   Expr *getCalcLastIteration() const {
+    assert(!isOpenMPLoopTransformationDirective(getDirectiveKind()));
     return const_cast<Expr *>(reinterpret_cast<const Expr *>(
         *std::next(child_begin(), CalcLastIterationOffset)));
   }
   Expr *getPreCond() const {
+    assert(!isOpenMPLoopTransformationDirective(getDirectiveKind()));
     return const_cast<Expr *>(reinterpret_cast<const Expr *>(
         *std::next(child_begin(), PreConditionOffset)));
   }
   Expr *getCond() const {
+    assert(!isOpenMPLoopTransformationDirective(getDirectiveKind()));
     return const_cast<Expr *>(
         reinterpret_cast<const Expr *>(*std::next(child_begin(), CondOffset)));
   }
   Expr *getInit() const {
+    assert(!isOpenMPLoopTransformationDirective(getDirectiveKind()));
     return const_cast<Expr *>(
         reinterpret_cast<const Expr *>(*std::next(child_begin(), InitOffset)));
   }
   Expr *getInc() const {
+    assert(!isOpenMPLoopTransformationDirective(getDirectiveKind()));
     return const_cast<Expr *>(
         reinterpret_cast<const Expr *>(*std::next(child_begin(), IncOffset)));
   }
   const Stmt *getPreInits() const {
+    assert(!isOpenMPLoopTransformationDirective(getDirectiveKind()));
+    return *std::next(child_begin(), PreInitsOffset);
+  }
+  Stmt *getPreInits() {
+    assert(!isOpenMPLoopTransformationDirective(getDirectiveKind()));
     return *std::next(child_begin(), PreInitsOffset);
   }
-  Stmt *getPreInits() { return *std::next(child_begin(), PreInitsOffset); }
   Expr *getIsLastIterVariable() const {
     assert((isOpenMPWorksharingDirective(getDirectiveKind()) ||
             isOpenMPTaskLoopDirective(getDirectiveKind()) ||
             isOpenMPDistributeDirective(getDirectiveKind())) &&
            "expected worksharing loop directive");
     return const_cast<Expr *>(reinterpret_cast<const Expr *>(
         *std::next(child_begin(), IsLastIterVariableOffset)));
   }
assert(isOpenMPLoopBoundSharingDirective(getDirectiveKind()) &&
            "expected loop bound distribute sharing directive");
     return const_cast<Expr *>(reinterpret_cast<const Expr *>(
         *std::next(child_begin(), CombinedParForInDistConditionOffset)));
   }
   /// Try to find the next loop sub-statement in the specified statement \p
   /// CurStmt.
   /// \param TryImperfectlyNestedLoops true, if we need to try to look for the
   /// imperfectly nested loop.
-  static Stmt *tryToFindNextInnerLoop(Stmt *CurStmt,
-                                      bool TryImperfectlyNestedLoops);
-  static const Stmt *tryToFindNextInnerLoop(const Stmt *CurStmt,
-                                            bool TryImperfectlyNestedLoops) {
+  /// \param PreInits If set, collects additional statemants that have to be
+  ///                 executed before the loop; typically variable declarations.
+  static Stmt *
+  tryToFindNextInnerLoop(Stmt *CurStmt, bool TryImperfectlyNestedLoops,
+                         llvm::SmallVectorImpl<Stmt *> *PreInits = nullptr);
+  static const Stmt *
+  tryToFindNextInnerLoop(const Stmt *CurStmt, bool TryImperfectlyNestedLoops,
+                         llvm::SmallVectorImpl<Stmt *> *PreInits = nullptr) {
     return tryToFindNextInnerLoop(const_cast<Stmt *>(CurStmt),
-                                  TryImperfectlyNestedLoops);
+                                  TryImperfectlyNestedLoops, PreInits);
+  }
+
+  /// Collect the loops associated to this directive, taking loop
+  /// transformations into account.
+  ///
+  /// \param Loops    Receives the associated loops.
+  /// \param PreInits Receives statements that have to be executed before the
+  ///                 loop; typically variable declarations.
+  void collectAssociatedLoops(llvm::SmallVectorImpl<Stmt *> &Loops,
+                              llvm::SmallVectorImpl<Stmt *> &PreInits);
+  void collectAssociatedLoops(llvm::SmallVectorImpl<const Stmt *> &Loops,
+                              llvm::SmallVectorImpl<Stmt *> &PreInits) const {
+    const_cast<OMPLoopDirective *>(this)->collectAssociatedLoops(
+        *reinterpret_cast<llvm::SmallVectorImpl<Stmt *> *>(&Loops), PreInits);
   }
+
   Stmt *getBody();
   const Stmt *getBody() const {
     return const_cast<OMPLoopDirective *>(this)->getBody();
   }
 
   ArrayRef<Expr *> counters() { return getCounters(); }
 
   ArrayRef<Expr *> counters() const {
public:
 
   ArrayRef<Expr *> finals_conditions() const {
     return const_cast<OMPLoopDirective *>(this)->getFinalsConditions();
   }
 
   static bool classof(const Stmt *T) {
     return T->getStmtClass() == OMPSimdDirectiveClass ||
            T->getStmtClass() == OMPForDirectiveClass ||
+           T->getStmtClass() == OMPTileDirectiveClass ||
            T->getStmtClass() == OMPForSimdDirectiveClass ||
            T->getStmtClass() == OMPParallelForDirectiveClass ||
            T->getStmtClass() == OMPParallelForSimdDirectiveClass ||
            T->getStmtClass() == OMPTaskLoopDirectiveClass ||
            T->getStmtClass() == OMPTaskLoopSimdDirectiveClass ||
            T->getStmtClass() == OMPMasterTaskLoopDirectiveClass ||
            T->getStmtClass() == OMPMasterTaskLoopSimdDirectiveClass ||
            T->getStmtClass() == OMPParallelMasterTaskLoopDirectiveClass ||
public:
   CreateEmpty(const ASTContext &C, unsigned NumClauses, unsigned CollapsedNum,
               EmptyShell);
 
   static bool classof(const Stmt *T) {
     return T->getStmtClass() == OMPTargetTeamsDistributeSimdDirectiveClass;
   }
 };
 
+/// This represents the '#pragma omp tile' loop transformation directive.
+class OMPTileDirective final
+    : public OMPLoopDirective,
+      private llvm::TrailingObjects<OMPTileDirective, OMPClause *, Stmt *> {
+  friend class ASTStmtReader;
+  friend TrailingObjects;
+
+  /// The loop nest after being tiled.
+  Stmt *TransformedStmt = nullptr;
+
+  size_t numTrailingObjects(OverloadToken<OMPClause *>) const {
+    return getNumClauses();
+  }
+
+  size_t numTrailingObjects(OverloadToken<Stmt *>) const {
+    return getNumStmts();
+  }
+
+  explicit OMPTileDirective(SourceLocation StartLoc, SourceLocation EndLoc,
+                            unsigned NumClauses, unsigned NumLoops)
+      : OMPLoopDirective(this, OMPTileDirectiveClass, llvm::omp::OMPD_tile,
+                         StartLoc, EndLoc, NumLoops, NumClauses,
+                         /*NumSpecialChildren=*/0) {}
+
+public:
+  /// Create a new AST node represention '#pragma omp tile'.
+  ///
+  /// \param C         Context of the AST.
+  /// \param StartLoc  Location of the introducer (e.g. the 'omp' token).
+  /// \param EndLoc    Location of the directive's end (e.g. the tok::eod).
+  /// \param Clauses   The directive's clauses.
+  /// \param NumLoops  Number of associated loops (number of items in the
+  ///                  'sizes' clause). \param AssociatedStmt The outermost
+  ///                  associated loop.
+  /// \param TransformedStmt The loop nest after tiling, or nullptr in
+  ///                        dependenct contexts.
+  static OMPTileDirective *Create(const ASTContext &C, SourceLocation StartLoc,
+                                  SourceLocation EndLoc,
+                                  ArrayRef<OMPClause *> Clauses,
+                                  unsigned NumLoops, Stmt *AssociatedStmt,
+                                  Stmt *TransformedStmt);
+
+  /// Build an empty '#pragma omp tile' AST node for deserialization.
+  ///
+  /// \param C          Context of the AST.
+  /// \param NumClauses Number of clauses to allocate.
+  /// \param NumLoops   Number of associated loops to allocate.
+  static OMPTileDirective *CreateEmpty(const ASTContext &C, unsigned NumClauses,
+                                       unsigned NumLoops);
+
+  unsigned getNumAssociatedLoops() const { return getCollapsedNumber(); }
+
+  /// Gets/sets the associated loops after tiling.
+  ///
+  /// This is in de-sugared format stored as a CompoundStmt.
+  ///
+  /// \code
+  /// {
+  ///   PreInits;
+  ///   for (...)
+  ///     ...
+  /// }
+  /// \endcode
+  ///
+  /// Note that if the generated loops a become associated loops of another
+  /// directive, they may need to be hoisted before them.
+  Stmt *getTransformedStmt() const { return TransformedStmt; }
+  void setTransformedStmt(Stmt *S) {
+    assert(!S || isa<CompoundStmt>(S));
+    TransformedStmt = S;
+  }
+
+  /// Return the pre-init statements.
+  ///
+  /// These must be executed before the loop. Typically, these are declarations
+  /// for compiler-introduced variables.
+  auto getPreInits() const {
+    CompoundStmt *C = cast<CompoundStmt>(TransformedStmt);
+    return llvm::make_range(C->body_begin(), C->body_begin() + C->size() - 1);
+  }
+
+  /// Return the transformed (tiled) for-loop.
+  Stmt *getTransformedForStmt() {
+    CompoundStmt *C = cast<CompoundStmt>(TransformedStmt);
+    return C->body_back();
+  }
+
+  static bool classof(const Stmt *T) {
+    return T->getStmtClass() == OMPTileDirectiveClass;
+  }
+};
+
+/// Determine the first associated statement of a loop-associated directive. If
+/// the statement is a loop transformation directive, uses the result of the
+/// loop transformation.
+Stmt *
+getTopmostAssociatedStructuredBlock(Stmt *S,
+                                    llvm::SmallVectorImpl<Stmt *> *PreInits);
+static inline const Stmt *
+getTopmostAssociatedStructuredBlock(const Stmt *S,
+                                    llvm::SmallVectorImpl<Stmt *> *PreInits) {
+  return getTopmostAssociatedStructuredBlock(const_cast<Stmt *>(S), PreInits);
+}
+
 /// This represents '#pragma omp scan' directive.
 ///
 /// \code
 /// #pragma omp scan inclusive(a)
 /// \endcode
 /// In this example directive '#pragma omp scan' has clause 'inclusive' with
 /// list item 'a'.
 class OMPScanDirective final : public OMPExecutableDirective {

clang/include/clang/Basic/DiagnosticCommonKinds.td

 // HIP
 def warn_ignored_hip_only_option : Warning<
   "'%0' is ignored since it is only supported for HIP">,
   InGroup<HIPOnly>;
 
 // OpenMP
 def err_omp_more_one_clause : Error<
   "directive '#pragma omp %0' cannot contain more than one '%1' clause%select{| with '%3' name modifier| with 'source' dependence}2">;
+def err_omp_required_clause : Error<
+  "directive '#pragma omp %0' requires the '%1' clause">;
 
 // Static Analyzer Core
 def err_unknown_analyzer_checker_or_package : Error<
     "no analyzer checkers or packages are associated with '%0'">;
 def note_suggest_disabling_all_checkers : Note<
     "use -analyzer-disable-all-checks to disable all static analyzer checkers">;
 
 // Poison system directories.
 def warn_poison_system_directories : Warning <
   "include location '%0' is unsafe for cross-compilation">,
   InGroup<DiagGroup<"poison-system-directories">>, DefaultIgnore;
 }

clang/include/clang/Basic/DiagnosticSemaKinds.td

def err_expected_allocator_expression : Error<"expected an allocator expression "
   " directive with the 'dynamic_allocators' clause">;
 def warn_omp_allocate_thread_on_task_target_directive : Warning<
   "allocator with the 'thread' trait access has unspecified behavior on '%0' directive">,
   InGroup<OpenMPClauses>;
 def err_omp_expected_private_copy_for_allocate : Error<
   "the referenced item is not found in any private clause on the same directive">;
 def err_omp_stmt_depends_on_loop_counter : Error<
   "the loop %select{initializer|condition}0 expression depends on the current loop control variable">;
+def err_omp_invariant_dependency : Error<
+  "expected loop invariant expression">;
 def err_omp_invariant_or_linear_dependency : Error<
   "expected loop invariant expression or '<invariant1> * %0 + <invariant2>' kind of expression">;
 def err_omp_wrong_dependency_iterator_type : Error<
   "expected an integer or a pointer type of the outer loop counter '%0' for non-rectangular nests">;
 def err_omp_unsupported_type : Error <
   "host requires %0 bit size %1 type support, but device '%2' does not support it">;
 def err_omp_lambda_capture_in_declare_target_not_to : Error<
   "variable captured in declare target region must appear in a to clause">;

clang/include/clang/Basic/OpenMPKinds.h

 
 /// OpenMP attributes for 'order' clause.
 enum OpenMPOrderClauseKind {
 #define OPENMP_ORDER_KIND(Name) OMPC_ORDER_##Name,
 #include "clang/Basic/OpenMPKinds.def"
   OMPC_ORDER_unknown,
 };
 
+enum OpenMPTileClauseKind {
+#define OPENMP_TILE_KIND(Name) OMPC_TILE_##Name,
+#include "clang/Basic/OpenMPKinds.def"
+  OMPC_TILE_unknown,
+};
+
 /// Scheduling data for loop-based OpenMP directives.
 struct OpenMPScheduleTy final {
   OpenMPScheduleClauseKind Schedule = OMPC_SCHEDULE_unknown;
   OpenMPScheduleClauseModifier M1 = OMPC_SCHEDULE_MODIFIER_unknown;
   OpenMPScheduleClauseModifier M2 = OMPC_SCHEDULE_MODIFIER_unknown;
 };
 
 /// OpenMP modifiers for 'reduction' clause.

 /// taskloop simd, or parallel master taskloop simd.
 bool isOpenMPTaskingDirective(OpenMPDirectiveKind Kind);
 
 /// Checks if the specified directive kind is one of the composite or combined
 /// directives that need loop bound sharing across loops outlined in nested
 /// functions
 bool isOpenMPLoopBoundSharingDirective(OpenMPDirectiveKind Kind);
 
+/// Checks if the specified directive is a loop transformation directive.
+/// \param DKind Specified directive.
+/// \return True iff the directive is a loop transformation.
+bool isOpenMPLoopTransformationDirective(OpenMPDirectiveKind DKind);
+
 /// Return the captured regions of an OpenMP directive.
 void getOpenMPCaptureRegions(
     llvm::SmallVectorImpl<OpenMPDirectiveKind> &CaptureRegions,
     OpenMPDirectiveKind DKind);
 }
 
 #endif
 

clang/include/clang/Basic/OpenMPKinds.def

 #define OPENMP_DEVICE_MODIFIER(Name)
 #endif
 #ifndef OPENMP_SCAN_CLAUSE
 #define OPENMP_SCAN_CLAUSE(Name)
 #endif
 #ifndef OPENMP_REDUCTION_MODIFIER
 #define OPENMP_REDUCTION_MODIFIER(Name)
 #endif
+#ifndef OPENMP_TILE_CLAUSE
+#  define OPENMP_TILE_CLAUSE(Name)
+#endif
 
 // OpenMP clauses.
 OPENMP_CLAUSE(allocator, OMPAllocatorClause)
 OPENMP_CLAUSE(if, OMPIfClause)
 OPENMP_CLAUSE(final, OMPFinalClause)
 OPENMP_CLAUSE(num_threads, OMPNumThreadsClause)
 OPENMP_CLAUSE(safelen, OMPSafelenClause)
 OPENMP_CLAUSE(simdlen, OMPSimdlenClause)

 OPENMP_CLAUSE(allocate, OMPAllocateClause)
 OPENMP_CLAUSE(nontemporal, OMPNontemporalClause)
 OPENMP_CLAUSE(order, OMPOrderClause)
 OPENMP_CLAUSE(depobj, OMPDepobjClause)
 OPENMP_CLAUSE(destroy, OMPDestroyClause)
 OPENMP_CLAUSE(detach, OMPDetachClause)
 OPENMP_CLAUSE(inclusive, OMPInclusiveClause)
 OPENMP_CLAUSE(exclusive, OMPExclusiveClause)
+OPENMP_CLAUSE(sizes, OMPSizesClause)
 
 // Clauses allowed for OpenMP directive 'scan'.
 OPENMP_SCAN_CLAUSE(inclusive)
 OPENMP_SCAN_CLAUSE(exclusive)
 
 // Clauses allowed for OpenMP directive 'parallel'.
 OPENMP_PARALLEL_CLAUSE(if)
 OPENMP_PARALLEL_CLAUSE(num_threads)

 // Clauses allowed for OpenMP directive 'depobj'.
 OPENMP_DEPOBJ_CLAUSE(depend)
 OPENMP_DEPOBJ_CLAUSE(destroy)
 OPENMP_DEPOBJ_CLAUSE(update)
 
 // Modifiers for 'reduction' clause.
 OPENMP_REDUCTION_MODIFIER(default)
 
+// Clauses allowed for directive 'omp tile'.
+OPENMP_TILE_CLAUSE(sizes)
+
 #undef OPENMP_REDUCTION_MODIFIER
 #undef OPENMP_SCAN_CLAUSE
 #undef OPENMP_DEVICE_MODIFIER
 #undef OPENMP_DEPOBJ_CLAUSE
 #undef OPENMP_FLUSH_CLAUSE
 #undef OPENMP_ORDER_KIND
 #undef OPENMP_LASTPRIVATE_KIND
 #undef OPENMP_DECLARE_VARIANT_CLAUSE

 #undef OPENMP_TEAMS_DISTRIBUTE_SIMD_CLAUSE
 #undef OPENMP_TEAMS_DISTRIBUTE_PARALLEL_FOR_SIMD_CLAUSE
 #undef OPENMP_TEAMS_DISTRIBUTE_PARALLEL_FOR_CLAUSE
 #undef OPENMP_TARGET_TEAMS_CLAUSE
 #undef OPENMP_TARGET_TEAMS_DISTRIBUTE_CLAUSE
 #undef OPENMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_FOR_CLAUSE
 #undef OPENMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_FOR_SIMD_CLAUSE
 #undef OPENMP_TARGET_TEAMS_DISTRIBUTE_SIMD_CLAUSE
+#undef OPENMP_TILE_CLAUSE

clang/include/clang/Basic/StmtNodes.td

 def OMPTeamsDistributeSimdDirective : StmtNode<OMPLoopDirective>;
 def OMPTeamsDistributeParallelForSimdDirective : StmtNode<OMPLoopDirective>;
 def OMPTeamsDistributeParallelForDirective : StmtNode<OMPLoopDirective>;
 def OMPTargetTeamsDirective : StmtNode<OMPExecutableDirective>;
 def OMPTargetTeamsDistributeDirective : StmtNode<OMPLoopDirective>;
 def OMPTargetTeamsDistributeParallelForDirective : StmtNode<OMPLoopDirective>;
 def OMPTargetTeamsDistributeParallelForSimdDirective : StmtNode<OMPLoopDirective>;
 def OMPTargetTeamsDistributeSimdDirective : StmtNode<OMPLoopDirective>;
+def OMPTileDirective : StmtNode<OMPLoopDirective>;

clang/include/clang/Parse/Parser.h

private:
   /// \param DKind Directive kind.
   /// \param Kind Kind of current clause.
   /// \param ParseOnly true to skip the clause's semantic actions and return
   /// nullptr.
   ///
   OMPClause *ParseOpenMPSingleExprWithArgClause(OpenMPDirectiveKind DKind,
                                                 OpenMPClauseKind Kind,
                                                 bool ParseOnly);
+
+  /// Parses the 'sizes' clause of a '#pragma omp tile' directive.
+  OMPClause *ParseOpenMPSizesClause();
+
   /// Parses clause without any additional arguments.
   ///
   /// \param Kind Kind of current clause.
   /// \param ParseOnly true to skip the clause's semantic actions and return
   /// nullptr.
   ///
   OMPClause *ParseOpenMPClause(OpenMPClauseKind Kind, bool ParseOnly = false);
   /// Parses clause with the list of variables of a kind \a Kind.

clang/include/clang/Sema/Sema.h

StmtResult ActOnOpenMPTargetTeamsDistributeParallelForSimdDirective(
       ArrayRef<OMPClause *> Clauses, Stmt *AStmt, SourceLocation StartLoc,
       SourceLocation EndLoc, VarsWithInheritedDSAType &VarsWithImplicitDSA);
   /// Called on well-formed '\#pragma omp target teams distribute simd' after
   /// parsing of the associated statement.
   StmtResult ActOnOpenMPTargetTeamsDistributeSimdDirective(
       ArrayRef<OMPClause *> Clauses, Stmt *AStmt, SourceLocation StartLoc,
       SourceLocation EndLoc, VarsWithInheritedDSAType &VarsWithImplicitDSA);
 
+  /// Called on well-formed '#pragma omp tile' after parsing of its clauses and
+  /// the associated statement.
+  StmtResult
+  ActOnOpenMPTileDirective(ArrayRef<OMPClause *> Clauses, Stmt *AStmt,
+                           SourceLocation StartLoc, SourceLocation EndLoc,
+                           VarsWithInheritedDSAType &VarsWithImplicitDSA);
+
   /// Checks correctness of linear modifiers.
   bool CheckOpenMPLinearModifier(OpenMPLinearClauseKind LinKind,
                                  SourceLocation LinLoc);
   /// Checks that the specified declaration matches requirements for the linear
   /// decls.
   bool CheckOpenMPLinearDecl(const ValueDecl *D, SourceLocation ELoc,
                              OpenMPLinearClauseKind LinKind, QualType Type,
                              bool IsDeclareSimd = false);
public:
   OMPClause *ActOnOpenMPIsDevicePtrClause(ArrayRef<Expr *> VarList,
                                           const OMPVarListLocTy &Locs);
   /// Called on well-formed 'nontemporal' clause.
   OMPClause *ActOnOpenMPNontemporalClause(ArrayRef<Expr *> VarList,
                                           SourceLocation StartLoc,
                                           SourceLocation LParenLoc,
                                           SourceLocation EndLoc);
 
+  /// Called on well-form 'size' clause.
+  OMPClause *ActOnOpenMPSizesClause(ArrayRef<Expr *> SizeExprs,
+                                    SourceLocation StartLoc,
+                                    SourceLocation LParenLoc,
+                                    SourceLocation EndLoc);
+
   /// The kind of conversion being performed.
   enum CheckedConversionKind {
     /// An implicit conversion.
     CCK_ImplicitConversion,
     /// A C-style cast.
     CCK_CStyleCast,
     /// A functional-style cast.
     CCK_FunctionalCast,

clang/include/clang/Serialization/ASTBitCodes.h

enum StmtCode {
       EXPR_USER_DEFINED_LITERAL,
 
       /// A CXXStdInitializerListExpr record.
       EXPR_CXX_STD_INITIALIZER_LIST,
 
       /// A CXXBoolLiteralExpr record.
       EXPR_CXX_BOOL_LITERAL,
 
-      EXPR_CXX_NULL_PTR_LITERAL,  // CXXNullPtrLiteralExpr
-      EXPR_CXX_TYPEID_EXPR,       // CXXTypeidExpr (of expr).
-      EXPR_CXX_TYPEID_TYPE,       // CXXTypeidExpr (of type).
-      EXPR_CXX_THIS,              // CXXThisExpr
-      EXPR_CXX_THROW,             // CXXThrowExpr
-      EXPR_CXX_DEFAULT_ARG,       // CXXDefaultArgExpr
-      EXPR_CXX_DEFAULT_INIT,      // CXXDefaultInitExpr
-      EXPR_CXX_BIND_TEMPORARY,    // CXXBindTemporaryExpr
+      EXPR_CXX_NULL_PTR_LITERAL, // CXXNullPtrLiteralExpr
+      EXPR_CXX_TYPEID_EXPR,      // CXXTypeidExpr (of expr).
+      EXPR_CXX_TYPEID_TYPE,      // CXXTypeidExpr (of type).
+      EXPR_CXX_THIS,             // CXXThisExpr
+      EXPR_CXX_THROW,            // CXXThrowExpr
+      EXPR_CXX_DEFAULT_ARG,      // CXXDefaultArgExpr
+      EXPR_CXX_DEFAULT_INIT,     // CXXDefaultInitExpr
+      EXPR_CXX_BIND_TEMPORARY,   // CXXBindTemporaryExpr
 
       EXPR_CXX_SCALAR_VALUE_INIT, // CXXScalarValueInitExpr
       EXPR_CXX_NEW,               // CXXNewExpr
       EXPR_CXX_DELETE,            // CXXDeleteExpr
       EXPR_CXX_PSEUDO_DESTRUCTOR, // CXXPseudoDestructorExpr
 
-      EXPR_EXPR_WITH_CLEANUPS,    // ExprWithCleanups
+      EXPR_EXPR_WITH_CLEANUPS, // ExprWithCleanups
 
       EXPR_CXX_DEPENDENT_SCOPE_MEMBER,   // CXXDependentScopeMemberExpr
       EXPR_CXX_DEPENDENT_SCOPE_DECL_REF, // DependentScopeDeclRefExpr
       EXPR_CXX_UNRESOLVED_CONSTRUCT,     // CXXUnresolvedConstructExpr
       EXPR_CXX_UNRESOLVED_MEMBER,        // UnresolvedMemberExpr
       EXPR_CXX_UNRESOLVED_LOOKUP,        // UnresolvedLookupExpr
 
-      EXPR_CXX_EXPRESSION_TRAIT,  // ExpressionTraitExpr
-      EXPR_CXX_NOEXCEPT,          // CXXNoexceptExpr
+      EXPR_CXX_EXPRESSION_TRAIT, // ExpressionTraitExpr
+      EXPR_CXX_NOEXCEPT,         // CXXNoexceptExpr
 
-      EXPR_OPAQUE_VALUE,          // OpaqueValueExpr
-      EXPR_BINARY_CONDITIONAL_OPERATOR,  // BinaryConditionalOperator
-      EXPR_TYPE_TRAIT,            // TypeTraitExpr
-      EXPR_ARRAY_TYPE_TRAIT,      // ArrayTypeTraitIntExpr
+      EXPR_OPAQUE_VALUE,                // OpaqueValueExpr
+      EXPR_BINARY_CONDITIONAL_OPERATOR, // BinaryConditionalOperator
+      EXPR_TYPE_TRAIT,                  // TypeTraitExpr
+      EXPR_ARRAY_TYPE_TRAIT,            // ArrayTypeTraitIntExpr
 
-      EXPR_PACK_EXPANSION,        // PackExpansionExpr
-      EXPR_SIZEOF_PACK,           // SizeOfPackExpr
-      EXPR_SUBST_NON_TYPE_TEMPLATE_PARM, // SubstNonTypeTemplateParmExpr
-      EXPR_SUBST_NON_TYPE_TEMPLATE_PARM_PACK,// SubstNonTypeTemplateParmPackExpr
-      EXPR_FUNCTION_PARM_PACK,    // FunctionParmPackExpr
-      EXPR_MATERIALIZE_TEMPORARY, // MaterializeTemporaryExpr
-      EXPR_CXX_FOLD,              // CXXFoldExpr
-      EXPR_CONCEPT_SPECIALIZATION,// ConceptSpecializationExpr
-      EXPR_REQUIRES,              // RequiresExpr
+      EXPR_PACK_EXPANSION,                    // PackExpansionExpr
+      EXPR_SIZEOF_PACK,                       // SizeOfPackExpr
+      EXPR_SUBST_NON_TYPE_TEMPLATE_PARM,      // SubstNonTypeTemplateParmExpr
+      EXPR_SUBST_NON_TYPE_TEMPLATE_PARM_PACK, // SubstNonTypeTemplateParmPackExpr
+      EXPR_FUNCTION_PARM_PACK,                // FunctionParmPackExpr
+      EXPR_MATERIALIZE_TEMPORARY,             // MaterializeTemporaryExpr
+      EXPR_CXX_FOLD,                          // CXXFoldExpr
+      EXPR_CONCEPT_SPECIALIZATION,            // ConceptSpecializationExpr
+      EXPR_REQUIRES,                          // RequiresExpr
 
       // CUDA
-      EXPR_CUDA_KERNEL_CALL,       // CUDAKernelCallExpr
+      EXPR_CUDA_KERNEL_CALL, // CUDAKernelCallExpr
 
       // OpenCL
-      EXPR_ASTYPE,                 // AsTypeExpr
+      EXPR_ASTYPE, // AsTypeExpr
 
       // Microsoft
-      EXPR_CXX_PROPERTY_REF_EXPR, // MSPropertyRefExpr
+      EXPR_CXX_PROPERTY_REF_EXPR,       // MSPropertyRefExpr
       EXPR_CXX_PROPERTY_SUBSCRIPT_EXPR, // MSPropertySubscriptExpr
-      EXPR_CXX_UUIDOF_EXPR,       // CXXUuidofExpr (of expr).
-      EXPR_CXX_UUIDOF_TYPE,       // CXXUuidofExpr (of type).
-      STMT_SEH_LEAVE,             // SEHLeaveStmt
-      STMT_SEH_EXCEPT,            // SEHExceptStmt
-      STMT_SEH_FINALLY,           // SEHFinallyStmt
-      STMT_SEH_TRY,               // SEHTryStmt
+      EXPR_CXX_UUIDOF_EXPR,             // CXXUuidofExpr (of expr).
+      EXPR_CXX_UUIDOF_TYPE,             // CXXUuidofExpr (of type).
+      STMT_SEH_LEAVE,                   // SEHLeaveStmt
+      STMT_SEH_EXCEPT,                  // SEHExceptStmt
+      STMT_SEH_FINALLY,                 // SEHFinallyStmt
+      STMT_SEH_TRY,                     // SEHTryStmt
 
       // OpenMP directives
       STMT_OMP_PARALLEL_DIRECTIVE,
       STMT_OMP_SIMD_DIRECTIVE,
       STMT_OMP_FOR_DIRECTIVE,
       STMT_OMP_FOR_SIMD_DIRECTIVE,
       STMT_OMP_SECTIONS_DIRECTIVE,
       STMT_OMP_SECTION_DIRECTIVE,
enum StmtCode {
       STMT_OMP_TEAMS_DISTRIBUTE_SIMD_DIRECTIVE,
       STMT_OMP_TEAMS_DISTRIBUTE_PARALLEL_FOR_SIMD_DIRECTIVE,
       STMT_OMP_TEAMS_DISTRIBUTE_PARALLEL_FOR_DIRECTIVE,
       STMT_OMP_TARGET_TEAMS_DIRECTIVE,
       STMT_OMP_TARGET_TEAMS_DISTRIBUTE_DIRECTIVE,
       STMT_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_FOR_DIRECTIVE,
       STMT_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_FOR_SIMD_DIRECTIVE,
       STMT_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD_DIRECTIVE,
+      STMT_OMP_TILE_DIRECTIVE,
       EXPR_OMP_ARRAY_SECTION,
 
       // ARC
-      EXPR_OBJC_BRIDGED_CAST,     // ObjCBridgedCastExpr
+      EXPR_OBJC_BRIDGED_CAST, // ObjCBridgedCastExpr
 
-      STMT_MS_DEPENDENT_EXISTS,   // MSDependentExistsStmt
-      EXPR_LAMBDA,                // LambdaExpr
+      STMT_MS_DEPENDENT_EXISTS, // MSDependentExistsStmt
+      EXPR_LAMBDA,              // LambdaExpr
       STMT_COROUTINE_BODY,
       STMT_CORETURN,
       EXPR_COAWAIT,
       EXPR_COYIELD,
       EXPR_DEPENDENT_COAWAIT,
     };
 
     /// The kinds of designators that can occur in a

clang/lib/AST/ASTTypeTraits.cpp

 #include "clang/AST/ASTContext.h"
 #include "clang/AST/DeclCXX.h"
 #include "clang/AST/NestedNameSpecifier.h"
 #include "clang/AST/OpenMPClause.h"
 
 using namespace clang;
 
 const ASTNodeKind::KindInfo ASTNodeKind::AllKindInfo[] = {
-  { NKI_None, "<None>" },
-  { NKI_None, "TemplateArgument" },
-  { NKI_None, "TemplateName" },
-  { NKI_None, "NestedNameSpecifierLoc" },
-  { NKI_None, "QualType" },
-  { NKI_None, "TypeLoc" },
-  { NKI_None, "CXXCtorInitializer" },
-  { NKI_None, "NestedNameSpecifier" },
-  { NKI_None, "Decl" },
+    {NKI_None, "<None>"},
+    {NKI_None, "TemplateArgument"},
+    {NKI_None, "TemplateName"},
+    {NKI_None, "NestedNameSpecifierLoc"},
+    {NKI_None, "QualType"},
+    {NKI_None, "TypeLoc"},
+    {NKI_None, "CXXCtorInitializer"},
+    {NKI_None, "NestedNameSpecifier"},
+    {NKI_None, "Decl"},
 #define DECL(DERIVED, BASE) { NKI_##BASE, #DERIVED "Decl" },
 #include "clang/AST/DeclNodes.inc"
-  { NKI_None, "Stmt" },
+    {NKI_None, "Stmt"},
 #define STMT(DERIVED, BASE) { NKI_##BASE, #DERIVED },
 #include "clang/AST/StmtNodes.inc"
-  { NKI_None, "Type" },
+    {NKI_None, "Type"},
 #define TYPE(DERIVED, BASE) { NKI_##BASE, #DERIVED "Type" },
 #include "clang/AST/TypeNodes.inc"
-  { NKI_None, "OMPClause" },
+    {NKI_None, "OMPClause"},
 #define OPENMP_CLAUSE(TextualSpelling, Class) {NKI_OMPClause, #Class},
 #include "clang/Basic/OpenMPKinds.def"
 };
 
 bool ASTNodeKind::isBaseOf(ASTNodeKind Other, unsigned *Distance) const {
   return isBaseOf(KindId, Other.KindId, Distance);
 }
 

clang/lib/AST/OpenMPClause.cpp

const OMPClauseWithPreInit *OMPClauseWithPreInit::get(const OMPClause *C) {
   case OMPC_device_type:
   case OMPC_match:
   case OMPC_nontemporal:
   case OMPC_order:
   case OMPC_destroy:
   case OMPC_detach:
   case OMPC_inclusive:
   case OMPC_exclusive:
+  case OMPC_sizes:
     break;
   }
 
   return nullptr;
 }
 
 OMPClauseWithPostUpdate *OMPClauseWithPostUpdate::get(OMPClause *C) {
   auto *Res = OMPClauseWithPostUpdate::get(const_cast<const OMPClause *>(C));
const OMPClauseWithPostUpdate *OMPClauseWithPostUpdate::get(const OMPClause *C) {
   case OMPC_device_type:
   case OMPC_match:
   case OMPC_nontemporal:
   case OMPC_order:
   case OMPC_destroy:
   case OMPC_detach:
   case OMPC_inclusive:
   case OMPC_exclusive:
+  case OMPC_sizes:
     break;
   }
 
   return nullptr;
 }
 
 /// Gets the address of the original, non-captured, expression used in the
 /// clause as the preinitializer.

 }
 
 OMPExclusiveClause *OMPExclusiveClause::CreateEmpty(const ASTContext &C,
                                                     unsigned N) {
   void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(N));
   return new (Mem) OMPExclusiveClause(N);
 }
 
+OMPSizesClause *OMPSizesClause::Create(const ASTContext &C,
+                                       SourceLocation StartLoc,
+                                       SourceLocation LParenLoc,
+                                       SourceLocation EndLoc,
+                                       ArrayRef<Expr *> Sizes) {
+  OMPSizesClause *Clause = CreateEmpty(C, Sizes.size());
+  Clause->setLocStart(StartLoc);
+  Clause->setLParenLoc(LParenLoc);
+  Clause->setLocEnd(EndLoc);
+  Clause->setSizesRefs(Sizes);
+  return Clause;
+}
+
+OMPSizesClause *OMPSizesClause::CreateEmpty(const ASTContext &C,
+                                            unsigned NumSizes) {
+  void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(NumSizes));
+  return new (Mem) OMPSizesClause(NumSizes);
+}
+
 //===----------------------------------------------------------------------===//
 //  OpenMP clauses printing methods
 //===----------------------------------------------------------------------===//
 
 void OMPClausePrinter::VisitOMPIfClause(OMPIfClause *Node) {
   OS << "if(";
   if (Node->getNameModifier() != llvm::omp::OMPD_unknown)
     OS << getOpenMPDirectiveName(Node->getNameModifier()) << ": ";
void OMPClausePrinter::VisitOMPNontemporalClause(OMPNontemporalClause *Node) {
   }
 }
 
 void OMPClausePrinter::VisitOMPOrderClause(OMPOrderClause *Node) {
   OS << "order(" << getOpenMPSimpleClauseTypeName(OMPC_order, Node->getKind())
      << ")";
 }
 
+void OMPClausePrinter::VisitOMPSizesClause(OMPSizesClause *Node) {
+  OS << "sizes(";
+  bool First = true;
+  for (auto Size : Node->getSizesRefs()) {
+    if (!First)
+      OS << ", ";
+    Size->printPretty(OS, nullptr, Policy, 0);
+    First = false;
+  }
+  OS << ")";
+}
+
 void OMPClausePrinter::VisitOMPInclusiveClause(OMPInclusiveClause *Node) {
   if (!Node->varlist_empty()) {
     OS << "inclusive";
     VisitOMPClauseList(Node, '(');
     OS << ")";
   }
 }
 

clang/lib/AST/StmtOpenMP.cpp

bool OMPExecutableDirective::isStandaloneDirective() const {
   // reasons they have empty synthetic structured block, to simplify codegen.
   if (isa<OMPTargetEnterDataDirective>(this) ||
       isa<OMPTargetExitDataDirective>(this) ||
       isa<OMPTargetUpdateDirective>(this))
     return true;
   return !hasAssociatedStmt() || !getAssociatedStmt();
 }
 
+Stmt *OMPExecutableDirective::ignoreCaptures() {
+  SmallVector<OpenMPDirectiveKind, 4> CaptureRegions;
+  getOpenMPCaptureRegions(CaptureRegions, getDirectiveKind());
+  auto Result = getAssociatedStmt();
+  for (auto ThisCaptureRegion : CaptureRegions) {
+    (void)ThisCaptureRegion;
+    Result = cast<CapturedStmt>(Result)->getCapturedStmt();
+  }
+  return Result;
+}
+
 const Stmt *OMPExecutableDirective::getStructuredBlock() const {
   assert(!isStandaloneDirective() &&
          "Standalone Executable Directives don't have Structured Blocks.");
   if (auto *LD = dyn_cast<OMPLoopDirective>(this))
     return LD->getBody();
-  return getInnermostCapturedStmt()->getCapturedStmt();
+  return ignoreCaptures();
 }
 
-Stmt *OMPLoopDirective::tryToFindNextInnerLoop(Stmt *CurStmt,
-                                               bool TryImperfectlyNestedLoops) {
+Stmt *clang::getTopmostAssociatedStructuredBlock(
+    Stmt *S, llvm::SmallVectorImpl<Stmt *> *PreInits) {
+  assert(S && "Must be a valid statement");
+
+  while (true) {
+    S = S->IgnoreContainers(/* IgnoreCaptured */ true);
+    if (!isa<OMPExecutableDirective>(S) ||
+        !isOpenMPLoopTransformationDirective(
+            cast<OMPExecutableDirective>(S)->getDirectiveKind()))
+      break;
+
+    if (auto *D = dyn_cast<OMPTileDirective>(S)) {
+      if (PreInits) {
+        auto PreInit = D->getPreInits();
+        PreInits->append(PreInit.begin(), PreInit.end());
+      }
+      S = D->getTransformedForStmt();
+      continue;
+    }
+    llvm_unreachable("Unhandled loop transformation");
+  }
+  return S;
+}
+
+Stmt *OMPLoopDirective::tryToFindNextInnerLoop(
+    Stmt *CurStmt, bool TryImperfectlyNestedLoops,
+    llvm::SmallVectorImpl<Stmt *> *PreInits) {
   Stmt *OrigStmt = CurStmt;
-  CurStmt = CurStmt->IgnoreContainers();
+  CurStmt = getTopmostAssociatedStructuredBlock(CurStmt, PreInits);
+
   // Additional work for imperfectly nested loops, introduced in OpenMP 5.0.
   if (TryImperfectlyNestedLoops) {
     if (auto *CS = dyn_cast<CompoundStmt>(CurStmt)) {
       CurStmt = nullptr;
       SmallVector<CompoundStmt *, 4> Statements(1, CS);
       SmallVector<CompoundStmt *, 4> NextStatements;
       while (!Statements.empty()) {
         CS = Statements.pop_back_val();
if (auto *CS = dyn_cast<CompoundStmt>(CurStmt)) {
       }
       if (!CurStmt)
         CurStmt = OrigStmt;
     }
   }
   return CurStmt;
 }
 
+void OMPLoopDirective::collectAssociatedLoops(
+    llvm::SmallVectorImpl<Stmt *> &Loops,
+    llvm::SmallVectorImpl<Stmt *> &PreInits) {
+  Stmt *Body = ignoreCaptures()->IgnoreContainers();
+  unsigned NumLoops = getCollapsedNumber();
+  assert(NumLoops >= 1);
+
+  // For each nest level, find the loop and its body.
+  // The topmost loop must not be surrounded by other code.
+  for (unsigned Cnt = 0; Cnt < NumLoops; ++Cnt) {
+    Stmt *Loop = OMPLoopDirective::tryToFindNextInnerLoop(
+        Body, /*TryImperfectlyNestedLoops=*/(Cnt > 0), &PreInits);
+    Loops.push_back(Loop);
+
+    // Get body to look next loop in
+    if (auto *For = dyn_cast<ForStmt>(Loop)) {
+      Body = For->getBody();
+    } else if (auto CXXFor = dyn_cast<CXXForRangeStmt>(Loop)) {
+      Body = CXXFor->getBody();
+    } else
+      llvm_unreachable("Expected canonical for loop or range-based for loop.");
+  }
+}
+
 Stmt *OMPLoopDirective::getBody() {
   // This relies on the loop form is already checked by Sema.
-  Stmt *Body =
-      getInnermostCapturedStmt()->getCapturedStmt()->IgnoreContainers();
+  Stmt *Body = getTopmostAssociatedStructuredBlock(ignoreCaptures(), nullptr);
   if (auto *For = dyn_cast<ForStmt>(Body)) {
     Body = For->getBody();
   } else {
     assert(isa<CXXForRangeStmt>(Body) &&
            "Expected canonical for loop or range-based for loop.");
     Body = cast<CXXForRangeStmt>(Body)->getBody();
   }
   for (unsigned Cnt = 1; Cnt < CollapsedNum; ++Cnt) {
OMPForDirective *OMPForDirective::CreateEmpty(const ASTContext &C,
                                               EmptyShell) {
   unsigned Size = llvm::alignTo(sizeof(OMPForDirective), alignof(OMPClause *));
   void *Mem =
       C.Allocate(Size + sizeof(OMPClause *) * NumClauses +
                  sizeof(Stmt *) * numLoopChildren(CollapsedNum, OMPD_for));
   return new (Mem) OMPForDirective(CollapsedNum, NumClauses);
 }
 
+OMPTileDirective *
+OMPTileDirective::Create(const ASTContext &C, SourceLocation StartLoc,
+                         SourceLocation EndLoc, ArrayRef<OMPClause *> Clauses,
+                         unsigned NumLoops, Stmt *AssociatedStmt,
+                         Stmt *TransformedStmt) {
+  OMPTileDirective *Dir = CreateEmpty(C, Clauses.size(), NumLoops);
+  Dir->setLocStart(StartLoc);
+  Dir->setLocEnd(EndLoc);
+  Dir->setClauses(Clauses);
+  Dir->setAssociatedStmt(AssociatedStmt);
+  Dir->setTransformedStmt(TransformedStmt);
+  return Dir;
+}
+
+OMPTileDirective *OMPTileDirective::CreateEmpty(const ASTContext &C,
+                                                unsigned NumClauses,
+                                                unsigned NumLoops) {
+  void *Mem = C.Allocate(totalSizeToAlloc<OMPClause *, Stmt *>(
+      NumClauses, numLoopChildren(NumLoops, OMPD_tile)));
+  return new (Mem) OMPTileDirective({}, {}, NumClauses, NumLoops);
+}
+
 OMPForSimdDirective *
 OMPForSimdDirective::Create(const ASTContext &C, SourceLocation StartLoc,
                             SourceLocation EndLoc, unsigned CollapsedNum,
                             ArrayRef<OMPClause *> Clauses, Stmt *AssociatedStmt,
                             const HelperExprs &Exprs) {
   unsigned Size =
       llvm::alignTo(sizeof(OMPForSimdDirective), alignof(OMPClause *));
   void *Mem =

clang/lib/AST/StmtPrinter.cpp

void StmtPrinter::PrintOMPExecutableDirective(OMPExecutableDirective *S,
   ArrayRef<OMPClause *> Clauses = S->clauses();
   for (auto *Clause : Clauses)
     if (Clause && !Clause->isImplicit()) {
       OS << ' ';
       Printer.Visit(Clause);
     }
   OS << NL;
   if (!ForceNoStmt && S->hasAssociatedStmt())
-    PrintStmt(S->getInnermostCapturedStmt()->getCapturedStmt());
+    PrintStmt(S->ignoreCaptures());
 }
 
 void StmtPrinter::VisitOMPParallelDirective(OMPParallelDirective *Node) {
   Indent() << "#pragma omp parallel";
   PrintOMPExecutableDirective(Node);
 }
 
 void StmtPrinter::VisitOMPSimdDirective(OMPSimdDirective *Node) {

 }
 
 void StmtPrinter::VisitOMPTargetTeamsDistributeSimdDirective(
     OMPTargetTeamsDistributeSimdDirective *Node) {
   Indent() << "#pragma omp target teams distribute simd";
   PrintOMPExecutableDirective(Node);
 }
 
+void StmtPrinter::VisitOMPTileDirective(OMPTileDirective *Node) {
+  Indent() << "#pragma omp tile";
+  PrintOMPExecutableDirective(Node);
+}
+
 //===----------------------------------------------------------------------===//
 //  Expr printing methods.
 //===----------------------------------------------------------------------===//
 
 void StmtPrinter::VisitSourceLocExpr(SourceLocExpr *Node) {
   OS << Node->getBuiltinStr() << "()";
 }
 

clang/lib/AST/StmtProfile.cpp

 }
 void OMPClauseProfiler::VisitOMPInclusiveClause(const OMPInclusiveClause *C) {
   VisitOMPClauseList(C);
 }
 void OMPClauseProfiler::VisitOMPExclusiveClause(const OMPExclusiveClause *C) {
   VisitOMPClauseList(C);
 }
 void OMPClauseProfiler::VisitOMPOrderClause(const OMPOrderClause *C) {}
+
+void OMPClauseProfiler::VisitOMPSizesClause(const OMPSizesClause *C) {
+  for (auto E : C->getSizesRefs())
+    if (E)
+      Profiler->VisitExpr(E);
+}
 } // namespace
 
 void
 StmtProfiler::VisitOMPExecutableDirective(const OMPExecutableDirective *S) {
   VisitStmt(S);
   OMPClauseProfiler P(this);
   ArrayRef<OMPClause *> Clauses = S->clauses();
   for (ArrayRef<OMPClause *>::iterator I = Clauses.begin(), E = Clauses.end();
void StmtProfiler::VisitOMPTargetTeamsDistributeParallelForSimdDirective(
   VisitOMPLoopDirective(S);
 }
 
 void StmtProfiler::VisitOMPTargetTeamsDistributeSimdDirective(
     const OMPTargetTeamsDistributeSimdDirective *S) {
   VisitOMPLoopDirective(S);
 }
 
+void StmtProfiler::VisitOMPTileDirective(const OMPTileDirective *S) {
+  VisitOMPLoopDirective(S);
+}
+
 void StmtProfiler::VisitExpr(const Expr *S) {
   VisitStmt(S);
 }
 
 void StmtProfiler::VisitConstantExpr(const ConstantExpr *S) {
   VisitExpr(S);
 }
 

clang/lib/Analysis/CFG.cpp

SmallVector<Stmt *, 8> Used(
       OMPExecutableDirective::used_clauses_children(D->clauses()));
   for (Stmt *S : llvm::reverse(Used)) {
     assert(S && "Expected non-null used-in-clause child.");
     if (CFGBlock *R = Visit(S))
       B = R;
   }
   // Visit associated structured block if any.
   if (!D->isStandaloneDirective())
-    if (CapturedStmt *CS = D->getInnermostCapturedStmt()) {
-      Stmt *S = CS->getCapturedStmt();
+    if (Stmt *S = D->ignoreCaptures()) {
       if (!isa<CompoundStmt>(S))
         addLocalScopeAndDtors(S);
       if (CFGBlock *R = addStmt(S))
         B = R;
     }
 
   return B;
 }

clang/lib/Basic/OpenMPKinds.cpp

#include "clang/Basic/OpenMPKinds.def"
   case OMPC_reverse_offload:
   case OMPC_dynamic_allocators:
   case OMPC_match:
   case OMPC_nontemporal:
   case OMPC_destroy:
   case OMPC_detach:
   case OMPC_inclusive:
   case OMPC_exclusive:
+  case OMPC_sizes:
     break;
   }
   llvm_unreachable("Invalid OpenMP simple clause kind");
 }
 
 const char *clang::getOpenMPSimpleClauseTypeName(OpenMPClauseKind Kind,
                                                  unsigned Type) {
   switch (Kind) {
#include "clang/Basic/OpenMPKinds.def"
   case OMPC_reverse_offload:
   case OMPC_dynamic_allocators:
   case OMPC_match:
   case OMPC_nontemporal:
   case OMPC_destroy:
   case OMPC_detach:
   case OMPC_inclusive:
   case OMPC_exclusive:
+  case OMPC_sizes:
     break;
   }
   llvm_unreachable("Invalid OpenMP simple clause kind");
 }
 
 bool clang::isAllowedClauseForDirective(OpenMPDirectiveKind DKind,
                                         OpenMPClauseKind CKind,
                                         unsigned OpenMPVersion) {

 #define OPENMP_DECLARE_VARIANT_CLAUSE(Name)                                    \
   case OMPC_##Name:                                                            \
     return true;
 #include "clang/Basic/OpenMPKinds.def"
     default:
       break;
     }
     break;
+  case OMPD_tile:
+    switch (CKind) {
+#define OPENMP_TILE_CLAUSE(Name)                                               \
+  case OMPC_##Name:                                                            \
+    return true;
+#include "clang/Basic/OpenMPKinds.def"
+    default:
+      break;
+    }
+    break;
   case OMPD_declare_target:
   case OMPD_end_declare_target:
   case OMPD_unknown:
   case OMPD_threadprivate:
   case OMPD_section:
   case OMPD_master:
   case OMPD_taskyield:
   case OMPD_barrier:
return DKind == OMPD_simd || DKind == OMPD_for || DKind == OMPD_for_simd ||
          DKind == OMPD_target_parallel_for_simd || DKind == OMPD_target_simd ||
          DKind == OMPD_teams_distribute ||
          DKind == OMPD_teams_distribute_simd ||
          DKind == OMPD_teams_distribute_parallel_for_simd ||
          DKind == OMPD_teams_distribute_parallel_for ||
          DKind == OMPD_target_teams_distribute ||
          DKind == OMPD_target_teams_distribute_parallel_for ||
          DKind == OMPD_target_teams_distribute_parallel_for_simd ||
-         DKind == OMPD_target_teams_distribute_simd;
+         DKind == OMPD_target_teams_distribute_simd || DKind == OMPD_tile;
 }
 
 bool clang::isOpenMPWorksharingDirective(OpenMPDirectiveKind DKind) {
   return DKind == OMPD_for || DKind == OMPD_for_simd ||
          DKind == OMPD_sections || DKind == OMPD_section ||
          DKind == OMPD_single || DKind == OMPD_parallel_for ||
          DKind == OMPD_parallel_for_simd || DKind == OMPD_parallel_sections ||
          DKind == OMPD_target_parallel_for ||
bool clang::isOpenMPLoopBoundSharingDirective(OpenMPDirectiveKind Kind) {
   return Kind == OMPD_distribute_parallel_for ||
          Kind == OMPD_distribute_parallel_for_simd ||
          Kind == OMPD_teams_distribute_parallel_for_simd ||
          Kind == OMPD_teams_distribute_parallel_for ||
          Kind == OMPD_target_teams_distribute_parallel_for ||
          Kind == OMPD_target_teams_distribute_parallel_for_simd;
 }
 
+bool clang::isOpenMPLoopTransformationDirective(OpenMPDirectiveKind DKind) {
+  return DKind == OMPD_tile;
+}
+
 void clang::getOpenMPCaptureRegions(
     SmallVectorImpl<OpenMPDirectiveKind> &CaptureRegions,
     OpenMPDirectiveKind DKind) {
   assert(DKind <= OMPD_unknown);
   switch (DKind) {
   case OMPD_parallel:
   case OMPD_parallel_for:
   case OMPD_parallel_for_simd:
void clang::getOpenMPCaptureRegions(
   case OMPD_taskgroup:
   case OMPD_distribute:
   case OMPD_ordered:
   case OMPD_atomic:
   case OMPD_target_data:
   case OMPD_distribute_simd:
     CaptureRegions.push_back(OMPD_unknown);
     break;
+  case OMPD_tile:
+    // loop transformations do not introduce captures.
+    break;
   case OMPD_threadprivate:
   case OMPD_allocate:
   case OMPD_taskyield:
   case OMPD_barrier:
   case OMPD_taskwait:
   case OMPD_cancellation_point:
   case OMPD_cancel:
   case OMPD_flush:

clang/lib/CodeGen/CGOpenMPRuntime.cpp

emitNumTeamsForTargetDirective(CodeGenFunction &CGF,
   case OMPD_taskloop:
   case OMPD_taskloop_simd:
   case OMPD_master_taskloop:
   case OMPD_master_taskloop_simd:
   case OMPD_parallel_master_taskloop:
   case OMPD_parallel_master_taskloop_simd:
   case OMPD_requires:
   case OMPD_unknown:
+  case OMPD_tile:
     break;
   }
   llvm_unreachable("Unexpected directive kind.");
 }
 
 static llvm::Value *getNumThreads(CodeGenFunction &CGF, const CapturedStmt *CS,
                                   llvm::Value *DefaultThreadLimitVal) {
   const Stmt *Child = CGOpenMPRuntime::getSingleCompoundChild(
emitNumThreadsForTargetDirective(CodeGenFunction &CGF,
   case OMPD_taskloop:
   case OMPD_taskloop_simd:
   case OMPD_master_taskloop:
   case OMPD_master_taskloop_simd:
   case OMPD_parallel_master_taskloop:
   case OMPD_parallel_master_taskloop_simd:
   case OMPD_requires:
   case OMPD_unknown:
+  case OMPD_tile:
     break;
   }
   llvm_unreachable("Unsupported directive kind.");
 }
 
 namespace {
 LLVM_ENABLE_BITMASK_ENUMS_IN_NAMESPACE();
 
if (const auto *NestedDir =
     case OMPD_taskloop:
     case OMPD_taskloop_simd:
     case OMPD_master_taskloop:
     case OMPD_master_taskloop_simd:
     case OMPD_parallel_master_taskloop:
     case OMPD_parallel_master_taskloop_simd:
     case OMPD_requires:
     case OMPD_unknown:
+    case OMPD_tile:
       llvm_unreachable("Unexpected directive.");
     }
   }
 
   return nullptr;
 }
 
 /// Emit the user-defined mapper function. The code generation follows the
if (RequiresDeviceCodegen) {
     case OMPD_taskloop:
     case OMPD_taskloop_simd:
     case OMPD_master_taskloop:
     case OMPD_master_taskloop_simd:
     case OMPD_parallel_master_taskloop:
     case OMPD_parallel_master_taskloop_simd:
     case OMPD_requires:
     case OMPD_unknown:
+    case OMPD_tile:
       llvm_unreachable("Unknown target directive for OpenMP device codegen.");
     }
     return;
   }
 
   if (const auto *E = dyn_cast<OMPExecutableDirective>(S)) {
     if (!E->hasAssociatedStmt() || !E->getAssociatedStmt())
       return;
auto &&ThenGen = [this, &D, Device, &InputInfo,
     case OMPD_target_teams_distribute_simd:
     case OMPD_target_teams_distribute_parallel_for:
     case OMPD_target_teams_distribute_parallel_for_simd:
     case OMPD_target_teams:
     case OMPD_target_parallel:
     case OMPD_target_parallel_for:
     case OMPD_target_parallel_for_simd:
     case OMPD_requires:
+    case OMPD_tile:
     case OMPD_unknown:
       llvm_unreachable("Unexpected standalone target data directive.");
       break;
     }
     CGF.EmitRuntimeCall(createRuntimeFunction(RTLFn), OffloadingArgs);
   };
 
   auto &&TargetThenGen = [this, &ThenGen, &D, &InputInfo, &MapTypesArray](

clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp

if (const auto *NestedDir =
     case OMPD_taskloop:
     case OMPD_taskloop_simd:
     case OMPD_master_taskloop:
     case OMPD_master_taskloop_simd:
     case OMPD_parallel_master_taskloop:
     case OMPD_parallel_master_taskloop_simd:
     case OMPD_requires:
     case OMPD_unknown:
+    case OMPD_tile:
       llvm_unreachable("Unexpected directive.");
     }
   }
 
   return false;
 }
 
 static bool supportsSPMDExecutionMode(ASTContext &Ctx,
static bool supportsSPMDExecutionMode(ASTContext &Ctx,
   case OMPD_taskloop:
   case OMPD_taskloop_simd:
   case OMPD_master_taskloop:
   case OMPD_master_taskloop_simd:
   case OMPD_parallel_master_taskloop:
   case OMPD_parallel_master_taskloop_simd:
   case OMPD_requires:
   case OMPD_unknown:
+  case OMPD_tile:
     break;
   }
   llvm_unreachable(
       "Unknown programming model for OpenMP directive on NVPTX target.");
 }
 
 /// Check if the directive is loops based and has schedule clause at all or has
 /// static scheduling.
if (const auto *NestedDir =
     case OMPD_taskloop:
     case OMPD_taskloop_simd:
     case OMPD_master_taskloop:
     case OMPD_master_taskloop_simd:
     case OMPD_parallel_master_taskloop:
     case OMPD_parallel_master_taskloop_simd:
     case OMPD_requires:
     case OMPD_unknown:
+    case OMPD_tile:
       llvm_unreachable("Unexpected directive.");
     }
   }
 
   return false;
 }
 
 /// Checks if the construct supports lightweight runtime. It must be SPMD
static bool supportsLightweightRuntime(ASTContext &Ctx,
   case OMPD_taskloop:
   case OMPD_taskloop_simd:
   case OMPD_master_taskloop:
   case OMPD_master_taskloop_simd:
   case OMPD_parallel_master_taskloop:
   case OMPD_parallel_master_taskloop_simd:
   case OMPD_requires:
   case OMPD_unknown:
+  case OMPD_tile:
     break;
   }
   llvm_unreachable(
       "Unknown programming model for OpenMP directive on NVPTX target.");
 }
 
 void CGOpenMPRuntimeNVPTX::emitNonSPMDKernel(const OMPExecutableDirective &D,
                                              StringRef ParentName,

clang/lib/CodeGen/CGStmt.cpp

#include "clang/AST/StmtNodes.inc"
   case Stmt::OMPTargetTeamsDistributeParallelForSimdDirectiveClass:
     EmitOMPTargetTeamsDistributeParallelForSimdDirective(
         cast<OMPTargetTeamsDistributeParallelForSimdDirective>(*S));
     break;
   case Stmt::OMPTargetTeamsDistributeSimdDirectiveClass:
     EmitOMPTargetTeamsDistributeSimdDirective(
         cast<OMPTargetTeamsDistributeSimdDirective>(*S));
     break;
+  case Stmt::OMPTileDirectiveClass:
+    EmitOMPTileDirective(cast<OMPTileDirective>(*S));
+    break;
   }
 }
 
 bool CodeGenFunction::EmitSimpleStmt(const Stmt *S) {
   switch (S->getStmtClass()) {
   default: return false;
   case Stmt::NullStmtClass: break;
   case Stmt::CompoundStmtClass: EmitCompoundStmt(cast<CompoundStmt>(*S)); break;

clang/lib/CodeGen/CGStmtOpenMP.cpp

OMPTeamsScope(CodeGenFunction &CGF, const OMPExecutableDirective &S)
       : OMPLexicalScope(CGF, S, /*CapturedRegion=*/llvm::None,
                         EmitPreInitStmt(S)) {}
 };
 
 /// Private scope for OpenMP loop-based directives, that supports capturing
 /// of used expression from loop statement.
 class OMPLoopScope : public CodeGenFunction::RunCleanupsScope {
   void emitPreInitStmt(CodeGenFunction &CGF, const OMPLoopDirective &S) {
+
+    // Collect loop nests and the pre-inits of associated loops.
+    SmallVector<const Stmt *, 4> Loops;
+    SmallVector<Stmt *, 8> AssociatedPreInits;
+    S.collectAssociatedLoops(Loops, AssociatedPreInits);
+
+    // Emit statements required by nested loop transformations. Has to be done
+    // before PreCondVars.
+    for (Stmt *APreInit : AssociatedPreInits)
+      CGF.EmitStmt(APreInit);
+
     CodeGenFunction::OMPMapVars PreCondVars;
     llvm::DenseSet<const VarDecl *> EmittedAsPrivate;
     for (const auto *E : S.counters()) {
       const auto *VD = cast<VarDecl>(cast<DeclRefExpr>(E)->getDecl());
       EmittedAsPrivate.insert(VD->getCanonicalDecl());
       (void)PreCondVars.setVarAddr(
           CGF, VD, CGF.CreateMemTemp(VD->getType().getNonReferenceType()));
     }
     // Mark private vars as undefs.
     for (const auto *C : S.getClausesOfKind<OMPPrivateClause>()) {
       for (const Expr *IRef : C->varlists()) {
         const auto *OrigVD = cast<VarDecl>(cast<DeclRefExpr>(IRef)->getDecl());
         if (EmittedAsPrivate.insert(OrigVD->getCanonicalDecl()).second) {
           (void)PreCondVars.setVarAddr(
               CGF, OrigVD,
               Address(llvm::UndefValue::get(
                           CGF.ConvertTypeForMem(CGF.getContext().getPointerType(
                               OrigVD->getType().getNonReferenceType()))),
                       CGF.getContext().getDeclAlign(OrigVD)));
         }
       }
     }
     (void)PreCondVars.apply(CGF);
     // Emit init, __range and __end variables for C++ range loops.
-    const Stmt *Body =
-        S.getInnermostCapturedStmt()->getCapturedStmt()->IgnoreContainers();
-    for (unsigned Cnt = 0; Cnt < S.getCollapsedNumber(); ++Cnt) {
-      Body = OMPLoopDirective::tryToFindNextInnerLoop(
-          Body, /*TryImperfectlyNestedLoops=*/true);
-      if (auto *For = dyn_cast<ForStmt>(Body)) {
-        Body = For->getBody();
-      } else {
-        assert(isa<CXXForRangeStmt>(Body) &&
-               "Expected canonical for loop or range-based for loop.");
-        auto *CXXFor = cast<CXXForRangeStmt>(Body);
+    for (const Stmt *Body : Loops) {
+      if (auto *CXXFor = dyn_cast<CXXForRangeStmt>(Body)) {
         if (const Stmt *Init = CXXFor->getInit())
           CGF.EmitStmt(Init);
         CGF.EmitStmt(CXXFor->getRangeStmt());
         CGF.EmitStmt(CXXFor->getEndStmt());
-        Body = CXXFor->getBody();
       }
     }
+
+    // Emit captures.
     if (const auto *PreInits = cast_or_null<DeclStmt>(S.getPreInits())) {
       for (const auto *I : PreInits->decls())
         CGF.EmitVarDecl(cast<VarDecl>(*I));
     }
     PreCondVars.restore(CGF);
   }
 
 public:
void CodeGenFunction::EmitOMPParallelDirective(const OMPParallelDirective &S) {
   }
   // Check for outer lastprivate conditional update.
   checkForLastprivateConditionalUpdate(*this, S);
 }
 
 static void emitBody(CodeGenFunction &CGF, const Stmt *S, const Stmt *NextLoop,
                      int MaxLevel, int Level = 0) {
   assert(Level < MaxLevel && "Too deep lookup during loop body codegen.");
-  const Stmt *SimplifiedS = S->IgnoreContainers();
+  const Stmt *SimplifiedS =
+      getTopmostAssociatedStructuredBlock(S->IgnoreContainers(), nullptr);
   if (const auto *CS = dyn_cast<CompoundStmt>(SimplifiedS)) {
     PrettyStackTraceLoc CrashInfo(
         CGF.getContext().getSourceManager(), CS->getLBracLoc(),
         "LLVM IR generation of compound statement ('{}')");
 
     // Keep track of the current cleanup stack depth, including debug scopes.
     CodeGenFunction::LexicalScope Scope(CGF, S->getSourceRange());
     for (const Stmt *CurStmt : CS->body())
static void emitOMPAtomicExpr(CodeGenFunction &CGF, OpenMPClauseKind Kind,
   case OMPC_device_type:
   case OMPC_match:
   case OMPC_nontemporal:
   case OMPC_order:
   case OMPC_destroy:
   case OMPC_detach:
   case OMPC_inclusive:
   case OMPC_exclusive:
+  case OMPC_sizes:
     llvm_unreachable("Clause is not allowed in 'omp atomic'.");
   }
 }
 
 void CodeGenFunction::EmitOMPAtomicDirective(const OMPAtomicDirective &S) {
   llvm::AtomicOrdering AO = llvm::AtomicOrdering::Monotonic;
   bool MemOrderingSpecified = false;
   if (S.getSingleClause<OMPSeqCstClause>()) {

 void CodeGenFunction::EmitOMPTargetTeamsDistributeSimdDirective(
     const OMPTargetTeamsDistributeSimdDirective &S) {
   auto &&CodeGen = [&S](CodeGenFunction &CGF, PrePostActionTy &Action) {
     emitTargetTeamsDistributeSimdRegion(CGF, Action, S);
   };
   emitCommonOMPTargetDirective(*this, S, CodeGen);
 }
 
+void CodeGenFunction::EmitOMPTileDirective(const OMPTileDirective &S) {
+  // Emit the de-sugared statement.
+  EmitStmt(S.getTransformedStmt());
+}
+
 void CodeGenFunction::EmitOMPTeamsDistributeDirective(
     const OMPTeamsDistributeDirective &S) {
 
   auto &&CodeGenDistribute = [&S](CodeGenFunction &CGF, PrePostActionTy &) {
     CGF.EmitOMPDistributeLoop(S, emitOMPLoopBodyWithStopPoint, S.getInc());
   };
 
   // Emit teams region as a standalone region.

clang/lib/CodeGen/CodeGenFunction.h

private:
   const bool ShouldEmitLifetimeMarkers;
 
   /// Add OpenCL kernel arg metadata and the kernel attribute metadata to
   /// the function metadata.
   void EmitOpenCLKernelMetadata(const FunctionDecl *FD,
                                 llvm::Function *Fn);
 
 public:
-  CodeGenFunction(CodeGenModule &cgm, bool suppressNewContext=false);
+  CodeGenFunction(CodeGenModule &cgm, bool suppressNewContext = false);
   ~CodeGenFunction();
 
   CodeGenTypes &getTypes() const { return CGM.getTypes(); }
   ASTContext &getContext() const { return CGM.getContext(); }
   CGDebugInfo *getDebugInfo() {
     if (DisableDebugInfo)
       return nullptr;
     return DebugInfo;

   void EmitOMPTargetTeamsDistributeDirective(
       const OMPTargetTeamsDistributeDirective &S);
   void EmitOMPTargetTeamsDistributeParallelForDirective(
       const OMPTargetTeamsDistributeParallelForDirective &S);
   void EmitOMPTargetTeamsDistributeParallelForSimdDirective(
       const OMPTargetTeamsDistributeParallelForSimdDirective &S);
   void EmitOMPTargetTeamsDistributeSimdDirective(
       const OMPTargetTeamsDistributeSimdDirective &S);
+  void EmitOMPTileDirective(const OMPTileDirective &S);
 
   /// Emit device code for the target directive.
   static void EmitOMPTargetDeviceFunction(CodeGenModule &CGM,
                                           StringRef ParentName,
                                           const OMPTargetDirective &S);
   static void
   EmitOMPTargetParallelDeviceFunction(CodeGenModule &CGM, StringRef ParentName,
                                       const OMPTargetParallelDirective &S);

clang/lib/Parse/ParseOpenMP.cpp

Parser::DeclGroupPtrTy Parser::ParseOpenMPDeclarativeDirectiveWithExtDecl(
   case OMPD_teams_distribute_simd:
   case OMPD_teams_distribute_parallel_for_simd:
   case OMPD_teams_distribute_parallel_for:
   case OMPD_target_teams:
   case OMPD_target_teams_distribute:
   case OMPD_target_teams_distribute_parallel_for:
   case OMPD_target_teams_distribute_parallel_for_simd:
   case OMPD_target_teams_distribute_simd:
+  case OMPD_tile:
     Diag(Tok, diag::err_omp_unexpected_directive)
         << 1 << getOpenMPDirectiveName(DKind);
     break;
   }
   while (Tok.isNot(tok::annot_pragma_openmp_end))
     ConsumeAnyToken();
   ConsumeAnyToken();
   return nullptr;
unsigned ScopeFlags = Scope::FnScope | Scope::DeclScope |
                         Scope::CompoundStmtScope | Scope::OpenMPDirectiveScope;
   SourceLocation Loc = ConsumeAnnotationToken(), EndLoc;
   OpenMPDirectiveKind DKind = parseOpenMPDirectiveKind(*this);
   OpenMPDirectiveKind CancelRegion = OMPD_unknown;
   // Name of critical directive.
   DeclarationNameInfo DirName;
   StmtResult Directive = StmtError();
   bool HasAssociatedStatement = true;
+  bool ErrorFound = false;
 
   switch (DKind) {
   case OMPD_threadprivate: {
     // FIXME: Should this be permitted in C++?
     if ((StmtCtx & ParsedStmtContext::AllowDeclarationsInC) ==
         ParsedStmtContext()) {
       Diag(Tok, diag::err_omp_immediate_directive)
           << getOpenMPDirectiveName(DKind) << 0;
Parser::ParseOpenMPDeclarativeOrExecutableDirective(ParsedStmtContext StmtCtx) {
   case OMPD_teams_distribute:
   case OMPD_teams_distribute_simd:
   case OMPD_teams_distribute_parallel_for_simd:
   case OMPD_teams_distribute_parallel_for:
   case OMPD_target_teams:
   case OMPD_target_teams_distribute:
   case OMPD_target_teams_distribute_parallel_for:
   case OMPD_target_teams_distribute_parallel_for_simd:
-  case OMPD_target_teams_distribute_simd: {
+  case OMPD_target_teams_distribute_simd:
+  case OMPD_tile: {
     // Special processing for flush and depobj clauses.
     Token ImplicitTok;
     bool ImplicitClauseAllowed = false;
     if (DKind == OMPD_flush || DKind == OMPD_depobj) {
       ImplicitTok = Tok;
       ImplicitClauseAllowed = true;
     }
     ConsumeToken();
if (DKind == OMPD_ordered && FirstClauses[OMPC_depend].getInt()) {
           ParsedStmtContext()) {
         Diag(Loc, diag::err_omp_immediate_directive)
             << getOpenMPDirectiveName(DKind) << 1
             << getOpenMPClauseName(OMPC_depend);
       }
       HasAssociatedStatement = false;
     }
 
+    if (DKind == OMPD_tile && !FirstClauses[OMPC_sizes].getInt()) {
+      Diag(Loc, diag::err_omp_required_clause)
+          << getOpenMPDirectiveName(OMPD_tile) << "sizes";
+      ErrorFound = true;
+    }
+
     StmtResult AssociatedStmt;
     if (HasAssociatedStatement) {
       // The body is a block scope like in Lambdas and Blocks.
       Actions.ActOnOpenMPRegionStart(DKind, getCurScope());
       // FIXME: We create a bogus CompoundStmt scope to hold the contents of
       // the captured region. Code elsewhere assumes that any FunctionScopeInfo
       // should have at least one compound statement scope within it.
       AssociatedStmt = (Sema::CompoundScopeRAII(Actions), ParseStatement());
bool Parser::ParseOpenMPSimpleVarList(
   }
 
   // Parse ')'.
   IsCorrect = !T.consumeClose() && IsCorrect;
 
   return !IsCorrect;
 }
 
+OMPClause *Parser::ParseOpenMPSizesClause() {
+  SourceLocation ClauseNameLoc = ConsumeToken();
+  SmallVector<Expr *, 4> ValExprs;
+
+  BalancedDelimiterTracker T(*this, tok::l_paren, tok::annot_pragma_openmp_end);
+  if (T.consumeOpen()) {
+    Diag(Tok, diag::err_expected) << tok::l_paren;
+    return nullptr;
+  }
+
+  // SourceLocation Loc = Tok.getLocation();
+  // SourceLocation LOpen = ConsumeToken();
+  // SmallVector<Expr*, 4> Vals;
+  // OpenMPVarListDataTy Data;
+
+  while (true) {
+    // SourceLocation Loc = ConsumeToken();
+    // SourceLocation LLoc = Tok.getLocation();
+
+    ExprResult Val = ParseConstantExpression();
+    if (!Val.isUsable()) {
+      // TODO: emit diagnostics
+      T.skipToEnd();
+      return nullptr;
+    }
+    // Actions.ActOnFinishFullExpr(Val.get(), {}, true, true);
+
+    ValExprs.push_back(Val.get());
+
+    if (Tok.is(tok::r_paren) || Tok.is(tok::annot_pragma_openmp_end))
+      break;
+
+    ExpectAndConsume(tok::comma);
+  }
+
+  T.consumeClose();
+
+  return Actions.ActOnOpenMPSizesClause(
+      ValExprs, ClauseNameLoc, T.getOpenLocation(), T.getCloseLocation());
+}
+
 /// Parsing of OpenMP clauses.
 ///
 ///    clause:
 ///       if-clause | final-clause | num_threads-clause | safelen-clause |
 ///       default-clause | private-clause | firstprivate-clause | shared-clause
 ///       | linear-clause | aligned-clause | collapse-clause |
 ///       lastprivate-clause | reduction-clause | proc_bind-clause |
 ///       schedule-clause | copyin-clause | copyprivate-clause | untied-clause |
OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind,
   case OMPC_threadprivate:
   case OMPC_uniform:
   case OMPC_match:
     if (!WrongDirective)
       Diag(Tok, diag::err_omp_unexpected_clause)
           << getOpenMPClauseName(CKind) << getOpenMPDirectiveName(DKind);
     SkipUntil(tok::comma, tok::annot_pragma_openmp_end, StopBeforeMatch);
     break;
+  case OMPC_sizes:
+    if (!FirstClause) {
+      Diag(Tok, diag::err_omp_more_one_clause)
+          << getOpenMPDirectiveName(DKind) << getOpenMPClauseName(CKind) << 0;
+      ErrorFound = true;
+    }
+
+    Clause = ParseOpenMPSizesClause();
+    break;
   }
   return ErrorFound ? nullptr : Clause;
 }
 
 /// Parses simple expression in parens for single-expression clauses of OpenMP
 /// constructs.
 /// \param RLoc Returned location of right paren.
 ExprResult Parser::ParseOpenMPParensExpr(StringRef ClauseName,

clang/lib/Sema/SemaExceptionSpec.cpp

CanThrowResult Sema::canThrow(const Stmt *S) {
   case Stmt::OMPTaskLoopSimdDirectiveClass:
   case Stmt::OMPTaskwaitDirectiveClass:
   case Stmt::OMPTaskyieldDirectiveClass:
   case Stmt::OMPTeamsDirectiveClass:
   case Stmt::OMPTeamsDistributeDirectiveClass:
   case Stmt::OMPTeamsDistributeParallelForDirectiveClass:
   case Stmt::OMPTeamsDistributeParallelForSimdDirectiveClass:
   case Stmt::OMPTeamsDistributeSimdDirectiveClass:
+  case Stmt::OMPTileDirectiveClass:
   case Stmt::ReturnStmtClass:
   case Stmt::SEHExceptStmtClass:
   case Stmt::SEHFinallyStmtClass:
   case Stmt::SEHLeaveStmtClass:
   case Stmt::SEHTryStmtClass:
   case Stmt::SwitchStmtClass:
   case Stmt::WhileStmtClass:
     return canSubStmtsThrow(*this, S);

clang/lib/Sema/SemaOpenMP.cpp

class DSAAttrChecker final : public StmtVisitor<DSAAttrChecker, void> {
   CapturedStmt *CS = nullptr;
   llvm::SmallVector<Expr *, 4> ImplicitFirstprivate;
   llvm::SmallVector<Expr *, 4> ImplicitMap[OMPC_MAP_delete];
   Sema::VarsWithInheritedDSAType VarsWithInheritedDSA;
   llvm::SmallDenseSet<const ValueDecl *, 4> ImplicitDeclarations;
 
   void VisitSubCaptures(OMPExecutableDirective *S) {
     // Check implicitly captured variables.
-    if (!S->hasAssociatedStmt() || !S->getAssociatedStmt())
+    if (!S->hasAssociatedStmt() || !S->getAssociatedStmt() ||
+        isOpenMPLoopTransformationDirective(S->getDirectiveKind()))
       return;
     visitSubCaptures(S->getInnermostCapturedStmt());
     // Try to capture inner this->member references to generate correct mappings
     // and diagnostics.
     if (TryCaptureCXXThisMembers ||
         (isOpenMPTargetExecutionDirective(Stack->getCurrentDirective()) &&
          llvm::any_of(S->getInnermostCapturedStmt()->captures(),
                       [](const CapturedStmt::Capture &C) {
for (OMPClause *C : S->clauses()) {
           if (CC)
             Visit(CC);
         }
       }
     }
     // Check implicitly captured variables.
     VisitSubCaptures(S);
   }
+
+  void VisitOMPTileDirective(OMPTileDirective *S) {
+    // #pragma omp tile does not introduce data sharing
+    VisitStmt(S);
+  }
+
   void VisitStmt(Stmt *S) {
     for (Stmt *C : S->children()) {
       if (C) {
         // Check implicitly captured variables in the task-based directives to
         // check if they must be firstprivatized.
         Visit(C);
       }
     }
case OMPD_target_exit_data: {
     // Mark this captured region as inlined, because we don't use outlined
     // function directly.
     getCurCapturedRegion()->TheCapturedDecl->addAttr(
         AlwaysInlineAttr::CreateImplicit(
             Context, {}, AttributeCommonInfo::AS_Keyword,
             AlwaysInlineAttr::Keyword_forceinline));
     break;
   }
+  case OMPD_tile:
+    // loop transformations have no captures.
+    break;
   case OMPD_threadprivate:
   case OMPD_allocate:
   case OMPD_taskyield:
   case OMPD_barrier:
   case OMPD_taskwait:
   case OMPD_cancellation_point:
   case OMPD_cancel:
   case OMPD_flush:
if (isOpenMPPrivate(Clause->getClauseKind()) ||
       DSAStack->setForceVarCapturing(Clause->getClauseKind() == OMPC_copyin);
       // Mark all variables in private list clauses as used in inner region.
       for (Stmt *VarRef : Clause->children()) {
         if (auto *E = cast_or_null<Expr>(VarRef)) {
           MarkDeclarationsReferencedInExpr(E);
         }
       }
       DSAStack->setForceVarCapturing(/*V=*/false);
+    } else if (isOpenMPLoopTransformationDirective(
+                   DSAStack->getCurrentDirective())) {
+      assert(CaptureRegions.empty());
     } else if (CaptureRegions.size() > 1 ||
                CaptureRegions.back() != OMPD_unknown) {
       if (auto *C = OMPClauseWithPreInit::get(Clause))
         PICs.push_back(C);
       if (auto *C = OMPClauseWithPostUpdate::get(Clause)) {
         if (Expr *E = C->getPostUpdateExpr())
           MarkDeclarationsReferencedInExpr(E);
       }
if (checkCancelRegion(*this, Kind, CancelRegion, StartLoc) ||
       checkNestingOfRegions(*this, DSAStack, Kind, DirName, CancelRegion,
                             StartLoc))
     return StmtError();
 
   llvm::SmallVector<OMPClause *, 8> ClausesWithImplicit;
   VarsWithInheritedDSAType VarsWithInheritedDSA;
   bool ErrorFound = false;
   ClausesWithImplicit.append(Clauses.begin(), Clauses.end());
-  if (AStmt && !CurContext->isDependentContext()) {
+  if (AStmt && !CurContext->isDependentContext() &&
+      !isOpenMPLoopTransformationDirective(Kind)) {
     assert(isa<CapturedStmt>(AStmt) && "Captured statement expected");
 
     // Check default data sharing attributes for referenced variables.
     DSAAttrChecker DSAChecker(DSAStack, *this, cast<CapturedStmt>(AStmt));
     int ThisCaptureLevel = getOpenMPCaptureLevels(Kind);
     Stmt *S = AStmt;
     while (--ThisCaptureLevel >= 0)
       S = cast<CapturedStmt>(S)->getCapturedStmt();
case OMPD_target_teams_distribute_parallel_for_simd:
     break;
   case OMPD_target_teams_distribute_simd:
     Res = ActOnOpenMPTargetTeamsDistributeSimdDirective(
         ClausesWithImplicit, AStmt, StartLoc, EndLoc, VarsWithInheritedDSA);
     AllowedNameModifiers.push_back(OMPD_target);
     if (LangOpts.OpenMP >= 50)
       AllowedNameModifiers.push_back(OMPD_simd);
     break;
+  case OMPD_tile:
+    Res = ActOnOpenMPTileDirective(ClausesWithImplicit, AStmt, StartLoc, EndLoc,
+                                   VarsWithInheritedDSA);
+    break;
   case OMPD_declare_target:
   case OMPD_end_declare_target:
   case OMPD_threadprivate:
   case OMPD_allocate:
   case OMPD_declare_reduction:
   case OMPD_declare_mapper:
   case OMPD_declare_simd:
   case OMPD_requires:
for (OMPClause *C : Clauses) {
       case OMPC_from:
       case OMPC_use_device_ptr:
       case OMPC_is_device_ptr:
       case OMPC_nontemporal:
       case OMPC_order:
       case OMPC_destroy:
       case OMPC_inclusive:
       case OMPC_exclusive:
+      case OMPC_sizes:
         continue;
       case OMPC_allocator:
       case OMPC_flush:
       case OMPC_depobj:
       case OMPC_threadprivate:
       case OMPC_uniform:
       case OMPC_unknown:
       case OMPC_unified_address:
StmtResult Sema::ActOnOpenMPParallelDirective(ArrayRef<OMPClause *> Clauses,
 
   return OMPParallelDirective::Create(Context, StartLoc, EndLoc, Clauses, AStmt,
                                       DSAStack->isCancelRegion());
 }
 
 namespace {
 /// Iteration space of a single for loop.
 struct LoopIterationSpace final {
+  /// The loop that has been analyzed.
+  Stmt *Loop = nullptr;
+  /// The loop's body.
+  Stmt *Body = nullptr;
   /// True if the condition operator is the strict compare operator (<, > or
   /// !=).
   bool IsStrictCompare = false;
   /// Condition of the loop.
   Expr *PreCond = nullptr;
   /// This expression calculates the number of iterations in the loop.
   /// It is always possible to calculate it before starting the loop.
   Expr *NumIterations = nullptr;

 };
 
 /// Helper class for checking canonical form of the OpenMP loops and
 /// extracting iteration space of each loop in the loop nest, that will be used
 /// for IR generation.
 class OpenMPIterationSpaceChecker {
   /// Reference to Sema.
   Sema &SemaRef;
+  /// Is the loop associated directive capturing its referenced variables?
+  bool Capturing;
+  /// Does the loop associated directive support non-rectangular loops?
+  bool SupportsNonRectangular;
   /// Data-sharing stack.
   DSAStackTy &Stack;
   /// A location for diagnostics (when there is no some better location).
   SourceLocation DefaultLoc;
   /// A location for diagnostics (when increment is not compatible).
   SourceLocation ConditionLoc;
   /// A source location for referring to loop init later.
   SourceRange InitSrcRange;
class OpenMPIterationSpaceChecker {
   Optional<unsigned> CondDependOnLC;
   /// Checks if the provide statement depends on the loop counter.
   Optional<unsigned> doesDependOnLoopCounter(const Stmt *S, bool IsInitializer);
   /// Original condition required for checking of the exit condition for
   /// non-rectangular loop.
   Expr *Condition = nullptr;
 
 public:
-  OpenMPIterationSpaceChecker(Sema &SemaRef, DSAStackTy &Stack,
+  OpenMPIterationSpaceChecker(Sema &SemaRef, bool Capturing,
+                              bool SupportsNonRectangular, DSAStackTy &Stack,
                               SourceLocation DefaultLoc)
-      : SemaRef(SemaRef), Stack(Stack), DefaultLoc(DefaultLoc),
-        ConditionLoc(DefaultLoc) {}
+      : SemaRef(SemaRef), Capturing(Capturing),
+        SupportsNonRectangular(SupportsNonRectangular), Stack(Stack),
+        DefaultLoc(DefaultLoc), ConditionLoc(DefaultLoc) {}
   /// Check init-expr for canonical loop form and save loop counter
   /// variable - #Var and its initialization value - #LB.
   bool checkAndSetInit(Stmt *S, bool EmitDiags = true);
   /// Check test-expr for canonical form, save upper-bound (#UB), flags
   /// for less/greater and for strict/non-strict comparison.
-  bool checkAndSetCond(Expr *S);
+  bool checkAndSetCond(Expr *S, bool EmitDiags = true);
   /// Check incr-expr for canonical loop form and return true if it
   /// does not conform, otherwise save loop step (#Step).
   bool checkAndSetInc(Expr *S);
   /// Return the loop counter variable.
   ValueDecl *getLoopDecl() const { return LCDecl; }
   /// Return the reference expression to loop counter variable.
   Expr *getLoopDeclRefExpr() const { return LCRef; }
   /// Source range of the loop init.
private:
   /// Check the right-hand side of an assignment in the increment
   /// expression.
   bool checkAndSetIncRHS(Expr *RHS);
   /// Helper to set loop counter variable and its initializer.
   bool setLCDeclAndLB(ValueDecl *NewLCDecl, Expr *NewDeclRefExpr, Expr *NewLB,
                       bool EmitDiags);
   /// Helper to set upper bound.
   bool setUB(Expr *NewUB, llvm::Optional<bool> LessOp, bool StrictOp,
-             SourceRange SR, SourceLocation SL);
+             SourceRange SR, SourceLocation SL, bool EmitDiags = true);
   /// Helper to set loop increment.
   bool setStep(Expr *NewStep, bool Subtract);
 };
 
 bool OpenMPIterationSpaceChecker::dependent() const {
   if (!LCDecl) {
     assert(!LB && !UB && !Step);
     return false;
bool OpenMPIterationSpaceChecker::setLCDeclAndLB(ValueDecl *NewLCDecl,
   if (EmitDiags)
     InitDependOnLC = doesDependOnLoopCounter(LB, /*IsInitializer=*/true);
   return false;
 }
 
 bool OpenMPIterationSpaceChecker::setUB(Expr *NewUB,
                                         llvm::Optional<bool> LessOp,
                                         bool StrictOp, SourceRange SR,
-                                        SourceLocation SL) {
+                                        SourceLocation SL, bool EmitDiags) {
   // State consistency checking to ensure correct usage.
   assert(LCDecl != nullptr && LB != nullptr && UB == nullptr &&
          Step == nullptr && !TestIsLessOp && !TestIsStrictOp);
   if (!NewUB)
     return true;
   UB = NewUB;
   if (LessOp)
     TestIsLessOp = LessOp;
   TestIsStrictOp = StrictOp;
   ConditionSrcRange = SR;
   ConditionLoc = SL;
+  if (EmitDiags)
     CondDependOnLC = doesDependOnLoopCounter(UB, /*IsInitializer=*/false);
   return false;
 }
 
 bool OpenMPIterationSpaceChecker::setStep(Expr *NewStep, bool Subtract) {
   // State consistency checking to ensure correct usage.
   assert(LCDecl != nullptr && LB != nullptr && Step == nullptr);
   if (!NewStep)
     return true;

 class LoopCounterRefChecker final
     : public ConstStmtVisitor<LoopCounterRefChecker, bool> {
   Sema &SemaRef;
   DSAStackTy &Stack;
   const ValueDecl *CurLCDecl = nullptr;
   const ValueDecl *DepDecl = nullptr;
   const ValueDecl *PrevDepDecl = nullptr;
   bool IsInitializer = true;
+  bool SupportsNonRectangular;
   unsigned BaseLoopId = 0;
   bool checkDecl(const Expr *E, const ValueDecl *VD) {
     if (getCanonicalDecl(VD) == getCanonicalDecl(CurLCDecl)) {
       SemaRef.Diag(E->getExprLoc(), diag::err_omp_stmt_depends_on_loop_counter)
           << (IsInitializer ? 0 : 1);
       return false;
     }
     const auto &&Data = Stack.isLoopControlVariable(VD);
     // OpenMP, 2.9.1 Canonical Loop Form, Restrictions.
     // The type of the loop iterator on which we depend may not have a random
     // access iterator type.
     if (Data.first && VD->getType()->isRecordType()) {
       SmallString<128> Name;
       llvm::raw_svector_ostream OS(Name);
       VD->getNameForDiagnostic(OS, SemaRef.getPrintingPolicy(),
                                /*Qualified=*/true);
       SemaRef.Diag(E->getExprLoc(),
                    diag::err_omp_wrong_dependency_iterator_type)
           << OS.str();
       SemaRef.Diag(VD->getLocation(), diag::note_previous_decl) << VD;
       return false;
     }
+    if (Data.first && !SupportsNonRectangular) {
+      SemaRef.Diag(E->getExprLoc(), diag::err_omp_invariant_dependency);
+      return false;
+    }
     if (Data.first &&
         (DepDecl || (PrevDepDecl &&
                      getCanonicalDecl(VD) != getCanonicalDecl(PrevDepDecl)))) {
       if (!DepDecl && PrevDepDecl)
         DepDecl = PrevDepDecl;
       SmallString<128> Name;
       llvm::raw_svector_ostream OS(Name);
       DepDecl->getNameForDiagnostic(OS, SemaRef.getPrintingPolicy(),
public:
   bool VisitStmt(const Stmt *S) {
     bool Res = false;
     for (const Stmt *Child : S->children())
       Res = (Child && Visit(Child)) || Res;
     return Res;
   }
   explicit LoopCounterRefChecker(Sema &SemaRef, DSAStackTy &Stack,
                                  const ValueDecl *CurLCDecl, bool IsInitializer,
-                                 const ValueDecl *PrevDepDecl = nullptr)
+                                 const ValueDecl *PrevDepDecl = nullptr,
+                                 bool SupportsNonRectangular = true)
       : SemaRef(SemaRef), Stack(Stack), CurLCDecl(CurLCDecl),
-        PrevDepDecl(PrevDepDecl), IsInitializer(IsInitializer) {}
+        PrevDepDecl(PrevDepDecl), IsInitializer(IsInitializer),
+        SupportsNonRectangular(SupportsNonRectangular) {}
   unsigned getBaseLoopId() const {
     assert(CurLCDecl && "Expected loop dependency.");
     return BaseLoopId;
   }
   const ValueDecl *getDepDecl() const {
     assert(CurLCDecl && "Expected loop dependency.");
     return DepDecl;
   }
 };
 } // namespace
 
 Optional<unsigned>
 OpenMPIterationSpaceChecker::doesDependOnLoopCounter(const Stmt *S,
                                                      bool IsInitializer) {
   // Check for the non-rectangular loops.
   LoopCounterRefChecker LoopStmtChecker(SemaRef, Stack, LCDecl, IsInitializer,
-                                        DepDecl);
+                                        DepDecl, SupportsNonRectangular);
   if (LoopStmtChecker.Visit(S)) {
     DepDecl = LoopStmtChecker.getDepDecl();
     return LoopStmtChecker.getBaseLoopId();
   }
   return llvm::None;
 }
 
 bool OpenMPIterationSpaceChecker::checkAndSetInit(Stmt *S, bool EmitDiags) {
if (const auto *VD = dyn_cast<VarDecl>(DRE->getDecl()))
       return getCanonicalDecl(VD);
   }
   if (const auto *ME = dyn_cast_or_null<MemberExpr>(E))
     if (ME->isArrow() && isa<CXXThisExpr>(ME->getBase()->IgnoreParenImpCasts()))
       return getCanonicalDecl(ME->getMemberDecl());
   return nullptr;
 }
 
-bool OpenMPIterationSpaceChecker::checkAndSetCond(Expr *S) {
+bool OpenMPIterationSpaceChecker::checkAndSetCond(Expr *S, bool EmitDiags) {
   // Check test-expr for canonical form, save upper-bound UB, flags for
   // less/greater and for strict/non-strict comparison.
   // OpenMP [2.9] Canonical loop form. Test-expr may be one of the following:
   //   var relational-op b
   //   b relational-op var
   //
   bool IneqCondIsCanonical = SemaRef.getLangOpts().OpenMP >= 50;
   if (!S) {
     SemaRef.Diag(DefaultLoc, diag::err_omp_loop_not_canonical_cond)
         << (IneqCondIsCanonical ? 1 : 0) << LCDecl;
     return true;
   }
   Condition = S;
   S = getExprAsWritten(S);
   SourceLocation CondLoc = S->getBeginLoc();
   if (auto *BO = dyn_cast<BinaryOperator>(S)) {
     if (BO->isRelationalOp()) {
       if (getInitLCDecl(BO->getLHS()) == LCDecl)
         return setUB(BO->getRHS(),
                      (BO->getOpcode() == BO_LT || BO->getOpcode() == BO_LE),
                      (BO->getOpcode() == BO_LT || BO->getOpcode() == BO_GT),
-                     BO->getSourceRange(), BO->getOperatorLoc());
+                     BO->getSourceRange(), BO->getOperatorLoc(), EmitDiags);
       if (getInitLCDecl(BO->getRHS()) == LCDecl)
         return setUB(BO->getLHS(),
                      (BO->getOpcode() == BO_GT || BO->getOpcode() == BO_GE),
                      (BO->getOpcode() == BO_LT || BO->getOpcode() == BO_GT),
-                     BO->getSourceRange(), BO->getOperatorLoc());
+                     BO->getSourceRange(), BO->getOperatorLoc(), EmitDiags);
     } else if (IneqCondIsCanonical && BO->getOpcode() == BO_NE)
-      return setUB(
-          getInitLCDecl(BO->getLHS()) == LCDecl ? BO->getRHS() : BO->getLHS(),
+      return setUB(getInitLCDecl(BO->getLHS()) == LCDecl ? BO->getRHS()
+                                                         : BO->getLHS(),
                    /*LessOp=*/llvm::None,
-          /*StrictOp=*/true, BO->getSourceRange(), BO->getOperatorLoc());
+                   /*StrictOp=*/true, BO->getSourceRange(),
+                   BO->getOperatorLoc(), EmitDiags);
   } else if (auto *CE = dyn_cast<CXXOperatorCallExpr>(S)) {
     if (CE->getNumArgs() == 2) {
       auto Op = CE->getOperator();
       switch (Op) {
       case OO_Greater:
       case OO_GreaterEqual:
       case OO_Less:
       case OO_LessEqual:
         if (getInitLCDecl(CE->getArg(0)) == LCDecl)
           return setUB(CE->getArg(1), Op == OO_Less || Op == OO_LessEqual,
                        Op == OO_Less || Op == OO_Greater, CE->getSourceRange(),
-                       CE->getOperatorLoc());
+                       CE->getOperatorLoc(), EmitDiags);
         if (getInitLCDecl(CE->getArg(1)) == LCDecl)
           return setUB(CE->getArg(0), Op == OO_Greater || Op == OO_GreaterEqual,
                        Op == OO_Less || Op == OO_Greater, CE->getSourceRange(),
-                       CE->getOperatorLoc());
+                       CE->getOperatorLoc(), EmitDiags);
         break;
       case OO_ExclaimEqual:
         if (IneqCondIsCanonical)
           return setUB(getInitLCDecl(CE->getArg(0)) == LCDecl ? CE->getArg(1)
                                                               : CE->getArg(0),
                        /*LessOp=*/llvm::None,
                        /*StrictOp=*/true, CE->getSourceRange(),
-                       CE->getOperatorLoc());
+                       CE->getOperatorLoc(), EmitDiags);
         break;
       default:
         break;
       }
     }
   }
   if (dependent() || SemaRef.CurContext->isDependentContext())
     return false;
if (dependent() || SemaRef.CurContext->isDependentContext())
     return false;
   SemaRef.Diag(S->getBeginLoc(), diag::err_omp_loop_not_canonical_incr)
       << S->getSourceRange() << LCDecl;
   return true;
 }
 
 static ExprResult
 tryBuildCapture(Sema &SemaRef, Expr *Capture,
-                llvm::MapVector<const Expr *, DeclRefExpr *> &Captures) {
-  if (SemaRef.CurContext->isDependentContext())
+                llvm::MapVector<const Expr *, DeclRefExpr *> &Captures,
+                bool Capturing = true) {
+  if (!Capturing || SemaRef.CurContext->isDependentContext())
     return ExprResult(Capture);
   if (Capture->isEvaluatable(SemaRef.Context, Expr::SE_AllowSideEffects))
     return SemaRef.PerformImplicitConversion(
         Capture->IgnoreImpCasts(), Capture->getType(), Sema::AA_Converting,
         /*AllowExplicit=*/true);
   auto I = Captures.find(Capture);
   if (I != Captures.end())
     return buildCapture(SemaRef, Capture, I->second);
if (InitDependOnLC) {
           SemaRef.BuildBinOp(S, DefaultLoc, BO_Comma, LBMaxVal.get(), LBVal);
       if (!LBMaxVal.isUsable())
         return nullptr;
       // (OuterVar = Max, LBVal)
       LBMaxVal = SemaRef.ActOnParenExpr(DefaultLoc, DefaultLoc, LBMaxVal.get());
       if (!LBMaxVal.isUsable())
         return nullptr;
 
-      Expr *LBMin = tryBuildCapture(SemaRef, LBMinVal.get(), Captures).get();
-      Expr *LBMax = tryBuildCapture(SemaRef, LBMaxVal.get(), Captures).get();
+      Expr *LBMin =
+          tryBuildCapture(SemaRef, LBMinVal.get(), Captures, Capturing).get();
+      Expr *LBMax =
+          tryBuildCapture(SemaRef, LBMaxVal.get(), Captures, Capturing).get();
       if (!LBMin || !LBMax)
         return nullptr;
       // LB(MinVal) < LB(MaxVal)
       ExprResult MinLessMaxRes =
           SemaRef.BuildBinOp(S, DefaultLoc, BO_LT, LBMin, LBMax);
       if (!MinLessMaxRes.isUsable())
         return nullptr;
       Expr *MinLessMax =
-          tryBuildCapture(SemaRef, MinLessMaxRes.get(), Captures).get();
+          tryBuildCapture(SemaRef, MinLessMaxRes.get(), Captures, Capturing)
+              .get();
       if (!MinLessMax)
         return nullptr;
       if (TestIsLessOp.getValue()) {
         // LB(MinVal) < LB(MaxVal) ? LB(MinVal) : LB(MaxVal) - min(LB(MinVal),
         // LB(MaxVal))
         ExprResult MinLB = SemaRef.ActOnConditionalOp(DefaultLoc, DefaultLoc,
                                                       MinLessMax, LBMin, LBMax);
         if (!MinLB.isUsable())
if (CondDependOnLC) {
           SemaRef.BuildBinOp(S, DefaultLoc, BO_Comma, UBMaxVal.get(), UBVal);
       if (!UBMaxVal.isUsable())
         return nullptr;
       // (OuterVar = Max, UBVal)
       UBMaxVal = SemaRef.ActOnParenExpr(DefaultLoc, DefaultLoc, UBMaxVal.get());
       if (!UBMaxVal.isUsable())
         return nullptr;
 
-      Expr *UBMin = tryBuildCapture(SemaRef, UBMinVal.get(), Captures).get();
-      Expr *UBMax = tryBuildCapture(SemaRef, UBMaxVal.get(), Captures).get();
+      Expr *UBMin =
+          tryBuildCapture(SemaRef, UBMinVal.get(), Captures, Capturing).get();
+      Expr *UBMax =
+          tryBuildCapture(SemaRef, UBMaxVal.get(), Captures, Capturing).get();
       if (!UBMin || !UBMax)
         return nullptr;
       // UB(MinVal) > UB(MaxVal)
       ExprResult MinGreaterMaxRes =
           SemaRef.BuildBinOp(S, DefaultLoc, BO_GT, UBMin, UBMax);
       if (!MinGreaterMaxRes.isUsable())
         return nullptr;
       Expr *MinGreaterMax =
-          tryBuildCapture(SemaRef, MinGreaterMaxRes.get(), Captures).get();
+          tryBuildCapture(SemaRef, MinGreaterMaxRes.get(), Captures, Capturing)
+              .get();
       if (!MinGreaterMax)
         return nullptr;
       if (TestIsLessOp.getValue()) {
         // UB(MinVal) > UB(MaxVal) ? UB(MinVal) : UB(MaxVal) - max(UB(MinVal),
         // UB(MaxVal))
         ExprResult MaxUB = SemaRef.ActOnConditionalOp(
             DefaultLoc, DefaultLoc, MinGreaterMax, UBMin, UBMax);
         if (!MaxUB.isUsable())
           return nullptr;
         UBVal = MaxUB.get();
       } else {
         // UB(MinVal) > UB(MaxVal) ? UB(MaxVal) : UB(MinVal) - min(UB(MinVal),
         // UB(MaxVal))
         ExprResult MinUB = SemaRef.ActOnConditionalOp(
             DefaultLoc, DefaultLoc, MinGreaterMax, UBMax, UBMin);
         if (!MinUB.isUsable())
           return nullptr;
         UBVal = MinUB.get();
       }
     }
     // Upper - Lower
     Expr *UBExpr = TestIsLessOp.getValue() ? UBVal : LBVal;
     Expr *LBExpr = TestIsLessOp.getValue() ? LBVal : UBVal;
-    Expr *Upper = tryBuildCapture(SemaRef, UBExpr, Captures).get();
-    Expr *Lower = tryBuildCapture(SemaRef, LBExpr, Captures).get();
+    Expr *Upper = tryBuildCapture(SemaRef, UBExpr, Captures, Capturing).get();
+    Expr *Lower = tryBuildCapture(SemaRef, LBExpr, Captures, Capturing).get();
     if (!Upper || !Lower)
       return nullptr;
 
     Diff = SemaRef.BuildBinOp(S, DefaultLoc, BO_Sub, Upper, Lower);
 
     if (!Diff.isUsable() && VarType->getAsCXXRecordDecl()) {
       // BuildBinOp already emitted error, this one is to point user to upper
       // and lower bound, and to tell what is passed to 'operator-'.
Expr *OpenMPIterationSpaceChecker::buildNumIterations(
   if (TestIsStrictOp)
     Diff = SemaRef.BuildBinOp(
         S, DefaultLoc, BO_Sub, Diff.get(),
         SemaRef.ActOnIntegerConstant(SourceLocation(), 1).get());
   if (!Diff.isUsable())
     return nullptr;
 
   // Upper - Lower [- 1] + Step
-  ExprResult NewStep = tryBuildCapture(SemaRef, Step, Captures);
+  ExprResult NewStep = tryBuildCapture(SemaRef, Step, Captures, Capturing);
   if (!NewStep.isUsable())
     return nullptr;
   Diff = SemaRef.BuildBinOp(S, DefaultLoc, BO_Add, Diff.get(), NewStep.get());
   if (!Diff.isUsable())
     return nullptr;
 
   // Parentheses (for dumping/debugging purposes only).
   Diff = SemaRef.ActOnParenExpr(DefaultLoc, DefaultLoc, Diff.get());
std::pair<Expr *, Expr *> OpenMPIterationSpaceChecker::buildMinMaxValues(
   Expr *MaxExpr = nullptr;
   Expr *LBExpr = TestIsLessOp.getValue() ? LB : UB;
   Expr *UBExpr = TestIsLessOp.getValue() ? UB : LB;
   bool LBNonRect = TestIsLessOp.getValue() ? InitDependOnLC.hasValue()
                                            : CondDependOnLC.hasValue();
   bool UBNonRect = TestIsLessOp.getValue() ? CondDependOnLC.hasValue()
                                            : InitDependOnLC.hasValue();
   Expr *Lower =
-      LBNonRect ? LBExpr : tryBuildCapture(SemaRef, LBExpr, Captures).get();
+      LBNonRect ? LBExpr
+                : tryBuildCapture(SemaRef, LBExpr, Captures, Capturing).get();
   Expr *Upper =
-      UBNonRect ? UBExpr : tryBuildCapture(SemaRef, UBExpr, Captures).get();
+      UBNonRect ? UBExpr
+                : tryBuildCapture(SemaRef, UBExpr, Captures, Capturing).get();
   if (!Upper || !Lower)
     return std::make_pair(nullptr, nullptr);
 
   if (TestIsLessOp.getValue())
     MinExpr = Lower;
   else
     MaxExpr = Upper;
 
std::pair<Expr *, Expr *> OpenMPIterationSpaceChecker::buildMinMaxValues(
   if (TestIsStrictOp)
     Diff = SemaRef.BuildBinOp(
         S, DefaultLoc, BO_Sub, Diff.get(),
         SemaRef.ActOnIntegerConstant(SourceLocation(), 1).get());
   if (!Diff.isUsable())
     return std::make_pair(nullptr, nullptr);
 
   // Upper - Lower [- 1] + Step
-  ExprResult NewStep = tryBuildCapture(SemaRef, Step, Captures);
+  ExprResult NewStep = tryBuildCapture(SemaRef, Step, Captures, Capturing);
   if (!NewStep.isUsable())
     return std::make_pair(nullptr, nullptr);
 
   // Parentheses (for dumping/debugging purposes only).
   Diff = SemaRef.ActOnParenExpr(DefaultLoc, DefaultLoc, Diff.get());
   if (!Diff.isUsable())
     return std::make_pair(nullptr, nullptr);
 
if (CondDependOnLC || InitDependOnLC)
     return SemaRef.PerformImplicitConversion(
         SemaRef.ActOnIntegerConstant(SourceLocation(), 1).get(),
         SemaRef.Context.BoolTy, /*Action=*/Sema::AA_Casting,
         /*AllowExplicit=*/true).get();
 
   // Try to build LB <op> UB, where <op> is <, >, <=, or >=.
   Sema::TentativeAnalysisScope Trap(SemaRef);
 
-  ExprResult NewLB = tryBuildCapture(SemaRef, LB, Captures);
-  ExprResult NewUB = tryBuildCapture(SemaRef, UB, Captures);
+  ExprResult NewLB = tryBuildCapture(SemaRef, LB, Captures, Capturing);
+  ExprResult NewUB = tryBuildCapture(SemaRef, UB, Captures, Capturing);
   if (!NewLB.isUsable() || !NewUB.isUsable())
     return nullptr;
 
   ExprResult CondExpr =
       SemaRef.BuildBinOp(S, DefaultLoc,
                          TestIsLessOp.getValue() ?
                            (TestIsStrictOp ? BO_LT : BO_LE) :
                            (TestIsStrictOp ? BO_GT : BO_GE),
Expr *OpenMPIterationSpaceChecker::buildOrderedLoopData(
   }
   ExprResult Diff;
   QualType VarType = LCDecl->getType().getNonReferenceType();
   if (VarType->isIntegerType() || VarType->isPointerType() ||
       SemaRef.getLangOpts().CPlusPlus) {
     // Upper - Lower
     Expr *Upper = TestIsLessOp.getValue()
                       ? Cnt
-                      : tryBuildCapture(SemaRef, LB, Captures).get();
+                      : tryBuildCapture(SemaRef, LB, Captures, Capturing).get();
     Expr *Lower = TestIsLessOp.getValue()
-                      ? tryBuildCapture(SemaRef, LB, Captures).get()
+                      ? tryBuildCapture(SemaRef, LB, Captures, Capturing).get()
                       : Cnt;
     if (!Upper || !Lower)
       return nullptr;
 
     Diff = SemaRef.BuildBinOp(S, DefaultLoc, BO_Sub, Upper, Lower);
 
     if (!Diff.isUsable() && VarType->getAsCXXRecordDecl()) {
       // BuildBinOp already emitted error, this one is to point user to upper
       // and lower bound, and to tell what is passed to 'operator-'.
       SemaRef.Diag(Upper->getBeginLoc(), diag::err_omp_loop_diff_cxx)
           << Upper->getSourceRange() << Lower->getSourceRange();
       return nullptr;
     }
   }
 
   if (!Diff.isUsable())
     return nullptr;
 
   // Parentheses (for dumping/debugging purposes only).
   Diff = SemaRef.ActOnParenExpr(DefaultLoc, DefaultLoc, Diff.get());
   if (!Diff.isUsable())
     return nullptr;
 
-  ExprResult NewStep = tryBuildCapture(SemaRef, Step, Captures);
+  ExprResult NewStep = tryBuildCapture(SemaRef, Step, Captures, Capturing);
   if (!NewStep.isUsable())
     return nullptr;
   // (Upper - Lower) / Step
   Diff = SemaRef.BuildBinOp(S, DefaultLoc, BO_Div, Diff.get(), NewStep.get());
   if (!Diff.isUsable())
     return nullptr;
 
   return Diff.get();
 }
 } // namespace
 
 void Sema::ActOnOpenMPLoopInitialization(SourceLocation ForLoc, Stmt *Init) {
   assert(getLangOpts().OpenMP && "OpenMP is not active.");
   assert(Init && "Expected loop in canonical form.");
   unsigned AssociatedLoops = DSAStack->getAssociatedLoops();
   if (AssociatedLoops > 0 &&
       isOpenMPLoopDirective(DSAStack->getCurrentDirective())) {
     DSAStack->loopStart();
-    OpenMPIterationSpaceChecker ISC(*this, *DSAStack, ForLoc);
+    OpenMPIterationSpaceChecker ISC(*this, true, true, *DSAStack, ForLoc);
     if (!ISC.checkAndSetInit(Init, /*EmitDiags=*/false)) {
       if (ValueDecl *D = ISC.getLoopDecl()) {
         auto *VD = dyn_cast<VarDecl>(D);
         DeclRefExpr *PrivateRef = nullptr;
         if (!VD) {
           if (VarDecl *Private = isOpenMPCapturedDecl(D)) {
             VD = Private;
           } else {

 /// for further processing (such as collapsing).
 static bool checkOpenMPIterationSpace(
     OpenMPDirectiveKind DKind, Stmt *S, Sema &SemaRef, DSAStackTy &DSA,
     unsigned CurrentNestedLoopCount, unsigned NestedLoopCount,
     unsigned TotalNestedLoopCount, Expr *CollapseLoopCountExpr,
     Expr *OrderedLoopCountExpr,
     Sema::VarsWithInheritedDSAType &VarsWithImplicitDSA,
     llvm::MutableArrayRef<LoopIterationSpace> ResultIterSpaces,
-    llvm::MapVector<const Expr *, DeclRefExpr *> &Captures) {
+    llvm::MapVector<const Expr *, DeclRefExpr *> &Captures, bool Capturing,
+    bool SupportsNonRectangular, bool EmitDiags) {
   // OpenMP [2.9.1, Canonical Loop Form]
   //   for (init-expr; test-expr; incr-expr) structured-block
   //   for (range-decl: range-expr) structured-block
   auto *For = dyn_cast_or_null<ForStmt>(S);
   auto *CXXFor = dyn_cast_or_null<CXXForRangeStmt>(S);
   // Ranged for is supported only in OpenMP 5.0.
   if (!For && (SemaRef.LangOpts.OpenMP <= 45 || !CXXFor)) {
     SemaRef.Diag(S->getBeginLoc(), diag::err_omp_not_for)
         << (CollapseLoopCountExpr != nullptr || OrderedLoopCountExpr != nullptr)
         << getOpenMPDirectiveName(DKind) << TotalNestedLoopCount
         << (CurrentNestedLoopCount > 0) << CurrentNestedLoopCount;
-    if (TotalNestedLoopCount > 1) {
+    if (TotalNestedLoopCount > 1 &&
+        (CollapseLoopCountExpr || OrderedLoopCountExpr)) {
       if (CollapseLoopCountExpr && OrderedLoopCountExpr)
         SemaRef.Diag(DSA.getConstructLoc(),
                      diag::note_omp_collapse_ordered_expr)
             << 2 << CollapseLoopCountExpr->getSourceRange()
             << OrderedLoopCountExpr->getSourceRange();
       else if (CollapseLoopCountExpr)
         SemaRef.Diag(CollapseLoopCountExpr->getExprLoc(),
                      diag::note_omp_collapse_ordered_expr)
             << 0 << CollapseLoopCountExpr->getSourceRange();
       else
         SemaRef.Diag(OrderedLoopCountExpr->getExprLoc(),
                      diag::note_omp_collapse_ordered_expr)
             << 1 << OrderedLoopCountExpr->getSourceRange();
     }
     return true;
   }
-  assert(((For && For->getBody()) || (CXXFor && CXXFor->getBody())) &&
-         "No loop body.");
+  ResultIterSpaces[CurrentNestedLoopCount].Loop = S;
+  Stmt *&Body = ResultIterSpaces[CurrentNestedLoopCount].Body;
+  Body = For ? For->getBody() : CXXFor->getBody();
+  assert(Body && "No loop body.");
 
-  OpenMPIterationSpaceChecker ISC(SemaRef, DSA,
+  OpenMPIterationSpaceChecker ISC(SemaRef, Capturing, SupportsNonRectangular,
+                                  DSA,
                                   For ? For->getForLoc() : CXXFor->getForLoc());
 
   // Check init.
   Stmt *Init = For ? For->getInit() : CXXFor->getBeginStmt();
-  if (ISC.checkAndSetInit(Init))
+  if (ISC.checkAndSetInit(Init, EmitDiags))
     return true;
 
   bool HasErrors = false;
 
   // Check loop variable's type.
   if (ValueDecl *LCDecl = ISC.getLoopDecl()) {
     // OpenMP [2.6, Canonical Loop Form]
     // Var is one of the following:
if (ValueDecl *LCDecl = ISC.getLoopDecl()) {
     // constant-linear-step that is the increment of the associated for-loop.
     // Exclude loop var from the list of variables with implicitly defined data
     // sharing attributes.
     VarsWithImplicitDSA.erase(LCDecl);
 
     assert(isOpenMPLoopDirective(DKind) && "DSA for non-loop vars");
 
     // Check test-expr.
-    HasErrors |= ISC.checkAndSetCond(For ? For->getCond() : CXXFor->getCond());
+    HasErrors |= ISC.checkAndSetCond(For ? For->getCond() : CXXFor->getCond(),
+                                     EmitDiags);
 
     // Check incr-expr.
     HasErrors |= ISC.checkAndSetInc(For ? For->getInc() : CXXFor->getInc());
   }
 
   if (ISC.dependent() || SemaRef.CurContext->isDependentContext() || HasErrors)
     return HasErrors;
 
   // Build the loop's iteration space representation.
   ResultIterSpaces[CurrentNestedLoopCount].PreCond = ISC.buildPreCond(
       DSA.getCurScope(), For ? For->getCond() : CXXFor->getCond(), Captures);
   ResultIterSpaces[CurrentNestedLoopCount].NumIterations =
       ISC.buildNumIterations(DSA.getCurScope(), ResultIterSpaces,
                              (isOpenMPWorksharingDirective(DKind) ||
                               isOpenMPTaskLoopDirective(DKind) ||
-                              isOpenMPDistributeDirective(DKind)),
+                              isOpenMPDistributeDirective(DKind) ||
+                              isOpenMPLoopTransformationDirective(DKind)),
                              Captures);
   ResultIterSpaces[CurrentNestedLoopCount].CounterVar =
       ISC.buildCounterVar(Captures, DSA);
   ResultIterSpaces[CurrentNestedLoopCount].PrivateCounterVar =
       ISC.buildPrivateCounterVar();
   ResultIterSpaces[CurrentNestedLoopCount].CounterInit = ISC.buildCounterInit();
   ResultIterSpaces[CurrentNestedLoopCount].CounterStep = ISC.buildCounterStep();
   ResultIterSpaces[CurrentNestedLoopCount].InitSrcRange = ISC.getInitSrcRange();
static bool checkOpenMPIterationSpace(
   ResultIterSpaces[CurrentNestedLoopCount].IsNonRectangularLB =
       ISC.doesInitDependOnLC();
   ResultIterSpaces[CurrentNestedLoopCount].IsNonRectangularUB =
       ISC.doesCondDependOnLC();
   ResultIterSpaces[CurrentNestedLoopCount].LoopDependentIdx =
       ISC.getLoopDependentIdx();
 
   HasErrors |=
-      (ResultIterSpaces[CurrentNestedLoopCount].PreCond == nullptr ||
+      (ResultIterSpaces[CurrentNestedLoopCount].Loop == nullptr ||
+       ResultIterSpaces[CurrentNestedLoopCount].Body == nullptr ||
+       ResultIterSpaces[CurrentNestedLoopCount].PreCond == nullptr ||
        ResultIterSpaces[CurrentNestedLoopCount].NumIterations == nullptr ||
        ResultIterSpaces[CurrentNestedLoopCount].CounterVar == nullptr ||
        ResultIterSpaces[CurrentNestedLoopCount].PrivateCounterVar == nullptr ||
        ResultIterSpaces[CurrentNestedLoopCount].CounterInit == nullptr ||
        ResultIterSpaces[CurrentNestedLoopCount].CounterStep == nullptr);
   if (!HasErrors && DSA.isOrderedRegion()) {
     if (DSA.getOrderedRegionParam().second->getNumForLoops()) {
       if (CurrentNestedLoopCount <
static bool checkOpenMPIterationSpace(
 
   return HasErrors;
 }
 
 /// Build 'VarRef = Start.
 static ExprResult
 buildCounterInit(Sema &SemaRef, Scope *S, SourceLocation Loc, ExprResult VarRef,
                  ExprResult Start, bool IsNonRectangularLB,
-                 llvm::MapVector<const Expr *, DeclRefExpr *> &Captures) {
+                 llvm::MapVector<const Expr *, DeclRefExpr *> &Captures,
+                 bool Capturing) {
   // Build 'VarRef = Start.
-  ExprResult NewStart = IsNonRectangularLB
-                            ? Start.get()
-                            : tryBuildCapture(SemaRef, Start.get(), Captures);
+  ExprResult NewStart =
+      IsNonRectangularLB
+          : tryBuildCapture(SemaRef, Start.get(), Captures, Capturing);
   if (!NewStart.isUsable())
     return ExprError();
   if (!SemaRef.Context.hasSameType(NewStart.get()->getType(),
                                    VarRef.get()->getType())) {
     NewStart = SemaRef.PerformImplicitConversion(
         NewStart.get(), VarRef.get()->getType(), Sema::AA_Converting,
         /*AllowExplicit=*/true);
     if (!NewStart.isUsable())

 /// Called on a for stmt to check itself and nested loops (if any).
 /// \return Returns 0 if one of the collapsed stmts is not canonical for loop,
 /// number of collapsed loops otherwise.
 static unsigned
 checkOpenMPLoop(OpenMPDirectiveKind DKind, Expr *CollapseLoopCountExpr,
                 Expr *OrderedLoopCountExpr, Stmt *AStmt, Sema &SemaRef,
                 DSAStackTy &DSA,
                 Sema::VarsWithInheritedDSAType &VarsWithImplicitDSA,
-                OMPLoopDirective::HelperExprs &Built) {
-  unsigned NestedLoopCount = 1;
+                OMPLoopDirective::HelperExprs &Built, unsigned MinLoopCount = 1,
+                bool EmitDiags = true) {
+  assert(MinLoopCount == 1 ||
+         (!CollapseLoopCountExpr && !OrderedLoopCountExpr));
+
+  unsigned NestedLoopCount = MinLoopCount;
+  bool Capturing = !isOpenMPLoopTransformationDirective(DKind);
+  bool SupportsNonPerfectlyNested = (SemaRef.LangOpts.OpenMP >= 50) &&
+                                    !isOpenMPLoopTransformationDirective(DKind);
+  bool SupportsNonRectangular = !isOpenMPLoopTransformationDirective(DKind);
+
   if (CollapseLoopCountExpr) {
     // Found 'collapse' clause - calculate collapse number.
     Expr::EvalResult Result;
     if (!CollapseLoopCountExpr->isValueDependent() &&
         CollapseLoopCountExpr->EvaluateAsInt(Result, SemaRef.getASTContext())) {
       NestedLoopCount = Result.Val.getInt().getLimitedValue();
     } else {
       Built.clear(/*Size=*/1);
if (!OrderedLoopCountExpr->isValueDependent() &&
             << CollapseLoopCountExpr->getSourceRange();
       }
       OrderedLoopCount = Result.getLimitedValue();
     } else {
       Built.clear(/*Size=*/1);
       return 1;
     }
   }
+  unsigned NumAssocociatedLoops = std::max(OrderedLoopCount, NestedLoopCount);
+
   // This is helper routine for loop directives (e.g., 'for', 'simd',
   // 'for simd', etc.).
   llvm::MapVector<const Expr *, DeclRefExpr *> Captures;
-  SmallVector<LoopIterationSpace, 4> IterSpaces(
-      std::max(OrderedLoopCount, NestedLoopCount));
-  Stmt *CurStmt = AStmt->IgnoreContainers(/* IgnoreCaptured */ true);
+  SmallVector<LoopIterationSpace, 4> IterSpaces(NumAssocociatedLoops);
+  Stmt *CurStmt = AStmt;
   for (unsigned Cnt = 0; Cnt < NestedLoopCount; ++Cnt) {
-    if (checkOpenMPIterationSpace(
-            DKind, CurStmt, SemaRef, DSA, Cnt, NestedLoopCount,
-            std::max(OrderedLoopCount, NestedLoopCount), CollapseLoopCountExpr,
-            OrderedLoopCountExpr, VarsWithImplicitDSA, IterSpaces, Captures))
+    CurStmt = OMPLoopDirective::tryToFindNextInnerLoop(
+        CurStmt, CurStmt > 0 && SupportsNonPerfectlyNested);
+
+    if (checkOpenMPIterationSpace(DKind, CurStmt, SemaRef, DSA, Cnt,
+                                  NestedLoopCount, NumAssocociatedLoops,
+                                  CollapseLoopCountExpr, OrderedLoopCountExpr,
+                                  VarsWithImplicitDSA, IterSpaces, Captures,
+                                  Capturing, SupportsNonRectangular, EmitDiags))
       return 0;
     // Move on to the next nested for loop, or to the loop body.
     // OpenMP [2.8.1, simd construct, Restrictions]
     // All loops associated with the construct must be perfectly nested; that
     // is, there must be no intervening code nor any OpenMP directive between
     // any two loops.
     if (auto *For = dyn_cast<ForStmt>(CurStmt)) {
       CurStmt = For->getBody();
     } else {
       assert(isa<CXXForRangeStmt>(CurStmt) &&
              "Expected canonical for or range-based for loops.");
       CurStmt = cast<CXXForRangeStmt>(CurStmt)->getBody();
     }
-    CurStmt = OMPLoopDirective::tryToFindNextInnerLoop(
-        CurStmt, SemaRef.LangOpts.OpenMP >= 50);
   }
   for (unsigned Cnt = NestedLoopCount; Cnt < OrderedLoopCount; ++Cnt) {
+    CurStmt = OMPLoopDirective::tryToFindNextInnerLoop(
+        CurStmt, SupportsNonPerfectlyNested);
+
     if (checkOpenMPIterationSpace(
             DKind, CurStmt, SemaRef, DSA, Cnt, NestedLoopCount,
             std::max(OrderedLoopCount, NestedLoopCount), CollapseLoopCountExpr,
-            OrderedLoopCountExpr, VarsWithImplicitDSA, IterSpaces, Captures))
+            OrderedLoopCountExpr, VarsWithImplicitDSA, IterSpaces, Captures,
+            Capturing, SupportsNonRectangular, EmitDiags))
       return 0;
     if (Cnt > 0 && IterSpaces[Cnt].CounterVar) {
       // Handle initialization of captured loop iterator variables.
       auto *DRE = cast<DeclRefExpr>(IterSpaces[Cnt].CounterVar);
       if (isa<OMPCapturedExprDecl>(DRE->getDecl())) {
         Captures[DRE] = DRE;
       }
     }
     // Move on to the next nested for loop, or to the loop body.
     // OpenMP [2.8.1, simd construct, Restrictions]
     // All loops associated with the construct must be perfectly nested; that
     // is, there must be no intervening code nor any OpenMP directive between
     // any two loops.
     if (auto *For = dyn_cast<ForStmt>(CurStmt)) {
       CurStmt = For->getBody();
     } else {
       assert(isa<CXXForRangeStmt>(CurStmt) &&
              "Expected canonical for or range-based for loops.");
       CurStmt = cast<CXXForRangeStmt>(CurStmt)->getBody();
     }
-    CurStmt = OMPLoopDirective::tryToFindNextInnerLoop(
-        CurStmt, SemaRef.LangOpts.OpenMP >= 50);
   }
 
   Built.clear(/* size */ NestedLoopCount);
 
   if (SemaRef.CurContext->isDependentContext())
     return NestedLoopCount;
 
   // An example of what is generated for the following code:
checkOpenMPLoop(OpenMPDirectiveKind DKind, Expr *CollapseLoopCountExpr,
   // Calculate the last iteration number beforehand instead of doing this on
   // each iteration. Do not do this if the number of iterations may be kfold-ed.
   llvm::APSInt Result;
   bool IsConstant =
       LastIteration.get()->isIntegerConstantExpr(Result, SemaRef.Context);
   ExprResult CalcLastIteration;
   if (!IsConstant) {
     ExprResult SaveRef =
-        tryBuildCapture(SemaRef, LastIteration.get(), Captures);
+        tryBuildCapture(SemaRef, LastIteration.get(), Captures, Capturing);
     LastIteration = SaveRef;
 
     // Prepare SaveRef + 1.
     NumIterations = SemaRef.BuildBinOp(
         CurScope, SaveRef.get()->getExprLoc(), BO_Add, SaveRef.get(),
         SemaRef.ActOnIntegerConstant(SourceLocation(), 1).get());
     if (!NumIterations.isUsable())
       return 0;
   }
 
   SourceLocation InitLoc = IterSpaces[0].InitSrcRange.getBegin();
 
   // Build variables passed into runtime, necessary for worksharing directives.
   ExprResult LB, UB, IL, ST, EUB, CombLB, CombUB, PrevLB, PrevUB, CombEUB;
   if (isOpenMPWorksharingDirective(DKind) || isOpenMPTaskLoopDirective(DKind) ||
-      isOpenMPDistributeDirective(DKind)) {
+      isOpenMPDistributeDirective(DKind) ||
+      isOpenMPLoopTransformationDirective(DKind)) {
     // Lower bound variable, initialized with zero.
     VarDecl *LBDecl = buildVarDecl(SemaRef, InitLoc, VType, ".omp.lb");
     LB = buildDeclRefExpr(SemaRef, LBDecl, VType, InitLoc);
     SemaRef.AddInitializerToDecl(LBDecl,
                                  SemaRef.ActOnIntegerConstant(InitLoc, 0).get(),
                                  /*DirectInit*/ false);
 
     // Upper bound variable, initialized with last iteration number.
checkOpenMPLoop(OpenMPDirectiveKind DKind, Expr *CollapseLoopCountExpr,
   }
 
   // Build the iteration variable and its initialization before loop.
   ExprResult IV;
   ExprResult Init, CombInit;
   {
     VarDecl *IVDecl = buildVarDecl(SemaRef, InitLoc, RealVType, ".omp.iv");
     IV = buildDeclRefExpr(SemaRef, IVDecl, RealVType, InitLoc);
-    Expr *RHS =
-        (isOpenMPWorksharingDirective(DKind) ||
-         isOpenMPTaskLoopDirective(DKind) || isOpenMPDistributeDirective(DKind))
+    Expr *RHS = (isOpenMPWorksharingDirective(DKind) ||
+                 isOpenMPTaskLoopDirective(DKind) ||
+                 isOpenMPDistributeDirective(DKind) ||
+                 isOpenMPLoopTransformationDirective(DKind))
                     ? LB.get()
                     : SemaRef.ActOnIntegerConstant(SourceLocation(), 0).get();
     Init = SemaRef.BuildBinOp(CurScope, InitLoc, BO_Assign, IV.get(), RHS);
     Init = SemaRef.ActOnFinishFullExpr(Init.get(), /*DiscardedValue*/ false);
 
     if (isOpenMPLoopBoundSharingDirective(DKind)) {
       Expr *CombRHS =
           (isOpenMPWorksharingDirective(DKind) ||
            isOpenMPTaskLoopDirective(DKind) ||
            isOpenMPDistributeDirective(DKind))
BoundUB =
             .BuildBinOp(CurScope, CondLoc, BO_Add, BoundUB,
                         SemaRef.ActOnIntegerConstant(SourceLocation(), 1).get())
             .get();
     BoundUB =
         SemaRef.ActOnFinishFullExpr(BoundUB, /*DiscardedValue*/ false).get();
   }
   ExprResult Cond =
       (isOpenMPWorksharingDirective(DKind) ||
-       isOpenMPTaskLoopDirective(DKind) || isOpenMPDistributeDirective(DKind))
+       isOpenMPTaskLoopDirective(DKind) || isOpenMPDistributeDirective(DKind) ||
+       isOpenMPLoopTransformationDirective(DKind))
           ? SemaRef.BuildBinOp(CurScope, CondLoc,
                                UseStrictCompare ? BO_LT : BO_LE, IV.get(),
                                BoundUB)
           : SemaRef.BuildBinOp(CurScope, CondLoc, BO_LT, IV.get(),
                                NumIterations.get());
   ExprResult CombDistCond;
   if (isOpenMPLoopBoundSharingDirective(DKind)) {
     CombDistCond = SemaRef.BuildBinOp(CurScope, CondLoc, BO_LT, IV.get(),
if (!Inc.isUsable())
     return 0;
 
   // Increments for worksharing loops (LB = LB + ST; UB = UB + ST).
   // Used for directives with static scheduling.
   // In combined construct, add combined version that use CombLB and CombUB
   // base variables for the update
   ExprResult NextLB, NextUB, CombNextLB, CombNextUB;
   if (isOpenMPWorksharingDirective(DKind) || isOpenMPTaskLoopDirective(DKind) ||
-      isOpenMPDistributeDirective(DKind)) {
+      isOpenMPDistributeDirective(DKind) ||
+      isOpenMPLoopTransformationDirective(DKind)) {
     // LB + ST
     NextLB = SemaRef.BuildBinOp(CurScope, IncLoc, BO_Add, LB.get(), ST.get());
     if (!NextLB.isUsable())
       return 0;
     // LB = LB + ST
     NextLB =
         SemaRef.BuildBinOp(CurScope, IncLoc, BO_Assign, LB.get(), NextLB.get());
     NextLB =
checkOpenMPLoop(OpenMPDirectiveKind DKind, Expr *CollapseLoopCountExpr,
   bool HasErrors = false;
   Built.Counters.resize(NestedLoopCount);
   Built.Inits.resize(NestedLoopCount);
   Built.Updates.resize(NestedLoopCount);
   Built.Finals.resize(NestedLoopCount);
   Built.DependentCounters.resize(NestedLoopCount);
   Built.DependentInits.resize(NestedLoopCount);
   Built.FinalsConditions.resize(NestedLoopCount);
+  Built.Loops.resize(NestedLoopCount);
+  Built.Bodys.resize(NestedLoopCount);
   {
     // We implement the following algorithm for obtaining the
     // original loop iteration variable values based on the
     // value of the collapsed loop iteration variable IV.
     //
     // Let n+1 be the number of collapsed loops in the nest.
     // Iteration variables (I0, I1, .... In)
     // Iteration counts (N0, N1, ... Nn)
for (unsigned int Cnt = 0; Cnt < NestedLoopCount; ++Cnt) {
         Prod = Iter;
       Acc = SemaRef.BuildBinOp(CurScope, UpdLoc, BO_Sub,
                                Acc.get(), Prod.get());
 
       // Build update: IS.CounterVar(Private) = IS.Start + Iter * IS.Step
       auto *VD = cast<VarDecl>(cast<DeclRefExpr>(IS.CounterVar)->getDecl());
       DeclRefExpr *CounterVar = buildDeclRefExpr(
           SemaRef, VD, IS.CounterVar->getType(), IS.CounterVar->getExprLoc(),
-          /*RefersToCapture=*/true);
-      ExprResult Init =
-          buildCounterInit(SemaRef, CurScope, UpdLoc, CounterVar,
-                           IS.CounterInit, IS.IsNonRectangularLB, Captures);
+          /*RefersToCapture=*/Capturing);
+      ExprResult Init = buildCounterInit(SemaRef, CurScope, UpdLoc, CounterVar,
+                                         IS.CounterInit, IS.IsNonRectangularLB,
+                                         Captures, Capturing);
       if (!Init.isUsable()) {
         HasErrors = true;
         break;
       }
       ExprResult Update = buildCounterUpdate(
           SemaRef, CurScope, UpdLoc, CounterVar, IS.CounterInit, Iter,
-          IS.CounterStep, IS.Subtract, IS.IsNonRectangularLB, &Captures);
+          IS.CounterStep, IS.Subtract, IS.IsNonRectangularLB,
+          Capturing ? &Captures : nullptr);
       if (!Update.isUsable()) {
         HasErrors = true;
         break;
       }
 
       // Build final: IS.CounterVar = IS.Start + IS.NumIters * IS.Step
-      ExprResult Final =
-          buildCounterUpdate(SemaRef, CurScope, UpdLoc, CounterVar,
-                             IS.CounterInit, IS.NumIterations, IS.CounterStep,
-                             IS.Subtract, IS.IsNonRectangularLB, &Captures);
+      ExprResult Final = buildCounterUpdate(
+          SemaRef, CurScope, UpdLoc, CounterVar, IS.CounterInit,
+          IS.NumIterations, IS.CounterStep, IS.Subtract, IS.IsNonRectangularLB,
+          Capturing ? &Captures : nullptr);
       if (!Final.isUsable()) {
         HasErrors = true;
         break;
       }
 
       if (!Update.isUsable() || !Final.isUsable()) {
         HasErrors = true;
         break;
       }
       // Save results
+      Built.Loops[Cnt] = IS.Loop;
+      Built.Bodys[Cnt] = IS.Body;
       Built.Counters[Cnt] = IS.CounterVar;
       Built.PrivateCounters[Cnt] = IS.PrivateCounterVar;
       Built.Inits[Cnt] = Init.get();
       Built.Updates[Cnt] = Update.get();
       Built.Finals[Cnt] = Final.get();
       Built.DependentCounters[Cnt] = nullptr;
       Built.DependentInits[Cnt] = nullptr;
       Built.FinalsConditions[Cnt] = nullptr;
checkOpenMPLoop(OpenMPDirectiveKind DKind, Expr *CollapseLoopCountExpr,
   Built.DistCombinedFields.EUB = CombEUB.get();
   Built.DistCombinedFields.Init = CombInit.get();
   Built.DistCombinedFields.Cond = CombCond.get();
   Built.DistCombinedFields.NLB = CombNextLB.get();
   Built.DistCombinedFields.NUB = CombNextUB.get();
   Built.DistCombinedFields.DistCond = CombDistCond.get();
   Built.DistCombinedFields.ParForInDistCond = ParForInDistCond.get();
 
+  assert(Capturing || Captures.empty());
   return NestedLoopCount;
 }
 
 static Expr *getCollapseNumberExpr(ArrayRef<OMPClause *> Clauses) {
   auto CollapseClauses =
       OMPExecutableDirective::getClausesOfKind<OMPCollapseClause>(Clauses);
   if (CollapseClauses.begin() != CollapseClauses.end())
     return (*CollapseClauses.begin())->getNumForLoops();
StmtResult Sema::ActOnOpenMPTargetTeamsDistributeSimdDirective(
   if (checkSimdlenSafelenSpecified(*this, Clauses))
     return StmtError();
 
   setFunctionHasBranchProtectedScope();
   return OMPTargetTeamsDistributeSimdDirective::Create(
       Context, StartLoc, EndLoc, NestedLoopCount, Clauses, AStmt, B);
 }
 
+StmtResult
+Sema::ActOnOpenMPTileDirective(ArrayRef<OMPClause *> Clauses, Stmt *AStmt,
+                               SourceLocation StartLoc, SourceLocation EndLoc,
+                               VarsWithInheritedDSAType &VarsWithImplicitDSA) {
+  auto SizesClauses =
+      OMPExecutableDirective::getClausesOfKind<OMPSizesClause>(Clauses);
+  if (SizesClauses.begin() == SizesClauses.end()) {
+    // A missing 'sizes' clause is already reported by the parser.
+    return StmtError();
+  }
+  const OMPSizesClause *SizesClause = *SizesClauses.begin();
+  unsigned NumLoops = SizesClause->getNumSizes();
+
+  // Empty statement should only be possible if there already was an error.
+  if (!AStmt)
+    return StmtError();
+
+  // Verify and diagnose loop nest.
+  OMPLoopDirective::HelperExprs NestHelper;
+  unsigned NestedLoopCount = checkOpenMPLoop(
+      OMPD_tile, nullptr, nullptr, AStmt, *this, *DSAStack, VarsWithImplicitDSA,
+      NestHelper, NumLoops, /*EmitDiags=*/true);
+  if (NestedLoopCount != NumLoops)
+    return StmtError();
+
+  // Delay tiling to when template is completely instantiated.
+  if (CurContext->isDependentContext())
+    return OMPTileDirective::Create(Context, StartLoc, EndLoc, Clauses,
+                                    NumLoops, AStmt, nullptr);
+
+  // Derive per-loop logical iteration spaces.
+  SmallVector<OMPLoopDirective::HelperExprs, 4> LoopHelpers;
+  LoopHelpers.resize(NumLoops);
+  for (unsigned i = 0; i < NumLoops; ++i) {
+    Stmt *LoopStmt = NestHelper.Loops[i];
+    VarsWithInheritedDSAType TmpDSA;
+    unsigned SingleNumLoops =
+        checkOpenMPLoop(OMPD_tile, nullptr, nullptr, LoopStmt, *this, *DSAStack,
+                        TmpDSA, LoopHelpers[i], 1, /*EmitDiags=*/false);
+    (void)SingleNumLoops;
+    assert(SingleNumLoops == 1 && "Expect single loop iteration space");
+  }
+
+  // Collection of generated variable declaration.
+  SmallVector<Stmt *, 9> PreInits;
+
+  // Create iteration variables for the generated loops.
+  SmallVector<VarDecl *, 4> FloorIndVars;
+  SmallVector<VarDecl *, 4> TileIndVars;
+  SmallVector<Stmt *, 9> BodyParts;
+  FloorIndVars.resize(NumLoops);
+  TileIndVars.resize(NumLoops);
+  BodyParts.reserve(2 * NumLoops + 1);
+  for (unsigned i = 0; i < NumLoops; i += 1) {
+    OMPLoopDirective::HelperExprs &LoopHelper = LoopHelpers[i];
+    assert(LoopHelper.Counters.size() == 1 &&
+           "Expect single-dimensional loop iteration space");
+    auto *OrigCntVar = cast<DeclRefExpr>(LoopHelper.Counters[0]);
+    std::string OrigVarName = OrigCntVar->getNameInfo().getAsString();
+    DeclRefExpr *IterVarRef = cast<DeclRefExpr>(LoopHelper.IterationVarRef);
+    QualType CntTy = IterVarRef->getType();
+
+    // Iteration variable for the floor (i.e. outer) loop.
+    {
+      std::string FloorCntName =
+          (Twine(".floor_") + llvm::utostr(i) + ".iv." + OrigVarName).str();
+      VarDecl *FloorCntDecl =
+          buildVarDecl(*this, {}, CntTy, FloorCntName, nullptr, OrigCntVar);
+      FloorIndVars[i] = FloorCntDecl;
+
+      // Create a declaration statement for the new iteration variable.
+      Decl *D = FloorCntDecl;
+      auto *DeclS =
+          new (Context) DeclStmt(DeclGroupRef::Create(Context, &D, 1), {}, {});
+      PreInits.push_back(DeclS);
+    }
+
+    // Iteration variable for the tile (i.e. inner) loop.
+    {
+      std::string TileCntName =
+          (Twine(".tile_") + llvm::utostr(i) + ".iv." + OrigVarName).str();
+
+      // Reuse the iteration variable created by checkOpenMPLoop. It is also
+      // used by the expressions to derive the original iteration variable's
+      // value from the logical iteration number.
+      auto *TileCntDecl = cast<VarDecl>(IterVarRef->getDecl());
+      TileCntDecl->setDeclName(&PP.getIdentifierTable().get(TileCntName));
+      TileIndVars[i] = TileCntDecl;
+
+      // Create a declaration statement for the new iteration variable.
+      Decl *D = TileCntDecl;
+      auto DeclS =
+          new (Context) DeclStmt(DeclGroupRef::Create(Context, &D, 1), {}, {});
+      PreInits.push_back(DeclS);
+    }
+
+    // Statements to set the original iteration variable's value from the
+    // logical iteration number.
+    // FIXME: If the innermost body is an loop itself, inserting these
+    // statements stops it being recognized  as a perfectly nested loop (e.g.
+    // for applying tiling again). If this is the case, sink the expressions
+    // further into the inner loop.
+    {
+      // Declaration of the original loop iteration variable.
+      Decl *CounterDecl = OrigCntVar->getDecl();
+      auto *CounterDeclStmt = new (Context)
+          DeclStmt(DeclGroupRef::Create(Context, &CounterDecl, 1), {}, {});
+
+      // Update expression generated by checkOpenMPLoop.
+      Stmt *Upd = LoopHelper.Updates[0];
+
+      BodyParts.push_back(CounterDeclStmt);
+      BodyParts.push_back(Upd);
+    }
+  }
+
+  // Once the original iteration values are set, append the innermost body.
+  BodyParts.push_back(NestHelper.Bodys.back());
+  Stmt *Inner = CompoundStmt::Create(Context, BodyParts, AStmt->getBeginLoc(),
+                                     AStmt->getEndLoc());
+
+  // Create tile loops from the inside to the outside.
+  for (int i = NumLoops - 1; i >= 0; --i) {
+    OMPLoopDirective::HelperExprs &LoopHelper = LoopHelpers[i];
+    Expr *NumIterations = LoopHelper.NumIterations;
+    auto *OrigCntVar = cast<DeclRefExpr>(LoopHelper.Counters[0]);
+    QualType CntTy = OrigCntVar->getType();
+    Expr *DimTileSize = SizesClause->getSizesRefs()[i];
+    Scope *CurScope = getCurScope();
+
+    // Commonly used variables.
+    DeclRefExpr *TileIV = buildDeclRefExpr(*this, TileIndVars[i], CntTy, {});
+    DeclRefExpr *FloorIV = buildDeclRefExpr(*this, FloorIndVars[i], CntTy, {});
+
+    // For init-statement: .tile.iv = .floor.iv
+    ExprResult InitStmt = BuildBinOp(CurScope, {}, BO_Assign, TileIV, FloorIV);
+    if (!InitStmt.isUsable())
+      return StmtError();
+
+    // For cond-expression: .tile.iv < min(.floor.iv + DimTileSize,
+    // NumIterations)
+    ExprResult EndOfTile =
+        BuildBinOp(CurScope, {}, BO_Add, FloorIV, DimTileSize);
+    if (!EndOfTile.isUsable())
+      return StmtError();
+    ExprResult IsPartialTile =
+        BuildBinOp(CurScope, {}, BO_LT, NumIterations, EndOfTile.get());
+    if (!IsPartialTile.isUsable())
+      return StmtError();
+    ExprResult MinTileAndIterSpace = ActOnConditionalOp(
+        {}, {}, IsPartialTile.get(), NumIterations, EndOfTile.get());
+    if (!MinTileAndIterSpace.isUsable())
+      return StmtError();
+    ExprResult CondExpr =
+        BuildBinOp(CurScope, {}, BO_LT, TileIV, MinTileAndIterSpace.get());
+    if (!CondExpr.isUsable())
+      return StmtError();
+
+    // For incr-statement: ++.tile.iv
+    ExprResult IncrStmt = BuildUnaryOp(CurScope, {}, UO_PreInc, TileIV);
+    if (!IncrStmt.isUsable())
+      return StmtError();
+
+    Inner = new (Context) ForStmt(Context, InitStmt.get(), CondExpr.get(),
+                                  nullptr, IncrStmt.get(), Inner, {}, {}, {});
+  }
+
+  // Create floor loops from the inside to the outside.
+  for (int i = NumLoops - 1; i >= 0; --i) {
+    auto &LoopHelper = LoopHelpers[i];
+    Expr *NumIterations = LoopHelper.NumIterations;
+    DeclRefExpr *OrigCntVar = cast<DeclRefExpr>(LoopHelper.Counters[0]);
+    QualType CntTy = OrigCntVar->getType();
+    Expr *DimTileSize = SizesClause->getSizesRefs()[i];
+    Scope *CurScope = getCurScope();
+
+    // Commonly used variables.
+    DeclRefExpr *FloorIV = buildDeclRefExpr(*this, FloorIndVars[i], CntTy, {});
+
+    // For init-statement: .floor.iv = 0
+    auto *IVStart = IntegerLiteral::Create(
+        Context, llvm::APInt::getNullValue(Context.getTypeSize(CntTy)), CntTy,
+        {});
+    ExprResult InitStmt = BuildBinOp(CurScope, {}, BO_Assign, FloorIV, IVStart);
+    if (!InitStmt.isUsable())
+      return StmtError();
+
+    // For cond-expression: .floor.iv < NumIterations
+    ExprResult CondExpr =
+        BuildBinOp(CurScope, {}, BO_LT, FloorIV, NumIterations);
+    if (!CondExpr.isUsable())
+      return StmtError();
+
+    // For incr-statement: .floor.iv += DimTileSize
+    ExprResult IncrStmt =
+        BuildBinOp(CurScope, {}, BO_AddAssign, FloorIV, DimTileSize);
+    if (!IncrStmt.isUsable())
+      return StmtError();
+
+    Inner = new (Context) ForStmt(Context, InitStmt.get(), CondExpr.get(),
+                                  nullptr, IncrStmt.get(), Inner, {}, {}, {});
+  }
+
+  // Create the de-sugared tile loop nest including pre-inits.
+  PreInits.push_back(Inner);
+  auto *TransformedStmt = CompoundStmt::Create(Context, PreInits, {}, {});
+
+  return OMPTileDirective::Create(Context, StartLoc, EndLoc, Clauses, NumLoops,
+                                  AStmt, TransformedStmt);
+}
+
 OMPClause *Sema::ActOnOpenMPSingleExprClause(OpenMPClauseKind Kind, Expr *Expr,
                                              SourceLocation StartLoc,
                                              SourceLocation LParenLoc,
                                              SourceLocation EndLoc) {
   OMPClause *Res = nullptr;
   switch (Kind) {
   case OMPC_final:
     Res = ActOnOpenMPFinalClause(Expr, StartLoc, LParenLoc, EndLoc);
OMPClause *Sema::ActOnOpenMPSingleExprClause(OpenMPClauseKind Kind, Expr *Expr,
   case OMPC_atomic_default_mem_order:
   case OMPC_device_type:
   case OMPC_match:
   case OMPC_nontemporal:
   case OMPC_order:
   case OMPC_destroy:
   case OMPC_inclusive:
   case OMPC_exclusive:
+  case OMPC_sizes:
     llvm_unreachable("Clause is not allowed.");
   }
   return Res;
 }
 
 // An OpenMP directive such as 'target parallel' has two captured regions:
 // for the 'target' and 'parallel' respectively.  This function returns
 // the region in which to capture expressions associated with a clause.
case OMPC_if:
     case OMPD_master:
     case OMPD_critical:
     case OMPD_taskgroup:
     case OMPD_distribute:
     case OMPD_ordered:
     case OMPD_atomic:
     case OMPD_teams_distribute:
     case OMPD_requires:
+    case OMPD_tile:
       llvm_unreachable("Unexpected OpenMP directive with if-clause");
     case OMPD_unknown:
       llvm_unreachable("Unknown OpenMP directive");
     }
     break;
   case OMPC_num_threads:
     switch (DKind) {
     case OMPD_target_parallel:
case OMPC_num_threads:
     case OMPD_taskgroup:
     case OMPD_distribute:
     case OMPD_ordered:
     case OMPD_atomic:
     case OMPD_distribute_simd:
     case OMPD_teams_distribute:
     case OMPD_teams_distribute_simd:
     case OMPD_requires:
+    case OMPD_tile:
       llvm_unreachable("Unexpected OpenMP directive with num_threads-clause");
     case OMPD_unknown:
       llvm_unreachable("Unknown OpenMP directive");
     }
     break;
   case OMPC_num_teams:
     switch (DKind) {
     case OMPD_target_teams:
case OMPC_num_teams:
     case OMPD_master:
     case OMPD_critical:
     case OMPD_taskgroup:
     case OMPD_distribute:
     case OMPD_ordered:
     case OMPD_atomic:
     case OMPD_distribute_simd:
     case OMPD_requires:
+    case OMPD_tile:
       llvm_unreachable("Unexpected OpenMP directive with num_teams-clause");
     case OMPD_unknown:
       llvm_unreachable("Unknown OpenMP directive");
     }
     break;
   case OMPC_thread_limit:
     switch (DKind) {
     case OMPD_target_teams:
case OMPC_thread_limit:
     case OMPD_master:
     case OMPD_critical:
     case OMPD_taskgroup:
     case OMPD_distribute:
     case OMPD_ordered:
     case OMPD_atomic:
     case OMPD_distribute_simd:
     case OMPD_requires:
+    case OMPD_tile:
       llvm_unreachable("Unexpected OpenMP directive with thread_limit-clause");
     case OMPD_unknown:
       llvm_unreachable("Unknown OpenMP directive");
     }
     break;
   case OMPC_schedule:
     switch (DKind) {
     case OMPD_parallel_for:
case OMPC_schedule:
     case OMPD_critical:
     case OMPD_taskgroup:
     case OMPD_distribute:
     case OMPD_ordered:
     case OMPD_atomic:
     case OMPD_distribute_simd:
     case OMPD_target_teams:
     case OMPD_requires:
+    case OMPD_tile:
       llvm_unreachable("Unexpected OpenMP directive with schedule clause");
     case OMPD_unknown:
       llvm_unreachable("Unknown OpenMP directive");
     }
     break;
   case OMPC_dist_schedule:
     switch (DKind) {
     case OMPD_teams_distribute_parallel_for:
case OMPC_dist_schedule:
     case OMPD_single:
     case OMPD_master:
     case OMPD_critical:
     case OMPD_taskgroup:
     case OMPD_ordered:
     case OMPD_atomic:
     case OMPD_target_teams:
     case OMPD_requires:
+    case OMPD_tile:
       llvm_unreachable("Unexpected OpenMP directive with schedule clause");
     case OMPD_unknown:
       llvm_unreachable("Unknown OpenMP directive");
     }
     break;
   case OMPC_device:
     switch (DKind) {
     case OMPD_target_update:
case OMPC_device:
     case OMPD_master:
     case OMPD_critical:
     case OMPD_taskgroup:
     case OMPD_distribute:
     case OMPD_ordered:
     case OMPD_atomic:
     case OMPD_distribute_simd:
     case OMPD_requires:
+    case OMPD_tile:
       llvm_unreachable("Unexpected OpenMP directive with num_teams-clause");
     case OMPD_unknown:
       llvm_unreachable("Unknown OpenMP directive");
     }
     break;
   case OMPC_grainsize:
   case OMPC_num_tasks:
   case OMPC_final:
case OMPC_priority:
     case OMPD_master:
     case OMPD_critical:
     case OMPD_taskgroup:
     case OMPD_distribute:
     case OMPD_ordered:
     case OMPD_atomic:
     case OMPD_distribute_simd:
     case OMPD_requires:
+    case OMPD_tile:
       llvm_unreachable("Unexpected OpenMP directive with grainsize-clause");
     case OMPD_unknown:
       llvm_unreachable("Unknown OpenMP directive");
     }
     break;
   case OMPC_firstprivate:
   case OMPC_lastprivate:
   case OMPC_reduction:
static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
   case OMPC_device_type:
   case OMPC_match:
   case OMPC_nontemporal:
   case OMPC_order:
   case OMPC_destroy:
   case OMPC_detach:
   case OMPC_inclusive:
   case OMPC_exclusive:
+  case OMPC_sizes:
     llvm_unreachable("Unexpected OpenMP clause.");
   }
   return CaptureRegion;
 }
 
 OMPClause *Sema::ActOnOpenMPIfClause(OpenMPDirectiveKind NameModifier,
                                      Expr *Condition, SourceLocation StartLoc,
                                      SourceLocation LParenLoc,
OMPClause *Sema::ActOnOpenMPSimpleClause(
   case OMPC_dynamic_allocators:
   case OMPC_device_type:
   case OMPC_match:
   case OMPC_nontemporal:
   case OMPC_destroy:
   case OMPC_detach:
   case OMPC_inclusive:
   case OMPC_exclusive:
+  case OMPC_sizes:
     llvm_unreachable("Clause is not allowed.");
   }
   return Res;
 }
 
 static std::string
 getListOfPossibleValues(OpenMPClauseKind K, unsigned First, unsigned Last,
                         ArrayRef<unsigned> Exclude = llvm::None) {
Diag(KindKwLoc, diag::err_omp_unexpected_clause_value)
                                    /*Last=*/OMPC_DEPEND_unknown, Except)
         << getOpenMPClauseName(OMPC_update);
     return nullptr;
   }
   return OMPUpdateClause::Create(Context, StartLoc, LParenLoc, KindKwLoc, Kind,
                                  EndLoc);
 }
 
+OMPClause *Sema::ActOnOpenMPSizesClause(ArrayRef<Expr *> SizeExprs,
+                                        SourceLocation StartLoc,
+                                        SourceLocation LParenLoc,
+                                        SourceLocation EndLoc) {
+  for (auto SizeExpr : SizeExprs) {
+    ExprResult NumForLoopsResult =
+        VerifyPositiveIntegerConstantInClause(SizeExpr, OMPC_sizes, true);
+    if (!NumForLoopsResult.isUsable())
+      return nullptr;
+  }
+
+  DSAStack->setAssociatedLoops(SizeExprs.size());
+  return OMPSizesClause::Create(Context, StartLoc, LParenLoc, EndLoc,
+                                SizeExprs);
+}
+
 OMPClause *Sema::ActOnOpenMPSingleExprWithArgClause(
     OpenMPClauseKind Kind, ArrayRef<unsigned> Argument, Expr *Expr,
     SourceLocation StartLoc, SourceLocation LParenLoc,
     ArrayRef<SourceLocation> ArgumentLoc, SourceLocation DelimLoc,
     SourceLocation EndLoc) {
   OMPClause *Res = nullptr;
   switch (Kind) {
   case OMPC_schedule:
OMPClause *Sema::ActOnOpenMPSingleExprWithArgClause(
   case OMPC_device_type:
   case OMPC_match:
   case OMPC_nontemporal:
   case OMPC_order:
   case OMPC_destroy:
   case OMPC_detach:
   case OMPC_inclusive:
   case OMPC_exclusive:
+  case OMPC_sizes:
     llvm_unreachable("Clause is not allowed.");
   }
   return Res;
 }
 
 static bool checkScheduleModifiers(Sema &S, OpenMPScheduleClauseModifier M1,
                                    OpenMPScheduleClauseModifier M2,
                                    SourceLocation M1Loc, SourceLocation M2Loc) {
OMPClause *Sema::ActOnOpenMPClause(OpenMPClauseKind Kind,
   case OMPC_atomic_default_mem_order:
   case OMPC_device_type:
   case OMPC_match:
   case OMPC_nontemporal:
   case OMPC_order:
   case OMPC_detach:
   case OMPC_inclusive:
   case OMPC_exclusive:
+  case OMPC_sizes:
     llvm_unreachable("Clause is not allowed.");
   }
   return Res;
 }
 
 OMPClause *Sema::ActOnOpenMPNowaitClause(SourceLocation StartLoc,
                                          SourceLocation EndLoc) {
   DSAStack->setNowaitRegion();
OMPClause *Sema::ActOnOpenMPVarListClause(
   case OMPC_reverse_offload:
   case OMPC_dynamic_allocators:
   case OMPC_atomic_default_mem_order:
   case OMPC_device_type:
   case OMPC_match:
   case OMPC_order:
   case OMPC_destroy:
   case OMPC_detach:
+  case OMPC_sizes:
     llvm_unreachable("Clause is not allowed.");
   }
   return Res;
 }
 
 ExprResult Sema::getOpenMPCapturedExpr(VarDecl *Capture, ExprValueKind VK,
                                        ExprObjectKind OK, SourceLocation Loc) {
   ExprResult Res = BuildDeclRefExpr(

clang/lib/Sema/TreeTransform.h

OMPClause *RebuildOMPOrderClause(OpenMPOrderClauseKind Kind,
                                    SourceLocation KindKwLoc,
                                    SourceLocation StartLoc,
                                    SourceLocation LParenLoc,
                                    SourceLocation EndLoc) {
     return getSema().ActOnOpenMPOrderClause(Kind, KindKwLoc, StartLoc,
                                             LParenLoc, EndLoc);
   }
 
+  OMPClause *RebuildOMPSizesClause(ArrayRef<Expr *> Sizes,
+                                   SourceLocation StartLoc,
+                                   SourceLocation LParenLoc,
+                                   SourceLocation EndLoc) {
+    return getSema().ActOnOpenMPSizesClause(Sizes, StartLoc, LParenLoc, EndLoc);
+  }
+
   /// Rebuild the operand to an Objective-C \@synchronized statement.
   ///
   /// By default, performs semantic analysis to build the new statement.
   /// Subclasses may override this routine to provide different behavior.
   ExprResult RebuildObjCAtSynchronizedOperand(SourceLocation atLoc,
                                               Expr *object) {
     return getSema().ActOnObjCAtSynchronizedOperand(atLoc, object);
   }
StmtResult TreeTransform<Derived>::TransformOMPExecutableDirective(
   }
   StmtResult AssociatedStmt;
   if (D->hasAssociatedStmt() && D->getAssociatedStmt()) {
     getDerived().getSema().ActOnOpenMPRegionStart(D->getDirectiveKind(),
                                                   /*CurScope=*/nullptr);
     StmtResult Body;
     {
       Sema::CompoundScopeRAII CompoundScope(getSema());
-      Stmt *CS = D->getInnermostCapturedStmt()->getCapturedStmt();
+      Stmt *CS = D->ignoreCaptures();
       Body = getDerived().TransformStmt(CS);
     }
     AssociatedStmt =
         getDerived().getSema().ActOnOpenMPRegionEnd(Body, TClauses);
     if (AssociatedStmt.isInvalid()) {
       return StmtError();
     }
   }
TreeTransform<Derived>::TransformOMPTargetTeamsDistributeSimdDirective(
   DeclarationNameInfo DirName;
   getDerived().getSema().StartOpenMPDSABlock(
       OMPD_target_teams_distribute_simd, DirName, nullptr, D->getBeginLoc());
   auto Res = getDerived().TransformOMPExecutableDirective(D);
   getDerived().getSema().EndOpenMPDSABlock(Res.get());
   return Res;
 }
 
+template <typename Derived>
+StmtResult
+TreeTransform<Derived>::TransformOMPTileDirective(OMPTileDirective *D) {
+  DeclarationNameInfo DirName;
+  getDerived().getSema().StartOpenMPDSABlock(D->getDirectiveKind(), DirName,
+                                             nullptr, D->getBeginLoc());
+  StmtResult Res = getDerived().TransformOMPExecutableDirective(D);
+  getDerived().getSema().EndOpenMPDSABlock(Res.get());
+  return Res;
+}
 
 //===----------------------------------------------------------------------===//
 // OpenMP clause transformation
 //===----------------------------------------------------------------------===//
 template <typename Derived>
 OMPClause *TreeTransform<Derived>::TransformOMPIfClause(OMPIfClause *C) {
   ExprResult Cond = getDerived().TransformExpr(C->getCondition());
   if (Cond.isInvalid())

 template <typename Derived>
 OMPClause *
 TreeTransform<Derived>::TransformOMPOrderClause(OMPOrderClause *C) {
   return getDerived().RebuildOMPOrderClause(C->getKind(), C->getKindKwLoc(),
                                             C->getBeginLoc(), C->getLParenLoc(),
                                             C->getEndLoc());
 }
 
+template <typename Derived>
+OMPClause *TreeTransform<Derived>::TransformOMPSizesClause(OMPSizesClause *C) {
+  SmallVector<Expr *, 4> TransformedSizes;
+  TransformedSizes.reserve(C->getNumSizes());
+  bool Changed = false;
+  for (auto E : C->getSizesRefs()) {
+    if (!E) {
+      TransformedSizes.push_back(nullptr);
+      continue;
+    }
+
+    auto T = getDerived().TransformExpr(E);
+    if (T.isInvalid())
+      return nullptr;
+    if (E != T.get())
+      Changed = true;
+    TransformedSizes.push_back(T.get());
+  }
+
+  if (!Changed && !getDerived().AlwaysRebuild())
+    return C;
+  return RebuildOMPSizesClause(TransformedSizes, C->getBeginLoc(),
+                               C->getLParenLoc(), C->getEndLoc());
+}
+
 //===----------------------------------------------------------------------===//
 // Expression transformation
 //===----------------------------------------------------------------------===//
 template<typename Derived>
 ExprResult
 TreeTransform<Derived>::TransformConstantExpr(ConstantExpr *E) {
   return TransformExpr(E->getSubExpr());
 }

clang/lib/Serialization/ASTReader.cpp

case OMPC_reverse_offload:
     C = new (Context) OMPReverseOffloadClause();
     break;
   case OMPC_dynamic_allocators:
     C = new (Context) OMPDynamicAllocatorsClause();
     break;
   case OMPC_atomic_default_mem_order:
     C = new (Context) OMPAtomicDefaultMemOrderClause();
     break;
   case OMPC_private:
     C = OMPPrivateClause::CreateEmpty(Context, Record.readInt());
     break;
   case OMPC_firstprivate:
     C = OMPFirstprivateClause::CreateEmpty(Context, Record.readInt());
     break;
   case OMPC_lastprivate:
     C = OMPLastprivateClause::CreateEmpty(Context, Record.readInt());
     break;
case OMPC_order:
     C = new (Context) OMPOrderClause();
     break;
   case OMPC_destroy:
     C = new (Context) OMPDestroyClause();
     break;
   case OMPC_detach:
     C = new (Context) OMPDetachClause();
     break;
+  case OMPC_sizes: {
+    unsigned NumSizes = Record.readInt();
+    C = OMPSizesClause::CreateEmpty(Context, NumSizes);
+    break;
+  }
   }
   assert(C && "Unknown OMPClause type");
 
   Visit(C);
   C->setLocStart(Record.readSourceLocation());
   C->setLocEnd(Record.readSourceLocation());
 
   return C;
void OMPClauseReader::VisitOMPAllocatorClause(OMPAllocatorClause *C) {
   C->setLParenLoc(Record.readSourceLocation());
 }
 
 void OMPClauseReader::VisitOMPCollapseClause(OMPCollapseClause *C) {
   C->setNumForLoops(Record.readSubExpr());
   C->setLParenLoc(Record.readSourceLocation());
 }
 
+void OMPClauseReader::VisitOMPSizesClause(OMPSizesClause *C) {
+  for (auto &E : C->getSizesRefs())
+    E = Record.readSubExpr();
+  C->setLParenLoc(Record.readSourceLocation());
+}
+
 void OMPClauseReader::VisitOMPDefaultClause(OMPDefaultClause *C) {
   C->setDefaultKind(static_cast<llvm::omp::DefaultKind>(Record.readInt()));
   C->setLParenLoc(Record.readSourceLocation());
   C->setDefaultKindKwLoc(Record.readSourceLocation());
 }
 
 void OMPClauseReader::VisitOMPProcBindClause(OMPProcBindClause *C) {
   C->setProcBindKind(static_cast<llvm::omp::ProcBindKind>(Record.readInt()));

clang/lib/Serialization/ASTReaderStmt.cpp

if (E->hasAssociatedStmt())
     E->setAssociatedStmt(Record.readSubStmt());
 }
 
 void ASTStmtReader::VisitOMPLoopDirective(OMPLoopDirective *D) {
   VisitStmt(D);
   // Two fields (NumClauses and CollapsedNum) were read in ReadStmtFromStream.
   Record.skipInts(2);
   VisitOMPExecutableDirective(D);
+
+  if (isOpenMPLoopTransformationDirective(D->getDirectiveKind()))
+    return;
+
   D->setIterationVariable(Record.readSubExpr());
   D->setLastIteration(Record.readSubExpr());
   D->setCalcLastIteration(Record.readSubExpr());
   D->setPreCond(Record.readSubExpr());
   D->setCond(Record.readSubExpr());
   D->setInit(Record.readSubExpr());
   D->setInc(Record.readSubExpr());
   D->setPreInits(Record.readSubStmt());
void ASTStmtReader::VisitOMPTargetTeamsDistributeParallelForSimdDirective(
   VisitOMPLoopDirective(D);
 }
 
 void ASTStmtReader::VisitOMPTargetTeamsDistributeSimdDirective(
     OMPTargetTeamsDistributeSimdDirective *D) {
   VisitOMPLoopDirective(D);
 }
 
+void ASTStmtReader::VisitOMPTileDirective(OMPTileDirective *D) {
+  VisitOMPLoopDirective(D);
+  D->setTransformedStmt(Record.readStmt());
+}
+
 //===----------------------------------------------------------------------===//
 // ASTReader Implementation
 //===----------------------------------------------------------------------===//
 
 Stmt *ASTReader::ReadStmt(ModuleFile &F) {
   switch (ReadingKind) {
   case Read_None:
     llvm_unreachable("should not call this when not reading anything");
while (true) {
     case STMT_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD_DIRECTIVE: {
       auto NumClauses = Record[ASTStmtReader::NumStmtFields];
       auto CollapsedNum = Record[ASTStmtReader::NumStmtFields + 1];
       S = OMPTargetTeamsDistributeSimdDirective::CreateEmpty(
           Context, NumClauses, CollapsedNum, Empty);
       break;
     }
 
+    case STMT_OMP_TILE_DIRECTIVE: {
+      unsigned NumClauses = Record[ASTStmtReader::NumStmtFields];
+      unsigned NumLoops = Record[ASTStmtReader::NumStmtFields + 1];
+      S = OMPTileDirective::CreateEmpty(Context, NumClauses, NumLoops);
+      break;
+    }
+
     case EXPR_CXX_OPERATOR_CALL:
       S = CXXOperatorCallExpr::CreateEmpty(
           Context, /*NumArgs=*/Record[ASTStmtReader::NumExprFields], Empty);
       break;
 
     case EXPR_CXX_MEMBER_CALL:
       S = CXXMemberCallExpr::CreateEmpty(
           Context, /*NumArgs=*/Record[ASTStmtReader::NumExprFields], Empty);

clang/lib/Serialization/ASTWriter.cpp

 }
 
 void OMPClauseWriter::VisitOMPOrderClause(OMPOrderClause *C) {
   Record.writeEnum(C->getKind());
   Record.AddSourceLocation(C->getLParenLoc());
   Record.AddSourceLocation(C->getKindKwLoc());
 }
 
+void OMPClauseWriter::VisitOMPSizesClause(OMPSizesClause *C) {
+  Record.push_back(C->getNumSizes());
+  for (auto Size : C->getSizesRefs())
+    Record.AddStmt(Size);
+  Record.AddSourceLocation(C->getLParenLoc());
+}
+
 void ASTRecordWriter::writeOMPTraitInfo(const OMPTraitInfo *TI) {
   writeUInt32(TI->Sets.size());
   for (const auto &Set : TI->Sets) {
     writeEnum(Set.Kind);
     writeUInt32(Set.Selectors.size());
     for (const auto &Selector : Set.Selectors) {
       writeEnum(Selector.Kind);
       writeBool(Selector.ScoreOrCondition);
       if (Selector.ScoreOrCondition)
         writeExprRef(Selector.ScoreOrCondition);
       writeUInt32(Selector.Properties.size());
       for (const auto &Property : Selector.Properties)
         writeEnum(Property.Kind);
     }
   }
 }

clang/lib/Serialization/ASTWriterStmt.cpp

if (E->hasAssociatedStmt())
     Record.AddStmt(E->getAssociatedStmt());
 }
 
 void ASTStmtWriter::VisitOMPLoopDirective(OMPLoopDirective *D) {
   VisitStmt(D);
   Record.push_back(D->getNumClauses());
   Record.push_back(D->getCollapsedNumber());
   VisitOMPExecutableDirective(D);
+
+  if (isOpenMPLoopTransformationDirective(D->getDirectiveKind()))
+    return;
+
   Record.AddStmt(D->getIterationVariable());
   Record.AddStmt(D->getLastIteration());
   Record.AddStmt(D->getCalcLastIteration());
   Record.AddStmt(D->getPreCond());
   Record.AddStmt(D->getCond());
   Record.AddStmt(D->getInit());
   Record.AddStmt(D->getInc());
   Record.AddStmt(D->getPreInits());

 }
 
 void ASTStmtWriter::VisitOMPTargetTeamsDistributeSimdDirective(
     OMPTargetTeamsDistributeSimdDirective *D) {
   VisitOMPLoopDirective(D);
   Code = serialization::STMT_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD_DIRECTIVE;
 }
 
+void ASTStmtWriter::VisitOMPTileDirective(OMPTileDirective *D) {
+  VisitOMPLoopDirective(D);
+  Record.AddStmt(D->getTransformedStmt());
+  Code = serialization::STMT_OMP_TILE_DIRECTIVE;
+}
+
 //===----------------------------------------------------------------------===//
 // ASTWriter Implementation
 //===----------------------------------------------------------------------===//
 
 unsigned ASTWriter::RecordSwitchCaseID(SwitchCase *S) {
   assert(SwitchCaseIDs.find(S) == SwitchCaseIDs.end() &&
          "SwitchCase recorded twice");
   unsigned NextID = SwitchCaseIDs.size();

clang/test/Index/openmp-tile.c

+// RUN: c-index-test -test-load-source local %s -fopenmp | FileCheck %s
+
+void test() {
+#pragma omp tile sizes(5)
+  for (int i = 0; i < 65; i += 1)
+    ;
+}
+
+// CHECK: openmp-tile.c:4:1: OMPTileDirective= Extent=[4:1 - 4:26]
+// CHECK: openmp-tile.c:4:24: IntegerLiteral= Extent=[4:24 - 4:25]
+// CHECK: openmp-tile.c:5:3: ForStmt= Extent=[5:3 - 6:6]
+

clang/test/OpenMP/tile_ast_print.cpp

+// Check no warnings/errors
+// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -fopenmp -fopenmp-version=51 -fsyntax-only -verify %s
+// expected-no-diagnostics
+
+// Check AST and unparsing
+// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -fopenmp -fopenmp-version=51 -ast-dump  %s | FileCheck %s --check-prefix=DUMP
+// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -fopenmp -fopenmp-version=51 -ast-print %s | FileCheck %s --check-prefix=PRINT
+
+// Check same results after serialization round-trip
+// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -fopenmp -fopenmp-version=51 -emit-pch -o %t %s
+// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -fopenmp -fopenmp-version=51 -include-pch %t -ast-dump-all %s | FileCheck %s --check-prefix=DUMP
+// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -fopenmp -fopenmp-version=51 -include-pch %t -ast-print    %s | FileCheck %s --check-prefix=PRINT
+
+#ifndef HEADER
+#define HEADER
+
+// placeholder for loop body code.
+extern "C" void body(...);
+
+
+// PRINT-LABEL: void foo1(
+// DUMP-LABEL:  FunctionDecl {{.*}} foo1
+void foo1() {
+  // PRINT:     #pragma omp tile sizes(5, 5)
+  // DUMP:      OMPTileDirective
+  // DUMP-NEXT:   OMPSizesClause
+  // DUMP-NEXT:     IntegerLiteral {{.*}} 5
+  // DUMP-NEXT:     IntegerLiteral {{.*}} 5
+  #pragma omp tile sizes(5,5)
+  // PRINT: for (int i = 7; i < 17; i += 3)
+  // DUMP-NEXT: ForStmt
+  for (int i = 7; i < 17; i += 3)
+    // PRINT: for (int j = 7; j < 17; j += 3)
+    // DUMP:  ForStmt
+    for (int j = 7; j < 17; j += 3)
+    // PRINT: body(i, j);
+    // DUMP:  CallExpr
+      body(i, j);
+}
+
+
+// PRINT-LABEL: void foo2(
+// DUMP-LABEL:  FunctionDecl {{.*}} foo2
+void foo2(int start1, int start2, int end1, int end2) {
+  // PRINT:     #pragma omp tile sizes(5, 5)
+  // DUMP:      OMPTileDirective
+  // DUMP-NEXT:   OMPSizesClause
+  // DUMP-NEXT:     IntegerLiteral {{.*}} 5
+  // DUMP-NEXT:     IntegerLiteral {{.*}} 5
+  #pragma omp tile sizes(5,5)
+  // PRINT: for (int i = start1; i < end1; i += 1)
+  // DUMP-NEXT: ForStmt
+  for (int i = start1; i < end1; i += 1)
+    // PRINT: for (int j = start2; j < end2; j += 1)
+    // DUMP:  ForStmt
+    for (int j = start2; j < end2; j += 1)
+      // PRINT: body(i, j);
+      // DUMP:  CallExpr
+      body(i, j);
+}
+
+
+// PRINT-LABEL: void foo3(
+// DUMP-LABEL:  FunctionDecl {{.*}} foo3
+void foo3() {
+  // PRINT: #pragma omp for
+  // DUMP:  OMPForDirective
+  // DUMP-NEXT:    CapturedStmt
+  // DUMP-NEXT:      CapturedDecl
+  #pragma omp for
+  // PRINT:     #pragma omp tile sizes(5)
+  // DUMP-NEXT:      OMPTileDirective
+  // DUMP-NEXT:   OMPSizesClause
+  // DUMP-NEXT:     IntegerLiteral {{.*}} 5
+  #pragma omp tile sizes(5)
+  for (int i = 7; i < 17; i += 3)
+    // PRINT: body(i);
+    // DUMP:  CallExpr
+    body(i);
+}
+
+
+// PRINT-LABEL: void foo4(
+// DUMP-LABEL:  FunctionDecl {{.*}} foo4
+void foo4() {
+  // PRINT: #pragma omp for collapse(3)
+  // DUMP: OMPForDirective
+  // DUMP-NEXT: OMPCollapseClause
+  // DUMP-NEXT:  ConstantExpr
+  // DUMP-NEXT:  IntegerLiteral {{.*}} 3
+  // DUMP-NEXT:    CapturedStmt
+  // DUMP-NEXT:      CapturedDecl
+  #pragma omp for collapse(3)
+  // PRINT:     #pragma omp tile sizes(5, 5)
+  // DUMP:      OMPTileDirective
+  // DUMP-NEXT:   OMPSizesClause
+  // DUMP-NEXT:     IntegerLiteral {{.*}} 5
+  // DUMP-NEXT:     IntegerLiteral {{.*}} 5
+  #pragma omp tile sizes(5, 5)
+  // PRINT: for (int i = 7; i < 17; i += 1)
+  // DUMP-NEXT: ForStmt
+  for (int i = 7; i < 17; i += 1)
+    // PRINT: for (int j = 7; j < 17; j += 1)
+    // DUMP:  ForStmt
+    for (int j = 7; j < 17; j += 1)
+      // PRINT: body(i, j);
+      // DUMP:  CallExpr
+      body(i, j);
+}
+
+
+// PRINT-LABEL: void foo5(
+// DUMP-LABEL:  FunctionDecl {{.*}} foo5
+void foo5(int start, int end, int step) {
+  // PRINT: #pragma omp for collapse(2)
+  // DUMP:      OMPForDirective
+  // DUMP-NEXT:   OMPCollapseClause
+  // DUMP-NEXT:    ConstantExpr
+  // DUMP-NEXT:    IntegerLiteral {{.*}} 2
+  // DUMP-NEXT:  CapturedStmt
+  // DUMP-NEXT:    CapturedDecl
+  #pragma omp for collapse(2)
+  // PRINT: for (int i = 7; i < 17; i += 1)
+  // DUMP-NEXT: ForStmt
+  for (int i = 7; i < 17; i += 1)
+    // PRINT:     #pragma omp tile sizes(5)
+    // DUMP:      OMPTileDirective
+    // DUMP-NEXT:   OMPSizesClause
+    // DUMP-NEXT:     IntegerLiteral {{.*}} 5
+    #pragma omp tile sizes(5)
+    // PRINT: for (int j = 7; j < 17; j += 1)
+    // DUMP-NEXT: ForStmt
+    for (int j = 7; j < 17; j += 1)
+      // PRINT: body(i, j);
+      // DUMP:  CallExpr
+      body(i, j);
+}
+
+
+// PRINT-LABEL: void foo6(
+// DUMP-LABEL: FunctionTemplateDecl {{.*}} foo6
+template<typename T, T Step, T Tile>
+void foo6(T start, T end) {
+  // PRINT: #pragma omp tile sizes(Tile)
+  // DUMP:      OMPTileDirective
+  // DUMP-NEXT:   OMPSizesClause
+  // DUMP-NEXT:     DeclRefExpr {{.*}} 'Tile' 'T'
+  #pragma omp tile sizes(Tile)
+    // PRINT-NEXT:  for (T i = start; i < end; i += Step)
+    // DUMP-NEXT: ForStmt
+    for (T i = start; i < end; i += Step)
+      // PRINT-NEXT: body(i);
+      // DUMP:  CallExpr
+      body(i);
+}
+
+// Also test instantiating the template.
+void tfoo6() {
+  foo6<int,3,5>(0, 42);
+}
+
+
+#endif
+

clang/test/OpenMP/tile_codegen.cpp

+// Check no warnings/errors
+// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -fopenmp -fopenmp-version=51 -fsyntax-only -verify %s
+// expected-no-diagnostics
+
+// Check code generation
+// RUN: %clang_cc1 -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 -triple x86_64-pc-linux-gnu -fopenmp -fopenmp-version=51 -emit-pch -o %t %s
+// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -fopenmp -fopenmp-version=51 -include-pch %t -emit-llvm %s -o - | FileCheck %s --check-prefix=IR
+
+
+#ifndef HEADER
+#define HEADER
+
+// placeholder for loop body code.
+extern "C" void body(...) {}
+
+
+// IR-LABEL: @foo1(
+// 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:    [[DOTFLOOR_0_IV_I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTTILE_0_IV_I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[I:%.*]] = alloca i32, align 4
+// 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:    store i32 0, i32* [[DOTFLOOR_0_IV_I]], align 4
+// IR-NEXT:    br label [[FOR_COND:%.*]]
+// IR:       for.cond:
+// IR-NEXT:    [[TMP0:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I]], align 4
+// IR-NEXT:    [[TMP1:%.*]] = load i32, i32* [[END_ADDR]], align 4
+// IR-NEXT:    [[TMP2:%.*]] = load i32, i32* [[START_ADDR]], align 4
+// IR-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP1]], [[TMP2]]
+// IR-NEXT:    [[SUB1:%.*]] = sub nsw i32 [[SUB]], 1
+// IR-NEXT:    [[TMP3:%.*]] = load i32, i32* [[STEP_ADDR]], align 4
+// IR-NEXT:    [[ADD:%.*]] = add nsw i32 [[SUB1]], [[TMP3]]
+// IR-NEXT:    [[TMP4:%.*]] = load i32, i32* [[STEP_ADDR]], align 4
+// IR-NEXT:    [[DIV:%.*]] = sdiv i32 [[ADD]], [[TMP4]]
+// IR-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1
+// IR-NEXT:    [[ADD3:%.*]] = add nsw i32 [[SUB2]], 1
+// IR-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP0]], [[ADD3]]
+// IR-NEXT:    br i1 [[CMP]], label [[FOR_BODY:%.*]], label [[FOR_END25:%.*]]
+// IR:       for.body:
+// IR-NEXT:    [[TMP5:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I]], align 4
+// IR-NEXT:    store i32 [[TMP5]], i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    br label [[FOR_COND4:%.*]]
+// IR:       for.cond4:
+// IR-NEXT:    [[TMP6:%.*]] = load i32, i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    [[TMP7:%.*]] = load i32, i32* [[END_ADDR]], align 4
+// IR-NEXT:    [[TMP8:%.*]] = load i32, i32* [[START_ADDR]], align 4
+// IR-NEXT:    [[SUB5:%.*]] = sub nsw i32 [[TMP7]], [[TMP8]]
+// IR-NEXT:    [[SUB6:%.*]] = sub nsw i32 [[SUB5]], 1
+// IR-NEXT:    [[TMP9:%.*]] = load i32, i32* [[STEP_ADDR]], align 4
+// IR-NEXT:    [[ADD7:%.*]] = add nsw i32 [[SUB6]], [[TMP9]]
+// IR-NEXT:    [[TMP10:%.*]] = load i32, i32* [[STEP_ADDR]], align 4
+// IR-NEXT:    [[DIV8:%.*]] = sdiv i32 [[ADD7]], [[TMP10]]
+// IR-NEXT:    [[SUB9:%.*]] = sub nsw i32 [[DIV8]], 1
+// IR-NEXT:    [[ADD10:%.*]] = add nsw i32 [[SUB9]], 1
+// IR-NEXT:    [[TMP11:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I]], align 4
+// IR-NEXT:    [[ADD11:%.*]] = add nsw i32 [[TMP11]], 5
+// IR-NEXT:    [[CMP12:%.*]] = icmp slt i32 [[ADD10]], [[ADD11]]
+// IR-NEXT:    br i1 [[CMP12]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
+// IR:       cond.true:
+// IR-NEXT:    [[TMP12:%.*]] = load i32, i32* [[END_ADDR]], align 4
+// IR-NEXT:    [[TMP13:%.*]] = load i32, i32* [[START_ADDR]], align 4
+// IR-NEXT:    [[SUB13:%.*]] = sub nsw i32 [[TMP12]], [[TMP13]]
+// IR-NEXT:    [[SUB14:%.*]] = sub nsw i32 [[SUB13]], 1
+// IR-NEXT:    [[TMP14:%.*]] = load i32, i32* [[STEP_ADDR]], align 4
+// IR-NEXT:    [[ADD15:%.*]] = add nsw i32 [[SUB14]], [[TMP14]]
+// IR-NEXT:    [[TMP15:%.*]] = load i32, i32* [[STEP_ADDR]], align 4
+// IR-NEXT:    [[DIV16:%.*]] = sdiv i32 [[ADD15]], [[TMP15]]
+// IR-NEXT:    [[SUB17:%.*]] = sub nsw i32 [[DIV16]], 1
+// IR-NEXT:    [[ADD18:%.*]] = add nsw i32 [[SUB17]], 1
+// IR-NEXT:    br label [[COND_END:%.*]]
+// IR:       cond.false:
+// IR-NEXT:    [[TMP16:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I]], align 4
+// IR-NEXT:    [[ADD19:%.*]] = add nsw i32 [[TMP16]], 5
+// IR-NEXT:    br label [[COND_END]]
+// IR:       cond.end:
+// IR-NEXT:    [[COND:%.*]] = phi i32 [ [[ADD18]], [[COND_TRUE]] ], [ [[ADD19]], [[COND_FALSE]] ]
+// IR-NEXT:    [[CMP20:%.*]] = icmp slt i32 [[TMP6]], [[COND]]
+// IR-NEXT:    br i1 [[CMP20]], label [[FOR_BODY21:%.*]], label [[FOR_END:%.*]]
+// IR:       for.body21:
+// IR-NEXT:    [[TMP17:%.*]] = load i32, i32* [[START_ADDR]], align 4
+// IR-NEXT:    store i32 [[TMP17]], i32* [[I]], align 4
+// IR-NEXT:    [[TMP19:%.*]] = load i32, i32* [[START_ADDR]], align 4
+// IR-NEXT:    [[TMP20:%.*]] = load i32, i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    [[TMP21:%.*]] = load i32, i32* [[STEP_ADDR]], align 4
+// IR-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP20]], [[TMP21]]
+// IR-NEXT:    [[ADD22:%.*]] = add nsw i32 [[TMP19]], [[MUL]]
+// IR-NEXT:    store i32 [[ADD22]], i32* [[I]], align 4
+// IR-NEXT:    [[TMP22:%.*]] = load i32, i32* [[I]], align 4
+// IR-NEXT:    call void (...) @body(i32 [[TMP22]])
+// IR-NEXT:    br label [[FOR_INC:%.*]]
+// IR:       for.inc:
+// IR-NEXT:    [[TMP23:%.*]] = load i32, i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP23]], 1
+// IR-NEXT:    store i32 [[INC]], i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    br label [[FOR_COND4]]
+// IR:       for.end:
+// IR-NEXT:    br label [[FOR_INC23:%.*]]
+// IR:       for.inc23:
+// IR-NEXT:    [[TMP24:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I]], align 4
+// IR-NEXT:    [[ADD24:%.*]] = add nsw i32 [[TMP24]], 5
+// IR-NEXT:    store i32 [[ADD24]], i32* [[DOTFLOOR_0_IV_I]], align 4
+// IR-NEXT:    br label [[FOR_COND]]
+// IR:       for.end25:
+// IR-NEXT:    ret void
+//
+extern "C" void foo1(int start, int end, int step) {
+#pragma omp tile sizes(5)
+  for (int i = start; i < end; i += step)
+    body(i);
+}
+
+
+// IR-LABEL: @foo2(
+// 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:    [[DOTFLOOR_0_IV_I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTTILE_0_IV_I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTFLOOR_1_IV_J:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTTILE_1_IV_J:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[J:%.*]] = alloca i32, align 4
+// 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:    store i32 0, i32* [[DOTFLOOR_0_IV_I]], align 4
+// IR-NEXT:    br label [[FOR_COND:%.*]]
+// IR:       for.cond:
+// IR-NEXT:    [[TMP0:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I]], align 4
+// IR-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP0]], 4
+// IR-NEXT:    br i1 [[CMP]], label [[FOR_BODY:%.*]], label [[FOR_END30:%.*]]
+// IR:       for.body:
+// IR-NEXT:    store i32 0, i32* [[DOTFLOOR_1_IV_J]], align 4
+// IR-NEXT:    br label [[FOR_COND1:%.*]]
+// IR:       for.cond1:
+// IR-NEXT:    [[TMP1:%.*]] = load i32, i32* [[DOTFLOOR_1_IV_J]], align 4
+// IR-NEXT:    [[CMP2:%.*]] = icmp slt i32 [[TMP1]], 4
+// IR-NEXT:    br i1 [[CMP2]], label [[FOR_BODY3:%.*]], label [[FOR_END27:%.*]]
+// IR:       for.body3:
+// IR-NEXT:    [[TMP2:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I]], align 4
+// IR-NEXT:    store i32 [[TMP2]], i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    br label [[FOR_COND4:%.*]]
+// IR:       for.cond4:
+// IR-NEXT:    [[TMP3:%.*]] = load i32, i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    [[TMP4:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I]], align 4
+// IR-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP4]], 5
+// IR-NEXT:    [[CMP5:%.*]] = icmp slt i32 4, [[ADD]]
+// IR-NEXT:    br i1 [[CMP5]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
+// IR:       cond.true:
+// IR-NEXT:    br label [[COND_END:%.*]]
+// IR:       cond.false:
+// IR-NEXT:    [[TMP5:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I]], align 4
+// IR-NEXT:    [[ADD6:%.*]] = add nsw i32 [[TMP5]], 5
+// IR-NEXT:    br label [[COND_END]]
+// IR:       cond.end:
+// IR-NEXT:    [[COND:%.*]] = phi i32 [ 4, [[COND_TRUE]] ], [ [[ADD6]], [[COND_FALSE]] ]
+// IR-NEXT:    [[CMP7:%.*]] = icmp slt i32 [[TMP3]], [[COND]]
+// IR-NEXT:    br i1 [[CMP7]], label [[FOR_BODY8:%.*]], label [[FOR_END24:%.*]]
+// IR:       for.body8:
+// IR-NEXT:    [[TMP6:%.*]] = load i32, i32* [[DOTFLOOR_1_IV_J]], align 4
+// IR-NEXT:    store i32 [[TMP6]], i32* [[DOTTILE_1_IV_J]], align 4
+// IR-NEXT:    br label [[FOR_COND9:%.*]]
+// IR:       for.cond9:
+// IR-NEXT:    [[TMP7:%.*]] = load i32, i32* [[DOTTILE_1_IV_J]], align 4
+// IR-NEXT:    [[TMP8:%.*]] = load i32, i32* [[DOTFLOOR_1_IV_J]], align 4
+// IR-NEXT:    [[ADD10:%.*]] = add nsw i32 [[TMP8]], 5
+// IR-NEXT:    [[CMP11:%.*]] = icmp slt i32 4, [[ADD10]]
+// IR-NEXT:    br i1 [[CMP11]], label [[COND_TRUE12:%.*]], label [[COND_FALSE13:%.*]]
+// IR:       cond.true12:
+// IR-NEXT:    br label [[COND_END15:%.*]]
+// IR:       cond.false13:
+// IR-NEXT:    [[TMP9:%.*]] = load i32, i32* [[DOTFLOOR_1_IV_J]], align 4
+// IR-NEXT:    [[ADD14:%.*]] = add nsw i32 [[TMP9]], 5
+// IR-NEXT:    br label [[COND_END15]]
+// IR:       cond.end15:
+// IR-NEXT:    [[COND16:%.*]] = phi i32 [ 4, [[COND_TRUE12]] ], [ [[ADD14]], [[COND_FALSE13]] ]
+// IR-NEXT:    [[CMP17:%.*]] = icmp slt i32 [[TMP7]], [[COND16]]
+// IR-NEXT:    br i1 [[CMP17]], label [[FOR_BODY18:%.*]], label [[FOR_END:%.*]]
+// IR:       for.body18:
+// IR-NEXT:    store i32 7, i32* [[I]], align 4
+// IR-NEXT:    [[TMP11:%.*]] = load i32, i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP11]], 3
+// IR-NEXT:    [[ADD19:%.*]] = add nsw i32 7, [[MUL]]
+// IR-NEXT:    store i32 [[ADD19]], i32* [[I]], align 4
+// IR-NEXT:    store i32 7, i32* [[J]], align 4
+// IR-NEXT:    [[TMP13:%.*]] = load i32, i32* [[DOTTILE_1_IV_J]], align 4
+// IR-NEXT:    [[MUL20:%.*]] = mul nsw i32 [[TMP13]], 3
+// IR-NEXT:    [[ADD21:%.*]] = add nsw i32 7, [[MUL20]]
+// IR-NEXT:    store i32 [[ADD21]], i32* [[J]], align 4
+// IR-NEXT:    [[TMP14:%.*]] = load i32, i32* [[I]], align 4
+// IR-NEXT:    [[TMP15:%.*]] = load i32, i32* [[J]], align 4
+// IR-NEXT:    call void (...) @body(i32 [[TMP14]], i32 [[TMP15]])
+// IR-NEXT:    br label [[FOR_INC:%.*]]
+// IR:       for.inc:
+// IR-NEXT:    [[TMP16:%.*]] = load i32, i32* [[DOTTILE_1_IV_J]], align 4
+// IR-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP16]], 1
+// IR-NEXT:    store i32 [[INC]], i32* [[DOTTILE_1_IV_J]], align 4
+// IR-NEXT:    br label [[FOR_COND9]]
+// IR:       for.end:
+// IR-NEXT:    br label [[FOR_INC22:%.*]]
+// IR:       for.inc22:
+// IR-NEXT:    [[TMP17:%.*]] = load i32, i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    [[INC23:%.*]] = add nsw i32 [[TMP17]], 1
+// IR-NEXT:    store i32 [[INC23]], i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    br label [[FOR_COND4]]
+// IR:       for.end24:
+// IR-NEXT:    br label [[FOR_INC25:%.*]]
+// IR:       for.inc25:
+// IR-NEXT:    [[TMP18:%.*]] = load i32, i32* [[DOTFLOOR_1_IV_J]], align 4
+// IR-NEXT:    [[ADD26:%.*]] = add nsw i32 [[TMP18]], 5
+// IR-NEXT:    store i32 [[ADD26]], i32* [[DOTFLOOR_1_IV_J]], align 4
+// IR-NEXT:    br label [[FOR_COND1]]
+// IR:       for.end27:
+// IR-NEXT:    br label [[FOR_INC28:%.*]]
+// IR:       for.inc28:
+// IR-NEXT:    [[TMP19:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I]], align 4
+// IR-NEXT:    [[ADD29:%.*]] = add nsw i32 [[TMP19]], 5
+// IR-NEXT:    store i32 [[ADD29]], i32* [[DOTFLOOR_0_IV_I]], align 4
+// IR-NEXT:    br label [[FOR_COND]]
+// IR:       for.end30:
+// IR-NEXT:    ret void
+//
+extern "C" void foo2(int start, int end, int step) {
+#pragma omp tile sizes(5,5)
+  for (int i = 7; i < 17; i+=3)
+    for (int j = 7; j < 17; j+=3)
+      body(i,j);
+}
+
+
+// IR-LABEL: @foo3(
+// IR-NEXT:  entry:
+// IR-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTFLOOR_0_IV_I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTTILE_0_IV_I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTFLOOR_1_IV_J:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTTILE_1_IV_J:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[TMP:%.*]] = 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:    [[DOTFLOOR_0_IV_I1:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[J:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @1)
+// IR-NEXT:    store i32 0, i32* [[DOTOMP_LB]], align 4
+// IR-NEXT:    store i32 0, 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_4(%struct.ident_t* @0, i32 [[TMP0]], i32 34, i32* [[DOTOMP_IS_LAST]], i32* [[DOTOMP_LB]], i32* [[DOTOMP_UB]], i32* [[DOTOMP_STRIDE]], i32 1, i32 1)
+// IR-NEXT:    [[TMP1:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
+// IR-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP1]], 0
+// IR-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
+// IR:       cond.true:
+// IR-NEXT:    br label [[COND_END:%.*]]
+// IR:       cond.false:
+// IR-NEXT:    [[TMP2:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
+// IR-NEXT:    br label [[COND_END]]
+// IR:       cond.end:
+// IR-NEXT:    [[COND:%.*]] = phi i32 [ 0, [[COND_TRUE]] ], [ [[TMP2]], [[COND_FALSE]] ]
+// IR-NEXT:    store i32 [[COND]], i32* [[DOTOMP_UB]], align 4
+// IR-NEXT:    [[TMP3:%.*]] = load i32, i32* [[DOTOMP_LB]], align 4
+// IR-NEXT:    store i32 [[TMP3]], i32* [[DOTOMP_IV]], align 4
+// IR-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// IR:       omp.inner.for.cond:
+// IR-NEXT:    [[TMP4:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
+// IR-NEXT:    [[TMP5:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
+// IR-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP4]], [[TMP5]]
+// IR-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// IR:       omp.inner.for.body:
+// IR-NEXT:    [[TMP6:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
+// IR-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP6]], 5
+// IR-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
+// IR-NEXT:    store i32 [[ADD]], i32* [[DOTFLOOR_0_IV_I1]], align 4
+// IR-NEXT:    store i32 0, i32* [[DOTFLOOR_1_IV_J]], align 4
+// IR-NEXT:    br label [[FOR_COND:%.*]]
+// IR:       for.cond:
+// IR-NEXT:    [[TMP7:%.*]] = load i32, i32* [[DOTFLOOR_1_IV_J]], align 4
+// IR-NEXT:    [[CMP3:%.*]] = icmp slt i32 [[TMP7]], 4
+// IR-NEXT:    br i1 [[CMP3]], label [[FOR_BODY:%.*]], label [[FOR_END33:%.*]]
+// IR:       for.body:
+// IR-NEXT:    [[TMP8:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I1]], align 4
+// IR-NEXT:    store i32 [[TMP8]], i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    br label [[FOR_COND4:%.*]]
+// IR:       for.cond4:
+// IR-NEXT:    [[TMP9:%.*]] = load i32, i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    [[TMP10:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I1]], align 4
+// IR-NEXT:    [[ADD5:%.*]] = add nsw i32 [[TMP10]], 5
+// IR-NEXT:    [[CMP6:%.*]] = icmp slt i32 4, [[ADD5]]
+// IR-NEXT:    br i1 [[CMP6]], label [[COND_TRUE7:%.*]], label [[COND_FALSE8:%.*]]
+// IR:       cond.true7:
+// IR-NEXT:    br label [[COND_END10:%.*]]
+// IR:       cond.false8:
+// IR-NEXT:    [[TMP11:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I1]], align 4
+// IR-NEXT:    [[ADD9:%.*]] = add nsw i32 [[TMP11]], 5
+// IR-NEXT:    br label [[COND_END10]]
+// IR:       cond.end10:
+// IR-NEXT:    [[COND11:%.*]] = phi i32 [ 4, [[COND_TRUE7]] ], [ [[ADD9]], [[COND_FALSE8]] ]
+// IR-NEXT:    [[CMP12:%.*]] = icmp slt i32 [[TMP9]], [[COND11]]
+// IR-NEXT:    br i1 [[CMP12]], label [[FOR_BODY13:%.*]], label [[FOR_END30:%.*]]
+// IR:       for.body13:
+// IR-NEXT:    [[TMP12:%.*]] = load i32, i32* [[DOTFLOOR_1_IV_J]], align 4
+// IR-NEXT:    store i32 [[TMP12]], i32* [[DOTTILE_1_IV_J]], align 4
+// IR-NEXT:    br label [[FOR_COND14:%.*]]
+// IR:       for.cond14:
+// IR-NEXT:    [[TMP13:%.*]] = load i32, i32* [[DOTTILE_1_IV_J]], align 4
+// IR-NEXT:    [[TMP14:%.*]] = load i32, i32* [[DOTFLOOR_1_IV_J]], align 4
+// IR-NEXT:    [[ADD15:%.*]] = add nsw i32 [[TMP14]], 5
+// IR-NEXT:    [[CMP16:%.*]] = icmp slt i32 4, [[ADD15]]
+// IR-NEXT:    br i1 [[CMP16]], label [[COND_TRUE17:%.*]], label [[COND_FALSE18:%.*]]
+// IR:       cond.true17:
+// IR-NEXT:    br label [[COND_END20:%.*]]
+// IR:       cond.false18:
+// IR-NEXT:    [[TMP15:%.*]] = load i32, i32* [[DOTFLOOR_1_IV_J]], align 4
+// IR-NEXT:    [[ADD19:%.*]] = add nsw i32 [[TMP15]], 5
+// IR-NEXT:    br label [[COND_END20]]
+// IR:       cond.end20:
+// IR-NEXT:    [[COND21:%.*]] = phi i32 [ 4, [[COND_TRUE17]] ], [ [[ADD19]], [[COND_FALSE18]] ]
+// IR-NEXT:    [[CMP22:%.*]] = icmp slt i32 [[TMP13]], [[COND21]]
+// IR-NEXT:    br i1 [[CMP22]], label [[FOR_BODY23:%.*]], label [[FOR_END:%.*]]
+// IR:       for.body23:
+// IR-NEXT:    store i32 7, i32* [[I]], align 4
+// IR-NEXT:    [[TMP17:%.*]] = load i32, i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    [[MUL24:%.*]] = mul nsw i32 [[TMP17]], 3
+// IR-NEXT:    [[ADD25:%.*]] = add nsw i32 7, [[MUL24]]
+// IR-NEXT:    store i32 [[ADD25]], i32* [[I]], align 4
+// IR-NEXT:    store i32 7, i32* [[J]], align 4
+// IR-NEXT:    [[TMP19:%.*]] = load i32, i32* [[DOTTILE_1_IV_J]], align 4
+// IR-NEXT:    [[MUL26:%.*]] = mul nsw i32 [[TMP19]], 3
+// IR-NEXT:    [[ADD27:%.*]] = add nsw i32 7, [[MUL26]]
+// IR-NEXT:    store i32 [[ADD27]], i32* [[J]], align 4
+// IR-NEXT:    br label [[FOR_INC:%.*]]
+// IR:       for.inc:
+// IR-NEXT:    [[TMP20:%.*]] = load i32, i32* [[DOTTILE_1_IV_J]], align 4
+// IR-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP20]], 1
+// IR-NEXT:    store i32 [[INC]], i32* [[DOTTILE_1_IV_J]], align 4
+// IR-NEXT:    br label [[FOR_COND14]]
+// IR:       for.end:
+// IR-NEXT:    br label [[FOR_INC28:%.*]]
+// IR:       for.inc28:
+// IR-NEXT:    [[TMP21:%.*]] = load i32, i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    [[INC29:%.*]] = add nsw i32 [[TMP21]], 1
+// IR-NEXT:    store i32 [[INC29]], i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    br label [[FOR_COND4]]
+// IR:       for.end30:
+// IR-NEXT:    br label [[FOR_INC31:%.*]]
+// IR:       for.inc31:
+// IR-NEXT:    [[TMP22:%.*]] = load i32, i32* [[DOTFLOOR_1_IV_J]], align 4
+// IR-NEXT:    [[ADD32:%.*]] = add nsw i32 [[TMP22]], 5
+// IR-NEXT:    store i32 [[ADD32]], i32* [[DOTFLOOR_1_IV_J]], align 4
+// IR-NEXT:    br label [[FOR_COND]]
+// IR:       for.end33:
+// IR-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// IR:       omp.body.continue:
+// IR-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// IR:       omp.inner.for.inc:
+// IR-NEXT:    [[TMP23:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
+// IR-NEXT:    [[ADD34:%.*]] = add nsw i32 [[TMP23]], 1
+// IR-NEXT:    store i32 [[ADD34]], i32* [[DOTOMP_IV]], align 4
+// IR-NEXT:    br label [[OMP_INNER_FOR_COND]]
+// IR:       omp.inner.for.end:
+// IR-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]
+// IR:       omp.loop.exit:
+// IR-NEXT:    call void @__kmpc_for_static_fini(%struct.ident_t* @0, i32 [[TMP0]])
+// IR-NEXT:    call void @__kmpc_barrier(%struct.ident_t* @2, i32 [[TMP0]])
+// IR-NEXT:    ret void
+//
+extern "C" void foo3() {
+#pragma omp for
+#pragma omp tile sizes(5,5)
+    for (int i = 7; i < 17; i += 3)
+      for (int j = 7; j < 17; j += 3)
+        ;
+}
+
+
+// IR-LABEL: @foo4(
+// IR-NEXT:  entry:
+// IR-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTFLOOR_0_IV_I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTTILE_0_IV_I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTFLOOR_1_IV_J:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTTILE_1_IV_J:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[TMP1:%.*]] = 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:    [[K:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTFLOOR_0_IV_I2:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[J:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @1)
+// IR-NEXT:    store i32 0, i32* [[DOTOMP_LB]], align 4
+// IR-NEXT:    store i32 3, 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_4(%struct.ident_t* @0, i32 [[TMP0]], i32 34, i32* [[DOTOMP_IS_LAST]], i32* [[DOTOMP_LB]], i32* [[DOTOMP_UB]], i32* [[DOTOMP_STRIDE]], i32 1, i32 1)
+// IR-NEXT:    [[TMP1:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
+// IR-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP1]], 3
+// IR-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
+// IR:       cond.true:
+// IR-NEXT:    br label [[COND_END:%.*]]
+// IR:       cond.false:
+// IR-NEXT:    [[TMP2:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
+// IR-NEXT:    br label [[COND_END]]
+// IR:       cond.end:
+// IR-NEXT:    [[COND:%.*]] = phi i32 [ 3, [[COND_TRUE]] ], [ [[TMP2]], [[COND_FALSE]] ]
+// IR-NEXT:    store i32 [[COND]], i32* [[DOTOMP_UB]], align 4
+// IR-NEXT:    [[TMP3:%.*]] = load i32, i32* [[DOTOMP_LB]], align 4
+// IR-NEXT:    store i32 [[TMP3]], i32* [[DOTOMP_IV]], align 4
+// IR-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// IR:       omp.inner.for.cond:
+// IR-NEXT:    [[TMP4:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
+// IR-NEXT:    [[TMP5:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
+// IR-NEXT:    [[CMP3:%.*]] = icmp sle i32 [[TMP4]], [[TMP5]]
+// IR-NEXT:    br i1 [[CMP3]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// IR:       omp.inner.for.body:
+// IR-NEXT:    [[TMP6:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
+// IR-NEXT:    [[DIV:%.*]] = sdiv i32 [[TMP6]], 1
+// IR-NEXT:    [[MUL:%.*]] = mul nsw i32 [[DIV]], 3
+// IR-NEXT:    [[ADD:%.*]] = add nsw i32 7, [[MUL]]
+// IR-NEXT:    store i32 [[ADD]], i32* [[K]], align 4
+// IR-NEXT:    [[TMP7:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
+// IR-NEXT:    [[TMP8:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
+// IR-NEXT:    [[DIV4:%.*]] = sdiv i32 [[TMP8]], 1
+// IR-NEXT:    [[MUL5:%.*]] = mul nsw i32 [[DIV4]], 1
+// IR-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP7]], [[MUL5]]
+// IR-NEXT:    [[MUL6:%.*]] = mul nsw i32 [[SUB]], 5
+// IR-NEXT:    [[ADD7:%.*]] = add nsw i32 0, [[MUL6]]
+// IR-NEXT:    store i32 [[ADD7]], i32* [[DOTFLOOR_0_IV_I2]], align 4
+// IR-NEXT:    store i32 0, i32* [[DOTFLOOR_1_IV_J]], align 4
+// IR-NEXT:    br label [[FOR_COND:%.*]]
+// IR:       for.cond:
+// IR-NEXT:    [[TMP9:%.*]] = load i32, i32* [[DOTFLOOR_1_IV_J]], align 4
+// IR-NEXT:    [[CMP8:%.*]] = icmp slt i32 [[TMP9]], 4
+// IR-NEXT:    br i1 [[CMP8]], label [[FOR_BODY:%.*]], label [[FOR_END38:%.*]]
+// IR:       for.body:
+// IR-NEXT:    [[TMP10:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I2]], align 4
+// IR-NEXT:    store i32 [[TMP10]], i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    br label [[FOR_COND9:%.*]]
+// IR:       for.cond9:
+// IR-NEXT:    [[TMP11:%.*]] = load i32, i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    [[TMP12:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I2]], align 4
+// IR-NEXT:    [[ADD10:%.*]] = add nsw i32 [[TMP12]], 5
+// IR-NEXT:    [[CMP11:%.*]] = icmp slt i32 4, [[ADD10]]
+// IR-NEXT:    br i1 [[CMP11]], label [[COND_TRUE12:%.*]], label [[COND_FALSE13:%.*]]
+// IR:       cond.true12:
+// IR-NEXT:    br label [[COND_END15:%.*]]
+// IR:       cond.false13:
+// IR-NEXT:    [[TMP13:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I2]], align 4
+// IR-NEXT:    [[ADD14:%.*]] = add nsw i32 [[TMP13]], 5
+// IR-NEXT:    br label [[COND_END15]]
+// IR:       cond.end15:
+// IR-NEXT:    [[COND16:%.*]] = phi i32 [ 4, [[COND_TRUE12]] ], [ [[ADD14]], [[COND_FALSE13]] ]
+// IR-NEXT:    [[CMP17:%.*]] = icmp slt i32 [[TMP11]], [[COND16]]
+// IR-NEXT:    br i1 [[CMP17]], label [[FOR_BODY18:%.*]], label [[FOR_END35:%.*]]
+// IR:       for.body18:
+// IR-NEXT:    [[TMP14:%.*]] = load i32, i32* [[DOTFLOOR_1_IV_J]], align 4
+// IR-NEXT:    store i32 [[TMP14]], i32* [[DOTTILE_1_IV_J]], align 4
+// IR-NEXT:    br label [[FOR_COND19:%.*]]
+// IR:       for.cond19:
+// IR-NEXT:    [[TMP15:%.*]] = load i32, i32* [[DOTTILE_1_IV_J]], align 4
+// IR-NEXT:    [[TMP16:%.*]] = load i32, i32* [[DOTFLOOR_1_IV_J]], align 4
+// IR-NEXT:    [[ADD20:%.*]] = add nsw i32 [[TMP16]], 5
+// IR-NEXT:    [[CMP21:%.*]] = icmp slt i32 4, [[ADD20]]
+// IR-NEXT:    br i1 [[CMP21]], label [[COND_TRUE22:%.*]], label [[COND_FALSE23:%.*]]
+// IR:       cond.true22:
+// IR-NEXT:    br label [[COND_END25:%.*]]
+// IR:       cond.false23:
+// IR-NEXT:    [[TMP17:%.*]] = load i32, i32* [[DOTFLOOR_1_IV_J]], align 4
+// IR-NEXT:    [[ADD24:%.*]] = add nsw i32 [[TMP17]], 5
+// IR-NEXT:    br label [[COND_END25]]
+// IR:       cond.end25:
+// IR-NEXT:    [[COND26:%.*]] = phi i32 [ 4, [[COND_TRUE22]] ], [ [[ADD24]], [[COND_FALSE23]] ]
+// IR-NEXT:    [[CMP27:%.*]] = icmp slt i32 [[TMP15]], [[COND26]]
+// IR-NEXT:    br i1 [[CMP27]], label [[FOR_BODY28:%.*]], label [[FOR_END:%.*]]
+// IR:       for.body28:
+// IR-NEXT:    store i32 7, i32* [[I]], align 4
+// IR-NEXT:    [[TMP19:%.*]] = load i32, i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    [[MUL29:%.*]] = mul nsw i32 [[TMP19]], 3
+// IR-NEXT:    [[ADD30:%.*]] = add nsw i32 7, [[MUL29]]
+// IR-NEXT:    store i32 [[ADD30]], i32* [[I]], align 4
+// IR-NEXT:    store i32 7, i32* [[J]], align 4
+// IR-NEXT:    [[TMP21:%.*]] = load i32, i32* [[DOTTILE_1_IV_J]], align 4
+// IR-NEXT:    [[MUL31:%.*]] = mul nsw i32 [[TMP21]], 3
+// IR-NEXT:    [[ADD32:%.*]] = add nsw i32 7, [[MUL31]]
+// IR-NEXT:    store i32 [[ADD32]], i32* [[J]], align 4
+// IR-NEXT:    br label [[FOR_INC:%.*]]
+// IR:       for.inc:
+// IR-NEXT:    [[TMP22:%.*]] = load i32, i32* [[DOTTILE_1_IV_J]], align 4
+// IR-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP22]], 1
+// IR-NEXT:    store i32 [[INC]], i32* [[DOTTILE_1_IV_J]], align 4
+// IR-NEXT:    br label [[FOR_COND19]]
+// IR:       for.end:
+// IR-NEXT:    br label [[FOR_INC33:%.*]]
+// IR:       for.inc33:
+// IR-NEXT:    [[TMP23:%.*]] = load i32, i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    [[INC34:%.*]] = add nsw i32 [[TMP23]], 1
+// IR-NEXT:    store i32 [[INC34]], i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    br label [[FOR_COND9]]
+// IR:       for.end35:
+// IR-NEXT:    br label [[FOR_INC36:%.*]]
+// IR:       for.inc36:
+// IR-NEXT:    [[TMP24:%.*]] = load i32, i32* [[DOTFLOOR_1_IV_J]], align 4
+// IR-NEXT:    [[ADD37:%.*]] = add nsw i32 [[TMP24]], 5
+// IR-NEXT:    store i32 [[ADD37]], i32* [[DOTFLOOR_1_IV_J]], align 4
+// IR-NEXT:    br label [[FOR_COND]]
+// IR:       for.end38:
+// IR-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// IR:       omp.body.continue:
+// IR-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// IR:       omp.inner.for.inc:
+// IR-NEXT:    [[TMP25:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
+// IR-NEXT:    [[ADD39:%.*]] = add nsw i32 [[TMP25]], 1
+// IR-NEXT:    store i32 [[ADD39]], i32* [[DOTOMP_IV]], align 4
+// IR-NEXT:    br label [[OMP_INNER_FOR_COND]]
+// IR:       omp.inner.for.end:
+// IR-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]
+// IR:       omp.loop.exit:
+// IR-NEXT:    call void @__kmpc_for_static_fini(%struct.ident_t* @0, i32 [[TMP0]])
+// IR-NEXT:    call void @__kmpc_barrier(%struct.ident_t* @2, i32 [[TMP0]])
+// IR-NEXT:    ret void
+//
+extern "C" void foo4() {
+#pragma omp for collapse(2)
+  for (int k = 7; k < 17; k += 3)
+#pragma omp tile sizes(5,5)
+  for (int i = 7; i < 17; i += 3)
+    for (int j = 7; j < 17; j += 3)
+      ;
+}
+
+
+// IR-LABEL: @foo5(
+// IR-NEXT:  entry:
+// IR-NEXT:    [[DOTOMP_IV:%.*]] = alloca i64, align 8
+// IR-NEXT:    [[DOTFLOOR_0_IV_I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTTILE_0_IV_I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[TMP:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[TMP1:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[TMP2:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTCAPTURE_EXPR_:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTCAPTURE_EXPR_3:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTCAPTURE_EXPR_5:%.*]] = alloca i64, align 8
+// IR-NEXT:    [[DOTFLOOR_0_IV_I10:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTTILE_0_IV_I11:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[J:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTOMP_LB:%.*]] = alloca i64, align 8
+// IR-NEXT:    [[DOTOMP_UB:%.*]] = alloca i64, align 8
+// IR-NEXT:    [[DOTOMP_STRIDE:%.*]] = alloca i64, align 8
+// IR-NEXT:    [[DOTOMP_IS_LAST:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTFLOOR_0_IV_I13:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTTILE_0_IV_I14:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[J15:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* @1)
+// IR-NEXT:    [[TMP1:%.*]] = load i32, i32* [[TMP]], align 4
+// IR-NEXT:    store i32 [[TMP1]], i32* [[DOTCAPTURE_EXPR_]], align 4
+// IR-NEXT:    [[TMP2:%.*]] = load i32, i32* [[TMP]], align 4
+// IR-NEXT:    [[ADD:%.*]] = add nsw i32 [[TMP2]], 5
+// IR-NEXT:    [[CMP:%.*]] = icmp slt i32 4, [[ADD]]
+// IR-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
+// IR:       cond.true:
+// IR-NEXT:    br label [[COND_END:%.*]]
+// IR:       cond.false:
+// IR-NEXT:    [[TMP3:%.*]] = load i32, i32* [[TMP]], align 4
+// IR-NEXT:    [[ADD4:%.*]] = add nsw i32 [[TMP3]], 5
+// IR-NEXT:    br label [[COND_END]]
+// IR:       cond.end:
+// IR-NEXT:    [[COND:%.*]] = phi i32 [ 4, [[COND_TRUE]] ], [ [[ADD4]], [[COND_FALSE]] ]
+// IR-NEXT:    store i32 [[COND]], i32* [[DOTCAPTURE_EXPR_3]], align 4
+// IR-NEXT:    [[TMP4:%.*]] = load i32, i32* [[DOTCAPTURE_EXPR_3]], align 4
+// IR-NEXT:    [[TMP5:%.*]] = load i32, i32* [[DOTCAPTURE_EXPR_]], align 4
+// IR-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP4]], [[TMP5]]
+// IR-NEXT:    [[SUB6:%.*]] = sub nsw i32 [[SUB]], 1
+// IR-NEXT:    [[ADD7:%.*]] = add nsw i32 [[SUB6]], 1
+// IR-NEXT:    [[DIV:%.*]] = sdiv i32 [[ADD7]], 1
+// IR-NEXT:    [[CONV:%.*]] = sext i32 [[DIV]] to i64
+// IR-NEXT:    [[MUL:%.*]] = mul nsw i64 1, [[CONV]]
+// IR-NEXT:    [[MUL8:%.*]] = mul nsw i64 [[MUL]], 4
+// IR-NEXT:    [[SUB9:%.*]] = sub nsw i64 [[MUL8]], 1
+// IR-NEXT:    store i64 [[SUB9]], i64* [[DOTCAPTURE_EXPR_5]], align 8
+// IR-NEXT:    store i32 0, i32* [[DOTFLOOR_0_IV_I10]], align 4
+// IR-NEXT:    [[TMP6:%.*]] = load i32, i32* [[DOTCAPTURE_EXPR_]], align 4
+// IR-NEXT:    store i32 [[TMP6]], i32* [[DOTTILE_0_IV_I11]], align 4
+// IR-NEXT:    store i32 7, i32* [[J]], align 4
+// IR-NEXT:    [[TMP7:%.*]] = load i32, i32* [[DOTCAPTURE_EXPR_]], align 4
+// IR-NEXT:    [[TMP8:%.*]] = load i32, i32* [[DOTCAPTURE_EXPR_3]], align 4
+// IR-NEXT:    [[CMP12:%.*]] = icmp slt i32 [[TMP7]], [[TMP8]]
+// IR-NEXT:    br i1 [[CMP12]], label [[OMP_PRECOND_THEN:%.*]], label [[OMP_PRECOND_END:%.*]]
+// IR:       omp.precond.then:
+// IR-NEXT:    store i64 0, i64* [[DOTOMP_LB]], align 8
+// IR-NEXT:    [[TMP9:%.*]] = load i64, i64* [[DOTCAPTURE_EXPR_5]], align 8
+// IR-NEXT:    store i64 [[TMP9]], i64* [[DOTOMP_UB]], align 8
+// IR-NEXT:    store i64 1, i64* [[DOTOMP_STRIDE]], align 8
+// IR-NEXT:    store i32 0, i32* [[DOTOMP_IS_LAST]], align 4
+// IR-NEXT:    call void @__kmpc_for_static_init_8(%struct.ident_t* @0, i32 [[TMP0]], i32 34, i32* [[DOTOMP_IS_LAST]], i64* [[DOTOMP_LB]], i64* [[DOTOMP_UB]], i64* [[DOTOMP_STRIDE]], i64 1, i64 1)
+// IR-NEXT:    [[TMP10:%.*]] = load i64, i64* [[DOTOMP_UB]], align 8
+// IR-NEXT:    [[TMP11:%.*]] = load i64, i64* [[DOTCAPTURE_EXPR_5]], align 8
+// IR-NEXT:    [[CMP16:%.*]] = icmp sgt i64 [[TMP10]], [[TMP11]]
+// IR-NEXT:    br i1 [[CMP16]], label [[COND_TRUE17:%.*]], label [[COND_FALSE18:%.*]]
+// IR:       cond.true17:
+// IR-NEXT:    [[TMP12:%.*]] = load i64, i64* [[DOTCAPTURE_EXPR_5]], align 8
+// IR-NEXT:    br label [[COND_END19:%.*]]
+// IR:       cond.false18:
+// IR-NEXT:    [[TMP13:%.*]] = load i64, i64* [[DOTOMP_UB]], align 8
+// IR-NEXT:    br label [[COND_END19]]
+// IR:       cond.end19:
+// IR-NEXT:    [[COND20:%.*]] = phi i64 [ [[TMP12]], [[COND_TRUE17]] ], [ [[TMP13]], [[COND_FALSE18]] ]
+// IR-NEXT:    store i64 [[COND20]], i64* [[DOTOMP_UB]], align 8
+// IR-NEXT:    [[TMP14:%.*]] = load i64, i64* [[DOTOMP_LB]], align 8
+// IR-NEXT:    store i64 [[TMP14]], i64* [[DOTOMP_IV]], align 8
+// IR-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// IR:       omp.inner.for.cond:
+// IR-NEXT:    [[TMP15:%.*]] = load i64, i64* [[DOTOMP_IV]], align 8
+// IR-NEXT:    [[TMP16:%.*]] = load i64, i64* [[DOTOMP_UB]], align 8
+// IR-NEXT:    [[CMP21:%.*]] = icmp sle i64 [[TMP15]], [[TMP16]]
+// IR-NEXT:    br i1 [[CMP21]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// IR:       omp.inner.for.body:
+// IR-NEXT:    [[TMP17:%.*]] = load i64, i64* [[DOTOMP_IV]], align 8
+// IR-NEXT:    [[TMP18:%.*]] = load i32, i32* [[DOTCAPTURE_EXPR_3]], align 4
+// IR-NEXT:    [[TMP19:%.*]] = load i32, i32* [[DOTCAPTURE_EXPR_]], align 4
+// IR-NEXT:    [[SUB22:%.*]] = sub nsw i32 [[TMP18]], [[TMP19]]
+// IR-NEXT:    [[SUB23:%.*]] = sub nsw i32 [[SUB22]], 1
+// IR-NEXT:    [[ADD24:%.*]] = add nsw i32 [[SUB23]], 1
+// IR-NEXT:    [[DIV25:%.*]] = sdiv i32 [[ADD24]], 1
+// IR-NEXT:    [[MUL26:%.*]] = mul nsw i32 1, [[DIV25]]
+// IR-NEXT:    [[MUL27:%.*]] = mul nsw i32 [[MUL26]], 4
+// IR-NEXT:    [[CONV28:%.*]] = sext i32 [[MUL27]] to i64
+// IR-NEXT:    [[DIV29:%.*]] = sdiv i64 [[TMP17]], [[CONV28]]
+// IR-NEXT:    [[MUL30:%.*]] = mul nsw i64 [[DIV29]], 5
+// IR-NEXT:    [[ADD31:%.*]] = add nsw i64 0, [[MUL30]]
+// IR-NEXT:    [[CONV32:%.*]] = trunc i64 [[ADD31]] to i32
+// IR-NEXT:    store i32 [[CONV32]], i32* [[DOTFLOOR_0_IV_I13]], align 4
+// IR-NEXT:    [[TMP20:%.*]] = load i32, i32* [[DOTCAPTURE_EXPR_]], align 4
+// IR-NEXT:    [[CONV33:%.*]] = sext i32 [[TMP20]] to i64
+// IR-NEXT:    [[TMP21:%.*]] = load i64, i64* [[DOTOMP_IV]], align 8
+// IR-NEXT:    [[TMP22:%.*]] = load i64, i64* [[DOTOMP_IV]], align 8
+// IR-NEXT:    [[TMP23:%.*]] = load i32, i32* [[DOTCAPTURE_EXPR_3]], align 4
+// IR-NEXT:    [[TMP24:%.*]] = load i32, i32* [[DOTCAPTURE_EXPR_]], align 4
+// IR-NEXT:    [[SUB34:%.*]] = sub nsw i32 [[TMP23]], [[TMP24]]
+// IR-NEXT:    [[SUB35:%.*]] = sub nsw i32 [[SUB34]], 1
+// IR-NEXT:    [[ADD36:%.*]] = add nsw i32 [[SUB35]], 1
+// IR-NEXT:    [[DIV37:%.*]] = sdiv i32 [[ADD36]], 1
+// IR-NEXT:    [[MUL38:%.*]] = mul nsw i32 1, [[DIV37]]
+// IR-NEXT:    [[MUL39:%.*]] = mul nsw i32 [[MUL38]], 4
+// IR-NEXT:    [[CONV40:%.*]] = sext i32 [[MUL39]] to i64
+// IR-NEXT:    [[DIV41:%.*]] = sdiv i64 [[TMP22]], [[CONV40]]
+// IR-NEXT:    [[TMP25:%.*]] = load i32, i32* [[DOTCAPTURE_EXPR_3]], align 4
+// IR-NEXT:    [[TMP26:%.*]] = load i32, i32* [[DOTCAPTURE_EXPR_]], align 4
+// IR-NEXT:    [[SUB42:%.*]] = sub nsw i32 [[TMP25]], [[TMP26]]
+// IR-NEXT:    [[SUB43:%.*]] = sub nsw i32 [[SUB42]], 1
+// IR-NEXT:    [[ADD44:%.*]] = add nsw i32 [[SUB43]], 1
+// IR-NEXT:    [[DIV45:%.*]] = sdiv i32 [[ADD44]], 1
+// IR-NEXT:    [[MUL46:%.*]] = mul nsw i32 1, [[DIV45]]
+// IR-NEXT:    [[MUL47:%.*]] = mul nsw i32 [[MUL46]], 4
+// IR-NEXT:    [[CONV48:%.*]] = sext i32 [[MUL47]] to i64
+// IR-NEXT:    [[MUL49:%.*]] = mul nsw i64 [[DIV41]], [[CONV48]]
+// IR-NEXT:    [[SUB50:%.*]] = sub nsw i64 [[TMP21]], [[MUL49]]
+// IR-NEXT:    [[DIV51:%.*]] = sdiv i64 [[SUB50]], 4
+// IR-NEXT:    [[MUL52:%.*]] = mul nsw i64 [[DIV51]], 1
+// IR-NEXT:    [[ADD53:%.*]] = add nsw i64 [[CONV33]], [[MUL52]]
+// IR-NEXT:    [[CONV54:%.*]] = trunc i64 [[ADD53]] to i32
+// IR-NEXT:    store i32 [[CONV54]], i32* [[DOTTILE_0_IV_I14]], align 4
+// IR-NEXT:    [[TMP27:%.*]] = load i64, i64* [[DOTOMP_IV]], align 8
+// IR-NEXT:    [[TMP28:%.*]] = load i64, i64* [[DOTOMP_IV]], align 8
+// IR-NEXT:    [[TMP29:%.*]] = load i32, i32* [[DOTCAPTURE_EXPR_3]], align 4
+// IR-NEXT:    [[TMP30:%.*]] = load i32, i32* [[DOTCAPTURE_EXPR_]], align 4
+// IR-NEXT:    [[SUB55:%.*]] = sub nsw i32 [[TMP29]], [[TMP30]]
+// IR-NEXT:    [[SUB56:%.*]] = sub nsw i32 [[SUB55]], 1
+// IR-NEXT:    [[ADD57:%.*]] = add nsw i32 [[SUB56]], 1
+// IR-NEXT:    [[DIV58:%.*]] = sdiv i32 [[ADD57]], 1
+// IR-NEXT:    [[MUL59:%.*]] = mul nsw i32 1, [[DIV58]]
+// IR-NEXT:    [[MUL60:%.*]] = mul nsw i32 [[MUL59]], 4
+// IR-NEXT:    [[CONV61:%.*]] = sext i32 [[MUL60]] to i64
+// IR-NEXT:    [[DIV62:%.*]] = sdiv i64 [[TMP28]], [[CONV61]]
+// IR-NEXT:    [[TMP31:%.*]] = load i32, i32* [[DOTCAPTURE_EXPR_3]], align 4
+// IR-NEXT:    [[TMP32:%.*]] = load i32, i32* [[DOTCAPTURE_EXPR_]], align 4
+// IR-NEXT:    [[SUB63:%.*]] = sub nsw i32 [[TMP31]], [[TMP32]]
+// IR-NEXT:    [[SUB64:%.*]] = sub nsw i32 [[SUB63]], 1
+// IR-NEXT:    [[ADD65:%.*]] = add nsw i32 [[SUB64]], 1
+// IR-NEXT:    [[DIV66:%.*]] = sdiv i32 [[ADD65]], 1
+// IR-NEXT:    [[MUL67:%.*]] = mul nsw i32 1, [[DIV66]]
+// IR-NEXT:    [[MUL68:%.*]] = mul nsw i32 [[MUL67]], 4
+// IR-NEXT:    [[CONV69:%.*]] = sext i32 [[MUL68]] to i64
+// IR-NEXT:    [[MUL70:%.*]] = mul nsw i64 [[DIV62]], [[CONV69]]
+// IR-NEXT:    [[SUB71:%.*]] = sub nsw i64 [[TMP27]], [[MUL70]]
+// IR-NEXT:    [[TMP33:%.*]] = load i64, i64* [[DOTOMP_IV]], align 8
+// IR-NEXT:    [[TMP34:%.*]] = load i64, i64* [[DOTOMP_IV]], align 8
+// IR-NEXT:    [[TMP35:%.*]] = load i32, i32* [[DOTCAPTURE_EXPR_3]], align 4
+// IR-NEXT:    [[TMP36:%.*]] = load i32, i32* [[DOTCAPTURE_EXPR_]], align 4
+// IR-NEXT:    [[SUB72:%.*]] = sub nsw i32 [[TMP35]], [[TMP36]]
+// IR-NEXT:    [[SUB73:%.*]] = sub nsw i32 [[SUB72]], 1
+// IR-NEXT:    [[ADD74:%.*]] = add nsw i32 [[SUB73]], 1
+// IR-NEXT:    [[DIV75:%.*]] = sdiv i32 [[ADD74]], 1
+// IR-NEXT:    [[MUL76:%.*]] = mul nsw i32 1, [[DIV75]]
+// IR-NEXT:    [[MUL77:%.*]] = mul nsw i32 [[MUL76]], 4
+// IR-NEXT:    [[CONV78:%.*]] = sext i32 [[MUL77]] to i64
+// IR-NEXT:    [[DIV79:%.*]] = sdiv i64 [[TMP34]], [[CONV78]]
+// IR-NEXT:    [[TMP37:%.*]] = load i32, i32* [[DOTCAPTURE_EXPR_3]], align 4
+// IR-NEXT:    [[TMP38:%.*]] = load i32, i32* [[DOTCAPTURE_EXPR_]], align 4
+// IR-NEXT:    [[SUB80:%.*]] = sub nsw i32 [[TMP37]], [[TMP38]]
+// IR-NEXT:    [[SUB81:%.*]] = sub nsw i32 [[SUB80]], 1
+// IR-NEXT:    [[ADD82:%.*]] = add nsw i32 [[SUB81]], 1
+// IR-NEXT:    [[DIV83:%.*]] = sdiv i32 [[ADD82]], 1
+// IR-NEXT:    [[MUL84:%.*]] = mul nsw i32 1, [[DIV83]]
+// IR-NEXT:    [[MUL85:%.*]] = mul nsw i32 [[MUL84]], 4
+// IR-NEXT:    [[CONV86:%.*]] = sext i32 [[MUL85]] to i64
+// IR-NEXT:    [[MUL87:%.*]] = mul nsw i64 [[DIV79]], [[CONV86]]
+// IR-NEXT:    [[SUB88:%.*]] = sub nsw i64 [[TMP33]], [[MUL87]]
+// IR-NEXT:    [[DIV89:%.*]] = sdiv i64 [[SUB88]], 4
+// IR-NEXT:    [[MUL90:%.*]] = mul nsw i64 [[DIV89]], 4
+// IR-NEXT:    [[SUB91:%.*]] = sub nsw i64 [[SUB71]], [[MUL90]]
+// IR-NEXT:    [[MUL92:%.*]] = mul nsw i64 [[SUB91]], 3
+// IR-NEXT:    [[ADD93:%.*]] = add nsw i64 7, [[MUL92]]
+// IR-NEXT:    [[CONV94:%.*]] = trunc i64 [[ADD93]] to i32
+// IR-NEXT:    store i32 [[CONV94]], i32* [[J15]], align 4
+// IR-NEXT:    store i32 7, i32* [[I]], align 4
+// IR-NEXT:    [[TMP40:%.*]] = load i32, i32* [[DOTTILE_0_IV_I14]], align 4
+// IR-NEXT:    [[MUL95:%.*]] = mul nsw i32 [[TMP40]], 3
+// IR-NEXT:    [[ADD96:%.*]] = add nsw i32 7, [[MUL95]]
+// IR-NEXT:    store i32 [[ADD96]], i32* [[I]], align 4
+// IR-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// IR:       omp.body.continue:
+// IR-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// IR:       omp.inner.for.inc:
+// IR-NEXT:    [[TMP41:%.*]] = load i64, i64* [[DOTOMP_IV]], align 8
+// IR-NEXT:    [[ADD97:%.*]] = add nsw i64 [[TMP41]], 1
+// IR-NEXT:    store i64 [[ADD97]], i64* [[DOTOMP_IV]], align 8
+// IR-NEXT:    br label [[OMP_INNER_FOR_COND]]
+// IR:       omp.inner.for.end:
+// IR-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]
+// IR:       omp.loop.exit:
+// IR-NEXT:    call void @__kmpc_for_static_fini(%struct.ident_t* @0, i32 [[TMP0]])
+// IR-NEXT:    br label [[OMP_PRECOND_END]]
+// IR:       omp.precond.end:
+// IR-NEXT:    call void @__kmpc_barrier(%struct.ident_t* @2, i32 [[TMP0]])
+// IR-NEXT:    ret void
+//
+extern "C" void foo5() {
+#pragma omp for collapse(3)
+#pragma omp tile sizes(5)
+  for (int i = 7; i < 17; i += 3)
+    for (int j = 7; j < 17; j += 3)
+      ;
+}
+
+
+// IR-LABEL: define internal void @.omp_outlined.(
+// IR-NEXT:  entry:
+// IR-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca i32*, align 8
+// IR-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca i32*, align 8
+// IR-NEXT:    [[DOTOMP_IV:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTFLOOR_0_IV_I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTTILE_0_IV_I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[TMP:%.*]] = 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:    [[DOTFLOOR_0_IV_I1:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[I:%.*]] = alloca i32, align 4
+// IR-NEXT:    store i32* [[DOTGLOBAL_TID_:%.*]], i32** [[DOTGLOBAL_TID__ADDR]], align 8
+// IR-NEXT:    store i32* [[DOTBOUND_TID_:%.*]], i32** [[DOTBOUND_TID__ADDR]], align 8
+// IR-NEXT:    store i32 0, i32* [[DOTOMP_LB]], align 4
+// IR-NEXT:    store i32 0, 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:    [[TMP0:%.*]] = load i32*, i32** [[DOTGLOBAL_TID__ADDR]], align 8
+// IR-NEXT:    [[TMP1:%.*]] = load i32, i32* [[TMP0]], align 4
+// IR-NEXT:    call void @__kmpc_for_static_init_4(%struct.ident_t* @0, i32 [[TMP1]], i32 34, i32* [[DOTOMP_IS_LAST]], i32* [[DOTOMP_LB]], i32* [[DOTOMP_UB]], i32* [[DOTOMP_STRIDE]], i32 1, i32 1)
+// IR-NEXT:    [[TMP2:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
+// IR-NEXT:    [[CMP:%.*]] = icmp sgt i32 [[TMP2]], 0
+// IR-NEXT:    br i1 [[CMP]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
+// IR:       cond.true:
+// IR-NEXT:    br label [[COND_END:%.*]]
+// IR:       cond.false:
+// IR-NEXT:    [[TMP3:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
+// IR-NEXT:    br label [[COND_END]]
+// IR:       cond.end:
+// IR-NEXT:    [[COND:%.*]] = phi i32 [ 0, [[COND_TRUE]] ], [ [[TMP3]], [[COND_FALSE]] ]
+// IR-NEXT:    store i32 [[COND]], i32* [[DOTOMP_UB]], align 4
+// IR-NEXT:    [[TMP4:%.*]] = load i32, i32* [[DOTOMP_LB]], align 4
+// IR-NEXT:    store i32 [[TMP4]], i32* [[DOTOMP_IV]], align 4
+// IR-NEXT:    br label [[OMP_INNER_FOR_COND:%.*]]
+// IR:       omp.inner.for.cond:
+// IR-NEXT:    [[TMP5:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
+// IR-NEXT:    [[TMP6:%.*]] = load i32, i32* [[DOTOMP_UB]], align 4
+// IR-NEXT:    [[CMP2:%.*]] = icmp sle i32 [[TMP5]], [[TMP6]]
+// IR-NEXT:    br i1 [[CMP2]], label [[OMP_INNER_FOR_BODY:%.*]], label [[OMP_INNER_FOR_END:%.*]]
+// IR:       omp.inner.for.body:
+// IR-NEXT:    [[TMP7:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
+// IR-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP7]], 5
+// IR-NEXT:    [[ADD:%.*]] = add nsw i32 0, [[MUL]]
+// IR-NEXT:    store i32 [[ADD]], i32* [[DOTFLOOR_0_IV_I1]], align 4
+// IR-NEXT:    [[TMP8:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I1]], align 4
+// IR-NEXT:    store i32 [[TMP8]], i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    br label [[FOR_COND:%.*]]
+// IR:       for.cond:
+// IR-NEXT:    [[TMP9:%.*]] = load i32, i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    [[TMP10:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I1]], align 4
+// IR-NEXT:    [[ADD3:%.*]] = add nsw i32 [[TMP10]], 5
+// IR-NEXT:    [[CMP4:%.*]] = icmp slt i32 4, [[ADD3]]
+// IR-NEXT:    br i1 [[CMP4]], label [[COND_TRUE5:%.*]], label [[COND_FALSE6:%.*]]
+// IR:       cond.true5:
+// IR-NEXT:    br label [[COND_END8:%.*]]
+// IR:       cond.false6:
+// IR-NEXT:    [[TMP11:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I1]], align 4
+// IR-NEXT:    [[ADD7:%.*]] = add nsw i32 [[TMP11]], 5
+// IR-NEXT:    br label [[COND_END8]]
+// IR:       cond.end8:
+// IR-NEXT:    [[COND9:%.*]] = phi i32 [ 4, [[COND_TRUE5]] ], [ [[ADD7]], [[COND_FALSE6]] ]
+// IR-NEXT:    [[CMP10:%.*]] = icmp slt i32 [[TMP9]], [[COND9]]
+// IR-NEXT:    br i1 [[CMP10]], label [[FOR_BODY:%.*]], label [[FOR_END:%.*]]
+// IR:       for.body:
+// IR-NEXT:    store i32 7, i32* [[I]], align 4
+// IR-NEXT:    [[TMP13:%.*]] = load i32, i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    [[MUL11:%.*]] = mul nsw i32 [[TMP13]], 3
+// IR-NEXT:    [[ADD12:%.*]] = add nsw i32 7, [[MUL11]]
+// IR-NEXT:    store i32 [[ADD12]], i32* [[I]], align 4
+// IR-NEXT:    br label [[FOR_INC:%.*]]
+// IR:       for.inc:
+// IR-NEXT:    [[TMP14:%.*]] = load i32, i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP14]], 1
+// IR-NEXT:    store i32 [[INC]], i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    br label [[FOR_COND]]
+// IR:       for.end:
+// IR-NEXT:    br label [[OMP_BODY_CONTINUE:%.*]]
+// IR:       omp.body.continue:
+// IR-NEXT:    br label [[OMP_INNER_FOR_INC:%.*]]
+// IR:       omp.inner.for.inc:
+// IR-NEXT:    [[TMP15:%.*]] = load i32, i32* [[DOTOMP_IV]], align 4
+// IR-NEXT:    [[ADD13:%.*]] = add nsw i32 [[TMP15]], 1
+// IR-NEXT:    store i32 [[ADD13]], i32* [[DOTOMP_IV]], align 4
+// IR-NEXT:    br label [[OMP_INNER_FOR_COND]]
+// IR:       omp.inner.for.end:
+// IR-NEXT:    br label [[OMP_LOOP_EXIT:%.*]]
+// IR:       omp.loop.exit:
+// IR-NEXT:    call void @__kmpc_for_static_fini(%struct.ident_t* @0, i32 [[TMP1]])
+// IR-NEXT:    ret void
+//
+extern "C" void foo6() {
+#pragma omp parallel for
+#pragma omp tile sizes(5)
+  for (int i = 7; i < 17; i += 3)
+    ;
+}
+
+
+// IR-LABEL: define linkonce_odr void @_Z4foo7IiLi3ELi5EEvT_S0_(
+// IR-NEXT:  entry:
+// IR-NEXT:    [[START_ADDR:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[END_ADDR:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTFLOOR_0_IV_I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTTILE_0_IV_I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[I:%.*]] = alloca i32, align 4
+// IR-NEXT:    store i32 [[START:%.*]], i32* [[START_ADDR]], align 4
+// IR-NEXT:    store i32 [[END:%.*]], i32* [[END_ADDR]], align 4
+// IR-NEXT:    store i32 0, i32* [[DOTFLOOR_0_IV_I]], align 4
+// IR-NEXT:    br label [[FOR_COND:%.*]]
+// IR:       for.cond:
+// IR-NEXT:    [[TMP0:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I]], align 4
+// IR-NEXT:    [[TMP1:%.*]] = load i32, i32* [[END_ADDR]], align 4
+// IR-NEXT:    [[TMP2:%.*]] = load i32, i32* [[START_ADDR]], align 4
+// IR-NEXT:    [[SUB:%.*]] = sub nsw i32 [[TMP1]], [[TMP2]]
+// IR-NEXT:    [[SUB1:%.*]] = sub nsw i32 [[SUB]], 1
+// IR-NEXT:    [[ADD:%.*]] = add nsw i32 [[SUB1]], 3
+// IR-NEXT:    [[DIV:%.*]] = sdiv i32 [[ADD]], 3
+// IR-NEXT:    [[SUB2:%.*]] = sub nsw i32 [[DIV]], 1
+// IR-NEXT:    [[ADD3:%.*]] = add nsw i32 [[SUB2]], 1
+// IR-NEXT:    [[CMP:%.*]] = icmp slt i32 [[TMP0]], [[ADD3]]
+// IR-NEXT:    br i1 [[CMP]], label [[FOR_BODY:%.*]], label [[FOR_END25:%.*]]
+// IR:       for.body:
+// IR-NEXT:    [[TMP3:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I]], align 4
+// IR-NEXT:    store i32 [[TMP3]], i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    br label [[FOR_COND4:%.*]]
+// IR:       for.cond4:
+// IR-NEXT:    [[TMP4:%.*]] = load i32, i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    [[TMP5:%.*]] = load i32, i32* [[END_ADDR]], align 4
+// IR-NEXT:    [[TMP6:%.*]] = load i32, i32* [[START_ADDR]], align 4
+// IR-NEXT:    [[SUB5:%.*]] = sub nsw i32 [[TMP5]], [[TMP6]]
+// IR-NEXT:    [[SUB6:%.*]] = sub nsw i32 [[SUB5]], 1
+// IR-NEXT:    [[ADD7:%.*]] = add nsw i32 [[SUB6]], 3
+// IR-NEXT:    [[DIV8:%.*]] = sdiv i32 [[ADD7]], 3
+// IR-NEXT:    [[SUB9:%.*]] = sub nsw i32 [[DIV8]], 1
+// IR-NEXT:    [[ADD10:%.*]] = add nsw i32 [[SUB9]], 1
+// IR-NEXT:    [[TMP7:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I]], align 4
+// IR-NEXT:    [[ADD11:%.*]] = add nsw i32 [[TMP7]], 5
+// IR-NEXT:    [[CMP12:%.*]] = icmp slt i32 [[ADD10]], [[ADD11]]
+// IR-NEXT:    br i1 [[CMP12]], label [[COND_TRUE:%.*]], label [[COND_FALSE:%.*]]
+// IR:       cond.true:
+// IR-NEXT:    [[TMP8:%.*]] = load i32, i32* [[END_ADDR]], align 4
+// IR-NEXT:    [[TMP9:%.*]] = load i32, i32* [[START_ADDR]], align 4
+// IR-NEXT:    [[SUB13:%.*]] = sub nsw i32 [[TMP8]], [[TMP9]]
+// IR-NEXT:    [[SUB14:%.*]] = sub nsw i32 [[SUB13]], 1
+// IR-NEXT:    [[ADD15:%.*]] = add nsw i32 [[SUB14]], 3
+// IR-NEXT:    [[DIV16:%.*]] = sdiv i32 [[ADD15]], 3
+// IR-NEXT:    [[SUB17:%.*]] = sub nsw i32 [[DIV16]], 1
+// IR-NEXT:    [[ADD18:%.*]] = add nsw i32 [[SUB17]], 1
+// IR-NEXT:    br label [[COND_END:%.*]]
+// IR:       cond.false:
+// IR-NEXT:    [[TMP10:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I]], align 4
+// IR-NEXT:    [[ADD19:%.*]] = add nsw i32 [[TMP10]], 5
+// IR-NEXT:    br label [[COND_END]]
+// IR:       cond.end:
+// IR-NEXT:    [[COND:%.*]] = phi i32 [ [[ADD18]], [[COND_TRUE]] ], [ [[ADD19]], [[COND_FALSE]] ]
+// IR-NEXT:    [[CMP20:%.*]] = icmp slt i32 [[TMP4]], [[COND]]
+// IR-NEXT:    br i1 [[CMP20]], label [[FOR_BODY21:%.*]], label [[FOR_END:%.*]]
+// IR:       for.body21:
+// IR-NEXT:    [[TMP11:%.*]] = load i32, i32* [[START_ADDR]], align 4
+// IR-NEXT:    store i32 [[TMP11]], i32* [[I]], align 4
+// IR-NEXT:    [[TMP13:%.*]] = load i32, i32* [[START_ADDR]], align 4
+// IR-NEXT:    [[TMP14:%.*]] = load i32, i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    [[MUL:%.*]] = mul nsw i32 [[TMP14]], 3
+// IR-NEXT:    [[ADD22:%.*]] = add nsw i32 [[TMP13]], [[MUL]]
+// IR-NEXT:    store i32 [[ADD22]], i32* [[I]], align 4
+// IR-NEXT:    [[TMP15:%.*]] = load i32, i32* [[I]], align 4
+// IR-NEXT:    call void (...) @body(i32 [[TMP15]])
+// IR-NEXT:    br label [[FOR_INC:%.*]]
+// IR:       for.inc:
+// IR-NEXT:    [[TMP16:%.*]] = load i32, i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP16]], 1
+// IR-NEXT:    store i32 [[INC]], i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    br label [[FOR_COND4]]
+// IR:       for.end:
+// IR-NEXT:    br label [[FOR_INC23:%.*]]
+// IR:       for.inc23:
+// IR-NEXT:    [[TMP17:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I]], align 4
+// IR-NEXT:    [[ADD24:%.*]] = add nsw i32 [[TMP17]], 5
+// IR-NEXT:    store i32 [[ADD24]], i32* [[DOTFLOOR_0_IV_I]], align 4
+// IR-NEXT:    br label [[FOR_COND]]
+// IR:       for.end25:
+// IR-NEXT:    ret void
+//
+template<typename T, T Step, T Tile>
+void foo7(T start, T end) {
+#pragma omp tile sizes(Tile)
+  for (T i = start; i < end; i += Step)
+    body(i);
+}
+extern "C" void tfoo7() {
+  foo7<int,3,5>(0, 42);
+}
+
+
+#endif /* HEADER */
+

clang/test/OpenMP/tile_messages.cpp

+// RUN: %clang_cc1 -triple x86_64-pc-linux-gnu -std=c++17 -fopenmp -fopenmp-version=51 -fsyntax-only -Wuninitialized -verify %s
+
+void func() {
+
+  // expected-error@+1 {{expected '('}}
+  #pragma omp tile sizes
+    ;
+
+  // expected-error@+2 {{expected expression}}
+  // expected-error@+1 {{expected ')'}} expected-note@+1 {{to match this '('}}
+  #pragma omp tile  sizes(
+    ;
+
+  // expected-error@+1 {{expected expression}}
+  #pragma omp tile sizes()
+    ;
+
+  // expected-error@+1 {{expected ')'}} expected-note@+1 {{to match this '('}}
+  #pragma omp tile sizes(5
+    for (int i = 0; i < 7; ++i);
+
+  // expected-error@+2 {{expected expression}}
+  // expected-error@+1 {{expected ')'}} expected-note@+1 {{to match this '('}}
+  #pragma omp tile sizes(5,
+    ;
+
+  // expected-error@+1 {{expected expression}}
+  #pragma omp tile sizes(5,)
+    ;
+
+  // expected-error@+2 {{expected expression}}
+  // expected-error@+1 {{expected ')'}} expected-note@+1 {{to match this '('}}
+  #pragma omp tile sizes(5+
+    ;
+
+  // expected-error@+1 {{expected expression}}
+  #pragma omp tile sizes(5+)
+    ;
+
+  // expected-error@+1 {{expected expression}}
+  #pragma omp tile sizes(for)
+    ;
+
+  // expected-error@+1 {{argument to 'sizes' clause must be a strictly positive integer value}}
+  #pragma omp tile sizes(0)
+    ;
+
+  // expected-error@+4 {{expression is not an integral constant expression}}
+  // expected-note@+3 {{read of non-const variable 'a' is not allowed in a constant expression}}
+  // expected-note@+1 {{declared here}}
+  int a;
+  #pragma omp tile sizes(a)
+    ;
+
+  // expected-warning@+2 {{extra tokens at the end of '#pragma omp tile' are ignored}}
+  // expected-error@+1 {{directive '#pragma omp tile' requires the 'sizes' clause}}
+  #pragma omp tile foo
+    ;
+
+  // expected-error@+1 {{directive '#pragma omp tile' cannot contain more than one 'sizes' clause}}
+  #pragma omp tile sizes(5) sizes(5)
+  for (int i = 0; i < 7; ++i)
+    ;
+
+  // expected-error@+1 {{unexpected OpenMP clause 'collapse' in directive '#pragma omp tile'}}
+  #pragma omp tile sizes(5) collapse(2)
+  for (int i = 0; i < 7; ++i)
+    ;
+
+  {
+    // expected-error@+2 {{expected statement}}
+    #pragma omp tile sizes(5)
+  }
+
+  // expected-error@+2 {{statement after '#pragma omp tile' must be a for loop}}
+  #pragma omp tile sizes(5)
+  int b = 0;
+
+  // expected-error@+3 {{statement after '#pragma omp tile' must be a for loop}}
+  #pragma omp tile sizes(5,5)
+  for (int i = 0; i < 7; ++i)
+    ;
+
+  // expected-error@+2 {{statement after '#pragma omp tile' must be a for loop}}
+  #pragma omp tile sizes(5,5)
+  for (int i = 0; i < 7; ++i) {
+    int k = 3;
+    for (int j = 0; j < 7; ++j)
+      ;
+  }
+
+  // expected-error@+3 {{expected loop invariant expression}}
+  #pragma omp tile sizes(5,5)
+  for (int i = 0; i < 7; ++i)
+    for (int j = i; j < 7; ++j)
+      ;
+
+  // expected-error@+3 {{expected loop invariant expression}}
+  #pragma omp tile sizes(5,5)
+  for (int i = 0; i < 7; ++i)
+    for (int j = 0; j < i; ++j)
+      ;
+
+  // expected-error@+3 {{expected loop invariant expression}}
+  #pragma omp tile sizes(5,5)
+  for (int i = 0; i < 7; ++i)
+    for (int j = 0; j < i; ++j)
+      ;
+
+  // expected-error@+4 {{expected 3 for loops after '#pragma omp for', but found only 2}}
+  // expected-note@+1 {{as specified in 'collapse' clause}}
+  #pragma omp for collapse(3)
+  #pragma omp tile sizes(5)
+  for (int i = 0; i < 7; ++i)
+    ;
+
+  // expected-error@+2 {{statement after '#pragma omp tile' must be a for loop}}
+  #pragma omp tile sizes(5)
+  #pragma omp for
+  for (int i = 0; i < 7; ++i)
+    ;
+
+  // expected-error@+2 {{condition of OpenMP for loop must be a relational comparison ('<', '<=', '>', '>=', or '!=') of loop variable 'i'}}
+  #pragma omp tile sizes(5)
+  for (int i = 0; i/3<7; ++i)
+    ;
+}
+

clang/tools/libclang/CIndex.cpp

public:
   void VisitOMPTargetTeamsDistributeDirective(
       const OMPTargetTeamsDistributeDirective *D);
   void VisitOMPTargetTeamsDistributeParallelForDirective(
       const OMPTargetTeamsDistributeParallelForDirective *D);
   void VisitOMPTargetTeamsDistributeParallelForSimdDirective(
       const OMPTargetTeamsDistributeParallelForSimdDirective *D);
   void VisitOMPTargetTeamsDistributeSimdDirective(
       const OMPTargetTeamsDistributeSimdDirective *D);
+  void VisitOMPTileDirective(const OMPTileDirective *D);
 
 private:
   void AddDeclarationNameInfo(const Stmt *S);
   void AddNestedNameSpecifierLoc(NestedNameSpecifierLoc Qualifier);
   void AddExplicitTemplateArgs(const TemplateArgumentLoc *A,
                                unsigned NumTemplateArgs);
   void AddMemberRef(const FieldDecl *D, SourceLocation L);
   void AddStmt(const Stmt *S);
void OMPClauseEnqueue::VisitOMPNontemporalClause(
     const OMPNontemporalClause *C) {
   VisitOMPClauseList(C);
   for (const auto *E : C->private_refs())
     Visitor->AddStmt(E);
 }
 void OMPClauseEnqueue::VisitOMPOrderClause(const OMPOrderClause *C) {}
 } // namespace
 
+void OMPClauseEnqueue::VisitOMPSizesClause(const OMPSizesClause *C) {
+  for (auto E : C->getSizesRefs())
+    Visitor->AddStmt(E);
+}
+
 void EnqueueVisitor::EnqueueChildren(const OMPClause *S) {
   unsigned size = WL.size();
   OMPClauseEnqueue Visitor(this);
   Visitor.Visit(S);
   if (size == WL.size())
     return;
   // Now reverse the entries we just added.  This will match the DFS
   // ordering performed by the worklist.
void EnqueueVisitor::VisitOMPTargetTeamsDistributeParallelForSimdDirective(
   VisitOMPLoopDirective(D);
 }
 
 void EnqueueVisitor::VisitOMPTargetTeamsDistributeSimdDirective(
     const OMPTargetTeamsDistributeSimdDirective *D) {
   VisitOMPLoopDirective(D);
 }
 
+void EnqueueVisitor::VisitOMPTileDirective(const OMPTileDirective *D) {
+  VisitOMPLoopDirective(D);
+}
+
 void CursorVisitor::EnqueueWorkList(VisitorWorkList &WL, const Stmt *S) {
   EnqueueVisitor(WL, MakeCXCursor(S, StmtParent, TU, RegionOfInterest))
       .Visit(S);
 }
 
 bool CursorVisitor::IsInRegionOfInterest(CXCursor C) {
   if (RegionOfInterest.isValid()) {
     SourceRange Range = getRawCursorExtent(C);
case CXCursor_OMPTargetTeamsDirective:
     return cxstring::createRef("OMPTargetTeamsDirective");
   case CXCursor_OMPTargetTeamsDistributeDirective:
     return cxstring::createRef("OMPTargetTeamsDistributeDirective");
   case CXCursor_OMPTargetTeamsDistributeParallelForDirective:
     return cxstring::createRef("OMPTargetTeamsDistributeParallelForDirective");
   case CXCursor_OMPTargetTeamsDistributeParallelForSimdDirective:
     return cxstring::createRef(
         "OMPTargetTeamsDistributeParallelForSimdDirective");
+  case CXCursor_OMPTileDirective:
+    return cxstring::createRef("OMPTileDirective");
   case CXCursor_OMPTargetTeamsDistributeSimdDirective:
     return cxstring::createRef("OMPTargetTeamsDistributeSimdDirective");
   case CXCursor_OverloadCandidate:
     return cxstring::createRef("OverloadCandidate");
   case CXCursor_TypeAliasTemplateDecl:
     return cxstring::createRef("TypeAliasTemplateDecl");
   case CXCursor_StaticAssert:
     return cxstring::createRef("StaticAssert");

clang/tools/libclang/CXCursor.cpp

case Stmt::OMPTargetTeamsDistributeParallelForDirectiveClass:
     K = CXCursor_OMPTargetTeamsDistributeParallelForDirective;
     break;
   case Stmt::OMPTargetTeamsDistributeParallelForSimdDirectiveClass:
     K = CXCursor_OMPTargetTeamsDistributeParallelForSimdDirective;
     break;
   case Stmt::OMPTargetTeamsDistributeSimdDirectiveClass:
     K = CXCursor_OMPTargetTeamsDistributeSimdDirective;
     break;
+  case Stmt::OMPTileDirectiveClass:
+    K = CXCursor_OMPTileDirective;
+    break;
   case Stmt::BuiltinBitCastExprClass:
     K = CXCursor_BuiltinBitCastExpr;
   }
 
   CXCursor C = {K, 0, {Parent, S, TU}};
   return C;
 }
 

llvm/include/llvm/Frontend/OpenMP/OMPKinds.def

 #define __OMP_DIRECTIVE_EXT(Name, Str) OMP_DIRECTIVE(OMPD_##Name, Str)
 #define __OMP_DIRECTIVE(Name) __OMP_DIRECTIVE_EXT(Name, #Name)
 
 __OMP_DIRECTIVE(threadprivate)
 __OMP_DIRECTIVE(parallel)
 __OMP_DIRECTIVE(task)
 __OMP_DIRECTIVE(simd)
 __OMP_DIRECTIVE(for)
+__OMP_DIRECTIVE(tile)
 __OMP_DIRECTIVE(sections)
 __OMP_DIRECTIVE(section)
 __OMP_DIRECTIVE(single)
 __OMP_DIRECTIVE(master)
 __OMP_DIRECTIVE(critical)
 __OMP_DIRECTIVE(taskyield)
 __OMP_DIRECTIVE(barrier)
 __OMP_DIRECTIVE(taskwait)