[OpenMP] Basic support for a parallel directive in a target region on an NVPTX device.

[arpith-jacob]

Summary

This patch introduces support for the execution of parallel constructs in a target
region on the NVPTX device. Parallel regions must be in the lexical scope of the
target directive.

The master thread in the master warp signals parallel work for worker threads in worker
warps on encountering a parallel region.

Note: The patch does not yet support capture of arguments in a parallel region so
the test cases are simple.

[arpith-jacob]

lib/CodeGen/CGOpenMPRuntime.cpp

114

On the nvptx device, it is illegal for an identifier to contain a dot ('.') so I've modified it here. If there is a better way to do this, please let me know.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

114

Could you just override this function in CGOpenMPRuntimeNVPTX?

lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp

127

I believe this FIXME is not needed as there is already no reminder operation.

130

s/it/if/g?

325–326

Why you need double casting, could you just use Bld.CreatePointerBitCastOrAddrSpaceCast()?

339

auto *

343–345

Could you use just an array here instead of SmallVector?

358

Use llvm::None instead of default constructor.

561

use llvm::ConstantPointerNull::get() instead.

563

use llvm::ConstantPointerNull::get() instead.

573–575

No need for braces in one-line substatement

[arpith-jacob]

lib/CodeGen/CGOpenMPRuntime.cpp

114

Alexey, thank you for your review of this patch.

To override getHelperName() in CGOpenMPRuntime.cpp, I will have to move the two classes CGOpenMPRegionInfo and CGOpenMPOutlinedRegionInfo from the anonymous namespace to the header file in CGOpenMPRuntime.h (under protected mode). I would prefer to do this since I will need to use these classes for implementing codegen of other directives in the future.

Is that okay with you?

[ABataev]

lib/CodeGen/CGOpenMPRuntime.cpp

114

Yes, go ahead

[arpith-jacob]

Alexey, thank you for your review. I've updated the patch addressing your comments.

• I experimented with various ways of changing the name of the outlined function. In the end I decided against moving the two classes to the header file. Instead I pass the name to the class via the constructor. The name is returned by a virtual function.

• I also removed the overloaded implementation of getThreadID(). I'm fine using the default implementation. We do not use the value returned by getThreadID() in our runtime and prefer to recompute it using PTX reserved registers.

[ABataev]

lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp

567–570

It is better to emit this code as PrePostAction, so it is called upon exit of any cleanup scope

[arpith-jacob]

lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp

567–570

Alexey, do you mean clean up during the execution of the serialized parallel region? Is something like this what you have in mind? Thanks.

auto &&SeqGen = [this, Fn, &CapturedVars, &RTLoc, &Loc](CodeGenFunction &CGF,
                                                        PrePostActionTy &) {
  auto &&CodeGen = [..](..) {
      llvm::Value *Args[] = {RTLoc, ThreadID};
      CGF.EmitRuntimeCall(
          createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_serialized_parallel),
          Args);

      llvm::SmallVector<llvm::Value *, 16> OutlinedFnArgs;
      OutlinedFnArgs.push_back(
          llvm::ConstantPointerNull::get(CGM.Int32Ty->getPointerTo()));
      OutlinedFnArgs.push_back(
          llvm::ConstantPointerNull::get(CGM.Int32Ty->getPointerTo()));
      OutlinedFnArgs.append(CapturedVars.begin(), CapturedVars.end());
      CGF.EmitCallOrInvoke(Fn, OutlinedFnArgs);
  };

  RegionCodeGenTy RCG(CodeGen);
  CommonActionTy Action(
      nullptr, llvm::None,
      createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_end_serialized_parallel),
      {emitUpdateLocation(CGF, Loc), ThreadID});
  RCG.setAction(Action);
  RCG(CGF);
};

[tra]

lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp

539–544

Are two back-to-back syncCTAThreads() intentional or did you mean to call something else to activate workers?

[arpith-jacob]

lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp

539–544

Intentional. The first wakes up worker warps to execute the parallel region. The second is for the master warp to wait until the region has completed execution. This is to implement the semantics of an implicit barrier at the end of an OpenMP parallel region.

[tra]

lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp

539–544

OK. It may be good to add few comments documenting these assumptions (or point to existing docs) here.

[arpith-jacob]

Updated patch based on reviews.

The serialized parallel region executed by the master was modified to call the 'end' runtime call with a PrePostActionTy so that it is called upon exit of any cleanup scope.

