[OPENMP] Initial codegen for '#pragma omp parallel'

[hfinkel]

lib/CodeGen/CGOpenMPRuntime.cpp

29

Please document the meaning of this string.

lib/CodeGen/CGOpenMPRuntime.h

46

This should have some overall comment explaining what this is the location of.

49

c-style?

73

This needs to be documented (including the fields too).

81

What's this?

lib/Sema/SemaOpenMP.cpp

675

Could these names conflict with user-provided names? (Should they be in the implementation namespace?)

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

29

Ok.

lib/CodeGen/CGOpenMPRuntime.h

46

Ok, I'll add comment

49

Yes, all these comments are taken from http://llvm.org/svn/llvm-project/openmp/trunk/runtime/src/kmp.h

73

Ok, will be commented.

81

Default position used for initialization of all ident_t objects. I'll described it.

lib/Sema/SemaOpenMP.cpp

675

Yes, that's possible. I'll rename them.

[ABataev]

Fix after Hal's review

[hfinkel]

• Original Message -----

From: "Alexey Bataev" <[e-mail]>
To: [e-mail], [e-mail], [e-mail], [e-mail], [e-mail],
[e-mail], "a bataev" <[e-mail]>
Cc: [e-mail]
Sent: Friday, March 7, 2014 5:56:04 AM
Subject: Re: [PATCH] [OPENMP] Initial codegen for '#pragma omp parallel'

