[Clang][OpenMP] Support for Code Generation of loop bind clause

[koops]

Summary

Support for Code Generation of "#pragma loop bind" clause.

1. bind(parallel)
2. bind(teams)
3. bind(thread)

[koops]

bind(thread) is not working at present. I have uploaded this patch to obtain feedback mainly on this.

[jdoerfert]

We need tests.

clang/lib/CodeGen/CGStmtOpenMP.cpp

7808

style, also below.

[ronlieb]

LIT Test ?

[ddpagan]

How is bind(thread) not working? Runtime issue?

[carlo.bertolli]

clang/lib/CodeGen/CGStmtOpenMP.cpp

7809

What if you have something like this:

void foo(..) {

#pragma omp loop
for (..) {
}

}

int main() {

#pragma omp target teams
{
  ...
  foo();
 }

#pragma omp target
 {
    #pragma omp parallel
    {
       foo();
    }

}

In the first invocation, loop is bound to teams. In the second, loop is bound to parallel.
This is a runtime condition.

I *believe* that not even OpenMP 6.0 TR1 allows us to decide at compile time if loop should be treated as worksharing or workdistribution....but I might be wrong.

Thanks!

[koops]

clang/lib/CodeGen/CGStmtOpenMP.cpp

7809

The loop directive does not have a bind clause. If you examine the current patch it takes care to preserve the old structure as is and the binding is done by default during runtime.

[carlo.bertolli]

clang/lib/CodeGen/CGStmtOpenMP.cpp

7809

Do you mean that this patch adds up to existing support - adding support for bind - and that existing support is already based on runtime support to decide on the case I mentioned?

[carlo.bertolli]

clang/lib/CodeGen/CGStmtOpenMP.cpp

7809

Disregard my comment: essentially, loop cannot be orphaned in the way I described.

[koops]

1. formatting
2. Adding lit test
3. Removing bind clause from the set of clauses passed during bind(parallel) to the OMPForDirective and bind(teams) to the OMPDistributeDirective.

[ABataev]

clang/lib/AST/StmtOpenMP.cpp

2362

We use tail allocation, this is definetely a mem leak.

clang/lib/CodeGen/CGStmtOpenMP.cpp

7815–7851

This is a bad place for this kind of operation here, it should be handled in Sema

[koops]

Moving the code to SemaOpenMP.cpp from CodeGen.

[koops]

Removing changes from :

1. clang/include/clang/AST/StmtOpenMP.h
2. clang/lib/AST/StmtOpenMP.cpp
3. clang/include/clang/Parse/Parser.h

These were useful when the code was in CodeGen to handle the bind clause.

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

6124 ↗

(On Diff #512448)

Remove this new line

6181–6189 ↗

(On Diff #512448)

1. You're overriding the directive kind here and do not restore it then. It may cause the compiler crash.
2. I think you need to here to create a new OpenMP region rather than overriding the existing one.

clang/test/OpenMP/loop_bind_codegen.cpp

1 ↗

(On Diff #512448)

Remove this

51–56 ↗

(On Diff #512448)

I think it should trigger the assert in setCurrentDirective

[koops]

clang/test/OpenMP/loop_bind_codegen.cpp

51–56 ↗

(On Diff #512448)

I did not see any crash at this point. Can you please give me other examples where I can see this crash?

[ABataev]

clang/include/clang/Basic/DiagnosticSemaKinds.td

9792 ↗

(On Diff #512448)

err_omp_...

clang/lib/Sema/SemaOpenMP.cpp

6146 ↗

(On Diff #512448)

ParentDirective

6167–6168 ↗

(On Diff #512448)

Remove extra parens

6173–6175 ↗

(On Diff #512448)

Remove extra braces

6181–6189 ↗

(On Diff #512448)

I see why you're doing this now. Would be good to see some semantic tests, like nesting of regions, use of unsupported clauses (e.g. linear on loop bind(teams)), etc.

6199–6202 ↗

(On Diff #512448)

Why need to drop bind clause here?

clang/test/OpenMP/loop_bind_codegen.cpp

51–56 ↗

(On Diff #512448)

Ah, you're overriding the loop directive kind.

[koops]

1. Adding semantic test clang/test/OpenMP/loop_bind_messages.cpp.
2. Changes suggested by Alexey.
3. >Why need to drop bind clause here? The new Directives to which loop directive is being mapped to, do not contain the bind clause. (a) omp loop bind(parallel) ----> omp for (b) omp loop bind(teams) -----> omp distribute (c) omp loop bind(thread) ------> omp simd

[ABataev]

In D144634#4272857, @koops wrote:

1. Adding semantic test clang/test/OpenMP/loop_bind_messages.cpp.
2. Changes suggested by Alexey.
3. >Why need to drop bind clause here? The new Directives to which loop directive is being mapped to, do not contain the bind clause. (a) omp loop bind(parallel) ----> omp for (b) omp loop bind(teams) -----> omp distribute (c) omp loop bind(thread) ------> omp simd

But why need to drop it? It makes processing more complex. The bind clause itself should not be a problem, no?

clang/lib/Sema/SemaOpenMP.cpp

638 ↗

(On Diff #514096)

NewDK

640 ↗

(On Diff #514096)

Add message to the assert.

clang/test/OpenMP/loop_bind_codegen.cpp

2–3 ↗

(On Diff #514096)

In other tests, this is used for testing with precompiled headers. You don't have this in your test runs currently.

8–13 ↗

(On Diff #514096)

Remove this commented code too

clang/test/OpenMP/loop_bind_messages.cpp

1–2 ↗

(On Diff #514096)

1. No need for this in semantic test

2, More tests required, nesting of regions, more tests for incompatible clauses.

[ABataev]

Also, add AST printing tests

[koops]

1. Taken care of some of the comments by Alexey.
2. Added extra tests of loop_bind_messages.cpp

However, generic_loop_messages.cpp & generic_loop_ast_print.cpp are present and provide a good coverage.

[ABataev]

In D144634#4308872, @koops wrote:

1. Taken care of some of the comments by Alexey.
2. Added extra tests of loop_bind_messages.cpp

However, generic_loop_messages.cpp & generic_loop_ast_print.cpp are present and provide a good coverage.

I rather doubt that these tests provide good coverage, since you're changing the directive kind here on the fly. This is a very new functionality, which was not tested before.
Add the tests for nesting of the regions, i.e. loop bind directive, enclosed in different regions, and diagnostics.

clang/lib/Sema/SemaOpenMP.cpp

339 ↗

(On Diff #518341)

What Ma1D stands for, what is the meaning of this abbrev? Add a comment.

clang/test/OpenMP/loop_bind_codegen.cpp

2 ↗

(On Diff #518341)

Add tests with save/load precompiled code.

clang/test/OpenMP/loop_bind_messages.cpp

3 ↗

(On Diff #518341)

Add tests with save/load precompiled code.

[koops]

However, generic_loop_messages.cpp & generic_loop_ast_print.cpp are present and provide a good coverage.

I rather doubt that these tests provide good coverage, since you're changing the directive kind here on the fly. This is a very new functionality, which was not tested before.
Add the tests for nesting of the regions, i.e. loop bind directive, enclosed in different regions, and diagnostics.

In https://reviews.llvm.org/D145823 : Add codegen for combined 'loop' directives, I see the following code
void CodeGenFunction::EmitOMPParallelGenericLoopDirective() {
....
....

emitCommonOMPParallelDirective(*this, S, OMPD_for, CodeGen, emitEmptyBoundParameters);

....
....
}
Here, the GenericLoopDirective is being changed to OMPD_for because of the Parallel Directive. Similarly,

void CodeGenFunction::EmitOMPTeamsGenericLoopDirective() {
....
....

CGF.CGM.getOpenMPRuntime().emitInlinedDirective(CGF, OMPD_distribute, CodeGenDistribute);

....
....
}
Here, GenericLoopDirective is being mapped to OMPD_distribute because of the Teams construct. These changes are very similar to mine. I still do not understand why you are saying "This is a very new functionality". If you can at least provide me some directions on which are all the "enclosed in different regions" it will be helpful.

[koops]

1. Addition of extra test case loop_bind_enclosed.cpp.
2. Changing of name from Map1D to MappedDirective

[koops]

This comment was removed by koops.

[ABataev]

clang/include/clang/Sema/Sema.h

11083 ↗

(On Diff #527127)

checkLastPrivateForMappedDirectives

clang/lib/AST/StmtPrinter.cpp

745–746 ↗

(On Diff #527127)

Just MappedDirective == llvm::omp::OMPD_loop

746–748 ↗

(On Diff #527127)

Remove braces

765–766 ↗

(On Diff #527127)

Just MappedDirective == llvm::omp::OMPD_loop

766–768 ↗

(On Diff #527127)

Remove braces

1005–1006 ↗

(On Diff #527127)

Just MappedDirective == llvm::omp::OMPD_loop

1006–1008 ↗

(On Diff #527127)

Remove braces

clang/lib/Sema/SemaOpenMP.cpp

343 ↗

(On Diff #527127)

I assume it shall be a member of SharingMapType, otherwise there might be problems with handling of enclosed constructs.

646 ↗

(On Diff #527127)

Why need a cast here?

647–648 ↗

(On Diff #527127)

use nullptr or just assert(Top &&

6140–6201 ↗

(On Diff #527127)

Could you outline this new code as a separate function?

10419 ↗

(On Diff #527127)

tmp is bad name. Give better name and follow LLVM coding style recommendations.

10461 ↗

(On Diff #527127)

Same

10474 ↗

(On Diff #527127)

Use c++ style of comments

10475–10476 ↗

(On Diff #527127)

if (DSAStack->getMappedDirective() == OMPD_loop && checkGenericLoopLastprivate(*this, Clauses, OMPD_loop, DSAStack)). Avoid structured complexity.

13993 ↗

(On Diff #527127)

Same, c++

13994–13995 ↗

(On Diff #527127)

Same, avoid complexity

13994–13998 ↗

(On Diff #527127)

`return DSAStack->getMappedDirective() == OMPD_loop && checkGenericLoopLastprivate(*this, Clauses, OMPD_loop, DSAStack);

[koops]

1. Taken care of Alexy's comments.
2. Reverting back changes to StmtPrinter.cpp because state of AST needs to be shown as is to the developer of clang when -ast-print/-ast-dump is used. This also meant some changes to the test cases to reflect the Mapped directives (from "omp loop" to "omp for", "omp distribute" or "omp simd").

[jdoerfert]

@ABataev wdyt?

[ddpagan]

clang/include/clang/Basic/DiagnosticSemaKinds.td

9823 ↗

(On Diff #529532)

Should be "allowed".

clang/lib/Sema/SemaOpenMP.cpp

6127 ↗

(On Diff #529532)

This function only applies to OMPD_loop. If *Kind != OMPD_loop then should just return UseClauseWithoutBind.

clang/test/OpenMP/generic_loop_ast_print.cpp

26 ↗

(On Diff #529532)

The AST that is printed should be what was originally specified in the source.

clang/test/OpenMP/loop_bind_codegen.cpp

1 ↗

(On Diff #529532)

Should auto-generate the the IR checking for this test using llvm/utils/update_cc_test_checks.py after verifying that what's generated is correct.

[ABataev]

clang/include/clang/Sema/Sema.h

11094 ↗

(On Diff #529532)

const member function?
Add a comment

11095 ↗

(On Diff #529532)

Add a comment

11095 ↗

(On Diff #529532)

SmallVectorImpl<OMPClause *> & instead of pointer to fully specified SmallVector

[koops]

clang/test/OpenMP/generic_loop_ast_print.cpp

26 ↗

(On Diff #529532)

Since I am mapping the loop directive to new directives (in the Sema itself) I am trying to print the state of AST as is to the developer of clang when -ast-print/-ast-dump is used. Without this correct printing of the state of the AST (new directives)) the developer may wonder why the code is not behaving as he/she specified and also in which phase the code gets changed.

[koops]

clang/include/clang/Sema/Sema.h

11094 ↗

(On Diff #529532)

"checkLastPrivateForMappedDirectives" is calling a non-const function "checkGenericLoopLastprivate". Since, this is a pre-existing function and is being called from other locations, I cannot make "checkLastPrivateForMappedDirectives" a const.

[koops]

1. Taking care of Alexy & David suggestions: a) Using update_cc_test_checks.py to generate CHECK statements. b) Change in error message from "handled" to "allowed". c) Adding comments for the bind clause. d) Mangled names of the functions in the CHECK statements are more generic with regular expressions.

[ABataev]

clang/include/clang/Basic/DiagnosticSemaKinds.td

9838 ↗

(On Diff #539017)

'reduction'

clang/lib/Sema/SemaOpenMP.cpp

6105 ↗

(On Diff #539017)

Why ClausesWithoutBind passing by pointer? Use llvm::SmallVectorImpl<OMPClause *> &ClausesWithoutBind instead

6107 ↗

(On Diff #539017)

Same, pass by reference

14088–14089 ↗

(On Diff #539017)

Restore formatting

[koops]

Addressing Alexey's comments.
The name matching in loop_bind_messages.cpp & generic_loop_codegen.cpp tests changed to take care of the failures.

[ABataev]

clang/include/clang/Basic/DiagnosticSemaKinds.td

9837 ↗

(On Diff #540373)

Also, 'loop'

clang/include/clang/Sema/Sema.h

11136 ↗

(On Diff #540373)

llvm::SmallVectorImpl<OMPClause *> &ClausesWithoutBind

[koops]

1. Removing the size for the template constructor, llvm::SmallVectorImpl<OMPClause *> &ClausesWithoutBind.
2. In the diagnostic message of err_omp_bind_required_on_loop, using single quotes for the name 'loop' construct.
3. Minor changes to the tests clang/test/OpenMP/loop_bind_codegen.cpp & clang/test/OpenMP/loop_bind_enclosed.cpp to use a generic pattern matching expression for the names & reducing some checks.

[ABataev]

clang/lib/Sema/SemaOpenMP.cpp

343 ↗

(On Diff #527127)

Not done

[koops]

Moving the variable MappedDirective into structure SharingMapTy. Enough comments have been put in place to explain.
Modified some of the loop_bind tests to cater to windows and debian platform.

[ABataev]

clang/include/clang/AST/StmtOpenMP.h

291

OpenMPDirectiveKind PrevMappedDirective = llvm::omp::OMPD_unknown;

299

Remove initialization of prevMappedDirective with the default value from the constructor

606

const member function

clang/include/clang/Sema/Sema.h

11476 ↗

(On Diff #544428)

All these changed functions are called with the default value of this new prevMappedDirective paramer. Remove.

clang/lib/Sema/SemaOpenMP.cpp

6391 ↗

(On Diff #544428)

Can you just use DSAStack->getMappedDirective() instead of prevMappedDirective?

10544 ↗

(On Diff #544428)

auto *ForDirective

[koops]

Fixing Alexey's suggestions.

[koops]

Correction for clang-format errors.

[ABataev]

clang/lib/Sema/TreeTransform.h

1652 ↗

(On Diff #544688)

PrevMappedDirective

[koops]

1. Converting prevMapped to PrevMappedDirective
2. Reducing the CHECK patterns in loop_bind_enclosed.cpp to the most essential ones to succeed matching for the IR generation on windows.

[koops]

1. clang-format error in clang/include/clang/AST/StmtOpenMP.h corrected.
2. Correcting clang/test/OpenMP/loop_bind_enclosed.cpp to use regular expressions for names of functions to accomodate names generated on Microsoft Windows.

[ABataev]

clang/include/clang/AST/StmtOpenMP.h

291

I don't see where this field is stored/loaded for PCH support. You need add a support for this in Serialization.

602–605

Better to make it a part of Create member function rather than having a separate function for this.

[koops]

1. clang/test/OpenMP/loop_bind_enclosed.cpp : Converting main and function names within it to use regular expressions to accomodate windows platform.

i) Instead of calling setMappedDirective() after the the creation of the Directive, made the MappedDirective a parameter of the Create() method.
ii) setMappedDirective() is called from the Directive::Create()
iii) The inheritance is : OMPExecutableDirective --> OMPLoopBasedDirective --> OMPLoopDirective --> OMPSimdDirective, from OMPSimdDirective constructor it involves a lot of changes to initialize PrevMappedDirective variable in OMPExecutableDirective starting from OMPSimdDirective, even if the variable is "protected". Hence we call setMappedDirective() from OMPSimdDirective::Create().

[koops]

clang/include/clang/AST/StmtOpenMP.h

291

Since the ASTReader and ASTWriter are used by the developer of clang, it is preferable to show the internal state of the compiler as is, e.g. if "#pragma omp loop bind(parallel)" is changed to "#pragma omp for" then the ASTWriter has to show it as "#pragma omp for". I can change it to "#pragma omp loop bind(parallel)" using the MappedDirective that is stored however, that would be misleading to the developer of clang.

[ABataev]

clang/include/clang/AST/StmtOpenMP.h

291

How it is related to PCH? Try to serialize/deserialize the template and then try to create a new instance. There will be an issue since there is no serialization support, when you try to create new instance in the Rebuild... function (mapLoopConstruct will use OMPD_unknown as Prev value instead of oringal previous value, since it is not serialized/deserialized)

[koops]

1. Support for PCH Serilization/deserilization for the PrevMapLoopConstruct variable in OMPExecutableDirective.
2. Modification of clang/test/PCH/pragma-loop.cpp to include "#pragma omp loop bind" with different parameters for bind clause.

[ABataev]

LG

[koops]

In clang/test/OpenMP/loop_bind_enclosed.cpp Generalizing the CHECK pattern for aarch64-linux, s390x-linux and ppc64le-linux.

[koops]

Making the CHECK pattern generic to match platforms tested after committing changes.

[koops]

Can someone please check for MacOS? Yesterday when this support was committed, the CHECK statements in tests loop_bind_codegen.cpp and loop_bind_enclosed.cpp had failed on Mac and I received a comment as follows:

When relanding, please remember to put a link to the review in the commit message.

I have made the CHECK statements generic enough to match those on Mac with CHECK for the functions that the test were expected.

[thakis]

In D144634#4573483, @koops wrote:

Can someone please check for MacOS? Yesterday when this support was committed, the CHECK statements in tests loop_bind_codegen.cpp and loop_bind_enclosed.cpp had failed on Mac

I patched this in and verified that check-clang now passes on macOS. Thanks!

and I received a comment as follows:

When relanding, please remember to put a link to the review in the commit message.

To do this, run git commit --amend and put Differential Revision: https://reviews.llvm.org/D144634 at the bottom of the commit message. git log has many, many examples of this :)

clang/include/clang/AST/StmtOpenMP.h

 
 protected:
   /// Data, associated with the directive.
   OMPChildren *Data = nullptr;
 
   /// Build instance of directive of class \a K.
   ///
   /// \param SC Statement class.
   /// \param K Kind of OpenMP directive.
   /// \param StartLoc Starting location of the directive (directive keyword).
   /// \param EndLoc Ending location of the directive.
   ///
   OMPExecutableDirective(StmtClass SC, OpenMPDirectiveKind K,
                          SourceLocation StartLoc, SourceLocation EndLoc)
       : Stmt(SC), Kind(K), StartLoc(std::move(StartLoc)),
         EndLoc(std::move(EndLoc)) {}
 
   template <typename T, typename... Params>
   static T *createDirective(const ASTContext &C, ArrayRef<OMPClause *> Clauses,
                             Stmt *AssociatedStmt, unsigned NumChildren,
                             Params &&... P) {
     void *Mem =
         C.Allocate(sizeof(T) + OMPChildren::size(Clauses.size(), AssociatedStmt,
                                                  NumChildren),
                    alignof(T));
const Stmt *getRawStmt() const {
     return const_cast<OMPExecutableDirective *>(this)->getRawStmt();
   }
   Stmt *getRawStmt() {
     assert(hasAssociatedStmt() &&
            "Expected directive with the associated statement.");
     return Data->getRawStmt();
   }
 };
 
 /// This represents '#pragma omp parallel' directive.
 ///
 /// \code
 /// #pragma omp parallel private(a,b) reduction(+: c,d)
 /// \endcode
 /// In this example directive '#pragma omp parallel' has clauses 'private'
 /// with the variables 'a' and 'b' and 'reduction' with operator '+' and
 /// variables 'c' and 'd'.
 ///
 class OMPParallelDirective : public OMPExecutableDirective {
   friend class ASTStmtReader;
   friend class OMPExecutableDirective;
public:
   /// \param AssociatedStmt Statement, associated with the directive.
   /// \param Exprs Helper expressions for CodeGen.
   ///
   static OMPGenericLoopDirective *
   Create(const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
          unsigned CollapsedNum, ArrayRef<OMPClause *> Clauses,
          Stmt *AssociatedStmt, const HelperExprs &Exprs);
 
+  struct LoopDirCrParam {
+    const ASTContext *C;
+    SourceLocation StartLoc;
+    SourceLocation EndLoc;
+    unsigned CollapsedNum;
+    ArrayRef<OMPClause *> Clauses;
+    Stmt *AssociatedStmt;
+    const HelperExprs Exprs;
+
+    LoopDirCrParam(const ASTContext &C, SourceLocation StartLoc,
+                   SourceLocation EndLoc, unsigned CollapsedNum,
+                   ArrayRef<OMPClause *> Clauses, Stmt *AssociatedStmt,
+                   const HelperExprs &Exprs)
+        : C(&C), StartLoc(StartLoc), EndLoc(EndLoc), CollapsedNum(CollapsedNum),
+          Clauses(Clauses), AssociatedStmt(AssociatedStmt), Exprs(Exprs) {}
+  };
+
+  void LoopParamInit(const ASTContext &C, SourceLocation StartLoc,
+                     SourceLocation EndLoc, unsigned CollapsedNum,
+                     ArrayRef<OMPClause *> Clauses, Stmt *AssociatedStmt,
+                     const HelperExprs &Exprs);
+
+  struct LoopDirCrParam *LoopDirCrParmV;
+
   /// Creates an empty directive with a place for \a NumClauses clauses.
   ///
   /// \param C AST context.
   /// \param NumClauses Number of clauses.
   /// \param CollapsedNum Number of collapsed nested loops.
   ///
   static OMPGenericLoopDirective *CreateEmpty(const ASTContext &C,
                                               unsigned NumClauses,

clang/lib/AST/StmtOpenMP.cpp

OMPGenericLoopDirective *OMPGenericLoopDirective::Create(
   Dir->setPrivateCounters(Exprs.PrivateCounters);
   Dir->setInits(Exprs.Inits);
   Dir->setUpdates(Exprs.Updates);
   Dir->setFinals(Exprs.Finals);
   Dir->setDependentCounters(Exprs.DependentCounters);
   Dir->setDependentInits(Exprs.DependentInits);
   Dir->setFinalsConditions(Exprs.FinalsConditions);
   Dir->setPreInits(Exprs.PreInits);
+
+  Dir->LoopParamInit(C, StartLoc, EndLoc, CollapsedNum, Clauses, AssociatedStmt,
+                     Exprs);
+
   return Dir;
 }
 
 OMPGenericLoopDirective *
 OMPGenericLoopDirective::CreateEmpty(const ASTContext &C, unsigned NumClauses,
                                      unsigned CollapsedNum, EmptyShell) {
   return createEmptyDirective<OMPGenericLoopDirective>(
       C, NumClauses, /*HasAssociatedStmt=*/true,
       numLoopChildren(CollapsedNum, OMPD_loop), CollapsedNum);
 }
 
+void OMPGenericLoopDirective::LoopParamInit(
+    const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
+    unsigned CollapsedNum, ArrayRef<OMPClause *> Clauses, Stmt *AssociatedStmt,
+    const HelperExprs &Exprs) {
+  this->LoopDirCrParmV = new LoopDirCrParam(C, StartLoc, EndLoc, CollapsedNum,
+                                            Clauses, AssociatedStmt, Exprs);
+}
+
 OMPTeamsGenericLoopDirective *OMPTeamsGenericLoopDirective::Create(
     const ASTContext &C, SourceLocation StartLoc, SourceLocation EndLoc,
     unsigned CollapsedNum, ArrayRef<OMPClause *> Clauses, Stmt *AssociatedStmt,
     const HelperExprs &Exprs) {
   auto *Dir = createDirective<OMPTeamsGenericLoopDirective>(
       C, Clauses, AssociatedStmt,
       numLoopChildren(CollapsedNum, OMPD_teams_loop), StartLoc, EndLoc,
       CollapsedNum);

clang/lib/CodeGen/CGStmtOpenMP.cpp

if (ForS && !isa<DeclStmt>(ForS->getInit())) {
       CGF.EmitOMPPrivateLoopCounters(S, LoopScope);
       (void)LoopScope.Privatize();
       CGF.EmitStmt(CS);
       LoopScope.restoreMap();
     } else {
       CGF.EmitStmt(CS);
     }
   };
+  /* Bind */
+  SmallVector<OMPClause *, 6> newClauses;
+  for (auto clausePtr : S.clauses()) {
+    if (clausePtr->getClauseKind() != llvm::omp::Clause::OMPC_bind) {
+      newClauses.push_back(clausePtr);
+    }
+  }
+  ArrayRef<clang::OMPClause *> newAClauses(newClauses);
+
+  for (const auto *C : S.getClausesOfKind<OMPBindClause>()) {
+    OpenMPBindClauseKind bindParam = C->getBindKind();
+    switch (bindParam) {
+    case OMPC_BIND_parallel: {
+      OMPForDirective *ompForD = OMPForDirective::Create(
+          *(S.LoopDirCrParmV->C), S.LoopDirCrParmV->StartLoc,
+          S.LoopDirCrParmV->EndLoc, S.LoopDirCrParmV->CollapsedNum, newAClauses,
+          S.LoopDirCrParmV->AssociatedStmt, S.LoopDirCrParmV->Exprs, NULL, 0);
+      EmitOMPForDirective(*ompForD);
+      return;
+      break;
+    }
+    case OMPC_BIND_teams: {
+      OMPDistributeDirective *ompDistD = OMPDistributeDirective::Create(
+          *(S.LoopDirCrParmV->C), S.LoopDirCrParmV->StartLoc,
+          S.LoopDirCrParmV->EndLoc, S.LoopDirCrParmV->CollapsedNum, newAClauses,
+          S.LoopDirCrParmV->AssociatedStmt, S.LoopDirCrParmV->Exprs);
+      EmitOMPDistributeDirective(*ompDistD);
+      return;
+      break;
+    }
+    case OMPC_BIND_thread: {
+      /*
+      OMPSimdDirective *ompSimdD = OMPSimdDirective::Create(
+          *(S.LoopDirCrParmV->C), S.LoopDirCrParmV->StartLoc,
+          S.LoopDirCrParmV->EndLoc, S.LoopDirCrParmV->CollapsedNum,
+          newAClauses, S.LoopDirCrParmV->AssociatedStmt,
+          S.LoopDirCrParmV->Exprs);
+      EmitOMPSimdDirective(*ompSimdD);
+      return;
+      */
+      break;
+    }
+    case OMPC_BIND_unknown:
+      return;
+    }
+  }
   OMPLexicalScope Scope(*this, S, OMPD_unknown);
   CGM.getOpenMPRuntime().emitInlinedDirective(*this, OMPD_loop, CodeGen);
 }
 
 void CodeGenFunction::EmitSimpleOMPExecutableDirective(
     const OMPExecutableDirective &D) {
   if (const auto *SD = dyn_cast<OMPScanDirective>(&D)) {
     EmitOMPScanDirective(*SD);