[OpenMP] Specialize default schedule on a worksharing loop on the NVPTX device.

[arpith-jacob]

Summary

The default schedule type on a worksharing loop is implementation
defined according to the OpenMP specifications. Currently, the
compiler codegens a doubly nested loop that effectively implements
a schedule of type (static). This is ideal for threads on CPUs.

On the NVPTX and other SIMT GPUs, this schedule provides very poor
performance because consecutive threads in a warp access loop arrays
in a non-coalesced manner. That is, to achieve coalescing, and good
performance, the best schedule is static with a chunk size of 1.

This patch adds support for target devices to select the best default
schedule depending on their architecture. It modifies loop codegen
to generate optimized code for (static,1) on the NVPTX device, i.e.,
by using a single loop instead of a doubly nested loop as is
currently the case.

[ABataev]

General comment: do you really need to define default schedule kind in Sema? Could you do it during codegen phase?

include/clang/Basic/OpenMPKinds.h

22–24

No way, these classes should not be used here

247–253

You can't use Sema and OMPClause in Basic, it is not allowed

lib/CodeGen/CGStmtOpenMP.cpp

2215–2218

What's the difference between this code and next else branch?

[arpith-jacob]

Hi Alexey,

Thank you for your review. The main difference in the specialized codegen (if vs. else part in CGStmtOpenMP.cpp).

If-part: emitForStaticInit uses the Chunk parameter (else has it set to null).
If-part: does not use EmitIgnoredExpr()

I can combine if- and else- part with appropriate guards over the above two statements if you like.

General comment: do you really need to define default schedule kind in Sema? Could you do it during codegen phase?

Yes, we need to define default schedule in Sema because it has to set getCond() and getInc() in CheckOpenMPLoop(), which is in Sema. For example, the Inc expression is calculated as follows based on the default schedule:

// Loop increment (IV = IV + 1) or (IV = IV + ST) if (static,1) scheduling.
ExprResult Inc =
  DefaultScheduleKind == OMPDSK_static_chunkone
      ? SemaRef.BuildBinOp(CurScope, IncLoc, BO_Add, IV.get(), ST.get())
      : SemaRef.BuildBinOp(CurScope, IncLoc, BO_Add, IV.get(),
                           SemaRef.ActOnIntegerConstant(IncLoc, 1).get());

You can't use Sema and OMPClause in Basic, it is not allowed

Ok, I can move getDefaultSchedule() from OpenMPKinds.cpp and make it a static function in SemaOpenMP.cpp. Then, CheckOpenMPLoop() can directly call this static function. Does this sound okay?

Thanks.

[ABataev]

In D29910#676186, @arpith-jacob wrote:

Hi Alexey,

Thank you for your review. The main difference in the specialized codegen (if vs. else part in CGStmtOpenMP.cpp).

If-part: emitForStaticInit uses the Chunk parameter (else has it set to null).
If-part: does not use EmitIgnoredExpr()

I can combine if- and else- part with appropriate guards over the above two statements if you like.

Ok, please try to do it.

General comment: do you really need to define default schedule kind in Sema? Could you do it during codegen phase?

Yes, we need to define default schedule in Sema because it has to set getCond() and getInc() in CheckOpenMPLoop(), which is in Sema. For example, the Inc expression is calculated as follows based on the default schedule:

// Loop increment (IV = IV + 1) or (IV = IV + ST) if (static,1) scheduling.
ExprResult Inc =
  DefaultScheduleKind == OMPDSK_static_chunkone
      ? SemaRef.BuildBinOp(CurScope, IncLoc, BO_Add, IV.get(), ST.get())
      : SemaRef.BuildBinOp(CurScope, IncLoc, BO_Add, IV.get(),
                           SemaRef.ActOnIntegerConstant(IncLoc, 1).get());

The maybe it is better to define a variable with loop increment and define this variable during codegen with 1 or stride (for GPU)? I don't like the idea of adding some kind of default scheduling, that is not defined in standard in Sema, preferer to define it during codegen and somewhere in NVPTX specific runtime support class.

You can't use Sema and OMPClause in Basic, it is not allowed

Ok, I can move getDefaultSchedule() from OpenMPKinds.cpp and make it a static function in SemaOpenMP.cpp. Then, CheckOpenMPLoop() can directly call this static function. Does this sound okay?

Thanks.

[arpith-jacob]

Hi Alexey,

Thank you for reviewing this patch.

I don't like the idea of adding some kind of default scheduling, that is not defined in standard in Sema

Actually, "default scheduling" is defined in the OpenMP spec. It is called "def-sched-var" and controls the scheduling of loops. It's value is implementation (compiler) defined. So why not allow the target device to choose this value in the compiler?

http://www.openmp.org/wp-content/uploads/openmp-4.5.pdf

Section 2.3.1: ICV Descriptions, pg 46:
def-sched-var - controls the implementation defined default scheduling of loop regions. There is one copy of this ICV per device.

Section 2.3.2: ICV Initialization, pg 47:
Table 2.1:
def-sched-var   No environment variable      Initial value is implementation defined

Section 2.7.1.1: Determining the Schedule of a Worksharing Loop
When execution encounters a loop directive, the schedule clause (if any) on the directive, and the run-sched-var and def-sched-var ICVs are used to determine how loop iterations are assigned to threads. See Section 2.3 on page 36 for details of how the values of the ICVs are determined. If the loop directive does not have a schedule clause then the current value of the def-sched-var ICV determines the schedule.

I've reworked the patch to handle the default scheduling in Sema and removed the function from OpenMPKind.cpp. Please let me know if this looks good.

I can rewrite the patch as you suggested, involving NVPTX specific RT, but I think the code will look quite ugly.

[ABataev]

In D29910#676777, @arpith-jacob wrote:

Hi Alexey,

Thank you for reviewing this patch.

I don't like the idea of adding some kind of default scheduling, that is not defined in standard in Sema

Actually, "default scheduling" is defined in the OpenMP spec. It is called "def-sched-var" and controls the scheduling of loops. It's value is implementation (compiler) defined. So why not allow the target device to choose this value in the compiler?

http://www.openmp.org/wp-content/uploads/openmp-4.5.pdf

Section 2.3.1: ICV Descriptions, pg 46:
def-sched-var - controls the implementation defined default scheduling of loop regions. There is one copy of this ICV per device.

Section 2.3.2: ICV Initialization, pg 47:
Table 2.1:
def-sched-var   No environment variable      Initial value is implementation defined

Section 2.7.1.1: Determining the Schedule of a Worksharing Loop
When execution encounters a loop directive, the schedule clause (if any) on the directive, and the run-sched-var and def-sched-var ICVs are used to determine how loop iterations are assigned to threads. See Section 2.3 on page 36 for details of how the values of the ICVs are determined. If the loop directive does not have a schedule clause then the current value of the def-sched-var ICV determines the schedule.

I've reworked the patch to handle the default scheduling in Sema and removed the function from OpenMPKind.cpp. Please let me know if this looks good.

I can rewrite the patch as you suggested, involving NVPTX specific RT, but I think the code will look quite ugly.

Arpith, yes, there is such variable. But also standard says, that it is device specific. My opinion, all device-specific things must be defined by runtime library or at least runtime-support class, not by Sema. Sema must be as much platform independent as possible