Comment at: lib/CodeGen/CGOpenMPRuntime.h:45
@@ +44,3 @@
+public:
+ enum OpenMPLocationFlags {

+ /// \brief Use tramponline for internal microtask.

[e-mail] wrote:

This should have some overall comment explaining what this is the
location of.

Ok, I'll add comment

Comment at: lib/CodeGen/CGOpenMPRuntime.h:48
@@ +47,3 @@
+ OMP_IDENT_IMD = 0x01,
+ /// \brief Use c-stile ident structure.

+ OMP_IDENT_KMPC = 0x02,

[e-mail] wrote:

c-style?

Yes, all these comments are taken from
http://llvm.org/svn/llvm-project/openmp/trunk/runtime/src/kmp.h

I was wondering whether 'c-stile' was a typo, and should have been c-style?

-Hal

Comment at: lib/CodeGen/CGOpenMPRuntime.h:72
@@ +71,3 @@
+private:
+ enum Ident_TFields {

+ Ident_T_Reserved_1,

[e-mail] wrote:

This needs to be documented (including the fields too).

Ok, will be commented.

Comment at: lib/CodeGen/CGOpenMPRuntime.h:80
@@ +79,3 @@
+ CodeGenModule &CGM;
+ llvm::Constant *DefaultOpenMPPSource;

+ llvm::StructType *Ident_TTy;

[e-mail] wrote:

What's this?

Default position used for initialization of all ident_t objects. I'll
described it.

Comment at: lib/CodeGen/CGOpenMPRuntime.cpp:28
@@ +27,3 @@
+CGOpenMPRuntime::CGOpenMPRuntime(CodeGenModule &CGM) : CGM(CGM) {
+ DefaultOpenMPPSource =
CGM.GetAddrOfConstantCString(";unknown;unknown;0;0;;");

+ DefaultOpenMPPSource =

[e-mail] wrote:

Please document the meaning of this string.

Ok.

Comment at: lib/Sema/SemaOpenMP.cpp:675
@@ +674,3 @@
+ { std::make_pair("global_tid", KmpInt32PtrTy),
+ std::make_pair("bound_tid", KmpInt32PtrTy),
+ std::make_pair(StringRef(), QualType()) // __context with

shared vars

[e-mail] wrote:

Could these names conflict with user-provided names? (Should they
be in the implementation namespace?)

Yes, that's possible. I'll rename them.

http://llvm-reviews.chandlerc.com/D2883

[ABataev]

Oh, yes, of course "style". I'll fix that

[ABataev]

Fixed typo in comment

[ABataev]

Ping

[hfinkel]

I think that this generally looks good. I still have a concern about the names: we generate names like context. Now I understand that a user should not generate names like this, but what if they do? Does it crash? We need to make sure that have reasonable error detection and reporting for these things. In the clang-omp branch the names are generated with a preceding ''.' -- ".kmpc_global_thread_num." for example -- and this can never conflict with a user-provided identifier. This seems safer, so why are you not doing that here?

[ABataev]

Update after latest Hal's review

[rsmith]

I agree that this is looking really good.

lib/CodeGen/CGOpenMPRuntime.cpp

54–69

Looks like this can be replaced by IRBuilder::InsertPointGuard?

175

We put the break; inside the braces.

lib/CodeGen/CGOpenMPRuntime.h

50–51

Typo of 'trampoline'.

50–65

Our coding convention doesn't use ALL_CAPS here. These should probably be OLF_IdentImd, OLF_IdentKmpc, ... to follow that convention (though I'm not really sure what these names mean, and they don't seem to directly correspond to the comments. What does IMD mean, or KMPC, for instance?)

Alternatively, you could explicitly state that these are named to match the names from kmp.h, rather than just saying the descriptions are from there.

lib/CodeGen/CGStmtOpenMP.cpp

32–36

Why are you creating the CGCapturedStmtInfo on the heap? Seems weird, since it looks like you delete it two lines afterwards.

lib/Sema/SemaOpenMP.cpp

673–680

Formatting here doesn't match our usual conventions. For the braced initialization, put the { on the previous line, dedent the };. Dedent the } for the case, and move the break; before it.

685–686

Yuck. NUM_OPENMP_DIRECTIVES shouldn't be a member of the enum, to remove the need to do this.

lib/Sema/TreeTransform.h

9784

Capital I please.

lib/Serialization/ASTReaderDecl.cpp

1080–1088

Do you have test coverage for this?

[ABataev]

Richard, thanks for the review!!! I do appreciate it.

lib/CodeGen/CGOpenMPRuntime.cpp

54–69

Yes, thanks, missed this one!

175

Thanks, will be fixed.

lib/CodeGen/CGOpenMPRuntime.h

50–51

Will be fixed

50–65

Richard, ok, I'll add required comment

lib/CodeGen/CGStmtOpenMP.cpp

32–36

I just followed the standard pattern for CapturedStmt (see CGStmt.cpp, CodeGenFunction::EmitCapturedStmt). But I'll rework it.

lib/Sema/SemaOpenMP.cpp

673–680

Ok. I'll reformat it.

685–686

This one is needed for checking that the directive is correctly specified.

lib/Sema/TreeTransform.h

9784

Ok

lib/Serialization/ASTReaderDecl.cpp

1080–1088

Yes, parallel_codegen.cpp and parallel_ast_print.cpp.

[ABataev]

Rework after Richard's review

[rsmith]

LGTM, but this is a significant enough change to CodeGen that I'd like rjmccall to look at it too.

lib/Serialization/ASTReaderDecl.cpp

1080–1088

Do those really test the serialization/deserialization code?

[ABataev]

lib/Serialization/ASTReaderDecl.cpp

1080–1088

Yes, take a look at these tests. There are 3 passes: 1) initial test "as is"; 2) "-emit-pch" pass; 3) "-include-pch" pass. The 2-nd and the 3-rd passes check serialization/deserialization of AST nodes.

[ABataev]

Ping!

[zinovy.nis]

lib/CodeGen/CodeGenModule.h

258

Probably not a comment for your line, but it'd be safer to use std::unique_ptr instead of raw pointers. That would also eliminate explicit calls to delete.

[rjmccall]

include/clang/Sema/Sema.h

2984

This typedef has a very generic-sounding name and a very specific purpose.

lib/CodeGen/CGOpenMPRuntime.cpp

32

Please create this string lazily.

48

Please create this global lazily.

49

Please /*comment*/ arguments like "true".

54

I would not expect this method to cause code to be emitted at the current insertion point. Maybe call it emitUpdateLocation or something like that?

63

Remove this entire thing and use an llvm::Twine inline in the argument to CreateTempAlloca.

92

This check is always true; you bailed out earlier if it wasn't.

126

I'm not sure why you're using an alloca if you're going to unconditionally do the call in a place that dominates the entire function. Just do the call and use the result wherever you need it.

lib/CodeGen/CGOpenMPRuntime.h

79

The _T suffix makes this pretty confusing to read — it's ident_t in the original, presumably, and capitalizing it does it no favors.

I would rename:

Ident_TTy -> IdentTy
Ident_TFields -> IdentFieldIndex
Ident_T_* -> IdentField_*

Also, put this enum and its comment next to the declaration of IdentTy, and include a C-style struct definition in the comment.

98

Please add a comment to the end of the line describing this function type in C terms, if reasonable.

101

Does this need to be as specific as an AllocaInst?

Also, these maps seem to contain information that's really specific to a CodeGenFunction; you should get notified when a CGF is torn down so that you can remove entries as appropriate.

lib/CodeGen/CGStmtOpenMP.cpp

2

There's no reason to make this a separate file. Just keep all the OpenMP code in one place as much as possible.

33

You should declare CGF inside this block, too.

43

Please comment this magic number.

lib/Sema/SemaOpenMP.cpp

685

Please define your enums so that things like this case aren't necessary. If you really need a NUM_OPENMP_DIRECTIVES constant, you can define it right after the enum.

745

This comment is useless. Please instead cite *why* you can mark the captured decl as no throw.

[ABataev]

John,
Thanks for review! See my comments. Updated patch will be prepared soon.

Best regards,

Alexey Bataev

Software Engineer
Intel Compiler Team
Intel Corp.

include/clang/Sema/Sema.h

2984

Ok, I'll rename it.

lib/CodeGen/CGOpenMPRuntime.cpp

32

Ok

48

Ok

49

Ok

54

Ok, I'll rename it.

63

Ok

92

Agree, I'll remove it

126

Agree

lib/CodeGen/CGOpenMPRuntime.h

79

Ok

98

Ok

101

I'll try to rework it.

lib/CodeGen/CGStmtOpenMP.cpp

2

I think it would be better to keep a separate file because there should be a lot of code for OpenMP.

33

Agree

43

Ok

lib/Sema/SemaOpenMP.cpp

685

Ok

745

Will be removed

[ABataev]

Update after John McCall review

[ABataev]

Ping!

[ABataev]

Ping

[ABataev]

LGTM (C. Bergström, Mon, 5 May 2014 23:48:34)

[ABataev]

Committed revision 208077.

include/clang/AST/Decl.h

 };
 
 /// \brief This represents the body of a CapturedStmt, and serves as its
 /// DeclContext.
 class CapturedDecl : public Decl, public DeclContext {
 private:
   /// \brief The number of parameters to the outlined function.
   unsigned NumParams;
+  /// \brief The position of context parameter in list of parameters.
+  unsigned ContextParam;
   /// \brief The body of the outlined function.
-  Stmt *Body;
+  llvm::PointerIntPair<Stmt *, 1, bool> BodyAndNothrow;
 
   explicit CapturedDecl(DeclContext *DC, unsigned NumParams)
     : Decl(Captured, DC, SourceLocation()), DeclContext(Captured),
-      NumParams(NumParams), Body(0) { }
+      NumParams(NumParams), ContextParam(0), BodyAndNothrow(0, false) { }
 
   ImplicitParamDecl **getParams() const {
     return reinterpret_cast<ImplicitParamDecl **>(
              const_cast<CapturedDecl *>(this) + 1);
   }
 
 public:
-  static CapturedDecl *Create(ASTContext &C, DeclContext *DC, unsigned NumParams);
+  static CapturedDecl *Create(ASTContext &C, DeclContext *DC,
+                              unsigned NumParams);
   static CapturedDecl *CreateDeserialized(ASTContext &C, unsigned ID,
                                           unsigned NumParams);
 
-  Stmt *getBody() const { return Body; }
-  void setBody(Stmt *B) { Body = B; }
+  Stmt *getBody() const { return BodyAndNothrow.getPointer(); }
+  void setBody(Stmt *B) { BodyAndNothrow.setPointer(B); }
+
+  bool isNothrow() const { return BodyAndNothrow.getInt(); }
+  void setNothrow(bool Nothrow = true) { BodyAndNothrow.setInt(Nothrow); }
 
   unsigned getNumParams() const { return NumParams; }
 
   ImplicitParamDecl *getParam(unsigned i) const {
     assert(i < NumParams);
     return getParams()[i];
   }
   void setParam(unsigned i, ImplicitParamDecl *P) {
     assert(i < NumParams);
     getParams()[i] = P;
   }
 
   /// \brief Retrieve the parameter containing captured variables.
-  ImplicitParamDecl *getContextParam() const { return getParam(0); }
-  void setContextParam(ImplicitParamDecl *P) { setParam(0, P); }
+  ImplicitParamDecl *getContextParam() const {
+    assert(ContextParam < NumParams);
+    return getParam(ContextParam);
+  }
+  void setContextParam(unsigned i, ImplicitParamDecl *P) {
+    assert(i < NumParams);
+    ContextParam = i;
+    setParam(i, P);
+  }
+  unsigned getContextParamPosition() const { return ContextParam; }
 
   typedef ImplicitParamDecl **param_iterator;
   /// \brief Retrieve an iterator pointing to the first parameter decl.
   param_iterator param_begin() const { return getParams(); }
   /// \brief Retrieve an iterator one past the last parameter decl.
   param_iterator param_end() const { return getParams() + NumParams; }
 
   // Implement isa/cast/dyncast/etc.

include/clang/Sema/Sema.h

public:
   StmtResult ActOnIndirectGotoStmt(SourceLocation GotoLoc,
                                    SourceLocation StarLoc,
                                    Expr *DestExp);
   StmtResult ActOnContinueStmt(SourceLocation ContinueLoc, Scope *CurScope);
   StmtResult ActOnBreakStmt(SourceLocation BreakLoc, Scope *CurScope);
 
   void ActOnCapturedRegionStart(SourceLocation Loc, Scope *CurScope,
                                 CapturedRegionKind Kind, unsigned NumParams);
+  typedef std::pair<StringRef, QualType> ParamNameType;
+  void ActOnCapturedRegionStart(SourceLocation Loc, Scope *CurScope,
+                                CapturedRegionKind Kind,
+                                ArrayRef<ParamNameType> Params);
   StmtResult ActOnCapturedRegionEnd(Stmt *S);
   void ActOnCapturedRegionError();
   RecordDecl *CreateCapturedStmtRecordDecl(CapturedDecl *&CD,
                                            SourceLocation Loc,
                                            unsigned NumParams);
   const VarDecl *getCopyElisionCandidate(QualType ReturnType, Expr *E,
                                          bool AllowFunctionParameters);
 
public:
   DeclGroupPtrTy ActOnOpenMPThreadprivateDirective(
                                      SourceLocation Loc,
                                      ArrayRef<Expr *> VarList);
   // \brief Builds a new OpenMPThreadPrivateDecl and checks its correctness.
   OMPThreadPrivateDecl *CheckOMPThreadPrivateDecl(
                                      SourceLocation Loc,
                                      ArrayRef<Expr *> VarList);
 
+  // brief Initialization of captured region for OpenMP parallel region.
+  void ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, SourceLocation Loc,
+                              Scope *CurScope);
   StmtResult ActOnOpenMPExecutableDirective(OpenMPDirectiveKind Kind,
                                             ArrayRef<OMPClause *> Clauses,
                                             Stmt *AStmt,
                                             SourceLocation StartLoc,
                                             SourceLocation EndLoc);
   /// \brief Called on well-formed '\#pragma omp parallel' after parsing
   /// of the  associated statement.
   StmtResult ActOnOpenMPParallelDirective(ArrayRef<OMPClause *> Clauses,

lib/CodeGen/CGException.cpp

if (KeepInsertionPoint)
     EmitBlock(createBasicBlock("throw.cont"));
 }
 
 void CodeGenFunction::EmitStartEHSpec(const Decl *D) {
   if (!CGM.getLangOpts().CXXExceptions)
     return;
   
   const FunctionDecl* FD = dyn_cast_or_null<FunctionDecl>(D);
-  if (FD == 0)
+  if (FD == 0) {
+    // Check if CapturedDecl is nothrow and create terminate scope for it.
+    if (const CapturedDecl* CD = dyn_cast_or_null<CapturedDecl>(D)) {
+      if (CD->isNothrow())
+        EHStack.pushTerminate();
+    }
     return;
+  }
   const FunctionProtoType *Proto = FD->getType()->getAs<FunctionProtoType>();
   if (Proto == 0)
     return;
 
   ExceptionSpecificationType EST = Proto->getExceptionSpecType();
   if (isNoexceptExceptionSpec(EST)) {
     if (Proto->getNoexceptSpec(getContext()) == FunctionProtoType::NR_Nothrow) {
       // noexcept functions are simple terminate scopes.
static void emitFilterDispatchBlock(CodeGenFunction &CGF,
   CGF.Builder.CreateUnreachable();
 }
 
 void CodeGenFunction::EmitEndEHSpec(const Decl *D) {
   if (!CGM.getLangOpts().CXXExceptions)
     return;
   
   const FunctionDecl* FD = dyn_cast_or_null<FunctionDecl>(D);
-  if (FD == 0)
+  if (FD == 0) {
+    // Check if CapturedDecl is nothrow and pop terminate scope for it.
+    if (const CapturedDecl* CD = dyn_cast_or_null<CapturedDecl>(D)) {
+      if (CD->isNothrow())
+        EHStack.popTerminate();
+    }
     return;
+  }
   const FunctionProtoType *Proto = FD->getType()->getAs<FunctionProtoType>();
   if (Proto == 0)
     return;
 
   ExceptionSpecificationType EST = Proto->getExceptionSpecType();
   if (isNoexceptExceptionSpec(EST)) {
     if (Proto->getNoexceptSpec(getContext()) == FunctionProtoType::NR_Nothrow) {
       EHStack.popTerminate();

lib/CodeGen/CGOpenMPRuntime.h

+//===----- CGOpenMPRuntime.h - Interface to OpenMP Runtimes -----*- C++ -*-===//
+//
+//                     The LLVM Compiler Infrastructure
+//
+// This file is distributed under the University of Illinois Open Source
+// License. See LICENSE.TXT for details.
+//
+//===----------------------------------------------------------------------===//
+//
+// This provides a class for OpenMP runtime code generation.
+//
+//===----------------------------------------------------------------------===//
+
+#ifndef CLANG_CODEGEN_OPENMPRUNTIME_H
+#define CLANG_CODEGEN_OPENMPRUNTIME_H
+
+#include "clang/AST/Type.h"
+#include "llvm/ADT/DenseMap.h"
+#include "llvm/IR/Type.h"
+#include "llvm/IR/Value.h"
+
+namespace llvm {
+class AllocaInst;
+class CallInst;
+class GlobalVariable;
+class Constant;
+class Function;
+class Module;
+class StructLayout;
+class FunctionType;
+class StructType;
+class Type;
+class Value;
+}
+
+namespace clang {
+
+namespace CodeGen {
+
+class CodeGenFunction;
+class CodeGenModule;
+
+class CGOpenMPRuntime {
+public:
+  /// \brief Values for bit flags used in the ident_t to describe the fields.
+  /// All descriptions are taken from
+  /// http://llvm.org/svn/llvm-project/openmp/trunk/runtime/src/kmp.h
+  enum OpenMPLocationFlags {
+    /// \brief Use tramponline for internal microtask.
+    OMP_IDENT_IMD = 0x01,
+    /// \brief Use c-stile ident structure.
+    OMP_IDENT_KMPC = 0x02,
+    /// \brief Atomic reduction option for kmpc_reduce.
+    OMP_ATOMIC_REDUCE = 0x10,
+    /// \brief Explicit 'barrier' directive.
+    OMP_IDENT_BARRIER_EXPL = 0x20,
+    /// \brief Implicit barrier in code.
+    OMP_IDENT_BARRIER_IMPL = 0x40,
+    /// \brief Implicit barrier in 'for' directive.
+    OMP_IDENT_BARRIER_IMPL_FOR = 0x40,
+    /// \brief Implicit barrier in 'sections' directive.
+    OMP_IDENT_BARRIER_IMPL_SECTIONS = 0xC0,
+    /// \brief Implicit barrier in 'single' directive.
+    OMP_IDENT_BARRIER_IMPL_SINGLE = 0x140
+  };
+  enum OpenMPRTLFunction {
+    // Call to void __kmpc_fork_call(ident_t *loc, kmp_int32 argc, kmpc_micro
+    // microtask, ...);
+    OMPRTL__kmpc_fork_call,
+    // Call to kmp_int32 kmpc_global_thread_num(ident_t *loc);
+    OMPRTL__kmpc_global_thread_num
+  };
+
+private:
+  // \brief Describes ident structure that describes a source location.
+  /// All descriptions are taken from
+  /// http://llvm.org/svn/llvm-project/openmp/trunk/runtime/src/kmp.h
+  enum Ident_TFields {
+    /// \brief might be used in Fortran
+    Ident_T_Reserved_1,
+    /// \brief OMP_IDENT_xxx flags; OMP_IDENT_KMPC identifies this union member.
+    Ident_T_Flags,
+    /// \brief Not really used in Fortran any more
+    Ident_T_Reserved_2,
+    /// \brief Source[4] in Fortran, do not use for C++
+    Ident_T_Reserved_3,
+    /// \brief String describing the source location. The string is composed of
+    /// semi-colon separated fields which describe the source file, the function
+    /// and a pair of line numbers that delimit the construct.
+    Ident_T_PSource
+  };
+  CodeGenModule &CGM;
+  /// \brief Default const ident_t object used for initialization of all other
+  /// ident_t objects.
+  llvm::Constant *DefaultOpenMPPSource;
+  llvm::StructType *Ident_TTy;
+  llvm::FunctionType *Kmpc_MicroTy;
+  llvm::GlobalVariable *DefaultOpenMPLocation;
+  /// \brief Map of local debug location and functions.
+  typedef llvm::DenseMap<llvm::Function *, llvm::AllocaInst *> OpenMPLocMapTy;
+  OpenMPLocMapTy OpenMPLocMap;
+  /// \brief Map of local gtid and functions.
+  typedef llvm::DenseMap<llvm::Function *, llvm::AllocaInst *> OpenMPGtidMapTy;
+  OpenMPGtidMapTy OpenMPGtidMap;
+
+public:
+  CGOpenMPRuntime(CodeGenModule &CGM);
+  ~CGOpenMPRuntime() {}
+
+  /// \brief Generates object of ident_t type with info for source location.
+  /// \param CGF Reference to current CodeGenFunction.
+  /// \param Loc Clang source location.
+  /// \param Flags Flags for OpenMP location.
+  ///
+  llvm::Value *GetOpenMPLocation(CodeGenFunction &CGF, SourceLocation Loc,
+                                 OpenMPLocationFlags Flags = OMP_IDENT_KMPC);
+
+  /// \brief Generates global thread number value.
+  /// \param CGF Reference to current CodeGenFunction.
+  /// \param Loc Clang source location.
+  ///
+  llvm::Value *GetOpenMPGlobalThreadNum(CodeGenFunction &CGF,
+                                        SourceLocation Loc);
+
+  /// \brief Returns pointer to ident_t type;
+  llvm::Type *getIdent_TPointerTy();
+
+  /// \brief Returns pointer to kmpc_micro type;
+  llvm::Type *getKmpc_MicroPointerTy();
+
+  /// \brief Returns specified OpenMP runtime function.
+  /// \param Function OpenMP runtime function.
+  /// \return Specified function.
+  llvm::Constant *CreateRuntimeFunction(OpenMPRTLFunction Function);
+};
+}
+}
+
+#endif

lib/CodeGen/CGOpenMPRuntime.cpp

+//===----- CGOpenMPRuntime.cpp - Interface to OpenMP Runtimes -------------===//
+//
+//                     The LLVM Compiler Infrastructure
+//
+// This file is distributed under the University of Illinois Open Source
+// License. See LICENSE.TXT for details.
+//
+//===----------------------------------------------------------------------===//
+//
+// This provides a class for OpenMP runtime code generation.
+//
+//===----------------------------------------------------------------------===//
+
+#include "CGOpenMPRuntime.h"
+#include "CodeGenFunction.h"
+#include "clang/AST/Decl.h"
+#include "llvm/ADT/ArrayRef.h"
+#include "llvm/IR/DerivedTypes.h"
+#include "llvm/IR/GlobalValue.h"
+#include "llvm/IR/Value.h"
+#include "llvm/Support/raw_ostream.h"
+#include <assert.h>
+
+using namespace clang;
+using namespace CodeGen;
+
+CGOpenMPRuntime::CGOpenMPRuntime(CodeGenModule &CGM) : CGM(CGM) {
+  // Initialize default location for psource field of ident_t structure of all
+  // ident_t objects. Format is ";file;function;line;column;;". Taken from
+  // http://llvm.org/svn/llvm-project/openmp/trunk/runtime/src/kmp_str.c
+  DefaultOpenMPPSource = CGM.GetAddrOfConstantCString(";unknown;unknown;0;0;;");
+  DefaultOpenMPPSource =
+      llvm::ConstantExpr::getBitCast(DefaultOpenMPPSource, CGM.Int8PtrTy);
+  Ident_TTy = llvm::StructType::create(
+      "ident_t", CGM.Int32Ty /* reserved_1 */, CGM.Int32Ty /* flags */,
+      CGM.Int32Ty /* reserved_2 */, CGM.Int32Ty /* reserved_3 */,
+      CGM.Int8PtrTy /* psource */, NULL);
+  // Build void (*kmpc_micro)(kmp_int32 *global_tid, kmp_int32 *bound_tid,...)
+  llvm::Type *MicroParams[] = { llvm::PointerType::getUnqual(CGM.Int32Ty),
+                                llvm::PointerType::getUnqual(CGM.Int32Ty) };
+  Kmpc_MicroTy = llvm::FunctionType::get(CGM.VoidTy, MicroParams, true);
+  llvm::Constant *Zero = llvm::ConstantInt::get(CGM.Int32Ty, 0, true);
+  llvm::Constant *Values[] = { Zero, llvm::ConstantInt::get(CGM.Int32Ty,
+                                                            OMP_IDENT_KMPC),
+                               Zero, Zero, DefaultOpenMPPSource };
+  llvm::Constant *Init = llvm::ConstantStruct::get(Ident_TTy, Values);
+  DefaultOpenMPLocation = new llvm::GlobalVariable(
+      CGM.getModule(), Ident_TTy, true, llvm::GlobalValue::PrivateLinkage, Init,
+      "__kmpc_default_loc.addr");
+  DefaultOpenMPLocation->setUnnamedAddr(true);
+}
+
+namespace {
+/// \brief RAII object that save current insert position and then restores it.
+class BuilderInsertPositionRAII {
+  CGBuilderTy &Builder;
+  CGBuilderTy::InsertPoint SavedIP;
+
+public:
+  BuilderInsertPositionRAII(CGBuilderTy &Builder,
+                            llvm::Instruction *NewInsertPoint)
+      : Builder(Builder), SavedIP(Builder.saveIP()) {
+    assert(SavedIP.isSet() && "No insertion point is set!");
+    Builder.SetInsertPoint(NewInsertPoint);
+  }
+  ~BuilderInsertPositionRAII() { Builder.restoreIP(SavedIP); }
+};
+}
+
+llvm::Value *CGOpenMPRuntime::GetOpenMPLocation(CodeGenFunction &CGF,
+                                                SourceLocation Loc,
+                                                OpenMPLocationFlags Flags) {
+  // If no debug info is generated - return global default location.
+  if (CGM.getCodeGenOpts().getDebugInfo() == CodeGenOptions::NoDebugInfo)
+    return DefaultOpenMPLocation;
+
+  assert(CGF.CurFn && "No function in current CodeGenFunction.");
+
+  SmallString<32> Buffer;
+  llvm::raw_svector_ostream OS(Buffer);
+  OS << "__kmpc_loc_" << Flags << ".addr";
+
+  llvm::AllocaInst *AI = 0;
+  OpenMPLocMapTy::iterator I = OpenMPLocMap.find(CGF.CurFn);
+  if (I != OpenMPLocMap.end()) {
+    AI = I->second;
+  } else {
+    // Generate "ident_t __kmpc_loc_<flags>.addr;"
+    AI = CGF.CreateTempAlloca(Ident_TTy, OS.str());
+    AI->setAlignment(CGM.getDataLayout().getPrefTypeAlignment(Ident_TTy));
+    OpenMPLocMap[CGF.CurFn] = AI;
+
+    BuilderInsertPositionRAII BIP(CGF.Builder, CGF.AllocaInsertPt);
+    CGF.Builder.CreateMemCpy(AI, DefaultOpenMPLocation,
+                             llvm::ConstantExpr::getSizeOf(Ident_TTy),
+                             CGM.PointerAlignInBytes);
+    // int32 *flags = &(__kmpc_loc_<flags>.addr).flags;
+    llvm::Value *FlagsField =
+        CGF.Builder.CreateConstInBoundsGEP2_32(AI, 0, Ident_T_Flags);
+    // *flags = <Flags>;
+    CGF.Builder.CreateStore(CGF.Builder.getInt32(Flags), FlagsField);
+  }
+
+  // char **psource = &__kmpc_loc_<flags>.addr.psource;
+  llvm::Value *PSource =
+      CGF.Builder.CreateConstInBoundsGEP2_32(AI, 0, Ident_T_PSource);
+  if (CGM.getCodeGenOpts().getDebugInfo() != CodeGenOptions::NoDebugInfo) {
+    if (Loc.isValid()) {
+      SmallString<128> Buffer2;
+      llvm::raw_svector_ostream OS2(Buffer2);
+      // Build debug location
+      PresumedLoc PLoc =
+          CGF.getContext().getSourceManager().getPresumedLoc(Loc);
+      OS2 << ";" << PLoc.getFilename() << ";";
+      if (const FunctionDecl *FD =
+              dyn_cast_or_null<FunctionDecl>(CGF.CurFuncDecl)) {
+        OS2 << FD->getQualifiedNameAsString();
+      }
+      OS2 << ";" << PLoc.getLine() << ";" << PLoc.getColumn() << ";;";
+      // *psource = ";<File>;<Function>;<Line>;<Column>;;";
+      CGF.Builder.CreateStore(CGF.Builder.CreateGlobalStringPtr(OS2.str()),
+                              PSource);
+    } else {
+      // *psource = ";unknown;unknown;0;0;;";
+      CGF.Builder.CreateStore(DefaultOpenMPPSource, PSource);
+    }
+  }
+  return AI;
+}
+
+llvm::Value *CGOpenMPRuntime::GetOpenMPGlobalThreadNum(CodeGenFunction &CGF,
+                                                       SourceLocation Loc) {
+  assert(CGF.CurFn && "No function in current CodeGenFunction.");
+
+  llvm::AllocaInst *AI = 0;
+  OpenMPGtidMapTy::iterator I = OpenMPGtidMap.find(CGF.CurFn);
+  if (I != OpenMPGtidMap.end()) {
+    AI = I->second;
+  } else {
+    // Generate "int32 __kmpc_global_thread_num.addr;"
+    AI = CGF.CreateTempAlloca(CGF.Int32Ty, "__kmpc_global_thread_num.addr");
+    AI->setAlignment(CGM.getDataLayout().getPrefTypeAlignment(CGF.Int32Ty));
+
+    BuilderInsertPositionRAII BIP(CGF.Builder, CGF.AllocaInsertPt);
+    llvm::Value *Args[] = { GetOpenMPLocation(CGF, Loc) };
+    CGF.Builder.CreateStore(
+        CGF.EmitRuntimeCall(
+            CreateRuntimeFunction(OMPRTL__kmpc_global_thread_num), Args),
+        AI);
+  }
+  return CGF.Builder.CreateLoad(AI, "__kmpc_global_thread_num");
+}
+
+llvm::Type *CGOpenMPRuntime::getIdent_TPointerTy() {
+  return llvm::PointerType::getUnqual(Ident_TTy);
+}
+
+llvm::Type *CGOpenMPRuntime::getKmpc_MicroPointerTy() {
+  return llvm::PointerType::getUnqual(Kmpc_MicroTy);
+}
+
+llvm::Constant *
+CGOpenMPRuntime::CreateRuntimeFunction(OpenMPRTLFunction Function) {
+  llvm::Constant *RTLFn = 0;
+  switch (Function) {
+  case OMPRTL__kmpc_fork_call: {
+    // Build void __kmpc_fork_call(ident_t *loc, kmp_int32 argc, kmpc_micro
+    // microtask, ...);
+    llvm::Type *TypeParams[] = { getIdent_TPointerTy(), CGM.Int32Ty,
+                                 getKmpc_MicroPointerTy() };
+    llvm::FunctionType *FnTy =
+        llvm::FunctionType::get(CGM.VoidTy, TypeParams, true);
+    RTLFn = CGM.CreateRuntimeFunction(FnTy, "__kmpc_fork_call");
+  } break;
+  case OMPRTL__kmpc_global_thread_num: {
+    // Build kmp_int32 kmpc_global_thread_num(ident_t *loc);
+    llvm::Type *TypeParams[] = { getIdent_TPointerTy() };
+    llvm::FunctionType *FnTy =
+        llvm::FunctionType::get(CGM.Int32Ty, TypeParams, false);
+    RTLFn = CGM.CreateRuntimeFunction(FnTy, "__kmpc_global_thread_num");
+  } break;
+  }
+  return RTLFn;
+}

lib/CodeGen/CGStmt.cpp

void CodeGenFunction::EmitStmt(const Stmt *S) {
   EmitStopPoint(S);
 
   switch (S->getStmtClass()) {
   case Stmt::NoStmtClass:
   case Stmt::CXXCatchStmtClass:
   case Stmt::SEHExceptStmtClass:
   case Stmt::SEHFinallyStmtClass:
   case Stmt::MSDependentExistsStmtClass:
-  case Stmt::OMPParallelDirectiveClass:
     llvm_unreachable("invalid statement class to emit generically");
   case Stmt::NullStmtClass:
   case Stmt::CompoundStmtClass:
   case Stmt::DeclStmtClass:
   case Stmt::LabelStmtClass:
   case Stmt::AttributedStmtClass:
   case Stmt::GotoStmtClass:
   case Stmt::BreakStmtClass:
case Stmt::CXXTryStmtClass:
     EmitCXXTryStmt(cast<CXXTryStmt>(*S));
     break;
   case Stmt::CXXForRangeStmtClass:
     EmitCXXForRangeStmt(cast<CXXForRangeStmt>(*S));
     break;
   case Stmt::SEHTryStmtClass:
     EmitSEHTryStmt(cast<SEHTryStmt>(*S));
     break;
+  case Stmt::OMPParallelDirectiveClass:
+    EmitOMPParallelDirective(cast<OMPParallelDirective>(*S));
+    break;
   }
 }
 
 bool CodeGenFunction::EmitSimpleStmt(const Stmt *S) {
   switch (S->getStmtClass()) {
   default: return false;
   case Stmt::NullStmtClass: break;
   case Stmt::CompoundStmtClass: EmitCompoundStmt(cast<CompoundStmt>(*S)); break;
CodeGenFunction::EmitCapturedStmt(const CapturedStmt &S, CapturedRegionKind K) {
   delete CGF.CapturedStmtInfo;
 
   // Emit call to the helper function.
   EmitCallOrInvoke(F, CapStruct.getAddress());
 
   return F;
 }
 
+llvm::Value *
+CodeGenFunction::GenerateCapturedStmtArgument(const CapturedStmt &S) {
+  LValue CapStruct = InitCapturedStruct(*this, S);
+  return CapStruct.getAddress();
+}
+
 /// Creates the outlined function for a CapturedStmt.
 llvm::Function *
 CodeGenFunction::GenerateCapturedStmtFunction(const CapturedDecl *CD,
                                               const RecordDecl *RD,
                                               SourceLocation Loc) {
   assert(CapturedStmtInfo &&
     "CapturedStmtInfo should be set when generating the captured function");
 

lib/CodeGen/CGStmtOpenMP.cpp

+//===--- CGStmtOpenMP.cpp - Emit LLVM Code from Statements ----------------===//
+//
+//                     The LLVM Compiler Infrastructure
+//
+// This file is distributed under the University of Illinois Open Source
+// License. See LICENSE.TXT for details.
+//
+//===----------------------------------------------------------------------===//
+//
+// This contains code to emit OpenMP nodes as LLVM code.
+//
+//===----------------------------------------------------------------------===//
+
+#include "CGOpenMPRuntime.h"
+#include "CodeGenFunction.h"
+#include "CodeGenModule.h"
+#include "clang/AST/Stmt.h"
+#include "clang/AST/StmtOpenMP.h"
+using namespace clang;
+using namespace CodeGen;
+
+//===----------------------------------------------------------------------===//
+//                              OpenMP Directive Emission
+//===----------------------------------------------------------------------===//
+
+void CodeGenFunction::EmitOMPParallelDirective(const OMPParallelDirective &S) {
+  const CapturedStmt *CS = cast<CapturedStmt>(S.getAssociatedStmt());
+  llvm::Value *CapturedStruct = GenerateCapturedStmtArgument(*CS);
+
+  CodeGenFunction CGF(CGM, true);
+  CGF.CapturedStmtInfo =
+      new CGCapturedStmtInfo(*CS, CS->getCapturedRegionKind());
+  llvm::Value *OutlinedFn = CGF.GenerateCapturedStmtFunction(
+      CS->getCapturedDecl(), CS->getCapturedRecordDecl(), CS->getLocStart());
+  delete CGF.CapturedStmtInfo;
+
+  // Build call __kmpc_fork_call(loc, 1, microtask, captured_struct/*context*/)
+  llvm::Value *Args[] = {
+    CGM.getOpenMPRuntime().GetOpenMPLocation(*this, S.getLocStart()),
+    Builder.getInt32(1),
+    Builder.CreateBitCast(OutlinedFn,
+                          CGM.getOpenMPRuntime().getKmpc_MicroPointerTy()),
+    EmitCastToVoidPtr(CapturedStruct)
+  };
+  llvm::Constant *RTLFn = CGM.getOpenMPRuntime().CreateRuntimeFunction(
+      CGOpenMPRuntime::OMPRTL__kmpc_fork_call);
+  EmitRuntimeCall(RTLFn, Args);
+}

lib/CodeGen/CMakeLists.txt

add_clang_library(clangCodeGen
   CGExprComplex.cpp
   CGExprConstant.cpp
   CGExprScalar.cpp
   CGObjC.cpp
   CGObjCGNU.cpp
   CGObjCMac.cpp
   CGObjCRuntime.cpp
   CGOpenCLRuntime.cpp
+  CGOpenMPRuntime.cpp
   CGRTTI.cpp
   CGRecordLayoutBuilder.cpp
   CGStmt.cpp
+  CGStmtOpenMP.cpp
   CGVTT.cpp
   CGVTables.cpp
   CodeGenABITypes.cpp
   CodeGenAction.cpp
   CodeGenFunction.cpp
   CodeGenModule.cpp
   CodeGenPGO.cpp
   CodeGenTBAA.cpp

lib/CodeGen/CodeGenFunction.h

#endif
   void EmitCXXTryStmt(const CXXTryStmt &S);
   void EmitSEHTryStmt(const SEHTryStmt &S);
   void EmitCXXForRangeStmt(const CXXForRangeStmt &S);
 
   llvm::Function *EmitCapturedStmt(const CapturedStmt &S, CapturedRegionKind K);
   llvm::Function *GenerateCapturedStmtFunction(const CapturedDecl *CD,
                                                const RecordDecl *RD,
                                                SourceLocation Loc);
+  llvm::Value *GenerateCapturedStmtArgument(const CapturedStmt &S);
+
+  void EmitOMPParallelDirective(const OMPParallelDirective &S);
 
   //===--------------------------------------------------------------------===//
   //                         LValue Expression Emission
   //===--------------------------------------------------------------------===//
 
   /// GetUndefRValue - Get an appropriate 'undef' rvalue for the given type.
   RValue GetUndefRValue(QualType Ty);
 

lib/CodeGen/CodeGenModule.h

namespace CodeGen {
 
   class CallArgList;
   class CodeGenFunction;
   class CodeGenTBAA;
   class CGCXXABI;
   class CGDebugInfo;
   class CGObjCRuntime;
   class CGOpenCLRuntime;
+  class CGOpenMPRuntime;
   class CGCUDARuntime;
   class BlockFieldFlags;
   class FunctionArgList;
   class PGOProfileData;
 
   struct OrderGlobalInits {
     unsigned int priority;
     unsigned int lex_order;
class CodeGenModule : public CodeGenTypeCache {
   // if TheTargetCodeGenInfo is NULL
   CodeGenTypes Types;
  
   /// VTables - Holds information about C++ vtables.
   CodeGenVTables VTables;
 
   CGObjCRuntime* ObjCRuntime;
   CGOpenCLRuntime* OpenCLRuntime;
+  CGOpenMPRuntime* OpenMPRuntime;
   CGCUDARuntime* CUDARuntime;
   CGDebugInfo* DebugInfo;
   ARCEntrypoints *ARCData;
   llvm::MDNode *NoObjCARCExceptionsMetadata;
   RREntrypoints *RRData;
   PGOProfileData *PGOData;
 
   // WeakRefReferences - A set of references that have only been seen via
class CodeGenModule : public CodeGenTypeCache {
   QualType ObjCFastEnumerationStateType;
   
   /// @}
 
   /// Lazily create the Objective-C runtime
   void createObjCRuntime();
 
   void createOpenCLRuntime();
+  void createOpenMPRuntime();
   void createCUDARuntime();
 
   bool isTriviallyRecursive(const FunctionDecl *F);
   bool shouldEmitFunction(GlobalDecl GD);
 
   /// @name Cache for Blocks Runtime Globals
   /// @{
 
public:
   bool hasObjCRuntime() { return !!ObjCRuntime; }
 
   /// getOpenCLRuntime() - Return a reference to the configured OpenCL runtime.
   CGOpenCLRuntime &getOpenCLRuntime() {
     assert(OpenCLRuntime != 0);
     return *OpenCLRuntime;
   }
 
+  /// getOpenMPRuntime() - Return a reference to the configured OpenMP runtime.
+  CGOpenMPRuntime &getOpenMPRuntime() {
+    assert(OpenMPRuntime != 0);
+    return *OpenMPRuntime;
+  }
+
   /// getCUDARuntime() - Return a reference to the configured CUDA runtime.
   CGCUDARuntime &getCUDARuntime() {
     assert(CUDARuntime != 0);
     return *CUDARuntime;
   }
 
   ARCEntrypoints &getARCEntrypoints() const {
     assert(getLangOpts().ObjCAutoRefCount && ARCData != 0);

lib/CodeGen/CodeGenModule.cpp

 
 #include "CodeGenModule.h"
 #include "CGCUDARuntime.h"
 #include "CGCXXABI.h"
 #include "CGCall.h"
 #include "CGDebugInfo.h"
 #include "CGObjCRuntime.h"
 #include "CGOpenCLRuntime.h"
+#include "CGOpenMPRuntime.h"
 #include "CodeGenFunction.h"
 #include "CodeGenPGO.h"
 #include "CodeGenTBAA.h"
 #include "TargetInfo.h"
 #include "clang/AST/ASTContext.h"
 #include "clang/AST/CharUnits.h"
 #include "clang/AST/DeclCXX.h"
 #include "clang/AST/DeclObjC.h"

 
 CodeGenModule::CodeGenModule(ASTContext &C, const CodeGenOptions &CGO,
                              llvm::Module &M, const llvm::DataLayout &TD,
                              DiagnosticsEngine &diags)
     : Context(C), LangOpts(C.getLangOpts()), CodeGenOpts(CGO), TheModule(M),
       Diags(diags), TheDataLayout(TD), Target(C.getTargetInfo()),
       ABI(createCXXABI(*this)), VMContext(M.getContext()), TBAA(0),
       TheTargetCodeGenInfo(0), Types(*this), VTables(*this), ObjCRuntime(0),
-      OpenCLRuntime(0), CUDARuntime(0), DebugInfo(0), ARCData(0),
-      NoObjCARCExceptionsMetadata(0), RRData(0), PGOData(0),
+      OpenCLRuntime(0), OpenMPRuntime(0), CUDARuntime(0), DebugInfo(0),
+      ARCData(0), NoObjCARCExceptionsMetadata(0), RRData(0), PGOData(0),
       CFConstantStringClassRef(0),
       ConstantStringClassRef(0), NSConstantStringType(0),
       NSConcreteGlobalBlock(0), NSConcreteStackBlock(0), BlockObjectAssign(0),
       BlockObjectDispose(0), BlockDescriptorType(0), GenericBlockLiteralType(0),
       LifetimeStartFn(0), LifetimeEndFn(0),
       SanitizerBlacklist(
           llvm::SpecialCaseList::createOrDie(CGO.SanitizerBlacklistFile)),
       SanOpts(SanitizerBlacklist->isIn(M) ? SanitizerOptions::Disabled
CodeGenModule::CodeGenModule(ASTContext &C, const CodeGenOptions &CGO,
   Int8PtrPtrTy = Int8PtrTy->getPointerTo(0);
 
   RuntimeCC = getTargetCodeGenInfo().getABIInfo().getRuntimeCC();
 
   if (LangOpts.ObjC1)
     createObjCRuntime();
   if (LangOpts.OpenCL)
     createOpenCLRuntime();
+  if (LangOpts.OpenMP)
+    createOpenMPRuntime();
   if (LangOpts.CUDA)
     createCUDARuntime();
 
   // Enable TBAA unless it's suppressed. ThreadSanitizer needs TBAA even at O0.
   if (SanOpts.Thread ||
       (!CodeGenOpts.RelaxedAliasing && CodeGenOpts.OptimizationLevel > 0))
     TBAA = new CodeGenTBAA(Context, VMContext, CodeGenOpts, getLangOpts(),
                            getCXXABI().getMangleContext());
CodeGenModule::CodeGenModule(ASTContext &C, const CodeGenOptions &CGO,
 
   if (!CodeGenOpts.InstrProfileInput.empty())
     PGOData = new PGOProfileData(*this, CodeGenOpts.InstrProfileInput);
 }
 
 CodeGenModule::~CodeGenModule() {
   delete ObjCRuntime;
   delete OpenCLRuntime;
+  delete OpenMPRuntime;
   delete CUDARuntime;
   delete TheTargetCodeGenInfo;
   delete TBAA;
   delete DebugInfo;
   delete ARCData;
   delete RRData;
 }
 
void CodeGenModule::createObjCRuntime() {
   }
   llvm_unreachable("bad runtime kind");
 }
 
 void CodeGenModule::createOpenCLRuntime() {
   OpenCLRuntime = new CGOpenCLRuntime(*this);
 }
 
+void CodeGenModule::createOpenMPRuntime() {
+  OpenMPRuntime = new CGOpenMPRuntime(*this);
+}
+
 void CodeGenModule::createCUDARuntime() {
   CUDARuntime = CreateNVCUDARuntime(*this);
 }
 
 void CodeGenModule::applyReplacements() {
   for (ReplacementsTy::iterator I = Replacements.begin(),
                                 E = Replacements.end();
        I != E; ++I) {

lib/Parse/ParseOpenMP.cpp

 //===--- ParseOpenMP.cpp - OpenMP directives parsing ----------------------===//
 //
 //                     The LLVM Compiler Infrastructure
 //
 // This file is distributed under the University of Illinois Open Source
 // License. See LICENSE.TXT for details.
 //
 //===----------------------------------------------------------------------===//
 /// \file
 /// \brief This file implements parsing of all OpenMP directives and clauses.
 ///
 //===----------------------------------------------------------------------===//
 
-#include "clang/AST/ASTConsumer.h"
 #include "RAIIObjectsForParser.h"
+#include "clang/AST/ASTConsumer.h"
+#include "clang/AST/ASTContext.h"
 #include "clang/AST/StmtOpenMP.h"
 #include "clang/Parse/ParseDiagnostic.h"
 #include "clang/Parse/Parser.h"
 #include "clang/Sema/Scope.h"
 #include "llvm/ADT/PointerIntPair.h"
 using namespace clang;
 
 //===----------------------------------------------------------------------===//
case OMPD_parallel: {
     ConsumeToken();
 
     StmtResult AssociatedStmt;
     bool CreateDirective = true;
     ParseScope OMPDirectiveScope(this, ScopeFlags);
     {
       // The body is a block scope like in Lambdas and Blocks.
       Sema::CompoundScopeRAII CompoundScope(Actions);
-      Actions.ActOnCapturedRegionStart(Loc, getCurScope(), CR_OpenMP, 1);
+      Actions.ActOnOpenMPRegionStart(DKind, Loc, getCurScope());
       Actions.ActOnStartOfCompoundStmt();
       // Parse statement
       AssociatedStmt = ParseStatement();
       Actions.ActOnFinishOfCompoundStmt();
       if (!AssociatedStmt.isUsable()) {
         Actions.ActOnCapturedRegionError();
         CreateDirective = false;
       } else {

lib/Sema/SemaOpenMP.cpp

-//===--- SemaOpenMP.cpp - Semantic Analysis for OpenMP constructs ----------===//
+//===--- SemaOpenMP.cpp - Semantic Analysis for OpenMP constructs ---------===//
 //
 //                     The LLVM Compiler Infrastructure
 //
 // This file is distributed under the University of Illinois Open Source
 // License. See LICENSE.TXT for details.
 //
 //===----------------------------------------------------------------------===//
 /// \file
 /// \brief This file implements semantic analysis for OpenMP directives and
 /// clauses.
 ///
 //===----------------------------------------------------------------------===//
 
-#include "clang/Basic/OpenMPKinds.h"
+#include "clang/AST/ASTContext.h"
 #include "clang/AST/Decl.h"
 #include "clang/AST/DeclCXX.h"
 #include "clang/AST/DeclOpenMP.h"
 #include "clang/AST/StmtCXX.h"
 #include "clang/AST/StmtOpenMP.h"
 #include "clang/AST/StmtVisitor.h"
+#include "clang/Basic/OpenMPKinds.h"
 #include "clang/Lex/Preprocessor.h"
 #include "clang/Sema/Initialization.h"
 #include "clang/Sema/Lookup.h"
 #include "clang/Sema/Scope.h"
 #include "clang/Sema/ScopeInfo.h"
 #include "clang/Sema/SemaInternal.h"
 using namespace clang;
 
public:
   bool isErrorFound() { return ErrorFound; }
   ArrayRef<Expr *> getImplicitFirstprivate() { return ImplicitFirstprivate; }
 
   DSAAttrChecker(DSAStackTy *S, Sema &Actions, CapturedStmt *CS)
     : Stack(S), Actions(Actions), ErrorFound(false), CS(CS) { }
 };
 }
 
+void Sema::ActOnOpenMPRegionStart(OpenMPDirectiveKind DKind, SourceLocation Loc,
+                                  Scope *CurScope) {
+  switch (DKind) {
+  case OMPD_parallel: {
+    QualType KmpInt32Ty = Context.getIntTypeForBitwidth(32, 1);
+    QualType KmpInt32PtrTy = Context.getPointerType(KmpInt32Ty);
+    Sema::ParamNameType Params[3] =
+        { std::make_pair("__global_tid", KmpInt32PtrTy),
+          std::make_pair("__bound_tid", KmpInt32PtrTy),
+          std::make_pair(StringRef(), QualType()) // __context with shared vars
+        };
+    ActOnCapturedRegionStart(Loc, CurScope, CR_OpenMP, Params);
+    }
+    break;
+  case OMPD_threadprivate:
+  case OMPD_task:
+    llvm_unreachable("OpenMP Directive is not allowed");
+  case OMPD_unknown:
+  case NUM_OPENMP_DIRECTIVES:
+    llvm_unreachable("Unknown OpenMP directive");
+  }
+}
+
 StmtResult Sema::ActOnOpenMPExecutableDirective(OpenMPDirectiveKind Kind,
                                                 ArrayRef<OMPClause *> Clauses,
                                                 Stmt *AStmt,
                                                 SourceLocation StartLoc,
                                                 SourceLocation EndLoc) {
   assert(AStmt && isa<CapturedStmt>(AStmt) && "Captured statement expected");
 
   StmtResult Res = StmtError();
StmtResult Sema::ActOnOpenMPExecutableDirective(OpenMPDirectiveKind Kind,
   if (ErrorFound) return StmtError();
   return Res;
 }
 
 StmtResult Sema::ActOnOpenMPParallelDirective(ArrayRef<OMPClause *> Clauses,
                                               Stmt *AStmt,
                                               SourceLocation StartLoc,
                                               SourceLocation EndLoc) {
+  assert(AStmt && isa<CapturedStmt>(AStmt) && "Captured statement expected");
+  CapturedStmt *CS = cast<CapturedStmt>(AStmt);
+  // Mark captured decl as nothrow
+  CS->getCapturedDecl()->setNothrow();
+
   getCurFunction()->setHasBranchProtectedScope();
 
   return Owned(OMPParallelDirective::Create(Context, StartLoc, EndLoc,
                                             Clauses, AStmt));
 }
 
 OMPClause *Sema::ActOnOpenMPSingleExprClause(OpenMPClauseKind Kind,
                                              Expr *Expr,

lib/Sema/SemaStmt.cpp

if (getLangOpts().CPlusPlus)
     RD = CXXRecordDecl::Create(Context, TTK_Struct, DC, Loc, Loc, /*Id=*/0);
   else
     RD = RecordDecl::Create(Context, TTK_Struct, DC, Loc, Loc, /*Id=*/0);
 
   DC->addDecl(RD);
   RD->setImplicit();
   RD->startDefinition();
 
+  assert(NumParams > 0 && "CapturedStmt requires context parameter");
   CD = CapturedDecl::Create(Context, CurContext, NumParams);
   DC->addDecl(CD);
-
-  // Build the context parameter
-  assert(NumParams > 0 && "CapturedStmt requires context parameter");
-  DC = CapturedDecl::castToDeclContext(CD);
-  IdentifierInfo *VarName = &Context.Idents.get("__context");
-  QualType ParamType = Context.getPointerType(Context.getTagDeclType(RD));
-  ImplicitParamDecl *Param
-    = ImplicitParamDecl::Create(Context, DC, Loc, VarName, ParamType);
-  DC->addDecl(Param);
-
-  CD->setContextParam(Param);
-
   return RD;
 }
 
 static void buildCapturedStmtCaptureList(
     SmallVectorImpl<CapturedStmt::Capture> &Captures,
     SmallVectorImpl<Expr *> &CaptureInits,
     ArrayRef<CapturingScopeInfo::Capture> Candidates) {
 

 }
 
 void Sema::ActOnCapturedRegionStart(SourceLocation Loc, Scope *CurScope,
                                     CapturedRegionKind Kind,
                                     unsigned NumParams) {
   CapturedDecl *CD = 0;
   RecordDecl *RD = CreateCapturedStmtRecordDecl(CD, Loc, NumParams);
 
+  // Build the context parameter
+  DeclContext *DC = CapturedDecl::castToDeclContext(CD);
+  IdentifierInfo *ParamName = &Context.Idents.get("__context");
+  QualType ParamType = Context.getPointerType(Context.getTagDeclType(RD));
+  ImplicitParamDecl *Param
+    = ImplicitParamDecl::Create(Context, DC, Loc, ParamName, ParamType);
+  DC->addDecl(Param);
+
+  CD->setContextParam(0, Param);
+
+  // Enter the capturing scope for this captured region.
+  PushCapturedRegionScope(CurScope, CD, RD, Kind);
+
+  if (CurScope)
+    PushDeclContext(CurScope, CD);
+  else
+    CurContext = CD;
+
+  PushExpressionEvaluationContext(PotentiallyEvaluated);
+}
+
+void Sema::ActOnCapturedRegionStart(SourceLocation Loc, Scope *CurScope,
+                                    CapturedRegionKind Kind,
+                                    ArrayRef<ParamNameType> Params) {
+  CapturedDecl *CD = 0;
+  RecordDecl *RD = CreateCapturedStmtRecordDecl(CD, Loc, Params.size());
+
+  // Build the context parameter
+  DeclContext *DC = CapturedDecl::castToDeclContext(CD);
+  bool ContextIsFound = false;
+  unsigned ParamNum = 0;
+  for (ArrayRef<ParamNameType>::iterator I = Params.begin(), E = Params.end();
+       I != E; ++I, ++ParamNum) {
+    if (I->second.isNull()) {
+      assert(!ContextIsFound &&
+             "null type has been found already for '__context' parameter");
+      IdentifierInfo *ParamName = &Context.Idents.get("__context");
+      QualType ParamType = Context.getPointerType(Context.getTagDeclType(RD));
+      ImplicitParamDecl *Param
+        = ImplicitParamDecl::Create(Context, DC, Loc, ParamName, ParamType);
+      DC->addDecl(Param);
+      CD->setContextParam(ParamNum, Param);
+      ContextIsFound = true;
+    } else {
+      IdentifierInfo *ParamName = &Context.Idents.get(I->first);
+      ImplicitParamDecl *Param
+        = ImplicitParamDecl::Create(Context, DC, Loc, ParamName, I->second);
+      DC->addDecl(Param);
+      CD->setParam(ParamNum, Param);
+    }
+  }
+  assert(ContextIsFound && "no null type for '__context' parameter");
   // Enter the capturing scope for this captured region.
   PushCapturedRegionScope(CurScope, CD, RD, Kind);
 
   if (CurScope)
     PushDeclContext(CurScope, CD);
   else
     CurContext = CD;
 

lib/Sema/TreeTransform.h

return getSema().BuildMemberReferenceExpr(Base, BaseType,
                                             NameInfo,
                                             /*TemplateArgs*/ 0);
 }
 
 template<typename Derived>
 StmtResult
 TreeTransform<Derived>::TransformCapturedStmt(CapturedStmt *S) {
   SourceLocation Loc = S->getLocStart();
-  unsigned NumParams = S->getCapturedDecl()->getNumParams();
+  CapturedDecl *CD = S->getCapturedDecl();
+  unsigned NumParams = CD->getNumParams();
+  unsigned ContextParamPos = CD->getContextParamPosition();
+  SmallVector<Sema::ParamNameType, 4> Params;
+  for (unsigned i = 0; i < NumParams; ++i) {
+    if (i != ContextParamPos) {
+      Params.push_back(
+             std::make_pair(
+                  CD->getParam(i)->getName(),
+                  getDerived().TransformType(CD->getParam(i)->getType())));
+    } else {
+      Params.push_back(std::make_pair(StringRef(), QualType()));
+    }
+  }
   getSema().ActOnCapturedRegionStart(Loc, /*CurScope*/0,
-                                     S->getCapturedRegionKind(), NumParams);
+                                     S->getCapturedRegionKind(), Params);
   StmtResult Body = getDerived().TransformStmt(S->getCapturedStmt());
 
   if (Body.isInvalid()) {
     getSema().ActOnCapturedRegionError();
     return StmtError();
   }
 
   return getSema().ActOnCapturedRegionEnd(Body.take());
 }
 
 } // end namespace clang
 
 #endif // LLVM_CLANG_SEMA_TREETRANSFORM_H

lib/Serialization/ASTReaderDecl.cpp

for (unsigned i = 0; i != numCaptures; ++i) {
     captures.push_back(BlockDecl::Capture(decl, byRef, nested, copyExpr));
   }
   BD->setCaptures(Reader.getContext(), captures.begin(),
                   captures.end(), capturesCXXThis);
 }
 
 void ASTDeclReader::VisitCapturedDecl(CapturedDecl *CD) {
   VisitDecl(CD);
+  unsigned ContextParamPos = Record[Idx++];
+  CD->setNothrow(Record[Idx++] != 0);
   // Body is set by VisitCapturedStmt.
-  for (unsigned i = 0; i < CD->NumParams; ++i)
+  for (unsigned i = 0; i < CD->NumParams; ++i) {
+    if (i != ContextParamPos)
       CD->setParam(i, ReadDeclAs<ImplicitParamDecl>(Record, Idx));
+    else
+      CD->setContextParam(i, ReadDeclAs<ImplicitParamDecl>(Record, Idx));
+  }
 }
 
 void ASTDeclReader::VisitLinkageSpecDecl(LinkageSpecDecl *D) {
   VisitDecl(D);
   D->setLanguage((LinkageSpecDecl::LanguageIDs)Record[Idx++]);
   D->setExternLoc(ReadSourceLocation(Record, Idx));
   D->setRBraceLoc(ReadSourceLocation(Record, Idx));
 }

lib/Serialization/ASTWriterDecl.cpp

void ASTDeclWriter::VisitBlockDecl(BlockDecl *D) {
   }
 
   Code = serialization::DECL_BLOCK;
 }
 
 void ASTDeclWriter::VisitCapturedDecl(CapturedDecl *CD) {
   Record.push_back(CD->getNumParams());
   VisitDecl(CD);
+  Record.push_back(CD->getContextParamPosition());
+  Record.push_back(CD->isNothrow() ? 1 : 0);
   // Body is stored by VisitCapturedStmt.
   for (unsigned i = 0; i < CD->getNumParams(); ++i)
     Writer.AddDeclRef(CD->getParam(i), Record);
   Code = serialization::DECL_CAPTURED;
 }
 
 void ASTDeclWriter::VisitLinkageSpecDecl(LinkageSpecDecl *D) {
   VisitDecl(D);

test/OpenMP/parallel_codegen.cpp

+// RUN: %clang_cc1 -verify -fopenmp -x c++ -emit-llvm %s -fexceptions -fcxx-exceptions -o - | FileCheck %s
+// RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -emit-pch -o %t %s
+// RUN: %clang_cc1 -fopenmp -x c++ -triple x86_64-unknown-unknown -fexceptions -fcxx-exceptions -g -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix=CHECK-DEBUG %s
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+// CHECK-DAG: %ident_t = type { i32, i32, i32, i32, i8* }
+// CHECK-DAG: %struct.anon = type { i32* }
+// CHECK-DAG: %struct.anon.0 = type { i8*** }
+// CHECK-DAG: @.str = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
+// CHECK-DAG: @__kmpc_default_loc.addr = private unnamed_addr constant %ident_t { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8]* @.str, i32 0, i32 0) }
+// CHECK-DEBUG-DAG: %ident_t = type { i32, i32, i32, i32, i8* }
+// CHECK-DEBUG-DAG: %struct.anon = type { i32* }
+// CHECK-DEBUG-DAG: %struct.anon.0 = type { i8*** }
+// CHECK-DEBUG-DAG: @.str = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00"
+// CHECK-DEBUG-DAG: @__kmpc_default_loc.addr = private unnamed_addr constant %ident_t { i32 0, i32 2, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8]* @.str, i32 0, i32 0) }
+// CHECK-DEBUG-DAG: [[LOC1:@.+]] = private unnamed_addr constant [{{.+}} x i8] c";{{.*}}parallel_codegen.cpp;main;[[@LINE+14]];9;;\00"
+// CHECK-DEBUG-DAG: [[LOC2:@.+]] = private unnamed_addr constant [{{.+}} x i8] c";{{.*}}parallel_codegen.cpp;tmain;[[@LINE+7]];9;;\00"
+
+template <class T>
+void foo(T argc) {}
+
+template <typename T>
+int tmain(T argc) {
+#pragma omp parallel
+  foo(argc);
+  return 0;
+}
+
+int main (int argc, char **argv) {
+#pragma omp parallel
+  foo(argc);
+  return tmain(argv);
+}
+
+// CHECK-LABEL: define i32 @main(i32 %argc, i8** %argv)
+// CHECK:       %agg.captured = alloca %struct.anon
+// CHECK:       [[ARGC_REF:%.+]] = getelementptr inbounds %struct.anon* %agg.captured, i32 0, i32 0
+// CHECK-NEXT:  store i32* %argc.addr, i32** [[ARGC_REF]]
+// CHECK-NEXT:  [[BITCAST:%.+]] = bitcast %struct.anon* %agg.captured to i8*
+// CHECK-NEXT:  call void (%ident_t*, i32, void (i32*, i32*, ...)*, ...)* @__kmpc_fork_call(%ident_t* @__kmpc_default_loc.addr, i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, %struct.anon*)* @__captured_stmt to void (i32*, i32*, ...)*), i8* [[BITCAST]])
+// CHECK-NEXT:  [[ARGV:%.+]] = load i8*** %argv.addr, align 8
+// CHECK-NEXT:  [[RET:%.+]] = call i32 @_Z5tmainIPPcEiT_(i8** [[ARGV]])
+// CHECK-NEXT:  ret i32 [[RET]]
+// CHECK-NEXT:  }
+// CHECK-DEBUG-LABEL: define i32 @main(i32 %argc, i8** %argv)
+// CHECK-DEBUG-DAG:   %agg.captured = alloca %struct.anon
+// CHECK-DEBUG-DAG:   %__kmpc_loc_2.addr = alloca %ident_t
+// CHECK-DEBUG:       [[KMPC_LOC_VOIDPTR:%.+]] = bitcast %ident_t* %__kmpc_loc_2.addr to i8*
+// CHECK-DEBUG-NEXT:  [[KMPC_DEFAULT_LOC_VOIDPTR:%.+]] = bitcast %ident_t* @__kmpc_default_loc.addr to i8*
+// CHECK-DEBUG-NEXT:   call void @llvm.memcpy.p0i8.p0i8.i64(i8* [[KMPC_LOC_VOIDPTR]], i8* [[KMPC_DEFAULT_LOC_VOIDPTR]], i64 ptrtoint (%ident_t* getelementptr (%ident_t* null, i32 1) to i64), i32 8, i1 false)
+// CHECK-DEBUG:       [[KMPC_LOC_FLAGS_REF:%.+]] = getelementptr inbounds %ident_t* %__kmpc_loc_2.addr, i32 0, i32 1
+// CHECK-DEBUG-NEXT:  store i32 2, i32* [[KMPC_LOC_FLAGS_REF]]
+// CHECK-DEBUG:       [[ARGC_REF:%.+]] = getelementptr inbounds %struct.anon* %agg.captured, i32 0, i32 0
+// CHECK-DEBUG-NEXT:  store i32* %argc.addr, i32** [[ARGC_REF]]
+// CHECK-DEBUG-NEXT:  [[KMPC_LOC_PSOURCE_REF:%.+]] = getelementptr inbounds %ident_t* %__kmpc_loc_2.addr, i32 0, i32 4
+// CHECK-DEBUG-NEXT:  store i8* getelementptr inbounds ([{{.+}} x i8]* [[LOC1]], i32 0, i32 0), i8** [[KMPC_LOC_PSOURCE_REF]]
+// CHECK-DEBUG-NEXT:  [[BITCAST:%.+]] = bitcast %struct.anon* %agg.captured to i8*
+// CHECK-DEBUG-NEXT:  call void (%ident_t*, i32, void (i32*, i32*, ...)*, ...)* @__kmpc_fork_call(%ident_t* %__kmpc_loc_2.addr, i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, %struct.anon*)* @__captured_stmt to void (i32*, i32*, ...)*), i8* [[BITCAST]])
+// CHECK-DEBUG-NEXT:  [[ARGV:%.+]] = load i8*** %argv.addr, align 8
+// CHECK-DEBUG-NEXT:  [[RET:%.+]] = call i32 @_Z5tmainIPPcEiT_(i8** [[ARGV]])
+// CHECK-DEBUG-NEXT:  ret i32 [[RET]]
+// CHECK-DEBUG-NEXT:  }
+
+// CHECK-LABEL: define internal void @__captured_stmt(i32* %__global_tid, i32* %__bound_tid, %struct.anon* %__context)
+// CHECK:       %__context.addr = alloca %struct.anon*
+// CHECK:       store %struct.anon* %__context, %struct.anon** %__context.addr
+// CHECK-NEXT:  [[CONTEXT_PTR:%.+]] = load %struct.anon** %__context.addr
+// CHECK-NEXT:  [[ARGC_PTR_REF:%.+]] = getelementptr inbounds %struct.anon* [[CONTEXT_PTR]], i32 0, i32 0
+// CHECK-NEXT:  [[ARGC_REF:%.+]] = load i32** [[ARGC_PTR_REF]]
+// CHECK-NEXT:  [[ARGC:%.+]] = load i32* [[ARGC_REF]]
+// CHECK-NEXT:  invoke void @_Z3fooIiEvT_(i32 [[ARGC]])
+// CHECK:       ret void
+// CHECK:       call void @__clang_call_terminate(i8*
+// CHECK-NEXT:  unreachable
+// CHECK-NEXT:  }
+// CHECK-DEBUG-LABEL: define internal void @__captured_stmt(i32* %__global_tid, i32* %__bound_tid, %struct.anon* %__context)
+// CHECK-DEBUG:       %__context.addr = alloca %struct.anon*
+// CHECK-DEBUG:       store %struct.anon* %__context, %struct.anon** %__context.addr
+// CHECK-DEBUG:       [[CONTEXT_PTR:%.+]] = load %struct.anon** %__context.addr
+// CHECK-DEBUG-NEXT:  [[ARGC_PTR_REF:%.+]] = getelementptr inbounds %struct.anon* [[CONTEXT_PTR]], i32 0, i32 0
+// CHECK-DEBUG-NEXT:  [[ARGC_REF:%.+]] = load i32** [[ARGC_PTR_REF]]
+// CHECK-DEBUG-NEXT:  [[ARGC:%.+]] = load i32* [[ARGC_REF]]
+// CHECK-DEBUG-NEXT:  invoke void @_Z3fooIiEvT_(i32 [[ARGC]])
+// CHECK-DEBUG:       ret void
+// CHECK-DEBUG:       call void @__clang_call_terminate(i8*
+// CHECK-DEBUG-NEXT:  unreachable
+// CHECK-DEBUG-NEXT:  }
+
+// CHECK-DAG: define linkonce_odr void @_Z3fooIiEvT_(i32 %argc)
+// CHECK-DAG: declare void @__kmpc_fork_call(%ident_t*, i32, void (i32*, i32*, ...)*, ...)
+// CHECK-DEBUG-DAG: define linkonce_odr void @_Z3fooIiEvT_(i32 %argc)
+// CHECK-DEBUG-DAG: declare void @__kmpc_fork_call(%ident_t*, i32, void (i32*, i32*, ...)*, ...)
+
+// CHECK-LABEL: define linkonce_odr i32 @_Z5tmainIPPcEiT_(i8** %argc)
+// CHECK:       %agg.captured = alloca %struct.anon.0
+// CHECK:       [[ARGC_REF:%.+]] = getelementptr inbounds %struct.anon.0* %agg.captured, i32 0, i32 0
+// CHECK-NEXT:  store i8*** %argc.addr, i8**** [[ARGC_REF]]
+// CHECK-NEXT:  [[BITCAST:%.+]] = bitcast %struct.anon.0* %agg.captured to i8*
+// CHECK-NEXT:  call void (%ident_t*, i32, void (i32*, i32*, ...)*, ...)* @__kmpc_fork_call(%ident_t* @__kmpc_default_loc.addr, i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, %struct.anon.0*)* @__captured_stmt1 to void (i32*, i32*, ...)*), i8* [[BITCAST]])
+// CHECK-NEXT:  ret i32 0
+// CHECK-NEXT:  }
+// CHECK-DEBUG-LABEL: define linkonce_odr i32 @_Z5tmainIPPcEiT_(i8** %argc)
+// CHECK-DEBUG-DAG:   %agg.captured = alloca %struct.anon.0
+// CHECK-DEBUG-DAG:   %__kmpc_loc_2.addr = alloca %ident_t
+// CHECK-DEBUG:       [[KMPC_LOC_VOIDPTR:%.+]] = bitcast %ident_t* %__kmpc_loc_2.addr to i8*
+// CHECK-DEBUG-NEXT:  [[KMPC_DEFAULT_LOC_VOIDPTR:%.+]] = bitcast %ident_t* @__kmpc_default_loc.addr to i8*
+// CHECK-DEBUG-NEXT:   call void @llvm.memcpy.p0i8.p0i8.i64(i8* [[KMPC_LOC_VOIDPTR]], i8* [[KMPC_DEFAULT_LOC_VOIDPTR]], i64 ptrtoint (%ident_t* getelementptr (%ident_t* null, i32 1) to i64), i32 8, i1 false)
+// CHECK-DEBUG:       [[KMPC_LOC_FLAGS_REF:%.+]] = getelementptr inbounds %ident_t* %__kmpc_loc_2.addr, i32 0, i32 1
+// CHECK-DEBUG-NEXT:  store i32 2, i32* [[KMPC_LOC_FLAGS_REF]]
+// CHECK-DEBUG:       [[ARGC_REF:%.+]] = getelementptr inbounds %struct.anon.0* %agg.captured, i32 0, i32 0
+// CHECK-DEBUG-NEXT:  store i8*** %argc.addr, i8**** [[ARGC_REF]]
+// CHECK-DEBUG-NEXT:  [[KMPC_LOC_PSOURCE_REF:%.+]] = getelementptr inbounds %ident_t* %__kmpc_loc_2.addr, i32 0, i32 4
+// CHECK-DEBUG-NEXT:  store i8* getelementptr inbounds ([{{.+}} x i8]* [[LOC2]], i32 0, i32 0), i8** [[KMPC_LOC_PSOURCE_REF]]
+// CHECK-DEBUG-NEXT:  [[BITCAST:%.+]] = bitcast %struct.anon.0* %agg.captured to i8*
+// CHECK-DEBUG-NEXT:  call void (%ident_t*, i32, void (i32*, i32*, ...)*, ...)* @__kmpc_fork_call(%ident_t* %__kmpc_loc_2.addr, i32 1, void (i32*, i32*, ...)* bitcast (void (i32*, i32*, %struct.anon.0*)* @__captured_stmt1 to void (i32*, i32*, ...)*), i8* [[BITCAST]])
+// CHECK-DEBUG-NEXT:  ret i32 0
+// CHECK-DEBUG-NEXT:  }
+
+// CHECK-LABEL: define internal void @__captured_stmt1(i32* %__global_tid, i32* %__bound_tid, %struct.anon.0* %__context)
+// CHECK:       %__context.addr = alloca %struct.anon.0*, align 8
+// CHECK:       store %struct.anon.0* %__context, %struct.anon.0** %__context.addr, align 8
+// CHECK-NEXT:  [[CONTEXT_PTR:%.+]] = load %struct.anon.0** %__context.addr
+// CHECK-NEXT:  [[ARGC_PTR_REF:%.+]] = getelementptr inbounds %struct.anon.0* [[CONTEXT_PTR]], i32 0, i32 0
+// CHECK-NEXT:  [[ARGC_REF:%.+]] = load i8**** [[ARGC_PTR_REF]]
+// CHECK-NEXT:  [[ARGC:%.+]] = load i8*** [[ARGC_REF]]
+// CHECK-NEXT:  invoke void @_Z3fooIPPcEvT_(i8** [[ARGC]])
+// CHECK:       ret void
+// CHECK:       call void @__clang_call_terminate(i8*
+// CHECK-NEXT:  unreachable
+// CHECK-NEXT:  }
+// CHECK-DEBUG-LABEL: define internal void @__captured_stmt1(i32* %__global_tid, i32* %__bound_tid, %struct.anon.0* %__context)
+// CHECK-DEBUG:       %__context.addr = alloca %struct.anon.0*, align 8
+// CHECK-DEBUG:       store %struct.anon.0* %__context, %struct.anon.0** %__context.addr, align 8
+// CHECK-DEBUG:       [[CONTEXT_PTR:%.+]] = load %struct.anon.0** %__context.addr
+// CHECK-DEBUG-NEXT:  [[ARGC_PTR_REF:%.+]] = getelementptr inbounds %struct.anon.0* [[CONTEXT_PTR]], i32 0, i32 0
+// CHECK-DEBUG-NEXT:  [[ARGC_REF:%.+]] = load i8**** [[ARGC_PTR_REF]]
+// CHECK-DEBUG-NEXT:  [[ARGC:%.+]] = load i8*** [[ARGC_REF]]
+// CHECK-DEBUG-NEXT:  invoke void @_Z3fooIPPcEvT_(i8** [[ARGC]])
+// CHECK-DEBUG:       ret void
+// CHECK-DEBUG:       call void @__clang_call_terminate(i8*
+// CHECK-DEBUG-NEXT:  unreachable
+// CHECK-DEBUG-NEXT:  }
+
+// CHECK: define linkonce_odr void @_Z3fooIPPcEvT_(i8** %argc)
+// CHECK-DEBUG: define linkonce_odr void @_Z3fooIPPcEvT_(i8** %argc)
+
+#endif