[OpenMP] Add OpenMP data sharing infrastructure using global memory

[gtbercea]

Summary

This patch handles the Clang code generation phase for the OpenMP data sharing infrastructure.

TODO: add a more detailed description.

[ABataev]

lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp

555 ↗

(On Diff #135565)

Do we need this?

640–641 ↗

(On Diff #135565)

auto->llvm::Value *

863–864 ↗

(On Diff #135565)

1. Update it to the latest version. And use clang-format!
2. This must be in a separate patch

933 ↗

(On Diff #135565)

Separate patch

1101 ↗

(On Diff #135565)

{}->llvm::None

1203–1210 ↗

(On Diff #135565)

Should it be in this patch?

lib/CodeGen/CGOpenMPRuntimeNVPTX.h

324–331 ↗

(On Diff #135565)

Do we need this function at all?

[ABataev]

LG

[gtbercea]

Add init stack function.

[gtbercea]

Patch splitting: limit support in this patch to standalone target regions only. Support for combined directives will be fully covered in a subsequent patch.

[ABataev]

LG

cfe/trunk/lib/CodeGen/CGDecl.cpp

if (D.getInit() && (Ty->isArrayType() || Ty->isRecordType()) &&
         return emission;
       }
 
       // Otherwise, tell the initialization code that we're in this case.
       emission.IsConstantAggregate = true;
     }
 
     // A normal fixed sized variable becomes an alloca in the entry block,
-    // unless it's an NRVO variable.
-
-    if (NRVO) {
+    // unless:
+    // - it's an NRVO variable.
+    // - we are compiling OpenMP and it's an OpenMP local variable.
+
+    Address OpenMPLocalAddr =
+        getLangOpts().OpenMP
+            ? CGM.getOpenMPRuntime().getAddressOfLocalVariable(*this, &D)
+            : Address::invalid();
+    if (getLangOpts().OpenMP && OpenMPLocalAddr.isValid()) {
+      address = OpenMPLocalAddr;
+    } else if (NRVO) {
       // The named return value optimization: allocate this variable in the
       // return slot, so that we can elide the copy when returning this
       // variable (C++0x [class.copy]p34).
       address = ReturnValue;
 
       if (const RecordType *RecordTy = Ty->getAs<RecordType>()) {
         if (!cast<CXXRecordDecl>(RecordTy->getDecl())->hasTrivialDestructor()) {
           // Create a flag that is used to indicate when the NRVO was applied
if (!IsScalar && !CurFuncIsThunk &&
       if (QualType::DestructionKind DtorKind = Ty.isDestructedType()) {
         assert((DtorKind == QualType::DK_cxx_destructor ||
                 DtorKind == QualType::DK_nontrivial_c_struct) &&
                "unexpected destructor type");
         pushDestroy(DtorKind, DeclPtr, Ty);
       }
     }
   } else {
+    // Check if the parameter address is controlled by OpenMP runtime.
+    Address OpenMPLocalAddr =
+        getLangOpts().OpenMP
+            ? CGM.getOpenMPRuntime().getAddressOfLocalVariable(*this, &D)
+            : Address::invalid();
+    if (getLangOpts().OpenMP && OpenMPLocalAddr.isValid()) {
+      DeclPtr = OpenMPLocalAddr;
+    } else {
       // Otherwise, create a temporary to hold the value.
       DeclPtr = CreateMemTemp(Ty, getContext().getDeclAlign(&D),
                               D.getName() + ".addr");
+    }
     DoStore = true;
   }
 
   llvm::Value *ArgVal = (DoStore ? Arg.getDirectValue() : nullptr);
 
   LValue lv = MakeAddrLValue(DeclPtr, Ty);
   if (IsScalar) {
     Qualifiers qs = Ty.getQualifiers();

cfe/trunk/lib/CodeGen/CGOpenMPRuntime.h

public:
   virtual llvm::Value *emitTaskOutlinedFunction(
       const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
       const VarDecl *PartIDVar, const VarDecl *TaskTVar,
       OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen,
       bool Tied, unsigned &NumberOfParts);
 
   /// \brief Cleans up references to the objects in finished function.
   ///
-  void functionFinished(CodeGenFunction &CGF);
+  virtual void functionFinished(CodeGenFunction &CGF);
 
   /// \brief Emits code for parallel or serial call of the \a OutlinedFn with
   /// variables captured in a record which address is stored in \a
   /// CapturedStruct.
   /// \param OutlinedFn Outlined function to be run in parallel threads. Type of
   /// this function is void(*)(kmp_int32 *, kmp_int32, struct context_vars*).
   /// \param CapturedVars A pointer to the record with the references to
   /// variables used in \a OutlinedFn function.
virtual Address getParameterAddress(CodeGenFunction &CGF,
                                       const VarDecl *TargetParam) const;
 
   /// Emits call of the outlined function with the provided arguments,
   /// translating these arguments to correct target-specific arguments.
   virtual void
   emitOutlinedFunctionCall(CodeGenFunction &CGF, SourceLocation Loc,
                            llvm::Value *OutlinedFn,
                            ArrayRef<llvm::Value *> Args = llvm::None) const;
+
+  /// Emits OpenMP-specific function prolog.
+  /// Required for device constructs.
+  virtual void emitFunctionProlog(CodeGenFunction &CGF, const Decl *D) {}
+
+  /// Gets the OpenMP-specific address of the local variable.
+  virtual Address getAddressOfLocalVariable(CodeGenFunction &CGF,
+                                            const VarDecl *VD);
 };
 
 /// Class supports emissionof SIMD-only code.
 class CGOpenMPSIMDRuntime final : public CGOpenMPRuntime {
 public:
   explicit CGOpenMPSIMDRuntime(CodeGenModule &CGM) : CGOpenMPRuntime(CGM) {}
   ~CGOpenMPSIMDRuntime() override {}
 

cfe/trunk/lib/CodeGen/CGOpenMPRuntime.cpp

 }
 
 Address CGOpenMPRuntime::getParameterAddress(CodeGenFunction &CGF,
                                              const VarDecl *NativeParam,
                                              const VarDecl *TargetParam) const {
   return CGF.GetAddrOfLocalVar(NativeParam);
 }
 
+Address CGOpenMPRuntime::getAddressOfLocalVariable(CodeGenFunction &CGF,
+                                                   const VarDecl *VD) {
+  return Address::invalid();
+}
+
 llvm::Value *CGOpenMPSIMDRuntime::emitParallelOutlinedFunction(
     const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
     OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen) {
   llvm_unreachable("Not supported in SIMD-only mode");
 }
 
 llvm::Value *CGOpenMPSIMDRuntime::emitTeamsOutlinedFunction(
     const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,

cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.h

Address getParameterAddress(CodeGenFunction &CGF, const VarDecl *NativeParam,
                               const VarDecl *TargetParam) const override;
 
   /// Emits call of the outlined function with the provided arguments,
   /// translating these arguments to correct target-specific arguments.
   void emitOutlinedFunctionCall(
       CodeGenFunction &CGF, SourceLocation Loc, llvm::Value *OutlinedFn,
       ArrayRef<llvm::Value *> Args = llvm::None) const override;
 
+  /// Emits OpenMP-specific function prolog.
+  /// Required for device constructs.
+  void emitFunctionProlog(CodeGenFunction &CGF, const Decl *D) override;
+
+  /// Gets the OpenMP-specific address of the local variable.
+  Address getAddressOfLocalVariable(CodeGenFunction &CGF,
+                                    const VarDecl *VD) override;
+
   /// Target codegen is specialized based on two programming models: the
   /// 'generic' fork-join model of OpenMP, and a more GPU efficient 'spmd'
   /// model for constructs like 'target parallel' that support it.
   enum ExecutionMode {
     /// Single Program Multiple Data.
     Spmd,
     /// Generic codegen to support fork-join model.
     Generic,
     Unknown,
   };
 
+  /// Cleans up references to the objects in finished function.
+  ///
+  void functionFinished(CodeGenFunction &CGF) override;
+
 private:
   // Track the execution mode when codegening directives within a target
   // region. The appropriate mode (generic/spmd) is set on entry to the
   // target region and used by containing directives such as 'parallel'
   // to emit optimized code.
   ExecutionMode CurrentExecutionMode;
+
+  /// Map between an outlined function and its wrapper.
+  llvm::DenseMap<llvm::Function *, llvm::Function *> WrapperFunctionsMap;
+
+  /// Emit function which wraps the outline parallel region
+  /// and controls the parameters which are passed to this function.
+  /// The wrapper ensures that the outlined function is called
+  /// with the correct arguments when data is shared.
+  llvm::Function *createParallelDataSharingWrapper(
+      llvm::Function *OutlinedParallelFn, const OMPExecutableDirective &D);
+
+  /// The map of local variables to their addresses in the global memory.
+  using DeclToAddrMapTy = llvm::MapVector<const Decl *,
+      std::pair<const FieldDecl *, Address>>;
+  /// Maps the function to the list of the globalized variables with their
+  /// addresses.
+  llvm::DenseMap<llvm::Function *,
+                 std::pair<const RecordDecl *, DeclToAddrMapTy>>
+      FunctionGlobalizedDecls;
+  /// Map from function to global record pointer.
+  llvm::DenseMap<llvm::Function *, llvm::Value *> FunctionToGlobalRecPtr;
 };
 
 } // CodeGen namespace.
 } // clang namespace.
 
 #endif // LLVM_CLANG_LIB_CODEGEN_CGOPENMPRUNTIMENVPTX_H

cfe/trunk/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp

 //===---- CGOpenMPRuntimeNVPTX.cpp - Interface to OpenMP NVPTX 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 specialized to NVPTX
 // targets.
 //
 //===----------------------------------------------------------------------===//
 
 #include "CGOpenMPRuntimeNVPTX.h"
-#include "clang/AST/DeclOpenMP.h"
 #include "CodeGenFunction.h"
+#include "clang/AST/DeclOpenMP.h"
 #include "clang/AST/StmtOpenMP.h"
+#include "clang/AST/StmtVisitor.h"
+#include "llvm/ADT/SmallPtrSet.h"
 
 using namespace clang;
 using namespace CodeGen;
 
 namespace {
 enum OpenMPRTLFunctionNVPTX {
   /// \brief Call to void __kmpc_kernel_init(kmp_int32 thread_limit,
   /// int16_t RequiresOMPRuntime);
enum OpenMPRTLFunctionNVPTX {
   /// lane_offset, int16_t shortCircuit),
   /// void (*kmp_InterWarpCopyFctPtr)(void* src, int32_t warp_num),
   /// void (*kmp_CopyToScratchpadFctPtr)(void *reduce_data, void * scratchpad,
   /// int32_t index, int32_t width),
   /// void (*kmp_LoadReduceFctPtr)(void *reduce_data, void * scratchpad, int32_t
   /// index, int32_t width, int32_t reduce))
   OMPRTL_NVPTX__kmpc_teams_reduce_nowait,
   /// \brief Call to __kmpc_nvptx_end_reduce_nowait(int32_t global_tid);
-  OMPRTL_NVPTX__kmpc_end_reduce_nowait
+  OMPRTL_NVPTX__kmpc_end_reduce_nowait,
+  /// \brief Call to void __kmpc_data_sharing_init_stack();
+  OMPRTL_NVPTX__kmpc_data_sharing_init_stack,
+  /// \brief Call to void* __kmpc_data_sharing_push_stack(size_t size,
+  /// int16_t UseSharedMemory);
+  OMPRTL_NVPTX__kmpc_data_sharing_push_stack,
+  /// \brief Call to void __kmpc_data_sharing_pop_stack(void *a);
+  OMPRTL_NVPTX__kmpc_data_sharing_pop_stack,
+  /// \brief Call to void __kmpc_begin_sharing_variables(void ***args,
+  /// size_t n_args);
+  OMPRTL_NVPTX__kmpc_begin_sharing_variables,
+  /// \brief Call to void __kmpc_end_sharing_variables();
+  OMPRTL_NVPTX__kmpc_end_sharing_variables,
+  /// \brief Call to void __kmpc_get_shared_variables(void ***GlobalArgs)
+  OMPRTL_NVPTX__kmpc_get_shared_variables,
 };
 
 /// Pre(post)-action for different OpenMP constructs specialized for NVPTX.
 class NVPTXActionTy final : public PrePostActionTy {
   llvm::Value *EnterCallee;
   ArrayRef<llvm::Value *> EnterArgs;
   llvm::Value *ExitCallee;
   ArrayRef<llvm::Value *> ExitArgs;

 };
 
 enum NamedBarrier : unsigned {
   /// Synchronize on this barrier #ID using a named barrier primitive.
   /// Only the subset of active threads in a parallel region arrive at the
   /// barrier.
   NB_Parallel = 1,
 };
+
+/// Get the list of variables that can escape their declaration context.
+class CheckVarsEscapingDeclContext final
+    : public ConstStmtVisitor<CheckVarsEscapingDeclContext> {
+  CodeGenFunction &CGF;
+  llvm::SetVector<const ValueDecl *> EscapedDecls;
+  llvm::SmallPtrSet<const ValueDecl *, 4> IgnoredDecls;
+  bool AllEscaped = false;
+  RecordDecl *GlobalizedRD = nullptr;
+  llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *> MappedDeclsFields;
+
+  void markAsEscaped(const ValueDecl *VD) {
+    if (IgnoredDecls.count(VD) ||
+        (CGF.CapturedStmtInfo &&
+         CGF.CapturedStmtInfo->lookup(cast<VarDecl>(VD))))
+      return;
+    EscapedDecls.insert(VD);
+  }
+
+  void VisitValueDecl(const ValueDecl *VD) {
+    if (VD->getType()->isLValueReferenceType()) {
+      markAsEscaped(VD);
+      if (const auto *VarD = dyn_cast<VarDecl>(VD)) {
+        if (!isa<ParmVarDecl>(VarD) && VarD->hasInit()) {
+          const bool SavedAllEscaped = AllEscaped;
+          AllEscaped = true;
+          Visit(VarD->getInit());
+          AllEscaped = SavedAllEscaped;
+        }
+      }
+    }
+  }
+  void VisitOpenMPCapturedStmt(const CapturedStmt *S) {
+    if (!S)
+      return;
+    for (const auto &C : S->captures()) {
+      if (C.capturesVariable() && !C.capturesVariableByCopy()) {
+        const ValueDecl *VD = C.getCapturedVar();
+        markAsEscaped(VD);
+        if (isa<OMPCapturedExprDecl>(VD))
+          VisitValueDecl(VD);
+      }
+    }
+  }
+
+  typedef std::pair<CharUnits /*Align*/, const ValueDecl *> VarsDataTy;
+  static bool stable_sort_comparator(const VarsDataTy P1, const VarsDataTy P2) {
+    return P1.first > P2.first;
+  }
+
+  void buildRecordForGlobalizedVars() {
+    assert(!GlobalizedRD &&
+           "Record for globalized variables is built already.");
+    if (EscapedDecls.empty())
+      return;
+    ASTContext &C = CGF.getContext();
+    SmallVector<VarsDataTy, 4> GlobalizedVars;
+    for (const auto *D : EscapedDecls)
+      GlobalizedVars.emplace_back(C.getDeclAlign(D), D);
+    std::stable_sort(GlobalizedVars.begin(), GlobalizedVars.end(),
+                     stable_sort_comparator);
+    // Build struct _globalized_locals_ty {
+    //         /*  globalized vars  */
+    //       };
+    GlobalizedRD = C.buildImplicitRecord("_globalized_locals_ty");
+    GlobalizedRD->startDefinition();
+    for (const auto &Pair : GlobalizedVars) {
+      const ValueDecl *VD = Pair.second;
+      QualType Type = VD->getType();
+      if (Type->isLValueReferenceType())
+        Type = C.getPointerType(Type.getNonReferenceType());
+      else
+        Type = Type.getNonReferenceType();
+      SourceLocation Loc = VD->getLocation();
+      auto *Field = FieldDecl::Create(
+          C, GlobalizedRD, Loc, Loc, VD->getIdentifier(), Type,
+          C.getTrivialTypeSourceInfo(Type, SourceLocation()),
+          /*BW=*/nullptr, /*Mutable=*/false,
+          /*InitStyle=*/ICIS_NoInit);
+      Field->setAccess(AS_public);
+      GlobalizedRD->addDecl(Field);
+      if (VD->hasAttrs()) {
+        for (specific_attr_iterator<AlignedAttr> I(VD->getAttrs().begin()),
+             E(VD->getAttrs().end());
+             I != E; ++I)
+          Field->addAttr(*I);
+      }
+      MappedDeclsFields.try_emplace(VD, Field);
+    }
+    GlobalizedRD->completeDefinition();
+  }
+
+public:
+  CheckVarsEscapingDeclContext(CodeGenFunction &CGF,
+                               ArrayRef<const ValueDecl *> IgnoredDecls)
+      : CGF(CGF), IgnoredDecls(IgnoredDecls.begin(), IgnoredDecls.end()) {}
+  virtual ~CheckVarsEscapingDeclContext() = default;
+  void VisitDeclStmt(const DeclStmt *S) {
+    if (!S)
+      return;
+    for (const auto *D : S->decls())
+      if (const auto *VD = dyn_cast_or_null<ValueDecl>(D))
+        VisitValueDecl(VD);
+  }
+  void VisitOMPExecutableDirective(const OMPExecutableDirective *D) {
+    if (!D)
+      return;
+    if (D->hasAssociatedStmt()) {
+      if (const auto *S =
+              dyn_cast_or_null<CapturedStmt>(D->getAssociatedStmt()))
+        VisitOpenMPCapturedStmt(S);
+    }
+  }
+  void VisitCapturedStmt(const CapturedStmt *S) {
+    if (!S)
+      return;
+    for (const auto &C : S->captures()) {
+      if (C.capturesVariable() && !C.capturesVariableByCopy()) {
+        const ValueDecl *VD = C.getCapturedVar();
+        markAsEscaped(VD);
+        if (isa<OMPCapturedExprDecl>(VD))
+          VisitValueDecl(VD);
+      }
+    }
+  }
+  void VisitLambdaExpr(const LambdaExpr *E) {
+    if (!E)
+      return;
+    for (const auto &C : E->captures()) {
+      if (C.capturesVariable()) {
+        if (C.getCaptureKind() == LCK_ByRef) {
+          const ValueDecl *VD = C.getCapturedVar();
+          markAsEscaped(VD);
+          if (E->isInitCapture(&C) || isa<OMPCapturedExprDecl>(VD))
+            VisitValueDecl(VD);
+        }
+      }
+    }
+  }
+  void VisitBlockExpr(const BlockExpr *E) {
+    if (!E)
+      return;
+    for (const auto &C : E->getBlockDecl()->captures()) {
+      if (C.isByRef()) {
+        const VarDecl *VD = C.getVariable();
+        markAsEscaped(VD);
+        if (isa<OMPCapturedExprDecl>(VD) || VD->isInitCapture())
+          VisitValueDecl(VD);
+      }
+    }
+  }
+  void VisitCallExpr(const CallExpr *E) {
+    if (!E)
+      return;
+    for (const Expr *Arg : E->arguments()) {
+      if (!Arg)
+        continue;
+      if (Arg->isLValue()) {
+        const bool SavedAllEscaped = AllEscaped;
+        AllEscaped = true;
+        Visit(Arg);
+        AllEscaped = SavedAllEscaped;
+      } else
+        Visit(Arg);
+    }
+    Visit(E->getCallee());
+  }
+  void VisitDeclRefExpr(const DeclRefExpr *E) {
+    if (!E)
+      return;
+    const ValueDecl *VD = E->getDecl();
+    if (AllEscaped)
+      markAsEscaped(VD);
+    if (isa<OMPCapturedExprDecl>(VD))
+      VisitValueDecl(VD);
+    else if (const auto *VarD = dyn_cast<VarDecl>(VD))
+      if (VarD->isInitCapture())
+        VisitValueDecl(VD);
+  }
+  void VisitUnaryOperator(const UnaryOperator *E) {
+    if (!E)
+      return;
+    if (E->getOpcode() == UO_AddrOf) {
+      const bool SavedAllEscaped = AllEscaped;
+      AllEscaped = true;
+      Visit(E->getSubExpr());
+      AllEscaped = SavedAllEscaped;
+    } else
+      Visit(E->getSubExpr());
+  }
+  void VisitImplicitCastExpr(const ImplicitCastExpr *E) {
+    if (!E)
+      return;
+    if (E->getCastKind() == CK_ArrayToPointerDecay) {
+      const bool SavedAllEscaped = AllEscaped;
+      AllEscaped = true;
+      Visit(E->getSubExpr());
+      AllEscaped = SavedAllEscaped;
+    } else
+      Visit(E->getSubExpr());
+  }
+  void VisitExpr(const Expr *E) {
+    if (!E)
+      return;
+    bool SavedAllEscaped = AllEscaped;
+    if (!E->isLValue())
+      AllEscaped = false;
+    for (const auto *Child : E->children())
+      if (Child)
+        Visit(Child);
+    AllEscaped = SavedAllEscaped;
+  }
+  void VisitStmt(const Stmt *S) {
+    if (!S)
+      return;
+    for (const auto *Child : S->children())
+      if (Child)
+        Visit(Child);
+  }
+
+  const RecordDecl *getGlobalizedRecord() {
+    if (!GlobalizedRD)
+      buildRecordForGlobalizedVars();
+    return GlobalizedRD;
+  }
+
+  const FieldDecl *getFieldForGlobalizedVar(const ValueDecl *VD) const {
+    assert(GlobalizedRD &&
+           "Record for globalized variables must be generated already.");
+    auto I = MappedDeclsFields.find(VD);
+    if (I == MappedDeclsFields.end())
+      return nullptr;
+    return I->getSecond();
+  }
+
+  ArrayRef<const ValueDecl *> getEscapedDecls() const {
+    return EscapedDecls.getArrayRef();
+  }
+};
 } // anonymous namespace
 
 /// Get the GPU warp size.
 static llvm::Value *getNVPTXWarpSize(CodeGenFunction &CGF) {
   return CGF.EmitRuntimeCall(
       llvm::Intrinsic::getDeclaration(
           &CGF.CGM.getModule(), llvm::Intrinsic::nvvm_read_ptx_sreg_warpsize),
       "nvptx_warp_size");
void CGOpenMPRuntimeNVPTX::emitGenericKernel(const OMPExecutableDirective &D,
                                              llvm::Constant *&OutlinedFnID,
                                              bool IsOffloadEntry,
                                              const RegionCodeGenTy &CodeGen) {
   ExecutionModeRAII ModeRAII(CurrentExecutionMode,
                              CGOpenMPRuntimeNVPTX::ExecutionMode::Generic);
   EntryFunctionState EST;
   WorkerFunctionState WST(CGM, D.getLocStart());
   Work.clear();
+  WrapperFunctionsMap.clear();
 
   // Emit target region as a standalone region.
   class NVPTXPrePostActionTy : public PrePostActionTy {
     CGOpenMPRuntimeNVPTX::EntryFunctionState &EST;
     CGOpenMPRuntimeNVPTX::WorkerFunctionState &WST;
 
   public:
     NVPTXPrePostActionTy(CGOpenMPRuntimeNVPTX::EntryFunctionState &EST,
void CGOpenMPRuntimeNVPTX::emitGenericEntryHeader(CodeGenFunction &CGF,
   CGF.EmitBranch(EST.ExitBB);
 
   CGF.EmitBlock(MasterCheckBB);
   auto *IsMaster =
       Bld.CreateICmpEQ(getNVPTXThreadID(CGF), getMasterThreadID(CGF));
   Bld.CreateCondBr(IsMaster, MasterBB, EST.ExitBB);
 
   CGF.EmitBlock(MasterBB);
+  // SEQUENTIAL (MASTER) REGION START
   // First action in sequential region:
   // Initialize the state of the OpenMP runtime library on the GPU.
   // TODO: Optimize runtime initialization and pass in correct value.
   llvm::Value *Args[] = {getThreadLimit(CGF),
                          Bld.getInt16(/*RequiresOMPRuntime=*/1)};
   CGF.EmitRuntimeCall(
       createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_kernel_init), Args);
+
+  // For data sharing, we need to initialize the stack.
+  CGF.EmitRuntimeCall(
+      createNVPTXRuntimeFunction(
+          OMPRTL_NVPTX__kmpc_data_sharing_init_stack));
+
+  const auto I = FunctionGlobalizedDecls.find(CGF.CurFn);
+  if (I == FunctionGlobalizedDecls.end())
+    return;
+  const RecordDecl *GlobalizedVarsRecord = I->getSecond().first;
+  QualType RecTy = CGM.getContext().getRecordType(GlobalizedVarsRecord);
+
+  // Recover pointer to this function's global record. The runtime will
+  // handle the specifics of the allocation of the memory.
+  // Use actual memory size of the record including the padding
+  // for alignment purposes.
+  unsigned Alignment =
+      CGM.getContext().getTypeAlignInChars(RecTy).getQuantity();
+  unsigned GlobalRecordSize =
+      CGM.getContext().getTypeSizeInChars(RecTy).getQuantity();
+  GlobalRecordSize = llvm::alignTo(GlobalRecordSize, Alignment);
+  // TODO: allow the usage of shared memory to be controlled by
+  // the user, for now, default to global.
+  llvm::Value *GlobalRecordSizeArg[] = {
+      llvm::ConstantInt::get(CGM.SizeTy, GlobalRecordSize),
+      CGF.Builder.getInt16(/*UseSharedMemory=*/0)};
+  llvm::Value *GlobalRecValue = CGF.EmitRuntimeCall(
+      createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_data_sharing_push_stack),
+      GlobalRecordSizeArg);
+  llvm::Value *GlobalRecCastAddr = Bld.CreatePointerBitCastOrAddrSpaceCast(
+      GlobalRecValue, CGF.ConvertTypeForMem(RecTy)->getPointerTo());
+  FunctionToGlobalRecPtr.try_emplace(CGF.CurFn, GlobalRecValue);
+
+  // Emit the "global alloca" which is a GEP from the global declaration record
+  // using the pointer returned by the runtime.
+  LValue Base = CGF.MakeNaturalAlignPointeeAddrLValue(GlobalRecCastAddr, RecTy);
+  auto &Res = I->getSecond().second;
+  for (auto &Rec : Res) {
+    const FieldDecl *FD = Rec.second.first;
+    LValue VarAddr = CGF.EmitLValueForField(Base, FD);
+    Rec.second.second = VarAddr.getAddress();
+  }
 }
 
 void CGOpenMPRuntimeNVPTX::emitGenericEntryFooter(CodeGenFunction &CGF,
                                                   EntryFunctionState &EST) {
+  const auto I = FunctionGlobalizedDecls.find(CGF.CurFn);
+  if (I != FunctionGlobalizedDecls.end()) {
+    if (!CGF.HaveInsertPoint())
+      return;
+    auto I = FunctionToGlobalRecPtr.find(CGF.CurFn);
+    if (I != FunctionToGlobalRecPtr.end()) {
+      llvm::Value *Args[] = {I->getSecond()};
+      CGF.EmitRuntimeCall(
+          createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_data_sharing_pop_stack),
+          Args);
+    }
+  }
+
   if (!EST.ExitBB)
     EST.ExitBB = CGF.createBasicBlock(".exit");
 
   llvm::BasicBlock *TerminateBB = CGF.createBasicBlock(".termination.notifier");
   CGF.EmitBranch(TerminateBB);
 
   CGF.EmitBlock(TerminateBB);
   // Signal termination condition.
for (auto *W : Work) {
 
     llvm::BasicBlock *ExecuteFNBB = CGF.createBasicBlock(".execute.fn");
     llvm::BasicBlock *CheckNextBB = CGF.createBasicBlock(".check.next");
     Bld.CreateCondBr(WorkFnMatch, ExecuteFNBB, CheckNextBB);
 
     // Execute this outlined function.
     CGF.EmitBlock(ExecuteFNBB);
 
-    // Insert call to work function.
-    // FIXME: Pass arguments to outlined function from master thread.
-    auto *Fn = cast<llvm::Function>(W);
-    Address ZeroAddr =
-        CGF.CreateDefaultAlignTempAlloca(CGF.Int32Ty, /*Name=*/".zero.addr");
-    CGF.InitTempAlloca(ZeroAddr, CGF.Builder.getInt32(/*C=*/0));
-    llvm::Value *FnArgs[] = {ZeroAddr.getPointer(), ZeroAddr.getPointer()};
-    emitCall(CGF, WST.Loc, Fn, FnArgs);
+    // Insert call to work function via shared wrapper. The shared
+    // wrapper takes two arguments:
+    //   - the parallelism level;
+    //   - the master thread ID;
+    emitOutlinedFunctionCall(CGF, WST.Loc, W,
+                             {Bld.getInt16(/*ParallelLevel=*/0),
+                              getMasterThreadID(CGF)});
 
     // Go to end of parallel region.
     CGF.EmitBranch(TerminateBB);
 
     CGF.EmitBlock(CheckNextBB);
   }
 
   // Signal end of parallel region.
case OMPRTL_NVPTX__kmpc_spmd_kernel_deinit: {
     llvm::FunctionType *FnTy =
         llvm::FunctionType::get(CGM.VoidTy, llvm::None, /*isVarArg*/ false);
     RTLFn = CGM.CreateRuntimeFunction(FnTy, "__kmpc_spmd_kernel_deinit");
     break;
   }
   case OMPRTL_NVPTX__kmpc_kernel_prepare_parallel: {
     /// Build void __kmpc_kernel_prepare_parallel(
     /// void *outlined_function, int16_t IsOMPRuntimeInitialized);
-    llvm::Type *TypeParams[] = {CGM.Int8PtrTy,
-                                CGM.Int16Ty};
+    llvm::Type *TypeParams[] = {CGM.Int8PtrTy, CGM.Int16Ty};
     llvm::FunctionType *FnTy =
         llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg*/ false);
     RTLFn = CGM.CreateRuntimeFunction(FnTy, "__kmpc_kernel_prepare_parallel");
     break;
   }
   case OMPRTL_NVPTX__kmpc_kernel_parallel: {
     /// Build bool __kmpc_kernel_parallel(void **outlined_function,
     /// int16_t IsOMPRuntimeInitialized);
case OMPRTL_NVPTX__kmpc_end_reduce_nowait: {
     // Build __kmpc_end_reduce_nowait(kmp_int32 global_tid);
     llvm::Type *TypeParams[] = {CGM.Int32Ty};
     llvm::FunctionType *FnTy =
         llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg=*/false);
     RTLFn = CGM.CreateRuntimeFunction(
         FnTy, /*Name=*/"__kmpc_nvptx_end_reduce_nowait");
     break;
   }
+  case OMPRTL_NVPTX__kmpc_data_sharing_init_stack: {
+    /// Build void __kmpc_data_sharing_init_stack();
+    llvm::FunctionType *FnTy =
+        llvm::FunctionType::get(CGM.VoidTy, llvm::None, /*isVarArg*/ false);
+    RTLFn = CGM.CreateRuntimeFunction(FnTy, "__kmpc_data_sharing_init_stack");
+    break;
+  }
+  case OMPRTL_NVPTX__kmpc_data_sharing_push_stack: {
+    // Build void *__kmpc_data_sharing_push_stack(size_t size,
+    // int16_t UseSharedMemory);
+    llvm::Type *TypeParams[] = {CGM.SizeTy, CGM.Int16Ty};
+    llvm::FunctionType *FnTy =
+        llvm::FunctionType::get(CGM.VoidPtrTy, TypeParams, /*isVarArg=*/false);
+    RTLFn = CGM.CreateRuntimeFunction(
+        FnTy, /*Name=*/"__kmpc_data_sharing_push_stack");
+    break;
+  }
+  case OMPRTL_NVPTX__kmpc_data_sharing_pop_stack: {
+    // Build void __kmpc_data_sharing_pop_stack(void *a);
+    llvm::Type *TypeParams[] = {CGM.VoidPtrTy};
+    llvm::FunctionType *FnTy =
+        llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg=*/false);
+    RTLFn = CGM.CreateRuntimeFunction(FnTy,
+                                      /*Name=*/"__kmpc_data_sharing_pop_stack");
+    break;
+  }
+  case OMPRTL_NVPTX__kmpc_begin_sharing_variables: {
+    /// Build void __kmpc_begin_sharing_variables(void ***args,
+    /// size_t n_args);
+    llvm::Type *TypeParams[] = {CGM.Int8PtrPtrTy->getPointerTo(), CGM.SizeTy};
+    llvm::FunctionType *FnTy =
+        llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg*/ false);
+    RTLFn = CGM.CreateRuntimeFunction(FnTy, "__kmpc_begin_sharing_variables");
+    break;
+  }
+  case OMPRTL_NVPTX__kmpc_end_sharing_variables: {
+    /// Build void __kmpc_end_sharing_variables();
+    llvm::FunctionType *FnTy =
+        llvm::FunctionType::get(CGM.VoidTy, llvm::None, /*isVarArg*/ false);
+    RTLFn = CGM.CreateRuntimeFunction(FnTy, "__kmpc_end_sharing_variables");
+    break;
+  }
+  case OMPRTL_NVPTX__kmpc_get_shared_variables: {
+    /// Build void __kmpc_get_shared_variables(void ***GlobalArgs);
+    llvm::Type *TypeParams[] = {CGM.Int8PtrPtrTy->getPointerTo()};
+    llvm::FunctionType *FnTy =
+        llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg*/ false);
+    RTLFn = CGM.CreateRuntimeFunction(FnTy, "__kmpc_get_shared_variables");
+    break;
+  }
   }
   return RTLFn;
 }
 
 void CGOpenMPRuntimeNVPTX::createOffloadEntry(llvm::Constant *ID,
                                               llvm::Constant *Addr,
                                               uint64_t Size, int32_t) {
   auto *F = dyn_cast<llvm::Function>(Addr);

 void CGOpenMPRuntimeNVPTX::emitNumTeamsClause(CodeGenFunction &CGF,
                                               const Expr *NumTeams,
                                               const Expr *ThreadLimit,
                                               SourceLocation Loc) {}
 
 llvm::Value *CGOpenMPRuntimeNVPTX::emitParallelOutlinedFunction(
     const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
     OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen) {
-  return CGOpenMPRuntime::emitParallelOutlinedFunction(D, ThreadIDVar,
-      InnermostKind, CodeGen);
+  auto *OutlinedFun =
+      cast<llvm::Function>(CGOpenMPRuntime::emitParallelOutlinedFunction(
+          D, ThreadIDVar, InnermostKind, CodeGen));
+  if (!isInSpmdExecutionMode()) {
+    llvm::Function *WrapperFun =
+        createParallelDataSharingWrapper(OutlinedFun, D);
+    WrapperFunctionsMap[OutlinedFun] = WrapperFun;
+  }
+
+  return OutlinedFun;
 }
 
 llvm::Value *CGOpenMPRuntimeNVPTX::emitTeamsOutlinedFunction(
     const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
     OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen) {
 
   llvm::Value *OutlinedFunVal = CGOpenMPRuntime::emitTeamsOutlinedFunction(
       D, ThreadIDVar, InnermostKind, CodeGen);
void CGOpenMPRuntimeNVPTX::emitParallelCall(
   else
     emitGenericParallelCall(CGF, Loc, OutlinedFn, CapturedVars, IfCond);
 }
 
 void CGOpenMPRuntimeNVPTX::emitGenericParallelCall(
     CodeGenFunction &CGF, SourceLocation Loc, llvm::Value *OutlinedFn,
     ArrayRef<llvm::Value *> CapturedVars, const Expr *IfCond) {
   llvm::Function *Fn = cast<llvm::Function>(OutlinedFn);
+  llvm::Function *WFn = WrapperFunctionsMap[Fn];
+
+  assert(WFn && "Wrapper function does not exist!");
+
+  // Force inline this outlined function at its call site.
+  Fn->setLinkage(llvm::GlobalValue::InternalLinkage);
 
-  auto &&L0ParallelGen = [this, Fn](CodeGenFunction &CGF, PrePostActionTy &) {
+  auto &&L0ParallelGen = [this, WFn, &CapturedVars](CodeGenFunction &CGF,
+                                                    PrePostActionTy &) {
     CGBuilderTy &Bld = CGF.Builder;
 
-    // TODO: Optimize runtime initialization.
-    llvm::Value *Args[] = {Bld.CreateBitOrPointerCast(Fn, CGM.Int8PtrTy),
-                           /*RequiresOMPRuntime=*/Bld.getInt16(1)};
-    CGF.EmitRuntimeCall(
-        createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_kernel_prepare_parallel),
+    llvm::Value *ID = Bld.CreateBitOrPointerCast(WFn, CGM.Int8PtrTy);
+
+    // Prepare for parallel region. Indicate the outlined function.
+    llvm::Value *Args[] = {ID, /*RequiresOMPRuntime=*/Bld.getInt16(1)};
+    CGF.EmitRuntimeCall(createNVPTXRuntimeFunction(
+                            OMPRTL_NVPTX__kmpc_kernel_prepare_parallel),
                         Args);
 
+    // Create a private scope that will globalize the arguments
+    // passed from the outside of the target region.
+    CodeGenFunction::OMPPrivateScope PrivateArgScope(CGF);
+
+    // There's somehting to share.
+    if (!CapturedVars.empty()) {
+      // Prepare for parallel region. Indicate the outlined function.
+      Address SharedArgs =
+          CGF.CreateDefaultAlignTempAlloca(CGF.VoidPtrPtrTy, "shared_arg_refs");
+      llvm::Value *SharedArgsPtr = SharedArgs.getPointer();
+
+      llvm::Value *DataSharingArgs[] = {
+          SharedArgsPtr,
+          llvm::ConstantInt::get(CGM.SizeTy, CapturedVars.size())};
+      CGF.EmitRuntimeCall(createNVPTXRuntimeFunction(
+                              OMPRTL_NVPTX__kmpc_begin_sharing_variables),
+                          DataSharingArgs);
+
+      // Store variable address in a list of references to pass to workers.
+      unsigned Idx = 0;
+      ASTContext &Ctx = CGF.getContext();
+      Address SharedArgListAddress = CGF.EmitLoadOfPointer(SharedArgs,
+          Ctx.getPointerType(Ctx.getPointerType(Ctx.VoidPtrTy))
+              .castAs<PointerType>());
+      for (llvm::Value *V : CapturedVars) {
+        Address Dst = Bld.CreateConstInBoundsGEP(
+            SharedArgListAddress, Idx, CGF.getPointerSize());
+        llvm::Value *PtrV = Bld.CreateBitCast(V, CGF.VoidPtrTy);
+        CGF.EmitStoreOfScalar(PtrV, Dst, /*Volatile=*/false,
+                              Ctx.getPointerType(Ctx.VoidPtrTy));
+        Idx++;
+      }
+    }
+
     // Activate workers. This barrier is used by the master to signal
     // work for the workers.
     syncCTAThreads(CGF);
 
     // OpenMP [2.5, Parallel Construct, p.49]
     // There is an implied barrier at the end of a parallel region. After the
     // end of a parallel region, only the master thread of the team resumes
     // execution of the enclosing task region.
     //
     // The master waits at this barrier until all workers are done.
     syncCTAThreads(CGF);
 
+    if (!CapturedVars.empty())
+      CGF.EmitRuntimeCall(
+          createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_end_sharing_variables));
+
     // Remember for post-processing in worker loop.
-    Work.emplace_back(Fn);
+    Work.emplace_back(WFn);
   };
 
   auto *RTLoc = emitUpdateLocation(CGF, Loc);
   auto *ThreadID = getThreadID(CGF, Loc);
   llvm::Value *Args[] = {RTLoc, ThreadID};
 
   auto &&SeqGen = [this, Fn, &CapturedVars, &Args, Loc](CodeGenFunction &CGF,
                                                         PrePostActionTy &) {
for (unsigned I = 0, E = Args.size(); I < E; ++I) {
     llvm::Value *TargetArg = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
         NativeArg, NativeArg->getType()->getPointerElementType()->getPointerTo(
                        /*AddrSpace=*/0));
     TargetArgs.emplace_back(
         CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(TargetArg, TargetType));
   }
   CGOpenMPRuntime::emitOutlinedFunctionCall(CGF, Loc, OutlinedFn, TargetArgs);
 }
+
+/// Emit function which wraps the outline parallel region
+/// and controls the arguments which are passed to this function.
+/// The wrapper ensures that the outlined function is called
+/// with the correct arguments when data is shared.
+llvm::Function *CGOpenMPRuntimeNVPTX::createParallelDataSharingWrapper(
+    llvm::Function *OutlinedParallelFn, const OMPExecutableDirective &D) {
+  ASTContext &Ctx = CGM.getContext();
+  const auto &CS = *D.getCapturedStmt(OMPD_parallel);
+
+  // Create a function that takes as argument the source thread.
+  FunctionArgList WrapperArgs;
+  QualType Int16QTy =
+      Ctx.getIntTypeForBitwidth(/*DestWidth=*/16, /*Signed=*/false);
+  QualType Int32QTy =
+      Ctx.getIntTypeForBitwidth(/*DestWidth=*/32, /*Signed=*/false);
+  ImplicitParamDecl ParallelLevelArg(Ctx, /*DC=*/nullptr, D.getLocStart(),
+                                     /*Id=*/nullptr, Int16QTy,
+                                     ImplicitParamDecl::Other);
+  ImplicitParamDecl WrapperArg(Ctx, /*DC=*/nullptr, D.getLocStart(),
+                               /*Id=*/nullptr, Int32QTy,
+                               ImplicitParamDecl::Other);
+  WrapperArgs.emplace_back(&ParallelLevelArg);
+  WrapperArgs.emplace_back(&WrapperArg);
+
+  auto &CGFI =
+      CGM.getTypes().arrangeBuiltinFunctionDeclaration(Ctx.VoidTy, WrapperArgs);
+
+  auto *Fn = llvm::Function::Create(
+      CGM.getTypes().GetFunctionType(CGFI), llvm::GlobalValue::InternalLinkage,
+      OutlinedParallelFn->getName() + "_wrapper", &CGM.getModule());
+  CGM.SetInternalFunctionAttributes(/*D=*/GlobalDecl(), Fn, CGFI);
+  Fn->setLinkage(llvm::GlobalValue::InternalLinkage);
+
+  CodeGenFunction CGF(CGM, /*suppressNewContext=*/true);
+  CGF.StartFunction(GlobalDecl(), Ctx.VoidTy, Fn, CGFI, WrapperArgs,
+                    D.getLocStart(), D.getLocStart());
+
+  const auto *RD = CS.getCapturedRecordDecl();
+  auto CurField = RD->field_begin();
+
+  // Get the array of arguments.
+  SmallVector<llvm::Value *, 8> Args;
+
+  // TODO: suppport SIMD and pass actual values
+  Args.emplace_back(
+      llvm::ConstantPointerNull::get(CGM.Int32Ty->getPointerTo()));
+  Args.emplace_back(
+      llvm::ConstantPointerNull::get(CGM.Int32Ty->getPointerTo()));
+
+  CGBuilderTy &Bld = CGF.Builder;
+  auto CI = CS.capture_begin();
+
+  // Use global memory for data sharing.
+  // Handle passing of global args to workers.
+  Address GlobalArgs =
+      CGF.CreateDefaultAlignTempAlloca(CGF.VoidPtrPtrTy, "global_args");
+  llvm::Value *GlobalArgsPtr = GlobalArgs.getPointer();
+  llvm::Value *DataSharingArgs[] = {GlobalArgsPtr};
+  CGF.EmitRuntimeCall(
+      createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_get_shared_variables),
+      DataSharingArgs);
+
+  // Retrieve the shared variables from the list of references returned
+  // by the runtime. Pass the variables to the outlined function.
+  if (CS.capture_size() > 0) {
+    ASTContext &CGFContext = CGF.getContext();
+    Address SharedArgListAddress = CGF.EmitLoadOfPointer(GlobalArgs,
+        CGFContext
+            .getPointerType(CGFContext.getPointerType(CGFContext.VoidPtrTy))
+            .castAs<PointerType>());
+    for (unsigned I = 0, E = CS.capture_size(); I < E; ++I, ++CI, ++CurField) {
+      QualType ElemTy = CurField->getType();
+      Address Src = Bld.CreateConstInBoundsGEP(
+          SharedArgListAddress, I, CGF.getPointerSize());
+      Address TypedAddress = Bld.CreateBitCast(
+          Src, CGF.ConvertTypeForMem(CGFContext.getPointerType(ElemTy)));
+      llvm::Value *Arg = CGF.EmitLoadOfScalar(TypedAddress,
+                                              /*Volatile=*/false,
+                                              CGFContext.getPointerType(ElemTy),
+                                              CI->getLocation());
+      Args.emplace_back(Arg);
+    }
+  }
+
+  emitOutlinedFunctionCall(CGF, D.getLocStart(), OutlinedParallelFn, Args);
+  CGF.FinishFunction();
+  return Fn;
+}
+
+void CGOpenMPRuntimeNVPTX::emitFunctionProlog(CodeGenFunction &CGF,
+                                              const Decl *D) {
+  assert(D && "Expected function or captured|block decl.");
+  assert(FunctionGlobalizedDecls.count(CGF.CurFn) == 0 &&
+         "Function is registered already.");
+  SmallVector<const ValueDecl *, 4> IgnoredDecls;
+  const Stmt *Body = nullptr;
+  if (const auto *FD = dyn_cast<FunctionDecl>(D)) {
+    Body = FD->getBody();
+  } else if (const auto *BD = dyn_cast<BlockDecl>(D)) {
+    Body = BD->getBody();
+  } else if (const auto *CD = dyn_cast<CapturedDecl>(D)) {
+    Body = CD->getBody();
+    if (CGF.CapturedStmtInfo->getKind() == CR_OpenMP) {
+      if (const auto *CS = dyn_cast<CapturedStmt>(Body)) {
+        IgnoredDecls.reserve(CS->capture_size());
+        for (const auto &Capture : CS->captures())
+          if (Capture.capturesVariable())
+            IgnoredDecls.emplace_back(Capture.getCapturedVar());
+      }
+    }
+  }
+  if (!Body)
+    return;
+  CheckVarsEscapingDeclContext VarChecker(CGF, IgnoredDecls);
+  VarChecker.Visit(Body);
+  const RecordDecl *GlobalizedVarsRecord = VarChecker.getGlobalizedRecord();
+  if (!GlobalizedVarsRecord)
+    return;
+  auto &Res =
+      FunctionGlobalizedDecls
+          .try_emplace(CGF.CurFn, GlobalizedVarsRecord, DeclToAddrMapTy())
+          .first->getSecond()
+          .second;
+  for (const ValueDecl *VD : VarChecker.getEscapedDecls()) {
+    const FieldDecl *FD = VarChecker.getFieldForGlobalizedVar(VD);
+    Res.insert(std::make_pair(VD, std::make_pair(FD, Address::invalid())));
+  }
+}
+
+Address CGOpenMPRuntimeNVPTX::getAddressOfLocalVariable(CodeGenFunction &CGF,
+                                                        const VarDecl *VD) {
+  auto I = FunctionGlobalizedDecls.find(CGF.CurFn);
+  if (I == FunctionGlobalizedDecls.end())
+    return Address::invalid();
+  auto VDI = I->getSecond().second.find(VD);
+  if (VDI == I->getSecond().second.end())
+    return Address::invalid();
+  return VDI->second.second;
+}
+
+void CGOpenMPRuntimeNVPTX::functionFinished(CodeGenFunction &CGF) {
+  FunctionToGlobalRecPtr.erase(CGF.CurFn);
+  FunctionGlobalizedDecls.erase(CGF.CurFn);
+  CGOpenMPRuntime::functionFinished(CGF);
+}

cfe/trunk/lib/CodeGen/CGStmtOpenMP.cpp

CodeGenFunction::GenerateOpenMPCapturedStmtFunction(const CapturedStmt &S) {
   FinishFunction(CD->getBodyRBrace());
   if (!NeedWrapperFunction)
     return F;
 
   FunctionOptions WrapperFO(&S, /*UIntPtrCastRequired=*/true,
                             /*RegisterCastedArgsOnly=*/true,
                             CapturedStmtInfo->getHelperName());
   CodeGenFunction WrapperCGF(CGM, /*suppressNewContext=*/true);
+  WrapperCGF.CapturedStmtInfo = CapturedStmtInfo;
   Args.clear();
   LocalAddrs.clear();
   VLASizes.clear();
   llvm::Function *WrapperF =
       emitOutlinedFunctionPrologue(WrapperCGF, Args, LocalAddrs, VLASizes,
                                    WrapperCGF.CXXThisValue, WrapperFO);
   llvm::SmallVector<llvm::Value *, 4> CallArgs;
   for (const auto *Arg : Args) {

cfe/trunk/lib/CodeGen/CodeGenFunction.cpp

if (getLangOpts().ObjCAutoRefCount &&
         !CurFnInfo->isReturnsRetained() &&
         RetTy->isObjCRetainableType())
       AutoreleaseResult = true;
   }
 
   EmitStartEHSpec(CurCodeDecl);
 
   PrologueCleanupDepth = EHStack.stable_begin();
+
+  // Emit OpenMP specific initialization of the device functions.
+  if (getLangOpts().OpenMP && CurCodeDecl)
+    CGM.getOpenMPRuntime().emitFunctionProlog(*this, CurCodeDecl);
+
   EmitFunctionProlog(*CurFnInfo, CurFn, Args);
 
   if (D && isa<CXXMethodDecl>(D) && cast<CXXMethodDecl>(D)->isInstance()) {
     CGM.getCXXABI().EmitInstanceFunctionProlog(*this);
     const CXXMethodDecl *MD = cast<CXXMethodDecl>(D);
     if (MD->getParent()->isLambda() &&
         MD->getOverloadedOperator() == OO_Call) {
       // We're in a lambda; figure out the captures.

cfe/trunk/test/OpenMP/nvptx_data_sharing.cpp

+// Test device global memory data sharing codegen.
+///==========================================================================///
+
+// RUN: %clang_cc1 -verify -fopenmp -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 -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 CK1
+
+// expected-no-diagnostics
+
+#ifndef HEADER
+#define HEADER
+
+void test_ds(){
+  #pragma omp target
+  {
+    int a = 10;
+    #pragma omp parallel
+    {
+      a = 1000;
+    }
+    int b = 100;
+    #pragma omp parallel
+    {
+      b = a + 10000;
+    }
+  }
+}
+
+/// ========= In the kernel function ========= ///
+
+// CK1: {{.*}}define void @__omp_offloading{{.*}}test_ds{{.*}}()
+// CK1: [[SHAREDARGS1:%.+]] = alloca i8**
+// CK1: [[SHAREDARGS2:%.+]] = alloca i8**
+// CK1: call void @__kmpc_kernel_init
+// CK1: call void @__kmpc_data_sharing_init_stack
+// CK1: [[GLOBALSTACK:%.+]] = call i8* @__kmpc_data_sharing_push_stack(i64 8, i16 0)
+// CK1: [[GLOBALSTACK2:%.+]] = bitcast i8* [[GLOBALSTACK]] to %struct._globalized_locals_ty*
+// CK1: [[A:%.+]] = getelementptr inbounds %struct._globalized_locals_ty, %struct._globalized_locals_ty* [[GLOBALSTACK2]], i32 0, i32 0
+// CK1: [[B:%.+]] = getelementptr inbounds %struct._globalized_locals_ty, %struct._globalized_locals_ty* [[GLOBALSTACK2]], i32 0, i32 1
+// CK1: store i32 10, i32* [[A]]
+// CK1: call void @__kmpc_kernel_prepare_parallel({{.*}}, i16 1)
+// CK1: call void @__kmpc_begin_sharing_variables(i8*** [[SHAREDARGS1]], i64 1)
+// CK1: [[SHARGSTMP1:%.+]] = load i8**, i8*** [[SHAREDARGS1]]
+// CK1: [[SHARGSTMP2:%.+]] = getelementptr inbounds i8*, i8** [[SHARGSTMP1]], i64 0
+// CK1: [[SHAREDVAR:%.+]] = bitcast i32* [[A]] to i8*
+// CK1: store i8* [[SHAREDVAR]], i8** [[SHARGSTMP2]]
+// CK1: call void @llvm.nvvm.barrier0()
+// CK1: call void @llvm.nvvm.barrier0()
+// CK1: call void @__kmpc_end_sharing_variables()
+// CK1: store i32 100, i32* [[B]]
+// CK1: call void @__kmpc_kernel_prepare_parallel({{.*}}, i16 1)
+// CK1: call void @__kmpc_begin_sharing_variables(i8*** [[SHAREDARGS2]], i64 2)
+// CK1: [[SHARGSTMP3:%.+]] = load i8**, i8*** [[SHAREDARGS2]]
+// CK1: [[SHARGSTMP4:%.+]] = getelementptr inbounds i8*, i8** [[SHARGSTMP3]], i64 0
+// CK1: [[SHAREDVAR1:%.+]] = bitcast i32* [[B]] to i8*
+// CK1: store i8* [[SHAREDVAR1]], i8** [[SHARGSTMP4]]
+// CK1: [[SHARGSTMP12:%.+]] = getelementptr inbounds i8*, i8** [[SHARGSTMP3]], i64 1
+// CK1: [[SHAREDVAR2:%.+]] = bitcast i32* [[A]] to i8*
+// CK1: store i8* [[SHAREDVAR2]], i8** [[SHARGSTMP12]]
+// CK1: call void @llvm.nvvm.barrier0()
+// CK1: call void @llvm.nvvm.barrier0()
+// CK1: call void @__kmpc_end_sharing_variables()
+// CK1: call void @__kmpc_data_sharing_pop_stack(i8* [[GLOBALSTACK]])
+// CK1: call void @__kmpc_kernel_deinit(i16 1)
+
+/// ========= In the data sharing wrapper function ========= ///
+
+// CK1: {{.*}}define internal void @__omp_outlined{{.*}}wrapper({{.*}})
+// CK1: [[SHAREDARGS4:%.+]] = alloca i8**
+// CK1: call void @__kmpc_get_shared_variables(i8*** [[SHAREDARGS4]])
+// CK1: [[SHARGSTMP13:%.+]] = load i8**, i8*** [[SHAREDARGS4]]
+// CK1: [[SHARGSTMP14:%.+]] = getelementptr inbounds i8*, i8** [[SHARGSTMP13]], i64 0
+// CK1: [[SHARGSTMP15:%.+]] = bitcast i8** [[SHARGSTMP14]] to i32**
+// CK1: [[SHARGSTMP16:%.+]] = load i32*, i32** [[SHARGSTMP15]]
+// CK1: call void @__omp_outlined__{{.*}}({{.*}}, i32* [[SHARGSTMP16]])
+
+/// ========= In the data sharing wrapper function ========= ///
+
+// CK1: {{.*}}define internal void @__omp_outlined{{.*}}wrapper({{.*}})
+// CK1: [[SHAREDARGS3:%.+]] = alloca i8**
+// CK1: call void @__kmpc_get_shared_variables(i8*** [[SHAREDARGS3]])
+// CK1: [[SHARGSTMP5:%.+]] = load i8**, i8*** [[SHAREDARGS3]]
+// CK1: [[SHARGSTMP6:%.+]] = getelementptr inbounds i8*, i8** [[SHARGSTMP5]], i64 0
+// CK1: [[SHARGSTMP7:%.+]] = bitcast i8** [[SHARGSTMP6]] to i32**
+// CK1: [[SHARGSTMP8:%.+]] = load i32*, i32** [[SHARGSTMP7]]
+// CK1: [[SHARGSTMP9:%.+]] = getelementptr inbounds i8*, i8** [[SHARGSTMP5]], i64 1
+// CK1: [[SHARGSTMP10:%.+]] = bitcast i8** [[SHARGSTMP9]] to i32**
+// CK1: [[SHARGSTMP11:%.+]] = load i32*, i32** [[SHARGSTMP10]]
+// CK1: call void @__omp_outlined__{{.*}}({{.*}}, i32* [[SHARGSTMP8]], i32* [[SHARGSTMP11]])
+
+#endif
+

cfe/trunk/test/OpenMP/nvptx_parallel_codegen.cpp

int bar(int n){
 
   a += ftemplate<int>(n);
 
   return a;
 }
 
   // CHECK-NOT: define {{.*}}void {{@__omp_offloading_.+template.+l17}}_worker()
 
-
-
-
-
-
-  // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l26}}_worker()
-  // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
-  // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
-  // CHECK: store i8* null, i8** [[OMP_WORK_FN]],
-  // CHECK: store i8 0, i8* [[OMP_EXEC_STATUS]],
-  // CHECK: br label {{%?}}[[AWAIT_WORK:.+]]
-  //
-  // CHECK: [[AWAIT_WORK]]
-  // CHECK: call void @llvm.nvvm.barrier0()
-  // CHECK: [[KPR:%.+]] = call i1 @__kmpc_kernel_parallel(i8** [[OMP_WORK_FN]]
-  // CHECK: [[KPRB:%.+]] = zext i1 [[KPR]] to i8
-  // store i8 [[KPRB]], i8* [[OMP_EXEC_STATUS]], align 1
-  // CHECK: [[WORK:%.+]] = load i8*, i8** [[OMP_WORK_FN]],
-  // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i8* [[WORK]], null
-  // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]]
-  //
-  // CHECK: [[SEL_WORKERS]]
-  // CHECK: [[ST:%.+]] = load i8, i8* [[OMP_EXEC_STATUS]]
-  // CHECK: [[IS_ACTIVE:%.+]] = icmp ne i8 [[ST]], 0
-  // CHECK: br i1 [[IS_ACTIVE]], label {{%?}}[[EXEC_PARALLEL:.+]], label {{%?}}[[BAR_PARALLEL:.+]]
-  //
-  // CHECK: [[EXEC_PARALLEL]]
-  // CHECK: [[WF1:%.+]] = load i8*, i8** [[OMP_WORK_FN]],
-  // CHECK: [[WM1:%.+]] = icmp eq i8* [[WF1]], bitcast (void (i32*, i32*)* [[PARALLEL_FN1:@.+]] to i8*)
-  // CHECK: br i1 [[WM1]], label {{%?}}[[EXEC_PFN1:.+]], label {{%?}}[[CHECK_NEXT1:.+]]
-  //
-  // CHECK: [[EXEC_PFN1]]
-  // CHECK: call void [[PARALLEL_FN1]](
-  // CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
-  //
-  // CHECK: [[CHECK_NEXT1]]
-  // CHECK: [[WF2:%.+]] = load i8*, i8** [[OMP_WORK_FN]],
-  // CHECK: [[WM2:%.+]] = icmp eq i8* [[WF2]], bitcast (void (i32*, i32*)* [[PARALLEL_FN2:@.+]] to i8*)
-  // CHECK: br i1 [[WM2]], label {{%?}}[[EXEC_PFN2:.+]], label {{%?}}[[CHECK_NEXT2:.+]]
-  //
-  // CHECK: [[EXEC_PFN2]]
-  // CHECK: call void [[PARALLEL_FN2]](
-  // CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
-  //
-  // CHECK: [[CHECK_NEXT2]]
-  // CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
-  //
-  // CHECK: [[TERM_PARALLEL]]
-  // CHECK: call void @__kmpc_kernel_end_parallel()
-  // CHECK: br label {{%?}}[[BAR_PARALLEL]]
-  //
-  // CHECK: [[BAR_PARALLEL]]
-  // CHECK: call void @llvm.nvvm.barrier0()
-  // CHECK: br label {{%?}}[[AWAIT_WORK]]
-  //
-  // CHECK: [[EXIT]]
-  // CHECK: ret void
+// CHECK: [[WM1:%.+]] = icmp eq i8* [[WF1]], bitcast (void (i16, i32)* [[PARALLEL_FN1:@.+]]_wrapper to i8*)
+// CHECK: call void [[PARALLEL_FN1]]_wrapper(
+// CHECK: [[WM2:%.+]] = icmp eq i8* [[WF2]], bitcast (void (i16, i32)* [[PARALLEL_FN2:@.+]]_wrapper to i8*)
+// CHECK: call void [[PARALLEL_FN2]]_wrapper(
 
-  // CHECK: define {{.*}}void [[T6:@__omp_offloading_.+template.+l26]](i[[SZ:32|64]]
-  // Create local storage for each capture.
-  // CHECK:  [[LOCAL_A:%.+]] = alloca i[[SZ]],
-  // CHECK-DAG:  store i[[SZ]] [[ARG_A:%.+]], i[[SZ]]* [[LOCAL_A]]
-  // Store captures in the context.
-  // CHECK-64-DAG:[[REF_A:%.+]] = bitcast i[[SZ]]* [[LOCAL_A]] to i32*
-  //
-  // CHECK-DAG: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
-  // CHECK-DAG: [[NTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
-  // CHECK-DAG: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
-  // CHECK-DAG: [[TH_LIMIT:%.+]] = sub i32 [[NTH]], [[WS]]
-  // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[TH_LIMIT]]
-  // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[CHECK_MASTER:.+]]
-  //
-  // CHECK: [[WORKER]]
-  // CHECK: {{call|invoke}} void [[T6]]_worker()
-  // CHECK: br label {{%?}}[[EXIT:.+]]
-  //
-  // CHECK: [[CHECK_MASTER]]
-  // CHECK-DAG: [[CMTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
-  // CHECK-DAG: [[CMNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
-  // CHECK-DAG: [[CMWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
-  // CHECK: [[IS_MASTER:%.+]] = icmp eq i32 [[CMTID]],
-  // CHECK: br i1 [[IS_MASTER]], label {{%?}}[[MASTER:.+]], label {{%?}}[[EXIT]]
-  //
-  // CHECK: [[MASTER]]
-  // CHECK-DAG: [[MNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
-  // CHECK-DAG: [[MWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
-  // CHECK: [[MTMP1:%.+]] = sub i32 [[MNTH]], [[MWS]]
-  // CHECK: call void @__kmpc_kernel_init(i32 [[MTMP1]]
-  // CHECK: call void @__kmpc_kernel_prepare_parallel(i8* bitcast (void (i32*, i32*)* [[PARALLEL_FN1]] to i8*),
-  // CHECK: call void @llvm.nvvm.barrier0()
-  // CHECK: call void @llvm.nvvm.barrier0()
-  // CHECK: call void @__kmpc_serialized_parallel(
-  // CHECK: {{call|invoke}} void [[PARALLEL_FN3:@.+]](
-  // CHECK: call void @__kmpc_end_serialized_parallel(
-  // CHECK: call void @__kmpc_kernel_prepare_parallel(i8* bitcast (void (i32*, i32*)* [[PARALLEL_FN2]] to i8*),
-  // CHECK: call void @llvm.nvvm.barrier0()
-  // CHECK: call void @llvm.nvvm.barrier0()
-  // CHECK-64-DAG: load i32, i32* [[REF_A]]
-  // CHECK-32-DAG: load i32, i32* [[LOCAL_A]]
-  // CHECK: br label {{%?}}[[TERMINATE:.+]]
-  //
-  // CHECK: [[TERMINATE]]
-  // CHECK: call void @__kmpc_kernel_deinit(
-  // CHECK: call void @llvm.nvvm.barrier0()
-  // CHECK: br label {{%?}}[[EXIT]]
-  //
-  // CHECK: [[EXIT]]
-  // CHECK: ret void
+// CHECK: call void @__kmpc_kernel_prepare_parallel(i8* bitcast (void (i16, i32)* [[PARALLEL_FN1]]_wrapper to i8*),
+// CHECK: call void @__kmpc_kernel_prepare_parallel(i8* bitcast (void (i16, i32)* [[PARALLEL_FN2]]_wrapper to i8*),
 
-  // CHECK-DAG: define internal void [[PARALLEL_FN1]](
-  // CHECK: [[A:%.+]] = alloca i[[SZ:32|64]],
-  // CHECK: store i[[SZ]] 42, i[[SZ]]* %a,
-  // CHECK: ret void
 
-  // CHECK-DAG: define internal void [[PARALLEL_FN3]](
-  // CHECK: [[A:%.+]] = alloca i[[SZ:32|64]],
-  // CHECK: store i[[SZ]] 43, i[[SZ]]* %a,
-  // CHECK: ret void
 
-  // CHECK-DAG: define internal void [[PARALLEL_FN2]](
-  // CHECK: [[A:%.+]] = alloca i[[SZ:32|64]],
-  // CHECK: store i[[SZ]] 44, i[[SZ]]* %a,
-  // CHECK: ret void
 
-
-
-
-
-
-
-  // CHECK-LABEL: define {{.*}}void {{@__omp_offloading_.+template.+l43}}_worker()
-  // CHECK-DAG: [[OMP_EXEC_STATUS:%.+]] = alloca i8,
-  // CHECK-DAG: [[OMP_WORK_FN:%.+]] = alloca i8*,
-  // CHECK: store i8* null, i8** [[OMP_WORK_FN]],
-  // CHECK: store i8 0, i8* [[OMP_EXEC_STATUS]],
-  // CHECK: br label {{%?}}[[AWAIT_WORK:.+]]
-  //
-  // CHECK: [[AWAIT_WORK]]
-  // CHECK: call void @llvm.nvvm.barrier0()
-  // CHECK: [[KPR:%.+]] = call i1 @__kmpc_kernel_parallel(i8** [[OMP_WORK_FN]]
-  // CHECK: [[KPRB:%.+]] = zext i1 [[KPR]] to i8
-  // store i8 [[KPRB]], i8* [[OMP_EXEC_STATUS]], align 1
-  // CHECK: [[WORK:%.+]] = load i8*, i8** [[OMP_WORK_FN]],
-  // CHECK: [[SHOULD_EXIT:%.+]] = icmp eq i8* [[WORK]], null
-  // CHECK: br i1 [[SHOULD_EXIT]], label {{%?}}[[EXIT:.+]], label {{%?}}[[SEL_WORKERS:.+]]
-  //
-  // CHECK: [[SEL_WORKERS]]
-  // CHECK: [[ST:%.+]] = load i8, i8* [[OMP_EXEC_STATUS]]
-  // CHECK: [[IS_ACTIVE:%.+]] = icmp ne i8 [[ST]], 0
-  // CHECK: br i1 [[IS_ACTIVE]], label {{%?}}[[EXEC_PARALLEL:.+]], label {{%?}}[[BAR_PARALLEL:.+]]
-  //
-  // CHECK: [[EXEC_PARALLEL]]
-  // CHECK: [[WF:%.+]] = load i8*, i8** [[OMP_WORK_FN]],
-  // CHECK: [[WM:%.+]] = icmp eq i8* [[WF]], bitcast (void (i32*, i32*)* [[PARALLEL_FN4:@.+]] to i8*)
-  // CHECK: br i1 [[WM]], label {{%?}}[[EXEC_PFN:.+]], label {{%?}}[[CHECK_NEXT:.+]]
-  //
-  // CHECK: [[EXEC_PFN]]
-  // CHECK: call void [[PARALLEL_FN4]](
-  // CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
-  //
-  // CHECK: [[CHECK_NEXT]]
-  // CHECK: br label {{%?}}[[TERM_PARALLEL:.+]]
-  //
-  // CHECK: [[TERM_PARALLEL]]
-  // CHECK: call void @__kmpc_kernel_end_parallel()
-  // CHECK: br label {{%?}}[[BAR_PARALLEL]]
-  //
-  // CHECK: [[BAR_PARALLEL]]
-  // CHECK: call void @llvm.nvvm.barrier0()
-  // CHECK: br label {{%?}}[[AWAIT_WORK]]
-  //
-  // CHECK: [[EXIT]]
-  // CHECK: ret void
+// CHECK: [[KPR:%.+]] = call i1 @__kmpc_kernel_parallel(i8** [[OMP_WORK_FN]],
+// CHECK: [[WM:%.+]] = icmp eq i8* [[WF]], bitcast (void (i16, i32)* [[PARALLEL_FN4:@.+]]_wrapper to i8*)
+// CHECK: call void [[PARALLEL_FN4]]_wrapper(
 
-  // CHECK: define {{.*}}void [[T6:@__omp_offloading_.+template.+l43]](i[[SZ:32|64]]
-  // Create local storage for each capture.
-  // CHECK:  [[LOCAL_N:%.+]] = alloca i[[SZ]],
-  // CHECK:  [[LOCAL_A:%.+]] = alloca i[[SZ]],
-  // CHECK:  [[LOCAL_AA:%.+]] = alloca i[[SZ]],
-  // CHECK:  [[LOCAL_B:%.+]] = alloca [10 x i32]*
-  // CHECK-DAG:  store i[[SZ]] [[ARG_N:%.+]], i[[SZ]]* [[LOCAL_N]]
-  // CHECK-DAG:  store i[[SZ]] [[ARG_A:%.+]], i[[SZ]]* [[LOCAL_A]]
-  // CHECK-DAG:  store i[[SZ]] [[ARG_AA:%.+]], i[[SZ]]* [[LOCAL_AA]]
-  // CHECK-DAG:   store [10 x i32]* [[ARG_B:%.+]], [10 x i32]** [[LOCAL_B]]
-  // Store captures in the context.
-  // CHECK-64-DAG:[[REF_N:%.+]] = bitcast i[[SZ]]* [[LOCAL_N]] to i32*
-  // CHECK-64-DAG:[[REF_A:%.+]] = bitcast i[[SZ]]* [[LOCAL_A]] to i32*
-  // CHECK-DAG:   [[REF_AA:%.+]] = bitcast i[[SZ]]* [[LOCAL_AA]] to i16*
-  // CHECK-DAG:   [[REF_B:%.+]] = load [10 x i32]*, [10 x i32]** [[LOCAL_B]],
-  //
-  // CHECK-DAG: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
-  // CHECK-DAG: [[NTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
-  // CHECK-DAG: [[WS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
-  // CHECK-DAG: [[TH_LIMIT:%.+]] = sub i32 [[NTH]], [[WS]]
-  // CHECK: [[IS_WORKER:%.+]] = icmp ult i32 [[TID]], [[TH_LIMIT]]
-  // CHECK: br i1 [[IS_WORKER]], label {{%?}}[[WORKER:.+]], label {{%?}}[[CHECK_MASTER:.+]]
-  //
-  // CHECK: [[WORKER]]
-  // CHECK: {{call|invoke}} void [[T6]]_worker()
-  // CHECK: br label {{%?}}[[EXIT:.+]]
-  //
-  // CHECK: [[CHECK_MASTER]]
-  // CHECK-DAG: [[CMTID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x()
-  // CHECK-DAG: [[CMNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
-  // CHECK-DAG: [[CMWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
-  // CHECK: [[IS_MASTER:%.+]] = icmp eq i32 [[CMTID]],
-  // CHECK: br i1 [[IS_MASTER]], label {{%?}}[[MASTER:.+]], label {{%?}}[[EXIT]]
-  //
-  // CHECK: [[MASTER]]
-  // CHECK-DAG: [[MNTH:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.ntid.x()
-  // CHECK-DAG: [[MWS:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.warpsize()
-  // CHECK: [[MTMP1:%.+]] = sub i32 [[MNTH]], [[MWS]]
-  // CHECK: call void @__kmpc_kernel_init(i32 [[MTMP1]]
-  // CHECK-64: [[N:%.+]] = load i32, i32* [[REF_N]],
-  // CHECK-32: [[N:%.+]] = load i32, i32* [[LOCAL_N]],
-  // CHECK: [[CMP:%.+]] = icmp sgt i32 [[N]], 1000
-  // CHECK: br i1 [[CMP]], label {{%?}}[[IF_THEN:.+]], label {{%?}}[[IF_ELSE:.+]]
-  //
-  // CHECK: [[IF_THEN]]
-  // CHECK: call void @__kmpc_kernel_prepare_parallel(i8* bitcast (void (i32*, i32*)* [[PARALLEL_FN4]] to i8*),
-  // CHECK: call void @llvm.nvvm.barrier0()
-  // CHECK: call void @llvm.nvvm.barrier0()
-  // CHECK: br label {{%?}}[[IF_END:.+]]
-  //
-  // CHECK: [[IF_ELSE]]
-  // CHECK: call void @__kmpc_serialized_parallel(
-  // CHECK: {{call|invoke}} void [[PARALLEL_FN4]](
-  // CHECK: call void @__kmpc_end_serialized_parallel(
-  // br label [[IF_END]]
-  //
-  // CHECK: [[IF_END]]
-  // CHECK-64-DAG: load i32, i32* [[REF_A]]
-  // CHECK-32-DAG: load i32, i32* [[LOCAL_A]]
-  // CHECK-DAG:    load i16, i16* [[REF_AA]]
-  // CHECK-DAG:    getelementptr inbounds [10 x i32], [10 x i32]* [[REF_B]], i[[SZ]] 0, i[[SZ]] 2
-  //
-  // CHECK: br label {{%?}}[[TERMINATE:.+]]
-  //
-  // CHECK: [[TERMINATE]]
-  // CHECK: call void @__kmpc_kernel_deinit(
-  // CHECK: call void @llvm.nvvm.barrier0()
-  // CHECK: br label {{%?}}[[EXIT]]
-  //
-  // CHECK: [[EXIT]]
-  // CHECK: ret void
+// CHECK: call void @__kmpc_kernel_prepare_parallel(i8* bitcast (void (i16, i32)* [[PARALLEL_FN4]]_wrapper to i8*),
 
-  // CHECK: define internal void [[PARALLEL_FN4]](
-  // CHECK: [[A:%.+]] = alloca i[[SZ:32|64]],
-  // CHECK: store i[[SZ]] 45, i[[SZ]]* %a,
-  // CHECK: ret void
 #endif