include/clang/AST/StmtOpenMP.h

 
 /// \brief This is a common base class for loop directives ('omp simd', 'omp
 /// for', 'omp for simd' etc.). It is responsible for the loop code generation.
 ///
 class OMPLoopDirective : public OMPExecutableDirective {
   friend class ASTStmtReader;
   /// \brief Number of collapsed loops as specified by 'collapse' clause.
   unsigned CollapsedNum;
+  /// \brief DefaultScheduleKind - Schedule type to use for a given target
+  /// if no 'schedule' clause or a 'schedule' type 'auto' is specified.
+  OpenMPDefaultScheduleKind DefaultScheduleKind;
 
   /// \brief Offsets to the stored exprs.
   /// This enumeration contains offsets to all the pointers to children
   /// expressions stored in OMPLoopDirective.
   /// The first 9 children are nesessary for all the loop directives, and
   /// the next 10 are specific to the worksharing ones.
   /// After the fixed children, three arrays of length CollapsedNum are
   /// allocated: loop counters, their updates and final values.
protected:
   template <typename T>
   OMPLoopDirective(const T *That, StmtClass SC, OpenMPDirectiveKind Kind,
                    SourceLocation StartLoc, SourceLocation EndLoc,
                    unsigned CollapsedNum, unsigned NumClauses,
                    unsigned NumSpecialChildren = 0)
       : OMPExecutableDirective(That, SC, Kind, StartLoc, EndLoc, NumClauses,
                                numLoopChildren(CollapsedNum, Kind) +
                                    NumSpecialChildren),
-        CollapsedNum(CollapsedNum) {}
+        CollapsedNum(CollapsedNum), DefaultScheduleKind(OMPDSK_unknown) {}
 
   /// \brief Offset to the start of children expression arrays.
   static unsigned getArraysOffset(OpenMPDirectiveKind Kind) {
     return (isOpenMPWorksharingDirective(Kind) ||
             isOpenMPTaskLoopDirective(Kind) ||
             isOpenMPDistributeDirective(Kind))
                ? WorksharingEnd
                : DefaultEnd;
protected:
   }
   void setPrevUpperBoundVariable(Expr *PrevUB) {
     assert((isOpenMPWorksharingDirective(getDirectiveKind()) ||
             isOpenMPTaskLoopDirective(getDirectiveKind()) ||
             isOpenMPDistributeDirective(getDirectiveKind())) &&
            "expected worksharing loop directive");
     *std::next(child_begin(), PrevUpperBoundVariableOffset) = PrevUB;
   }
+  void setDefaultSchedule(OpenMPDefaultScheduleKind SK) {
+    DefaultScheduleKind = SK;
+  }
   void setCounters(ArrayRef<Expr *> A);
   void setPrivateCounters(ArrayRef<Expr *> A);
   void setInits(ArrayRef<Expr *> A);
   void setUpdates(ArrayRef<Expr *> A);
   void setFinals(ArrayRef<Expr *> A);
 
 public:
   /// \brief The expressions built for the OpenMP loop CodeGen for the
struct HelperExprs {
     /// \brief Update of UpperBound for statically sheduled 'omp for' loops.
     Expr *NUB;
     /// \brief PreviousLowerBound - local variable passed to runtime in the
     /// enclosing schedule or null if that does not apply.
     Expr *PrevLB;
     /// \brief PreviousUpperBound - local variable passed to runtime in the
     /// enclosing schedule or null if that does not apply.
     Expr *PrevUB;
+    /// \brief DefaultScheduleKind - Schedule type to use for the given target
+    /// if no 'schedule' clause or a 'schedule' type 'auto' is specified.
+    OpenMPDefaultScheduleKind DefaultScheduleKind;
     /// \brief Counters Loop counters.
     SmallVector<Expr *, 4> Counters;
     /// \brief PrivateCounters Loop counters.
     SmallVector<Expr *, 4> PrivateCounters;
     /// \brief Expressions for loop counters inits for CodeGen.
     SmallVector<Expr *, 4> Inits;
     /// \brief Expressions for loop counters update for CodeGen.
     SmallVector<Expr *, 4> Updates;
void clear(unsigned Size) {
       UB = nullptr;
       ST = nullptr;
       EUB = nullptr;
       NLB = nullptr;
       NUB = nullptr;
       NumIterations = nullptr;
       PrevLB = nullptr;
       PrevUB = nullptr;
+      DefaultScheduleKind = OMPDSK_unknown;
       Counters.resize(Size);
       PrivateCounters.resize(Size);
       Inits.resize(Size);
       Updates.resize(Size);
       Finals.resize(Size);
       for (unsigned i = 0; i < Size; ++i) {
         Counters[i] = nullptr;
         PrivateCounters[i] = nullptr;
public:
   Expr *getPrevUpperBoundVariable() const {
     assert((isOpenMPWorksharingDirective(getDirectiveKind()) ||
             isOpenMPTaskLoopDirective(getDirectiveKind()) ||
             isOpenMPDistributeDirective(getDirectiveKind())) &&
            "expected worksharing loop directive");
     return const_cast<Expr *>(reinterpret_cast<const Expr *>(
         *std::next(child_begin(), PrevUpperBoundVariableOffset)));
   }
+  OpenMPDefaultScheduleKind getDefaultSchedule() const {
+    return DefaultScheduleKind;
+  }
   const Stmt *getBody() const {
     // This relies on the loop form is already checked by Sema.
     Stmt *Body = getAssociatedStmt()->IgnoreContainers(true);
     Body = cast<ForStmt>(Body)->getBody();
     for (unsigned Cnt = 1; Cnt < CollapsedNum; ++Cnt) {
       Body = Body->IgnoreContainers();
       Body = cast<ForStmt>(Body)->getBody();
     }

include/clang/Basic/OpenMPKinds.h

 
 #ifndef LLVM_CLANG_BASIC_OPENMPKINDS_H
 #define LLVM_CLANG_BASIC_OPENMPKINDS_H
 
 #include "llvm/ADT/StringRef.h"
 
 namespace clang {
 
 /// \brief OpenMP directives.
 enum OpenMPDirectiveKind {
 #define OPENMP_DIRECTIVE(Name) \
   OMPD_##Name,
 #define OPENMP_DIRECTIVE_EXT(Name, Str) \
   OMPD_##Name,
 #include "clang/Basic/OpenMPKinds.def"
   OMPD_unknown
 };
 
 /// \brief OpenMP clauses.

 
 /// Scheduling data for loop-based OpenMP directives.
 struct OpenMPScheduleTy final {
   OpenMPScheduleClauseKind Schedule = OMPC_SCHEDULE_unknown;
   OpenMPScheduleClauseModifier M1 = OMPC_SCHEDULE_MODIFIER_unknown;
   OpenMPScheduleClauseModifier M2 = OMPC_SCHEDULE_MODIFIER_unknown;
 };
 
+/// Default schedule type for any loop-based (#for) OpenMP directive.
+enum OpenMPDefaultScheduleKind { OMPDSK_static_chunkone, OMPDSK_unknown };
+
 OpenMPDirectiveKind getOpenMPDirectiveKind(llvm::StringRef Str);
 const char *getOpenMPDirectiveName(OpenMPDirectiveKind Kind);
 
 OpenMPClauseKind getOpenMPClauseKind(llvm::StringRef Str);
 const char *getOpenMPClauseName(OpenMPClauseKind Kind);
 
 unsigned getOpenMPSimpleClauseType(OpenMPClauseKind Kind, llvm::StringRef Str);
 const char *getOpenMPSimpleClauseTypeName(OpenMPClauseKind Kind, unsigned Type);

 bool isOpenMPLoopBoundSharingDirective(OpenMPDirectiveKind Kind);
 
 /// Return the captured regions of an OpenMP directive.
 void getOpenMPCaptureRegions(
     llvm::SmallVectorImpl<OpenMPDirectiveKind> &CaptureRegions,
     OpenMPDirectiveKind DKind);
 }
 
 #endif
 

lib/AST/StmtOpenMP.cpp

OMPSimdDirective::Create(const ASTContext &C, SourceLocation StartLoc,
   Dir->setInit(Exprs.Init);
   Dir->setInc(Exprs.Inc);
   Dir->setCounters(Exprs.Counters);
   Dir->setPrivateCounters(Exprs.PrivateCounters);
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
   Dir->setPreInits(Exprs.PreInits);
+  // TODO: Set default schedule.
   return Dir;
 }
 
 OMPSimdDirective *OMPSimdDirective::CreateEmpty(const ASTContext &C,
                                                 unsigned NumClauses,
                                                 unsigned CollapsedNum,
                                                 EmptyShell) {
   unsigned Size = llvm::alignTo(sizeof(OMPSimdDirective), alignof(OMPClause *));
OMPForDirective::Create(const ASTContext &C, SourceLocation StartLoc,
   Dir->setPrevUpperBoundVariable(Exprs.PrevUB);
   Dir->setCounters(Exprs.Counters);
   Dir->setPrivateCounters(Exprs.PrivateCounters);
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
   Dir->setPreInits(Exprs.PreInits);
   Dir->setHasCancel(HasCancel);
+  Dir->setDefaultSchedule(Exprs.DefaultScheduleKind);
   return Dir;
 }
 
 OMPForDirective *OMPForDirective::CreateEmpty(const ASTContext &C,
                                               unsigned NumClauses,
                                               unsigned CollapsedNum,
                                               EmptyShell) {
   unsigned Size = llvm::alignTo(sizeof(OMPForDirective), alignof(OMPClause *));
OMPForSimdDirective::Create(const ASTContext &C, SourceLocation StartLoc,
   Dir->setPrevLowerBoundVariable(Exprs.PrevLB);
   Dir->setPrevUpperBoundVariable(Exprs.PrevUB);
   Dir->setCounters(Exprs.Counters);
   Dir->setPrivateCounters(Exprs.PrivateCounters);
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
   Dir->setPreInits(Exprs.PreInits);
+  // TODO: Set default schedule.
   return Dir;
 }
 
 OMPForSimdDirective *OMPForSimdDirective::CreateEmpty(const ASTContext &C,
                                                       unsigned NumClauses,
                                                       unsigned CollapsedNum,
                                                       EmptyShell) {
   unsigned Size =
OMPParallelForDirective *OMPParallelForDirective::Create(
   Dir->setPrevUpperBoundVariable(Exprs.PrevUB);
   Dir->setCounters(Exprs.Counters);
   Dir->setPrivateCounters(Exprs.PrivateCounters);
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
   Dir->setPreInits(Exprs.PreInits);
   Dir->setHasCancel(HasCancel);
+  // TODO: Set default schedule.
   return Dir;
 }
 
 OMPParallelForDirective *
 OMPParallelForDirective::CreateEmpty(const ASTContext &C, unsigned NumClauses,
                                      unsigned CollapsedNum, EmptyShell) {
   unsigned Size =
       llvm::alignTo(sizeof(OMPParallelForDirective), alignof(OMPClause *));
OMPParallelForSimdDirective *OMPParallelForSimdDirective::Create(
   Dir->setPrevLowerBoundVariable(Exprs.PrevLB);
   Dir->setPrevUpperBoundVariable(Exprs.PrevUB);
   Dir->setCounters(Exprs.Counters);
   Dir->setPrivateCounters(Exprs.PrivateCounters);
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
   Dir->setPreInits(Exprs.PreInits);
+  // TODO: Set default schedule.
   return Dir;
 }
 
 OMPParallelForSimdDirective *
 OMPParallelForSimdDirective::CreateEmpty(const ASTContext &C,
                                          unsigned NumClauses,
                                          unsigned CollapsedNum, EmptyShell) {
   unsigned Size =
OMPTargetParallelForDirective *OMPTargetParallelForDirective::Create(
   Dir->setPrevUpperBoundVariable(Exprs.PrevUB);
   Dir->setCounters(Exprs.Counters);
   Dir->setPrivateCounters(Exprs.PrivateCounters);
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
   Dir->setPreInits(Exprs.PreInits);
   Dir->setHasCancel(HasCancel);
+  // TODO: Set default schedule.
   return Dir;
 }
 
 OMPTargetParallelForDirective *
 OMPTargetParallelForDirective::CreateEmpty(const ASTContext &C,
                                            unsigned NumClauses,
                                            unsigned CollapsedNum, EmptyShell) {
   unsigned Size = llvm::alignTo(sizeof(OMPTargetParallelForDirective),
OMPDistributeDirective *OMPDistributeDirective::Create(
   Dir->setPrevLowerBoundVariable(Exprs.PrevLB);
   Dir->setPrevUpperBoundVariable(Exprs.PrevUB);
   Dir->setCounters(Exprs.Counters);
   Dir->setPrivateCounters(Exprs.PrivateCounters);
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
   Dir->setPreInits(Exprs.PreInits);
+  // TODO: Set default schedule.
   return Dir;
 }
 
 OMPDistributeDirective *
 OMPDistributeDirective::CreateEmpty(const ASTContext &C, unsigned NumClauses,
                                     unsigned CollapsedNum, EmptyShell) {
   unsigned Size =
       llvm::alignTo(sizeof(OMPDistributeDirective), alignof(OMPClause *));
OMPDistributeParallelForDirective *OMPDistributeParallelForDirective::Create(
   Dir->setPrevLowerBoundVariable(Exprs.PrevLB);
   Dir->setPrevUpperBoundVariable(Exprs.PrevUB);
   Dir->setCounters(Exprs.Counters);
   Dir->setPrivateCounters(Exprs.PrivateCounters);
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
   Dir->setPreInits(Exprs.PreInits);
+  // TODO: Set default schedule.
   return Dir;
 }
 
 OMPDistributeParallelForDirective *
 OMPDistributeParallelForDirective::CreateEmpty(const ASTContext &C,
                                                unsigned NumClauses,
                                                unsigned CollapsedNum,
                                                EmptyShell) {
OMPDistributeParallelForSimdDirective::Create(
   Dir->setPrevLowerBoundVariable(Exprs.PrevLB);
   Dir->setPrevUpperBoundVariable(Exprs.PrevUB);
   Dir->setCounters(Exprs.Counters);
   Dir->setPrivateCounters(Exprs.PrivateCounters);
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
   Dir->setPreInits(Exprs.PreInits);
+  // TODO: Set default schedule.
   return Dir;
 }
 
 OMPDistributeParallelForSimdDirective *
 OMPDistributeParallelForSimdDirective::CreateEmpty(const ASTContext &C,
                                                    unsigned NumClauses,
                                                    unsigned CollapsedNum,
                                                    EmptyShell) {
OMPDistributeSimdDirective *OMPDistributeSimdDirective::Create(
   Dir->setPrevLowerBoundVariable(Exprs.PrevLB);
   Dir->setPrevUpperBoundVariable(Exprs.PrevUB);
   Dir->setCounters(Exprs.Counters);
   Dir->setPrivateCounters(Exprs.PrivateCounters);
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
   Dir->setPreInits(Exprs.PreInits);
+  // TODO: Set default schedule.
   return Dir;
 }
 
 OMPDistributeSimdDirective *
 OMPDistributeSimdDirective::CreateEmpty(const ASTContext &C,
                                         unsigned NumClauses,
                                         unsigned CollapsedNum, EmptyShell) {
   unsigned Size =
OMPTargetParallelForSimdDirective *OMPTargetParallelForSimdDirective::Create(
   Dir->setPrevLowerBoundVariable(Exprs.PrevLB);
   Dir->setPrevUpperBoundVariable(Exprs.PrevUB);
   Dir->setCounters(Exprs.Counters);
   Dir->setPrivateCounters(Exprs.PrivateCounters);
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
   Dir->setPreInits(Exprs.PreInits);
+  // TODO: Set default schedule.
   return Dir;
 }
 
 OMPTargetParallelForSimdDirective *
 OMPTargetParallelForSimdDirective::CreateEmpty(const ASTContext &C,
                                                unsigned NumClauses,
                                                unsigned CollapsedNum,
                                                EmptyShell) {
OMPTeamsDistributeDirective *OMPTeamsDistributeDirective::Create(
   Dir->setPrevLowerBoundVariable(Exprs.PrevLB);
   Dir->setPrevUpperBoundVariable(Exprs.PrevUB);
   Dir->setCounters(Exprs.Counters);
   Dir->setPrivateCounters(Exprs.PrivateCounters);
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
   Dir->setPreInits(Exprs.PreInits);
+  // TODO: Set default schedule.
   return Dir;
 }
 
 OMPTeamsDistributeDirective *
 OMPTeamsDistributeDirective::CreateEmpty(const ASTContext &C,
                                          unsigned NumClauses,
                                          unsigned CollapsedNum, EmptyShell) {
   unsigned Size =
OMPTeamsDistributeSimdDirective *OMPTeamsDistributeSimdDirective::Create(
   Dir->setPrevLowerBoundVariable(Exprs.PrevLB);
   Dir->setPrevUpperBoundVariable(Exprs.PrevUB);
   Dir->setCounters(Exprs.Counters);
   Dir->setPrivateCounters(Exprs.PrivateCounters);
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
   Dir->setPreInits(Exprs.PreInits);
+  // TODO: Set default schedule.
   return Dir;
 }
 
 OMPTeamsDistributeSimdDirective *OMPTeamsDistributeSimdDirective::CreateEmpty(
     const ASTContext &C, unsigned NumClauses, unsigned CollapsedNum,
     EmptyShell) {
   unsigned Size = llvm::alignTo(sizeof(OMPTeamsDistributeSimdDirective),
                                 alignof(OMPClause *));
OMPTeamsDistributeParallelForSimdDirective::Create(
   Dir->setPrevLowerBoundVariable(Exprs.PrevLB);
   Dir->setPrevUpperBoundVariable(Exprs.PrevUB);
   Dir->setCounters(Exprs.Counters);
   Dir->setPrivateCounters(Exprs.PrivateCounters);
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
   Dir->setPreInits(Exprs.PreInits);
+  // TODO: Set default schedule.
   return Dir;
 }
 
 OMPTeamsDistributeParallelForSimdDirective *
 OMPTeamsDistributeParallelForSimdDirective::CreateEmpty(const ASTContext &C,
                                                         unsigned NumClauses,
                                                         unsigned CollapsedNum,
                                                         EmptyShell) {
OMPTeamsDistributeParallelForDirective::Create(
   Dir->setPrevLowerBoundVariable(Exprs.PrevLB);
   Dir->setPrevUpperBoundVariable(Exprs.PrevUB);
   Dir->setCounters(Exprs.Counters);
   Dir->setPrivateCounters(Exprs.PrivateCounters);
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
   Dir->setPreInits(Exprs.PreInits);
+  // TODO: Set default schedule.
   return Dir;
 }
 
 OMPTeamsDistributeParallelForDirective *
 OMPTeamsDistributeParallelForDirective::CreateEmpty(const ASTContext &C,
                                                     unsigned NumClauses,
                                                     unsigned CollapsedNum,
                                                     EmptyShell) {
OMPTargetTeamsDistributeDirective *OMPTargetTeamsDistributeDirective::Create(
   Dir->setPrevLowerBoundVariable(Exprs.PrevLB);
   Dir->setPrevUpperBoundVariable(Exprs.PrevUB);
   Dir->setCounters(Exprs.Counters);
   Dir->setPrivateCounters(Exprs.PrivateCounters);
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
   Dir->setPreInits(Exprs.PreInits);
+  // TODO: Set default schedule.
   return Dir;
 }
 
 OMPTargetTeamsDistributeDirective *
 OMPTargetTeamsDistributeDirective::CreateEmpty(const ASTContext &C,
                                                unsigned NumClauses,
                                                unsigned CollapsedNum,
                                                EmptyShell) {
OMPTargetTeamsDistributeParallelForDirective::Create(
   Dir->setPrevLowerBoundVariable(Exprs.PrevLB);
   Dir->setPrevUpperBoundVariable(Exprs.PrevUB);
   Dir->setCounters(Exprs.Counters);
   Dir->setPrivateCounters(Exprs.PrivateCounters);
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
   Dir->setPreInits(Exprs.PreInits);
+  // TODO: Set default schedule.
   return Dir;
 }
 
 OMPTargetTeamsDistributeParallelForDirective *
 OMPTargetTeamsDistributeParallelForDirective::CreateEmpty(const ASTContext &C,
                                                           unsigned NumClauses,
                                                           unsigned CollapsedNum,
                                                           EmptyShell) {
OMPTargetTeamsDistributeParallelForSimdDirective::Create(
   Dir->setPrevLowerBoundVariable(Exprs.PrevLB);
   Dir->setPrevUpperBoundVariable(Exprs.PrevUB);
   Dir->setCounters(Exprs.Counters);
   Dir->setPrivateCounters(Exprs.PrivateCounters);
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
   Dir->setPreInits(Exprs.PreInits);
+  // TODO: Set default schedule.
   return Dir;
 }
 
 OMPTargetTeamsDistributeParallelForSimdDirective *
 OMPTargetTeamsDistributeParallelForSimdDirective::CreateEmpty(
     const ASTContext &C, unsigned NumClauses, unsigned CollapsedNum,
     EmptyShell) {
   auto Size =
OMPTargetTeamsDistributeSimdDirective::Create(
   Dir->setPrevLowerBoundVariable(Exprs.PrevLB);
   Dir->setPrevUpperBoundVariable(Exprs.PrevUB);
   Dir->setCounters(Exprs.Counters);
   Dir->setPrivateCounters(Exprs.PrivateCounters);
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
   Dir->setPreInits(Exprs.PreInits);
+  // TODO: Set default schedule.
   return Dir;
 }
 
 OMPTargetTeamsDistributeSimdDirective *
 OMPTargetTeamsDistributeSimdDirective::CreateEmpty(const ASTContext &C,
                                                    unsigned NumClauses,
                                                    unsigned CollapsedNum,
                                                    EmptyShell) {

lib/CodeGen/CGStmtOpenMP.cpp

// Emit 'then' code.
           Chunk = EmitScalarExpr(Ch);
           Chunk = EmitScalarConversion(Chunk, Ch->getType(),
                                        S.getIterationVariable()->getType(),
                                        S.getLocStart());
         }
       }
       const unsigned IVSize = getContext().getTypeSize(IVExpr->getType());
       const bool IVSigned = IVExpr->getType()->hasSignedIntegerRepresentation();
+      // For NVPTX and other GPU targets high performance is often achieved
+      // if adjacent threads access memory in a coalesced manner.  This is
+      // true for loops that access memory with stride one if a static
+      // schedule with chunk size of 1 is used.  We generate such code
+      // whenever the OpenMP standard gives us freedom to do so.
+      //
+      // This case is called if there is no schedule clause, with a
+      // schedule(auto), or with a schedule(static,1).
+      //
+      // Codegen is optimized for this case.  Since chunk size is 1 we do not
+      // need to generate the inner loop, i.e., the chunk iterator can be
+      // removed.
+      // while(idx < GlobalUB) {
+      //   BODY;
+      //   idx += ST;
+      // }
+      if (S.getDefaultSchedule() == OMPDSK_static_chunkone) {
+        ScheduleKind.Schedule = OMPC_SCHEDULE_static;
+        if (!Chunk) // Force use of chunk=1
+          Chunk = Builder.getIntN(IVSize, 1);
+      }
       // OpenMP 4.5, 2.7.1 Loop Construct, Description.
       // If the static schedule kind is specified or if the ordered clause is
       // specified, and if no monotonic modifier is specified, the effect will
       // be as if the monotonic modifier was specified.
-      if (RT.isStaticNonchunked(ScheduleKind.Schedule,
+      if (S.getDefaultSchedule() == OMPDSK_static_chunkone ||
+          (RT.isStaticNonchunked(ScheduleKind.Schedule,
                                  /* Chunked */ Chunk != nullptr) &&
-          !Ordered) {
+           !Ordered)) {
         if (isOpenMPSimdDirective(S.getDirectiveKind()))
           EmitOMPSimdInit(S, /*IsMonotonic=*/true);
         // OpenMP [2.7.1, Loop Construct, Description, table 2-1]
         // When no chunk_size is specified, the iteration space is divided into
         // chunks that are approximately equal in size, and at most one chunk is
         // distributed to each thread. Note that the size of the chunks is
         // unspecified in this case.
-        RT.emitForStaticInit(*this, S.getLocStart(), ScheduleKind,
-                             IVSize, IVSigned, Ordered,
-                             IL.getAddress(), LB.getAddress(),
-                             UB.getAddress(), ST.getAddress());
+        RT.emitForStaticInit(*this, S.getLocStart(), ScheduleKind, IVSize,
+                             IVSigned, Ordered, IL.getAddress(),
+                             LB.getAddress(), UB.getAddress(), ST.getAddress(),
+                             Chunk);
         auto LoopExit =
             getJumpDestInCurrentScope(createBasicBlock("omp.loop.exit"));
+        if (S.getDefaultSchedule() != OMPDSK_static_chunkone) {
           // UB = min(UB, GlobalUB);
           EmitIgnoredExpr(S.getEnsureUpperBound());
+        }
         // IV = LB;
         EmitIgnoredExpr(S.getInit());
         // while (idx <= UB) { BODY; ++idx; }
+        // if OMPDSK_static_chunkone:
+        //   while (idx <= GlobalUB) { BODY; idx += ST; }
         EmitOMPInnerLoop(S, LoopScope.requiresCleanups(), S.getCond(),
                          S.getInc(),
                          [&S, LoopExit](CodeGenFunction &CGF) {
                            CGF.EmitOMPLoopBody(S, LoopExit);
                            CGF.EmitStopPoint(&S);
                          },
                          [](CodeGenFunction &) {});
         EmitBlock(LoopExit.getBlock());

lib/Sema/SemaOpenMP.cpp

for (auto *E : PostUpdates) {
                                               PostUpdate, ConvE)
                              .get()
                        : ConvE;
     }
   }
   return PostUpdate;
 }
 
+/// Get the default schedule type for any loop-based OpenMP directive,
+/// specialized for a particular target.  This is used to guide codegen
+/// if a) no 'schedule' clause is specified, or b) a 'schedule' type of
+/// 'auto' is specified by the user.
+static OpenMPDefaultScheduleKind
+getDefaultSchedule(Sema &S, OpenMPDirectiveKind Kind,
+                   ArrayRef<OMPClause *> Clauses) {
+  OpenMPDefaultScheduleKind DefaultSchedule = OMPDSK_unknown;
+
+  if (S.getLangOpts().OpenMPIsDevice &&
+      S.Context.getTargetInfo().getTriple().isNVPTX()) {
+    // Force a schedule type of (static,1) if there is no schedule clause, or
+    // the user specifies schedule(auto) or schedule(static,1).
+    bool ChunkSizeOne = false;
+    auto ScheduleKind = OMPC_SCHEDULE_unknown;
+    auto ScheduleClause =
+        OMPExecutableDirective::getClausesOfKind<OMPScheduleClause>(Clauses);
+    if (ScheduleClause.begin() != ScheduleClause.end()) {
+      ScheduleKind = (*ScheduleClause.begin())->getScheduleKind();
+      if (const auto *Ch = (*ScheduleClause.begin())->getChunkSize()) {
+        if (!Ch->isValueDependent() && !Ch->isTypeDependent() &&
+            !Ch->isInstantiationDependent() &&
+            !Ch->containsUnexpandedParameterPack()) {
+          SourceLocation ChLoc = Ch->getLocStart();
+          ExprResult Val = S.PerformOpenMPImplicitIntegerConversion(
+              ChLoc, const_cast<Expr *>(Ch));
+          if (!Val.isInvalid()) {
+            Expr *ValExpr = Val.get();
+            llvm::APSInt Result;
+            ChunkSizeOne = ValExpr->isIntegerConstantExpr(Result, S.Context) &&
+                           Result == 1;
+          }
+        }
+      }
+    }
+
+    // Ordered clause requires dynamic dispatch.
+    auto OrderedClause =
+        OMPExecutableDirective::getClausesOfKind<OMPOrderedClause>(Clauses);
+    bool Ordered = OrderedClause.begin() != OrderedClause.end();
+
+    bool StaticOneSchedule =
+        (!Ordered && (ScheduleKind == OMPC_SCHEDULE_unknown ||
+                      ScheduleKind == OMPC_SCHEDULE_auto ||
+                      (ScheduleKind == OMPC_SCHEDULE_static && ChunkSizeOne)));
+
+    if (StaticOneSchedule)
+      DefaultSchedule = OMPDSK_static_chunkone;
+  }
+
+  return DefaultSchedule;
+}
+
 /// \brief 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,
+CheckOpenMPLoop(OpenMPDirectiveKind DKind, ArrayRef<OMPClause *> Clauses,
+                Expr *CollapseLoopCountExpr, Expr *OrderedLoopCountExpr,
+                Stmt *AStmt, Sema &SemaRef, DSAStackTy &DSA,
                 llvm::DenseMap<ValueDecl *, Expr *> &VarsWithImplicitDSA,
                 OMPLoopDirective::HelperExprs &Built) {
+  OpenMPDefaultScheduleKind DefaultScheduleKind =
+      getDefaultSchedule(SemaRef, DKind, Clauses);
+
   unsigned NestedLoopCount = 1;
   if (CollapseLoopCountExpr) {
     // Found 'collapse' clause - calculate collapse number.
     llvm::APSInt Result;
     if (CollapseLoopCountExpr->EvaluateAsInt(Result, SemaRef.getASTContext()))
       NestedLoopCount = Result.getLimitedValue();
   }
   if (OrderedLoopCountExpr) {
Expr *RHS =
             : SemaRef.ActOnIntegerConstant(SourceLocation(), 0).get();
     Init = SemaRef.BuildBinOp(CurScope, InitLoc, BO_Assign, IV.get(), RHS);
     Init = SemaRef.ActOnFinishFullExpr(Init.get());
   }
 
   // Loop condition (IV < NumIterations) or (IV <= UB) for worksharing loops.
   SourceLocation CondLoc;
   ExprResult Cond =
+      (DefaultScheduleKind != OMPDSK_static_chunkone &&
        (isOpenMPWorksharingDirective(DKind) ||
-       isOpenMPTaskLoopDirective(DKind) || isOpenMPDistributeDirective(DKind))
+        isOpenMPTaskLoopDirective(DKind) || isOpenMPDistributeDirective(DKind)))
           ? SemaRef.BuildBinOp(CurScope, CondLoc, BO_LE, IV.get(), UB.get())
           : SemaRef.BuildBinOp(CurScope, CondLoc, BO_LT, IV.get(),
                                NumIterations.get());
 
-  // Loop increment (IV = IV + 1)
+  // Loop increment (IV = IV + 1) or (IV = IV + ST) if (static,1) scheduling.
   SourceLocation IncLoc;
   ExprResult Inc =
-      SemaRef.BuildBinOp(CurScope, IncLoc, BO_Add, IV.get(),
+      DefaultScheduleKind == OMPDSK_static_chunkone
+          ? SemaRef.BuildBinOp(CurScope, IncLoc, BO_Add, IV.get(), ST.get())
+          : SemaRef.BuildBinOp(CurScope, IncLoc, BO_Add, IV.get(),
                                SemaRef.ActOnIntegerConstant(IncLoc, 1).get());
   if (!Inc.isUsable())
     return 0;
   Inc = SemaRef.BuildBinOp(CurScope, IncLoc, BO_Assign, IV.get(), Inc.get());
   Inc = SemaRef.ActOnFinishFullExpr(Inc.get());
   if (!Inc.isUsable())
     return 0;
 
   // Increments for worksharing loops (LB = LB + ST; UB = UB + ST).
CheckOpenMPLoop(OpenMPDirectiveKind DKind, ArrayRef<OMPClause *> Clauses,
   Built.UB = UB.get();
   Built.IL = IL.get();
   Built.ST = ST.get();
   Built.EUB = EUB.get();
   Built.NLB = NextLB.get();
   Built.NUB = NextUB.get();
   Built.PrevLB = PrevLB.get();
   Built.PrevUB = PrevUB.get();
+  Built.DefaultScheduleKind = DefaultScheduleKind;
 
   Expr *CounterVal = SemaRef.DefaultLvalueConversion(IV.get()).get();
   // Fill data for doacross depend clauses.
   for (auto Pair : DSA.getDoacrossDependClauses()) {
     if (Pair.first->getDependencyKind() == OMPC_DEPEND_source)
       Pair.first->setCounterValue(CounterVal);
     else {
       if (NestedLoopCount != Pair.second.size() ||
StmtResult Sema::ActOnOpenMPSimdDirective(
     llvm::DenseMap<ValueDecl *, Expr *> &VarsWithImplicitDSA) {
   if (!AStmt)
     return StmtError();
 
   assert(isa<CapturedStmt>(AStmt) && "Captured statement expected");
   OMPLoopDirective::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);
+  unsigned NestedLoopCount =
+      CheckOpenMPLoop(OMPD_simd, Clauses, getCollapseNumberExpr(Clauses),
+                      getOrderedNumberExpr(Clauses), AStmt, *this, *DSAStack,
+                      VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
 
   assert((CurContext->isDependentContext() || B.builtAll()) &&
          "omp simd loop exprs were not built");
 
   if (!CurContext->isDependentContext()) {
     // Finalize the clauses that need pre-built expressions for CodeGen.
StmtResult Sema::ActOnOpenMPForDirective(
     llvm::DenseMap<ValueDecl *, Expr *> &VarsWithImplicitDSA) {
   if (!AStmt)
     return StmtError();
 
   assert(isa<CapturedStmt>(AStmt) && "Captured statement expected");
   OMPLoopDirective::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);
+  unsigned NestedLoopCount =
+      CheckOpenMPLoop(OMPD_for, Clauses, getCollapseNumberExpr(Clauses),
+                      getOrderedNumberExpr(Clauses), AStmt, *this, *DSAStack,
+                      VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
 
   assert((CurContext->isDependentContext() || B.builtAll()) &&
          "omp for loop exprs were not built");
 
   if (!CurContext->isDependentContext()) {
     // Finalize the clauses that need pre-built expressions for CodeGen.
StmtResult Sema::ActOnOpenMPForSimdDirective(
   if (!AStmt)
     return StmtError();
 
   assert(isa<CapturedStmt>(AStmt) && "Captured statement expected");
   OMPLoopDirective::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),
+      CheckOpenMPLoop(OMPD_for_simd, Clauses, getCollapseNumberExpr(Clauses),
                       getOrderedNumberExpr(Clauses), AStmt, *this, *DSAStack,
                       VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
 
   assert((CurContext->isDependentContext() || B.builtAll()) &&
          "omp for simd loop exprs were not built");
 
StmtResult Sema::ActOnOpenMPParallelForDirective(
   // 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;
   // 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);
+  unsigned NestedLoopCount = CheckOpenMPLoop(
+      OMPD_parallel_for, Clauses, getCollapseNumberExpr(Clauses),
   if (NestedLoopCount == 0)
     return StmtError();
 
   assert((CurContext->isDependentContext() || B.builtAll()) &&
          "omp parallel for loop exprs were not built");
 
   if (!CurContext->isDependentContext()) {
     // Finalize the clauses that need pre-built expressions for CodeGen.
StmtResult Sema::ActOnOpenMPParallelForSimdDirective(
   // 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;
   // 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);
+  unsigned NestedLoopCount = CheckOpenMPLoop(
+      OMPD_parallel_for_simd, Clauses, getCollapseNumberExpr(Clauses),
   if (NestedLoopCount == 0)
     return StmtError();
 
   if (!CurContext->isDependentContext()) {
     // Finalize the clauses that need pre-built expressions for CodeGen.
     for (auto C : Clauses) {
       if (auto *LC = dyn_cast<OMPLinearClause>(C))
         if (FinishOpenMPLinearClause(*LC, cast<DeclRefExpr>(B.IterationVarRef),
StmtResult Sema::ActOnOpenMPTargetParallelForDirective(
   // 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;
   // 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), AStmt, *this, *DSAStack,
-                      VarsWithImplicitDSA, B);
+  unsigned NestedLoopCount = CheckOpenMPLoop(
+      OMPD_target_parallel_for, Clauses, getCollapseNumberExpr(Clauses),
   if (NestedLoopCount == 0)
     return StmtError();
 
   assert((CurContext->isDependentContext() || B.builtAll()) &&
          "omp target parallel for loop exprs were not built");
 
   if (!CurContext->isDependentContext()) {
     // Finalize the clauses that need pre-built expressions for CodeGen.
StmtResult Sema::ActOnOpenMPTaskLoopDirective(
   if (!AStmt)
     return StmtError();
 
   assert(isa<CapturedStmt>(AStmt) && "Captured statement expected");
   OMPLoopDirective::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),
+      CheckOpenMPLoop(OMPD_taskloop, Clauses, getCollapseNumberExpr(Clauses),
                       /*OrderedLoopCountExpr=*/nullptr, AStmt, *this, *DSAStack,
                       VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
 
   assert((CurContext->isDependentContext() || B.builtAll()) &&
          "omp for loop exprs were not built");
 
StmtResult Sema::ActOnOpenMPTaskLoopSimdDirective(
     llvm::DenseMap<ValueDecl *, Expr *> &VarsWithImplicitDSA) {
   if (!AStmt)
     return StmtError();
 
   assert(isa<CapturedStmt>(AStmt) && "Captured statement expected");
   OMPLoopDirective::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);
+  unsigned NestedLoopCount = CheckOpenMPLoop(
+      OMPD_taskloop_simd, Clauses, getCollapseNumberExpr(Clauses),
   if (NestedLoopCount == 0)
     return StmtError();
 
   assert((CurContext->isDependentContext() || B.builtAll()) &&
          "omp for loop exprs were not built");
 
   if (!CurContext->isDependentContext()) {
     // Finalize the clauses that need pre-built expressions for CodeGen.
StmtResult Sema::ActOnOpenMPDistributeDirective(
   if (!AStmt)
     return StmtError();
 
   assert(isa<CapturedStmt>(AStmt) && "Captured statement expected");
   OMPLoopDirective::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),
+      CheckOpenMPLoop(OMPD_distribute, Clauses, getCollapseNumberExpr(Clauses),
                       nullptr /*ordered not a clause on distribute*/, AStmt,
                       *this, *DSAStack, VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
 
   assert((CurContext->isDependentContext() || B.builtAll()) &&
          "omp for loop exprs were not built");
 
StmtResult Sema::ActOnOpenMPDistributeParallelForDirective(
   // 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;
   // 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),
+      OMPD_distribute_parallel_for, Clauses, getCollapseNumberExpr(Clauses),
       nullptr /*ordered not a clause on distribute*/, AStmt, *this, *DSAStack,
       VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
 
   assert((CurContext->isDependentContext() || B.builtAll()) &&
          "omp for loop exprs were not built");
 
StmtResult Sema::ActOnOpenMPDistributeParallelForSimdDirective(
   // 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;
   // 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*/, AStmt, *this, *DSAStack,
-      VarsWithImplicitDSA, B);
+  unsigned NestedLoopCount =
+      CheckOpenMPLoop(OMPD_distribute_parallel_for_simd, Clauses,
+                      getCollapseNumberExpr(Clauses),
+                      nullptr /*ordered not a clause on distribute*/, AStmt,
+                      *this, *DSAStack, VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
 
   assert((CurContext->isDependentContext() || B.builtAll()) &&
          "omp for loop exprs were not built");
 
   if (checkSimdlenSafelenSpecified(*this, Clauses))
     return StmtError();
StmtResult Sema::ActOnOpenMPDistributeSimdDirective(
   // 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;
   // 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*/, AStmt,
-                      *this, *DSAStack, VarsWithImplicitDSA, B);
+  unsigned NestedLoopCount = CheckOpenMPLoop(
+      OMPD_distribute_simd, Clauses, getCollapseNumberExpr(Clauses),
+      nullptr /*ordered not a clause on distribute*/, AStmt, *this, *DSAStack,
+      VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
 
   assert((CurContext->isDependentContext() || B.builtAll()) &&
          "omp for loop exprs were not built");
 
   if (checkSimdlenSafelenSpecified(*this, Clauses))
     return StmtError();
StmtResult Sema::ActOnOpenMPTargetParallelForSimdDirective(
   // 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;
   // 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),
+      OMPD_target_parallel_for_simd, Clauses, getCollapseNumberExpr(Clauses),
       getOrderedNumberExpr(Clauses), AStmt, *this, *DSAStack,
       VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
 
   assert((CurContext->isDependentContext() || B.builtAll()) &&
          "omp target parallel for simd loop exprs were not built");
 
StmtResult Sema::ActOnOpenMPTargetSimdDirective(
   // 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;
   // In presence of clause 'collapse' with number of loops, it will define the
   // nested loops number.
   unsigned NestedLoopCount =
-      CheckOpenMPLoop(OMPD_target_simd, getCollapseNumberExpr(Clauses),
+      CheckOpenMPLoop(OMPD_target_simd, Clauses, getCollapseNumberExpr(Clauses),
                       getOrderedNumberExpr(Clauses), AStmt, *this, *DSAStack,
                       VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
 
   assert((CurContext->isDependentContext() || B.builtAll()) &&
          "omp target simd loop exprs were not built");
 
StmtResult Sema::ActOnOpenMPTeamsDistributeDirective(
   // 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;
   // 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*/, AStmt,
-                      *this, *DSAStack, VarsWithImplicitDSA, B);
+  unsigned NestedLoopCount = CheckOpenMPLoop(
+      OMPD_teams_distribute, Clauses, getCollapseNumberExpr(Clauses),
+      nullptr /*ordered not a clause on distribute*/, AStmt, *this, *DSAStack,
+      VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
 
   assert((CurContext->isDependentContext() || B.builtAll()) &&
          "omp teams distribute loop exprs were not built");
 
   getCurFunction()->setHasBranchProtectedScope();
   return OMPTeamsDistributeDirective::Create(
StmtResult Sema::ActOnOpenMPTeamsDistributeSimdDirective(
   // 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;
   // 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),
+      OMPD_teams_distribute_simd, Clauses, getCollapseNumberExpr(Clauses),
       nullptr /*ordered not a clause on distribute*/, AStmt, *this, *DSAStack,
       VarsWithImplicitDSA, B);
 
   if (NestedLoopCount == 0)
     return StmtError();
 
   assert((CurContext->isDependentContext() || B.builtAll()) &&
          "omp teams distribute simd loop exprs were not built");
StmtResult Sema::ActOnOpenMPTeamsDistributeParallelForSimdDirective(
   // 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;
   // In presence of clause 'collapse' with number of loops, it will
   // define the nested loops number.
-  auto NestedLoopCount = CheckOpenMPLoop(
-      OMPD_teams_distribute_parallel_for_simd, getCollapseNumberExpr(Clauses),
-      nullptr /*ordered not a clause on distribute*/, AStmt, *this, *DSAStack,
-      VarsWithImplicitDSA, B);
+  auto NestedLoopCount =
+      CheckOpenMPLoop(OMPD_teams_distribute_parallel_for_simd, Clauses,
+                      getCollapseNumberExpr(Clauses),
+                      nullptr /*ordered not a clause on distribute*/, AStmt,
+                      *this, *DSAStack, VarsWithImplicitDSA, B);
 
   if (NestedLoopCount == 0)
     return StmtError();
 
   assert((CurContext->isDependentContext() || B.builtAll()) &&
          "omp for loop exprs were not built");
 
   if (!CurContext->isDependentContext()) {
StmtResult Sema::ActOnOpenMPTeamsDistributeParallelForDirective(
   // 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;
   // 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*/, AStmt, *this, *DSAStack,
-      VarsWithImplicitDSA, B);
+  unsigned NestedLoopCount =
+      CheckOpenMPLoop(OMPD_teams_distribute_parallel_for, Clauses,
+                      getCollapseNumberExpr(Clauses),
+                      nullptr /*ordered not a clause on distribute*/, AStmt,
+                      *this, *DSAStack, VarsWithImplicitDSA, B);
 
   if (NestedLoopCount == 0)
     return StmtError();
 
   assert((CurContext->isDependentContext() || B.builtAll()) &&
          "omp for loop exprs were not built");
 
   if (!CurContext->isDependentContext()) {
StmtResult Sema::ActOnOpenMPTargetTeamsDistributeDirective(
   // 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;
   // In presence of clause 'collapse' with number of loops, it will
   // define the nested loops number.
   auto NestedLoopCount = CheckOpenMPLoop(
-      OMPD_target_teams_distribute,
-      getCollapseNumberExpr(Clauses),
+      OMPD_target_teams_distribute, Clauses, getCollapseNumberExpr(Clauses),
       nullptr /*ordered not a clause on distribute*/, AStmt, *this, *DSAStack,
       VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
 
   assert((CurContext->isDependentContext() || B.builtAll()) &&
          "omp target teams distribute loop exprs were not built");
 
StmtResult Sema::ActOnOpenMPTargetTeamsDistributeParallelForDirective(
   // 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;
   // In presence of clause 'collapse' with number of loops, it will
   // define the nested loops number.
-  auto NestedLoopCount = CheckOpenMPLoop(
-      OMPD_target_teams_distribute_parallel_for,
+  auto NestedLoopCount =
+      CheckOpenMPLoop(OMPD_target_teams_distribute_parallel_for, Clauses,
                       getCollapseNumberExpr(Clauses),
-      nullptr /*ordered not a clause on distribute*/, AStmt, *this, *DSAStack,
-      VarsWithImplicitDSA, B);
+                      nullptr /*ordered not a clause on distribute*/, AStmt,
+                      *this, *DSAStack, VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
 
   assert((CurContext->isDependentContext() || B.builtAll()) &&
          "omp target teams distribute parallel for loop exprs were not built");
 
   if (!CurContext->isDependentContext()) {
     // Finalize the clauses that need pre-built expressions for CodeGen.
StmtResult Sema::ActOnOpenMPTargetTeamsDistributeParallelForSimdDirective(
   // 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;
   // In presence of clause 'collapse' with number of loops, it will
   // define the nested loops number.
-  auto NestedLoopCount = CheckOpenMPLoop(
-      OMPD_target_teams_distribute_parallel_for_simd,
+  auto NestedLoopCount =
+      CheckOpenMPLoop(OMPD_target_teams_distribute_parallel_for_simd, Clauses,
                       getCollapseNumberExpr(Clauses),
-      nullptr /*ordered not a clause on distribute*/, AStmt, *this, *DSAStack,
-      VarsWithImplicitDSA, B);
+                      nullptr /*ordered not a clause on distribute*/, AStmt,
+                      *this, *DSAStack, VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
 
   assert((CurContext->isDependentContext() || B.builtAll()) &&
          "omp target teams distribute parallel for simd loop exprs were not "
          "built");
 
   if (!CurContext->isDependentContext()) {
StmtResult Sema::ActOnOpenMPTargetTeamsDistributeSimdDirective(
   // 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;
   // In presence of clause 'collapse' with number of loops, it will
   // define the nested loops number.
-  auto NestedLoopCount = CheckOpenMPLoop(
-      OMPD_target_teams_distribute_simd, getCollapseNumberExpr(Clauses),
-      nullptr /*ordered not a clause on distribute*/, AStmt, *this, *DSAStack,
-      VarsWithImplicitDSA, B);
+  auto NestedLoopCount =
+      CheckOpenMPLoop(OMPD_target_teams_distribute_simd, Clauses,
+                      getCollapseNumberExpr(Clauses),
+                      nullptr /*ordered not a clause on distribute*/, AStmt,
+                      *this, *DSAStack, VarsWithImplicitDSA, B);
   if (NestedLoopCount == 0)
     return StmtError();
 
   assert((CurContext->isDependentContext() || B.builtAll()) &&
          "omp target teams distribute simd loop exprs were not built");
 
   getCurFunction()->setHasBranchProtectedScope();
   return OMPTargetTeamsDistributeSimdDirective::Create(

test/OpenMP/nvptx_coalesced_scheduling_codegen.cpp

+// Test target codegen - host bc file has to be created first.
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple nvptx64-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
+// RUN: %clang_cc1 -verify -fopenmp -fopenmp-version=45 -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+// Check that the execution mode of the target regions on the gpu is set to the right mode.
+// CHECK-DAG: {{@__omp_offloading_.+l19}}_exec_mode = weak constant i8 0
+
+template<typename tx>
+tx ftemplate() {
+  tx a[100];
+  tx b[10][10];
+
+  #pragma omp target parallel
+  {
+    #pragma omp for
+    for (int i = 0; i < 99; i++) {
+      a[i] = 1;
+    }
+
+    #pragma omp for schedule(auto)
+    for (int i = 0; i < 98; i++) {
+      a[i] = 2;
+    }
+
+    #pragma omp for schedule(static,1)
+    for (int i = 0; i < 97; i++) {
+      a[i] = 3;
+    }
+
+    #pragma omp for schedule(static,2)
+    for (int i = 0; i < 96; i++) {
+      a[i] = 1;
+    }
+
+    #pragma omp for schedule(static)
+    for (int i = 0; i < 95; i++) {
+      a[i] = 1;
+    }
+
+    #pragma omp for schedule(auto) ordered
+    for (int i = 0; i < 94; i++) {
+      a[i] = 1;
+    }
+
+    #pragma omp for schedule(runtime)
+    for (int i = 0; i < 93; i++) {
+      a[i] = 1;
+    }
+
+    #pragma omp for schedule(dynamic)
+    for (int i = 0; i < 92; i++) {
+      a[i] = 1;
+    }
+
+    #pragma omp for schedule(guided)
+    for (int i = 0; i < 91; i++) {
+      a[i] = 1;
+    }
+  }
+
+  return a[0] + b[9][9];
+}
+
+int bar(){
+  int a = 0;
+
+  a += ftemplate<int>();
+
+  return a;
+}
+
+  // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l19}}(
+  // CHECK: call void @__kmpc_spmd_kernel_init(
+  // CHECK: br label {{%?}}[[EXEC:.+]]
+  //
+  // CHECK: [[EXEC]]
+  // CHECK: {{call|invoke}} void [[OP1:@.+]](i32*
+  // CHECK: br label {{%?}}[[DONE:.+]]
+  //
+  // CHECK: [[DONE]]
+  // CHECK: call void @__kmpc_spmd_kernel_deinit()
+  // CHECK: br label {{%?}}[[EXIT:.+]]
+  //
+  // CHECK: [[EXIT]]
+  // CHECK: ret void
+  // CHECK: }
+
+  // CHECK: define internal void [[OP1]](
+
+  // No schedule clause.
+  //
+  // CHECK: store i32 0, i32* [[LB_PTR:%.+]], align
+  // CHECK: store i32 98, i32* [[UB_PTR:%.+]], align
+  // CHECK: store i32 1, i32* [[ST_PTR:%.+]], align
+  // CHECK: call void @__kmpc_for_static_init_4(%ident_t* {{@.+}}, i32 {{%.+}}, i32 33, i32* {{%.+}}, i32* [[LB_PTR]], i32* [[UB_PTR]], i32* [[ST_PTR]], i32 1, i32 1)
+  // CHECK: [[LB:%.+]] = load i32, i32* [[LB_PTR]], align
+  // CHECK: store i32 [[LB]], i32* [[IV_PTR:%.+]], align
+  // CHECK: br label {{%?}}[[FOR_COND:.+]]
+  //
+  // CHECK: [[FOR_COND]]
+  // CHECK: [[IV:%.+]] = load i32, i32* [[IV_PTR]], align
+  // CHECK: [[CMP:%.+]] = icmp slt i32 [[IV]], 99
+  // CHECK: br i1 [[CMP]], label {{%?}}[[FOR_BODY:.+]], label {{%?}}[[FOR_END:.+]]
+  //
+  // [[FOR_BODY]]
+  // CHECK: [[IV:%.+]] = load i32, i32* [[IV_PTR]], align
+  // CHECK: [[MUL:%.+]] = mul nsw i32 [[IV]], 1
+  // CHECK: [[ADD:%.+]] = add nsw i32 0, [[MUL]]
+  // CHECK: store i32 [[ADD]], i32* [[I_PTR:%.+]], align
+  // CHECK: [[I:%.+]] = load i32, i32* [[I_PTR]], align
+  // CHECK-32: [[ELEM_PTR:%.+]] = getelementptr inbounds [100 x i32], [100 x i32]* {{%.+}}, i32 0, i32 [[I]]
+  // CHECK-64: [[IDX:%.+]] = sext i32 [[I]] to i64
+  // CHECK-64: [[ELEM_PTR:%.+]] = getelementptr inbounds [100 x i32], [100 x i32]* {{%.+}}, i64 0, i64 [[IDX]]
+  // CHECK: store i32 1, i32* [[ELEM_PTR]], align
+  // CHECK: br label {{%?}}[[FOR_CONT:.+]]
+  //
+  // CHECK: [[FOR_CONT]]
+  // CHECK: br label {{%?}}[[FOR_INC:.+]]
+  //
+  // CHECK: [[FOR_INC]]
+  // CHECK: [[IV:%.+]] = load i32, i32* [[IV_PTR]], align
+  // CHECK: [[ST:%.+]] = load i32, i32* [[ST_PTR]], align
+  // CHECK: [[ADD:%.+]] = add nsw i32 [[IV]], [[ST]]
+  // CHECK: store i32 [[ADD]], i32* [[IV_PTR]], align
+  // CHECK: br label {{%?}}[[FOR_COND]]
+  //
+  // CHECK: [[FOR_END]]
+
+
+
+  // schedule(auto) clause.
+  //
+  // CHECK: store i32 0, i32* [[LB_PTR:%.+]], align
+  // CHECK: store i32 97, i32* [[UB_PTR:%.+]], align
+  // CHECK: store i32 1, i32* [[ST_PTR:%.+]], align
+  // CHECK: call void @__kmpc_for_static_init_4(%ident_t* {{@.+}}, i32 {{%.+}}, i32 33, i32* {{%.+}}, i32* [[LB_PTR]], i32* [[UB_PTR]], i32* [[ST_PTR]], i32 1, i32 1)
+  // CHECK: [[LB:%.+]] = load i32, i32* [[LB_PTR]], align
+  // CHECK: store i32 [[LB]], i32* [[IV_PTR:%.+]], align
+  // CHECK: br label {{%?}}[[FOR_COND:.+]]
+  //
+  // CHECK: [[FOR_COND]]
+  // CHECK: [[IV:%.+]] = load i32, i32* [[IV_PTR]], align
+  // CHECK: [[CMP:%.+]] = icmp slt i32 [[IV]], 98
+  // CHECK: br i1 [[CMP]], label {{%?}}[[FOR_BODY:.+]], label {{%?}}[[FOR_END:.+]]
+  //
+  // [[FOR_BODY]]
+  // CHECK: [[IV:%.+]] = load i32, i32* [[IV_PTR]], align
+  // CHECK: [[MUL:%.+]] = mul nsw i32 [[IV]], 1
+  // CHECK: [[ADD:%.+]] = add nsw i32 0, [[MUL]]
+  // CHECK: store i32 [[ADD]], i32* [[I_PTR:%.+]], align
+  // CHECK: [[I:%.+]] = load i32, i32* [[I_PTR]], align
+  // CHECK-32: [[ELEM_PTR:%.+]] = getelementptr inbounds [100 x i32], [100 x i32]* {{%.+}}, i32 0, i32 [[I]]
+  // CHECK-64: [[IDX:%.+]] = sext i32 [[I]] to i64
+  // CHECK-64: [[ELEM_PTR:%.+]] = getelementptr inbounds [100 x i32], [100 x i32]* {{%.+}}, i64 0, i64 [[IDX]]
+  // CHECK: store i32 2, i32* [[ELEM_PTR]], align
+  // CHECK: br label {{%?}}[[FOR_CONT:.+]]
+  //
+  // CHECK: [[FOR_CONT]]
+  // CHECK: br label {{%?}}[[FOR_INC:.+]]
+  //
+  // CHECK: [[FOR_INC]]
+  // CHECK: [[IV:%.+]] = load i32, i32* [[IV_PTR]], align
+  // CHECK: [[ST:%.+]] = load i32, i32* [[ST_PTR]], align
+  // CHECK: [[ADD:%.+]] = add nsw i32 [[IV]], [[ST]]
+  // CHECK: store i32 [[ADD]], i32* [[IV_PTR]], align
+  // CHECK: br label {{%?}}[[FOR_COND]]
+  //
+  // CHECK: [[FOR_END]]
+
+
+
+  // schedule(static,1) clause.
+  //
+  // CHECK: store i32 0, i32* [[LB_PTR:%.+]], align
+  // CHECK: store i32 96, i32* [[UB_PTR:%.+]], align
+  // CHECK: store i32 1, i32* [[ST_PTR:%.+]], align
+  // CHECK: call void @__kmpc_for_static_init_4(%ident_t* {{@.+}}, i32 {{%.+}}, i32 33, i32* {{%.+}}, i32* [[LB_PTR]], i32* [[UB_PTR]], i32* [[ST_PTR]], i32 1, i32 1)
+  // CHECK: [[LB:%.+]] = load i32, i32* [[LB_PTR]], align
+  // CHECK: store i32 [[LB]], i32* [[IV_PTR:%.+]], align
+  // CHECK: br label {{%?}}[[FOR_COND:.+]]
+  //
+  // CHECK: [[FOR_COND]]
+  // CHECK: [[IV:%.+]] = load i32, i32* [[IV_PTR]], align
+  // CHECK: [[CMP:%.+]] = icmp slt i32 [[IV]], 97
+  // CHECK: br i1 [[CMP]], label {{%?}}[[FOR_BODY:.+]], label {{%?}}[[FOR_END:.+]]
+  //
+  // [[FOR_BODY]]
+  // CHECK: [[IV:%.+]] = load i32, i32* [[IV_PTR]], align
+  // CHECK: [[MUL:%.+]] = mul nsw i32 [[IV]], 1
+  // CHECK: [[ADD:%.+]] = add nsw i32 0, [[MUL]]
+  // CHECK: store i32 [[ADD]], i32* [[I_PTR:%.+]], align
+  // CHECK: [[I:%.+]] = load i32, i32* [[I_PTR]], align
+  // CHECK-32: [[ELEM_PTR:%.+]] = getelementptr inbounds [100 x i32], [100 x i32]* {{%.+}}, i32 0, i32 [[I]]
+  // CHECK-64: [[IDX:%.+]] = sext i32 [[I]] to i64
+  // CHECK-64: [[ELEM_PTR:%.+]] = getelementptr inbounds [100 x i32], [100 x i32]* {{%.+}}, i64 0, i64 [[IDX]]
+  // CHECK: store i32 3, i32* [[ELEM_PTR]], align
+  // CHECK: br label {{%?}}[[FOR_CONT:.+]]
+  //
+  // CHECK: [[FOR_CONT]]
+  // CHECK: br label {{%?}}[[FOR_INC:.+]]
+  //
+  // CHECK: [[FOR_INC]]
+  // CHECK: [[IV:%.+]] = load i32, i32* [[IV_PTR]], align
+  // CHECK: [[ST:%.+]] = load i32, i32* [[ST_PTR]], align
+  // CHECK: [[ADD:%.+]] = add nsw i32 [[IV]], [[ST]]
+  // CHECK: store i32 [[ADD]], i32* [[IV_PTR]], align
+  // CHECK: br label {{%?}}[[FOR_COND]]
+  //
+  // CHECK: [[FOR_END]]
+
+
+
+  // schedule(static,2) clause.  Non-coalesced codegen.
+  //
+  // CHECK: store i32 0, i32* [[LB_PTR:%.+]], align
+  // CHECK: store i32 95, i32* [[UB_PTR:%.+]], align
+  // CHECK: store i32 1, i32* [[ST_PTR:%.+]], align
+  // CHECK: call void @__kmpc_for_static_init_4(%ident_t* {{@.+}}, i32 {{%.+}}, i32 33, i32* {{%.+}}, i32* [[LB_PTR]], i32* [[UB_PTR]], i32* [[ST_PTR]], i32 1, i32 2)
+  // CHECK: br label {{%?}}[[DISPATCH_COND:.+]]
+  //
+  // CHECK: [[DISPATCH_COND]]
+  // CHECK: [[UB:%.+]] = load i32, i32* [[UB_PTR]], align
+  // CHECK: = icmp sgt i32 [[UB]], 95
+  //
+  // CHECK: [[LB:%.+]] = load i32, i32* [[LB_PTR]], align
+  // CHECK: store i32 [[LB]], i32* [[IV_PTR:%.+]], align
+  //
+  // CHECK:  = getelementptr
+  //
+  // CHECK: [[IV:%.+]] = load i32, i32* [[IV_PTR]], align
+  // CHECK: [[ADD:%.+]] = add nsw i32 [[IV]], 1
+  // CHECK: store i32 [[ADD]], i32* [[IV_PTR]], align
+
+
+
+  // schedule(static) clause.  Non-coalesced codegen.
+  //
+  // CHECK: store i32 0, i32* [[LB_PTR:%.+]], align
+  // CHECK: store i32 94, i32* [[UB_PTR:%.+]], align
+  // CHECK: store i32 1, i32* [[ST_PTR:%.+]], align
+  // CHECK: call void @__kmpc_for_static_init_4(%ident_t* {{@.+}}, i32 {{%.+}}, i32 34, i32* {{%.+}}, i32* [[LB_PTR]], i32* [[UB_PTR]], i32* [[ST_PTR]], i32 1, i32 1)
+  // CHECK: [[UB:%.+]] = load i32, i32* [[UB_PTR]], align
+  // CHECK: = icmp sgt i32 [[UB]], 94
+  //
+  // CHECK: [[LB:%.+]] = load i32, i32* [[LB_PTR]], align
+  // CHECK: store i32 [[LB]], i32* [[IV_PTR:%.+]], align
+  //
+  // CHECK:  = getelementptr
+  //
+  // CHECK: [[IV:%.+]] = load i32, i32* [[IV_PTR]], align
+  // CHECK: [[ADD:%.+]] = add nsw i32 [[IV]], 1
+  // CHECK: store i32 [[ADD]], i32* [[IV_PTR]], align
+
+
+
+  // schedule(auto) ordered clause.  Non-coalesced codegen.
+  //
+  // CHECK: store i32 0, i32* [[LB_PTR:%.+]], align
+  // CHECK: store i32 93, i32* [[UB_PTR:%.+]], align
+  // CHECK: store i32 1, i32* [[ST_PTR:%.+]], align
+  // CHECK: call void @__kmpc_dispatch_init_4(%ident_t* {{@.+}}, i32 {{%.+}}, i32 70
+  // CHECK: call i32 @__kmpc_dispatch_next_4(%ident_t* {{@.+}}, i32 {{%.+}}, i32* {{%.+}}, i32* [[LB_PTR]], i32* [[UB_PTR]], i32* [[ST_PTR]])
+  //
+  // CHECK: [[LB:%.+]] = load i32, i32* [[LB_PTR]], align
+  // CHECK: store i32 [[LB]], i32* [[IV_PTR:%.+]], align
+  //
+  // CHECK:  = getelementptr
+  //
+  // CHECK: [[IV:%.+]] = load i32, i32* [[IV_PTR]], align
+  // CHECK: [[ADD:%.+]] = add nsw i32 [[IV]], 1
+  // CHECK: store i32 [[ADD]], i32* [[IV_PTR]], align
+
+
+
+  // schedule(runtime) clause.  Non-coalesced codegen.
+  //
+  // CHECK: store i32 0, i32* [[LB_PTR:%.+]], align
+  // CHECK: store i32 92, i32* [[UB_PTR:%.+]], align
+  // CHECK: store i32 1, i32* [[ST_PTR:%.+]], align
+  // CHECK: call void @__kmpc_dispatch_init_4(%ident_t* {{@.+}}, i32 {{%.+}}, i32 37
+  // CHECK: call i32 @__kmpc_dispatch_next_4(%ident_t* {{@.+}}, i32 {{%.+}}, i32* {{%.+}}, i32* [[LB_PTR]], i32* [[UB_PTR]], i32* [[ST_PTR]])
+  //
+  // CHECK: [[LB:%.+]] = load i32, i32* [[LB_PTR]], align
+  // CHECK: store i32 [[LB]], i32* [[IV_PTR:%.+]], align
+  //
+  // CHECK:  = getelementptr
+  //
+  // CHECK: [[IV:%.+]] = load i32, i32* [[IV_PTR]], align
+  // CHECK: [[ADD:%.+]] = add nsw i32 [[IV]], 1
+  // CHECK: store i32 [[ADD]], i32* [[IV_PTR]], align
+
+
+
+  // schedule(dynamic) clause.  Non-coalesced codegen.
+  //
+  // CHECK: store i32 0, i32* [[LB_PTR:%.+]], align
+  // CHECK: store i32 91, i32* [[UB_PTR:%.+]], align
+  // CHECK: store i32 1, i32* [[ST_PTR:%.+]], align
+  // CHECK: call void @__kmpc_dispatch_init_4(%ident_t* {{@.+}}, i32 {{%.+}}, i32 35
+  // CHECK: call i32 @__kmpc_dispatch_next_4(%ident_t* {{@.+}}, i32 {{%.+}}, i32* {{%.+}}, i32* [[LB_PTR]], i32* [[UB_PTR]], i32* [[ST_PTR]])
+  //
+  // CHECK: [[LB:%.+]] = load i32, i32* [[LB_PTR]], align
+  // CHECK: store i32 [[LB]], i32* [[IV_PTR:%.+]], align
+  //
+  // CHECK:  = getelementptr
+  //
+  // CHECK: [[IV:%.+]] = load i32, i32* [[IV_PTR]], align
+  // CHECK: [[ADD:%.+]] = add nsw i32 [[IV]], 1
+  // CHECK: store i32 [[ADD]], i32* [[IV_PTR]], align
+
+
+
+  // CHECK: ret void
+  // CHECK: }
+
+#endif