Added reference to OpenMP spec regarding implicit barrier at the end of a parallel region.

[ABataev]

lib/CodeGen/CGOpenMPRuntime.h

63

I don't think it is good to name it CommonActionTy, I think it is something specific to NVPTX codegen. Could you give it some specific name + move it somewhere to NVPTX-related part of code (to .cpp file, if possible)

lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp

407

Does it really return I1 type? Or I8?

567–570

Yes, exactly.

[arpith-jacob]

Moved CommonActionTy to CGOpenMPRuntimeNVPTX.cpp and renamed it to NVPTXActionTy, allowing us to customize the class in the future, if necessary.

[arpith-jacob]

lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp

407

Alexey, the runtime function is called by every worker thread. It returns a 'bool' that indicates if the thread has been activated in the parallel region, which is why I used an Int1Ty. Please let me know if you see problems with this. Thanks.

[arpith-jacob]

Using CGF.ConvertTypeForMem(Context.getBoolType()) to get the right type for 'bool' rather than using i1.

[arpith-jacob]

Use i1 type for bool after all. But this time use the api ConvertType().

[ABataev]

LG with a nit

lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp

415

llvm::None instead of {}

lib/CodeGen/CGOpenMPRuntime.h

 public:
   explicit PrePostActionTy() {}
   virtual void Enter(CodeGenFunction &CGF) {}
   virtual void Exit(CodeGenFunction &CGF) {}
   virtual ~PrePostActionTy() {}
 };
 
 /// Class provides a way to call simple version of codegen for OpenMP region, or
 /// an advanced with possible pre|post-actions in codegen.
 class RegionCodeGenTy final {
   intptr_t CodeGen;
   typedef void (*CodeGenTy)(intptr_t, CodeGenFunction &, PrePostActionTy &);
   CodeGenTy Callback;
   mutable PrePostActionTy *PrePostAction;
   RegionCodeGenTy() = delete;
   RegionCodeGenTy &operator=(const RegionCodeGenTy &) = delete;
   template <typename Callable>
protected:
   /// evaluates to false.
   virtual void emitTargetOutlinedFunctionHelper(const OMPExecutableDirective &D,
                                                 StringRef ParentName,
                                                 llvm::Function *&OutlinedFn,
                                                 llvm::Constant *&OutlinedFnID,
                                                 bool IsOffloadEntry,
                                                 const RegionCodeGenTy &CodeGen);
 
+  /// \brief Emits code for OpenMP 'if' clause using specified \a CodeGen
+  /// function. Here is the logic:
+  /// if (Cond) {
+  ///   ThenGen();
+  /// } else {
+  ///   ElseGen();
+  /// }
+  void emitOMPIfClause(CodeGenFunction &CGF, const Expr *Cond,
+                       const RegionCodeGenTy &ThenGen,
+                       const RegionCodeGenTy &ElseGen);
+
+  /// \brief Emits object of ident_t type with info for source location.
+  /// \param Flags Flags for OpenMP location.
+  ///
+  llvm::Value *emitUpdateLocation(CodeGenFunction &CGF, SourceLocation Loc,
+                                  unsigned Flags = 0);
+
+  /// \brief Returns pointer to ident_t type.
+  llvm::Type *getIdentTyPointerTy();
+
+  /// \brief Gets thread id value for the current thread.
+  ///
+  virtual llvm::Value *getThreadID(CodeGenFunction &CGF, SourceLocation Loc);
+
 private:
   /// \brief Default const ident_t object used for initialization of all other
   /// ident_t objects.
   llvm::Constant *DefaultOpenMPPSource = nullptr;
   /// \brief Map of flags and corresponding default locations.
   typedef llvm::DenseMap<unsigned, llvm::Value *> OpenMPDefaultLocMapTy;
   OpenMPDefaultLocMapTy OpenMPDefaultLocMap;
   Address getOrCreateDefaultLocation(unsigned Flags);
private:
   /// found along the way.
   /// \param S Starting statement.
   /// \param ParentName Name of the function declaration that is being scanned.
   void scanForTargetRegionsFunctions(const Stmt *S, StringRef ParentName);
 
   /// \brief Build type kmp_routine_entry_t (if not built yet).
   void emitKmpRoutineEntryT(QualType KmpInt32Ty);
 
-  /// \brief Emits object of ident_t type with info for source location.
-  /// \param Flags Flags for OpenMP location.
-  ///
-  llvm::Value *emitUpdateLocation(CodeGenFunction &CGF, SourceLocation Loc,
-                                  unsigned Flags = 0);
-
-  /// \brief Returns pointer to ident_t type.
-  llvm::Type *getIdentTyPointerTy();
-
   /// \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(unsigned Function);
 
private:
   /// \param VD Threadprivate variable.
   /// \return Cache variable for the specified threadprivate.
   llvm::Constant *getOrCreateThreadPrivateCache(const VarDecl *VD);
 
   /// \brief Emits address of the word in a memory where current thread id is
   /// stored.
   virtual Address emitThreadIDAddress(CodeGenFunction &CGF, SourceLocation Loc);
 
-  /// \brief Gets thread id value for the current thread.
-  ///
-  llvm::Value *getThreadID(CodeGenFunction &CGF, SourceLocation Loc);
-
   /// \brief Gets (if variable with the given name already exist) or creates
   /// internal global variable with the specified Name. The created variable has
   /// linkage CommonLinkage by default and is initialized by null value.
   /// \param Ty Type of the global variable. If it is exist already the type
   /// must be the same.
   /// \param Name Name of the variable.
   llvm::Constant *getOrCreateInternalVariable(llvm::Type *Ty,
                                               const llvm::Twine &Name);

lib/CodeGen/CGOpenMPRuntime.cpp

CGOpenMPOutlinedRegionInfo(const CapturedStmt &CS, const VarDecl *ThreadIDVar,
     assert(ThreadIDVar != nullptr && "No ThreadID in OpenMP region.");
   }
 
   /// \brief Get a variable or parameter for storing global thread id
   /// inside OpenMP construct.
   const VarDecl *getThreadIDVariable() const override { return ThreadIDVar; }
 
   /// \brief Get the name of the capture helper.
-  StringRef getHelperName() const override { return ".omp_outlined."; }
+  StringRef getHelperName() const override { return "__omp_outlined__"; }
 
   static bool classof(const CGCapturedStmtInfo *Info) {
     return CGOpenMPRegionInfo::classof(Info) &&
            cast<CGOpenMPRegionInfo>(Info)->getRegionKind() ==
                ParallelOutlinedRegion;
   }
 
 private:

 
 /// \brief Emits code for OpenMP 'if' clause using specified \a CodeGen
 /// function. Here is the logic:
 /// if (Cond) {
 ///   ThenGen();
 /// } else {
 ///   ElseGen();
 /// }
-static void emitOMPIfClause(CodeGenFunction &CGF, const Expr *Cond,
+void CGOpenMPRuntime::emitOMPIfClause(CodeGenFunction &CGF, const Expr *Cond,
                                       const RegionCodeGenTy &ThenGen,
                                       const RegionCodeGenTy &ElseGen) {
   CodeGenFunction::LexicalScope ConditionScope(CGF, Cond->getSourceRange());
 
   // If the condition constant folds and can be elided, try to avoid emitting
   // the condition and the dead arm of the if/else.
   bool CondConstant;
   if (CGF.ConstantFoldsToSimpleInteger(Cond, CondConstant)) {
     if (CondConstant)
       ThenGen(CGF);

lib/CodeGen/CGOpenMPRuntimeNVPTX.h

 #include "clang/AST/StmtOpenMP.h"
 #include "llvm/IR/CallSite.h"
 
 namespace clang {
 namespace CodeGen {
 
 class CGOpenMPRuntimeNVPTX : public CGOpenMPRuntime {
 private:
+  // Parallel outlined function work for workers to execute.
+  llvm::SmallVector<llvm::Function *, 16> Work;
+
   struct EntryFunctionState {
     llvm::BasicBlock *ExitBB = nullptr;
   };
 
   class WorkerFunctionState {
   public:
     llvm::Function *WorkerFn;
     const CGFunctionInfo *CGFI;
private:
   // Base class overrides.
   //
 
   /// \brief Creates offloading entry for the provided entry ID \a ID,
   /// address \a Addr and size \a Size.
   void createOffloadEntry(llvm::Constant *ID, llvm::Constant *Addr,
                           uint64_t Size) override;
 
+  /// \brief Gets thread id value for the current thread.
+  ///
+  llvm::Value *getThreadID(CodeGenFunction &CGF, SourceLocation Loc) override;
+
   /// \brief Emit outlined function specialized for the Fork-Join
   /// programming model for applicable target directives on the NVPTX device.
   /// \param D Directive to emit.
   /// \param ParentName Name of the function that encloses the target region.
   /// \param OutlinedFn Outlined function value to be defined by this call.
   /// \param OutlinedFnID Outlined function ID value to be defined by this call.
   /// \param IsOffloadEntry True if the outlined function is an offload entry.
   /// An outlined function may not be an entry if, e.g. the if clause always
private:
   /// evaluates to false.
   void emitTargetOutlinedFunction(const OMPExecutableDirective &D,
                                   StringRef ParentName,
                                   llvm::Function *&OutlinedFn,
                                   llvm::Constant *&OutlinedFnID,
                                   bool IsOffloadEntry,
                                   const RegionCodeGenTy &CodeGen) override;
 
+  /// \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.
+  /// This call is for the Generic Execution Mode.
+  /// \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.
+  /// \param IfCond Condition in the associated 'if' clause, if it was
+  /// specified, nullptr otherwise.
+  void emitGenericParallelCall(CodeGenFunction &CGF, SourceLocation Loc,
+                               llvm::Value *OutlinedFn,
+                               ArrayRef<llvm::Value *> CapturedVars,
+                               const Expr *IfCond);
+
 public:
   explicit CGOpenMPRuntimeNVPTX(CodeGenModule &CGM);
 
   /// \brief This function ought to emit, in the general case, a call to
   // the openmp runtime kmpc_push_num_teams. In NVPTX backend it is not needed
   // as these numbers are obtained through the PTX grid and block configuration.
   /// \param NumTeams An integer expression of teams.
   /// \param ThreadLimit An integer expression of threads.
public:
   /// \param OutlinedFn Outlined function to be run by team masters. 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.
   ///
   void emitTeamsCall(CodeGenFunction &CGF, const OMPExecutableDirective &D,
                      SourceLocation Loc, llvm::Value *OutlinedFn,
                      ArrayRef<llvm::Value *> CapturedVars) override;
+
+  /// \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.
+  /// \param IfCond Condition in the associated 'if' clause, if it was
+  /// specified, nullptr otherwise.
+  void emitParallelCall(CodeGenFunction &CGF, SourceLocation Loc,
+                        llvm::Value *OutlinedFn,
+                        ArrayRef<llvm::Value *> CapturedVars,
+                        const Expr *IfCond) override;
 };
 
 } // CodeGen namespace.
 } // clang namespace.
 
 #endif // LLVM_CLANG_LIB_CODEGEN_CGOPENMPRUNTIMENVPTX_H

lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp

 using namespace CodeGen;
 
 namespace {
 enum OpenMPRTLFunctionNVPTX {
   /// \brief Call to void __kmpc_kernel_init(kmp_int32 thread_limit);
   OMPRTL_NVPTX__kmpc_kernel_init,
   /// \brief Call to void __kmpc_kernel_deinit();
   OMPRTL_NVPTX__kmpc_kernel_deinit,
+  /// \brief Call to void __kmpc_kernel_prepare_parallel(void
+  /// *outlined_function);
+  OMPRTL_NVPTX__kmpc_kernel_prepare_parallel,
+  /// \brief Call to bool __kmpc_kernel_parallel(void **outlined_function);
+  OMPRTL_NVPTX__kmpc_kernel_parallel,
+  /// \brief Call to void __kmpc_kernel_end_parallel();
+  OMPRTL_NVPTX__kmpc_kernel_end_parallel,
+  /// Call to void __kmpc_serialized_parallel(ident_t *loc, kmp_int32
+  /// global_tid);
+  OMPRTL_NVPTX__kmpc_serialized_parallel,
+  /// Call to void __kmpc_end_serialized_parallel(ident_t *loc, kmp_int32
+  /// global_tid);
+  OMPRTL_NVPTX__kmpc_end_serialized_parallel,
 };
 } // namespace
 
 /// Get the GPU warp size.
 static llvm::Value *getNVPTXWarpSize(CodeGenFunction &CGF) {
   CGBuilderTy &Bld = CGF.Builder;
   return Bld.CreateCall(
       llvm::Intrinsic::getDeclaration(
static llvm::Value *getMasterThreadID(CodeGenFunction &CGF) {
 
   // We assume that the warp size is a power of 2.
   llvm::Value *Mask = Bld.CreateSub(getNVPTXWarpSize(CGF), Bld.getInt32(1));
 
   return Bld.CreateAnd(Bld.CreateSub(NumThreads, Bld.getInt32(1)),
                        Bld.CreateNot(Mask), "master_tid");
 }
 
+/// Get the id of the current block on the GPU.
+static llvm::Value *getNVPTXBlockID(CodeGenFunction &CGF) {
+  CGBuilderTy &Bld = CGF.Builder;
+  return Bld.CreateCall(
+      llvm::Intrinsic::getDeclaration(
+          &CGF.CGM.getModule(), llvm::Intrinsic::nvvm_read_ptx_sreg_ctaid_x),
+      llvm::None, "nvptx_block_id");
+}
+
+/// Get number of OMP workers for parallel region after subtracting
+/// the master warp.
+static llvm::Value *getNumWorkers(CodeGenFunction &CGF) {
+  CGBuilderTy &Bld = CGF.Builder;
+  return Bld.CreateNUWSub(getNVPTXNumThreads(CGF), Bld.getInt32(32),
+                          "num_workers");
+}
+
+/// Get thread id in team.
+/// FIXME: Remove the expensive remainder operation.
+static llvm::Value *getTeamThreadId(CodeGenFunction &CGF) {
+  CGBuilderTy &Bld = CGF.Builder;
+  // N % M = N & (M-1) it M is a power of 2. The master Id is expected to be a
+  // power of two in all cases.
+  auto *Mask = Bld.CreateNUWSub(getMasterThreadID(CGF), Bld.getInt32(1));
+  return Bld.CreateAnd(getNVPTXThreadID(CGF), Mask, "team_tid");
+}
+
+/// Get global thread id.
+static llvm::Value *getGlobalThreadId(CodeGenFunction &CGF) {
+  CGBuilderTy &Bld = CGF.Builder;
+  return Bld.CreateAdd(Bld.CreateMul(getNVPTXBlockID(CGF), getNumWorkers(CGF)),
+                       getTeamThreadId(CGF), "global_tid");
+}
+
+llvm::Value *CGOpenMPRuntimeNVPTX::getThreadID(CodeGenFunction &CGF,
+                                               SourceLocation Loc) {
+  assert(CGF.CurFn && "No function in current CodeGenFunction.");
+  return getGlobalThreadId(CGF);
+}
+
 CGOpenMPRuntimeNVPTX::WorkerFunctionState::WorkerFunctionState(
     CodeGenModule &CGM)
     : WorkerFn(nullptr), CGFI(nullptr) {
   createWorkerFunction(CGM);
 }
 
 void CGOpenMPRuntimeNVPTX::WorkerFunctionState::createWorkerFunction(
     CodeGenModule &CGM) {

 void CGOpenMPRuntimeNVPTX::emitGenericKernel(const OMPExecutableDirective &D,
                                              StringRef ParentName,
                                              llvm::Function *&OutlinedFn,
                                              llvm::Constant *&OutlinedFnID,
                                              bool IsOffloadEntry,
                                              const RegionCodeGenTy &CodeGen) {
   EntryFunctionState EST;
   WorkerFunctionState WST(CGM);
+  Work.clear();
 
   // Emit target region as a standalone region.
   class NVPTXPrePostActionTy : public PrePostActionTy {
     CGOpenMPRuntimeNVPTX &RT;
     CGOpenMPRuntimeNVPTX::EntryFunctionState &EST;
     CGOpenMPRuntimeNVPTX::WorkerFunctionState &WST;
 
   public:
void CGOpenMPRuntimeNVPTX::emitWorkerLoop(CodeGenFunction &CGF,
 
   Address WorkFn =
       CGF.CreateDefaultAlignTempAlloca(CGF.Int8PtrTy, /*Name=*/"work_fn");
   Address ExecStatus =
       CGF.CreateDefaultAlignTempAlloca(CGF.Int8Ty, /*Name=*/"exec_status");
   CGF.InitTempAlloca(ExecStatus, Bld.getInt8(/*C=*/0));
   CGF.InitTempAlloca(WorkFn, llvm::Constant::getNullValue(CGF.Int8PtrTy));
 
-  // TODO: Call into runtime to get parallel work.
+  llvm::Value *Args[] = {WorkFn.getPointer()};
+  llvm::Value *Ret = CGF.EmitRuntimeCall(
+      createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_kernel_parallel), Args);
+  Bld.CreateStore(Bld.CreateZExt(Ret, CGF.Int8Ty), ExecStatus);
 
   // On termination condition (workid == 0), exit loop.
   llvm::Value *ShouldTerminate =
       Bld.CreateIsNull(Bld.CreateLoad(WorkFn), "should_terminate");
   Bld.CreateCondBr(ShouldTerminate, ExitBB, SelectWorkersBB);
 
   // Activate requested workers.
   CGF.EmitBlock(SelectWorkersBB);
   llvm::Value *IsActive =
       Bld.CreateIsNotNull(Bld.CreateLoad(ExecStatus), "is_active");
   Bld.CreateCondBr(IsActive, ExecuteBB, BarrierBB);
 
   // Signal start of parallel region.
   CGF.EmitBlock(ExecuteBB);
-  // TODO: Add parallel work.
+
+  // Process work items: outlined parallel functions.
+  for (auto *W : Work) {
+    // Try to match this outlined function.
+    auto ID = Bld.CreatePtrToInt(W, CGM.Int64Ty);
+    ID = Bld.CreateIntToPtr(ID, CGM.Int8PtrTy);
+    llvm::Value *WorkFnMatch =
+        Bld.CreateICmpEQ(Bld.CreateLoad(WorkFn), ID, "work_match");
+
+    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::SmallVector<llvm::Value *, 16> FnArgs;
+    FnArgs.push_back(ZeroAddr.getPointer());
+    FnArgs.push_back(ZeroAddr.getPointer());
+    CGF.EmitCallOrInvoke(Fn, FnArgs);
+
+    // Go to end of parallel region.
+    CGF.EmitBranch(TerminateBB);
+
+    CGF.EmitBlock(CheckNextBB);
+  }
 
   // Signal end of parallel region.
   CGF.EmitBlock(TerminateBB);
+  CGF.EmitRuntimeCall(
+      createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_kernel_end_parallel),
+      ArrayRef<llvm::Value *>());
   CGF.EmitBranch(BarrierBB);
 
   // All active and inactive workers wait at a barrier after parallel region.
   CGF.EmitBlock(BarrierBB);
   // Barrier after parallel region.
   syncCTAThreads(CGF);
   CGF.EmitBranch(AwaitBB);
 
CGOpenMPRuntimeNVPTX::createNVPTXRuntimeFunction(unsigned Function) {
   }
   case OMPRTL_NVPTX__kmpc_kernel_deinit: {
     // Build void __kmpc_kernel_deinit();
     llvm::FunctionType *FnTy =
         llvm::FunctionType::get(CGM.VoidTy, {}, /*isVarArg*/ false);
     RTLFn = CGM.CreateRuntimeFunction(FnTy, "__kmpc_kernel_deinit");
     break;
   }
+  case OMPRTL_NVPTX__kmpc_kernel_prepare_parallel: {
+    /// Build void __kmpc_kernel_prepare_parallel(
+    /// void *outlined_function);
+    llvm::Type *TypeParams[] = {CGM.Int8PtrTy};
+    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);
+    llvm::Type *TypeParams[] = {CGM.Int8PtrPtrTy};
+    llvm::FunctionType *FnTy =
+        llvm::FunctionType::get(llvm::Type::getInt1Ty(CGM.getLLVMContext()),
+                                TypeParams, /*isVarArg*/ false);
+    RTLFn = CGM.CreateRuntimeFunction(FnTy, "__kmpc_kernel_parallel");
+    break;
+  }
+  case OMPRTL_NVPTX__kmpc_kernel_end_parallel: {
+    /// Build void __kmpc_kernel_end_parallel();
+    llvm::FunctionType *FnTy =
+        llvm::FunctionType::get(CGM.VoidTy, {}, /*isVarArg*/ false);
+    RTLFn = CGM.CreateRuntimeFunction(FnTy, "__kmpc_kernel_end_parallel");
+    break;
+  }
+  case OMPRTL_NVPTX__kmpc_serialized_parallel: {
+    // Build void __kmpc_serialized_parallel(ident_t *loc, kmp_int32
+    // global_tid);
+    llvm::Type *TypeParams[] = {getIdentTyPointerTy(), CGM.Int32Ty};
+    llvm::FunctionType *FnTy =
+        llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg*/ false);
+    RTLFn = CGM.CreateRuntimeFunction(FnTy, "__kmpc_serialized_parallel");
+    break;
+  }
+  case OMPRTL_NVPTX__kmpc_end_serialized_parallel: {
+    // Build void __kmpc_end_serialized_parallel(ident_t *loc, kmp_int32
+    // global_tid);
+    llvm::Type *TypeParams[] = {getIdentTyPointerTy(), CGM.Int32Ty};
+    llvm::FunctionType *FnTy =
+        llvm::FunctionType::get(CGM.VoidTy, TypeParams, /*isVarArg*/ false);
+    RTLFn = CGM.CreateRuntimeFunction(FnTy, "__kmpc_end_serialized_parallel");
+    break;
+  }
   }
   return RTLFn;
 }
 
 void CGOpenMPRuntimeNVPTX::createOffloadEntry(llvm::Constant *ID,
                                               llvm::Constant *Addr,
                                               uint64_t Size) {
   auto *F = dyn_cast<llvm::Function>(Addr);
void CGOpenMPRuntimeNVPTX::emitNumTeamsClause(CodeGenFunction &CGF,
                                               const Expr *NumTeams,
                                               const Expr *ThreadLimit,
                                               SourceLocation Loc) {}
 
 llvm::Value *CGOpenMPRuntimeNVPTX::emitParallelOrTeamsOutlinedFunction(
     const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
     OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen) {
 
-  llvm::Function *OutlinedFun = nullptr;
-  if (isa<OMPTeamsDirective>(D)) {
-    llvm::Value *OutlinedFunVal =
-        CGOpenMPRuntime::emitParallelOrTeamsOutlinedFunction(
-            D, ThreadIDVar, InnermostKind, CodeGen);
-    OutlinedFun = cast<llvm::Function>(OutlinedFunVal);
-    OutlinedFun->removeFnAttr(llvm::Attribute::NoInline);
-    OutlinedFun->addFnAttr(llvm::Attribute::AlwaysInline);
-  } else
-    llvm_unreachable("parallel directive is not yet supported for nvptx "
-                     "backend.");
-
-  return OutlinedFun;
+  return CGOpenMPRuntime::emitParallelOrTeamsOutlinedFunction(
 }
 
 void CGOpenMPRuntimeNVPTX::emitTeamsCall(CodeGenFunction &CGF,
                                          const OMPExecutableDirective &D,
                                          SourceLocation Loc,
                                          llvm::Value *OutlinedFn,
                                          ArrayRef<llvm::Value *> CapturedVars) {
   if (!CGF.HaveInsertPoint())
     return;
 
   Address ZeroAddr =
       CGF.CreateTempAlloca(CGF.Int32Ty, CharUnits::fromQuantity(4),
                            /*Name*/ ".zero.addr");
   CGF.InitTempAlloca(ZeroAddr, CGF.Builder.getInt32(/*C*/ 0));
   llvm::SmallVector<llvm::Value *, 16> OutlinedFnArgs;
   OutlinedFnArgs.push_back(ZeroAddr.getPointer());
   OutlinedFnArgs.push_back(ZeroAddr.getPointer());
   OutlinedFnArgs.append(CapturedVars.begin(), CapturedVars.end());
   CGF.EmitCallOrInvoke(OutlinedFn, OutlinedFnArgs);
 }
+
+void CGOpenMPRuntimeNVPTX::emitParallelCall(
+    CodeGenFunction &CGF, SourceLocation Loc, llvm::Value *OutlinedFn,
+    ArrayRef<llvm::Value *> CapturedVars, const Expr *IfCond) {
+  if (!CGF.HaveInsertPoint())
+    return;
+
+  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);
+
+  auto *RTLoc = emitUpdateLocation(CGF, Loc);
+  auto &&L0ParallelGen = [this, Fn, &CapturedVars](CodeGenFunction &CGF,
+                                                   PrePostActionTy &) {
+    CGBuilderTy &Bld = CGF.Builder;
+
+    // Prepare for parallel region. Indicate the outlined function.
+    llvm::Value *Args[] = {Bld.CreateBitOrPointerCast(Fn, CGM.Int8PtrTy)};
+    CGF.EmitRuntimeCall(
+        createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_kernel_prepare_parallel),
+        Args);
+
+    // Activate workers.
+    syncCTAThreads(CGF);
+
+    // Barrier at end of parallel region.
+    syncCTAThreads(CGF);
+
+    // Remember for post-processing in worker loop.
+    Work.push_back(Fn);
+  };
+
+  auto &&SeqGen = [this, Fn, &CapturedVars, &RTLoc, &Loc](CodeGenFunction &CGF,
+                                                          PrePostActionTy &) {
+    auto DL = CGM.getDataLayout();
+    auto ThreadID = getThreadID(CGF, Loc);
+
+    llvm::Value *Args[] = {RTLoc, ThreadID};
+    CGF.EmitRuntimeCall(
+        createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_serialized_parallel),
+        Args);
+
+    llvm::SmallVector<llvm::Value *, 16> OutlinedFnArgs;
+    OutlinedFnArgs.push_back(
+        llvm::Constant::getNullValue(CGM.Int32Ty->getPointerTo()));
+    OutlinedFnArgs.push_back(
+        llvm::Constant::getNullValue(CGM.Int32Ty->getPointerTo()));
+    OutlinedFnArgs.append(CapturedVars.begin(), CapturedVars.end());
+    CGF.EmitCallOrInvoke(Fn, OutlinedFnArgs);
+
+    llvm::Value *EndArgs[] = {emitUpdateLocation(CGF, Loc), ThreadID};
+    CGF.EmitRuntimeCall(
+        createNVPTXRuntimeFunction(OMPRTL_NVPTX__kmpc_end_serialized_parallel),
+        EndArgs);
+  };
+
+  if (IfCond) {
+    emitOMPIfClause(CGF, IfCond, L0ParallelGen, SeqGen);
+  } else {
+    CodeGenFunction::RunCleanupsScope Scope(CGF);
+    RegionCodeGenTy ThenRCG(L0ParallelGen);
+    ThenRCG(CGF);
+  }
+}

test/OpenMP/nvptx_parallel_codegen.cpp

+// Test target codegen - host bc file has to be created first.
+// 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 CHECK --check-prefix CHECK-64
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple i386-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm-bc %s -o %t-x86-host.bc
+// RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
+// RUN: %clang_cc1 -verify -fopenmp -fexceptions -fcxx-exceptions -x c++ -triple nvptx-unknown-unknown -fopenmp-targets=nvptx-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-x86-host.bc -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
+// expected-no-diagnostics
+#ifndef HEADER
+#define HEADER
+
+template<typename tx>
+tx ftemplate(int n) {
+  tx a = 0;
+  short aa = 0;
+  tx b[10];
+
+  #pragma omp target if(0)
+  {
+    #pragma omp parallel
+    {
+      int a = 41;
+    }
+    a += 1;
+  }
+
+  #pragma omp target
+  {
+    #pragma omp parallel
+    {
+      int a = 42;
+    }
+    #pragma omp parallel if(0)
+    {
+      int a = 43;
+    }
+    #pragma omp parallel if(1)
+    {
+      int a = 44;
+    }
+    a += 1;
+  }
+
+  #pragma omp target if(n>40)
+  {
+    #pragma omp parallel if(n>1000)
+    {
+      int a = 45;
+    }
+    a += 1;
+    aa += 1;
+    b[2] += 1;
+  }
+
+  return a;
+}
+
+int bar(int n){
+  int a = 0;
+
+  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: 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: [[CMTMP1:%.+]] = sub i32 [[CMWS]], 1
+  // CHECK: [[CMTMP2:%.+]] = sub i32 [[CMNTH]], 1
+  // CHECK: [[MID:%.+]] = and i32 [[CMTMP2]],
+  // CHECK: [[IS_MASTER:%.+]] = icmp eq i32 [[CMTID]], [[MID]]
+  // 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 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-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: 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: [[CMTMP1:%.+]] = sub i32 [[CMWS]], 1
+  // CHECK: [[CMTMP2:%.+]] = sub i32 [[CMNTH]], 1
+  // CHECK: [[MID:%.+]] = and i32 [[CMTMP2]],
+  // CHECK: [[IS_MASTER:%.+]] = icmp eq i32 [[CMTID]], [[MID]]
+  // 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 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: define internal void [[PARALLEL_FN4]](
+  // CHECK: [[A:%.+]] = alloca i[[SZ:32|64]],
+  // CHECK: store i[[SZ]] 45, i[[SZ]]* %a,
+  // CHECK: ret void
+#endif