[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

4751–4754

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

4751–4754

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

4751–4754

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

4751–4754

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

4751–4754

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

4751–4754

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

4751–4754

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

2790–2791

[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

443

Please add a documentation what this class represents.

448

Inaccurate for the tile directive.

456

Inaccurate for the tile directive.

651–677

Please add doxygen comments.

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

659

I do not see why making this a forwarding reference.

680–709

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

448

We can rename it to something like NumAssociatedLoops

651–677

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.

659

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

680–709

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

448

Sounds good.

651–677

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.

659

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

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

680–709

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

651–677

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

659

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

651–677

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
  
  ...

659

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

651–677

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.

659

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() == llvm::omp::OMPC_simdlen;
   }
 };
 
+/// 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(llvm::omp::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() == llvm::omp::OMPC_sizes;
+  }
+};
+
 /// This represents 'collapse' clause in the '#pragma omp ...'
 /// directive.
 ///
 /// \code
 /// #pragma omp simd collapse(3)
 /// \endcode
 /// In this example directive '#pragma omp simd' has clause 'collapse'
 /// with single expression '3'.

clang/include/clang/AST/RecursiveASTVisitor.h

 }
 
 DEF_TRAVERSE_STMT(OMPParallelDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
 DEF_TRAVERSE_STMT(OMPSimdDirective,
                   { TRY_TO(TraverseOMPExecutableDirective(S)); })
 
+DEF_TRAVERSE_STMT(OMPTileDirective,
+                  { 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)); })

 
 template <typename Derived>
 bool RecursiveASTVisitor<Derived>::VisitOMPSimdlenClause(OMPSimdlenClause *C) {
   TRY_TO(TraverseStmt(C->getSimdlen()));
   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>::VisitOMPCollapseClause(OMPCollapseClause *C) {
   TRY_TO(TraverseStmt(C->getNumForLoops()));
   return true;
 }
 
 template <typename Derived>
 bool RecursiveASTVisitor<Derived>::VisitOMPDefaultClause(OMPDefaultClause *) {

clang/include/clang/AST/StmtOpenMP.h

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;
   }
 };
 
+class OMPLoopBasedDirective : public OMPExecutableDirective {
+  friend class ASTStmtReader;
+
+protected:
+  /// Number of collapsed loops as specified by 'collapse' clause.
+  unsigned CollapsedNum = 0;
+
+  /// Build instance of loop directive of class \a Kind.
+  ///
+  /// \param SC Statement class.
+  /// \param Kind Kind of OpenMP directive.
+  /// \param StartLoc Starting location of the directive (directive keyword).
+  /// \param EndLoc Ending location of the directive.
+  /// \param CollapsedNum Number of collapsed loops from 'collapse' clause.
+  ///
+  OMPLoopBasedDirective(StmtClass SC, OpenMPDirectiveKind Kind,
+                        SourceLocation StartLoc, SourceLocation EndLoc,
+                        unsigned CollapsedNum)
+      : OMPExecutableDirective(SC, Kind, StartLoc, EndLoc),
+        CollapsedNum(CollapsedNum) {}
+
+public:
+  /// The expressions built to support OpenMP loops in combined/composite
+  /// pragmas (e.g. pragma omp distribute parallel for)
+  struct DistCombinedHelperExprs {
+    /// DistributeLowerBound - used when composing 'omp distribute' with
+    /// 'omp for' in a same construct.
+    Expr *LB;
+    /// DistributeUpperBound - used when composing 'omp distribute' with
+    /// 'omp for' in a same construct.
+    Expr *UB;
+    /// DistributeEnsureUpperBound - used when composing 'omp distribute'
+    ///  with 'omp for' in a same construct, EUB depends on DistUB
+    Expr *EUB;
+    /// Distribute loop iteration variable init used when composing 'omp
+    /// distribute'
+    ///  with 'omp for' in a same construct
+    Expr *Init;
+    /// Distribute Loop condition used when composing 'omp distribute'
+    ///  with 'omp for' in a same construct
+    Expr *Cond;
+    /// Update of LowerBound for statically scheduled omp loops for
+    /// outer loop in combined constructs (e.g. 'distribute parallel for')
+    Expr *NLB;
+    /// Update of UpperBound for statically scheduled omp loops for
+    /// outer loop in combined constructs (e.g. 'distribute parallel for')
+    Expr *NUB;
+    /// Distribute Loop condition used when composing 'omp distribute'
+    ///  with 'omp for' in a same construct when schedule is chunked.
+    Expr *DistCond;
+    /// 'omp parallel for' loop condition used when composed with
+    /// 'omp distribute' in the same construct and when schedule is
+    /// chunked and the chunk size is 1.
+    Expr *ParForInDistCond;
+  };
+
+  /// The expressions built for the OpenMP loop CodeGen for the
+  /// whole collapsed loop nest.
+  struct HelperExprs {
+    /// Loop iteration variable.
+    Expr *IterationVarRef;
+    /// Loop last iteration number.
+    Expr *LastIteration;
+    /// Loop number of iterations.
+    Expr *NumIterations;
+    /// Calculation of last iteration.
+    Expr *CalcLastIteration;
+    /// Loop pre-condition.
+    Expr *PreCond;
+    /// Loop condition.
+    Expr *Cond;
+    /// Loop iteration variable init.
+    Expr *Init;
+    /// Loop increment.
+    Expr *Inc;
+    /// IsLastIteration - local flag variable passed to runtime.
+    Expr *IL;
+    /// LowerBound - local variable passed to runtime.
+    Expr *LB;
+    /// UpperBound - local variable passed to runtime.
+    Expr *UB;
+    /// Stride - local variable passed to runtime.
+    Expr *ST;
+    /// EnsureUpperBound -- expression UB = min(UB, NumIterations).
+    Expr *EUB;
+    /// Update of LowerBound for statically scheduled 'omp for' loops.
+    Expr *NLB;
+    /// Update of UpperBound for statically scheduled 'omp for' loops.
+    Expr *NUB;
+    /// PreviousLowerBound - local variable passed to runtime in the
+    /// enclosing schedule or null if that does not apply.
+    Expr *PrevLB;
+    /// PreviousUpperBound - local variable passed to runtime in the
+    /// enclosing schedule or null if that does not apply.
+    Expr *PrevUB;
+    /// DistInc - increment expression for distribute loop when found
+    /// combined with a further loop level (e.g. in 'distribute parallel for')
+    /// 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;
+    /// 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;
+    /// Final loop counter values for GodeGen.
+    SmallVector<Expr *, 4> Finals;
+    /// List of counters required for the generation of the non-rectangular
+    /// loops.
+    SmallVector<Expr *, 4> DependentCounters;
+    /// List of initializers required for the generation of the non-rectangular
+    /// loops.
+    SmallVector<Expr *, 4> DependentInits;
+    /// List of final conditions required for the generation of the
+    /// non-rectangular loops.
+    SmallVector<Expr *, 4> FinalsConditions;
+    /// Init statement for all captured expressions.
+    Stmt *PreInits;
+
+    /// Expressions used when combining OpenMP loop pragmas
+    DistCombinedHelperExprs DistCombinedFields;
+
+    /// Check if all the expressions are built (does not check the
+    /// worksharing ones).
+    bool builtAll() {
+      return IterationVarRef != nullptr && LastIteration != nullptr &&
+             NumIterations != nullptr && PreCond != nullptr &&
+             Cond != nullptr && Init != nullptr && Inc != nullptr;
+    }
+
+    /// Initialize all the fields to null.
+    /// \param Size Number of elements in the
+    /// counters/finals/updates/dependent_counters/dependent_inits/finals_conditions
+    /// arrays.
+    void clear(unsigned Size) {
+      IterationVarRef = nullptr;
+      LastIteration = nullptr;
+      CalcLastIteration = nullptr;
+      PreCond = nullptr;
+      Cond = nullptr;
+      Init = nullptr;
+      Inc = nullptr;
+      IL = nullptr;
+      LB = nullptr;
+      UB = nullptr;
+      ST = nullptr;
+      EUB = nullptr;
+      NLB = nullptr;
+      NUB = nullptr;
+      NumIterations = nullptr;
+      PrevLB = nullptr;
+      PrevUB = nullptr;
+      DistInc = nullptr;
+      PrevEUB = nullptr;
+      Counters.resize(Size);
+      PrivateCounters.resize(Size);
+      Inits.resize(Size);
+      Updates.resize(Size);
+      Finals.resize(Size);
+      DependentCounters.resize(Size);
+      DependentInits.resize(Size);
+      FinalsConditions.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;
+      }
+      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; }
+
+  /// 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) {
+    return tryToFindNextInnerLoop(const_cast<Stmt *>(CurStmt),
+                                  TryImperfectlyNestedLoops);
+  }
+
+  static bool
+  doForAllLoops(Stmt *CurStmt, bool TryImperfectlyNestedLoops,
+                unsigned NumLoops,
+                llvm::function_ref<bool(unsigned, Stmt *)> Callback);
+  static bool
+  doForAllLoops(const Stmt *CurStmt, bool TryImperfectlyNestedLoops,
+                unsigned NumLoops,
+                llvm::function_ref<bool(unsigned, const Stmt *)> Callback) {
+    auto &&NewCallback = [Callback](unsigned Cnt, Stmt *CurStmt) {
+      return Callback(Cnt, CurStmt);
+    };
+    return doForAllLoops(const_cast<Stmt *>(CurStmt), TryImperfectlyNestedLoops,
+                         NumLoops, NewCallback);
+  }
+
+  static void doForAllLoopsBodies(
+      Stmt *CurStmt, bool TryImperfectlyNestedLoops, unsigned NumLoops,
+      llvm::function_ref<void(unsigned, Stmt *, Stmt *)> Callback);
+  static void doForAllLoopsBodies(
+      const Stmt *CurStmt, bool TryImperfectlyNestedLoops, unsigned NumLoops,
+      llvm::function_ref<void(unsigned, const Stmt *, const Stmt *)> Callback) {
+    auto &&NewCallback = [Callback](unsigned Cnt, Stmt *Loop, Stmt *Body) {
+      Callback(Cnt, Loop, Body);
+    };
+    doForAllLoopsBodies(const_cast<Stmt *>(CurStmt), TryImperfectlyNestedLoops,
+                        NumLoops, NewCallback);
+  }
+
+  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 ||
+           T->getStmtClass() == OMPParallelMasterTaskLoopSimdDirectiveClass ||
+           T->getStmtClass() == OMPDistributeDirectiveClass ||
+           T->getStmtClass() == OMPTargetParallelForDirectiveClass ||
+           T->getStmtClass() == OMPDistributeParallelForDirectiveClass ||
+           T->getStmtClass() == OMPDistributeParallelForSimdDirectiveClass ||
+           T->getStmtClass() == OMPDistributeSimdDirectiveClass ||
+           T->getStmtClass() == OMPTargetParallelForSimdDirectiveClass ||
+           T->getStmtClass() == OMPTargetSimdDirectiveClass ||
+           T->getStmtClass() == OMPTeamsDistributeDirectiveClass ||
+           T->getStmtClass() == OMPTeamsDistributeSimdDirectiveClass ||
+           T->getStmtClass() ==
+               OMPTeamsDistributeParallelForSimdDirectiveClass ||
+           T->getStmtClass() == OMPTeamsDistributeParallelForDirectiveClass ||
+           T->getStmtClass() ==
+               OMPTargetTeamsDistributeParallelForDirectiveClass ||
+           T->getStmtClass() ==
+               OMPTargetTeamsDistributeParallelForSimdDirectiveClass ||
+           T->getStmtClass() == OMPTargetTeamsDistributeDirectiveClass ||
+           T->getStmtClass() == OMPTargetTeamsDistributeSimdDirectiveClass;
+  }
+};
+
 /// 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 {
+class OMPLoopDirective : public OMPLoopBasedDirective {
   friend class ASTStmtReader;
-  /// Number of collapsed loops as specified by 'collapse' clause.
-  unsigned CollapsedNum = 0;
 
   /// 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
enum {
     // arrays) for combined distribute loop directives.
     CombinedDistributeEnd = 29,
   };
 
   /// Get the counters storage.
   MutableArrayRef<Expr *> getCounters() {
     auto **Storage = reinterpret_cast<Expr **>(
         &Data->getChildren()[getArraysOffset(getDirectiveKind())]);
-    return llvm::makeMutableArrayRef(Storage, CollapsedNum);
+    return llvm::makeMutableArrayRef(Storage, getCollapsedNumber());
   }
 
   /// Get the private counters storage.
   MutableArrayRef<Expr *> getPrivateCounters() {
     auto **Storage = reinterpret_cast<Expr **>(
         &Data->getChildren()[getArraysOffset(getDirectiveKind()) +
-                             CollapsedNum]);
-    return llvm::makeMutableArrayRef(Storage, CollapsedNum);
+                             getCollapsedNumber()]);
+    return llvm::makeMutableArrayRef(Storage, getCollapsedNumber());
   }
 
   /// Get the updates storage.
   MutableArrayRef<Expr *> getInits() {
     auto **Storage = reinterpret_cast<Expr **>(
         &Data->getChildren()[getArraysOffset(getDirectiveKind()) +
-                             2 * CollapsedNum]);
-    return llvm::makeMutableArrayRef(Storage, CollapsedNum);
+                             2 * getCollapsedNumber()]);
+    return llvm::makeMutableArrayRef(Storage, getCollapsedNumber());
   }
 
   /// Get the updates storage.
   MutableArrayRef<Expr *> getUpdates() {
     auto **Storage = reinterpret_cast<Expr **>(
         &Data->getChildren()[getArraysOffset(getDirectiveKind()) +
-                             3 * CollapsedNum]);
-    return llvm::makeMutableArrayRef(Storage, CollapsedNum);
+                             3 * getCollapsedNumber()]);
+    return llvm::makeMutableArrayRef(Storage, getCollapsedNumber());
   }
 
   /// Get the final counter updates storage.
   MutableArrayRef<Expr *> getFinals() {
     auto **Storage = reinterpret_cast<Expr **>(
         &Data->getChildren()[getArraysOffset(getDirectiveKind()) +
-                             4 * CollapsedNum]);
-    return llvm::makeMutableArrayRef(Storage, CollapsedNum);
+                             4 * getCollapsedNumber()]);
+    return llvm::makeMutableArrayRef(Storage, getCollapsedNumber());
   }
 
   /// Get the dependent counters storage.
   MutableArrayRef<Expr *> getDependentCounters() {
     auto **Storage = reinterpret_cast<Expr **>(
         &Data->getChildren()[getArraysOffset(getDirectiveKind()) +
-                             5 * CollapsedNum]);
-    return llvm::makeMutableArrayRef(Storage, CollapsedNum);
+                             5 * getCollapsedNumber()]);
+    return llvm::makeMutableArrayRef(Storage, getCollapsedNumber());
   }
 
   /// Get the dependent inits storage.
   MutableArrayRef<Expr *> getDependentInits() {
     auto **Storage = reinterpret_cast<Expr **>(
         &Data->getChildren()[getArraysOffset(getDirectiveKind()) +
-                             6 * CollapsedNum]);
-    return llvm::makeMutableArrayRef(Storage, CollapsedNum);
+                             6 * getCollapsedNumber()]);
+    return llvm::makeMutableArrayRef(Storage, getCollapsedNumber());
   }
 
   /// Get the finals conditions storage.
   MutableArrayRef<Expr *> getFinalsConditions() {
     auto **Storage = reinterpret_cast<Expr **>(
         &Data->getChildren()[getArraysOffset(getDirectiveKind()) +
-                             7 * CollapsedNum]);
-    return llvm::makeMutableArrayRef(Storage, CollapsedNum);
+                             7 * getCollapsedNumber()]);
+    return llvm::makeMutableArrayRef(Storage, getCollapsedNumber());
   }
 
 protected:
   /// Build instance of loop directive of class \a Kind.
   ///
   /// \param SC Statement class.
   /// \param Kind Kind of OpenMP directive.
   /// \param StartLoc Starting location of the directive (directive keyword).
   /// \param EndLoc Ending location of the directive.
   /// \param CollapsedNum Number of collapsed loops from 'collapse' clause.
   ///
   OMPLoopDirective(StmtClass SC, OpenMPDirectiveKind Kind,
                    SourceLocation StartLoc, SourceLocation EndLoc,
                    unsigned CollapsedNum)
-      : OMPExecutableDirective(SC, Kind, StartLoc, EndLoc),
-        CollapsedNum(CollapsedNum) {}
+      : OMPLoopBasedDirective(SC, Kind, StartLoc, EndLoc, CollapsedNum) {}
 
   /// Offset to the start of children expression arrays.
   static unsigned getArraysOffset(OpenMPDirectiveKind Kind) {
     if (isOpenMPLoopBoundSharingDirective(Kind))
       return CombinedDistributeEnd;
     if (isOpenMPWorksharingDirective(Kind) || isOpenMPTaskLoopDirective(Kind) ||
         isOpenMPDistributeDirective(Kind))
       return WorksharingEnd;
protected:
   void setInits(ArrayRef<Expr *> A);
   void setUpdates(ArrayRef<Expr *> A);
   void setFinals(ArrayRef<Expr *> A);
   void setDependentCounters(ArrayRef<Expr *> A);
   void setDependentInits(ArrayRef<Expr *> A);
   void setFinalsConditions(ArrayRef<Expr *> A);
 
 public:
-  /// The expressions built to support OpenMP loops in combined/composite
-  /// pragmas (e.g. pragma omp distribute parallel for)
-  struct DistCombinedHelperExprs {
-    /// DistributeLowerBound - used when composing 'omp distribute' with
-    /// 'omp for' in a same construct.
-    Expr *LB;
-    /// DistributeUpperBound - used when composing 'omp distribute' with
-    /// 'omp for' in a same construct.
-    Expr *UB;
-    /// DistributeEnsureUpperBound - used when composing 'omp distribute'
-    ///  with 'omp for' in a same construct, EUB depends on DistUB
-    Expr *EUB;
-    /// Distribute loop iteration variable init used when composing 'omp
-    /// distribute'
-    ///  with 'omp for' in a same construct
-    Expr *Init;
-    /// Distribute Loop condition used when composing 'omp distribute'
-    ///  with 'omp for' in a same construct
-    Expr *Cond;
-    /// Update of LowerBound for statically scheduled omp loops for
-    /// outer loop in combined constructs (e.g. 'distribute parallel for')
-    Expr *NLB;
-    /// Update of UpperBound for statically scheduled omp loops for
-    /// outer loop in combined constructs (e.g. 'distribute parallel for')
-    Expr *NUB;
-    /// Distribute Loop condition used when composing 'omp distribute'
-    ///  with 'omp for' in a same construct when schedule is chunked.
-    Expr *DistCond;
-    /// 'omp parallel for' loop condition used when composed with
-    /// 'omp distribute' in the same construct and when schedule is
-    /// chunked and the chunk size is 1.
-    Expr *ParForInDistCond;
-  };
-
-  /// The expressions built for the OpenMP loop CodeGen for the
-  /// whole collapsed loop nest.
-  struct HelperExprs {
-    /// Loop iteration variable.
-    Expr *IterationVarRef;
-    /// Loop last iteration number.
-    Expr *LastIteration;
-    /// Loop number of iterations.
-    Expr *NumIterations;
-    /// Calculation of last iteration.
-    Expr *CalcLastIteration;
-    /// Loop pre-condition.
-    Expr *PreCond;
-    /// Loop condition.
-    Expr *Cond;
-    /// Loop iteration variable init.
-    Expr *Init;
-    /// Loop increment.
-    Expr *Inc;
-    /// IsLastIteration - local flag variable passed to runtime.
-    Expr *IL;
-    /// LowerBound - local variable passed to runtime.
-    Expr *LB;
-    /// UpperBound - local variable passed to runtime.
-    Expr *UB;
-    /// Stride - local variable passed to runtime.
-    Expr *ST;
-    /// EnsureUpperBound -- expression UB = min(UB, NumIterations).
-    Expr *EUB;
-    /// Update of LowerBound for statically scheduled 'omp for' loops.
-    Expr *NLB;
-    /// Update of UpperBound for statically scheduled 'omp for' loops.
-    Expr *NUB;
-    /// PreviousLowerBound - local variable passed to runtime in the
-    /// enclosing schedule or null if that does not apply.
-    Expr *PrevLB;
-    /// PreviousUpperBound - local variable passed to runtime in the
-    /// enclosing schedule or null if that does not apply.
-    Expr *PrevUB;
-    /// DistInc - increment expression for distribute loop when found
-    /// combined with a further loop level (e.g. in 'distribute parallel for')
-    /// 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;
-    /// 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;
-    /// Final loop counter values for GodeGen.
-    SmallVector<Expr *, 4> Finals;
-    /// List of counters required for the generation of the non-rectangular
-    /// loops.
-    SmallVector<Expr *, 4> DependentCounters;
-    /// List of initializers required for the generation of the non-rectangular
-    /// loops.
-    SmallVector<Expr *, 4> DependentInits;
-    /// List of final conditions required for the generation of the
-    /// non-rectangular loops.
-    SmallVector<Expr *, 4> FinalsConditions;
-    /// Init statement for all captured expressions.
-    Stmt *PreInits;
-
-    /// Expressions used when combining OpenMP loop pragmas
-    DistCombinedHelperExprs DistCombinedFields;
-
-    /// Check if all the expressions are built (does not check the
-    /// worksharing ones).
-    bool builtAll() {
-      return IterationVarRef != nullptr && LastIteration != nullptr &&
-             NumIterations != nullptr && PreCond != nullptr &&
-             Cond != nullptr && Init != nullptr && Inc != nullptr;
-    }
-
-    /// Initialize all the fields to null.
-    /// \param Size Number of elements in the
-    /// counters/finals/updates/dependent_counters/dependent_inits/finals_conditions
-    /// arrays.
-    void clear(unsigned Size) {
-      IterationVarRef = nullptr;
-      LastIteration = nullptr;
-      CalcLastIteration = nullptr;
-      PreCond = nullptr;
-      Cond = nullptr;
-      Init = nullptr;
-      Inc = nullptr;
-      IL = nullptr;
-      LB = nullptr;
-      UB = nullptr;
-      ST = nullptr;
-      EUB = nullptr;
-      NLB = nullptr;
-      NUB = nullptr;
-      NumIterations = nullptr;
-      PrevLB = nullptr;
-      PrevUB = nullptr;
-      DistInc = nullptr;
-      PrevEUB = nullptr;
-      Counters.resize(Size);
-      PrivateCounters.resize(Size);
-      Inits.resize(Size);
-      Updates.resize(Size);
-      Finals.resize(Size);
-      DependentCounters.resize(Size);
-      DependentInits.resize(Size);
-      FinalsConditions.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;
-      }
-      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 {
     return cast<Expr>(Data->getChildren()[IterationVariableOffset]);
   }
   Expr *getLastIteration() const {
     return cast<Expr>(Data->getChildren()[LastIterationOffset]);
   }
   Expr *getCalcLastIteration() const {
     return cast<Expr>(Data->getChildren()[CalcLastIterationOffset]);
assert(isOpenMPLoopBoundSharingDirective(getDirectiveKind()) &&
            "expected loop bound distribute sharing directive");
     return cast<Expr>(Data->getChildren()[CombinedDistConditionOffset]);
   }
   Expr *getCombinedParForInDistCond() const {
     assert(isOpenMPLoopBoundSharingDirective(getDirectiveKind()) &&
            "expected loop bound distribute sharing directive");
     return cast<Expr>(Data->getChildren()[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) {
-    return tryToFindNextInnerLoop(const_cast<Stmt *>(CurStmt),
-                                  TryImperfectlyNestedLoops);
-  }
   Stmt *getBody();
   const Stmt *getBody() const {
     return const_cast<OMPLoopDirective *>(this)->getBody();
   }
 
   ArrayRef<Expr *> counters() { return getCounters(); }
 
   ArrayRef<Expr *> counters() const {
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 OMPLoopBasedDirective {
+  friend class ASTStmtReader;
+  friend class OMPExecutableDirective;
+
+  /// Default list of offsets.
+  enum {
+    PreInitsOffset = 0,
+    TransformedStmtOffset,
+  };
+
+  explicit OMPTileDirective(SourceLocation StartLoc, SourceLocation EndLoc,
+                            unsigned NumLoops)
+      : OMPLoopBasedDirective(OMPTileDirectiveClass, llvm::omp::OMPD_tile,
+                              StartLoc, EndLoc, NumLoops) {}
+
+  void setPreInits(Stmt *PreInits) {
+    Data->getChildren()[PreInitsOffset] = PreInits;
+  }
+
+  void setTransformedStmt(Stmt *S) {
+    Data->getChildren()[TransformedStmtOffset] = S;
+  }
+
+public:
+  /// Create a new AST node representation for '#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
+  ///                        dependent contexts.
+  /// \param PreInits Helper preinits statements for the loop nest.
+  static OMPTileDirective *Create(const ASTContext &C, SourceLocation StartLoc,
+                                  SourceLocation EndLoc,
+                                  ArrayRef<OMPClause *> Clauses,
+                                  unsigned NumLoops, Stmt *AssociatedStmt,
+                                  Stmt *TransformedStmt, Stmt *PreInits);
+
+  /// 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
+  ///   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 Data->getChildren()[TransformedStmtOffset];
+  }
+
+  /// Return preinits statement.
+  Stmt *getPreInits() const { return Data->getChildren()[PreInitsOffset]; }
+
+  static bool classof(const Stmt *T) {
+    return T->getStmtClass() == OMPTileDirectiveClass;
+  }
+};
+
 /// 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_device_unsupported_type : Error <
   "%0 requires %1 bit size %2 type support, but device '%3' 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

 /// 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/StmtNodes.td

 def SEHLeaveStmt : StmtNode<Stmt>;
 def MSDependentExistsStmt : StmtNode<Stmt>;
 
 // OpenCL Extensions.
 def AsTypeExpr : StmtNode<Expr>;
 
 // OpenMP Directives.
 def OMPExecutableDirective : StmtNode<Stmt, 1>;
-def OMPLoopDirective : StmtNode<OMPExecutableDirective, 1>;
+def OMPLoopBasedDirective : StmtNode<OMPExecutableDirective, 1>;
+def OMPLoopDirective : StmtNode<OMPLoopBasedDirective, 1>;
 def OMPParallelDirective : StmtNode<OMPExecutableDirective>;
 def OMPSimdDirective : StmtNode<OMPLoopDirective>;
+def OMPTileDirective : StmtNode<OMPLoopBasedDirective>;
 def OMPForDirective : StmtNode<OMPLoopDirective>;
 def OMPForSimdDirective : StmtNode<OMPLoopDirective>;
 def OMPSectionsDirective : StmtNode<OMPExecutableDirective>;
 def OMPSectionDirective : StmtNode<OMPExecutableDirective>;
 def OMPSingleDirective : StmtNode<OMPExecutableDirective>;
 def OMPMasterDirective : StmtNode<OMPExecutableDirective>;
 def OMPCriticalDirective : StmtNode<OMPExecutableDirective>;
 def OMPParallelForDirective : 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

public:
   using VarsWithInheritedDSAType =
       llvm::SmallDenseMap<const ValueDecl *, const Expr *, 4>;
   /// Called on well-formed '\#pragma omp simd' after parsing
   /// of the associated statement.
   StmtResult
   ActOnOpenMPSimdDirective(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);
   /// Called on well-formed '\#pragma omp for' after parsing
   /// of the associated statement.
   StmtResult
   ActOnOpenMPForDirective(ArrayRef<OMPClause *> Clauses, Stmt *AStmt,
                           SourceLocation StartLoc, SourceLocation EndLoc,
                           VarsWithInheritedDSAType &VarsWithImplicitDSA);
   /// Called on well-formed '\#pragma omp for simd' after parsing
   /// of the associated statement.
public:
   OMPClause *ActOnOpenMPSafelenClause(Expr *Length,
                                       SourceLocation StartLoc,
                                       SourceLocation LParenLoc,
                                       SourceLocation EndLoc);
   /// Called on well-formed 'simdlen' clause.
   OMPClause *ActOnOpenMPSimdlenClause(Expr *Length, SourceLocation StartLoc,
                                       SourceLocation LParenLoc,
                                       SourceLocation EndLoc);
+  /// Called on well-form 'sizes' clause.
+  OMPClause *ActOnOpenMPSizesClause(ArrayRef<Expr *> SizeExprs,
+                                    SourceLocation StartLoc,
+                                    SourceLocation LParenLoc,
+                                    SourceLocation EndLoc);
   /// Called on well-formed 'collapse' clause.
   OMPClause *ActOnOpenMPCollapseClause(Expr *NumForLoops,
                                        SourceLocation StartLoc,
                                        SourceLocation LParenLoc,
                                        SourceLocation EndLoc);
   /// Called on well-formed 'ordered' clause.
   OMPClause *
   ActOnOpenMPOrderedClause(SourceLocation StartLoc, SourceLocation EndLoc,

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_SCALAR_VALUE_INIT, // CXXScalarValueInitExpr
       EXPR_CXX_NEW,               // CXXNewExpr
       EXPR_CXX_DELETE,            // CXXDeleteExpr
       EXPR_CXX_PSEUDO_DESTRUCTOR, // CXXPseudoDestructorExpr
 
       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_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
 
       // CUDA
       EXPR_CUDA_KERNEL_CALL,       // CUDAKernelCallExpr
 
       // OpenCL
       EXPR_ASTYPE,                 // AsTypeExpr
 
       // Microsoft
       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
 
       // OpenMP directives
       STMT_OMP_PARALLEL_DIRECTIVE,
       STMT_OMP_SIMD_DIRECTIVE,
+      STMT_OMP_TILE_DIRECTIVE,
       STMT_OMP_FOR_DIRECTIVE,
       STMT_OMP_FOR_SIMD_DIRECTIVE,
       STMT_OMP_SECTIONS_DIRECTIVE,
       STMT_OMP_SECTION_DIRECTIVE,
       STMT_OMP_SINGLE_DIRECTIVE,
       STMT_OMP_MASTER_DIRECTIVE,
       STMT_OMP_CRITICAL_DIRECTIVE,
       STMT_OMP_PARALLEL_FOR_DIRECTIVE,
enum StmtCode {
       STMT_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_FOR_DIRECTIVE,
       STMT_OMP_TARGET_TEAMS_DISTRIBUTE_PARALLEL_FOR_SIMD_DIRECTIVE,
       STMT_OMP_TARGET_TEAMS_DISTRIBUTE_SIMD_DIRECTIVE,
       EXPR_OMP_ARRAY_SECTION,
       EXPR_OMP_ARRAY_SHAPING,
       EXPR_OMP_ITERATOR,
 
       // ARC
       EXPR_OBJC_BRIDGED_CAST,     // ObjCBridgedCastExpr
 
       STMT_MS_DEPENDENT_EXISTS,   // MSDependentExistsStmt
       EXPR_LAMBDA,                // LambdaExpr
       STMT_COROUTINE_BODY,
       STMT_CORETURN,
       EXPR_COAWAIT,
       EXPR_COYIELD,
       EXPR_DEPENDENT_COAWAIT,
 
       // FixedPointLiteral

clang/lib/AST/OpenMPClause.cpp

const OMPClauseWithPreInit *OMPClauseWithPreInit::get(const OMPClause *C) {
   case OMPC_final:
     return static_cast<const OMPFinalClause *>(C);
   case OMPC_priority:
     return static_cast<const OMPPriorityClause *>(C);
   case OMPC_default:
   case OMPC_proc_bind:
   case OMPC_safelen:
   case OMPC_simdlen:
+  case OMPC_sizes:
   case OMPC_allocator:
   case OMPC_allocate:
   case OMPC_collapse:
   case OMPC_private:
   case OMPC_shared:
   case OMPC_aligned:
   case OMPC_copyin:
   case OMPC_copyprivate:
const OMPClauseWithPostUpdate *OMPClauseWithPostUpdate::get(const OMPClause *C) {
   case OMPC_firstprivate:
   case OMPC_default:
   case OMPC_proc_bind:
   case OMPC_if:
   case OMPC_final:
   case OMPC_num_threads:
   case OMPC_safelen:
   case OMPC_simdlen:
+  case OMPC_sizes:
   case OMPC_allocator:
   case OMPC_allocate:
   case OMPC_collapse:
   case OMPC_private:
   case OMPC_shared:
   case OMPC_aligned:
   case OMPC_copyin:
   case OMPC_copyprivate:

 }
 
 OMPInReductionClause *OMPInReductionClause::CreateEmpty(const ASTContext &C,
                                                         unsigned N) {
   void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(6 * N));
   return new (Mem) OMPInReductionClause(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);
+}
+
 OMPAllocateClause *
 OMPAllocateClause::Create(const ASTContext &C, SourceLocation StartLoc,
                           SourceLocation LParenLoc, Expr *Allocator,
                           SourceLocation ColonLoc, SourceLocation EndLoc,
                           ArrayRef<Expr *> VL) {
   // Allocate space for private variables and initializer expressions.
   void *Mem = C.Allocate(totalSizeToAlloc<Expr *>(VL.size()));
   auto *Clause = new (Mem) OMPAllocateClause(StartLoc, LParenLoc, Allocator,

 }
 
 void OMPClausePrinter::VisitOMPSimdlenClause(OMPSimdlenClause *Node) {
   OS << "simdlen(";
   Node->getSimdlen()->printPretty(OS, nullptr, Policy, 0);
   OS << ")";
 }
 
+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::VisitOMPAllocatorClause(OMPAllocatorClause *Node) {
   OS << "allocator(";
   Node->getAllocator()->printPretty(OS, nullptr, Policy, 0);
   OS << ")";
 }
 
 void OMPClausePrinter::VisitOMPCollapseClause(OMPCollapseClause *Node) {
   OS << "collapse(";

clang/lib/AST/StmtOpenMP.cpp

 Stmt *OMPExecutableDirective::getStructuredBlock() {
   assert(!isStandaloneDirective() &&
          "Standalone Executable Directives don't have Structured Blocks.");
   if (auto *LD = dyn_cast<OMPLoopDirective>(this))
     return LD->getBody();
   return getRawStmt();
 }
 
-Stmt *OMPLoopDirective::tryToFindNextInnerLoop(Stmt *CurStmt,
-                                               bool TryImperfectlyNestedLoops) {
+Stmt *
+OMPLoopBasedDirective::tryToFindNextInnerLoop(Stmt *CurStmt,
   Stmt *OrigStmt = CurStmt;
   CurStmt = CurStmt->IgnoreContainers();
   // 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 (!CS)
           continue;
         for (Stmt *S : CS->body()) {
           if (!S)
             continue;
-          if (isa<ForStmt>(S) || isa<CXXForRangeStmt>(S)) {
+          if (isa<ForStmt>(S) || isa<CXXForRangeStmt>(S) ||
+              (isa<OMPLoopBasedDirective>(S) && !isa<OMPLoopDirective>(S))) {
             // Only single loop construct is allowed.
             if (CurStmt) {
               CurStmt = OrigStmt;
               break;
             }
             CurStmt = S;
             continue;
           }
if (auto *CS = dyn_cast<CompoundStmt>(CurStmt)) {
       }
       if (!CurStmt)
         CurStmt = OrigStmt;
     }
   }
   return CurStmt;
 }
 
-Stmt *OMPLoopDirective::getBody() {
-  // This relies on the loop form is already checked by Sema.
-  Stmt *Body = Data->getRawStmt()->IgnoreContainers();
-  if (auto *For = dyn_cast<ForStmt>(Body)) {
-    Body = For->getBody();
+bool OMPLoopBasedDirective::doForAllLoops(
+    Stmt *CurStmt, bool TryImperfectlyNestedLoops, unsigned NumLoops,
+    llvm::function_ref<bool(unsigned, Stmt *)> Callback) {
+  CurStmt = CurStmt->IgnoreContainers();
+  for (unsigned Cnt = 0; Cnt < NumLoops; ++Cnt) {
+    if (auto *Dir = dyn_cast<OMPTileDirective>(CurStmt))
+      CurStmt = Dir->getTransformedStmt();
+    if (Callback(Cnt, CurStmt))
+      return false;
+    // 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>(Body) &&
-           "Expected canonical for loop or range-based for loop.");
-    Body = cast<CXXForRangeStmt>(Body)->getBody();
+      assert(isa<CXXForRangeStmt>(CurStmt) &&
+             "Expected canonical for or range-based for loops.");
+      CurStmt = cast<CXXForRangeStmt>(CurStmt)->getBody();
     }
-  for (unsigned Cnt = 1; Cnt < CollapsedNum; ++Cnt) {
-    Body = tryToFindNextInnerLoop(Body, /*TryImperfectlyNestedLoops=*/true);
-    if (auto *For = dyn_cast<ForStmt>(Body)) {
+    CurStmt = OMPLoopBasedDirective::tryToFindNextInnerLoop(
+        CurStmt, TryImperfectlyNestedLoops);
+  }
+  return true;
+}
+
+void OMPLoopBasedDirective::doForAllLoopsBodies(
+    Stmt *CurStmt, bool TryImperfectlyNestedLoops, unsigned NumLoops,
+    llvm::function_ref<void(unsigned, Stmt *, Stmt *)> Callback) {
+  bool Res = OMPLoopBasedDirective::doForAllLoops(
+      CurStmt, TryImperfectlyNestedLoops, NumLoops,
+      [Callback](unsigned Cnt, Stmt *Loop) {
+        Stmt *Body = nullptr;
+        if (auto *For = dyn_cast<ForStmt>(Loop)) {
           Body = For->getBody();
         } else {
-      assert(isa<CXXForRangeStmt>(Body) &&
-             "Expected canonical for loop or range-based for loop.");
-      Body = cast<CXXForRangeStmt>(Body)->getBody();
+          assert(isa<CXXForRangeStmt>(Loop) &&
+                 "Expected canonical for or range-based for loops.");
+          Body = cast<CXXForRangeStmt>(Loop)->getBody();
         }
-  }
+        Callback(Cnt, Loop, Body);
+        return false;
+      });
+  assert(Res && "Expected only loops");
+  (void)Res;
+
+Stmt *OMPLoopDirective::getBody() {
+  // This relies on the loop form is already checked by Sema.
+  Stmt *Body = nullptr;
+  OMPLoopBasedDirective::doForAllLoopsBodies(
+      Data->getRawStmt(), /*TryImperfectlyNestedLoops=*/true, CollapsedNum,
+      [&Body](unsigned, Stmt *, Stmt *BodyStmt) { Body = BodyStmt; });
   return Body;
 }
 
 void OMPLoopDirective::setCounters(ArrayRef<Expr *> A) {
   assert(A.size() == getCollapsedNumber() &&
          "Number of loop counters is not the same as the collapsed number");
   llvm::copy(A, getCounters().begin());
 }
OMPForDirective *OMPForDirective::CreateEmpty(const ASTContext &C,
                                               unsigned NumClauses,
                                               unsigned CollapsedNum,
                                               EmptyShell) {
   return createEmptyDirective<OMPForDirective>(
       C, NumClauses, /*HasAssociatedStmt=*/true,
       numLoopChildren(CollapsedNum, OMPD_for) + 1, CollapsedNum);
 }
 
+OMPTileDirective *
+OMPTileDirective::Create(const ASTContext &C, SourceLocation StartLoc,
+                         SourceLocation EndLoc, ArrayRef<OMPClause *> Clauses,
+                         unsigned NumLoops, Stmt *AssociatedStmt,
+                         Stmt *TransformedStmt, Stmt *PreInits) {
+  OMPTileDirective *Dir = createDirective<OMPTileDirective>(
+      C, Clauses, AssociatedStmt, TransformedStmtOffset + 1, StartLoc, EndLoc,
+      NumLoops);
+  Dir->setTransformedStmt(TransformedStmt);
+  Dir->setPreInits(PreInits);
+  return Dir;
+}
+
+OMPTileDirective *OMPTileDirective::CreateEmpty(const ASTContext &C,
+                                                unsigned NumClauses,
+                                                unsigned NumLoops) {
+  return createEmptyDirective<OMPTileDirective>(
+      C, NumClauses, /*HasAssociatedStmt=*/true, TransformedStmtOffset + 1,
+      SourceLocation(), SourceLocation(), NumLoops);
+}
+
 OMPForSimdDirective *
 OMPForSimdDirective::Create(const ASTContext &C, SourceLocation StartLoc,
                             SourceLocation EndLoc, unsigned CollapsedNum,
                             ArrayRef<OMPClause *> Clauses, Stmt *AssociatedStmt,
                             const HelperExprs &Exprs) {
   auto *Dir = createDirective<OMPForSimdDirective>(
       C, Clauses, AssociatedStmt, numLoopChildren(CollapsedNum, OMPD_for_simd),
       StartLoc, EndLoc, CollapsedNum);

clang/lib/AST/StmtPrinter.cpp

void StmtPrinter::VisitOMPParallelDirective(OMPParallelDirective *Node) {
   PrintOMPExecutableDirective(Node);
 }
 
 void StmtPrinter::VisitOMPSimdDirective(OMPSimdDirective *Node) {
   Indent() << "#pragma omp simd";
   PrintOMPExecutableDirective(Node);
 }
 
+void StmtPrinter::VisitOMPTileDirective(OMPTileDirective *Node) {
+  Indent() << "#pragma omp tile";
+  PrintOMPExecutableDirective(Node);
+}
+
 void StmtPrinter::VisitOMPForDirective(OMPForDirective *Node) {
   Indent() << "#pragma omp for";
   PrintOMPExecutableDirective(Node);
 }
 
 void StmtPrinter::VisitOMPForSimdDirective(OMPForSimdDirective *Node) {
   Indent() << "#pragma omp for simd";
   PrintOMPExecutableDirective(Node);

clang/lib/AST/StmtProfile.cpp

if (C->getSafelen())
     Profiler->VisitStmt(C->getSafelen());
 }
 
 void OMPClauseProfiler::VisitOMPSimdlenClause(const OMPSimdlenClause *C) {
   if (C->getSimdlen())
     Profiler->VisitStmt(C->getSimdlen());
 }
 
+void OMPClauseProfiler::VisitOMPSizesClause(const OMPSizesClause *C) {
+  for (auto E : C->getSizesRefs())
+    if (E)
+      Profiler->VisitExpr(E);
+}
+
 void OMPClauseProfiler::VisitOMPAllocatorClause(const OMPAllocatorClause *C) {
   if (C->getAllocator())
     Profiler->VisitStmt(C->getAllocator());
 }
 
 void OMPClauseProfiler::VisitOMPCollapseClause(const OMPCollapseClause *C) {
   if (C->getNumForLoops())
     Profiler->VisitStmt(C->getNumForLoops());
StmtProfiler::VisitOMPExecutableDirective(const OMPExecutableDirective *S) {
   OMPClauseProfiler P(this);
   ArrayRef<OMPClause *> Clauses = S->clauses();
   for (ArrayRef<OMPClause *>::iterator I = Clauses.begin(), E = Clauses.end();
        I != E; ++I)
     if (*I)
       P.Visit(*I);
 }
 
-void StmtProfiler::VisitOMPLoopDirective(const OMPLoopDirective *S) {
+void StmtProfiler::VisitOMPLoopBasedDirective(const OMPLoopBasedDirective *S) {
   VisitOMPExecutableDirective(S);
 }
 
+void StmtProfiler::VisitOMPLoopDirective(const OMPLoopDirective *S) {
+  VisitOMPLoopBasedDirective(S);
+}
+
 void StmtProfiler::VisitOMPParallelDirective(const OMPParallelDirective *S) {
   VisitOMPExecutableDirective(S);
 }
 
 void StmtProfiler::VisitOMPSimdDirective(const OMPSimdDirective *S) {
   VisitOMPLoopDirective(S);
 }
 
+void StmtProfiler::VisitOMPTileDirective(const OMPTileDirective *S) {
+  VisitOMPLoopBasedDirective(S);
+}
+
 void StmtProfiler::VisitOMPForDirective(const OMPForDirective *S) {
   VisitOMPLoopDirective(S);
 }
 
 void StmtProfiler::VisitOMPForSimdDirective(const OMPForSimdDirective *S) {
   VisitOMPLoopDirective(S);
 }
 

clang/lib/Basic/OpenMPKinds.cpp

#include "clang/Basic/OpenMPKinds.def"
         .Default(OMPC_REDUCTION_unknown);
   case OMPC_unknown:
   case OMPC_threadprivate:
   case OMPC_if:
   case OMPC_final:
   case OMPC_num_threads:
   case OMPC_safelen:
   case OMPC_simdlen:
+  case OMPC_sizes:
   case OMPC_allocator:
   case OMPC_allocate:
   case OMPC_collapse:
   case OMPC_private:
   case OMPC_firstprivate:
   case OMPC_shared:
   case OMPC_task_reduction:
   case OMPC_in_reduction:
#include "clang/Basic/OpenMPKinds.def"
     llvm_unreachable("Invalid OpenMP 'reduction' clause modifier");
   case OMPC_unknown:
   case OMPC_threadprivate:
   case OMPC_if:
   case OMPC_final:
   case OMPC_num_threads:
   case OMPC_safelen:
   case OMPC_simdlen:
+  case OMPC_sizes:
   case OMPC_allocator:
   case OMPC_allocate:
   case OMPC_collapse:
   case OMPC_private:
   case OMPC_firstprivate:
   case OMPC_shared:
   case OMPC_task_reduction:
   case OMPC_in_reduction:
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(unsigned(DKind) < llvm::omp::Directive_enumSize);
   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_parallel_for_simd:
   case OMPD_cancel:
   case OMPD_cancellation_point:
   case OMPD_ordered:
   case OMPD_threadprivate:
   case OMPD_allocate:
   case OMPD_task:
   case OMPD_simd:
+  case OMPD_tile:
   case OMPD_sections:
   case OMPD_section:
   case OMPD_single:
   case OMPD_master:
   case OMPD_critical:
   case OMPD_taskyield:
   case OMPD_barrier:
   case OMPD_taskwait:
emitNumThreadsForTargetDirective(CodeGenFunction &CGF,
   case OMPD_parallel_for_simd:
   case OMPD_cancel:
   case OMPD_cancellation_point:
   case OMPD_ordered:
   case OMPD_threadprivate:
   case OMPD_allocate:
   case OMPD_task:
   case OMPD_simd:
+  case OMPD_tile:
   case OMPD_sections:
   case OMPD_section:
   case OMPD_single:
   case OMPD_master:
   case OMPD_critical:
   case OMPD_taskyield:
   case OMPD_barrier:
   case OMPD_taskwait:
if (const auto *NestedDir =
     case OMPD_parallel_for_simd:
     case OMPD_cancel:
     case OMPD_cancellation_point:
     case OMPD_ordered:
     case OMPD_threadprivate:
     case OMPD_allocate:
     case OMPD_task:
     case OMPD_simd:
+    case OMPD_tile:
     case OMPD_sections:
     case OMPD_section:
     case OMPD_single:
     case OMPD_master:
     case OMPD_critical:
     case OMPD_taskyield:
     case OMPD_barrier:
     case OMPD_taskwait:
if (RequiresDeviceCodegen) {
     case OMPD_parallel_for_simd:
     case OMPD_cancel:
     case OMPD_cancellation_point:
     case OMPD_ordered:
     case OMPD_threadprivate:
     case OMPD_allocate:
     case OMPD_task:
     case OMPD_simd:
+    case OMPD_tile:
     case OMPD_sections:
     case OMPD_section:
     case OMPD_single:
     case OMPD_master:
     case OMPD_critical:
     case OMPD_taskyield:
     case OMPD_barrier:
     case OMPD_taskwait:
auto &&ThenGen = [this, &D, Device, &InputInfo, &MapTypesArray,
     case OMPD_parallel_for_simd:
     case OMPD_cancel:
     case OMPD_cancellation_point:
     case OMPD_ordered:
     case OMPD_threadprivate:
     case OMPD_allocate:
     case OMPD_task:
     case OMPD_simd:
+    case OMPD_tile:
     case OMPD_sections:
     case OMPD_section:
     case OMPD_single:
     case OMPD_master:
     case OMPD_critical:
     case OMPD_taskyield:
     case OMPD_barrier:
     case OMPD_taskwait:

clang/lib/CodeGen/CGStmt.cpp

case Stmt::SEHTryStmtClass:
     EmitSEHTryStmt(cast<SEHTryStmt>(*S));
     break;
   case Stmt::OMPParallelDirectiveClass:
     EmitOMPParallelDirective(cast<OMPParallelDirective>(*S));
     break;
   case Stmt::OMPSimdDirectiveClass:
     EmitOMPSimdDirective(cast<OMPSimdDirective>(*S));
     break;
+  case Stmt::OMPTileDirectiveClass:
+    EmitOMPTileDirective(cast<OMPTileDirective>(*S));
+    break;
   case Stmt::OMPForDirectiveClass:
     EmitOMPForDirective(cast<OMPForDirective>(*S));
     break;
   case Stmt::OMPForSimdDirectiveClass:
     EmitOMPForSimdDirective(cast<OMPForSimdDirective>(*S));
     break;
   case Stmt::OMPSectionsDirectiveClass:
     EmitOMPSectionsDirective(cast<OMPSectionsDirective>(*S));

clang/lib/CodeGen/CGStmtOpenMP.cpp

public:
   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) {
+  void emitPreInitStmt(CodeGenFunction &CGF, const OMPLoopBasedDirective &S) {
+    const DeclStmt *PreInits;
     CodeGenFunction::OMPMapVars PreCondVars;
+    if (auto *LD = dyn_cast<OMPLoopDirective>(&S)) {
       llvm::DenseSet<const VarDecl *> EmittedAsPrivate;
-    for (const auto *E : S.counters()) {
+      for (const auto *E : LD->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 auto *C : LD->getClausesOfKind<OMPPrivateClause>()) {
         for (const Expr *IRef : C->varlists()) {
-        const auto *OrigVD = cast<VarDecl>(cast<DeclRefExpr>(IRef)->getDecl());
+          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(
+                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);
+      (void)OMPLoopBasedDirective::doForAllLoops(
+          LD->getInnermostCapturedStmt()->getCapturedStmt(),
+          /*TryImperfectlyNestedLoops=*/true, LD->getCollapsedNumber(),
+          [&CGF](unsigned Cnt, const Stmt *CurStmt) {
+            if (const auto *CXXFor = dyn_cast<CXXForRangeStmt>(CurStmt)) {
               if (const Stmt *Init = CXXFor->getInit())
                 CGF.EmitStmt(Init);
               CGF.EmitStmt(CXXFor->getRangeStmt());
               CGF.EmitStmt(CXXFor->getEndStmt());
-        Body = CXXFor->getBody();
             }
+            return false;
+          });
+      PreInits = cast_or_null<DeclStmt>(LD->getPreInits());
+    } else if (const auto *Tile = dyn_cast<OMPTileDirective>(&S)) {
+      PreInits = cast_or_null<DeclStmt>(Tile->getPreInits());
+    } else {
+      llvm_unreachable("Unknown loop-based directive kind.");
     }
-    if (const auto *PreInits = cast_or_null<DeclStmt>(S.getPreInits())) {
+    if (PreInits) {
       for (const auto *I : PreInits->decls())
         CGF.EmitVarDecl(cast<VarDecl>(*I));
     }
     PreCondVars.restore(CGF);
   }
 
 public:
-  OMPLoopScope(CodeGenFunction &CGF, const OMPLoopDirective &S)
+  OMPLoopScope(CodeGenFunction &CGF, const OMPLoopBasedDirective &S)
       : CodeGenFunction::RunCleanupsScope(CGF) {
     emitPreInitStmt(CGF, S);
   }
 };
 
 class OMPSimdLexicalScope : public CodeGenFunction::LexicalScope {
   CodeGenFunction::OMPPrivateScope InlinedShareds;
 
emitCommonOMPParallelDirective(*this, S, OMPD_parallel, CodeGen,
                                    emitEmptyBoundParameters);
     emitPostUpdateForReductionClause(*this, S,
                                      [](CodeGenFunction &) { return nullptr; });
   }
   // Check for outer lastprivate conditional update.
   checkForLastprivateConditionalUpdate(*this, S);
 }
 
+namespace {
+/// RAII to handle scopes for loop transformation directives.
+class OMPTransformDirectiveScopeRAII {
+  OMPLoopScope *Scope = nullptr;
+  CodeGenFunction::CGCapturedStmtInfo *CGSI = nullptr;
+  CodeGenFunction::CGCapturedStmtRAII *CapInfoRAII = nullptr;
+
+public:
+  OMPTransformDirectiveScopeRAII(CodeGenFunction &CGF, const Stmt *S) {
+    if (const auto *Dir = dyn_cast<OMPLoopBasedDirective>(S)) {
+      Scope = new OMPLoopScope(CGF, *Dir);
+      CGSI = new CodeGenFunction::CGCapturedStmtInfo(CR_OpenMP);
+      CapInfoRAII = new CodeGenFunction::CGCapturedStmtRAII(CGF, CGSI);
+    }
+  }
+  ~OMPTransformDirectiveScopeRAII() {
+    if (!Scope)
+      return;
+    delete CapInfoRAII;
+    delete CGSI;
+    delete Scope;
+  }
+};
+} // namespace
+
 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();
   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())
       emitBody(CGF, CurStmt, NextLoop, MaxLevel, Level);
     return;
   }
   if (SimplifiedS == NextLoop) {
+    OMPTransformDirectiveScopeRAII PossiblyTransformDirectiveScope(CGF,
+                                                                   SimplifiedS);
+    if (auto *Dir = dyn_cast<OMPTileDirective>(SimplifiedS))
+      SimplifiedS = Dir->getTransformedStmt();
     if (const auto *For = dyn_cast<ForStmt>(SimplifiedS)) {
       S = For->getBody();
     } else {
       assert(isa<CXXForRangeStmt>(SimplifiedS) &&
              "Expected canonical for loop or range-based for loop.");
       const auto *CXXFor = cast<CXXForRangeStmt>(SimplifiedS);
       CGF.EmitStmt(CXXFor->getLoopVarStmt());
       S = CXXFor->getBody();
if (IsInscanRegion) {
     EmitBlock(OMPBeforeScanBlock);
   }
 
   // Emit loop variables for C++ range loops.
   const Stmt *Body =
       D.getInnermostCapturedStmt()->getCapturedStmt()->IgnoreContainers();
   // Emit loop body.
   emitBody(*this, Body,
-           OMPLoopDirective::tryToFindNextInnerLoop(
+           OMPLoopBasedDirective::tryToFindNextInnerLoop(
                Body, /*TryImperfectlyNestedLoops=*/true),
            D.getCollapsedNumber());
 
   // Jump to the dispatcher at the end of the loop body.
   if (IsInscanRegion)
     EmitBranch(OMPScanExitBlock);
 
   // The end (updates/cleanups).
auto LPCRegion =
         CGOpenMPRuntime::LastprivateConditionalRAII::disable(*this, S);
     OMPLexicalScope Scope(*this, S, OMPD_unknown);
     CGM.getOpenMPRuntime().emitInlinedDirective(*this, OMPD_simd, CodeGen);
   }
   // Check for outer lastprivate conditional update.
   checkForLastprivateConditionalUpdate(*this, S);
 }
 
+void CodeGenFunction::EmitOMPTileDirective(const OMPTileDirective &S) {
+  // Emit the de-sugared statement.
+  OMPTransformDirectiveScopeRAII TileScope(*this, &S);
+  EmitStmt(S.getTransformedStmt());
+}
+
 void CodeGenFunction::EmitOMPOuterLoop(
     bool DynamicOrOrdered, bool IsMonotonic, const OMPLoopDirective &S,
     CodeGenFunction::OMPPrivateScope &LoopScope,
     const CodeGenFunction::OMPLoopArguments &LoopArgs,
     const CodeGenFunction::CodeGenLoopTy &CodeGenLoop,
     const CodeGenFunction::CodeGenOrderedTy &CodeGenOrdered) {
   CGOpenMPRuntime &RT = CGM.getOpenMPRuntime();
 
static void emitOMPAtomicExpr(CodeGenFunction &CGF, OpenMPClauseKind Kind,
   case OMPC_private:
   case OMPC_firstprivate:
   case OMPC_lastprivate:
   case OMPC_reduction:
   case OMPC_task_reduction:
   case OMPC_in_reduction:
   case OMPC_safelen:
   case OMPC_simdlen:
+  case OMPC_sizes:
   case OMPC_allocator:
   case OMPC_allocate:
   case OMPC_collapse:
   case OMPC_default:
   case OMPC_seq_cst:
   case OMPC_acq_rel:
   case OMPC_acquire:
   case OMPC_release:

clang/lib/CodeGen/CodeGenFunction.h

OMPTargetDataInfo(Address BasePointersArray, Address PointersArray,
           NumberOfTargetItems(NumberOfTargetItems) {}
   };
   void EmitOMPTargetTaskBasedDirective(const OMPExecutableDirective &S,
                                        const RegionCodeGenTy &BodyGen,
                                        OMPTargetDataInfo &InputInfo);
 
   void EmitOMPParallelDirective(const OMPParallelDirective &S);
   void EmitOMPSimdDirective(const OMPSimdDirective &S);
+  void EmitOMPTileDirective(const OMPTileDirective &S);
   void EmitOMPForDirective(const OMPForDirective &S);
   void EmitOMPForSimdDirective(const OMPForSimdDirective &S);
   void EmitOMPSectionsDirective(const OMPSectionsDirective &S);
   void EmitOMPSectionDirective(const OMPSectionDirective &S);
   void EmitOMPSingleDirective(const OMPSingleDirective &S);
   void EmitOMPMasterDirective(const OMPMasterDirective &S);
   void EmitOMPCriticalDirective(const OMPCriticalDirective &S);
   void EmitOMPParallelForDirective(const OMPParallelForDirective &S);

clang/lib/Parse/ParseOpenMP.cpp

case OMPD_declare_target: {
     Actions.ActOnFinishOpenMPDeclareTargetDirective();
     return Actions.BuildDeclaratorGroup(Decls);
   }
   case OMPD_unknown:
     Diag(Tok, diag::err_omp_unknown_directive);
     break;
   case OMPD_parallel:
   case OMPD_simd:
+  case OMPD_tile:
   case OMPD_task:
   case OMPD_taskyield:
   case OMPD_barrier:
   case OMPD_taskwait:
   case OMPD_taskgroup:
   case OMPD_flush:
   case OMPD_depobj:
   case OMPD_scan:
if ((StmtCtx & ParsedStmtContext::AllowStandaloneOpenMPDirectives) ==
       Diag(Tok, diag::err_omp_immediate_directive)
           << getOpenMPDirectiveName(DKind) << 0;
     }
     HasAssociatedStatement = false;
     // Fall through for further analysis.
     LLVM_FALLTHROUGH;
   case OMPD_parallel:
   case OMPD_simd:
+  case OMPD_tile:
   case OMPD_for:
   case OMPD_for_simd:
   case OMPD_sections:
   case OMPD_single:
   case OMPD_section:
   case OMPD_master:
   case OMPD_critical:
   case OMPD_parallel_for:
if (DKind == OMPD_ordered && FirstClauses[unsigned(OMPC_depend)].getInt()) {
           ParsedStmtContext()) {
         Diag(Loc, diag::err_omp_immediate_directive)
             << getOpenMPDirectiveName(DKind) << 1
             << getOpenMPClauseName(OMPC_depend);
       }
       HasAssociatedStatement = false;
     }
 
+    if (DKind == OMPD_tile && !FirstClauses[unsigned(OMPC_sizes)].getInt()) {
+      Diag(Loc, diag::err_omp_required_clause)
+          << getOpenMPDirectiveName(OMPD_tile) << "sizes";
+    }
+
     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.
       ParsingOpenMPDirectiveRAII NormalScope(*this, /*Value=*/false);
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;
+  }
+
+  while (true) {
+    ExprResult Val = ParseConstantExpression();
+    if (!Val.isUsable()) {
+      T.skipToEnd();
+      return nullptr;
+    }
+
+    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());
+}
+
 OMPClause *Parser::ParseOpenMPUsesAllocatorClause(OpenMPDirectiveKind DKind) {
   SourceLocation Loc = Tok.getLocation();
   ConsumeAnyToken();
 
   // Parse '('.
   BalancedDelimiterTracker T(*this, tok::l_paren, tok::annot_pragma_openmp_end);
   if (T.expectAndConsume(diag::err_expected_lparen_after, "uses_allocator"))
     return nullptr;
OMPClause *Parser::ParseOpenMPClause(OpenMPDirectiveKind DKind,
   case OMPC_is_device_ptr:
   case OMPC_allocate:
   case OMPC_nontemporal:
   case OMPC_inclusive:
   case OMPC_exclusive:
   case OMPC_affinity:
     Clause = ParseOpenMPVarListClause(DKind, CKind, WrongDirective);
     break;
+  case OMPC_sizes:
+    if (!FirstClause) {
+      Diag(Tok, diag::err_omp_more_one_clause)
+          << getOpenMPDirectiveName(DKind) << getOpenMPClauseName(CKind) << 0;
+      ErrorFound = true;
+    }
+
+    Clause = ParseOpenMPSizesClause();
+    break;
   case OMPC_uses_allocators:
     Clause = ParseOpenMPUsesAllocatorClause(DKind);
     break;
   case OMPC_device_type:
   case OMPC_unknown:
     skipUntilPragmaOpenMPEnd(DKind);
     break;
   case OMPC_threadprivate:

clang/lib/Sema/SemaExceptionSpec.cpp

CanThrowResult Sema::canThrow(const Stmt *S) {
   case Stmt::OMPParallelForSimdDirectiveClass:
   case Stmt::OMPParallelMasterDirectiveClass:
   case Stmt::OMPParallelMasterTaskLoopDirectiveClass:
   case Stmt::OMPParallelMasterTaskLoopSimdDirectiveClass:
   case Stmt::OMPParallelSectionsDirectiveClass:
   case Stmt::OMPSectionDirectiveClass:
   case Stmt::OMPSectionsDirectiveClass:
   case Stmt::OMPSimdDirectiveClass:
+  case Stmt::OMPTileDirectiveClass:
   case Stmt::OMPSingleDirectiveClass:
   case Stmt::OMPTargetDataDirectiveClass:
   case Stmt::OMPTargetDirectiveClass:
   case Stmt::OMPTargetEnterDataDirectiveClass:
   case Stmt::OMPTargetExitDataDirectiveClass:
   case Stmt::OMPTargetParallelDirectiveClass:
   case Stmt::OMPTargetParallelForDirectiveClass:
   case Stmt::OMPTargetParallelForSimdDirectiveClass:

clang/lib/Sema/SemaOpenMP.cpp

class DSAAttrChecker final : public StmtVisitor<DSAAttrChecker, void> {
 
   void VisitSubCaptures(OMPExecutableDirective *S) {
     // Check implicitly captured variables.
     if (!S->hasAssociatedStmt() || !S->getAssociatedStmt())
       return;
     if (S->getDirectiveKind() == OMPD_atomic ||
         S->getDirectiveKind() == OMPD_critical ||
         S->getDirectiveKind() == OMPD_section ||
-        S->getDirectiveKind() == OMPD_master) {
+        S->getDirectiveKind() == OMPD_master ||
+        isOpenMPLoopTransformationDirective(S->getDirectiveKind())) {
       Visit(S->getAssociatedStmt());
       return;
     }
     visitSubCaptures(S->getInnermostCapturedStmt());
     // Try to capture inner this->member references to generate correct mappings
     // and diagnostics.
     if (TryCaptureCXXThisMembers ||
         (isOpenMPTargetExecutionDirective(Stack->getCurrentDirective()) &&
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);
       }
     }
ActOnCapturedRegionStart(DSAStack->getConstructLoc(), CurScope, CR_OpenMP,
                              std::make_pair(StringRef(), QualType()),
                              /*OpenMPCaptureLevel=*/1);
     break;
   }
   case OMPD_atomic:
   case OMPD_critical:
   case OMPD_section:
   case OMPD_master:
+  case OMPD_tile:
     break;
   case OMPD_simd:
   case OMPD_for:
   case OMPD_for_simd:
   case OMPD_sections:
   case OMPD_single:
   case OMPD_taskgroup:
   case OMPD_distribute:
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() &&
+             "No captured regions in loop transformation directives.");
     } 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) ||
                             StartLoc))
     return StmtError();
 
   llvm::SmallVector<OMPClause *, 8> ClausesWithImplicit;
   VarsWithInheritedDSAType VarsWithInheritedDSA;
   bool ErrorFound = false;
   ClausesWithImplicit.append(Clauses.begin(), Clauses.end());
   if (AStmt && !CurContext->isDependentContext() && Kind != OMPD_atomic &&
-      Kind != OMPD_critical && Kind != OMPD_section && Kind != OMPD_master) {
+      Kind != OMPD_critical && Kind != OMPD_section && Kind != OMPD_master &&
+      !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_parallel:
     AllowedNameModifiers.push_back(OMPD_parallel);
     break;
   case OMPD_simd:
     Res = ActOnOpenMPSimdDirective(ClausesWithImplicit, AStmt, StartLoc, EndLoc,
                                    VarsWithInheritedDSA);
     if (LangOpts.OpenMP >= 50)
       AllowedNameModifiers.push_back(OMPD_simd);
     break;
+  case OMPD_tile:
+    Res =
+        ActOnOpenMPTileDirective(ClausesWithImplicit, AStmt, StartLoc, EndLoc);
+    break;
   case OMPD_for:
     Res = ActOnOpenMPForDirective(ClausesWithImplicit, AStmt, StartLoc, EndLoc,
                                   VarsWithInheritedDSA);
     break;
   case OMPD_for_simd:
     Res = ActOnOpenMPForSimdDirective(ClausesWithImplicit, AStmt, StartLoc,
                                       EndLoc, VarsWithInheritedDSA);
     if (LangOpts.OpenMP >= 50)
for (OMPClause *C : Clauses) {
       case OMPC_ordered:
       case OMPC_device:
       case OMPC_num_teams:
       case OMPC_thread_limit:
       case OMPC_hint:
       case OMPC_collapse:
       case OMPC_safelen:
       case OMPC_simdlen:
+      case OMPC_sizes:
       case OMPC_default:
       case OMPC_proc_bind:
       case OMPC_private:
       case OMPC_firstprivate:
       case OMPC_lastprivate:
       case OMPC_shared:
       case OMPC_reduction:
       case OMPC_task_reduction:

 };
 
 /// 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;
+  /// 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,
-                              SourceLocation DefaultLoc)
-      : SemaRef(SemaRef), Stack(Stack), DefaultLoc(DefaultLoc),
-        ConditionLoc(DefaultLoc) {}
+  OpenMPIterationSpaceChecker(Sema &SemaRef, bool SupportsNonRectangular,
+                              DSAStackTy &Stack, SourceLocation DefaultLoc)
+      : SemaRef(SemaRef), 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);
   /// Check incr-expr for canonical loop form and return true if it
   /// does not conform, otherwise save loop step (#Step).

 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) {
Expr *OpenMPIterationSpaceChecker::buildNumIterations(
   }
   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();
   if (!Upper || !Lower)
     return nullptr;
 
-  ExprResult Diff =
-      calculateNumIters(SemaRef, S, DefaultLoc, Lower, Upper, Step, VarType,
-                        TestIsStrictOp, /*RoundToStep=*/true, Captures);
+  ExprResult Diff = calculateNumIters(SemaRef, S, DefaultLoc, Lower, Upper,
+                                      Step, VarType, TestIsStrictOp,
+                                      /*RoundToStep=*/true, Captures);
   if (!Diff.isUsable())
     return nullptr;
 
   // OpenMP runtime requires 32-bit or 64-bit loop variables.
   QualType Type = Diff.get()->getType();
   ASTContext &C = SemaRef.Context;
   bool UseVarType = VarType->hasIntegerRepresentation() &&
                     C.getTypeSize(Type) > C.getTypeSize(VarType);
std::pair<Expr *, Expr *> OpenMPIterationSpaceChecker::buildMinMaxValues(
   if (TestIsLessOp.getValue())
     MinExpr = Lower;
   else
     MaxExpr = Upper;
 
   // Build minimum/maximum value based on number of iterations.
   QualType VarType = LCDecl->getType().getNonReferenceType();
 
-  ExprResult Diff =
-      calculateNumIters(SemaRef, S, DefaultLoc, Lower, Upper, Step, VarType,
-                        TestIsStrictOp, /*RoundToStep=*/false, Captures);
+  ExprResult Diff = calculateNumIters(SemaRef, S, DefaultLoc, Lower, Upper,
+                                      Step, VarType, TestIsStrictOp,
+                                      /*RoundToStep=*/false, Captures);
   if (!Diff.isUsable())
     return std::make_pair(nullptr, nullptr);
 
   // ((Upper - Lower [- 1]) / Step) * Step
   // Parentheses (for dumping/debugging purposes only).
   Diff = SemaRef.ActOnParenExpr(DefaultLoc, DefaultLoc, Diff.get());
   if (!Diff.isUsable())
     return std::make_pair(nullptr, nullptr);
Expr *Upper = TestIsLessOp.getValue()
                     ? Cnt
                     : tryBuildCapture(SemaRef, LB, Captures).get();
   Expr *Lower = TestIsLessOp.getValue()
                     ? tryBuildCapture(SemaRef, LB, Captures).get()
                     : Cnt;
   if (!Upper || !Lower)
     return nullptr;
 
-  ExprResult Diff = calculateNumIters(SemaRef, S, DefaultLoc, Lower, Upper,
-                                      Step, VarType, /*TestIsStrictOp=*/false,
-                                      /*RoundToStep=*/false, Captures);
+  ExprResult Diff = calculateNumIters(
+      SemaRef, S, DefaultLoc, Lower, Upper, Step, VarType,
+      /*TestIsStrictOp=*/false, /*RoundToStep=*/false, Captures);
   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, /*SupportsNonRectangular=*/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 {

 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) {
+  bool SupportsNonRectangular = !isOpenMPLoopTransformationDirective(DKind);
   // 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)
if (TotalNestedLoopCount > 1) {
                      diag::note_omp_collapse_ordered_expr)
             << 1 << OrderedLoopCountExpr->getSourceRange();
     }
     return true;
   }
   assert(((For && For->getBody()) || (CXXFor && CXXFor->getBody())) &&
          "No loop body.");
 
-  OpenMPIterationSpaceChecker ISC(SemaRef, DSA,
+  OpenMPIterationSpaceChecker ISC(SemaRef, SupportsNonRectangular, DSA,
                                   For ? For->getForLoc() : CXXFor->getForLoc());
 
   // Check init.
   Stmt *Init = For ? For->getInit() : CXXFor->getBeginStmt();
   if (ISC.checkAndSetInit(Init))
     return true;
 
   bool HasErrors = false;
static bool checkOpenMPIterationSpace(
 
   // 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();

 /// 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) {
+                OMPLoopBasedDirective::HelperExprs &Built) {
   unsigned NestedLoopCount = 1;
+  bool SupportsNonPerfectlyNested = (SemaRef.LangOpts.OpenMP >= 50) &&
+                                    !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) {
     } else {
       Built.clear(/*Size=*/1);
       return 1;
     }
   }
   // 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);
-  for (unsigned Cnt = 0; Cnt < NestedLoopCount; ++Cnt) {
-    if (checkOpenMPIterationSpace(
-            DKind, CurStmt, SemaRef, DSA, Cnt, NestedLoopCount,
-            std::max(OrderedLoopCount, NestedLoopCount), CollapseLoopCountExpr,
-            OrderedLoopCountExpr, VarsWithImplicitDSA, IterSpaces, Captures))
-      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) {
+  unsigned NumLoops = std::max(OrderedLoopCount, NestedLoopCount);
+  SmallVector<LoopIterationSpace, 4> IterSpaces(NumLoops);
+  if (!OMPLoopBasedDirective::doForAllLoops(
+          AStmt->IgnoreContainers(!isOpenMPLoopTransformationDirective(DKind)),
+          SupportsNonPerfectlyNested, NumLoops,
+          [DKind, &SemaRef, &DSA, NumLoops, NestedLoopCount,
+           CollapseLoopCountExpr, OrderedLoopCountExpr, &VarsWithImplicitDSA,
+           &IterSpaces, &Captures](unsigned Cnt, Stmt *CurStmt) {
             if (checkOpenMPIterationSpace(
                     DKind, CurStmt, SemaRef, DSA, Cnt, NestedLoopCount,
-            std::max(OrderedLoopCount, NestedLoopCount), CollapseLoopCountExpr,
-            OrderedLoopCountExpr, VarsWithImplicitDSA, IterSpaces, Captures))
-      return 0;
-    if (Cnt > 0 && IterSpaces[Cnt].CounterVar) {
+                    NumLoops, CollapseLoopCountExpr, OrderedLoopCountExpr,
+                    VarsWithImplicitDSA, IterSpaces, Captures))
+              return true;
+            if (Cnt > 0 && Cnt >= NestedLoopCount &&
+                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);
-  }
+            return false;
+          }))
+    return 0;
 
   Built.clear(/* size */ NestedLoopCount);
 
   if (SemaRef.CurContext->isDependentContext())
     return NestedLoopCount;
 
   // An example of what is generated for the following code:
   //
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 =

 StmtResult
 Sema::ActOnOpenMPSimdDirective(ArrayRef<OMPClause *> Clauses, Stmt *AStmt,
                                SourceLocation StartLoc, SourceLocation EndLoc,
                                VarsWithInheritedDSAType &VarsWithImplicitDSA) {
   if (!AStmt)
     return StmtError();
 
   assert(isa<CapturedStmt>(AStmt) && "Captured statement expected");
-  OMPLoopDirective::HelperExprs B;
+  OMPLoopBasedDirective::HelperExprs B;
   // In presence of clause 'collapse' or 'ordered' with number of loops, it will
   // define the nested loops number.
   unsigned NestedLoopCount = checkOpenMPLoop(
       OMPD_simd, getCollapseNumberExpr(Clauses), getOrderedNumberExpr(Clauses),
       AStmt, *this, *DSAStack, VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
 

 StmtResult
 Sema::ActOnOpenMPForDirective(ArrayRef<OMPClause *> Clauses, Stmt *AStmt,
                               SourceLocation StartLoc, SourceLocation EndLoc,
                               VarsWithInheritedDSAType &VarsWithImplicitDSA) {
   if (!AStmt)
     return StmtError();
 
   assert(isa<CapturedStmt>(AStmt) && "Captured statement expected");
-  OMPLoopDirective::HelperExprs B;
+  OMPLoopBasedDirective::HelperExprs B;
   // In presence of clause 'collapse' or 'ordered' with number of loops, it will
   // define the nested loops number.
   unsigned NestedLoopCount = checkOpenMPLoop(
       OMPD_for, getCollapseNumberExpr(Clauses), getOrderedNumberExpr(Clauses),
       AStmt, *this, *DSAStack, VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
 

 
 StmtResult Sema::ActOnOpenMPForSimdDirective(
     ArrayRef<OMPClause *> Clauses, Stmt *AStmt, SourceLocation StartLoc,
     SourceLocation EndLoc, VarsWithInheritedDSAType &VarsWithImplicitDSA) {
   if (!AStmt)
     return StmtError();
 
   assert(isa<CapturedStmt>(AStmt) && "Captured statement expected");
-  OMPLoopDirective::HelperExprs B;
+  OMPLoopBasedDirective::HelperExprs B;
   // In presence of clause 'collapse' or 'ordered' with number of loops, it will
   // define the nested loops number.
   unsigned NestedLoopCount =
       checkOpenMPLoop(OMPD_for_simd, getCollapseNumberExpr(Clauses),
                       getOrderedNumberExpr(Clauses), AStmt, *this, *DSAStack,
                       VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
StmtResult Sema::ActOnOpenMPParallelForDirective(
   auto *CS = cast<CapturedStmt>(AStmt);
   // 1.2.2 OpenMP Language Terminology
   // Structured block - An executable statement with a single entry at the
   // top and a single exit at the bottom.
   // The point of exit cannot be a branch out of the structured block.
   // longjmp() and throw() must not violate the entry/exit criteria.
   CS->getCapturedDecl()->setNothrow();
 
-  OMPLoopDirective::HelperExprs B;
+  OMPLoopBasedDirective::HelperExprs B;
   // In presence of clause 'collapse' or 'ordered' with number of loops, it will
   // define the nested loops number.
   unsigned NestedLoopCount =
       checkOpenMPLoop(OMPD_parallel_for, getCollapseNumberExpr(Clauses),
                       getOrderedNumberExpr(Clauses), AStmt, *this, *DSAStack,
                       VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
StmtResult Sema::ActOnOpenMPParallelForSimdDirective(
   auto *CS = cast<CapturedStmt>(AStmt);
   // 1.2.2 OpenMP Language Terminology
   // Structured block - An executable statement with a single entry at the
   // top and a single exit at the bottom.
   // The point of exit cannot be a branch out of the structured block.
   // longjmp() and throw() must not violate the entry/exit criteria.
   CS->getCapturedDecl()->setNothrow();
 
-  OMPLoopDirective::HelperExprs B;
+  OMPLoopBasedDirective::HelperExprs B;
   // In presence of clause 'collapse' or 'ordered' with number of loops, it will
   // define the nested loops number.
   unsigned NestedLoopCount =
       checkOpenMPLoop(OMPD_parallel_for_simd, getCollapseNumberExpr(Clauses),
                       getOrderedNumberExpr(Clauses), AStmt, *this, *DSAStack,
                       VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
for (int ThisCaptureLevel = getOpenMPCaptureLevels(OMPD_target_parallel_for);
     // 1.2.2 OpenMP Language Terminology
     // Structured block - An executable statement with a single entry at the
     // top and a single exit at the bottom.
     // The point of exit cannot be a branch out of the structured block.
     // longjmp() and throw() must not violate the entry/exit criteria.
     CS->getCapturedDecl()->setNothrow();
   }
 
-  OMPLoopDirective::HelperExprs B;
+  OMPLoopBasedDirective::HelperExprs B;
   // In presence of clause 'collapse' or 'ordered' with number of loops, it will
   // define the nested loops number.
   unsigned NestedLoopCount =
       checkOpenMPLoop(OMPD_target_parallel_for, getCollapseNumberExpr(Clauses),
                       getOrderedNumberExpr(Clauses), CS, *this, *DSAStack,
                       VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();

 
 StmtResult Sema::ActOnOpenMPTaskLoopDirective(
     ArrayRef<OMPClause *> Clauses, Stmt *AStmt, SourceLocation StartLoc,
     SourceLocation EndLoc, VarsWithInheritedDSAType &VarsWithImplicitDSA) {
   if (!AStmt)
     return StmtError();
 
   assert(isa<CapturedStmt>(AStmt) && "Captured statement expected");
-  OMPLoopDirective::HelperExprs B;
+  OMPLoopBasedDirective::HelperExprs B;
   // In presence of clause 'collapse' or 'ordered' with number of loops, it will
   // define the nested loops number.
   unsigned NestedLoopCount =
       checkOpenMPLoop(OMPD_taskloop, getCollapseNumberExpr(Clauses),
                       /*OrderedLoopCountExpr=*/nullptr, AStmt, *this, *DSAStack,
                       VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();

 
 StmtResult Sema::ActOnOpenMPTaskLoopSimdDirective(
     ArrayRef<OMPClause *> Clauses, Stmt *AStmt, SourceLocation StartLoc,
     SourceLocation EndLoc, VarsWithInheritedDSAType &VarsWithImplicitDSA) {
   if (!AStmt)
     return StmtError();
 
   assert(isa<CapturedStmt>(AStmt) && "Captured statement expected");
-  OMPLoopDirective::HelperExprs B;
+  OMPLoopBasedDirective::HelperExprs B;
   // In presence of clause 'collapse' or 'ordered' with number of loops, it will
   // define the nested loops number.
   unsigned NestedLoopCount =
       checkOpenMPLoop(OMPD_taskloop_simd, getCollapseNumberExpr(Clauses),
                       /*OrderedLoopCountExpr=*/nullptr, AStmt, *this, *DSAStack,
                       VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();

 
 StmtResult Sema::ActOnOpenMPMasterTaskLoopDirective(
     ArrayRef<OMPClause *> Clauses, Stmt *AStmt, SourceLocation StartLoc,
     SourceLocation EndLoc, VarsWithInheritedDSAType &VarsWithImplicitDSA) {
   if (!AStmt)
     return StmtError();
 
   assert(isa<CapturedStmt>(AStmt) && "Captured statement expected");
-  OMPLoopDirective::HelperExprs B;
+  OMPLoopBasedDirective::HelperExprs B;
   // In presence of clause 'collapse' or 'ordered' with number of loops, it will
   // define the nested loops number.
   unsigned NestedLoopCount =
       checkOpenMPLoop(OMPD_master_taskloop, getCollapseNumberExpr(Clauses),
                       /*OrderedLoopCountExpr=*/nullptr, AStmt, *this, *DSAStack,
                       VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();

 
 StmtResult Sema::ActOnOpenMPMasterTaskLoopSimdDirective(
     ArrayRef<OMPClause *> Clauses, Stmt *AStmt, SourceLocation StartLoc,
     SourceLocation EndLoc, VarsWithInheritedDSAType &VarsWithImplicitDSA) {
   if (!AStmt)
     return StmtError();
 
   assert(isa<CapturedStmt>(AStmt) && "Captured statement expected");
-  OMPLoopDirective::HelperExprs B;
+  OMPLoopBasedDirective::HelperExprs B;
   // In presence of clause 'collapse' or 'ordered' with number of loops, it will
   // define the nested loops number.
   unsigned NestedLoopCount =
       checkOpenMPLoop(OMPD_master_taskloop_simd, getCollapseNumberExpr(Clauses),
                       /*OrderedLoopCountExpr=*/nullptr, AStmt, *this, *DSAStack,
                       VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
for (int ThisCaptureLevel =
     // 1.2.2 OpenMP Language Terminology
     // Structured block - An executable statement with a single entry at the
     // top and a single exit at the bottom.
     // The point of exit cannot be a branch out of the structured block.
     // longjmp() and throw() must not violate the entry/exit criteria.
     CS->getCapturedDecl()->setNothrow();
   }
 
-  OMPLoopDirective::HelperExprs B;
+  OMPLoopBasedDirective::HelperExprs B;
   // In presence of clause 'collapse' or 'ordered' with number of loops, it will
   // define the nested loops number.
   unsigned NestedLoopCount = checkOpenMPLoop(
       OMPD_parallel_master_taskloop, getCollapseNumberExpr(Clauses),
       /*OrderedLoopCountExpr=*/nullptr, CS, *this, *DSAStack,
       VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
for (int ThisCaptureLevel =
     // 1.2.2 OpenMP Language Terminology
     // Structured block - An executable statement with a single entry at the
     // top and a single exit at the bottom.
     // The point of exit cannot be a branch out of the structured block.
     // longjmp() and throw() must not violate the entry/exit criteria.
     CS->getCapturedDecl()->setNothrow();
   }
 
-  OMPLoopDirective::HelperExprs B;
+  OMPLoopBasedDirective::HelperExprs B;
   // In presence of clause 'collapse' or 'ordered' with number of loops, it will
   // define the nested loops number.
   unsigned NestedLoopCount = checkOpenMPLoop(
       OMPD_parallel_master_taskloop_simd, getCollapseNumberExpr(Clauses),
       /*OrderedLoopCountExpr=*/nullptr, CS, *this, *DSAStack,
       VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();

 
 StmtResult Sema::ActOnOpenMPDistributeDirective(
     ArrayRef<OMPClause *> Clauses, Stmt *AStmt, SourceLocation StartLoc,
     SourceLocation EndLoc, VarsWithInheritedDSAType &VarsWithImplicitDSA) {
   if (!AStmt)
     return StmtError();
 
   assert(isa<CapturedStmt>(AStmt) && "Captured statement expected");
-  OMPLoopDirective::HelperExprs B;
+  OMPLoopBasedDirective::HelperExprs B;
   // In presence of clause 'collapse' with number of loops, it will
   // define the nested loops number.
   unsigned NestedLoopCount =
       checkOpenMPLoop(OMPD_distribute, getCollapseNumberExpr(Clauses),
                       nullptr /*ordered not a clause on distribute*/, AStmt,
                       *this, *DSAStack, VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
for (int ThisCaptureLevel =
     // 1.2.2 OpenMP Language Terminology
     // Structured block - An executable statement with a single entry at the
     // top and a single exit at the bottom.
     // The point of exit cannot be a branch out of the structured block.
     // longjmp() and throw() must not violate the entry/exit criteria.
     CS->getCapturedDecl()->setNothrow();
   }
 
-  OMPLoopDirective::HelperExprs B;
+  OMPLoopBasedDirective::HelperExprs B;
   // In presence of clause 'collapse' with number of loops, it will
   // define the nested loops number.
   unsigned NestedLoopCount = checkOpenMPLoop(
       OMPD_distribute_parallel_for, getCollapseNumberExpr(Clauses),
       nullptr /*ordered not a clause on distribute*/, CS, *this, *DSAStack,
       VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
for (int ThisCaptureLevel =
     // 1.2.2 OpenMP Language Terminology
     // Structured block - An executable statement with a single entry at the
     // top and a single exit at the bottom.
     // The point of exit cannot be a branch out of the structured block.
     // longjmp() and throw() must not violate the entry/exit criteria.
     CS->getCapturedDecl()->setNothrow();
   }
 
-  OMPLoopDirective::HelperExprs B;
+  OMPLoopBasedDirective::HelperExprs B;
   // In presence of clause 'collapse' with number of loops, it will
   // define the nested loops number.
   unsigned NestedLoopCount = checkOpenMPLoop(
       OMPD_distribute_parallel_for_simd, getCollapseNumberExpr(Clauses),
       nullptr /*ordered not a clause on distribute*/, CS, *this, *DSAStack,
       VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
for (int ThisCaptureLevel = getOpenMPCaptureLevels(OMPD_distribute_simd);
     // 1.2.2 OpenMP Language Terminology
     // Structured block - An executable statement with a single entry at the
     // top and a single exit at the bottom.
     // The point of exit cannot be a branch out of the structured block.
     // longjmp() and throw() must not violate the entry/exit criteria.
     CS->getCapturedDecl()->setNothrow();
   }
 
-  OMPLoopDirective::HelperExprs B;
+  OMPLoopBasedDirective::HelperExprs B;
   // In presence of clause 'collapse' with number of loops, it will
   // define the nested loops number.
   unsigned NestedLoopCount =
       checkOpenMPLoop(OMPD_distribute_simd, getCollapseNumberExpr(Clauses),
                       nullptr /*ordered not a clause on distribute*/, CS, *this,
                       *DSAStack, VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
for (int ThisCaptureLevel = getOpenMPCaptureLevels(OMPD_target_parallel_for);
     // 1.2.2 OpenMP Language Terminology
     // Structured block - An executable statement with a single entry at the
     // top and a single exit at the bottom.
     // The point of exit cannot be a branch out of the structured block.
     // longjmp() and throw() must not violate the entry/exit criteria.
     CS->getCapturedDecl()->setNothrow();
   }
 
-  OMPLoopDirective::HelperExprs B;
+  OMPLoopBasedDirective::HelperExprs B;
   // In presence of clause 'collapse' or 'ordered' with number of loops, it will
   // define the nested loops number.
   unsigned NestedLoopCount = checkOpenMPLoop(
       OMPD_target_parallel_for_simd, getCollapseNumberExpr(Clauses),
       getOrderedNumberExpr(Clauses), CS, *this, *DSAStack,
       VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
for (int ThisCaptureLevel = getOpenMPCaptureLevels(OMPD_target_simd);
     // 1.2.2 OpenMP Language Terminology
     // Structured block - An executable statement with a single entry at the
     // top and a single exit at the bottom.
     // The point of exit cannot be a branch out of the structured block.
     // longjmp() and throw() must not violate the entry/exit criteria.
     CS->getCapturedDecl()->setNothrow();
   }
 
-  OMPLoopDirective::HelperExprs B;
+  OMPLoopBasedDirective::HelperExprs B;
   // In presence of clause 'collapse' with number of loops, it will define the
   // nested loops number.
   unsigned NestedLoopCount =
       checkOpenMPLoop(OMPD_target_simd, getCollapseNumberExpr(Clauses),
                       getOrderedNumberExpr(Clauses), CS, *this, *DSAStack,
                       VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
for (int ThisCaptureLevel = getOpenMPCaptureLevels(OMPD_teams_distribute);
     // 1.2.2 OpenMP Language Terminology
     // Structured block - An executable statement with a single entry at the
     // top and a single exit at the bottom.
     // The point of exit cannot be a branch out of the structured block.
     // longjmp() and throw() must not violate the entry/exit criteria.
     CS->getCapturedDecl()->setNothrow();
   }
 
-  OMPLoopDirective::HelperExprs B;
+  OMPLoopBasedDirective::HelperExprs B;
   // In presence of clause 'collapse' with number of loops, it will
   // define the nested loops number.
   unsigned NestedLoopCount =
       checkOpenMPLoop(OMPD_teams_distribute, getCollapseNumberExpr(Clauses),
                       nullptr /*ordered not a clause on distribute*/, CS, *this,
                       *DSAStack, VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
for (int ThisCaptureLevel =
     // 1.2.2 OpenMP Language Terminology
     // Structured block - An executable statement with a single entry at the
     // top and a single exit at the bottom.
     // The point of exit cannot be a branch out of the structured block.
     // longjmp() and throw() must not violate the entry/exit criteria.
     CS->getCapturedDecl()->setNothrow();
   }
 
-  OMPLoopDirective::HelperExprs B;
+  OMPLoopBasedDirective::HelperExprs B;
   // In presence of clause 'collapse' with number of loops, it will
   // define the nested loops number.
   unsigned NestedLoopCount = checkOpenMPLoop(
       OMPD_teams_distribute_simd, getCollapseNumberExpr(Clauses),
       nullptr /*ordered not a clause on distribute*/, CS, *this, *DSAStack,
       VarsWithImplicitDSA, B);
 
   if (NestedLoopCount == 0)
for (int ThisCaptureLevel =
     // 1.2.2 OpenMP Language Terminology
     // Structured block - An executable statement with a single entry at the
     // top and a single exit at the bottom.
     // The point of exit cannot be a branch out of the structured block.
     // longjmp() and throw() must not violate the entry/exit criteria.
     CS->getCapturedDecl()->setNothrow();
   }
 
-  OMPLoopDirective::HelperExprs B;
+  OMPLoopBasedDirective::HelperExprs B;
   // In presence of clause 'collapse' with number of loops, it will
   // define the nested loops number.
   unsigned NestedLoopCount = checkOpenMPLoop(
       OMPD_teams_distribute_parallel_for_simd, getCollapseNumberExpr(Clauses),
       nullptr /*ordered not a clause on distribute*/, CS, *this, *DSAStack,
       VarsWithImplicitDSA, B);
 
   if (NestedLoopCount == 0)
for (int ThisCaptureLevel =
     // 1.2.2 OpenMP Language Terminology
     // Structured block - An executable statement with a single entry at the
     // top and a single exit at the bottom.
     // The point of exit cannot be a branch out of the structured block.
     // longjmp() and throw() must not violate the entry/exit criteria.
     CS->getCapturedDecl()->setNothrow();
   }
 
-  OMPLoopDirective::HelperExprs B;
+  OMPLoopBasedDirective::HelperExprs B;
   // In presence of clause 'collapse' with number of loops, it will
   // define the nested loops number.
   unsigned NestedLoopCount = checkOpenMPLoop(
       OMPD_teams_distribute_parallel_for, getCollapseNumberExpr(Clauses),
       nullptr /*ordered not a clause on distribute*/, CS, *this, *DSAStack,
       VarsWithImplicitDSA, B);
 
   if (NestedLoopCount == 0)
for (int ThisCaptureLevel =
     // 1.2.2 OpenMP Language Terminology
     // Structured block - An executable statement with a single entry at the
     // top and a single exit at the bottom.
     // The point of exit cannot be a branch out of the structured block.
     // longjmp() and throw() must not violate the entry/exit criteria.
     CS->getCapturedDecl()->setNothrow();
   }
 
-  OMPLoopDirective::HelperExprs B;
+  OMPLoopBasedDirective::HelperExprs B;
   // In presence of clause 'collapse' with number of loops, it will
   // define the nested loops number.
   unsigned NestedLoopCount = checkOpenMPLoop(
       OMPD_target_teams_distribute, getCollapseNumberExpr(Clauses),
       nullptr /*ordered not a clause on distribute*/, CS, *this, *DSAStack,
       VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
for (int ThisCaptureLevel =
     // 1.2.2 OpenMP Language Terminology
     // Structured block - An executable statement with a single entry at the
     // top and a single exit at the bottom.
     // The point of exit cannot be a branch out of the structured block.
     // longjmp() and throw() must not violate the entry/exit criteria.
     CS->getCapturedDecl()->setNothrow();
   }
 
-  OMPLoopDirective::HelperExprs B;
+  OMPLoopBasedDirective::HelperExprs B;
   // In presence of clause 'collapse' with number of loops, it will
   // define the nested loops number.
   unsigned NestedLoopCount = checkOpenMPLoop(
       OMPD_target_teams_distribute_parallel_for, getCollapseNumberExpr(Clauses),
       nullptr /*ordered not a clause on distribute*/, CS, *this, *DSAStack,
       VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
for (int ThisCaptureLevel = getOpenMPCaptureLevels(
     // 1.2.2 OpenMP Language Terminology
     // Structured block - An executable statement with a single entry at the
     // top and a single exit at the bottom.
     // The point of exit cannot be a branch out of the structured block.
     // longjmp() and throw() must not violate the entry/exit criteria.
     CS->getCapturedDecl()->setNothrow();
   }
 
-  OMPLoopDirective::HelperExprs B;
+  OMPLoopBasedDirective::HelperExprs B;
   // In presence of clause 'collapse' with number of loops, it will
   // define the nested loops number.
   unsigned NestedLoopCount =
       checkOpenMPLoop(OMPD_target_teams_distribute_parallel_for_simd,
                       getCollapseNumberExpr(Clauses),
                       nullptr /*ordered not a clause on distribute*/, CS, *this,
                       *DSAStack, VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
for (int ThisCaptureLevel =
     // 1.2.2 OpenMP Language Terminology
     // Structured block - An executable statement with a single entry at the
     // top and a single exit at the bottom.
     // The point of exit cannot be a branch out of the structured block.
     // longjmp() and throw() must not violate the entry/exit criteria.
     CS->getCapturedDecl()->setNothrow();
   }
 
-  OMPLoopDirective::HelperExprs B;
+  OMPLoopBasedDirective::HelperExprs B;
   // In presence of clause 'collapse' with number of loops, it will
   // define the nested loops number.
   unsigned NestedLoopCount = checkOpenMPLoop(
       OMPD_target_teams_distribute_simd, getCollapseNumberExpr(Clauses),
       nullptr /*ordered not a clause on distribute*/, CS, *this, *DSAStack,
       VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
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) {
+  auto SizesClauses =
+      OMPExecutableDirective::getClausesOfKind<OMPSizesClause>(Clauses);
+  if (SizesClauses.empty()) {
+    // 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.
+  SmallVector<OMPLoopBasedDirective::HelperExprs, 4> LoopHelpers(NumLoops);
+  Stmt *Body = nullptr;
+  SmallVector<Stmt *, 4> OriginalInits;
+  if (!OMPLoopBasedDirective::doForAllLoops(
+          AStmt->IgnoreContainers(), /*TryImperfectlyNestedLoops=*/false,
+          NumLoops,
+          [this, &LoopHelpers, &Body, &OriginalInits](unsigned Cnt,
+                                                      Stmt *CurStmt) {
+            VarsWithInheritedDSAType TmpDSA;
+            unsigned SingleNumLoops =
+                checkOpenMPLoop(OMPD_tile, nullptr, nullptr, CurStmt, *this,
+                                *DSAStack, TmpDSA, LoopHelpers[Cnt]);
+            if (SingleNumLoops == 0)
+              return true;
+            assert(SingleNumLoops == 1 && "Expect single loop iteration space");
+            if (auto *For = dyn_cast<ForStmt>(CurStmt)) {
+              OriginalInits.push_back(For->getInit());
+              Body = For->getBody();
+            } else {
+              assert(isa<CXXForRangeStmt>(CurStmt) &&
+                     "Expected canonical for or range-based for loops.");
+              auto *CXXFor = cast<CXXForRangeStmt>(CurStmt);
+              OriginalInits.push_back(CXXFor->getBeginStmt());
+              Body = CXXFor->getBody();
+            }
+            return false;
+          }))
+    return StmtError();
+
+  // Delay tiling to when template is completely instantiated.
+  if (CurContext->isDependentContext())
+    return OMPTileDirective::Create(Context, StartLoc, EndLoc, Clauses,
+                                    NumLoops, AStmt, nullptr, nullptr);
+
+  // Collection of generated variable declaration.
+  SmallVector<Decl *, 4> PreInits;
+
+  // Create iteration variables for the generated loops.
+  SmallVector<VarDecl *, 4> FloorIndVars;
+  SmallVector<VarDecl *, 4> TileIndVars;
+  FloorIndVars.resize(NumLoops);
+  TileIndVars.resize(NumLoops);
+  for (unsigned I = 0; I < NumLoops; ++I) {
+    OMPLoopBasedDirective::HelperExprs &LoopHelper = LoopHelpers[I];
+    if (auto *PI = cast_or_null<DeclStmt>(LoopHelper.PreInits))
+      PreInits.append(PI->decl_begin(), PI->decl_end());
+    assert(LoopHelper.Counters.size() == 1 &&
+           "Expect single-dimensional loop iteration space");
+    auto *OrigCntVar = cast<DeclRefExpr>(LoopHelper.Counters.front());
+    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;
+    }
+
+    // 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;
+    }
+    if (auto *PI = dyn_cast_or_null<DeclStmt>(OriginalInits[I]))
+      PreInits.append(PI->decl_begin(), PI->decl_end());
+    // Gather declarations for the data members used as counters.
+    for (Expr *CounterRef : LoopHelper.Counters) {
+      auto *CounterDecl = cast<DeclRefExpr>(CounterRef)->getDecl();
+      if (isa<OMPCapturedExprDecl>(CounterDecl))
+        PreInits.push_back(CounterDecl);
+    }
+  }
+
+  // Once the original iteration values are set, append the innermost body.
+  Stmt *Inner = Body;
+
+  // Create tile loops from the inside to the outside.
+  for (int I = NumLoops - 1; I >= 0; --I) {
+    OMPLoopBasedDirective::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,
+                                           OrigCntVar->getExprLoc());
+    DeclRefExpr *FloorIV = buildDeclRefExpr(*this, FloorIndVars[I], CntTy,
+                                            OrigCntVar->getExprLoc());
+
+    // For init-statement: auto .tile.iv = .floor.iv
+    AddInitializerToDecl(TileIndVars[I], DefaultLvalueConversion(FloorIV).get(),
+                         /*DirectInit=*/false);
+    Decl *CounterDecl = TileIndVars[I];
+    StmtResult InitStmt = new (Context)
+        DeclStmt(DeclGroupRef::Create(Context, &CounterDecl, 1),
+                 OrigCntVar->getBeginLoc(), OrigCntVar->getEndLoc());
+    if (!InitStmt.isUsable())
+      return StmtError();
+
+    // For cond-expression: .tile.iv < min(.floor.iv + DimTileSize,
+    // NumIterations)
+    ExprResult EndOfTile = BuildBinOp(CurScope, LoopHelper.Cond->getExprLoc(),
+                                      BO_Add, FloorIV, DimTileSize);
+    if (!EndOfTile.isUsable())
+      return StmtError();
+    ExprResult IsPartialTile =
+        BuildBinOp(CurScope, LoopHelper.Cond->getExprLoc(), BO_LT,
+                   NumIterations, EndOfTile.get());
+    if (!IsPartialTile.isUsable())
+      return StmtError();
+    ExprResult MinTileAndIterSpace = ActOnConditionalOp(
+        LoopHelper.Cond->getBeginLoc(), LoopHelper.Cond->getEndLoc(),
+        IsPartialTile.get(), NumIterations, EndOfTile.get());
+    if (!MinTileAndIterSpace.isUsable())
+      return StmtError();
+    ExprResult CondExpr = BuildBinOp(CurScope, LoopHelper.Cond->getExprLoc(),
+                                     BO_LT, TileIV, MinTileAndIterSpace.get());
+    if (!CondExpr.isUsable())
+      return StmtError();
+
+    // For incr-statement: ++.tile.iv
+    ExprResult IncrStmt =
+        BuildUnaryOp(CurScope, LoopHelper.Inc->getExprLoc(), UO_PreInc, TileIV);
+    if (!IncrStmt.isUsable())
+      return StmtError();
+
+    // Statements to set the original iteration variable's value from the
+    // logical iteration number.
+    // Generated for loop is:
+    // Original_for_init;
+    // for (auto .tile.iv = .floor.iv; .tile.iv < min(.floor.iv + DimTileSize,
+    // NumIterations); ++.tile.iv) {
+    //   Original_Body;
+    //   Original_counter_update;
+    // }
+    // 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.
+    SmallVector<Stmt *, 4> BodyParts;
+    BodyParts.append(LoopHelper.Updates.begin(), LoopHelper.Updates.end());
+    BodyParts.push_back(Inner);
+    Inner = CompoundStmt::Create(Context, BodyParts, Inner->getBeginLoc(),
+                                 Inner->getEndLoc());
+    Inner = new (Context)
+        ForStmt(Context, InitStmt.get(), CondExpr.get(), nullptr,
+                IncrStmt.get(), Inner, LoopHelper.Init->getBeginLoc(),
+                LoopHelper.Init->getBeginLoc(), LoopHelper.Inc->getEndLoc());
+  }
+
+  // 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,
+                                            OrigCntVar->getExprLoc());
+
+    // For init-statement: auto .floor.iv = 0
+    AddInitializerToDecl(
+        FloorIndVars[I],
+        ActOnIntegerConstant(LoopHelper.Init->getExprLoc(), 0).get(),
+        /*DirectInit=*/false);
+    Decl *CounterDecl = FloorIndVars[I];
+    StmtResult InitStmt = new (Context)
+        DeclStmt(DeclGroupRef::Create(Context, &CounterDecl, 1),
+                 OrigCntVar->getBeginLoc(), OrigCntVar->getEndLoc());
+    if (!InitStmt.isUsable())
+      return StmtError();
+
+    // For cond-expression: .floor.iv < NumIterations
+    ExprResult CondExpr = BuildBinOp(CurScope, LoopHelper.Cond->getExprLoc(),
+                                     BO_LT, FloorIV, NumIterations);
+    if (!CondExpr.isUsable())
+      return StmtError();
+
+    // For incr-statement: .floor.iv += DimTileSize
+    ExprResult IncrStmt = BuildBinOp(CurScope, LoopHelper.Inc->getExprLoc(),
+                                     BO_AddAssign, FloorIV, DimTileSize);
+    if (!IncrStmt.isUsable())
+      return StmtError();
+
+    Inner = new (Context)
+        ForStmt(Context, InitStmt.get(), CondExpr.get(), nullptr,
+                IncrStmt.get(), Inner, LoopHelper.Init->getBeginLoc(),
+                LoopHelper.Init->getBeginLoc(), LoopHelper.Inc->getEndLoc());
+  }
+
+  return OMPTileDirective::Create(Context, StartLoc, EndLoc, Clauses, NumLoops,
+                                  AStmt, Inner,
+                                  buildPreInits(Context, PreInits));
+}
+
 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_linear:
   case OMPC_aligned:
   case OMPC_copyin:
   case OMPC_copyprivate:
   case OMPC_nowait:
   case OMPC_untied:
   case OMPC_mergeable:
   case OMPC_threadprivate:
+  case OMPC_sizes:
   case OMPC_allocate:
   case OMPC_flush:
   case OMPC_read:
   case OMPC_write:
   case OMPC_update:
   case OMPC_capture:
   case OMPC_seq_cst:
   case OMPC_acq_rel:
case OMPC_if:
     case OMPD_declare_mapper:
     case OMPD_declare_simd:
     case OMPD_declare_variant:
     case OMPD_begin_declare_variant:
     case OMPD_end_declare_variant:
     case OMPD_declare_target:
     case OMPD_end_declare_target:
     case OMPD_teams:
+    case OMPD_tile:
     case OMPD_for:
     case OMPD_sections:
     case OMPD_section:
     case OMPD_single:
     case OMPD_master:
     case OMPD_critical:
     case OMPD_taskgroup:
     case OMPD_distribute:
case OMPC_num_threads:
     case OMPD_declare_simd:
     case OMPD_declare_variant:
     case OMPD_begin_declare_variant:
     case OMPD_end_declare_variant:
     case OMPD_declare_target:
     case OMPD_end_declare_target:
     case OMPD_teams:
     case OMPD_simd:
+    case OMPD_tile:
     case OMPD_for:
     case OMPD_for_simd:
     case OMPD_sections:
     case OMPD_section:
     case OMPD_single:
     case OMPD_master:
     case OMPD_critical:
     case OMPD_taskgroup:
case OMPC_num_teams:
     case OMPD_declare_mapper:
     case OMPD_declare_simd:
     case OMPD_declare_variant:
     case OMPD_begin_declare_variant:
     case OMPD_end_declare_variant:
     case OMPD_declare_target:
     case OMPD_end_declare_target:
     case OMPD_simd:
+    case OMPD_tile:
     case OMPD_for:
     case OMPD_for_simd:
     case OMPD_sections:
     case OMPD_section:
     case OMPD_single:
     case OMPD_master:
     case OMPD_critical:
     case OMPD_taskgroup:
case OMPC_thread_limit:
     case OMPD_declare_mapper:
     case OMPD_declare_simd:
     case OMPD_declare_variant:
     case OMPD_begin_declare_variant:
     case OMPD_end_declare_variant:
     case OMPD_declare_target:
     case OMPD_end_declare_target:
     case OMPD_simd:
+    case OMPD_tile:
     case OMPD_for:
     case OMPD_for_simd:
     case OMPD_sections:
     case OMPD_section:
     case OMPD_single:
     case OMPD_master:
     case OMPD_critical:
     case OMPD_taskgroup:
case OMPC_schedule:
     case OMPD_declare_mapper:
     case OMPD_declare_simd:
     case OMPD_declare_variant:
     case OMPD_begin_declare_variant:
     case OMPD_end_declare_variant:
     case OMPD_declare_target:
     case OMPD_end_declare_target:
     case OMPD_simd:
+    case OMPD_tile:
     case OMPD_sections:
     case OMPD_section:
     case OMPD_single:
     case OMPD_master:
     case OMPD_critical:
     case OMPD_taskgroup:
     case OMPD_distribute:
     case OMPD_ordered:
case OMPC_dist_schedule:
     case OMPD_declare_mapper:
     case OMPD_declare_simd:
     case OMPD_declare_variant:
     case OMPD_begin_declare_variant:
     case OMPD_end_declare_variant:
     case OMPD_declare_target:
     case OMPD_end_declare_target:
     case OMPD_simd:
+    case OMPD_tile:
     case OMPD_for:
     case OMPD_for_simd:
     case OMPD_sections:
     case OMPD_section:
     case OMPD_single:
     case OMPD_master:
     case OMPD_critical:
     case OMPD_taskgroup:
case OMPC_device:
     case OMPD_declare_mapper:
     case OMPD_declare_simd:
     case OMPD_declare_variant:
     case OMPD_begin_declare_variant:
     case OMPD_end_declare_variant:
     case OMPD_declare_target:
     case OMPD_end_declare_target:
     case OMPD_simd:
+    case OMPD_tile:
     case OMPD_for:
     case OMPD_for_simd:
     case OMPD_sections:
     case OMPD_section:
     case OMPD_single:
     case OMPD_master:
     case OMPD_critical:
     case OMPD_taskgroup:
case OMPC_priority:
     case OMPD_declare_mapper:
     case OMPD_declare_simd:
     case OMPD_declare_variant:
     case OMPD_begin_declare_variant:
     case OMPD_end_declare_variant:
     case OMPD_declare_target:
     case OMPD_end_declare_target:
     case OMPD_simd:
+    case OMPD_tile:
     case OMPD_for:
     case OMPD_for_simd:
     case OMPD_sections:
     case OMPD_section:
     case OMPD_single:
     case OMPD_master:
     case OMPD_critical:
     case OMPD_taskgroup:
static OpenMPDirectiveKind getOpenMPCaptureRegionForClause(
   case OMPC_reduction:
   case OMPC_task_reduction:
   case OMPC_in_reduction:
   case OMPC_linear:
   case OMPC_default:
   case OMPC_proc_bind:
   case OMPC_safelen:
   case OMPC_simdlen:
+  case OMPC_sizes:
   case OMPC_allocator:
   case OMPC_collapse:
   case OMPC_private:
   case OMPC_shared:
   case OMPC_aligned:
   case OMPC_copyin:
   case OMPC_copyprivate:
   case OMPC_ordered:
case OMPC_update:
     Res = ActOnOpenMPUpdateClause(static_cast<OpenMPDependClauseKind>(Argument),
                                   ArgumentLoc, StartLoc, LParenLoc, EndLoc);
     break;
   case OMPC_if:
   case OMPC_final:
   case OMPC_num_threads:
   case OMPC_safelen:
   case OMPC_simdlen:
+  case OMPC_sizes:
   case OMPC_allocator:
   case OMPC_collapse:
   case OMPC_schedule:
   case OMPC_private:
   case OMPC_firstprivate:
   case OMPC_lastprivate:
   case OMPC_shared:
   case OMPC_reduction:
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 (Expr *SizeExpr : SizeExprs) {
+    ExprResult NumForLoopsResult = VerifyPositiveIntegerConstantInClause(
+        SizeExpr, OMPC_sizes, /*StrictlyPositive=*/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:
case OMPC_device:
     Res = ActOnOpenMPDeviceClause(
         static_cast<OpenMPDeviceClauseModifier>(Argument.back()), Expr,
         StartLoc, LParenLoc, ArgumentLoc.back(), EndLoc);
     break;
   case OMPC_final:
   case OMPC_num_threads:
   case OMPC_safelen:
   case OMPC_simdlen:
+  case OMPC_sizes:
   case OMPC_allocator:
   case OMPC_collapse:
   case OMPC_default:
   case OMPC_proc_bind:
   case OMPC_private:
   case OMPC_firstprivate:
   case OMPC_lastprivate:
   case OMPC_shared:
OMPClause *Sema::ActOnOpenMPClause(OpenMPClauseKind Kind,
   case OMPC_destroy:
     Res = ActOnOpenMPDestroyClause(StartLoc, EndLoc);
     break;
   case OMPC_if:
   case OMPC_final:
   case OMPC_num_threads:
   case OMPC_safelen:
   case OMPC_simdlen:
+  case OMPC_sizes:
   case OMPC_allocator:
   case OMPC_collapse:
   case OMPC_schedule:
   case OMPC_private:
   case OMPC_firstprivate:
   case OMPC_lastprivate:
   case OMPC_shared:
   case OMPC_reduction:
Res = ActOnOpenMPAffinityClause(StartLoc, LParenLoc, ColonLoc, EndLoc,
                                     DepModOrTailExpr, VarList);
     break;
   case OMPC_if:
   case OMPC_depobj:
   case OMPC_final:
   case OMPC_num_threads:
   case OMPC_safelen:
   case OMPC_simdlen:
+  case OMPC_sizes:
   case OMPC_allocator:
   case OMPC_collapse:
   case OMPC_default:
   case OMPC_proc_bind:
   case OMPC_schedule:
   case OMPC_ordered:
   case OMPC_nowait:
   case OMPC_untied:

clang/lib/Sema/TreeTransform.h

#include "llvm/Frontend/OpenMP/OMP.inc"
   /// By default, performs semantic analysis to build the new OpenMP clause.
   /// Subclasses may override this routine to provide different behavior.
   OMPClause *RebuildOMPSimdlenClause(Expr *Len, SourceLocation StartLoc,
                                      SourceLocation LParenLoc,
                                      SourceLocation EndLoc) {
     return getSema().ActOnOpenMPSimdlenClause(Len, StartLoc, LParenLoc, EndLoc);
   }
 
+  OMPClause *RebuildOMPSizesClause(ArrayRef<Expr *> Sizes,
+                                   SourceLocation StartLoc,
+                                   SourceLocation LParenLoc,
+                                   SourceLocation EndLoc) {
+    return getSema().ActOnOpenMPSizesClause(Sizes, StartLoc, LParenLoc, EndLoc);
+  }
+
   /// Build a new OpenMP 'allocator' clause.
   ///
   /// By default, performs semantic analysis to build the new OpenMP clause.
   /// Subclasses may override this routine to provide different behavior.
   OMPClause *RebuildOMPAllocatorClause(Expr *A, SourceLocation StartLoc,
                                        SourceLocation LParenLoc,
                                        SourceLocation EndLoc) {
     return getSema().ActOnOpenMPAllocatorClause(A, StartLoc, LParenLoc, EndLoc);
StmtResult Body;
       Sema::CompoundScopeRAII CompoundScope(getSema());
       Stmt *CS;
       if (D->getDirectiveKind() == OMPD_atomic ||
           D->getDirectiveKind() == OMPD_critical ||
           D->getDirectiveKind() == OMPD_section ||
           D->getDirectiveKind() == OMPD_master)
         CS = D->getAssociatedStmt();
       else
-        CS = D->getInnermostCapturedStmt()->getCapturedStmt();
+        CS = D->getRawStmt();
       Body = getDerived().TransformStmt(CS);
     }
     AssociatedStmt =
         getDerived().getSema().ActOnOpenMPRegionEnd(Body, TClauses);
     if (AssociatedStmt.isInvalid()) {
       return StmtError();
     }
   }
getDerived().getSema().StartOpenMPDSABlock(OMPD_simd, DirName, nullptr,
                                              D->getBeginLoc());
   StmtResult 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;
+}
+
+template <typename Derived>
+StmtResult
 TreeTransform<Derived>::TransformOMPForDirective(OMPForDirective *D) {
   DeclarationNameInfo DirName;
   getDerived().getSema().StartOpenMPDSABlock(OMPD_for, DirName, nullptr,
                                              D->getBeginLoc());
   StmtResult Res = getDerived().TransformOMPExecutableDirective(D);
   getDerived().getSema().EndOpenMPDSABlock(Res.get());
   return Res;
 }
TreeTransform<Derived>::TransformOMPSimdlenClause(OMPSimdlenClause *C) {
   ExprResult E = getDerived().TransformExpr(C->getSimdlen());
   if (E.isInvalid())
     return nullptr;
   return getDerived().RebuildOMPSimdlenClause(
       E.get(), 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 (Expr *E : C->getSizesRefs()) {
+    if (!E) {
+      TransformedSizes.push_back(nullptr);
+      continue;
+    }
+
+    ExprResult 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());
+}
+
+template <typename Derived>
 OMPClause *
 TreeTransform<Derived>::TransformOMPCollapseClause(OMPCollapseClause *C) {
   ExprResult E = getDerived().TransformExpr(C->getNumForLoops());
   if (E.isInvalid())
     return nullptr;
   return getDerived().RebuildOMPCollapseClause(
       E.get(), C->getBeginLoc(), C->getLParenLoc(), C->getEndLoc());
 }

clang/lib/Serialization/ASTReader.cpp

case llvm::omp::OMPC_num_threads:
     C = new (Context) OMPNumThreadsClause();
     break;
   case llvm::omp::OMPC_safelen:
     C = new (Context) OMPSafelenClause();
     break;
   case llvm::omp::OMPC_simdlen:
     C = new (Context) OMPSimdlenClause();
     break;
+  case llvm::omp::OMPC_sizes: {
+    unsigned NumSizes = Record.readInt();
+    C = OMPSizesClause::CreateEmpty(Context, NumSizes);
+    break;
+  }
   case llvm::omp::OMPC_allocator:
     C = new (Context) OMPAllocatorClause();
     break;
   case llvm::omp::OMPC_collapse:
     C = new (Context) OMPCollapseClause();
     break;
   case llvm::omp::OMPC_default:
     C = new (Context) OMPDefaultClause();
void OMPClauseReader::VisitOMPSafelenClause(OMPSafelenClause *C) {
   C->setLParenLoc(Record.readSourceLocation());
 }
 
 void OMPClauseReader::VisitOMPSimdlenClause(OMPSimdlenClause *C) {
   C->setSimdlen(Record.readSubExpr());
   C->setLParenLoc(Record.readSourceLocation());
 }
 
+void OMPClauseReader::VisitOMPSizesClause(OMPSizesClause *C) {
+  for (Expr *&E : C->getSizesRefs())
+    E = Record.readSubExpr();
+  C->setLParenLoc(Record.readSourceLocation());
+}
+
 void OMPClauseReader::VisitOMPAllocatorClause(OMPAllocatorClause *C) {
   C->setAllocator(Record.readExpr());
   C->setLParenLoc(Record.readSourceLocation());
 }
 
 void OMPClauseReader::VisitOMPCollapseClause(OMPCollapseClause *C) {
   C->setNumForLoops(Record.readSubExpr());
   C->setLParenLoc(Record.readSourceLocation());

clang/lib/Serialization/ASTReaderStmt.cpp

 //===----------------------------------------------------------------------===//
 
 void ASTStmtReader::VisitOMPExecutableDirective(OMPExecutableDirective *E) {
   Record.readOMPChildren(E->Data);
   E->setLocStart(readSourceLocation());
   E->setLocEnd(readSourceLocation());
 }
 
-void ASTStmtReader::VisitOMPLoopDirective(OMPLoopDirective *D) {
+void ASTStmtReader::VisitOMPLoopBasedDirective(OMPLoopBasedDirective *D) {
   VisitStmt(D);
   // Field CollapsedNum was read in ReadStmtFromStream.
   Record.skipInts(1);
   VisitOMPExecutableDirective(D);
 }
 
+void ASTStmtReader::VisitOMPLoopDirective(OMPLoopDirective *D) {
+  VisitOMPLoopBasedDirective(D);
+}
+
 void ASTStmtReader::VisitOMPParallelDirective(OMPParallelDirective *D) {
   VisitStmt(D);
   VisitOMPExecutableDirective(D);
   D->setHasCancel(Record.readBool());
 }
 
 void ASTStmtReader::VisitOMPSimdDirective(OMPSimdDirective *D) {
   VisitOMPLoopDirective(D);
 }
 
+void ASTStmtReader::VisitOMPTileDirective(OMPTileDirective *D) {
+  VisitOMPLoopBasedDirective(D);
+}
+
 void ASTStmtReader::VisitOMPForDirective(OMPForDirective *D) {
   VisitOMPLoopDirective(D);
   D->setHasCancel(Record.readBool());
 }
 
 void ASTStmtReader::VisitOMPForSimdDirective(OMPForSimdDirective *D) {
   VisitOMPLoopDirective(D);
 }
while (true) {
     case STMT_OMP_SIMD_DIRECTIVE: {
       unsigned CollapsedNum = Record[ASTStmtReader::NumStmtFields];
       unsigned NumClauses = Record[ASTStmtReader::NumStmtFields + 1];
       S = OMPSimdDirective::CreateEmpty(Context, NumClauses,
                                         CollapsedNum, Empty);
       break;
     }
 
+    case STMT_OMP_TILE_DIRECTIVE: {
+      unsigned NumLoops = Record[ASTStmtReader::NumStmtFields];
+      unsigned NumClauses = Record[ASTStmtReader::NumStmtFields + 1];
+      S = OMPTileDirective::CreateEmpty(Context, NumClauses, NumLoops);
+      break;
+    }
+
     case STMT_OMP_FOR_DIRECTIVE: {
       unsigned CollapsedNum = Record[ASTStmtReader::NumStmtFields];
       unsigned NumClauses = Record[ASTStmtReader::NumStmtFields + 1];
       S = OMPForDirective::CreateEmpty(Context, NumClauses, CollapsedNum,
                                        Empty);
       break;
     }
 

clang/lib/Serialization/ASTWriter.cpp

void OMPClauseWriter::VisitOMPSafelenClause(OMPSafelenClause *C) {
   Record.AddSourceLocation(C->getLParenLoc());
 }
 
 void OMPClauseWriter::VisitOMPSimdlenClause(OMPSimdlenClause *C) {
   Record.AddStmt(C->getSimdlen());
   Record.AddSourceLocation(C->getLParenLoc());
 }
 
+void OMPClauseWriter::VisitOMPSizesClause(OMPSizesClause *C) {
+  Record.push_back(C->getNumSizes());
+  for (Expr *Size : C->getSizesRefs())
+    Record.AddStmt(Size);
+  Record.AddSourceLocation(C->getLParenLoc());
+}
+
 void OMPClauseWriter::VisitOMPAllocatorClause(OMPAllocatorClause *C) {
   Record.AddStmt(C->getAllocator());
   Record.AddSourceLocation(C->getLParenLoc());
 }
 
 void OMPClauseWriter::VisitOMPCollapseClause(OMPCollapseClause *C) {
   Record.AddStmt(C->getNumForLoops());
   Record.AddSourceLocation(C->getLParenLoc());

clang/lib/Serialization/ASTWriterStmt.cpp

 //===----------------------------------------------------------------------===//
 
 void ASTStmtWriter::VisitOMPExecutableDirective(OMPExecutableDirective *E) {
   Record.writeOMPChildren(E->Data);
   Record.AddSourceLocation(E->getBeginLoc());
   Record.AddSourceLocation(E->getEndLoc());
 }
 
-void ASTStmtWriter::VisitOMPLoopDirective(OMPLoopDirective *D) {
+void ASTStmtWriter::VisitOMPLoopBasedDirective(OMPLoopBasedDirective *D) {
   VisitStmt(D);
   Record.writeUInt32(D->getCollapsedNumber());
   VisitOMPExecutableDirective(D);
 }
 
+void ASTStmtWriter::VisitOMPLoopDirective(OMPLoopDirective *D) {
+  VisitOMPLoopBasedDirective(D);
+}
+
 void ASTStmtWriter::VisitOMPParallelDirective(OMPParallelDirective *D) {
   VisitStmt(D);
   VisitOMPExecutableDirective(D);
   Record.writeBool(D->hasCancel());
   Code = serialization::STMT_OMP_PARALLEL_DIRECTIVE;
 }
 
 void ASTStmtWriter::VisitOMPSimdDirective(OMPSimdDirective *D) {
   VisitOMPLoopDirective(D);
   Code = serialization::STMT_OMP_SIMD_DIRECTIVE;
 }
 
+void ASTStmtWriter::VisitOMPTileDirective(OMPTileDirective *D) {
+  VisitOMPLoopBasedDirective(D);
+  Code = serialization::STMT_OMP_TILE_DIRECTIVE;
+}
+
 void ASTStmtWriter::VisitOMPForDirective(OMPForDirective *D) {
   VisitOMPLoopDirective(D);
   Record.writeBool(D->hasCancel());
   Code = serialization::STMT_OMP_FOR_DIRECTIVE;
 }
 
 void ASTStmtWriter::VisitOMPForSimdDirective(OMPForSimdDirective *D) {
   VisitOMPLoopDirective(D);

clang/lib/StaticAnalyzer/Core/ExprEngine.cpp

switch (S->getStmtClass()) {
     case Stmt::OMPTeamsDistributeSimdDirectiveClass:
     case Stmt::OMPTeamsDistributeParallelForSimdDirectiveClass:
     case Stmt::OMPTeamsDistributeParallelForDirectiveClass:
     case Stmt::OMPTargetTeamsDirectiveClass:
     case Stmt::OMPTargetTeamsDistributeDirectiveClass:
     case Stmt::OMPTargetTeamsDistributeParallelForDirectiveClass:
     case Stmt::OMPTargetTeamsDistributeParallelForSimdDirectiveClass:
     case Stmt::OMPTargetTeamsDistributeSimdDirectiveClass:
+    case Stmt::OMPTileDirectiveClass:
     case Stmt::CapturedStmtClass: {
       const ExplodedNode *node = Bldr.generateSink(S, Pred, Pred->getState());
       Engine.addAbortedBlock(node, currBldrCtx->getBlock());
       break;
     }
 
     case Stmt::ParenExprClass:
       llvm_unreachable("ParenExprs already handled.");

clang/test/Index/openmp-tile.c

+// RUN: c-index-test -test-load-source local %s -fopenmp -fopenmp-version=51 | 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:  value: Int 3
+  // 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:      value: Int 2
+  // 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 code generation
+// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp -fopenmp-version=51 -emit-llvm %s -o - | FileCheck %s --check-prefix=IR
+
+// Check same results after serialization round-trip
+// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp -fopenmp-version=51 -emit-pch -o %t %s
+// RUN: %clang_cc1 -verify -triple x86_64-pc-linux-gnu -fopenmp -fopenmp-version=51 -include-pch %t -emit-llvm %s -o - | FileCheck %s --check-prefix=IR
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+// placeholder for loop body code.
+extern "C" void body(...) {}
+
+// IR: define {{.*}}void @_ZN1SC2Ev(%struct.S*
+// IR:         [[THIS_ADDR:%.+]] = alloca %struct.S*, align 8
+// IR-NEXT:    [[I_REF:%.+]] = alloca i32*, align 8
+// IR-NEXT:    [[FLOOR:%.+]] = alloca i32, align 4
+// IR-NEXT:    [[TILE:%.+]] = alloca i32, align 4
+// IR-NEXT:    store %struct.S* %{{.+}}, %struct.S** [[THIS_ADDR]], align 8
+// IR-NEXT:    [[THIS:%.+]] = load %struct.S*, %struct.S** [[THIS_ADDR]], align 8
+// IR-NEXT:    [[I:%.+]] = getelementptr inbounds %struct.S, %struct.S* [[THIS]], i32 0, i32 0
+// IR-NEXT:    store i32* [[I]], i32** [[I_REF]], align 8
+// IR-NEXT:    store i32 0, i32* [[FLOOR]], align 4
+// IR-NEXT:    br label %[[FOR_COND:.+]]
+// IR:         [[FOR_COND]]:
+// IR-NEXT:    [[TMP0:%.+]] = load i32, i32* [[FLOOR]], align 4
+// IR-NEXT:    [[CMP:%.+]] = icmp slt i32 [[TMP0]], 4
+// IR-NEXT:    br i1 [[CMP]], label %[[FOR_BODY:.+]], label %[[FOR_END11:.+]]
+// IR:         [[FOR_BODY]]:
+// IR-NEXT:    [[TMP1:%.+]] = load i32, i32* [[FLOOR]], align 4
+// IR-NEXT:    store i32 [[TMP1]], i32* [[TILE]], align 4
+// IR-NEXT:    br label %[[FOR_COND3:.+]]
+// IR:         [[FOR_COND3]]:
+// IR-NEXT:    [[TMP2:%.+]] = load i32, i32* [[TILE]], align 4
+// IR-NEXT:    [[TMP3:%.+]] = load i32, i32* [[FLOOR]], align 4
+// IR-NEXT:    [[ADD:%.+]] = add nsw i32 [[TMP3]], 5
+// IR-NEXT:    [[CMP4:%.+]] = icmp slt i32 4, [[ADD]]
+// IR-NEXT:    br i1 [[CMP4]], label %[[COND_TRUE:.+]], label %[[COND_FALSE:.+]]
+// IR:         [[COND_TRUE]]:
+// IR-NEXT:    br label %[[COND_END:.+]]
+// IR:         [[COND_FALSE]]:
+// IR-NEXT:    [[TMP4:%.+]] = load i32, i32* [[FLOOR]], align 4
+// IR-NEXT:    [[ADD5:%.+]] = add nsw i32 [[TMP4]], 5
+// IR-NEXT:    br label %[[COND_END]]
+// IR:         [[COND_END]]:
+// IR-NEXT:    [[COND:%.+]] = phi i32 [ 4, %[[COND_TRUE]] ], [ [[ADD5]], %[[COND_FALSE]] ]
+// IR-NEXT:    [[CMP6:%.+]] = icmp slt i32 [[TMP2]], [[COND]]
+// IR-NEXT:    br i1 [[CMP6]], label %[[FOR_BODY7:.+]], label %[[FOR_END:.+]]
+// IR:         [[FOR_BODY7]]:
+// IR-NEXT:    [[TMP5:%.+]] = load i32, i32* [[TILE]], align 4
+// IR-NEXT:    [[MUL:%.+]] = mul nsw i32 [[TMP5]], 3
+// IR-NEXT:    [[ADD8:%.+]] = add nsw i32 7, [[MUL]]
+// IR-NEXT:    [[TMP6:%.+]] = load i32*, i32** [[I_REF]], align 8
+// IR-NEXT:    store i32 [[ADD8]], i32* [[TMP6]], align 4
+// IR-NEXT:    [[TMP7:%.+]] = load i32*, i32** [[I_REF]], align 8
+// IR-NEXT:    [[TMP8:%.+]] = load i32, i32* [[TMP7]], align 4
+// IR-NEXT:    call void (...) @body(i32 [[TMP8]])
+// IR-NEXT:    br label %[[FOR_INC:.+]]
+// IR:         [[FOR_INC]]:
+// IR-NEXT:    [[TMP9:%.+]] = load i32, i32* [[TILE]], align 4
+// IR-NEXT:    [[INC:%.+]] = add nsw i32 [[TMP9]], 1
+// IR-NEXT:    store i32 [[INC]], i32* [[TILE]], align 4
+// IR-NEXT:    br label %[[FOR_COND3]]
+// IR:         [[FOR_END]]:
+// IR-NEXT:    br label %[[FOR_INC9:.+]]
+// IR:         [[FOR_INC9]]:
+// IR-NEXT:    [[TMP10:%.+]] = load i32, i32* [[FLOOR]], align 4
+// IR-NEXT:    [[ADD10:%.+]] = add nsw i32 [[TMP10]], 5
+// IR-NEXT:    store i32 [[ADD10]], i32* [[FLOOR]], align 4
+// IR-NEXT:    br label %[[FOR_COND]]
+// IR:         [[FOR_END11]]:
+// IR-NEXT:    ret void
+
+struct S {
+  int i;
+  S() {
+#pragma omp tile sizes(5)
+    for (i = 7; i < 17; i += 3)
+      body(i);
+  }
+} s;
+
+// IR-LABEL: define {{.*}}void @foo1(
+// IR:         [[START_ADDR:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[END_ADDR:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[STEP_ADDR:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[CAP_EXPR:%.+]] = alloca i32, align 4
+// IR-NEXT:    [[CAP_EXPR1:%.+]] = alloca i32, align 4
+// IR-NEXT:    [[CAP_EXPR2:%.+]] = alloca i32, align 4
+// IR-NEXT:    [[CAP_EXPR3:%.+]] = 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:    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:    [[TMP0:%.+]] = load i32, i32* [[START_ADDR]], align 4
+// IR-NEXT:    store i32 [[TMP0]], i32* [[CAP_EXPR]], align 4
+// IR-NEXT:    [[TMP1:%.+]] = load i32, i32* [[END_ADDR]], align 4
+// IR-NEXT:    store i32 [[TMP1]], i32* [[CAP_EXPR1]], align 4
+// IR-NEXT:    [[TMP2:%.+]] = load i32, i32* [[STEP_ADDR]], align 4
+// IR-NEXT:    store i32 [[TMP2]], i32* [[CAP_EXPR2]], align 4
+// IR-NEXT:    [[TMP3:%.+]] = load i32, i32* [[CAP_EXPR1]], align 4
+// IR-NEXT:    [[TMP4:%.+]] = load i32, i32* [[CAP_EXPR]], align 4
+// IR-NEXT:    [[SUB:%.+]] = sub i32 [[TMP3]], [[TMP4]]
+// IR-NEXT:    [[SUB4:%.+]] = sub i32 [[SUB]], 1
+// IR-NEXT:    [[TMP5:%.+]] = load i32, i32* [[CAP_EXPR2]], align 4
+// IR-NEXT:    [[ADD:%.+]] = add i32 [[SUB4]], [[TMP5]]
+// IR-NEXT:    [[TMP6:%.+]] = load i32, i32* [[CAP_EXPR2]], align 4
+// IR-NEXT:    [[DIV:%.+]] = udiv i32 [[ADD]], [[TMP6]]
+// IR-NEXT:    [[SUB5:%.+]] = sub i32 [[DIV]], 1
+// IR-NEXT:    store i32 [[SUB5]], i32* [[CAP_EXPR3]], 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:    [[TMP8:%.+]] = load i32, i32* [[CAP_EXPR3]], align 4
+// IR-NEXT:    [[ADD3:%.*]] = add i32 [[TMP8]], 1
+// IR-NEXT:    [[CMP:%.*]] = icmp ult 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:    [[TMP11:%.+]] = load i32, i32* [[CAP_EXPR3]], align 4
+// IR-NEXT:    [[ADD10:%.*]] = add i32 [[TMP11]], 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 ult i32 [[ADD10]], [[ADD11]]
+// IR-NEXT:    br i1 [[CMP12]], label %[[COND_TRUE:.*]], label %[[COND_FALSE:.*]]
+// IR:         [[COND_TRUE]]:
+// IR-NEXT:    [[TMP13:%.+]] = load i32, i32* [[CAP_EXPR3]], align 4
+// IR-NEXT:    [[ADD18:%.*]] = add i32 [[TMP13]], 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 ult i32 [[TMP6]], [[COND]]
+// IR-NEXT:    br i1 [[CMP20]], label %[[FOR_BODY21:.*]], label %[[FOR_END:.*]]
+// IR:         [[FOR_BODY21]]:
+// IR-NEXT:    [[TMP15:%.+]] = load i32, i32* [[CAP_EXPR]], align 4
+// IR-NEXT:    [[TMP19:%.*]] = load i32, i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    [[TMP17:%.+]] = load i32, i32* [[CAP_EXPR2]], align 4
+// IR-NEXT:    [[MUL:%.*]] = mul i32 [[TMP19]], [[TMP17]]
+// IR-NEXT:    [[ADD22:%.*]] = add i32 [[TMP15]], [[MUL]]
+// IR-NEXT:    store i32 [[ADD22]], i32* [[I]], align 4
+// IR-NEXT:    [[TMP21:%.*]] = load i32, i32* [[I]], align 4
+// IR-NEXT:    call void (...) @body(i32 [[TMP21]])
+// IR-NEXT:    br label %[[FOR_INC:.*]]
+// IR:         [[FOR_INC]]:
+// IR-NEXT:    [[TMP22:%.*]] = load i32, i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP22]], 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:    [[TMP23:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I]], align 4
+// IR-NEXT:    [[ADD24:%.*]] = add nsw i32 [[TMP23]], 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) {
+  int i;
+#pragma omp tile sizes(5)
+  for (i = start; i < end; i += step)
+    body(i);
+}
+
+// IR-LABEL: define {{.*}}void @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:    [[I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[J:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTFLOOR_0_IV_I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTFLOOR_1_IV_J:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTTILE_0_IV_I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTTILE_1_IV_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 7, i32* [[I]], align 4
+// IR-NEXT:    store i32 7, i32* [[J]], 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* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    [[MUL:%.+]] = mul nsw i32 [[TMP6]], 3
+// IR-NEXT:    [[ADD9:%.+]] = add nsw i32 7, [[MUL]]
+// IR-NEXT:    store i32 [[ADD9]], i32* [[I]], align 4
+// IR-NEXT:    [[TMP7:%.+]] = load i32, i32* [[DOTFLOOR_1_IV_J]], align 4
+// IR-NEXT:    store i32 [[TMP7]], i32* [[DOTTILE_1_IV_J]], align 4
+// IR-NEXT:    br label %[[FOR_COND10:.+]]
+// IR:         [[FOR_COND10]]:
+// 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:    [[TMP11:%.*]] = load i32, i32* [[DOTTILE_1_IV_J]], align 4
+// IR-NEXT:    [[MUL20:%.*]] = mul nsw i32 [[TMP11]], 3
+// IR-NEXT:    [[ADD21:%.*]] = add nsw i32 7, [[MUL20]]
+// IR-NEXT:    store i32 [[ADD21]], i32* [[J]], align 4
+// IR-NEXT:    [[TMP12:%.*]] = load i32, i32* [[I]], align 4
+// IR-NEXT:    [[TMP13:%.*]] = load i32, i32* [[J]], align 4
+// IR-NEXT:    call void (...) @body(i32 [[TMP12]], i32 [[TMP13]])
+// IR-NEXT:    br label %[[FOR_INC:.*]]
+// IR:         [[FOR_INC]]:
+// IR-NEXT:    [[TMP14:%.*]] = load i32, i32* [[DOTTILE_1_IV_J]], align 4
+// IR-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP14]], 1
+// IR-NEXT:    store i32 [[INC]], i32* [[DOTTILE_1_IV_J]], align 4
+// IR-NEXT:    br label %[[FOR_COND10]]
+// IR:         [[FOR_END]]:
+// IR-NEXT:    br label %[[FOR_INC22:.*]]
+// IR:         [[FOR_INC22]]:
+// IR-NEXT:    [[TMP15:%.*]] = load i32, i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    [[INC23:%.*]] = add nsw i32 [[TMP15]], 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:    [[TMP16:%.*]] = load i32, i32* [[DOTFLOOR_1_IV_J]], align 4
+// IR-NEXT:    [[ADD26:%.*]] = add nsw i32 [[TMP16]], 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:    [[TMP17:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I]], align 4
+// IR-NEXT:    [[ADD29:%.*]] = add nsw i32 [[TMP17]], 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:    [[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_I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[J:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTFLOOR_1_IV_J:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTTILE_0_IV_I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTTILE_1_IV_J:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* [[GLOB2:@.*]])
+// 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* [[GLOB1:@.*]], 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_I]], align 4
+// IR-NEXT:    store i32 7, i32* [[I]], align 4
+// IR-NEXT:    store i32 7, i32* [[J]], 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_I]], 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_I]], 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_I]], 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* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    [[MUL13:%.+]] = mul nsw i32 [[TMP12]], 3
+// IR-NEXT:    [[ADD14:%.+]] = add nsw i32 7, [[MUL13]]
+// IR-NEXT:    store i32 [[ADD14]], i32* [[I]], align 4
+// 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:    [[TMP17:%.*]] = load i32, i32* [[DOTTILE_1_IV_J]], align 4
+// IR-NEXT:    [[MUL26:%.*]] = mul nsw i32 [[TMP17]], 3
+// IR-NEXT:    [[ADD27:%.*]] = add nsw i32 7, [[MUL26]]
+// IR-NEXT:    store i32 [[ADD27]], i32* [[J]], align 4
+// IR-NEXT:    [[TMP18:%.*]] = load i32, i32* [[I]], align 4
+// IR-NEXT:    [[TMP19:%.*]] = load i32, i32* [[J]], align 4
+// IR-NEXT:    call void (...) @body(i32 [[TMP18]], i32 [[TMP19]])
+// 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* [[GLOB1]], i32 [[TMP0]])
+// IR-NEXT:    call void @__kmpc_barrier(%struct.ident_t* [[GLOB3:@.*]], 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)
+        body(i, j);
+}
+
+// IR-LABEL: @foo4(
+// IR-NEXT:  entry:
+// IR-NEXT:    [[DOTOMP_IV:%.*]] = 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_I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[J:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTFLOOR_1_IV_J:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTTILE_0_IV_I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTTILE_1_IV_J:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[TMP0:%.*]] = call i32 @__kmpc_global_thread_num(%struct.ident_t* [[GLOB2]])
+// 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* [[GLOB1]], 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_I]], align 4
+// IR-NEXT:    store i32 7, i32* [[I]], align 4
+// IR-NEXT:    store i32 7, i32* [[J]], 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_I]], 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_I]], 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_I]], 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* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    [[MUL18:%.+]] = mul nsw i32 [[TMP14]], 3
+// IR-NEXT:    [[ADD19:%.+]] = add nsw i32 7, [[MUL18]]
+// IR-NEXT:    store i32 [[ADD19]], i32* [[I]], align 4
+// 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:    [[TMP19:%.*]] = load i32, i32* [[DOTTILE_1_IV_J]], align 4
+// IR-NEXT:    [[MUL31:%.*]] = mul nsw i32 [[TMP19]], 3
+// IR-NEXT:    [[ADD32:%.*]] = add nsw i32 7, [[MUL31]]
+// IR-NEXT:    store i32 [[ADD32]], i32* [[J]], align 4
+// IR-NEXT:    [[TMP20:%.*]] = load i32, i32* [[I]], align 4
+// IR-NEXT:    [[TMP21:%.*]] = load i32, i32* [[J]], align 4
+// IR-NEXT:    call void (...) @body(i32 [[TMP20]], i32 [[TMP21]])
+// 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* [[GLOB1]], i32 [[TMP0]])
+// IR-NEXT:    call void @__kmpc_barrier(%struct.ident_t* [[GLOB3]], 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)
+      body(i, j);
+}
+
+
+// IR-LABEL: @foo5(
+// IR-NEXT:  entry:
+// IR-NEXT:    [[DOTOMP_IV:%.*]] = alloca i64, align 8
+// 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_I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTTILE_0_IV_I:%.*]] = 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_I10:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTTILE_0_IV_I11:%.*]] = 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* [[GLOB2]])
+// 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 i32 [[TMP4]], [[TMP5]]
+// IR-NEXT:    [[SUB6:%.*]] = sub i32 [[SUB]], 1
+// IR-NEXT:    [[ADD7:%.*]] = add i32 [[SUB6]], 1
+// IR-NEXT:    [[DIV:%.*]] = udiv i32 [[ADD7]], 1
+// IR-NEXT:    [[CONV:%.*]] = zext 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_I]], align 4
+// IR-NEXT:    [[TMP6:%.*]] = load i32, i32* [[DOTCAPTURE_EXPR_]], align 4
+// IR-NEXT:    store i32 [[TMP6]], i32* [[DOTTILE_0_IV_I]], 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* [[GLOB1]], 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 i32 [[TMP18]], [[TMP19]]
+// IR-NEXT:    [[SUB23:%.*]] = sub i32 [[SUB22]], 1
+// IR-NEXT:    [[ADD24:%.*]] = add i32 [[SUB23]], 1
+// IR-NEXT:    [[DIV25:%.*]] = udiv i32 [[ADD24]], 1
+// IR-NEXT:    [[MUL26:%.*]] = mul i32 1, [[DIV25]]
+// IR-NEXT:    [[MUL27:%.*]] = mul i32 [[MUL26]], 4
+// IR-NEXT:    [[CONV28:%.*]] = zext 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_I10]], 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 i32 [[TMP23]], [[TMP24]]
+// IR-NEXT:    [[SUB35:%.*]] = sub i32 [[SUB34]], 1
+// IR-NEXT:    [[ADD36:%.*]] = add i32 [[SUB35]], 1
+// IR-NEXT:    [[DIV37:%.*]] = udiv i32 [[ADD36]], 1
+// IR-NEXT:    [[MUL38:%.*]] = mul i32 1, [[DIV37]]
+// IR-NEXT:    [[MUL39:%.*]] = mul i32 [[MUL38]], 4
+// IR-NEXT:    [[CONV40:%.*]] = zext 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 i32 [[TMP25]], [[TMP26]]
+// IR-NEXT:    [[SUB43:%.*]] = sub i32 [[SUB42]], 1
+// IR-NEXT:    [[ADD44:%.*]] = add i32 [[SUB43]], 1
+// IR-NEXT:    [[DIV45:%.*]] = udiv i32 [[ADD44]], 1
+// IR-NEXT:    [[MUL46:%.*]] = mul i32 1, [[DIV45]]
+// IR-NEXT:    [[MUL47:%.*]] = mul i32 [[MUL46]], 4
+// IR-NEXT:    [[CONV48:%.*]] = zext 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_I11]], 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 i32 [[TMP29]], [[TMP30]]
+// IR-NEXT:    [[SUB56:%.*]] = sub i32 [[SUB55]], 1
+// IR-NEXT:    [[ADD57:%.*]] = add i32 [[SUB56]], 1
+// IR-NEXT:    [[DIV58:%.*]] = udiv i32 [[ADD57]], 1
+// IR-NEXT:    [[MUL59:%.*]] = mul i32 1, [[DIV58]]
+// IR-NEXT:    [[MUL60:%.*]] = mul i32 [[MUL59]], 4
+// IR-NEXT:    [[CONV61:%.*]] = zext 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 i32 [[TMP31]], [[TMP32]]
+// IR-NEXT:    [[SUB64:%.*]] = sub i32 [[SUB63]], 1
+// IR-NEXT:    [[ADD65:%.*]] = add i32 [[SUB64]], 1
+// IR-NEXT:    [[DIV66:%.*]] = udiv i32 [[ADD65]], 1
+// IR-NEXT:    [[MUL67:%.*]] = mul i32 1, [[DIV66]]
+// IR-NEXT:    [[MUL68:%.*]] = mul i32 [[MUL67]], 4
+// IR-NEXT:    [[CONV69:%.*]] = zext 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 i32 [[TMP35]], [[TMP36]]
+// IR-NEXT:    [[SUB73:%.*]] = sub i32 [[SUB72]], 1
+// IR-NEXT:    [[ADD74:%.*]] = add i32 [[SUB73]], 1
+// IR-NEXT:    [[DIV75:%.*]] = udiv i32 [[ADD74]], 1
+// IR-NEXT:    [[MUL76:%.*]] = mul i32 1, [[DIV75]]
+// IR-NEXT:    [[MUL77:%.*]] = mul i32 [[MUL76]], 4
+// IR-NEXT:    [[CONV78:%.*]] = zext 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 i32 [[TMP37]], [[TMP38]]
+// IR-NEXT:    [[SUB81:%.*]] = sub i32 [[SUB80]], 1
+// IR-NEXT:    [[ADD82:%.*]] = add i32 [[SUB81]], 1
+// IR-NEXT:    [[DIV83:%.*]] = udiv i32 [[ADD82]], 1
+// IR-NEXT:    [[MUL84:%.*]] = mul i32 1, [[DIV83]]
+// IR-NEXT:    [[MUL85:%.*]] = mul i32 [[MUL84]], 4
+// IR-NEXT:    [[CONV86:%.*]] = zext 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:    [[TMP39:%.*]] = load i32, i32* [[DOTTILE_0_IV_I11]], align 4
+// IR-NEXT:    [[MUL95:%.*]] = mul nsw i32 [[TMP39]], 3
+// IR-NEXT:    [[ADD96:%.*]] = add nsw i32 7, [[MUL95]]
+// IR-NEXT:    store i32 [[ADD96]], i32* [[I]], align 4
+// IR-NEXT:    [[TMP40:%.*]] = load i32, i32* [[I]], align 4
+// IR-NEXT:    [[TMP41:%.*]] = load i32, i32* [[J15]], align 4
+// IR-NEXT:    call void (...) @body(i32 [[TMP40]], i32 [[TMP41]])
+// 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:    [[TMP42:%.*]] = load i64, i64* [[DOTOMP_IV]], align 8
+// IR-NEXT:    [[ADD97:%.*]] = add nsw i64 [[TMP42]], 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* [[GLOB1]], i32 [[TMP0]])
+// IR-NEXT:    br label %[[OMP_PRECOND_END]]
+// IR:         [[OMP_PRECOND_END]]:
+// IR-NEXT:    call void @__kmpc_barrier(%struct.ident_t* [[GLOB3]], 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)
+      body(i, j);
+}
+
+
+// IR-LABEL: @foo6(
+// IR-NEXT:  entry:
+// IR-NEXT:    call void (%struct.ident_t*, i32, void (i32*, i32*, ...)*, ...) @__kmpc_fork_call(%struct.ident_t* [[GLOB2]], i32 0, void (i32*, i32*, ...)* bitcast (void (i32*, i32*)* @.omp_outlined. to void (i32*, i32*, ...)*))
+// IR-NEXT:    ret void
+//
+// IR-LABEL: @.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:    [[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_I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[I:%.*]] = alloca i32, align 4
+// IR-NEXT:    [[DOTTILE_0_IV_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* [[GLOB1]], 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_I]], align 4
+// IR-NEXT:    store i32 7, i32* [[I]], align 4
+// IR-NEXT:    [[TMP8:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I]], 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_I]], 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_I]], 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:    [[TMP12:%.*]] = load i32, i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    [[MUL11:%.*]] = mul nsw i32 [[TMP12]], 3
+// IR-NEXT:    [[ADD12:%.*]] = add nsw i32 7, [[MUL11]]
+// IR-NEXT:    store i32 [[ADD12]], i32* [[I]], align 4
+// IR-NEXT:    [[TMP13:%.*]] = load i32, i32* [[I]], align 4
+// IR-NEXT:    call void (...) @body(i32 [[TMP13]])
+// 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* [[GLOB1]], 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)
+    body(i);
+}
+
+
+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);
+}
+
+// IR-LABEL: define {{.*}}void @tfoo7(
+// IR-NEXT:  entry:
+// IR-NEXT:    call void @_Z4foo7IiLi3ELi5EEvT_S0_(i32 0, i32 42)
+// IR-NEXT:    ret void
+//
+// 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:    [[CAPTURE_EXPR:%.+]] = alloca i32, align 4
+// IR-NEXT:    [[CAPTURE_EXPR1:%.+]] = alloca i32, align 4
+// IR-NEXT:    [[CAPTURE_EXPR2:%.+]] = alloca i32, align 4
+// IR-NEXT:    [[I:%.*]] = 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:    store i32 [[START:%.*]], i32* [[START_ADDR]], align 4
+// IR-NEXT:    store i32 [[END:%.*]], i32* [[END_ADDR]], align 4
+// IR-NEXT:    [[TMP0:%.+]] = load i32, i32* [[START_ADDR]], align 4
+// IR-NEXT:    store i32 [[TMP0]], i32* [[CAPTURE_EXPR]], align 4
+// IR-NEXT:    [[TMP1:%.+]] = load i32, i32* [[END_ADDR]], align 4
+// IR-NEXT:    store i32 [[TMP1]], i32* [[CAPTURE_EXPR1]], align 4
+// IR-NEXT:    [[TMP2:%.+]] = load i32, i32* [[CAPTURE_EXPR1]], align 4
+// IR-NEXT:    [[TMP3:%.+]] = load i32, i32* [[CAPTURE_EXPR]], align 4
+// IR-NEXT:    [[SUB:%.+]] = sub i32 [[TMP2]], [[TMP3]]
+// IR-NEXT:    [[SUB3:%.+]] = sub i32 [[SUB]], 1
+// IR-NEXT:    [[ADD:%.+]] = add i32 [[SUB3]], 3
+// IR-NEXT:    [[DIV:%.+]] = udiv i32 [[ADD]], 3
+// IR-NEXT:    [[SUB4:%.+]] = sub i32 [[DIV]], 1
+// IR-NEXT:    store i32 [[SUB4]], i32* [[CAPTURE_EXPR2]], align 4
+// IR-NEXT:    [[TMP4:%.+]] = load i32, i32* [[START_ADDR]], align 4
+// IR-NEXT:    store i32 [[TMP4]], i32* [[I]], 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:    [[TMP6:%.+]] = load i32, i32* [[CAPTURE_EXPR2]], align 4
+// IR-NEXT:    [[ADD3:%.*]] = add i32 [[TMP6]], 1
+// IR-NEXT:    [[CMP:%.*]] = icmp ult 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* [[CAPTURE_EXPR2]], align 4
+// IR-NEXT:    [[ADD10:%.*]] = add i32 [[TMP5]], 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 ult i32 [[ADD10]], [[ADD11]]
+// IR-NEXT:    br i1 [[CMP12]], label %[[COND_TRUE:.*]], label %[[COND_FALSE:.*]]
+// IR:         [[COND_TRUE]]:
+// IR-NEXT:    [[TMP8:%.*]] = load i32, i32* [[CAPTURE_EXPR2]], align 4
+// IR-NEXT:    [[ADD18:%.*]] = add i32 [[TMP8]], 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 ult i32 [[TMP4]], [[COND]]
+// IR-NEXT:    br i1 [[CMP20]], label %[[FOR_BODY21:.*]], label %[[FOR_END:.*]]
+// IR:         [[FOR_BODY21]]:
+// IR-NEXT:    [[TMP11:%.*]] = load i32, i32* [[CAPTURE_EXPR]], align 4
+// IR-NEXT:    [[TMP13:%.*]] = load i32, i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    [[MUL:%.*]] = mul i32 [[TMP13]], 3
+// IR-NEXT:    [[ADD22:%.*]] = add i32 [[TMP11]], [[MUL]]
+// IR-NEXT:    store i32 [[ADD22]], i32* [[I]], align 4
+// IR-NEXT:    [[TMP14:%.*]] = load i32, i32* [[I]], align 4
+// IR-NEXT:    call void (...) @body(i32 [[TMP14]])
+// IR-NEXT:    br label %[[FOR_INC:.*]]
+// IR:         [[FOR_INC]]:
+// IR-NEXT:    [[TMP15:%.*]] = load i32, i32* [[DOTTILE_0_IV_I]], align 4
+// IR-NEXT:    [[INC:%.*]] = add nsw i32 [[TMP15]], 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:    [[TMP16:%.*]] = load i32, i32* [[DOTFLOOR_0_IV_I]], align 4
+// IR-NEXT:    [[ADD24:%.*]] = add nsw i32 [[TMP16]], 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 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@+5 {{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 VisitExpressionTraitExpr(const ExpressionTraitExpr *E);
   void VisitUnresolvedMemberExpr(const UnresolvedMemberExpr *U);
   void VisitVAArgExpr(const VAArgExpr *E);
   void VisitSizeOfPackExpr(const SizeOfPackExpr *E);
   void VisitPseudoObjectExpr(const PseudoObjectExpr *E);
   void VisitOpaqueValueExpr(const OpaqueValueExpr *E);
   void VisitLambdaExpr(const LambdaExpr *E);
   void VisitOMPExecutableDirective(const OMPExecutableDirective *D);
+  void VisitOMPLoopBasedDirective(const OMPLoopBasedDirective *D);
   void VisitOMPLoopDirective(const OMPLoopDirective *D);
   void VisitOMPParallelDirective(const OMPParallelDirective *D);
   void VisitOMPSimdDirective(const OMPSimdDirective *D);
+  void VisitOMPTileDirective(const OMPTileDirective *D);
   void VisitOMPForDirective(const OMPForDirective *D);
   void VisitOMPForSimdDirective(const OMPForSimdDirective *D);
   void VisitOMPSectionsDirective(const OMPSectionsDirective *D);
   void VisitOMPSectionDirective(const OMPSectionDirective *D);
   void VisitOMPSingleDirective(const OMPSingleDirective *D);
   void VisitOMPMasterDirective(const OMPMasterDirective *D);
   void VisitOMPCriticalDirective(const OMPCriticalDirective *D);
   void VisitOMPParallelForDirective(const OMPParallelForDirective *D);

 void OMPClauseEnqueue::VisitOMPSafelenClause(const OMPSafelenClause *C) {
   Visitor->AddStmt(C->getSafelen());
 }
 
 void OMPClauseEnqueue::VisitOMPSimdlenClause(const OMPSimdlenClause *C) {
   Visitor->AddStmt(C->getSimdlen());
 }
 
+void OMPClauseEnqueue::VisitOMPSizesClause(const OMPSizesClause *C) {
+  for (auto E : C->getSizesRefs())
+    Visitor->AddStmt(E);
+}
+
 void OMPClauseEnqueue::VisitOMPAllocatorClause(const OMPAllocatorClause *C) {
   Visitor->AddStmt(C->getAllocator());
 }
 
 void OMPClauseEnqueue::VisitOMPCollapseClause(const OMPCollapseClause *C) {
   Visitor->AddStmt(C->getNumForLoops());
 }
 
void EnqueueVisitor::VisitOMPExecutableDirective(
     const OMPExecutableDirective *D) {
   EnqueueChildren(D);
   for (ArrayRef<OMPClause *>::iterator I = D->clauses().begin(),
                                        E = D->clauses().end();
        I != E; ++I)
     EnqueueChildren(*I);
 }
 
-void EnqueueVisitor::VisitOMPLoopDirective(const OMPLoopDirective *D) {
+void EnqueueVisitor::VisitOMPLoopBasedDirective(
+    const OMPLoopBasedDirective *D) {
   VisitOMPExecutableDirective(D);
 }
 
+void EnqueueVisitor::VisitOMPLoopDirective(const OMPLoopDirective *D) {
+  VisitOMPLoopBasedDirective(D);
+}
+
 void EnqueueVisitor::VisitOMPParallelDirective(const OMPParallelDirective *D) {
   VisitOMPExecutableDirective(D);
 }
 
 void EnqueueVisitor::VisitOMPSimdDirective(const OMPSimdDirective *D) {
   VisitOMPLoopDirective(D);
 }
 
+void EnqueueVisitor::VisitOMPTileDirective(const OMPTileDirective *D) {
+  VisitOMPLoopBasedDirective(D);
+}
+
 void EnqueueVisitor::VisitOMPForDirective(const OMPForDirective *D) {
   VisitOMPLoopDirective(D);
 }
 
 void EnqueueVisitor::VisitOMPForSimdDirective(const OMPForSimdDirective *D) {
   VisitOMPLoopDirective(D);
 }
 
CXString clang_getCursorKindSpelling(enum CXCursorKind Kind) {
   case CXCursor_CXXAccessSpecifier:
     return cxstring::createRef("CXXAccessSpecifier");
   case CXCursor_ModuleImportDecl:
     return cxstring::createRef("ModuleImport");
   case CXCursor_OMPParallelDirective:
     return cxstring::createRef("OMPParallelDirective");
   case CXCursor_OMPSimdDirective:
     return cxstring::createRef("OMPSimdDirective");
+  case CXCursor_OMPTileDirective:
+    return cxstring::createRef("OMPTileDirective");
   case CXCursor_OMPForDirective:
     return cxstring::createRef("OMPForDirective");
   case CXCursor_OMPForSimdDirective:
     return cxstring::createRef("OMPForSimdDirective");
   case CXCursor_OMPSectionsDirective:
     return cxstring::createRef("OMPSectionsDirective");
   case CXCursor_OMPSectionDirective:
     return cxstring::createRef("OMPSectionDirective");

clang/tools/libclang/CXCursor.cpp

case Stmt::MSDependentExistsStmtClass:
     K = CXCursor_UnexposedStmt;
     break;
   case Stmt::OMPParallelDirectiveClass:
     K = CXCursor_OMPParallelDirective;
     break;
   case Stmt::OMPSimdDirectiveClass:
     K = CXCursor_OMPSimdDirective;
     break;
+  case Stmt::OMPTileDirectiveClass:
+    K = CXCursor_OMPTileDirective;
+    break;
   case Stmt::OMPForDirectiveClass:
     K = CXCursor_OMPForDirective;
     break;
   case Stmt::OMPForSimdDirectiveClass:
     K = CXCursor_OMPForSimdDirective;
     break;
   case Stmt::OMPSectionsDirectiveClass:
     K = CXCursor_OMPSectionsDirective;

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

 def OMPC_Default : Clause<"default"> {
   let clangClass = "OMPDefaultClause";
   let flangClassValue = "OmpDefaultClause";
 }
 def OMPC_Private : Clause<"private"> {
   let clangClass = "OMPPrivateClause";
   let flangClassValue = "OmpObjectList";
 }
+def OMPC_Sizes: Clause<"sizes"> { let clangClass = "OMPSizesClause"; }
 def OMPC_FirstPrivate : Clause<"firstprivate"> {
   let clangClass = "OMPFirstprivateClause";
   let flangClassValue = "OmpObjectList";
 }
 def OMPC_LastPrivate : Clause<"lastprivate"> {
   let clangClass = "OMPLastprivateClause";
   let flangClassValue = "OmpObjectList";
 }
def OMP_Simd : Directive<"simd"> {
   ];
   let allowedOnceClauses = [
     VersionedClause<OMPC_Collapse>,
     VersionedClause<OMPC_SafeLen>,
     VersionedClause<OMPC_SimdLen>,
     VersionedClause<OMPC_If, 50>,
   ];
 }
+def OMP_Tile : Directive<"tile"> {
+  let allowedOnceClauses = [
+    VersionedClause<OMPC_Sizes, 51>,
+  ];
+}
 def OMP_For : Directive<"for"> {
   let allowedClauses = [
     VersionedClause<OMPC_Private>,
     VersionedClause<OMPC_LastPrivate>,
     VersionedClause<OMPC_FirstPrivate>,
     VersionedClause<OMPC_Reduction>,
     VersionedClause<OMPC_Collapse>,
     VersionedClause<OMPC_Schedule>,