1 //===---- CGOpenMPRuntimeGPU.cpp - Interface to OpenMP GPU Runtimes ----===//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 //
9 // This provides a generalized class for OpenMP runtime code generation
10 // specialized by GPU targets NVPTX and AMDGCN.
11 //
12 //===----------------------------------------------------------------------===//
13 
14 #include "CGOpenMPRuntimeGPU.h"
15 #include "CodeGenFunction.h"
16 #include "clang/AST/Attr.h"
17 #include "clang/AST/DeclOpenMP.h"
18 #include "clang/AST/StmtOpenMP.h"
19 #include "clang/AST/StmtVisitor.h"
20 #include "clang/Basic/Cuda.h"
21 #include "llvm/ADT/SmallPtrSet.h"
22 #include "llvm/Frontend/OpenMP/OMPGridValues.h"
23 #include "llvm/Support/MathExtras.h"
24 
25 using namespace clang;
26 using namespace CodeGen;
27 using namespace llvm::omp;
28 
29 namespace {
30 /// Pre(post)-action for different OpenMP constructs specialized for NVPTX.
31 class NVPTXActionTy final : public PrePostActionTy {
32   llvm::FunctionCallee EnterCallee = nullptr;
33   ArrayRef<llvm::Value *> EnterArgs;
34   llvm::FunctionCallee ExitCallee = nullptr;
35   ArrayRef<llvm::Value *> ExitArgs;
36   bool Conditional = false;
37   llvm::BasicBlock *ContBlock = nullptr;
38 
39 public:
40   NVPTXActionTy(llvm::FunctionCallee EnterCallee,
41                 ArrayRef<llvm::Value *> EnterArgs,
42                 llvm::FunctionCallee ExitCallee,
43                 ArrayRef<llvm::Value *> ExitArgs, bool Conditional = false)
44       : EnterCallee(EnterCallee), EnterArgs(EnterArgs), ExitCallee(ExitCallee),
45         ExitArgs(ExitArgs), Conditional(Conditional) {}
46   void Enter(CodeGenFunction &CGF) override {
47     llvm::Value *EnterRes = CGF.EmitRuntimeCall(EnterCallee, EnterArgs);
48     if (Conditional) {
49       llvm::Value *CallBool = CGF.Builder.CreateIsNotNull(EnterRes);
50       auto *ThenBlock = CGF.createBasicBlock("omp_if.then");
51       ContBlock = CGF.createBasicBlock("omp_if.end");
52       // Generate the branch (If-stmt)
53       CGF.Builder.CreateCondBr(CallBool, ThenBlock, ContBlock);
54       CGF.EmitBlock(ThenBlock);
55     }
56   }
57   void Done(CodeGenFunction &CGF) {
58     // Emit the rest of blocks/branches
59     CGF.EmitBranch(ContBlock);
60     CGF.EmitBlock(ContBlock, true);
61   }
62   void Exit(CodeGenFunction &CGF) override {
63     CGF.EmitRuntimeCall(ExitCallee, ExitArgs);
64   }
65 };
66 
67 /// A class to track the execution mode when codegening directives within
68 /// a target region. The appropriate mode (SPMD|NON-SPMD) is set on entry
69 /// to the target region and used by containing directives such as 'parallel'
70 /// to emit optimized code.
71 class ExecutionRuntimeModesRAII {
72 private:
73   CGOpenMPRuntimeGPU::ExecutionMode SavedExecMode =
74       CGOpenMPRuntimeGPU::EM_Unknown;
75   CGOpenMPRuntimeGPU::ExecutionMode &ExecMode;
76   bool SavedRuntimeMode = false;
77   bool *RuntimeMode = nullptr;
78 
79 public:
80   /// Constructor for Non-SPMD mode.
81   ExecutionRuntimeModesRAII(CGOpenMPRuntimeGPU::ExecutionMode &ExecMode)
82       : ExecMode(ExecMode) {
83     SavedExecMode = ExecMode;
84     ExecMode = CGOpenMPRuntimeGPU::EM_NonSPMD;
85   }
86   /// Constructor for SPMD mode.
87   ExecutionRuntimeModesRAII(CGOpenMPRuntimeGPU::ExecutionMode &ExecMode,
88                             bool &RuntimeMode, bool FullRuntimeMode)
89       : ExecMode(ExecMode), RuntimeMode(&RuntimeMode) {
90     SavedExecMode = ExecMode;
91     SavedRuntimeMode = RuntimeMode;
92     ExecMode = CGOpenMPRuntimeGPU::EM_SPMD;
93     RuntimeMode = FullRuntimeMode;
94   }
95   ~ExecutionRuntimeModesRAII() {
96     ExecMode = SavedExecMode;
97     if (RuntimeMode)
98       *RuntimeMode = SavedRuntimeMode;
99   }
100 };
101 
102 /// GPU Configuration:  This information can be derived from cuda registers,
103 /// however, providing compile time constants helps generate more efficient
104 /// code.  For all practical purposes this is fine because the configuration
105 /// is the same for all known NVPTX architectures.
106 enum MachineConfiguration : unsigned {
107   /// See "llvm/Frontend/OpenMP/OMPGridValues.h" for various related target
108   /// specific Grid Values like GV_Warp_Size, GV_Slot_Size
109 
110   /// Global memory alignment for performance.
111   GlobalMemoryAlignment = 128,
112 
113   /// Maximal size of the shared memory buffer.
114   SharedMemorySize = 128,
115 };
116 
117 static const ValueDecl *getPrivateItem(const Expr *RefExpr) {
118   RefExpr = RefExpr->IgnoreParens();
119   if (const auto *ASE = dyn_cast<ArraySubscriptExpr>(RefExpr)) {
120     const Expr *Base = ASE->getBase()->IgnoreParenImpCasts();
121     while (const auto *TempASE = dyn_cast<ArraySubscriptExpr>(Base))
122       Base = TempASE->getBase()->IgnoreParenImpCasts();
123     RefExpr = Base;
124   } else if (auto *OASE = dyn_cast<OMPArraySectionExpr>(RefExpr)) {
125     const Expr *Base = OASE->getBase()->IgnoreParenImpCasts();
126     while (const auto *TempOASE = dyn_cast<OMPArraySectionExpr>(Base))
127       Base = TempOASE->getBase()->IgnoreParenImpCasts();
128     while (const auto *TempASE = dyn_cast<ArraySubscriptExpr>(Base))
129       Base = TempASE->getBase()->IgnoreParenImpCasts();
130     RefExpr = Base;
131   }
132   RefExpr = RefExpr->IgnoreParenImpCasts();
133   if (const auto *DE = dyn_cast<DeclRefExpr>(RefExpr))
134     return cast<ValueDecl>(DE->getDecl()->getCanonicalDecl());
135   const auto *ME = cast<MemberExpr>(RefExpr);
136   return cast<ValueDecl>(ME->getMemberDecl()->getCanonicalDecl());
137 }
138 
139 
140 static RecordDecl *buildRecordForGlobalizedVars(
141     ASTContext &C, ArrayRef<const ValueDecl *> EscapedDecls,
142     ArrayRef<const ValueDecl *> EscapedDeclsForTeams,
143     llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>
144         &MappedDeclsFields, int BufSize) {
145   using VarsDataTy = std::pair<CharUnits /*Align*/, const ValueDecl *>;
146   if (EscapedDecls.empty() && EscapedDeclsForTeams.empty())
147     return nullptr;
148   SmallVector<VarsDataTy, 4> GlobalizedVars;
149   for (const ValueDecl *D : EscapedDecls)
150     GlobalizedVars.emplace_back(
151         CharUnits::fromQuantity(std::max(
152             C.getDeclAlign(D).getQuantity(),
153             static_cast<CharUnits::QuantityType>(GlobalMemoryAlignment))),
154         D);
155   for (const ValueDecl *D : EscapedDeclsForTeams)
156     GlobalizedVars.emplace_back(C.getDeclAlign(D), D);
157   llvm::stable_sort(GlobalizedVars, [](VarsDataTy L, VarsDataTy R) {
158     return L.first > R.first;
159   });
160 
161   // Build struct _globalized_locals_ty {
162   //         /*  globalized vars  */[WarSize] align (max(decl_align,
163   //         GlobalMemoryAlignment))
164   //         /*  globalized vars  */ for EscapedDeclsForTeams
165   //       };
166   RecordDecl *GlobalizedRD = C.buildImplicitRecord("_globalized_locals_ty");
167   GlobalizedRD->startDefinition();
168   llvm::SmallPtrSet<const ValueDecl *, 16> SingleEscaped(
169       EscapedDeclsForTeams.begin(), EscapedDeclsForTeams.end());
170   for (const auto &Pair : GlobalizedVars) {
171     const ValueDecl *VD = Pair.second;
172     QualType Type = VD->getType();
173     if (Type->isLValueReferenceType())
174       Type = C.getPointerType(Type.getNonReferenceType());
175     else
176       Type = Type.getNonReferenceType();
177     SourceLocation Loc = VD->getLocation();
178     FieldDecl *Field;
179     if (SingleEscaped.count(VD)) {
180       Field = FieldDecl::Create(
181           C, GlobalizedRD, Loc, Loc, VD->getIdentifier(), Type,
182           C.getTrivialTypeSourceInfo(Type, SourceLocation()),
183           /*BW=*/nullptr, /*Mutable=*/false,
184           /*InitStyle=*/ICIS_NoInit);
185       Field->setAccess(AS_public);
186       if (VD->hasAttrs()) {
187         for (specific_attr_iterator<AlignedAttr> I(VD->getAttrs().begin()),
188              E(VD->getAttrs().end());
189              I != E; ++I)
190           Field->addAttr(*I);
191       }
192     } else {
193       llvm::APInt ArraySize(32, BufSize);
194       Type = C.getConstantArrayType(Type, ArraySize, nullptr, ArrayType::Normal,
195                                     0);
196       Field = FieldDecl::Create(
197           C, GlobalizedRD, Loc, Loc, VD->getIdentifier(), Type,
198           C.getTrivialTypeSourceInfo(Type, SourceLocation()),
199           /*BW=*/nullptr, /*Mutable=*/false,
200           /*InitStyle=*/ICIS_NoInit);
201       Field->setAccess(AS_public);
202       llvm::APInt Align(32, std::max(C.getDeclAlign(VD).getQuantity(),
203                                      static_cast<CharUnits::QuantityType>(
204                                          GlobalMemoryAlignment)));
205       Field->addAttr(AlignedAttr::CreateImplicit(
206           C, /*IsAlignmentExpr=*/true,
207           IntegerLiteral::Create(C, Align,
208                                  C.getIntTypeForBitwidth(32, /*Signed=*/0),
209                                  SourceLocation()),
210           {}, AttributeCommonInfo::AS_GNU, AlignedAttr::GNU_aligned));
211     }
212     GlobalizedRD->addDecl(Field);
213     MappedDeclsFields.try_emplace(VD, Field);
214   }
215   GlobalizedRD->completeDefinition();
216   return GlobalizedRD;
217 }
218 
219 /// Get the list of variables that can escape their declaration context.
220 class CheckVarsEscapingDeclContext final
221     : public ConstStmtVisitor<CheckVarsEscapingDeclContext> {
222   CodeGenFunction &CGF;
223   llvm::SetVector<const ValueDecl *> EscapedDecls;
224   llvm::SetVector<const ValueDecl *> EscapedVariableLengthDecls;
225   llvm::SmallPtrSet<const Decl *, 4> EscapedParameters;
226   RecordDecl *GlobalizedRD = nullptr;
227   llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *> MappedDeclsFields;
228   bool AllEscaped = false;
229   bool IsForCombinedParallelRegion = false;
230 
231   void markAsEscaped(const ValueDecl *VD) {
232     // Do not globalize declare target variables.
233     if (!isa<VarDecl>(VD) ||
234         OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD))
235       return;
236     VD = cast<ValueDecl>(VD->getCanonicalDecl());
237     // Use user-specified allocation.
238     if (VD->hasAttrs() && VD->hasAttr<OMPAllocateDeclAttr>())
239       return;
240     // Variables captured by value must be globalized.
241     if (auto *CSI = CGF.CapturedStmtInfo) {
242       if (const FieldDecl *FD = CSI->lookup(cast<VarDecl>(VD))) {
243         // Check if need to capture the variable that was already captured by
244         // value in the outer region.
245         if (!IsForCombinedParallelRegion) {
246           if (!FD->hasAttrs())
247             return;
248           const auto *Attr = FD->getAttr<OMPCaptureKindAttr>();
249           if (!Attr)
250             return;
251           if (((Attr->getCaptureKind() != OMPC_map) &&
252                !isOpenMPPrivate(Attr->getCaptureKind())) ||
253               ((Attr->getCaptureKind() == OMPC_map) &&
254                !FD->getType()->isAnyPointerType()))
255             return;
256         }
257         if (!FD->getType()->isReferenceType()) {
258           assert(!VD->getType()->isVariablyModifiedType() &&
259                  "Parameter captured by value with variably modified type");
260           EscapedParameters.insert(VD);
261         } else if (!IsForCombinedParallelRegion) {
262           return;
263         }
264       }
265     }
266     if ((!CGF.CapturedStmtInfo ||
267          (IsForCombinedParallelRegion && CGF.CapturedStmtInfo)) &&
268         VD->getType()->isReferenceType())
269       // Do not globalize variables with reference type.
270       return;
271     if (VD->getType()->isVariablyModifiedType())
272       EscapedVariableLengthDecls.insert(VD);
273     else
274       EscapedDecls.insert(VD);
275   }
276 
277   void VisitValueDecl(const ValueDecl *VD) {
278     if (VD->getType()->isLValueReferenceType())
279       markAsEscaped(VD);
280     if (const auto *VarD = dyn_cast<VarDecl>(VD)) {
281       if (!isa<ParmVarDecl>(VarD) && VarD->hasInit()) {
282         const bool SavedAllEscaped = AllEscaped;
283         AllEscaped = VD->getType()->isLValueReferenceType();
284         Visit(VarD->getInit());
285         AllEscaped = SavedAllEscaped;
286       }
287     }
288   }
289   void VisitOpenMPCapturedStmt(const CapturedStmt *S,
290                                ArrayRef<OMPClause *> Clauses,
291                                bool IsCombinedParallelRegion) {
292     if (!S)
293       return;
294     for (const CapturedStmt::Capture &C : S->captures()) {
295       if (C.capturesVariable() && !C.capturesVariableByCopy()) {
296         const ValueDecl *VD = C.getCapturedVar();
297         bool SavedIsForCombinedParallelRegion = IsForCombinedParallelRegion;
298         if (IsCombinedParallelRegion) {
299           // Check if the variable is privatized in the combined construct and
300           // those private copies must be shared in the inner parallel
301           // directive.
302           IsForCombinedParallelRegion = false;
303           for (const OMPClause *C : Clauses) {
304             if (!isOpenMPPrivate(C->getClauseKind()) ||
305                 C->getClauseKind() == OMPC_reduction ||
306                 C->getClauseKind() == OMPC_linear ||
307                 C->getClauseKind() == OMPC_private)
308               continue;
309             ArrayRef<const Expr *> Vars;
310             if (const auto *PC = dyn_cast<OMPFirstprivateClause>(C))
311               Vars = PC->getVarRefs();
312             else if (const auto *PC = dyn_cast<OMPLastprivateClause>(C))
313               Vars = PC->getVarRefs();
314             else
315               llvm_unreachable("Unexpected clause.");
316             for (const auto *E : Vars) {
317               const Decl *D =
318                   cast<DeclRefExpr>(E)->getDecl()->getCanonicalDecl();
319               if (D == VD->getCanonicalDecl()) {
320                 IsForCombinedParallelRegion = true;
321                 break;
322               }
323             }
324             if (IsForCombinedParallelRegion)
325               break;
326           }
327         }
328         markAsEscaped(VD);
329         if (isa<OMPCapturedExprDecl>(VD))
330           VisitValueDecl(VD);
331         IsForCombinedParallelRegion = SavedIsForCombinedParallelRegion;
332       }
333     }
334   }
335 
336   void buildRecordForGlobalizedVars(bool IsInTTDRegion) {
337     assert(!GlobalizedRD &&
338            "Record for globalized variables is built already.");
339     ArrayRef<const ValueDecl *> EscapedDeclsForParallel, EscapedDeclsForTeams;
340     unsigned WarpSize = CGF.getTarget().getGridValue().GV_Warp_Size;
341     if (IsInTTDRegion)
342       EscapedDeclsForTeams = EscapedDecls.getArrayRef();
343     else
344       EscapedDeclsForParallel = EscapedDecls.getArrayRef();
345     GlobalizedRD = ::buildRecordForGlobalizedVars(
346         CGF.getContext(), EscapedDeclsForParallel, EscapedDeclsForTeams,
347         MappedDeclsFields, WarpSize);
348   }
349 
350 public:
351   CheckVarsEscapingDeclContext(CodeGenFunction &CGF,
352                                ArrayRef<const ValueDecl *> TeamsReductions)
353       : CGF(CGF), EscapedDecls(TeamsReductions.begin(), TeamsReductions.end()) {
354   }
355   virtual ~CheckVarsEscapingDeclContext() = default;
356   void VisitDeclStmt(const DeclStmt *S) {
357     if (!S)
358       return;
359     for (const Decl *D : S->decls())
360       if (const auto *VD = dyn_cast_or_null<ValueDecl>(D))
361         VisitValueDecl(VD);
362   }
363   void VisitOMPExecutableDirective(const OMPExecutableDirective *D) {
364     if (!D)
365       return;
366     if (!D->hasAssociatedStmt())
367       return;
368     if (const auto *S =
369             dyn_cast_or_null<CapturedStmt>(D->getAssociatedStmt())) {
370       // Do not analyze directives that do not actually require capturing,
371       // like `omp for` or `omp simd` directives.
372       llvm::SmallVector<OpenMPDirectiveKind, 4> CaptureRegions;
373       getOpenMPCaptureRegions(CaptureRegions, D->getDirectiveKind());
374       if (CaptureRegions.size() == 1 && CaptureRegions.back() == OMPD_unknown) {
375         VisitStmt(S->getCapturedStmt());
376         return;
377       }
378       VisitOpenMPCapturedStmt(
379           S, D->clauses(),
380           CaptureRegions.back() == OMPD_parallel &&
381               isOpenMPDistributeDirective(D->getDirectiveKind()));
382     }
383   }
384   void VisitCapturedStmt(const CapturedStmt *S) {
385     if (!S)
386       return;
387     for (const CapturedStmt::Capture &C : S->captures()) {
388       if (C.capturesVariable() && !C.capturesVariableByCopy()) {
389         const ValueDecl *VD = C.getCapturedVar();
390         markAsEscaped(VD);
391         if (isa<OMPCapturedExprDecl>(VD))
392           VisitValueDecl(VD);
393       }
394     }
395   }
396   void VisitLambdaExpr(const LambdaExpr *E) {
397     if (!E)
398       return;
399     for (const LambdaCapture &C : E->captures()) {
400       if (C.capturesVariable()) {
401         if (C.getCaptureKind() == LCK_ByRef) {
402           const ValueDecl *VD = C.getCapturedVar();
403           markAsEscaped(VD);
404           if (E->isInitCapture(&C) || isa<OMPCapturedExprDecl>(VD))
405             VisitValueDecl(VD);
406         }
407       }
408     }
409   }
410   void VisitBlockExpr(const BlockExpr *E) {
411     if (!E)
412       return;
413     for (const BlockDecl::Capture &C : E->getBlockDecl()->captures()) {
414       if (C.isByRef()) {
415         const VarDecl *VD = C.getVariable();
416         markAsEscaped(VD);
417         if (isa<OMPCapturedExprDecl>(VD) || VD->isInitCapture())
418           VisitValueDecl(VD);
419       }
420     }
421   }
422   void VisitCallExpr(const CallExpr *E) {
423     if (!E)
424       return;
425     for (const Expr *Arg : E->arguments()) {
426       if (!Arg)
427         continue;
428       if (Arg->isLValue()) {
429         const bool SavedAllEscaped = AllEscaped;
430         AllEscaped = true;
431         Visit(Arg);
432         AllEscaped = SavedAllEscaped;
433       } else {
434         Visit(Arg);
435       }
436     }
437     Visit(E->getCallee());
438   }
439   void VisitDeclRefExpr(const DeclRefExpr *E) {
440     if (!E)
441       return;
442     const ValueDecl *VD = E->getDecl();
443     if (AllEscaped)
444       markAsEscaped(VD);
445     if (isa<OMPCapturedExprDecl>(VD))
446       VisitValueDecl(VD);
447     else if (const auto *VarD = dyn_cast<VarDecl>(VD))
448       if (VarD->isInitCapture())
449         VisitValueDecl(VD);
450   }
451   void VisitUnaryOperator(const UnaryOperator *E) {
452     if (!E)
453       return;
454     if (E->getOpcode() == UO_AddrOf) {
455       const bool SavedAllEscaped = AllEscaped;
456       AllEscaped = true;
457       Visit(E->getSubExpr());
458       AllEscaped = SavedAllEscaped;
459     } else {
460       Visit(E->getSubExpr());
461     }
462   }
463   void VisitImplicitCastExpr(const ImplicitCastExpr *E) {
464     if (!E)
465       return;
466     if (E->getCastKind() == CK_ArrayToPointerDecay) {
467       const bool SavedAllEscaped = AllEscaped;
468       AllEscaped = true;
469       Visit(E->getSubExpr());
470       AllEscaped = SavedAllEscaped;
471     } else {
472       Visit(E->getSubExpr());
473     }
474   }
475   void VisitExpr(const Expr *E) {
476     if (!E)
477       return;
478     bool SavedAllEscaped = AllEscaped;
479     if (!E->isLValue())
480       AllEscaped = false;
481     for (const Stmt *Child : E->children())
482       if (Child)
483         Visit(Child);
484     AllEscaped = SavedAllEscaped;
485   }
486   void VisitStmt(const Stmt *S) {
487     if (!S)
488       return;
489     for (const Stmt *Child : S->children())
490       if (Child)
491         Visit(Child);
492   }
493 
494   /// Returns the record that handles all the escaped local variables and used
495   /// instead of their original storage.
496   const RecordDecl *getGlobalizedRecord(bool IsInTTDRegion) {
497     if (!GlobalizedRD)
498       buildRecordForGlobalizedVars(IsInTTDRegion);
499     return GlobalizedRD;
500   }
501 
502   /// Returns the field in the globalized record for the escaped variable.
503   const FieldDecl *getFieldForGlobalizedVar(const ValueDecl *VD) const {
504     assert(GlobalizedRD &&
505            "Record for globalized variables must be generated already.");
506     auto I = MappedDeclsFields.find(VD);
507     if (I == MappedDeclsFields.end())
508       return nullptr;
509     return I->getSecond();
510   }
511 
512   /// Returns the list of the escaped local variables/parameters.
513   ArrayRef<const ValueDecl *> getEscapedDecls() const {
514     return EscapedDecls.getArrayRef();
515   }
516 
517   /// Checks if the escaped local variable is actually a parameter passed by
518   /// value.
519   const llvm::SmallPtrSetImpl<const Decl *> &getEscapedParameters() const {
520     return EscapedParameters;
521   }
522 
523   /// Returns the list of the escaped variables with the variably modified
524   /// types.
525   ArrayRef<const ValueDecl *> getEscapedVariableLengthDecls() const {
526     return EscapedVariableLengthDecls.getArrayRef();
527   }
528 };
529 } // anonymous namespace
530 
531 /// Get the id of the warp in the block.
532 /// We assume that the warp size is 32, which is always the case
533 /// on the NVPTX device, to generate more efficient code.
534 static llvm::Value *getNVPTXWarpID(CodeGenFunction &CGF) {
535   CGBuilderTy &Bld = CGF.Builder;
536   unsigned LaneIDBits =
537       llvm::Log2_32(CGF.getTarget().getGridValue().GV_Warp_Size);
538   auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
539   return Bld.CreateAShr(RT.getGPUThreadID(CGF), LaneIDBits, "nvptx_warp_id");
540 }
541 
542 /// Get the id of the current lane in the Warp.
543 /// We assume that the warp size is 32, which is always the case
544 /// on the NVPTX device, to generate more efficient code.
545 static llvm::Value *getNVPTXLaneID(CodeGenFunction &CGF) {
546   CGBuilderTy &Bld = CGF.Builder;
547   unsigned LaneIDBits =
548       llvm::Log2_32(CGF.getTarget().getGridValue().GV_Warp_Size);
549   unsigned LaneIDMask = ~0u >> (32u - LaneIDBits);
550   auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
551   return Bld.CreateAnd(RT.getGPUThreadID(CGF), Bld.getInt32(LaneIDMask),
552                        "nvptx_lane_id");
553 }
554 
555 CGOpenMPRuntimeGPU::ExecutionMode
556 CGOpenMPRuntimeGPU::getExecutionMode() const {
557   return CurrentExecutionMode;
558 }
559 
560 static CGOpenMPRuntimeGPU::DataSharingMode
561 getDataSharingMode(CodeGenModule &CGM) {
562   return CGM.getLangOpts().OpenMPCUDAMode ? CGOpenMPRuntimeGPU::CUDA
563                                           : CGOpenMPRuntimeGPU::Generic;
564 }
565 
566 /// Check for inner (nested) SPMD construct, if any
567 static bool hasNestedSPMDDirective(ASTContext &Ctx,
568                                    const OMPExecutableDirective &D) {
569   const auto *CS = D.getInnermostCapturedStmt();
570   const auto *Body =
571       CS->getCapturedStmt()->IgnoreContainers(/*IgnoreCaptured=*/true);
572   const Stmt *ChildStmt = CGOpenMPRuntime::getSingleCompoundChild(Ctx, Body);
573 
574   if (const auto *NestedDir =
575           dyn_cast_or_null<OMPExecutableDirective>(ChildStmt)) {
576     OpenMPDirectiveKind DKind = NestedDir->getDirectiveKind();
577     switch (D.getDirectiveKind()) {
578     case OMPD_target:
579       if (isOpenMPParallelDirective(DKind))
580         return true;
581       if (DKind == OMPD_teams) {
582         Body = NestedDir->getInnermostCapturedStmt()->IgnoreContainers(
583             /*IgnoreCaptured=*/true);
584         if (!Body)
585           return false;
586         ChildStmt = CGOpenMPRuntime::getSingleCompoundChild(Ctx, Body);
587         if (const auto *NND =
588                 dyn_cast_or_null<OMPExecutableDirective>(ChildStmt)) {
589           DKind = NND->getDirectiveKind();
590           if (isOpenMPParallelDirective(DKind))
591             return true;
592         }
593       }
594       return false;
595     case OMPD_target_teams:
596       return isOpenMPParallelDirective(DKind);
597     case OMPD_target_simd:
598     case OMPD_target_parallel:
599     case OMPD_target_parallel_for:
600     case OMPD_target_parallel_for_simd:
601     case OMPD_target_teams_distribute:
602     case OMPD_target_teams_distribute_simd:
603     case OMPD_target_teams_distribute_parallel_for:
604     case OMPD_target_teams_distribute_parallel_for_simd:
605     case OMPD_parallel:
606     case OMPD_for:
607     case OMPD_parallel_for:
608     case OMPD_parallel_master:
609     case OMPD_parallel_sections:
610     case OMPD_for_simd:
611     case OMPD_parallel_for_simd:
612     case OMPD_cancel:
613     case OMPD_cancellation_point:
614     case OMPD_ordered:
615     case OMPD_threadprivate:
616     case OMPD_allocate:
617     case OMPD_task:
618     case OMPD_simd:
619     case OMPD_sections:
620     case OMPD_section:
621     case OMPD_single:
622     case OMPD_master:
623     case OMPD_critical:
624     case OMPD_taskyield:
625     case OMPD_barrier:
626     case OMPD_taskwait:
627     case OMPD_taskgroup:
628     case OMPD_atomic:
629     case OMPD_flush:
630     case OMPD_depobj:
631     case OMPD_scan:
632     case OMPD_teams:
633     case OMPD_target_data:
634     case OMPD_target_exit_data:
635     case OMPD_target_enter_data:
636     case OMPD_distribute:
637     case OMPD_distribute_simd:
638     case OMPD_distribute_parallel_for:
639     case OMPD_distribute_parallel_for_simd:
640     case OMPD_teams_distribute:
641     case OMPD_teams_distribute_simd:
642     case OMPD_teams_distribute_parallel_for:
643     case OMPD_teams_distribute_parallel_for_simd:
644     case OMPD_target_update:
645     case OMPD_declare_simd:
646     case OMPD_declare_variant:
647     case OMPD_begin_declare_variant:
648     case OMPD_end_declare_variant:
649     case OMPD_declare_target:
650     case OMPD_end_declare_target:
651     case OMPD_declare_reduction:
652     case OMPD_declare_mapper:
653     case OMPD_taskloop:
654     case OMPD_taskloop_simd:
655     case OMPD_master_taskloop:
656     case OMPD_master_taskloop_simd:
657     case OMPD_parallel_master_taskloop:
658     case OMPD_parallel_master_taskloop_simd:
659     case OMPD_requires:
660     case OMPD_unknown:
661     default:
662       llvm_unreachable("Unexpected directive.");
663     }
664   }
665 
666   return false;
667 }
668 
669 static bool supportsSPMDExecutionMode(ASTContext &Ctx,
670                                       const OMPExecutableDirective &D) {
671   OpenMPDirectiveKind DirectiveKind = D.getDirectiveKind();
672   switch (DirectiveKind) {
673   case OMPD_target:
674   case OMPD_target_teams:
675     return hasNestedSPMDDirective(Ctx, D);
676   case OMPD_target_parallel:
677   case OMPD_target_parallel_for:
678   case OMPD_target_parallel_for_simd:
679   case OMPD_target_teams_distribute_parallel_for:
680   case OMPD_target_teams_distribute_parallel_for_simd:
681   case OMPD_target_simd:
682   case OMPD_target_teams_distribute_simd:
683     return true;
684   case OMPD_target_teams_distribute:
685     return false;
686   case OMPD_parallel:
687   case OMPD_for:
688   case OMPD_parallel_for:
689   case OMPD_parallel_master:
690   case OMPD_parallel_sections:
691   case OMPD_for_simd:
692   case OMPD_parallel_for_simd:
693   case OMPD_cancel:
694   case OMPD_cancellation_point:
695   case OMPD_ordered:
696   case OMPD_threadprivate:
697   case OMPD_allocate:
698   case OMPD_task:
699   case OMPD_simd:
700   case OMPD_sections:
701   case OMPD_section:
702   case OMPD_single:
703   case OMPD_master:
704   case OMPD_critical:
705   case OMPD_taskyield:
706   case OMPD_barrier:
707   case OMPD_taskwait:
708   case OMPD_taskgroup:
709   case OMPD_atomic:
710   case OMPD_flush:
711   case OMPD_depobj:
712   case OMPD_scan:
713   case OMPD_teams:
714   case OMPD_target_data:
715   case OMPD_target_exit_data:
716   case OMPD_target_enter_data:
717   case OMPD_distribute:
718   case OMPD_distribute_simd:
719   case OMPD_distribute_parallel_for:
720   case OMPD_distribute_parallel_for_simd:
721   case OMPD_teams_distribute:
722   case OMPD_teams_distribute_simd:
723   case OMPD_teams_distribute_parallel_for:
724   case OMPD_teams_distribute_parallel_for_simd:
725   case OMPD_target_update:
726   case OMPD_declare_simd:
727   case OMPD_declare_variant:
728   case OMPD_begin_declare_variant:
729   case OMPD_end_declare_variant:
730   case OMPD_declare_target:
731   case OMPD_end_declare_target:
732   case OMPD_declare_reduction:
733   case OMPD_declare_mapper:
734   case OMPD_taskloop:
735   case OMPD_taskloop_simd:
736   case OMPD_master_taskloop:
737   case OMPD_master_taskloop_simd:
738   case OMPD_parallel_master_taskloop:
739   case OMPD_parallel_master_taskloop_simd:
740   case OMPD_requires:
741   case OMPD_unknown:
742   default:
743     break;
744   }
745   llvm_unreachable(
746       "Unknown programming model for OpenMP directive on NVPTX target.");
747 }
748 
749 /// Check if the directive is loops based and has schedule clause at all or has
750 /// static scheduling.
751 static bool hasStaticScheduling(const OMPExecutableDirective &D) {
752   assert(isOpenMPWorksharingDirective(D.getDirectiveKind()) &&
753          isOpenMPLoopDirective(D.getDirectiveKind()) &&
754          "Expected loop-based directive.");
755   return !D.hasClausesOfKind<OMPOrderedClause>() &&
756          (!D.hasClausesOfKind<OMPScheduleClause>() ||
757           llvm::any_of(D.getClausesOfKind<OMPScheduleClause>(),
758                        [](const OMPScheduleClause *C) {
759                          return C->getScheduleKind() == OMPC_SCHEDULE_static;
760                        }));
761 }
762 
763 /// Check for inner (nested) lightweight runtime construct, if any
764 static bool hasNestedLightweightDirective(ASTContext &Ctx,
765                                           const OMPExecutableDirective &D) {
766   assert(supportsSPMDExecutionMode(Ctx, D) && "Expected SPMD mode directive.");
767   const auto *CS = D.getInnermostCapturedStmt();
768   const auto *Body =
769       CS->getCapturedStmt()->IgnoreContainers(/*IgnoreCaptured=*/true);
770   const Stmt *ChildStmt = CGOpenMPRuntime::getSingleCompoundChild(Ctx, Body);
771 
772   if (const auto *NestedDir =
773           dyn_cast_or_null<OMPExecutableDirective>(ChildStmt)) {
774     OpenMPDirectiveKind DKind = NestedDir->getDirectiveKind();
775     switch (D.getDirectiveKind()) {
776     case OMPD_target:
777       if (isOpenMPParallelDirective(DKind) &&
778           isOpenMPWorksharingDirective(DKind) && isOpenMPLoopDirective(DKind) &&
779           hasStaticScheduling(*NestedDir))
780         return true;
781       if (DKind == OMPD_teams_distribute_simd || DKind == OMPD_simd)
782         return true;
783       if (DKind == OMPD_parallel) {
784         Body = NestedDir->getInnermostCapturedStmt()->IgnoreContainers(
785             /*IgnoreCaptured=*/true);
786         if (!Body)
787           return false;
788         ChildStmt = CGOpenMPRuntime::getSingleCompoundChild(Ctx, Body);
789         if (const auto *NND =
790                 dyn_cast_or_null<OMPExecutableDirective>(ChildStmt)) {
791           DKind = NND->getDirectiveKind();
792           if (isOpenMPWorksharingDirective(DKind) &&
793               isOpenMPLoopDirective(DKind) && hasStaticScheduling(*NND))
794             return true;
795         }
796       } else if (DKind == OMPD_teams) {
797         Body = NestedDir->getInnermostCapturedStmt()->IgnoreContainers(
798             /*IgnoreCaptured=*/true);
799         if (!Body)
800           return false;
801         ChildStmt = CGOpenMPRuntime::getSingleCompoundChild(Ctx, Body);
802         if (const auto *NND =
803                 dyn_cast_or_null<OMPExecutableDirective>(ChildStmt)) {
804           DKind = NND->getDirectiveKind();
805           if (isOpenMPParallelDirective(DKind) &&
806               isOpenMPWorksharingDirective(DKind) &&
807               isOpenMPLoopDirective(DKind) && hasStaticScheduling(*NND))
808             return true;
809           if (DKind == OMPD_parallel) {
810             Body = NND->getInnermostCapturedStmt()->IgnoreContainers(
811                 /*IgnoreCaptured=*/true);
812             if (!Body)
813               return false;
814             ChildStmt = CGOpenMPRuntime::getSingleCompoundChild(Ctx, Body);
815             if (const auto *NND =
816                     dyn_cast_or_null<OMPExecutableDirective>(ChildStmt)) {
817               DKind = NND->getDirectiveKind();
818               if (isOpenMPWorksharingDirective(DKind) &&
819                   isOpenMPLoopDirective(DKind) && hasStaticScheduling(*NND))
820                 return true;
821             }
822           }
823         }
824       }
825       return false;
826     case OMPD_target_teams:
827       if (isOpenMPParallelDirective(DKind) &&
828           isOpenMPWorksharingDirective(DKind) && isOpenMPLoopDirective(DKind) &&
829           hasStaticScheduling(*NestedDir))
830         return true;
831       if (DKind == OMPD_distribute_simd || DKind == OMPD_simd)
832         return true;
833       if (DKind == OMPD_parallel) {
834         Body = NestedDir->getInnermostCapturedStmt()->IgnoreContainers(
835             /*IgnoreCaptured=*/true);
836         if (!Body)
837           return false;
838         ChildStmt = CGOpenMPRuntime::getSingleCompoundChild(Ctx, Body);
839         if (const auto *NND =
840                 dyn_cast_or_null<OMPExecutableDirective>(ChildStmt)) {
841           DKind = NND->getDirectiveKind();
842           if (isOpenMPWorksharingDirective(DKind) &&
843               isOpenMPLoopDirective(DKind) && hasStaticScheduling(*NND))
844             return true;
845         }
846       }
847       return false;
848     case OMPD_target_parallel:
849       if (DKind == OMPD_simd)
850         return true;
851       return isOpenMPWorksharingDirective(DKind) &&
852              isOpenMPLoopDirective(DKind) && hasStaticScheduling(*NestedDir);
853     case OMPD_target_teams_distribute:
854     case OMPD_target_simd:
855     case OMPD_target_parallel_for:
856     case OMPD_target_parallel_for_simd:
857     case OMPD_target_teams_distribute_simd:
858     case OMPD_target_teams_distribute_parallel_for:
859     case OMPD_target_teams_distribute_parallel_for_simd:
860     case OMPD_parallel:
861     case OMPD_for:
862     case OMPD_parallel_for:
863     case OMPD_parallel_master:
864     case OMPD_parallel_sections:
865     case OMPD_for_simd:
866     case OMPD_parallel_for_simd:
867     case OMPD_cancel:
868     case OMPD_cancellation_point:
869     case OMPD_ordered:
870     case OMPD_threadprivate:
871     case OMPD_allocate:
872     case OMPD_task:
873     case OMPD_simd:
874     case OMPD_sections:
875     case OMPD_section:
876     case OMPD_single:
877     case OMPD_master:
878     case OMPD_critical:
879     case OMPD_taskyield:
880     case OMPD_barrier:
881     case OMPD_taskwait:
882     case OMPD_taskgroup:
883     case OMPD_atomic:
884     case OMPD_flush:
885     case OMPD_depobj:
886     case OMPD_scan:
887     case OMPD_teams:
888     case OMPD_target_data:
889     case OMPD_target_exit_data:
890     case OMPD_target_enter_data:
891     case OMPD_distribute:
892     case OMPD_distribute_simd:
893     case OMPD_distribute_parallel_for:
894     case OMPD_distribute_parallel_for_simd:
895     case OMPD_teams_distribute:
896     case OMPD_teams_distribute_simd:
897     case OMPD_teams_distribute_parallel_for:
898     case OMPD_teams_distribute_parallel_for_simd:
899     case OMPD_target_update:
900     case OMPD_declare_simd:
901     case OMPD_declare_variant:
902     case OMPD_begin_declare_variant:
903     case OMPD_end_declare_variant:
904     case OMPD_declare_target:
905     case OMPD_end_declare_target:
906     case OMPD_declare_reduction:
907     case OMPD_declare_mapper:
908     case OMPD_taskloop:
909     case OMPD_taskloop_simd:
910     case OMPD_master_taskloop:
911     case OMPD_master_taskloop_simd:
912     case OMPD_parallel_master_taskloop:
913     case OMPD_parallel_master_taskloop_simd:
914     case OMPD_requires:
915     case OMPD_unknown:
916     default:
917       llvm_unreachable("Unexpected directive.");
918     }
919   }
920 
921   return false;
922 }
923 
924 /// Checks if the construct supports lightweight runtime. It must be SPMD
925 /// construct + inner loop-based construct with static scheduling.
926 static bool supportsLightweightRuntime(ASTContext &Ctx,
927                                        const OMPExecutableDirective &D) {
928   if (!supportsSPMDExecutionMode(Ctx, D))
929     return false;
930   OpenMPDirectiveKind DirectiveKind = D.getDirectiveKind();
931   switch (DirectiveKind) {
932   case OMPD_target:
933   case OMPD_target_teams:
934   case OMPD_target_parallel:
935     return hasNestedLightweightDirective(Ctx, D);
936   case OMPD_target_parallel_for:
937   case OMPD_target_parallel_for_simd:
938   case OMPD_target_teams_distribute_parallel_for:
939   case OMPD_target_teams_distribute_parallel_for_simd:
940     // (Last|First)-privates must be shared in parallel region.
941     return hasStaticScheduling(D);
942   case OMPD_target_simd:
943   case OMPD_target_teams_distribute_simd:
944     return true;
945   case OMPD_target_teams_distribute:
946     return false;
947   case OMPD_parallel:
948   case OMPD_for:
949   case OMPD_parallel_for:
950   case OMPD_parallel_master:
951   case OMPD_parallel_sections:
952   case OMPD_for_simd:
953   case OMPD_parallel_for_simd:
954   case OMPD_cancel:
955   case OMPD_cancellation_point:
956   case OMPD_ordered:
957   case OMPD_threadprivate:
958   case OMPD_allocate:
959   case OMPD_task:
960   case OMPD_simd:
961   case OMPD_sections:
962   case OMPD_section:
963   case OMPD_single:
964   case OMPD_master:
965   case OMPD_critical:
966   case OMPD_taskyield:
967   case OMPD_barrier:
968   case OMPD_taskwait:
969   case OMPD_taskgroup:
970   case OMPD_atomic:
971   case OMPD_flush:
972   case OMPD_depobj:
973   case OMPD_scan:
974   case OMPD_teams:
975   case OMPD_target_data:
976   case OMPD_target_exit_data:
977   case OMPD_target_enter_data:
978   case OMPD_distribute:
979   case OMPD_distribute_simd:
980   case OMPD_distribute_parallel_for:
981   case OMPD_distribute_parallel_for_simd:
982   case OMPD_teams_distribute:
983   case OMPD_teams_distribute_simd:
984   case OMPD_teams_distribute_parallel_for:
985   case OMPD_teams_distribute_parallel_for_simd:
986   case OMPD_target_update:
987   case OMPD_declare_simd:
988   case OMPD_declare_variant:
989   case OMPD_begin_declare_variant:
990   case OMPD_end_declare_variant:
991   case OMPD_declare_target:
992   case OMPD_end_declare_target:
993   case OMPD_declare_reduction:
994   case OMPD_declare_mapper:
995   case OMPD_taskloop:
996   case OMPD_taskloop_simd:
997   case OMPD_master_taskloop:
998   case OMPD_master_taskloop_simd:
999   case OMPD_parallel_master_taskloop:
1000   case OMPD_parallel_master_taskloop_simd:
1001   case OMPD_requires:
1002   case OMPD_unknown:
1003   default:
1004     break;
1005   }
1006   llvm_unreachable(
1007       "Unknown programming model for OpenMP directive on NVPTX target.");
1008 }
1009 
1010 void CGOpenMPRuntimeGPU::emitNonSPMDKernel(const OMPExecutableDirective &D,
1011                                              StringRef ParentName,
1012                                              llvm::Function *&OutlinedFn,
1013                                              llvm::Constant *&OutlinedFnID,
1014                                              bool IsOffloadEntry,
1015                                              const RegionCodeGenTy &CodeGen) {
1016   ExecutionRuntimeModesRAII ModeRAII(CurrentExecutionMode);
1017   EntryFunctionState EST;
1018   WrapperFunctionsMap.clear();
1019 
1020   // Emit target region as a standalone region.
1021   class NVPTXPrePostActionTy : public PrePostActionTy {
1022     CGOpenMPRuntimeGPU::EntryFunctionState &EST;
1023 
1024   public:
1025     NVPTXPrePostActionTy(CGOpenMPRuntimeGPU::EntryFunctionState &EST)
1026         : EST(EST) {}
1027     void Enter(CodeGenFunction &CGF) override {
1028       auto &RT =
1029           static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
1030       RT.emitKernelInit(CGF, EST, /* IsSPMD */ false);
1031       // Skip target region initialization.
1032       RT.setLocThreadIdInsertPt(CGF, /*AtCurrentPoint=*/true);
1033     }
1034     void Exit(CodeGenFunction &CGF) override {
1035       auto &RT =
1036           static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
1037       RT.clearLocThreadIdInsertPt(CGF);
1038       RT.emitKernelDeinit(CGF, EST, /* IsSPMD */ false);
1039     }
1040   } Action(EST);
1041   CodeGen.setAction(Action);
1042   IsInTTDRegion = true;
1043   emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,
1044                                    IsOffloadEntry, CodeGen);
1045   IsInTTDRegion = false;
1046 }
1047 
1048 void CGOpenMPRuntimeGPU::emitKernelInit(CodeGenFunction &CGF,
1049                                         EntryFunctionState &EST, bool IsSPMD) {
1050   CGBuilderTy &Bld = CGF.Builder;
1051   Bld.restoreIP(OMPBuilder.createTargetInit(Bld, IsSPMD, requiresFullRuntime()));
1052   IsInTargetMasterThreadRegion = IsSPMD;
1053   if (!IsSPMD)
1054     emitGenericVarsProlog(CGF, EST.Loc);
1055 }
1056 
1057 void CGOpenMPRuntimeGPU::emitKernelDeinit(CodeGenFunction &CGF,
1058                                           EntryFunctionState &EST,
1059                                           bool IsSPMD) {
1060   if (!IsSPMD)
1061     emitGenericVarsEpilog(CGF);
1062 
1063   CGBuilderTy &Bld = CGF.Builder;
1064   OMPBuilder.createTargetDeinit(Bld, IsSPMD, requiresFullRuntime());
1065 }
1066 
1067 void CGOpenMPRuntimeGPU::emitSPMDKernel(const OMPExecutableDirective &D,
1068                                           StringRef ParentName,
1069                                           llvm::Function *&OutlinedFn,
1070                                           llvm::Constant *&OutlinedFnID,
1071                                           bool IsOffloadEntry,
1072                                           const RegionCodeGenTy &CodeGen) {
1073   ExecutionRuntimeModesRAII ModeRAII(
1074       CurrentExecutionMode, RequiresFullRuntime,
1075       CGM.getLangOpts().OpenMPCUDAForceFullRuntime ||
1076           !supportsLightweightRuntime(CGM.getContext(), D));
1077   EntryFunctionState EST;
1078 
1079   // Emit target region as a standalone region.
1080   class NVPTXPrePostActionTy : public PrePostActionTy {
1081     CGOpenMPRuntimeGPU &RT;
1082     CGOpenMPRuntimeGPU::EntryFunctionState &EST;
1083 
1084   public:
1085     NVPTXPrePostActionTy(CGOpenMPRuntimeGPU &RT,
1086                          CGOpenMPRuntimeGPU::EntryFunctionState &EST)
1087         : RT(RT), EST(EST) {}
1088     void Enter(CodeGenFunction &CGF) override {
1089       RT.emitKernelInit(CGF, EST, /* IsSPMD */ true);
1090       // Skip target region initialization.
1091       RT.setLocThreadIdInsertPt(CGF, /*AtCurrentPoint=*/true);
1092     }
1093     void Exit(CodeGenFunction &CGF) override {
1094       RT.clearLocThreadIdInsertPt(CGF);
1095       RT.emitKernelDeinit(CGF, EST, /* IsSPMD */ true);
1096     }
1097   } Action(*this, EST);
1098   CodeGen.setAction(Action);
1099   IsInTTDRegion = true;
1100   emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,
1101                                    IsOffloadEntry, CodeGen);
1102   IsInTTDRegion = false;
1103 }
1104 
1105 // Create a unique global variable to indicate the execution mode of this target
1106 // region. The execution mode is either 'generic', or 'spmd' depending on the
1107 // target directive. This variable is picked up by the offload library to setup
1108 // the device appropriately before kernel launch. If the execution mode is
1109 // 'generic', the runtime reserves one warp for the master, otherwise, all
1110 // warps participate in parallel work.
1111 static void setPropertyExecutionMode(CodeGenModule &CGM, StringRef Name,
1112                                      bool Mode) {
1113   auto *GVMode = new llvm::GlobalVariable(
1114       CGM.getModule(), CGM.Int8Ty, /*isConstant=*/true,
1115       llvm::GlobalValue::WeakAnyLinkage,
1116       llvm::ConstantInt::get(CGM.Int8Ty, Mode ? OMP_TGT_EXEC_MODE_SPMD
1117                                               : OMP_TGT_EXEC_MODE_GENERIC),
1118       Twine(Name, "_exec_mode"));
1119   CGM.addCompilerUsedGlobal(GVMode);
1120 }
1121 
1122 void CGOpenMPRuntimeGPU::createOffloadEntry(llvm::Constant *ID,
1123                                               llvm::Constant *Addr,
1124                                               uint64_t Size, int32_t,
1125                                               llvm::GlobalValue::LinkageTypes) {
1126   // TODO: Add support for global variables on the device after declare target
1127   // support.
1128   llvm::Function *Fn = dyn_cast<llvm::Function>(Addr);
1129   if (!Fn)
1130     return;
1131 
1132   llvm::Module &M = CGM.getModule();
1133   llvm::LLVMContext &Ctx = CGM.getLLVMContext();
1134 
1135   // Get "nvvm.annotations" metadata node.
1136   llvm::NamedMDNode *MD = M.getOrInsertNamedMetadata("nvvm.annotations");
1137 
1138   llvm::Metadata *MDVals[] = {
1139       llvm::ConstantAsMetadata::get(Fn), llvm::MDString::get(Ctx, "kernel"),
1140       llvm::ConstantAsMetadata::get(
1141           llvm::ConstantInt::get(llvm::Type::getInt32Ty(Ctx), 1))};
1142   // Append metadata to nvvm.annotations.
1143   MD->addOperand(llvm::MDNode::get(Ctx, MDVals));
1144 
1145   // Add a function attribute for the kernel.
1146   Fn->addFnAttr(llvm::Attribute::get(Ctx, "kernel"));
1147 }
1148 
1149 void CGOpenMPRuntimeGPU::emitTargetOutlinedFunction(
1150     const OMPExecutableDirective &D, StringRef ParentName,
1151     llvm::Function *&OutlinedFn, llvm::Constant *&OutlinedFnID,
1152     bool IsOffloadEntry, const RegionCodeGenTy &CodeGen) {
1153   if (!IsOffloadEntry) // Nothing to do.
1154     return;
1155 
1156   assert(!ParentName.empty() && "Invalid target region parent name!");
1157 
1158   bool Mode = supportsSPMDExecutionMode(CGM.getContext(), D);
1159   if (Mode)
1160     emitSPMDKernel(D, ParentName, OutlinedFn, OutlinedFnID, IsOffloadEntry,
1161                    CodeGen);
1162   else
1163     emitNonSPMDKernel(D, ParentName, OutlinedFn, OutlinedFnID, IsOffloadEntry,
1164                       CodeGen);
1165 
1166   setPropertyExecutionMode(CGM, OutlinedFn->getName(), Mode);
1167 }
1168 
1169 namespace {
1170 LLVM_ENABLE_BITMASK_ENUMS_IN_NAMESPACE();
1171 /// Enum for accesseing the reserved_2 field of the ident_t struct.
1172 enum ModeFlagsTy : unsigned {
1173   /// Bit set to 1 when in SPMD mode.
1174   KMP_IDENT_SPMD_MODE = 0x01,
1175   /// Bit set to 1 when a simplified runtime is used.
1176   KMP_IDENT_SIMPLE_RT_MODE = 0x02,
1177   LLVM_MARK_AS_BITMASK_ENUM(/*LargestValue=*/KMP_IDENT_SIMPLE_RT_MODE)
1178 };
1179 
1180 /// Special mode Undefined. Is the combination of Non-SPMD mode + SimpleRuntime.
1181 static const ModeFlagsTy UndefinedMode =
1182     (~KMP_IDENT_SPMD_MODE) & KMP_IDENT_SIMPLE_RT_MODE;
1183 } // anonymous namespace
1184 
1185 unsigned CGOpenMPRuntimeGPU::getDefaultLocationReserved2Flags() const {
1186   switch (getExecutionMode()) {
1187   case EM_SPMD:
1188     if (requiresFullRuntime())
1189       return KMP_IDENT_SPMD_MODE & (~KMP_IDENT_SIMPLE_RT_MODE);
1190     return KMP_IDENT_SPMD_MODE | KMP_IDENT_SIMPLE_RT_MODE;
1191   case EM_NonSPMD:
1192     assert(requiresFullRuntime() && "Expected full runtime.");
1193     return (~KMP_IDENT_SPMD_MODE) & (~KMP_IDENT_SIMPLE_RT_MODE);
1194   case EM_Unknown:
1195     return UndefinedMode;
1196   }
1197   llvm_unreachable("Unknown flags are requested.");
1198 }
1199 
1200 CGOpenMPRuntimeGPU::CGOpenMPRuntimeGPU(CodeGenModule &CGM)
1201     : CGOpenMPRuntime(CGM, "_", "$") {
1202   if (!CGM.getLangOpts().OpenMPIsDevice)
1203     llvm_unreachable("OpenMP can only handle device code.");
1204 
1205   llvm::OpenMPIRBuilder &OMPBuilder = getOMPBuilder();
1206   if (!CGM.getLangOpts().OMPHostIRFile.empty()) {
1207     OMPBuilder.createGlobalFlag(CGM.getLangOpts().OpenMPTargetDebug,
1208                                 "__omp_rtl_debug_kind");
1209     OMPBuilder.createGlobalFlag(CGM.getLangOpts().OpenMPTeamSubscription,
1210                                 "__omp_rtl_assume_teams_oversubscription");
1211     OMPBuilder.createGlobalFlag(CGM.getLangOpts().OpenMPThreadSubscription,
1212                                 "__omp_rtl_assume_threads_oversubscription");
1213   }
1214 }
1215 
1216 void CGOpenMPRuntimeGPU::emitProcBindClause(CodeGenFunction &CGF,
1217                                               ProcBindKind ProcBind,
1218                                               SourceLocation Loc) {
1219   // Do nothing in case of SPMD mode and L0 parallel.
1220   if (getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD)
1221     return;
1222 
1223   CGOpenMPRuntime::emitProcBindClause(CGF, ProcBind, Loc);
1224 }
1225 
1226 void CGOpenMPRuntimeGPU::emitNumThreadsClause(CodeGenFunction &CGF,
1227                                                 llvm::Value *NumThreads,
1228                                                 SourceLocation Loc) {
1229   // Nothing to do.
1230 }
1231 
1232 void CGOpenMPRuntimeGPU::emitNumTeamsClause(CodeGenFunction &CGF,
1233                                               const Expr *NumTeams,
1234                                               const Expr *ThreadLimit,
1235                                               SourceLocation Loc) {}
1236 
1237 llvm::Function *CGOpenMPRuntimeGPU::emitParallelOutlinedFunction(
1238     const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
1239     OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen) {
1240   // Emit target region as a standalone region.
1241   class NVPTXPrePostActionTy : public PrePostActionTy {
1242     bool &IsInParallelRegion;
1243     bool PrevIsInParallelRegion;
1244 
1245   public:
1246     NVPTXPrePostActionTy(bool &IsInParallelRegion)
1247         : IsInParallelRegion(IsInParallelRegion) {}
1248     void Enter(CodeGenFunction &CGF) override {
1249       PrevIsInParallelRegion = IsInParallelRegion;
1250       IsInParallelRegion = true;
1251     }
1252     void Exit(CodeGenFunction &CGF) override {
1253       IsInParallelRegion = PrevIsInParallelRegion;
1254     }
1255   } Action(IsInParallelRegion);
1256   CodeGen.setAction(Action);
1257   bool PrevIsInTTDRegion = IsInTTDRegion;
1258   IsInTTDRegion = false;
1259   bool PrevIsInTargetMasterThreadRegion = IsInTargetMasterThreadRegion;
1260   IsInTargetMasterThreadRegion = false;
1261   auto *OutlinedFun =
1262       cast<llvm::Function>(CGOpenMPRuntime::emitParallelOutlinedFunction(
1263           D, ThreadIDVar, InnermostKind, CodeGen));
1264   IsInTargetMasterThreadRegion = PrevIsInTargetMasterThreadRegion;
1265   IsInTTDRegion = PrevIsInTTDRegion;
1266   if (getExecutionMode() != CGOpenMPRuntimeGPU::EM_SPMD &&
1267       !IsInParallelRegion) {
1268     llvm::Function *WrapperFun =
1269         createParallelDataSharingWrapper(OutlinedFun, D);
1270     WrapperFunctionsMap[OutlinedFun] = WrapperFun;
1271   }
1272 
1273   return OutlinedFun;
1274 }
1275 
1276 /// Get list of lastprivate variables from the teams distribute ... or
1277 /// teams {distribute ...} directives.
1278 static void
1279 getDistributeLastprivateVars(ASTContext &Ctx, const OMPExecutableDirective &D,
1280                              llvm::SmallVectorImpl<const ValueDecl *> &Vars) {
1281   assert(isOpenMPTeamsDirective(D.getDirectiveKind()) &&
1282          "expected teams directive.");
1283   const OMPExecutableDirective *Dir = &D;
1284   if (!isOpenMPDistributeDirective(D.getDirectiveKind())) {
1285     if (const Stmt *S = CGOpenMPRuntime::getSingleCompoundChild(
1286             Ctx,
1287             D.getInnermostCapturedStmt()->getCapturedStmt()->IgnoreContainers(
1288                 /*IgnoreCaptured=*/true))) {
1289       Dir = dyn_cast_or_null<OMPExecutableDirective>(S);
1290       if (Dir && !isOpenMPDistributeDirective(Dir->getDirectiveKind()))
1291         Dir = nullptr;
1292     }
1293   }
1294   if (!Dir)
1295     return;
1296   for (const auto *C : Dir->getClausesOfKind<OMPLastprivateClause>()) {
1297     for (const Expr *E : C->getVarRefs())
1298       Vars.push_back(getPrivateItem(E));
1299   }
1300 }
1301 
1302 /// Get list of reduction variables from the teams ... directives.
1303 static void
1304 getTeamsReductionVars(ASTContext &Ctx, const OMPExecutableDirective &D,
1305                       llvm::SmallVectorImpl<const ValueDecl *> &Vars) {
1306   assert(isOpenMPTeamsDirective(D.getDirectiveKind()) &&
1307          "expected teams directive.");
1308   for (const auto *C : D.getClausesOfKind<OMPReductionClause>()) {
1309     for (const Expr *E : C->privates())
1310       Vars.push_back(getPrivateItem(E));
1311   }
1312 }
1313 
1314 llvm::Function *CGOpenMPRuntimeGPU::emitTeamsOutlinedFunction(
1315     const OMPExecutableDirective &D, const VarDecl *ThreadIDVar,
1316     OpenMPDirectiveKind InnermostKind, const RegionCodeGenTy &CodeGen) {
1317   SourceLocation Loc = D.getBeginLoc();
1318 
1319   const RecordDecl *GlobalizedRD = nullptr;
1320   llvm::SmallVector<const ValueDecl *, 4> LastPrivatesReductions;
1321   llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *> MappedDeclsFields;
1322   unsigned WarpSize = CGM.getTarget().getGridValue().GV_Warp_Size;
1323   // Globalize team reductions variable unconditionally in all modes.
1324   if (getExecutionMode() != CGOpenMPRuntimeGPU::EM_SPMD)
1325     getTeamsReductionVars(CGM.getContext(), D, LastPrivatesReductions);
1326   if (getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD) {
1327     getDistributeLastprivateVars(CGM.getContext(), D, LastPrivatesReductions);
1328     if (!LastPrivatesReductions.empty()) {
1329       GlobalizedRD = ::buildRecordForGlobalizedVars(
1330           CGM.getContext(), llvm::None, LastPrivatesReductions,
1331           MappedDeclsFields, WarpSize);
1332     }
1333   } else if (!LastPrivatesReductions.empty()) {
1334     assert(!TeamAndReductions.first &&
1335            "Previous team declaration is not expected.");
1336     TeamAndReductions.first = D.getCapturedStmt(OMPD_teams)->getCapturedDecl();
1337     std::swap(TeamAndReductions.second, LastPrivatesReductions);
1338   }
1339 
1340   // Emit target region as a standalone region.
1341   class NVPTXPrePostActionTy : public PrePostActionTy {
1342     SourceLocation &Loc;
1343     const RecordDecl *GlobalizedRD;
1344     llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>
1345         &MappedDeclsFields;
1346 
1347   public:
1348     NVPTXPrePostActionTy(
1349         SourceLocation &Loc, const RecordDecl *GlobalizedRD,
1350         llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>
1351             &MappedDeclsFields)
1352         : Loc(Loc), GlobalizedRD(GlobalizedRD),
1353           MappedDeclsFields(MappedDeclsFields) {}
1354     void Enter(CodeGenFunction &CGF) override {
1355       auto &Rt =
1356           static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
1357       if (GlobalizedRD) {
1358         auto I = Rt.FunctionGlobalizedDecls.try_emplace(CGF.CurFn).first;
1359         I->getSecond().MappedParams =
1360             std::make_unique<CodeGenFunction::OMPMapVars>();
1361         DeclToAddrMapTy &Data = I->getSecond().LocalVarData;
1362         for (const auto &Pair : MappedDeclsFields) {
1363           assert(Pair.getFirst()->isCanonicalDecl() &&
1364                  "Expected canonical declaration");
1365           Data.insert(std::make_pair(Pair.getFirst(), MappedVarData()));
1366         }
1367       }
1368       Rt.emitGenericVarsProlog(CGF, Loc);
1369     }
1370     void Exit(CodeGenFunction &CGF) override {
1371       static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime())
1372           .emitGenericVarsEpilog(CGF);
1373     }
1374   } Action(Loc, GlobalizedRD, MappedDeclsFields);
1375   CodeGen.setAction(Action);
1376   llvm::Function *OutlinedFun = CGOpenMPRuntime::emitTeamsOutlinedFunction(
1377       D, ThreadIDVar, InnermostKind, CodeGen);
1378 
1379   return OutlinedFun;
1380 }
1381 
1382 void CGOpenMPRuntimeGPU::emitGenericVarsProlog(CodeGenFunction &CGF,
1383                                                  SourceLocation Loc,
1384                                                  bool WithSPMDCheck) {
1385   if (getDataSharingMode(CGM) != CGOpenMPRuntimeGPU::Generic &&
1386       getExecutionMode() != CGOpenMPRuntimeGPU::EM_SPMD)
1387     return;
1388 
1389   CGBuilderTy &Bld = CGF.Builder;
1390 
1391   const auto I = FunctionGlobalizedDecls.find(CGF.CurFn);
1392   if (I == FunctionGlobalizedDecls.end())
1393     return;
1394 
1395   for (auto &Rec : I->getSecond().LocalVarData) {
1396     const auto *VD = cast<VarDecl>(Rec.first);
1397     bool EscapedParam = I->getSecond().EscapedParameters.count(Rec.first);
1398     QualType VarTy = VD->getType();
1399 
1400     // Get the local allocation of a firstprivate variable before sharing
1401     llvm::Value *ParValue;
1402     if (EscapedParam) {
1403       LValue ParLVal =
1404           CGF.MakeAddrLValue(CGF.GetAddrOfLocalVar(VD), VD->getType());
1405       ParValue = CGF.EmitLoadOfScalar(ParLVal, Loc);
1406     }
1407 
1408     // Allocate space for the variable to be globalized
1409     llvm::Value *AllocArgs[] = {CGF.getTypeSize(VD->getType())};
1410     llvm::CallBase *VoidPtr =
1411         CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
1412                                 CGM.getModule(), OMPRTL___kmpc_alloc_shared),
1413                             AllocArgs, VD->getName());
1414     // FIXME: We should use the variables actual alignment as an argument.
1415     VoidPtr->addRetAttr(llvm::Attribute::get(
1416         CGM.getLLVMContext(), llvm::Attribute::Alignment,
1417         CGM.getContext().getTargetInfo().getNewAlign() / 8));
1418 
1419     // Cast the void pointer and get the address of the globalized variable.
1420     llvm::PointerType *VarPtrTy = CGF.ConvertTypeForMem(VarTy)->getPointerTo();
1421     llvm::Value *CastedVoidPtr = Bld.CreatePointerBitCastOrAddrSpaceCast(
1422         VoidPtr, VarPtrTy, VD->getName() + "_on_stack");
1423     LValue VarAddr = CGF.MakeNaturalAlignAddrLValue(CastedVoidPtr, VarTy);
1424     Rec.second.PrivateAddr = VarAddr.getAddress(CGF);
1425     Rec.second.GlobalizedVal = VoidPtr;
1426 
1427     // Assign the local allocation to the newly globalized location.
1428     if (EscapedParam) {
1429       CGF.EmitStoreOfScalar(ParValue, VarAddr);
1430       I->getSecond().MappedParams->setVarAddr(CGF, VD, VarAddr.getAddress(CGF));
1431     }
1432     if (auto *DI = CGF.getDebugInfo())
1433       VoidPtr->setDebugLoc(DI->SourceLocToDebugLoc(VD->getLocation()));
1434   }
1435   for (const auto *VD : I->getSecond().EscapedVariableLengthDecls) {
1436     // Use actual memory size of the VLA object including the padding
1437     // for alignment purposes.
1438     llvm::Value *Size = CGF.getTypeSize(VD->getType());
1439     CharUnits Align = CGM.getContext().getDeclAlign(VD);
1440     Size = Bld.CreateNUWAdd(
1441         Size, llvm::ConstantInt::get(CGF.SizeTy, Align.getQuantity() - 1));
1442     llvm::Value *AlignVal =
1443         llvm::ConstantInt::get(CGF.SizeTy, Align.getQuantity());
1444 
1445     Size = Bld.CreateUDiv(Size, AlignVal);
1446     Size = Bld.CreateNUWMul(Size, AlignVal);
1447 
1448     // Allocate space for this VLA object to be globalized.
1449     llvm::Value *AllocArgs[] = {CGF.getTypeSize(VD->getType())};
1450     llvm::CallBase *VoidPtr =
1451         CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
1452                                 CGM.getModule(), OMPRTL___kmpc_alloc_shared),
1453                             AllocArgs, VD->getName());
1454     VoidPtr->addRetAttr(
1455         llvm::Attribute::get(CGM.getLLVMContext(), llvm::Attribute::Alignment,
1456                              CGM.getContext().getTargetInfo().getNewAlign()));
1457 
1458     I->getSecond().EscapedVariableLengthDeclsAddrs.emplace_back(
1459         std::pair<llvm::Value *, llvm::Value *>(
1460             {VoidPtr, CGF.getTypeSize(VD->getType())}));
1461     LValue Base = CGF.MakeAddrLValue(VoidPtr, VD->getType(),
1462                                      CGM.getContext().getDeclAlign(VD),
1463                                      AlignmentSource::Decl);
1464     I->getSecond().MappedParams->setVarAddr(CGF, cast<VarDecl>(VD),
1465                                             Base.getAddress(CGF));
1466   }
1467   I->getSecond().MappedParams->apply(CGF);
1468 }
1469 
1470 void CGOpenMPRuntimeGPU::emitGenericVarsEpilog(CodeGenFunction &CGF,
1471                                                  bool WithSPMDCheck) {
1472   if (getDataSharingMode(CGM) != CGOpenMPRuntimeGPU::Generic &&
1473       getExecutionMode() != CGOpenMPRuntimeGPU::EM_SPMD)
1474     return;
1475 
1476   const auto I = FunctionGlobalizedDecls.find(CGF.CurFn);
1477   if (I != FunctionGlobalizedDecls.end()) {
1478     // Deallocate the memory for each globalized VLA object
1479     for (auto AddrSizePair :
1480          llvm::reverse(I->getSecond().EscapedVariableLengthDeclsAddrs)) {
1481       CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
1482                               CGM.getModule(), OMPRTL___kmpc_free_shared),
1483                           {AddrSizePair.first, AddrSizePair.second});
1484     }
1485     // Deallocate the memory for each globalized value
1486     for (auto &Rec : llvm::reverse(I->getSecond().LocalVarData)) {
1487       const auto *VD = cast<VarDecl>(Rec.first);
1488       I->getSecond().MappedParams->restore(CGF);
1489 
1490       llvm::Value *FreeArgs[] = {Rec.second.GlobalizedVal,
1491                                  CGF.getTypeSize(VD->getType())};
1492       CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
1493                               CGM.getModule(), OMPRTL___kmpc_free_shared),
1494                           FreeArgs);
1495     }
1496   }
1497 }
1498 
1499 void CGOpenMPRuntimeGPU::emitTeamsCall(CodeGenFunction &CGF,
1500                                          const OMPExecutableDirective &D,
1501                                          SourceLocation Loc,
1502                                          llvm::Function *OutlinedFn,
1503                                          ArrayRef<llvm::Value *> CapturedVars) {
1504   if (!CGF.HaveInsertPoint())
1505     return;
1506 
1507   Address ZeroAddr = CGF.CreateDefaultAlignTempAlloca(CGF.Int32Ty,
1508                                                       /*Name=*/".zero.addr");
1509   CGF.Builder.CreateStore(CGF.Builder.getInt32(/*C*/ 0), ZeroAddr);
1510   llvm::SmallVector<llvm::Value *, 16> OutlinedFnArgs;
1511   OutlinedFnArgs.push_back(emitThreadIDAddress(CGF, Loc).getPointer());
1512   OutlinedFnArgs.push_back(ZeroAddr.getPointer());
1513   OutlinedFnArgs.append(CapturedVars.begin(), CapturedVars.end());
1514   emitOutlinedFunctionCall(CGF, Loc, OutlinedFn, OutlinedFnArgs);
1515 }
1516 
1517 void CGOpenMPRuntimeGPU::emitParallelCall(CodeGenFunction &CGF,
1518                                           SourceLocation Loc,
1519                                           llvm::Function *OutlinedFn,
1520                                           ArrayRef<llvm::Value *> CapturedVars,
1521                                           const Expr *IfCond,
1522                                           llvm::Value *NumThreads) {
1523   if (!CGF.HaveInsertPoint())
1524     return;
1525 
1526   auto &&ParallelGen = [this, Loc, OutlinedFn, CapturedVars, IfCond,
1527                         NumThreads](CodeGenFunction &CGF,
1528                                     PrePostActionTy &Action) {
1529     CGBuilderTy &Bld = CGF.Builder;
1530     llvm::Value *NumThreadsVal = NumThreads;
1531     llvm::Function *WFn = WrapperFunctionsMap[OutlinedFn];
1532     llvm::Value *ID = llvm::ConstantPointerNull::get(CGM.Int8PtrTy);
1533     if (WFn)
1534       ID = Bld.CreateBitOrPointerCast(WFn, CGM.Int8PtrTy);
1535     llvm::Value *FnPtr = Bld.CreateBitOrPointerCast(OutlinedFn, CGM.Int8PtrTy);
1536 
1537     // Create a private scope that will globalize the arguments
1538     // passed from the outside of the target region.
1539     // TODO: Is that needed?
1540     CodeGenFunction::OMPPrivateScope PrivateArgScope(CGF);
1541 
1542     Address CapturedVarsAddrs = CGF.CreateDefaultAlignTempAlloca(
1543         llvm::ArrayType::get(CGM.VoidPtrTy, CapturedVars.size()),
1544         "captured_vars_addrs");
1545     // There's something to share.
1546     if (!CapturedVars.empty()) {
1547       // Prepare for parallel region. Indicate the outlined function.
1548       ASTContext &Ctx = CGF.getContext();
1549       unsigned Idx = 0;
1550       for (llvm::Value *V : CapturedVars) {
1551         Address Dst = Bld.CreateConstArrayGEP(CapturedVarsAddrs, Idx);
1552         llvm::Value *PtrV;
1553         if (V->getType()->isIntegerTy())
1554           PtrV = Bld.CreateIntToPtr(V, CGF.VoidPtrTy);
1555         else
1556           PtrV = Bld.CreatePointerBitCastOrAddrSpaceCast(V, CGF.VoidPtrTy);
1557         CGF.EmitStoreOfScalar(PtrV, Dst, /*Volatile=*/false,
1558                               Ctx.getPointerType(Ctx.VoidPtrTy));
1559         ++Idx;
1560       }
1561     }
1562 
1563     llvm::Value *IfCondVal = nullptr;
1564     if (IfCond)
1565       IfCondVal = Bld.CreateIntCast(CGF.EvaluateExprAsBool(IfCond), CGF.Int32Ty,
1566                                     /* isSigned */ false);
1567     else
1568       IfCondVal = llvm::ConstantInt::get(CGF.Int32Ty, 1);
1569 
1570     if (!NumThreadsVal)
1571       NumThreadsVal = llvm::ConstantInt::get(CGF.Int32Ty, -1);
1572     else
1573       NumThreadsVal = Bld.CreateZExtOrTrunc(NumThreadsVal, CGF.Int32Ty),
1574 
1575       assert(IfCondVal && "Expected a value");
1576     llvm::Value *RTLoc = emitUpdateLocation(CGF, Loc);
1577     llvm::Value *Args[] = {
1578         RTLoc,
1579         getThreadID(CGF, Loc),
1580         IfCondVal,
1581         NumThreadsVal,
1582         llvm::ConstantInt::get(CGF.Int32Ty, -1),
1583         FnPtr,
1584         ID,
1585         Bld.CreateBitOrPointerCast(CapturedVarsAddrs.getPointer(),
1586                                    CGF.VoidPtrPtrTy),
1587         llvm::ConstantInt::get(CGM.SizeTy, CapturedVars.size())};
1588     CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
1589                             CGM.getModule(), OMPRTL___kmpc_parallel_51),
1590                         Args);
1591   };
1592 
1593   RegionCodeGenTy RCG(ParallelGen);
1594   RCG(CGF);
1595 }
1596 
1597 void CGOpenMPRuntimeGPU::syncCTAThreads(CodeGenFunction &CGF) {
1598   // Always emit simple barriers!
1599   if (!CGF.HaveInsertPoint())
1600     return;
1601   // Build call __kmpc_barrier_simple_spmd(nullptr, 0);
1602   // This function does not use parameters, so we can emit just default values.
1603   llvm::Value *Args[] = {
1604       llvm::ConstantPointerNull::get(
1605           cast<llvm::PointerType>(getIdentTyPointerTy())),
1606       llvm::ConstantInt::get(CGF.Int32Ty, /*V=*/0, /*isSigned=*/true)};
1607   CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
1608                           CGM.getModule(), OMPRTL___kmpc_barrier_simple_spmd),
1609                       Args);
1610 }
1611 
1612 void CGOpenMPRuntimeGPU::emitBarrierCall(CodeGenFunction &CGF,
1613                                            SourceLocation Loc,
1614                                            OpenMPDirectiveKind Kind, bool,
1615                                            bool) {
1616   // Always emit simple barriers!
1617   if (!CGF.HaveInsertPoint())
1618     return;
1619   // Build call __kmpc_cancel_barrier(loc, thread_id);
1620   unsigned Flags = getDefaultFlagsForBarriers(Kind);
1621   llvm::Value *Args[] = {emitUpdateLocation(CGF, Loc, Flags),
1622                          getThreadID(CGF, Loc)};
1623 
1624   CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
1625                           CGM.getModule(), OMPRTL___kmpc_barrier),
1626                       Args);
1627 }
1628 
1629 void CGOpenMPRuntimeGPU::emitCriticalRegion(
1630     CodeGenFunction &CGF, StringRef CriticalName,
1631     const RegionCodeGenTy &CriticalOpGen, SourceLocation Loc,
1632     const Expr *Hint) {
1633   llvm::BasicBlock *LoopBB = CGF.createBasicBlock("omp.critical.loop");
1634   llvm::BasicBlock *TestBB = CGF.createBasicBlock("omp.critical.test");
1635   llvm::BasicBlock *SyncBB = CGF.createBasicBlock("omp.critical.sync");
1636   llvm::BasicBlock *BodyBB = CGF.createBasicBlock("omp.critical.body");
1637   llvm::BasicBlock *ExitBB = CGF.createBasicBlock("omp.critical.exit");
1638 
1639   auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
1640 
1641   // Get the mask of active threads in the warp.
1642   llvm::Value *Mask = CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
1643       CGM.getModule(), OMPRTL___kmpc_warp_active_thread_mask));
1644   // Fetch team-local id of the thread.
1645   llvm::Value *ThreadID = RT.getGPUThreadID(CGF);
1646 
1647   // Get the width of the team.
1648   llvm::Value *TeamWidth = RT.getGPUNumThreads(CGF);
1649 
1650   // Initialize the counter variable for the loop.
1651   QualType Int32Ty =
1652       CGF.getContext().getIntTypeForBitwidth(/*DestWidth=*/32, /*Signed=*/0);
1653   Address Counter = CGF.CreateMemTemp(Int32Ty, "critical_counter");
1654   LValue CounterLVal = CGF.MakeAddrLValue(Counter, Int32Ty);
1655   CGF.EmitStoreOfScalar(llvm::Constant::getNullValue(CGM.Int32Ty), CounterLVal,
1656                         /*isInit=*/true);
1657 
1658   // Block checks if loop counter exceeds upper bound.
1659   CGF.EmitBlock(LoopBB);
1660   llvm::Value *CounterVal = CGF.EmitLoadOfScalar(CounterLVal, Loc);
1661   llvm::Value *CmpLoopBound = CGF.Builder.CreateICmpSLT(CounterVal, TeamWidth);
1662   CGF.Builder.CreateCondBr(CmpLoopBound, TestBB, ExitBB);
1663 
1664   // Block tests which single thread should execute region, and which threads
1665   // should go straight to synchronisation point.
1666   CGF.EmitBlock(TestBB);
1667   CounterVal = CGF.EmitLoadOfScalar(CounterLVal, Loc);
1668   llvm::Value *CmpThreadToCounter =
1669       CGF.Builder.CreateICmpEQ(ThreadID, CounterVal);
1670   CGF.Builder.CreateCondBr(CmpThreadToCounter, BodyBB, SyncBB);
1671 
1672   // Block emits the body of the critical region.
1673   CGF.EmitBlock(BodyBB);
1674 
1675   // Output the critical statement.
1676   CGOpenMPRuntime::emitCriticalRegion(CGF, CriticalName, CriticalOpGen, Loc,
1677                                       Hint);
1678 
1679   // After the body surrounded by the critical region, the single executing
1680   // thread will jump to the synchronisation point.
1681   // Block waits for all threads in current team to finish then increments the
1682   // counter variable and returns to the loop.
1683   CGF.EmitBlock(SyncBB);
1684   // Reconverge active threads in the warp.
1685   (void)CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
1686                                 CGM.getModule(), OMPRTL___kmpc_syncwarp),
1687                             Mask);
1688 
1689   llvm::Value *IncCounterVal =
1690       CGF.Builder.CreateNSWAdd(CounterVal, CGF.Builder.getInt32(1));
1691   CGF.EmitStoreOfScalar(IncCounterVal, CounterLVal);
1692   CGF.EmitBranch(LoopBB);
1693 
1694   // Block that is reached when  all threads in the team complete the region.
1695   CGF.EmitBlock(ExitBB, /*IsFinished=*/true);
1696 }
1697 
1698 /// Cast value to the specified type.
1699 static llvm::Value *castValueToType(CodeGenFunction &CGF, llvm::Value *Val,
1700                                     QualType ValTy, QualType CastTy,
1701                                     SourceLocation Loc) {
1702   assert(!CGF.getContext().getTypeSizeInChars(CastTy).isZero() &&
1703          "Cast type must sized.");
1704   assert(!CGF.getContext().getTypeSizeInChars(ValTy).isZero() &&
1705          "Val type must sized.");
1706   llvm::Type *LLVMCastTy = CGF.ConvertTypeForMem(CastTy);
1707   if (ValTy == CastTy)
1708     return Val;
1709   if (CGF.getContext().getTypeSizeInChars(ValTy) ==
1710       CGF.getContext().getTypeSizeInChars(CastTy))
1711     return CGF.Builder.CreateBitCast(Val, LLVMCastTy);
1712   if (CastTy->isIntegerType() && ValTy->isIntegerType())
1713     return CGF.Builder.CreateIntCast(Val, LLVMCastTy,
1714                                      CastTy->hasSignedIntegerRepresentation());
1715   Address CastItem = CGF.CreateMemTemp(CastTy);
1716   Address ValCastItem = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
1717       CastItem, Val->getType()->getPointerTo(CastItem.getAddressSpace()));
1718   CGF.EmitStoreOfScalar(Val, ValCastItem, /*Volatile=*/false, ValTy,
1719                         LValueBaseInfo(AlignmentSource::Type),
1720                         TBAAAccessInfo());
1721   return CGF.EmitLoadOfScalar(CastItem, /*Volatile=*/false, CastTy, Loc,
1722                               LValueBaseInfo(AlignmentSource::Type),
1723                               TBAAAccessInfo());
1724 }
1725 
1726 /// This function creates calls to one of two shuffle functions to copy
1727 /// variables between lanes in a warp.
1728 static llvm::Value *createRuntimeShuffleFunction(CodeGenFunction &CGF,
1729                                                  llvm::Value *Elem,
1730                                                  QualType ElemType,
1731                                                  llvm::Value *Offset,
1732                                                  SourceLocation Loc) {
1733   CodeGenModule &CGM = CGF.CGM;
1734   CGBuilderTy &Bld = CGF.Builder;
1735   CGOpenMPRuntimeGPU &RT =
1736       *(static_cast<CGOpenMPRuntimeGPU *>(&CGM.getOpenMPRuntime()));
1737   llvm::OpenMPIRBuilder &OMPBuilder = RT.getOMPBuilder();
1738 
1739   CharUnits Size = CGF.getContext().getTypeSizeInChars(ElemType);
1740   assert(Size.getQuantity() <= 8 &&
1741          "Unsupported bitwidth in shuffle instruction.");
1742 
1743   RuntimeFunction ShuffleFn = Size.getQuantity() <= 4
1744                                   ? OMPRTL___kmpc_shuffle_int32
1745                                   : OMPRTL___kmpc_shuffle_int64;
1746 
1747   // Cast all types to 32- or 64-bit values before calling shuffle routines.
1748   QualType CastTy = CGF.getContext().getIntTypeForBitwidth(
1749       Size.getQuantity() <= 4 ? 32 : 64, /*Signed=*/1);
1750   llvm::Value *ElemCast = castValueToType(CGF, Elem, ElemType, CastTy, Loc);
1751   llvm::Value *WarpSize =
1752       Bld.CreateIntCast(RT.getGPUWarpSize(CGF), CGM.Int16Ty, /*isSigned=*/true);
1753 
1754   llvm::Value *ShuffledVal = CGF.EmitRuntimeCall(
1755       OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(), ShuffleFn),
1756       {ElemCast, Offset, WarpSize});
1757 
1758   return castValueToType(CGF, ShuffledVal, CastTy, ElemType, Loc);
1759 }
1760 
1761 static void shuffleAndStore(CodeGenFunction &CGF, Address SrcAddr,
1762                             Address DestAddr, QualType ElemType,
1763                             llvm::Value *Offset, SourceLocation Loc) {
1764   CGBuilderTy &Bld = CGF.Builder;
1765 
1766   CharUnits Size = CGF.getContext().getTypeSizeInChars(ElemType);
1767   // Create the loop over the big sized data.
1768   // ptr = (void*)Elem;
1769   // ptrEnd = (void*) Elem + 1;
1770   // Step = 8;
1771   // while (ptr + Step < ptrEnd)
1772   //   shuffle((int64_t)*ptr);
1773   // Step = 4;
1774   // while (ptr + Step < ptrEnd)
1775   //   shuffle((int32_t)*ptr);
1776   // ...
1777   Address ElemPtr = DestAddr;
1778   Address Ptr = SrcAddr;
1779   Address PtrEnd = Bld.CreatePointerBitCastOrAddrSpaceCast(
1780       Bld.CreateConstGEP(SrcAddr, 1), CGF.VoidPtrTy);
1781   for (int IntSize = 8; IntSize >= 1; IntSize /= 2) {
1782     if (Size < CharUnits::fromQuantity(IntSize))
1783       continue;
1784     QualType IntType = CGF.getContext().getIntTypeForBitwidth(
1785         CGF.getContext().toBits(CharUnits::fromQuantity(IntSize)),
1786         /*Signed=*/1);
1787     llvm::Type *IntTy = CGF.ConvertTypeForMem(IntType);
1788     Ptr = Bld.CreatePointerBitCastOrAddrSpaceCast(Ptr, IntTy->getPointerTo());
1789     ElemPtr =
1790         Bld.CreatePointerBitCastOrAddrSpaceCast(ElemPtr, IntTy->getPointerTo());
1791     if (Size.getQuantity() / IntSize > 1) {
1792       llvm::BasicBlock *PreCondBB = CGF.createBasicBlock(".shuffle.pre_cond");
1793       llvm::BasicBlock *ThenBB = CGF.createBasicBlock(".shuffle.then");
1794       llvm::BasicBlock *ExitBB = CGF.createBasicBlock(".shuffle.exit");
1795       llvm::BasicBlock *CurrentBB = Bld.GetInsertBlock();
1796       CGF.EmitBlock(PreCondBB);
1797       llvm::PHINode *PhiSrc =
1798           Bld.CreatePHI(Ptr.getType(), /*NumReservedValues=*/2);
1799       PhiSrc->addIncoming(Ptr.getPointer(), CurrentBB);
1800       llvm::PHINode *PhiDest =
1801           Bld.CreatePHI(ElemPtr.getType(), /*NumReservedValues=*/2);
1802       PhiDest->addIncoming(ElemPtr.getPointer(), CurrentBB);
1803       Ptr = Address(PhiSrc, Ptr.getAlignment());
1804       ElemPtr = Address(PhiDest, ElemPtr.getAlignment());
1805       llvm::Value *PtrDiff = Bld.CreatePtrDiff(
1806           CGF.Int8Ty, PtrEnd.getPointer(),
1807           Bld.CreatePointerBitCastOrAddrSpaceCast(Ptr.getPointer(),
1808                                                   CGF.VoidPtrTy));
1809       Bld.CreateCondBr(Bld.CreateICmpSGT(PtrDiff, Bld.getInt64(IntSize - 1)),
1810                        ThenBB, ExitBB);
1811       CGF.EmitBlock(ThenBB);
1812       llvm::Value *Res = createRuntimeShuffleFunction(
1813           CGF,
1814           CGF.EmitLoadOfScalar(Ptr, /*Volatile=*/false, IntType, Loc,
1815                                LValueBaseInfo(AlignmentSource::Type),
1816                                TBAAAccessInfo()),
1817           IntType, Offset, Loc);
1818       CGF.EmitStoreOfScalar(Res, ElemPtr, /*Volatile=*/false, IntType,
1819                             LValueBaseInfo(AlignmentSource::Type),
1820                             TBAAAccessInfo());
1821       Address LocalPtr = Bld.CreateConstGEP(Ptr, 1);
1822       Address LocalElemPtr = Bld.CreateConstGEP(ElemPtr, 1);
1823       PhiSrc->addIncoming(LocalPtr.getPointer(), ThenBB);
1824       PhiDest->addIncoming(LocalElemPtr.getPointer(), ThenBB);
1825       CGF.EmitBranch(PreCondBB);
1826       CGF.EmitBlock(ExitBB);
1827     } else {
1828       llvm::Value *Res = createRuntimeShuffleFunction(
1829           CGF,
1830           CGF.EmitLoadOfScalar(Ptr, /*Volatile=*/false, IntType, Loc,
1831                                LValueBaseInfo(AlignmentSource::Type),
1832                                TBAAAccessInfo()),
1833           IntType, Offset, Loc);
1834       CGF.EmitStoreOfScalar(Res, ElemPtr, /*Volatile=*/false, IntType,
1835                             LValueBaseInfo(AlignmentSource::Type),
1836                             TBAAAccessInfo());
1837       Ptr = Bld.CreateConstGEP(Ptr, 1);
1838       ElemPtr = Bld.CreateConstGEP(ElemPtr, 1);
1839     }
1840     Size = Size % IntSize;
1841   }
1842 }
1843 
1844 namespace {
1845 enum CopyAction : unsigned {
1846   // RemoteLaneToThread: Copy over a Reduce list from a remote lane in
1847   // the warp using shuffle instructions.
1848   RemoteLaneToThread,
1849   // ThreadCopy: Make a copy of a Reduce list on the thread's stack.
1850   ThreadCopy,
1851   // ThreadToScratchpad: Copy a team-reduced array to the scratchpad.
1852   ThreadToScratchpad,
1853   // ScratchpadToThread: Copy from a scratchpad array in global memory
1854   // containing team-reduced data to a thread's stack.
1855   ScratchpadToThread,
1856 };
1857 } // namespace
1858 
1859 struct CopyOptionsTy {
1860   llvm::Value *RemoteLaneOffset;
1861   llvm::Value *ScratchpadIndex;
1862   llvm::Value *ScratchpadWidth;
1863 };
1864 
1865 /// Emit instructions to copy a Reduce list, which contains partially
1866 /// aggregated values, in the specified direction.
1867 static void emitReductionListCopy(
1868     CopyAction Action, CodeGenFunction &CGF, QualType ReductionArrayTy,
1869     ArrayRef<const Expr *> Privates, Address SrcBase, Address DestBase,
1870     CopyOptionsTy CopyOptions = {nullptr, nullptr, nullptr}) {
1871 
1872   CodeGenModule &CGM = CGF.CGM;
1873   ASTContext &C = CGM.getContext();
1874   CGBuilderTy &Bld = CGF.Builder;
1875 
1876   llvm::Value *RemoteLaneOffset = CopyOptions.RemoteLaneOffset;
1877   llvm::Value *ScratchpadIndex = CopyOptions.ScratchpadIndex;
1878   llvm::Value *ScratchpadWidth = CopyOptions.ScratchpadWidth;
1879 
1880   // Iterates, element-by-element, through the source Reduce list and
1881   // make a copy.
1882   unsigned Idx = 0;
1883   unsigned Size = Privates.size();
1884   for (const Expr *Private : Privates) {
1885     Address SrcElementAddr = Address::invalid();
1886     Address DestElementAddr = Address::invalid();
1887     Address DestElementPtrAddr = Address::invalid();
1888     // Should we shuffle in an element from a remote lane?
1889     bool ShuffleInElement = false;
1890     // Set to true to update the pointer in the dest Reduce list to a
1891     // newly created element.
1892     bool UpdateDestListPtr = false;
1893     // Increment the src or dest pointer to the scratchpad, for each
1894     // new element.
1895     bool IncrScratchpadSrc = false;
1896     bool IncrScratchpadDest = false;
1897 
1898     switch (Action) {
1899     case RemoteLaneToThread: {
1900       // Step 1.1: Get the address for the src element in the Reduce list.
1901       Address SrcElementPtrAddr = Bld.CreateConstArrayGEP(SrcBase, Idx);
1902       SrcElementAddr = CGF.EmitLoadOfPointer(
1903           SrcElementPtrAddr,
1904           C.getPointerType(Private->getType())->castAs<PointerType>());
1905 
1906       // Step 1.2: Create a temporary to store the element in the destination
1907       // Reduce list.
1908       DestElementPtrAddr = Bld.CreateConstArrayGEP(DestBase, Idx);
1909       DestElementAddr =
1910           CGF.CreateMemTemp(Private->getType(), ".omp.reduction.element");
1911       ShuffleInElement = true;
1912       UpdateDestListPtr = true;
1913       break;
1914     }
1915     case ThreadCopy: {
1916       // Step 1.1: Get the address for the src element in the Reduce list.
1917       Address SrcElementPtrAddr = Bld.CreateConstArrayGEP(SrcBase, Idx);
1918       SrcElementAddr = CGF.EmitLoadOfPointer(
1919           SrcElementPtrAddr,
1920           C.getPointerType(Private->getType())->castAs<PointerType>());
1921 
1922       // Step 1.2: Get the address for dest element.  The destination
1923       // element has already been created on the thread's stack.
1924       DestElementPtrAddr = Bld.CreateConstArrayGEP(DestBase, Idx);
1925       DestElementAddr = CGF.EmitLoadOfPointer(
1926           DestElementPtrAddr,
1927           C.getPointerType(Private->getType())->castAs<PointerType>());
1928       break;
1929     }
1930     case ThreadToScratchpad: {
1931       // Step 1.1: Get the address for the src element in the Reduce list.
1932       Address SrcElementPtrAddr = Bld.CreateConstArrayGEP(SrcBase, Idx);
1933       SrcElementAddr = CGF.EmitLoadOfPointer(
1934           SrcElementPtrAddr,
1935           C.getPointerType(Private->getType())->castAs<PointerType>());
1936 
1937       // Step 1.2: Get the address for dest element:
1938       // address = base + index * ElementSizeInChars.
1939       llvm::Value *ElementSizeInChars = CGF.getTypeSize(Private->getType());
1940       llvm::Value *CurrentOffset =
1941           Bld.CreateNUWMul(ElementSizeInChars, ScratchpadIndex);
1942       llvm::Value *ScratchPadElemAbsolutePtrVal =
1943           Bld.CreateNUWAdd(DestBase.getPointer(), CurrentOffset);
1944       ScratchPadElemAbsolutePtrVal =
1945           Bld.CreateIntToPtr(ScratchPadElemAbsolutePtrVal, CGF.VoidPtrTy);
1946       DestElementAddr = Address(ScratchPadElemAbsolutePtrVal,
1947                                 C.getTypeAlignInChars(Private->getType()));
1948       IncrScratchpadDest = true;
1949       break;
1950     }
1951     case ScratchpadToThread: {
1952       // Step 1.1: Get the address for the src element in the scratchpad.
1953       // address = base + index * ElementSizeInChars.
1954       llvm::Value *ElementSizeInChars = CGF.getTypeSize(Private->getType());
1955       llvm::Value *CurrentOffset =
1956           Bld.CreateNUWMul(ElementSizeInChars, ScratchpadIndex);
1957       llvm::Value *ScratchPadElemAbsolutePtrVal =
1958           Bld.CreateNUWAdd(SrcBase.getPointer(), CurrentOffset);
1959       ScratchPadElemAbsolutePtrVal =
1960           Bld.CreateIntToPtr(ScratchPadElemAbsolutePtrVal, CGF.VoidPtrTy);
1961       SrcElementAddr = Address(ScratchPadElemAbsolutePtrVal,
1962                                C.getTypeAlignInChars(Private->getType()));
1963       IncrScratchpadSrc = true;
1964 
1965       // Step 1.2: Create a temporary to store the element in the destination
1966       // Reduce list.
1967       DestElementPtrAddr = Bld.CreateConstArrayGEP(DestBase, Idx);
1968       DestElementAddr =
1969           CGF.CreateMemTemp(Private->getType(), ".omp.reduction.element");
1970       UpdateDestListPtr = true;
1971       break;
1972     }
1973     }
1974 
1975     // Regardless of src and dest of copy, we emit the load of src
1976     // element as this is required in all directions
1977     SrcElementAddr = Bld.CreateElementBitCast(
1978         SrcElementAddr, CGF.ConvertTypeForMem(Private->getType()));
1979     DestElementAddr = Bld.CreateElementBitCast(DestElementAddr,
1980                                                SrcElementAddr.getElementType());
1981 
1982     // Now that all active lanes have read the element in the
1983     // Reduce list, shuffle over the value from the remote lane.
1984     if (ShuffleInElement) {
1985       shuffleAndStore(CGF, SrcElementAddr, DestElementAddr, Private->getType(),
1986                       RemoteLaneOffset, Private->getExprLoc());
1987     } else {
1988       switch (CGF.getEvaluationKind(Private->getType())) {
1989       case TEK_Scalar: {
1990         llvm::Value *Elem = CGF.EmitLoadOfScalar(
1991             SrcElementAddr, /*Volatile=*/false, Private->getType(),
1992             Private->getExprLoc(), LValueBaseInfo(AlignmentSource::Type),
1993             TBAAAccessInfo());
1994         // Store the source element value to the dest element address.
1995         CGF.EmitStoreOfScalar(
1996             Elem, DestElementAddr, /*Volatile=*/false, Private->getType(),
1997             LValueBaseInfo(AlignmentSource::Type), TBAAAccessInfo());
1998         break;
1999       }
2000       case TEK_Complex: {
2001         CodeGenFunction::ComplexPairTy Elem = CGF.EmitLoadOfComplex(
2002             CGF.MakeAddrLValue(SrcElementAddr, Private->getType()),
2003             Private->getExprLoc());
2004         CGF.EmitStoreOfComplex(
2005             Elem, CGF.MakeAddrLValue(DestElementAddr, Private->getType()),
2006             /*isInit=*/false);
2007         break;
2008       }
2009       case TEK_Aggregate:
2010         CGF.EmitAggregateCopy(
2011             CGF.MakeAddrLValue(DestElementAddr, Private->getType()),
2012             CGF.MakeAddrLValue(SrcElementAddr, Private->getType()),
2013             Private->getType(), AggValueSlot::DoesNotOverlap);
2014         break;
2015       }
2016     }
2017 
2018     // Step 3.1: Modify reference in dest Reduce list as needed.
2019     // Modifying the reference in Reduce list to point to the newly
2020     // created element.  The element is live in the current function
2021     // scope and that of functions it invokes (i.e., reduce_function).
2022     // RemoteReduceData[i] = (void*)&RemoteElem
2023     if (UpdateDestListPtr) {
2024       CGF.EmitStoreOfScalar(Bld.CreatePointerBitCastOrAddrSpaceCast(
2025                                 DestElementAddr.getPointer(), CGF.VoidPtrTy),
2026                             DestElementPtrAddr, /*Volatile=*/false,
2027                             C.VoidPtrTy);
2028     }
2029 
2030     // Step 4.1: Increment SrcBase/DestBase so that it points to the starting
2031     // address of the next element in scratchpad memory, unless we're currently
2032     // processing the last one.  Memory alignment is also taken care of here.
2033     if ((IncrScratchpadDest || IncrScratchpadSrc) && (Idx + 1 < Size)) {
2034       llvm::Value *ScratchpadBasePtr =
2035           IncrScratchpadDest ? DestBase.getPointer() : SrcBase.getPointer();
2036       llvm::Value *ElementSizeInChars = CGF.getTypeSize(Private->getType());
2037       ScratchpadBasePtr = Bld.CreateNUWAdd(
2038           ScratchpadBasePtr,
2039           Bld.CreateNUWMul(ScratchpadWidth, ElementSizeInChars));
2040 
2041       // Take care of global memory alignment for performance
2042       ScratchpadBasePtr = Bld.CreateNUWSub(
2043           ScratchpadBasePtr, llvm::ConstantInt::get(CGM.SizeTy, 1));
2044       ScratchpadBasePtr = Bld.CreateUDiv(
2045           ScratchpadBasePtr,
2046           llvm::ConstantInt::get(CGM.SizeTy, GlobalMemoryAlignment));
2047       ScratchpadBasePtr = Bld.CreateNUWAdd(
2048           ScratchpadBasePtr, llvm::ConstantInt::get(CGM.SizeTy, 1));
2049       ScratchpadBasePtr = Bld.CreateNUWMul(
2050           ScratchpadBasePtr,
2051           llvm::ConstantInt::get(CGM.SizeTy, GlobalMemoryAlignment));
2052 
2053       if (IncrScratchpadDest)
2054         DestBase = Address(ScratchpadBasePtr, CGF.getPointerAlign());
2055       else /* IncrScratchpadSrc = true */
2056         SrcBase = Address(ScratchpadBasePtr, CGF.getPointerAlign());
2057     }
2058 
2059     ++Idx;
2060   }
2061 }
2062 
2063 /// This function emits a helper that gathers Reduce lists from the first
2064 /// lane of every active warp to lanes in the first warp.
2065 ///
2066 /// void inter_warp_copy_func(void* reduce_data, num_warps)
2067 ///   shared smem[warp_size];
2068 ///   For all data entries D in reduce_data:
2069 ///     sync
2070 ///     If (I am the first lane in each warp)
2071 ///       Copy my local D to smem[warp_id]
2072 ///     sync
2073 ///     if (I am the first warp)
2074 ///       Copy smem[thread_id] to my local D
2075 static llvm::Value *emitInterWarpCopyFunction(CodeGenModule &CGM,
2076                                               ArrayRef<const Expr *> Privates,
2077                                               QualType ReductionArrayTy,
2078                                               SourceLocation Loc) {
2079   ASTContext &C = CGM.getContext();
2080   llvm::Module &M = CGM.getModule();
2081 
2082   // ReduceList: thread local Reduce list.
2083   // At the stage of the computation when this function is called, partially
2084   // aggregated values reside in the first lane of every active warp.
2085   ImplicitParamDecl ReduceListArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
2086                                   C.VoidPtrTy, ImplicitParamDecl::Other);
2087   // NumWarps: number of warps active in the parallel region.  This could
2088   // be smaller than 32 (max warps in a CTA) for partial block reduction.
2089   ImplicitParamDecl NumWarpsArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
2090                                 C.getIntTypeForBitwidth(32, /* Signed */ true),
2091                                 ImplicitParamDecl::Other);
2092   FunctionArgList Args;
2093   Args.push_back(&ReduceListArg);
2094   Args.push_back(&NumWarpsArg);
2095 
2096   const CGFunctionInfo &CGFI =
2097       CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args);
2098   auto *Fn = llvm::Function::Create(CGM.getTypes().GetFunctionType(CGFI),
2099                                     llvm::GlobalValue::InternalLinkage,
2100                                     "_omp_reduction_inter_warp_copy_func", &M);
2101   CGM.SetInternalFunctionAttributes(GlobalDecl(), Fn, CGFI);
2102   Fn->setDoesNotRecurse();
2103   CodeGenFunction CGF(CGM);
2104   CGF.StartFunction(GlobalDecl(), C.VoidTy, Fn, CGFI, Args, Loc, Loc);
2105 
2106   CGBuilderTy &Bld = CGF.Builder;
2107 
2108   // This array is used as a medium to transfer, one reduce element at a time,
2109   // the data from the first lane of every warp to lanes in the first warp
2110   // in order to perform the final step of a reduction in a parallel region
2111   // (reduction across warps).  The array is placed in NVPTX __shared__ memory
2112   // for reduced latency, as well as to have a distinct copy for concurrently
2113   // executing target regions.  The array is declared with common linkage so
2114   // as to be shared across compilation units.
2115   StringRef TransferMediumName =
2116       "__openmp_nvptx_data_transfer_temporary_storage";
2117   llvm::GlobalVariable *TransferMedium =
2118       M.getGlobalVariable(TransferMediumName);
2119   unsigned WarpSize = CGF.getTarget().getGridValue().GV_Warp_Size;
2120   if (!TransferMedium) {
2121     auto *Ty = llvm::ArrayType::get(CGM.Int32Ty, WarpSize);
2122     unsigned SharedAddressSpace = C.getTargetAddressSpace(LangAS::cuda_shared);
2123     TransferMedium = new llvm::GlobalVariable(
2124         M, Ty, /*isConstant=*/false, llvm::GlobalVariable::WeakAnyLinkage,
2125         llvm::UndefValue::get(Ty), TransferMediumName,
2126         /*InsertBefore=*/nullptr, llvm::GlobalVariable::NotThreadLocal,
2127         SharedAddressSpace);
2128     CGM.addCompilerUsedGlobal(TransferMedium);
2129   }
2130 
2131   auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
2132   // Get the CUDA thread id of the current OpenMP thread on the GPU.
2133   llvm::Value *ThreadID = RT.getGPUThreadID(CGF);
2134   // nvptx_lane_id = nvptx_id % warpsize
2135   llvm::Value *LaneID = getNVPTXLaneID(CGF);
2136   // nvptx_warp_id = nvptx_id / warpsize
2137   llvm::Value *WarpID = getNVPTXWarpID(CGF);
2138 
2139   Address AddrReduceListArg = CGF.GetAddrOfLocalVar(&ReduceListArg);
2140   Address LocalReduceList(
2141       Bld.CreatePointerBitCastOrAddrSpaceCast(
2142           CGF.EmitLoadOfScalar(
2143               AddrReduceListArg, /*Volatile=*/false, C.VoidPtrTy, Loc,
2144               LValueBaseInfo(AlignmentSource::Type), TBAAAccessInfo()),
2145           CGF.ConvertTypeForMem(ReductionArrayTy)->getPointerTo()),
2146       CGF.getPointerAlign());
2147 
2148   unsigned Idx = 0;
2149   for (const Expr *Private : Privates) {
2150     //
2151     // Warp master copies reduce element to transfer medium in __shared__
2152     // memory.
2153     //
2154     unsigned RealTySize =
2155         C.getTypeSizeInChars(Private->getType())
2156             .alignTo(C.getTypeAlignInChars(Private->getType()))
2157             .getQuantity();
2158     for (unsigned TySize = 4; TySize > 0 && RealTySize > 0; TySize /=2) {
2159       unsigned NumIters = RealTySize / TySize;
2160       if (NumIters == 0)
2161         continue;
2162       QualType CType = C.getIntTypeForBitwidth(
2163           C.toBits(CharUnits::fromQuantity(TySize)), /*Signed=*/1);
2164       llvm::Type *CopyType = CGF.ConvertTypeForMem(CType);
2165       CharUnits Align = CharUnits::fromQuantity(TySize);
2166       llvm::Value *Cnt = nullptr;
2167       Address CntAddr = Address::invalid();
2168       llvm::BasicBlock *PrecondBB = nullptr;
2169       llvm::BasicBlock *ExitBB = nullptr;
2170       if (NumIters > 1) {
2171         CntAddr = CGF.CreateMemTemp(C.IntTy, ".cnt.addr");
2172         CGF.EmitStoreOfScalar(llvm::Constant::getNullValue(CGM.IntTy), CntAddr,
2173                               /*Volatile=*/false, C.IntTy);
2174         PrecondBB = CGF.createBasicBlock("precond");
2175         ExitBB = CGF.createBasicBlock("exit");
2176         llvm::BasicBlock *BodyBB = CGF.createBasicBlock("body");
2177         // There is no need to emit line number for unconditional branch.
2178         (void)ApplyDebugLocation::CreateEmpty(CGF);
2179         CGF.EmitBlock(PrecondBB);
2180         Cnt = CGF.EmitLoadOfScalar(CntAddr, /*Volatile=*/false, C.IntTy, Loc);
2181         llvm::Value *Cmp =
2182             Bld.CreateICmpULT(Cnt, llvm::ConstantInt::get(CGM.IntTy, NumIters));
2183         Bld.CreateCondBr(Cmp, BodyBB, ExitBB);
2184         CGF.EmitBlock(BodyBB);
2185       }
2186       // kmpc_barrier.
2187       CGM.getOpenMPRuntime().emitBarrierCall(CGF, Loc, OMPD_unknown,
2188                                              /*EmitChecks=*/false,
2189                                              /*ForceSimpleCall=*/true);
2190       llvm::BasicBlock *ThenBB = CGF.createBasicBlock("then");
2191       llvm::BasicBlock *ElseBB = CGF.createBasicBlock("else");
2192       llvm::BasicBlock *MergeBB = CGF.createBasicBlock("ifcont");
2193 
2194       // if (lane_id == 0)
2195       llvm::Value *IsWarpMaster = Bld.CreateIsNull(LaneID, "warp_master");
2196       Bld.CreateCondBr(IsWarpMaster, ThenBB, ElseBB);
2197       CGF.EmitBlock(ThenBB);
2198 
2199       // Reduce element = LocalReduceList[i]
2200       Address ElemPtrPtrAddr = Bld.CreateConstArrayGEP(LocalReduceList, Idx);
2201       llvm::Value *ElemPtrPtr = CGF.EmitLoadOfScalar(
2202           ElemPtrPtrAddr, /*Volatile=*/false, C.VoidPtrTy, SourceLocation());
2203       // elemptr = ((CopyType*)(elemptrptr)) + I
2204       Address ElemPtr = Address(ElemPtrPtr, Align);
2205       ElemPtr = Bld.CreateElementBitCast(ElemPtr, CopyType);
2206       if (NumIters > 1)
2207         ElemPtr = Bld.CreateGEP(ElemPtr, Cnt);
2208 
2209       // Get pointer to location in transfer medium.
2210       // MediumPtr = &medium[warp_id]
2211       llvm::Value *MediumPtrVal = Bld.CreateInBoundsGEP(
2212           TransferMedium->getValueType(), TransferMedium,
2213           {llvm::Constant::getNullValue(CGM.Int64Ty), WarpID});
2214       Address MediumPtr(MediumPtrVal, Align);
2215       // Casting to actual data type.
2216       // MediumPtr = (CopyType*)MediumPtrAddr;
2217       MediumPtr = Bld.CreateElementBitCast(MediumPtr, CopyType);
2218 
2219       // elem = *elemptr
2220       //*MediumPtr = elem
2221       llvm::Value *Elem = CGF.EmitLoadOfScalar(
2222           ElemPtr, /*Volatile=*/false, CType, Loc,
2223           LValueBaseInfo(AlignmentSource::Type), TBAAAccessInfo());
2224       // Store the source element value to the dest element address.
2225       CGF.EmitStoreOfScalar(Elem, MediumPtr, /*Volatile=*/true, CType,
2226                             LValueBaseInfo(AlignmentSource::Type),
2227                             TBAAAccessInfo());
2228 
2229       Bld.CreateBr(MergeBB);
2230 
2231       CGF.EmitBlock(ElseBB);
2232       Bld.CreateBr(MergeBB);
2233 
2234       CGF.EmitBlock(MergeBB);
2235 
2236       // kmpc_barrier.
2237       CGM.getOpenMPRuntime().emitBarrierCall(CGF, Loc, OMPD_unknown,
2238                                              /*EmitChecks=*/false,
2239                                              /*ForceSimpleCall=*/true);
2240 
2241       //
2242       // Warp 0 copies reduce element from transfer medium.
2243       //
2244       llvm::BasicBlock *W0ThenBB = CGF.createBasicBlock("then");
2245       llvm::BasicBlock *W0ElseBB = CGF.createBasicBlock("else");
2246       llvm::BasicBlock *W0MergeBB = CGF.createBasicBlock("ifcont");
2247 
2248       Address AddrNumWarpsArg = CGF.GetAddrOfLocalVar(&NumWarpsArg);
2249       llvm::Value *NumWarpsVal = CGF.EmitLoadOfScalar(
2250           AddrNumWarpsArg, /*Volatile=*/false, C.IntTy, Loc);
2251 
2252       // Up to 32 threads in warp 0 are active.
2253       llvm::Value *IsActiveThread =
2254           Bld.CreateICmpULT(ThreadID, NumWarpsVal, "is_active_thread");
2255       Bld.CreateCondBr(IsActiveThread, W0ThenBB, W0ElseBB);
2256 
2257       CGF.EmitBlock(W0ThenBB);
2258 
2259       // SrcMediumPtr = &medium[tid]
2260       llvm::Value *SrcMediumPtrVal = Bld.CreateInBoundsGEP(
2261           TransferMedium->getValueType(), TransferMedium,
2262           {llvm::Constant::getNullValue(CGM.Int64Ty), ThreadID});
2263       Address SrcMediumPtr(SrcMediumPtrVal, Align);
2264       // SrcMediumVal = *SrcMediumPtr;
2265       SrcMediumPtr = Bld.CreateElementBitCast(SrcMediumPtr, CopyType);
2266 
2267       // TargetElemPtr = (CopyType*)(SrcDataAddr[i]) + I
2268       Address TargetElemPtrPtr = Bld.CreateConstArrayGEP(LocalReduceList, Idx);
2269       llvm::Value *TargetElemPtrVal = CGF.EmitLoadOfScalar(
2270           TargetElemPtrPtr, /*Volatile=*/false, C.VoidPtrTy, Loc);
2271       Address TargetElemPtr = Address(TargetElemPtrVal, Align);
2272       TargetElemPtr = Bld.CreateElementBitCast(TargetElemPtr, CopyType);
2273       if (NumIters > 1)
2274         TargetElemPtr = Bld.CreateGEP(TargetElemPtr, Cnt);
2275 
2276       // *TargetElemPtr = SrcMediumVal;
2277       llvm::Value *SrcMediumValue =
2278           CGF.EmitLoadOfScalar(SrcMediumPtr, /*Volatile=*/true, CType, Loc);
2279       CGF.EmitStoreOfScalar(SrcMediumValue, TargetElemPtr, /*Volatile=*/false,
2280                             CType);
2281       Bld.CreateBr(W0MergeBB);
2282 
2283       CGF.EmitBlock(W0ElseBB);
2284       Bld.CreateBr(W0MergeBB);
2285 
2286       CGF.EmitBlock(W0MergeBB);
2287 
2288       if (NumIters > 1) {
2289         Cnt = Bld.CreateNSWAdd(Cnt, llvm::ConstantInt::get(CGM.IntTy, /*V=*/1));
2290         CGF.EmitStoreOfScalar(Cnt, CntAddr, /*Volatile=*/false, C.IntTy);
2291         CGF.EmitBranch(PrecondBB);
2292         (void)ApplyDebugLocation::CreateEmpty(CGF);
2293         CGF.EmitBlock(ExitBB);
2294       }
2295       RealTySize %= TySize;
2296     }
2297     ++Idx;
2298   }
2299 
2300   CGF.FinishFunction();
2301   return Fn;
2302 }
2303 
2304 /// Emit a helper that reduces data across two OpenMP threads (lanes)
2305 /// in the same warp.  It uses shuffle instructions to copy over data from
2306 /// a remote lane's stack.  The reduction algorithm performed is specified
2307 /// by the fourth parameter.
2308 ///
2309 /// Algorithm Versions.
2310 /// Full Warp Reduce (argument value 0):
2311 ///   This algorithm assumes that all 32 lanes are active and gathers
2312 ///   data from these 32 lanes, producing a single resultant value.
2313 /// Contiguous Partial Warp Reduce (argument value 1):
2314 ///   This algorithm assumes that only a *contiguous* subset of lanes
2315 ///   are active.  This happens for the last warp in a parallel region
2316 ///   when the user specified num_threads is not an integer multiple of
2317 ///   32.  This contiguous subset always starts with the zeroth lane.
2318 /// Partial Warp Reduce (argument value 2):
2319 ///   This algorithm gathers data from any number of lanes at any position.
2320 /// All reduced values are stored in the lowest possible lane.  The set
2321 /// of problems every algorithm addresses is a super set of those
2322 /// addressable by algorithms with a lower version number.  Overhead
2323 /// increases as algorithm version increases.
2324 ///
2325 /// Terminology
2326 /// Reduce element:
2327 ///   Reduce element refers to the individual data field with primitive
2328 ///   data types to be combined and reduced across threads.
2329 /// Reduce list:
2330 ///   Reduce list refers to a collection of local, thread-private
2331 ///   reduce elements.
2332 /// Remote Reduce list:
2333 ///   Remote Reduce list refers to a collection of remote (relative to
2334 ///   the current thread) reduce elements.
2335 ///
2336 /// We distinguish between three states of threads that are important to
2337 /// the implementation of this function.
2338 /// Alive threads:
2339 ///   Threads in a warp executing the SIMT instruction, as distinguished from
2340 ///   threads that are inactive due to divergent control flow.
2341 /// Active threads:
2342 ///   The minimal set of threads that has to be alive upon entry to this
2343 ///   function.  The computation is correct iff active threads are alive.
2344 ///   Some threads are alive but they are not active because they do not
2345 ///   contribute to the computation in any useful manner.  Turning them off
2346 ///   may introduce control flow overheads without any tangible benefits.
2347 /// Effective threads:
2348 ///   In order to comply with the argument requirements of the shuffle
2349 ///   function, we must keep all lanes holding data alive.  But at most
2350 ///   half of them perform value aggregation; we refer to this half of
2351 ///   threads as effective. The other half is simply handing off their
2352 ///   data.
2353 ///
2354 /// Procedure
2355 /// Value shuffle:
2356 ///   In this step active threads transfer data from higher lane positions
2357 ///   in the warp to lower lane positions, creating Remote Reduce list.
2358 /// Value aggregation:
2359 ///   In this step, effective threads combine their thread local Reduce list
2360 ///   with Remote Reduce list and store the result in the thread local
2361 ///   Reduce list.
2362 /// Value copy:
2363 ///   In this step, we deal with the assumption made by algorithm 2
2364 ///   (i.e. contiguity assumption).  When we have an odd number of lanes
2365 ///   active, say 2k+1, only k threads will be effective and therefore k
2366 ///   new values will be produced.  However, the Reduce list owned by the
2367 ///   (2k+1)th thread is ignored in the value aggregation.  Therefore
2368 ///   we copy the Reduce list from the (2k+1)th lane to (k+1)th lane so
2369 ///   that the contiguity assumption still holds.
2370 static llvm::Function *emitShuffleAndReduceFunction(
2371     CodeGenModule &CGM, ArrayRef<const Expr *> Privates,
2372     QualType ReductionArrayTy, llvm::Function *ReduceFn, SourceLocation Loc) {
2373   ASTContext &C = CGM.getContext();
2374 
2375   // Thread local Reduce list used to host the values of data to be reduced.
2376   ImplicitParamDecl ReduceListArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
2377                                   C.VoidPtrTy, ImplicitParamDecl::Other);
2378   // Current lane id; could be logical.
2379   ImplicitParamDecl LaneIDArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr, C.ShortTy,
2380                               ImplicitParamDecl::Other);
2381   // Offset of the remote source lane relative to the current lane.
2382   ImplicitParamDecl RemoteLaneOffsetArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
2383                                         C.ShortTy, ImplicitParamDecl::Other);
2384   // Algorithm version.  This is expected to be known at compile time.
2385   ImplicitParamDecl AlgoVerArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
2386                                C.ShortTy, ImplicitParamDecl::Other);
2387   FunctionArgList Args;
2388   Args.push_back(&ReduceListArg);
2389   Args.push_back(&LaneIDArg);
2390   Args.push_back(&RemoteLaneOffsetArg);
2391   Args.push_back(&AlgoVerArg);
2392 
2393   const CGFunctionInfo &CGFI =
2394       CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args);
2395   auto *Fn = llvm::Function::Create(
2396       CGM.getTypes().GetFunctionType(CGFI), llvm::GlobalValue::InternalLinkage,
2397       "_omp_reduction_shuffle_and_reduce_func", &CGM.getModule());
2398   CGM.SetInternalFunctionAttributes(GlobalDecl(), Fn, CGFI);
2399   Fn->setDoesNotRecurse();
2400 
2401   CodeGenFunction CGF(CGM);
2402   CGF.StartFunction(GlobalDecl(), C.VoidTy, Fn, CGFI, Args, Loc, Loc);
2403 
2404   CGBuilderTy &Bld = CGF.Builder;
2405 
2406   Address AddrReduceListArg = CGF.GetAddrOfLocalVar(&ReduceListArg);
2407   Address LocalReduceList(
2408       Bld.CreatePointerBitCastOrAddrSpaceCast(
2409           CGF.EmitLoadOfScalar(AddrReduceListArg, /*Volatile=*/false,
2410                                C.VoidPtrTy, SourceLocation()),
2411           CGF.ConvertTypeForMem(ReductionArrayTy)->getPointerTo()),
2412       CGF.getPointerAlign());
2413 
2414   Address AddrLaneIDArg = CGF.GetAddrOfLocalVar(&LaneIDArg);
2415   llvm::Value *LaneIDArgVal = CGF.EmitLoadOfScalar(
2416       AddrLaneIDArg, /*Volatile=*/false, C.ShortTy, SourceLocation());
2417 
2418   Address AddrRemoteLaneOffsetArg = CGF.GetAddrOfLocalVar(&RemoteLaneOffsetArg);
2419   llvm::Value *RemoteLaneOffsetArgVal = CGF.EmitLoadOfScalar(
2420       AddrRemoteLaneOffsetArg, /*Volatile=*/false, C.ShortTy, SourceLocation());
2421 
2422   Address AddrAlgoVerArg = CGF.GetAddrOfLocalVar(&AlgoVerArg);
2423   llvm::Value *AlgoVerArgVal = CGF.EmitLoadOfScalar(
2424       AddrAlgoVerArg, /*Volatile=*/false, C.ShortTy, SourceLocation());
2425 
2426   // Create a local thread-private variable to host the Reduce list
2427   // from a remote lane.
2428   Address RemoteReduceList =
2429       CGF.CreateMemTemp(ReductionArrayTy, ".omp.reduction.remote_reduce_list");
2430 
2431   // This loop iterates through the list of reduce elements and copies,
2432   // element by element, from a remote lane in the warp to RemoteReduceList,
2433   // hosted on the thread's stack.
2434   emitReductionListCopy(RemoteLaneToThread, CGF, ReductionArrayTy, Privates,
2435                         LocalReduceList, RemoteReduceList,
2436                         {/*RemoteLaneOffset=*/RemoteLaneOffsetArgVal,
2437                          /*ScratchpadIndex=*/nullptr,
2438                          /*ScratchpadWidth=*/nullptr});
2439 
2440   // The actions to be performed on the Remote Reduce list is dependent
2441   // on the algorithm version.
2442   //
2443   //  if (AlgoVer==0) || (AlgoVer==1 && (LaneId < Offset)) || (AlgoVer==2 &&
2444   //  LaneId % 2 == 0 && Offset > 0):
2445   //    do the reduction value aggregation
2446   //
2447   //  The thread local variable Reduce list is mutated in place to host the
2448   //  reduced data, which is the aggregated value produced from local and
2449   //  remote lanes.
2450   //
2451   //  Note that AlgoVer is expected to be a constant integer known at compile
2452   //  time.
2453   //  When AlgoVer==0, the first conjunction evaluates to true, making
2454   //    the entire predicate true during compile time.
2455   //  When AlgoVer==1, the second conjunction has only the second part to be
2456   //    evaluated during runtime.  Other conjunctions evaluates to false
2457   //    during compile time.
2458   //  When AlgoVer==2, the third conjunction has only the second part to be
2459   //    evaluated during runtime.  Other conjunctions evaluates to false
2460   //    during compile time.
2461   llvm::Value *CondAlgo0 = Bld.CreateIsNull(AlgoVerArgVal);
2462 
2463   llvm::Value *Algo1 = Bld.CreateICmpEQ(AlgoVerArgVal, Bld.getInt16(1));
2464   llvm::Value *CondAlgo1 = Bld.CreateAnd(
2465       Algo1, Bld.CreateICmpULT(LaneIDArgVal, RemoteLaneOffsetArgVal));
2466 
2467   llvm::Value *Algo2 = Bld.CreateICmpEQ(AlgoVerArgVal, Bld.getInt16(2));
2468   llvm::Value *CondAlgo2 = Bld.CreateAnd(
2469       Algo2, Bld.CreateIsNull(Bld.CreateAnd(LaneIDArgVal, Bld.getInt16(1))));
2470   CondAlgo2 = Bld.CreateAnd(
2471       CondAlgo2, Bld.CreateICmpSGT(RemoteLaneOffsetArgVal, Bld.getInt16(0)));
2472 
2473   llvm::Value *CondReduce = Bld.CreateOr(CondAlgo0, CondAlgo1);
2474   CondReduce = Bld.CreateOr(CondReduce, CondAlgo2);
2475 
2476   llvm::BasicBlock *ThenBB = CGF.createBasicBlock("then");
2477   llvm::BasicBlock *ElseBB = CGF.createBasicBlock("else");
2478   llvm::BasicBlock *MergeBB = CGF.createBasicBlock("ifcont");
2479   Bld.CreateCondBr(CondReduce, ThenBB, ElseBB);
2480 
2481   CGF.EmitBlock(ThenBB);
2482   // reduce_function(LocalReduceList, RemoteReduceList)
2483   llvm::Value *LocalReduceListPtr = Bld.CreatePointerBitCastOrAddrSpaceCast(
2484       LocalReduceList.getPointer(), CGF.VoidPtrTy);
2485   llvm::Value *RemoteReduceListPtr = Bld.CreatePointerBitCastOrAddrSpaceCast(
2486       RemoteReduceList.getPointer(), CGF.VoidPtrTy);
2487   CGM.getOpenMPRuntime().emitOutlinedFunctionCall(
2488       CGF, Loc, ReduceFn, {LocalReduceListPtr, RemoteReduceListPtr});
2489   Bld.CreateBr(MergeBB);
2490 
2491   CGF.EmitBlock(ElseBB);
2492   Bld.CreateBr(MergeBB);
2493 
2494   CGF.EmitBlock(MergeBB);
2495 
2496   // if (AlgoVer==1 && (LaneId >= Offset)) copy Remote Reduce list to local
2497   // Reduce list.
2498   Algo1 = Bld.CreateICmpEQ(AlgoVerArgVal, Bld.getInt16(1));
2499   llvm::Value *CondCopy = Bld.CreateAnd(
2500       Algo1, Bld.CreateICmpUGE(LaneIDArgVal, RemoteLaneOffsetArgVal));
2501 
2502   llvm::BasicBlock *CpyThenBB = CGF.createBasicBlock("then");
2503   llvm::BasicBlock *CpyElseBB = CGF.createBasicBlock("else");
2504   llvm::BasicBlock *CpyMergeBB = CGF.createBasicBlock("ifcont");
2505   Bld.CreateCondBr(CondCopy, CpyThenBB, CpyElseBB);
2506 
2507   CGF.EmitBlock(CpyThenBB);
2508   emitReductionListCopy(ThreadCopy, CGF, ReductionArrayTy, Privates,
2509                         RemoteReduceList, LocalReduceList);
2510   Bld.CreateBr(CpyMergeBB);
2511 
2512   CGF.EmitBlock(CpyElseBB);
2513   Bld.CreateBr(CpyMergeBB);
2514 
2515   CGF.EmitBlock(CpyMergeBB);
2516 
2517   CGF.FinishFunction();
2518   return Fn;
2519 }
2520 
2521 /// This function emits a helper that copies all the reduction variables from
2522 /// the team into the provided global buffer for the reduction variables.
2523 ///
2524 /// void list_to_global_copy_func(void *buffer, int Idx, void *reduce_data)
2525 ///   For all data entries D in reduce_data:
2526 ///     Copy local D to buffer.D[Idx]
2527 static llvm::Value *emitListToGlobalCopyFunction(
2528     CodeGenModule &CGM, ArrayRef<const Expr *> Privates,
2529     QualType ReductionArrayTy, SourceLocation Loc,
2530     const RecordDecl *TeamReductionRec,
2531     const llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>
2532         &VarFieldMap) {
2533   ASTContext &C = CGM.getContext();
2534 
2535   // Buffer: global reduction buffer.
2536   ImplicitParamDecl BufferArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
2537                               C.VoidPtrTy, ImplicitParamDecl::Other);
2538   // Idx: index of the buffer.
2539   ImplicitParamDecl IdxArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr, C.IntTy,
2540                            ImplicitParamDecl::Other);
2541   // ReduceList: thread local Reduce list.
2542   ImplicitParamDecl ReduceListArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
2543                                   C.VoidPtrTy, ImplicitParamDecl::Other);
2544   FunctionArgList Args;
2545   Args.push_back(&BufferArg);
2546   Args.push_back(&IdxArg);
2547   Args.push_back(&ReduceListArg);
2548 
2549   const CGFunctionInfo &CGFI =
2550       CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args);
2551   auto *Fn = llvm::Function::Create(
2552       CGM.getTypes().GetFunctionType(CGFI), llvm::GlobalValue::InternalLinkage,
2553       "_omp_reduction_list_to_global_copy_func", &CGM.getModule());
2554   CGM.SetInternalFunctionAttributes(GlobalDecl(), Fn, CGFI);
2555   Fn->setDoesNotRecurse();
2556   CodeGenFunction CGF(CGM);
2557   CGF.StartFunction(GlobalDecl(), C.VoidTy, Fn, CGFI, Args, Loc, Loc);
2558 
2559   CGBuilderTy &Bld = CGF.Builder;
2560 
2561   Address AddrReduceListArg = CGF.GetAddrOfLocalVar(&ReduceListArg);
2562   Address AddrBufferArg = CGF.GetAddrOfLocalVar(&BufferArg);
2563   Address LocalReduceList(
2564       Bld.CreatePointerBitCastOrAddrSpaceCast(
2565           CGF.EmitLoadOfScalar(AddrReduceListArg, /*Volatile=*/false,
2566                                C.VoidPtrTy, Loc),
2567           CGF.ConvertTypeForMem(ReductionArrayTy)->getPointerTo()),
2568       CGF.getPointerAlign());
2569   QualType StaticTy = C.getRecordType(TeamReductionRec);
2570   llvm::Type *LLVMReductionsBufferTy =
2571       CGM.getTypes().ConvertTypeForMem(StaticTy);
2572   llvm::Value *BufferArrPtr = Bld.CreatePointerBitCastOrAddrSpaceCast(
2573       CGF.EmitLoadOfScalar(AddrBufferArg, /*Volatile=*/false, C.VoidPtrTy, Loc),
2574       LLVMReductionsBufferTy->getPointerTo());
2575   llvm::Value *Idxs[] = {llvm::ConstantInt::getNullValue(CGF.Int32Ty),
2576                          CGF.EmitLoadOfScalar(CGF.GetAddrOfLocalVar(&IdxArg),
2577                                               /*Volatile=*/false, C.IntTy,
2578                                               Loc)};
2579   unsigned Idx = 0;
2580   for (const Expr *Private : Privates) {
2581     // Reduce element = LocalReduceList[i]
2582     Address ElemPtrPtrAddr = Bld.CreateConstArrayGEP(LocalReduceList, Idx);
2583     llvm::Value *ElemPtrPtr = CGF.EmitLoadOfScalar(
2584         ElemPtrPtrAddr, /*Volatile=*/false, C.VoidPtrTy, SourceLocation());
2585     // elemptr = ((CopyType*)(elemptrptr)) + I
2586     ElemPtrPtr = Bld.CreatePointerBitCastOrAddrSpaceCast(
2587         ElemPtrPtr, CGF.ConvertTypeForMem(Private->getType())->getPointerTo());
2588     Address ElemPtr =
2589         Address(ElemPtrPtr, C.getTypeAlignInChars(Private->getType()));
2590     const ValueDecl *VD = cast<DeclRefExpr>(Private)->getDecl();
2591     // Global = Buffer.VD[Idx];
2592     const FieldDecl *FD = VarFieldMap.lookup(VD);
2593     LValue GlobLVal = CGF.EmitLValueForField(
2594         CGF.MakeNaturalAlignAddrLValue(BufferArrPtr, StaticTy), FD);
2595     Address GlobAddr = GlobLVal.getAddress(CGF);
2596     llvm::Value *BufferPtr = Bld.CreateInBoundsGEP(
2597         GlobAddr.getElementType(), GlobAddr.getPointer(), Idxs);
2598     GlobLVal.setAddress(Address(BufferPtr, GlobAddr.getAlignment()));
2599     switch (CGF.getEvaluationKind(Private->getType())) {
2600     case TEK_Scalar: {
2601       llvm::Value *V = CGF.EmitLoadOfScalar(
2602           ElemPtr, /*Volatile=*/false, Private->getType(), Loc,
2603           LValueBaseInfo(AlignmentSource::Type), TBAAAccessInfo());
2604       CGF.EmitStoreOfScalar(V, GlobLVal);
2605       break;
2606     }
2607     case TEK_Complex: {
2608       CodeGenFunction::ComplexPairTy V = CGF.EmitLoadOfComplex(
2609           CGF.MakeAddrLValue(ElemPtr, Private->getType()), Loc);
2610       CGF.EmitStoreOfComplex(V, GlobLVal, /*isInit=*/false);
2611       break;
2612     }
2613     case TEK_Aggregate:
2614       CGF.EmitAggregateCopy(GlobLVal,
2615                             CGF.MakeAddrLValue(ElemPtr, Private->getType()),
2616                             Private->getType(), AggValueSlot::DoesNotOverlap);
2617       break;
2618     }
2619     ++Idx;
2620   }
2621 
2622   CGF.FinishFunction();
2623   return Fn;
2624 }
2625 
2626 /// This function emits a helper that reduces all the reduction variables from
2627 /// the team into the provided global buffer for the reduction variables.
2628 ///
2629 /// void list_to_global_reduce_func(void *buffer, int Idx, void *reduce_data)
2630 ///  void *GlobPtrs[];
2631 ///  GlobPtrs[0] = (void*)&buffer.D0[Idx];
2632 ///  ...
2633 ///  GlobPtrs[N] = (void*)&buffer.DN[Idx];
2634 ///  reduce_function(GlobPtrs, reduce_data);
2635 static llvm::Value *emitListToGlobalReduceFunction(
2636     CodeGenModule &CGM, ArrayRef<const Expr *> Privates,
2637     QualType ReductionArrayTy, SourceLocation Loc,
2638     const RecordDecl *TeamReductionRec,
2639     const llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>
2640         &VarFieldMap,
2641     llvm::Function *ReduceFn) {
2642   ASTContext &C = CGM.getContext();
2643 
2644   // Buffer: global reduction buffer.
2645   ImplicitParamDecl BufferArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
2646                               C.VoidPtrTy, ImplicitParamDecl::Other);
2647   // Idx: index of the buffer.
2648   ImplicitParamDecl IdxArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr, C.IntTy,
2649                            ImplicitParamDecl::Other);
2650   // ReduceList: thread local Reduce list.
2651   ImplicitParamDecl ReduceListArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
2652                                   C.VoidPtrTy, ImplicitParamDecl::Other);
2653   FunctionArgList Args;
2654   Args.push_back(&BufferArg);
2655   Args.push_back(&IdxArg);
2656   Args.push_back(&ReduceListArg);
2657 
2658   const CGFunctionInfo &CGFI =
2659       CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args);
2660   auto *Fn = llvm::Function::Create(
2661       CGM.getTypes().GetFunctionType(CGFI), llvm::GlobalValue::InternalLinkage,
2662       "_omp_reduction_list_to_global_reduce_func", &CGM.getModule());
2663   CGM.SetInternalFunctionAttributes(GlobalDecl(), Fn, CGFI);
2664   Fn->setDoesNotRecurse();
2665   CodeGenFunction CGF(CGM);
2666   CGF.StartFunction(GlobalDecl(), C.VoidTy, Fn, CGFI, Args, Loc, Loc);
2667 
2668   CGBuilderTy &Bld = CGF.Builder;
2669 
2670   Address AddrBufferArg = CGF.GetAddrOfLocalVar(&BufferArg);
2671   QualType StaticTy = C.getRecordType(TeamReductionRec);
2672   llvm::Type *LLVMReductionsBufferTy =
2673       CGM.getTypes().ConvertTypeForMem(StaticTy);
2674   llvm::Value *BufferArrPtr = Bld.CreatePointerBitCastOrAddrSpaceCast(
2675       CGF.EmitLoadOfScalar(AddrBufferArg, /*Volatile=*/false, C.VoidPtrTy, Loc),
2676       LLVMReductionsBufferTy->getPointerTo());
2677 
2678   // 1. Build a list of reduction variables.
2679   // void *RedList[<n>] = {<ReductionVars>[0], ..., <ReductionVars>[<n>-1]};
2680   Address ReductionList =
2681       CGF.CreateMemTemp(ReductionArrayTy, ".omp.reduction.red_list");
2682   auto IPriv = Privates.begin();
2683   llvm::Value *Idxs[] = {llvm::ConstantInt::getNullValue(CGF.Int32Ty),
2684                          CGF.EmitLoadOfScalar(CGF.GetAddrOfLocalVar(&IdxArg),
2685                                               /*Volatile=*/false, C.IntTy,
2686                                               Loc)};
2687   unsigned Idx = 0;
2688   for (unsigned I = 0, E = Privates.size(); I < E; ++I, ++IPriv, ++Idx) {
2689     Address Elem = CGF.Builder.CreateConstArrayGEP(ReductionList, Idx);
2690     // Global = Buffer.VD[Idx];
2691     const ValueDecl *VD = cast<DeclRefExpr>(*IPriv)->getDecl();
2692     const FieldDecl *FD = VarFieldMap.lookup(VD);
2693     LValue GlobLVal = CGF.EmitLValueForField(
2694         CGF.MakeNaturalAlignAddrLValue(BufferArrPtr, StaticTy), FD);
2695     Address GlobAddr = GlobLVal.getAddress(CGF);
2696     llvm::Value *BufferPtr = Bld.CreateInBoundsGEP(
2697         GlobAddr.getElementType(), GlobAddr.getPointer(), Idxs);
2698     llvm::Value *Ptr = CGF.EmitCastToVoidPtr(BufferPtr);
2699     CGF.EmitStoreOfScalar(Ptr, Elem, /*Volatile=*/false, C.VoidPtrTy);
2700     if ((*IPriv)->getType()->isVariablyModifiedType()) {
2701       // Store array size.
2702       ++Idx;
2703       Elem = CGF.Builder.CreateConstArrayGEP(ReductionList, Idx);
2704       llvm::Value *Size = CGF.Builder.CreateIntCast(
2705           CGF.getVLASize(
2706                  CGF.getContext().getAsVariableArrayType((*IPriv)->getType()))
2707               .NumElts,
2708           CGF.SizeTy, /*isSigned=*/false);
2709       CGF.Builder.CreateStore(CGF.Builder.CreateIntToPtr(Size, CGF.VoidPtrTy),
2710                               Elem);
2711     }
2712   }
2713 
2714   // Call reduce_function(GlobalReduceList, ReduceList)
2715   llvm::Value *GlobalReduceList =
2716       CGF.EmitCastToVoidPtr(ReductionList.getPointer());
2717   Address AddrReduceListArg = CGF.GetAddrOfLocalVar(&ReduceListArg);
2718   llvm::Value *ReducedPtr = CGF.EmitLoadOfScalar(
2719       AddrReduceListArg, /*Volatile=*/false, C.VoidPtrTy, Loc);
2720   CGM.getOpenMPRuntime().emitOutlinedFunctionCall(
2721       CGF, Loc, ReduceFn, {GlobalReduceList, ReducedPtr});
2722   CGF.FinishFunction();
2723   return Fn;
2724 }
2725 
2726 /// This function emits a helper that copies all the reduction variables from
2727 /// the team into the provided global buffer for the reduction variables.
2728 ///
2729 /// void list_to_global_copy_func(void *buffer, int Idx, void *reduce_data)
2730 ///   For all data entries D in reduce_data:
2731 ///     Copy buffer.D[Idx] to local D;
2732 static llvm::Value *emitGlobalToListCopyFunction(
2733     CodeGenModule &CGM, ArrayRef<const Expr *> Privates,
2734     QualType ReductionArrayTy, SourceLocation Loc,
2735     const RecordDecl *TeamReductionRec,
2736     const llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>
2737         &VarFieldMap) {
2738   ASTContext &C = CGM.getContext();
2739 
2740   // Buffer: global reduction buffer.
2741   ImplicitParamDecl BufferArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
2742                               C.VoidPtrTy, ImplicitParamDecl::Other);
2743   // Idx: index of the buffer.
2744   ImplicitParamDecl IdxArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr, C.IntTy,
2745                            ImplicitParamDecl::Other);
2746   // ReduceList: thread local Reduce list.
2747   ImplicitParamDecl ReduceListArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
2748                                   C.VoidPtrTy, ImplicitParamDecl::Other);
2749   FunctionArgList Args;
2750   Args.push_back(&BufferArg);
2751   Args.push_back(&IdxArg);
2752   Args.push_back(&ReduceListArg);
2753 
2754   const CGFunctionInfo &CGFI =
2755       CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args);
2756   auto *Fn = llvm::Function::Create(
2757       CGM.getTypes().GetFunctionType(CGFI), llvm::GlobalValue::InternalLinkage,
2758       "_omp_reduction_global_to_list_copy_func", &CGM.getModule());
2759   CGM.SetInternalFunctionAttributes(GlobalDecl(), Fn, CGFI);
2760   Fn->setDoesNotRecurse();
2761   CodeGenFunction CGF(CGM);
2762   CGF.StartFunction(GlobalDecl(), C.VoidTy, Fn, CGFI, Args, Loc, Loc);
2763 
2764   CGBuilderTy &Bld = CGF.Builder;
2765 
2766   Address AddrReduceListArg = CGF.GetAddrOfLocalVar(&ReduceListArg);
2767   Address AddrBufferArg = CGF.GetAddrOfLocalVar(&BufferArg);
2768   Address LocalReduceList(
2769       Bld.CreatePointerBitCastOrAddrSpaceCast(
2770           CGF.EmitLoadOfScalar(AddrReduceListArg, /*Volatile=*/false,
2771                                C.VoidPtrTy, Loc),
2772           CGF.ConvertTypeForMem(ReductionArrayTy)->getPointerTo()),
2773       CGF.getPointerAlign());
2774   QualType StaticTy = C.getRecordType(TeamReductionRec);
2775   llvm::Type *LLVMReductionsBufferTy =
2776       CGM.getTypes().ConvertTypeForMem(StaticTy);
2777   llvm::Value *BufferArrPtr = Bld.CreatePointerBitCastOrAddrSpaceCast(
2778       CGF.EmitLoadOfScalar(AddrBufferArg, /*Volatile=*/false, C.VoidPtrTy, Loc),
2779       LLVMReductionsBufferTy->getPointerTo());
2780 
2781   llvm::Value *Idxs[] = {llvm::ConstantInt::getNullValue(CGF.Int32Ty),
2782                          CGF.EmitLoadOfScalar(CGF.GetAddrOfLocalVar(&IdxArg),
2783                                               /*Volatile=*/false, C.IntTy,
2784                                               Loc)};
2785   unsigned Idx = 0;
2786   for (const Expr *Private : Privates) {
2787     // Reduce element = LocalReduceList[i]
2788     Address ElemPtrPtrAddr = Bld.CreateConstArrayGEP(LocalReduceList, Idx);
2789     llvm::Value *ElemPtrPtr = CGF.EmitLoadOfScalar(
2790         ElemPtrPtrAddr, /*Volatile=*/false, C.VoidPtrTy, SourceLocation());
2791     // elemptr = ((CopyType*)(elemptrptr)) + I
2792     ElemPtrPtr = Bld.CreatePointerBitCastOrAddrSpaceCast(
2793         ElemPtrPtr, CGF.ConvertTypeForMem(Private->getType())->getPointerTo());
2794     Address ElemPtr =
2795         Address(ElemPtrPtr, C.getTypeAlignInChars(Private->getType()));
2796     const ValueDecl *VD = cast<DeclRefExpr>(Private)->getDecl();
2797     // Global = Buffer.VD[Idx];
2798     const FieldDecl *FD = VarFieldMap.lookup(VD);
2799     LValue GlobLVal = CGF.EmitLValueForField(
2800         CGF.MakeNaturalAlignAddrLValue(BufferArrPtr, StaticTy), FD);
2801     Address GlobAddr = GlobLVal.getAddress(CGF);
2802     llvm::Value *BufferPtr = Bld.CreateInBoundsGEP(
2803         GlobAddr.getElementType(), GlobAddr.getPointer(), Idxs);
2804     GlobLVal.setAddress(Address(BufferPtr, GlobAddr.getAlignment()));
2805     switch (CGF.getEvaluationKind(Private->getType())) {
2806     case TEK_Scalar: {
2807       llvm::Value *V = CGF.EmitLoadOfScalar(GlobLVal, Loc);
2808       CGF.EmitStoreOfScalar(V, ElemPtr, /*Volatile=*/false, Private->getType(),
2809                             LValueBaseInfo(AlignmentSource::Type),
2810                             TBAAAccessInfo());
2811       break;
2812     }
2813     case TEK_Complex: {
2814       CodeGenFunction::ComplexPairTy V = CGF.EmitLoadOfComplex(GlobLVal, Loc);
2815       CGF.EmitStoreOfComplex(V, CGF.MakeAddrLValue(ElemPtr, Private->getType()),
2816                              /*isInit=*/false);
2817       break;
2818     }
2819     case TEK_Aggregate:
2820       CGF.EmitAggregateCopy(CGF.MakeAddrLValue(ElemPtr, Private->getType()),
2821                             GlobLVal, Private->getType(),
2822                             AggValueSlot::DoesNotOverlap);
2823       break;
2824     }
2825     ++Idx;
2826   }
2827 
2828   CGF.FinishFunction();
2829   return Fn;
2830 }
2831 
2832 /// This function emits a helper that reduces all the reduction variables from
2833 /// the team into the provided global buffer for the reduction variables.
2834 ///
2835 /// void global_to_list_reduce_func(void *buffer, int Idx, void *reduce_data)
2836 ///  void *GlobPtrs[];
2837 ///  GlobPtrs[0] = (void*)&buffer.D0[Idx];
2838 ///  ...
2839 ///  GlobPtrs[N] = (void*)&buffer.DN[Idx];
2840 ///  reduce_function(reduce_data, GlobPtrs);
2841 static llvm::Value *emitGlobalToListReduceFunction(
2842     CodeGenModule &CGM, ArrayRef<const Expr *> Privates,
2843     QualType ReductionArrayTy, SourceLocation Loc,
2844     const RecordDecl *TeamReductionRec,
2845     const llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>
2846         &VarFieldMap,
2847     llvm::Function *ReduceFn) {
2848   ASTContext &C = CGM.getContext();
2849 
2850   // Buffer: global reduction buffer.
2851   ImplicitParamDecl BufferArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
2852                               C.VoidPtrTy, ImplicitParamDecl::Other);
2853   // Idx: index of the buffer.
2854   ImplicitParamDecl IdxArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr, C.IntTy,
2855                            ImplicitParamDecl::Other);
2856   // ReduceList: thread local Reduce list.
2857   ImplicitParamDecl ReduceListArg(C, /*DC=*/nullptr, Loc, /*Id=*/nullptr,
2858                                   C.VoidPtrTy, ImplicitParamDecl::Other);
2859   FunctionArgList Args;
2860   Args.push_back(&BufferArg);
2861   Args.push_back(&IdxArg);
2862   Args.push_back(&ReduceListArg);
2863 
2864   const CGFunctionInfo &CGFI =
2865       CGM.getTypes().arrangeBuiltinFunctionDeclaration(C.VoidTy, Args);
2866   auto *Fn = llvm::Function::Create(
2867       CGM.getTypes().GetFunctionType(CGFI), llvm::GlobalValue::InternalLinkage,
2868       "_omp_reduction_global_to_list_reduce_func", &CGM.getModule());
2869   CGM.SetInternalFunctionAttributes(GlobalDecl(), Fn, CGFI);
2870   Fn->setDoesNotRecurse();
2871   CodeGenFunction CGF(CGM);
2872   CGF.StartFunction(GlobalDecl(), C.VoidTy, Fn, CGFI, Args, Loc, Loc);
2873 
2874   CGBuilderTy &Bld = CGF.Builder;
2875 
2876   Address AddrBufferArg = CGF.GetAddrOfLocalVar(&BufferArg);
2877   QualType StaticTy = C.getRecordType(TeamReductionRec);
2878   llvm::Type *LLVMReductionsBufferTy =
2879       CGM.getTypes().ConvertTypeForMem(StaticTy);
2880   llvm::Value *BufferArrPtr = Bld.CreatePointerBitCastOrAddrSpaceCast(
2881       CGF.EmitLoadOfScalar(AddrBufferArg, /*Volatile=*/false, C.VoidPtrTy, Loc),
2882       LLVMReductionsBufferTy->getPointerTo());
2883 
2884   // 1. Build a list of reduction variables.
2885   // void *RedList[<n>] = {<ReductionVars>[0], ..., <ReductionVars>[<n>-1]};
2886   Address ReductionList =
2887       CGF.CreateMemTemp(ReductionArrayTy, ".omp.reduction.red_list");
2888   auto IPriv = Privates.begin();
2889   llvm::Value *Idxs[] = {llvm::ConstantInt::getNullValue(CGF.Int32Ty),
2890                          CGF.EmitLoadOfScalar(CGF.GetAddrOfLocalVar(&IdxArg),
2891                                               /*Volatile=*/false, C.IntTy,
2892                                               Loc)};
2893   unsigned Idx = 0;
2894   for (unsigned I = 0, E = Privates.size(); I < E; ++I, ++IPriv, ++Idx) {
2895     Address Elem = CGF.Builder.CreateConstArrayGEP(ReductionList, Idx);
2896     // Global = Buffer.VD[Idx];
2897     const ValueDecl *VD = cast<DeclRefExpr>(*IPriv)->getDecl();
2898     const FieldDecl *FD = VarFieldMap.lookup(VD);
2899     LValue GlobLVal = CGF.EmitLValueForField(
2900         CGF.MakeNaturalAlignAddrLValue(BufferArrPtr, StaticTy), FD);
2901     Address GlobAddr = GlobLVal.getAddress(CGF);
2902     llvm::Value *BufferPtr = Bld.CreateInBoundsGEP(
2903         GlobAddr.getElementType(), GlobAddr.getPointer(), Idxs);
2904     llvm::Value *Ptr = CGF.EmitCastToVoidPtr(BufferPtr);
2905     CGF.EmitStoreOfScalar(Ptr, Elem, /*Volatile=*/false, C.VoidPtrTy);
2906     if ((*IPriv)->getType()->isVariablyModifiedType()) {
2907       // Store array size.
2908       ++Idx;
2909       Elem = CGF.Builder.CreateConstArrayGEP(ReductionList, Idx);
2910       llvm::Value *Size = CGF.Builder.CreateIntCast(
2911           CGF.getVLASize(
2912                  CGF.getContext().getAsVariableArrayType((*IPriv)->getType()))
2913               .NumElts,
2914           CGF.SizeTy, /*isSigned=*/false);
2915       CGF.Builder.CreateStore(CGF.Builder.CreateIntToPtr(Size, CGF.VoidPtrTy),
2916                               Elem);
2917     }
2918   }
2919 
2920   // Call reduce_function(ReduceList, GlobalReduceList)
2921   llvm::Value *GlobalReduceList =
2922       CGF.EmitCastToVoidPtr(ReductionList.getPointer());
2923   Address AddrReduceListArg = CGF.GetAddrOfLocalVar(&ReduceListArg);
2924   llvm::Value *ReducedPtr = CGF.EmitLoadOfScalar(
2925       AddrReduceListArg, /*Volatile=*/false, C.VoidPtrTy, Loc);
2926   CGM.getOpenMPRuntime().emitOutlinedFunctionCall(
2927       CGF, Loc, ReduceFn, {ReducedPtr, GlobalReduceList});
2928   CGF.FinishFunction();
2929   return Fn;
2930 }
2931 
2932 ///
2933 /// Design of OpenMP reductions on the GPU
2934 ///
2935 /// Consider a typical OpenMP program with one or more reduction
2936 /// clauses:
2937 ///
2938 /// float foo;
2939 /// double bar;
2940 /// #pragma omp target teams distribute parallel for \
2941 ///             reduction(+:foo) reduction(*:bar)
2942 /// for (int i = 0; i < N; i++) {
2943 ///   foo += A[i]; bar *= B[i];
2944 /// }
2945 ///
2946 /// where 'foo' and 'bar' are reduced across all OpenMP threads in
2947 /// all teams.  In our OpenMP implementation on the NVPTX device an
2948 /// OpenMP team is mapped to a CUDA threadblock and OpenMP threads
2949 /// within a team are mapped to CUDA threads within a threadblock.
2950 /// Our goal is to efficiently aggregate values across all OpenMP
2951 /// threads such that:
2952 ///
2953 ///   - the compiler and runtime are logically concise, and
2954 ///   - the reduction is performed efficiently in a hierarchical
2955 ///     manner as follows: within OpenMP threads in the same warp,
2956 ///     across warps in a threadblock, and finally across teams on
2957 ///     the NVPTX device.
2958 ///
2959 /// Introduction to Decoupling
2960 ///
2961 /// We would like to decouple the compiler and the runtime so that the
2962 /// latter is ignorant of the reduction variables (number, data types)
2963 /// and the reduction operators.  This allows a simpler interface
2964 /// and implementation while still attaining good performance.
2965 ///
2966 /// Pseudocode for the aforementioned OpenMP program generated by the
2967 /// compiler is as follows:
2968 ///
2969 /// 1. Create private copies of reduction variables on each OpenMP
2970 ///    thread: 'foo_private', 'bar_private'
2971 /// 2. Each OpenMP thread reduces the chunk of 'A' and 'B' assigned
2972 ///    to it and writes the result in 'foo_private' and 'bar_private'
2973 ///    respectively.
2974 /// 3. Call the OpenMP runtime on the GPU to reduce within a team
2975 ///    and store the result on the team master:
2976 ///
2977 ///     __kmpc_nvptx_parallel_reduce_nowait_v2(...,
2978 ///        reduceData, shuffleReduceFn, interWarpCpyFn)
2979 ///
2980 ///     where:
2981 ///       struct ReduceData {
2982 ///         double *foo;
2983 ///         double *bar;
2984 ///       } reduceData
2985 ///       reduceData.foo = &foo_private
2986 ///       reduceData.bar = &bar_private
2987 ///
2988 ///     'shuffleReduceFn' and 'interWarpCpyFn' are pointers to two
2989 ///     auxiliary functions generated by the compiler that operate on
2990 ///     variables of type 'ReduceData'.  They aid the runtime perform
2991 ///     algorithmic steps in a data agnostic manner.
2992 ///
2993 ///     'shuffleReduceFn' is a pointer to a function that reduces data
2994 ///     of type 'ReduceData' across two OpenMP threads (lanes) in the
2995 ///     same warp.  It takes the following arguments as input:
2996 ///
2997 ///     a. variable of type 'ReduceData' on the calling lane,
2998 ///     b. its lane_id,
2999 ///     c. an offset relative to the current lane_id to generate a
3000 ///        remote_lane_id.  The remote lane contains the second
3001 ///        variable of type 'ReduceData' that is to be reduced.
3002 ///     d. an algorithm version parameter determining which reduction
3003 ///        algorithm to use.
3004 ///
3005 ///     'shuffleReduceFn' retrieves data from the remote lane using
3006 ///     efficient GPU shuffle intrinsics and reduces, using the
3007 ///     algorithm specified by the 4th parameter, the two operands
3008 ///     element-wise.  The result is written to the first operand.
3009 ///
3010 ///     Different reduction algorithms are implemented in different
3011 ///     runtime functions, all calling 'shuffleReduceFn' to perform
3012 ///     the essential reduction step.  Therefore, based on the 4th
3013 ///     parameter, this function behaves slightly differently to
3014 ///     cooperate with the runtime to ensure correctness under
3015 ///     different circumstances.
3016 ///
3017 ///     'InterWarpCpyFn' is a pointer to a function that transfers
3018 ///     reduced variables across warps.  It tunnels, through CUDA
3019 ///     shared memory, the thread-private data of type 'ReduceData'
3020 ///     from lane 0 of each warp to a lane in the first warp.
3021 /// 4. Call the OpenMP runtime on the GPU to reduce across teams.
3022 ///    The last team writes the global reduced value to memory.
3023 ///
3024 ///     ret = __kmpc_nvptx_teams_reduce_nowait(...,
3025 ///             reduceData, shuffleReduceFn, interWarpCpyFn,
3026 ///             scratchpadCopyFn, loadAndReduceFn)
3027 ///
3028 ///     'scratchpadCopyFn' is a helper that stores reduced
3029 ///     data from the team master to a scratchpad array in
3030 ///     global memory.
3031 ///
3032 ///     'loadAndReduceFn' is a helper that loads data from
3033 ///     the scratchpad array and reduces it with the input
3034 ///     operand.
3035 ///
3036 ///     These compiler generated functions hide address
3037 ///     calculation and alignment information from the runtime.
3038 /// 5. if ret == 1:
3039 ///     The team master of the last team stores the reduced
3040 ///     result to the globals in memory.
3041 ///     foo += reduceData.foo; bar *= reduceData.bar
3042 ///
3043 ///
3044 /// Warp Reduction Algorithms
3045 ///
3046 /// On the warp level, we have three algorithms implemented in the
3047 /// OpenMP runtime depending on the number of active lanes:
3048 ///
3049 /// Full Warp Reduction
3050 ///
3051 /// The reduce algorithm within a warp where all lanes are active
3052 /// is implemented in the runtime as follows:
3053 ///
3054 /// full_warp_reduce(void *reduce_data,
3055 ///                  kmp_ShuffleReductFctPtr ShuffleReduceFn) {
3056 ///   for (int offset = WARPSIZE/2; offset > 0; offset /= 2)
3057 ///     ShuffleReduceFn(reduce_data, 0, offset, 0);
3058 /// }
3059 ///
3060 /// The algorithm completes in log(2, WARPSIZE) steps.
3061 ///
3062 /// 'ShuffleReduceFn' is used here with lane_id set to 0 because it is
3063 /// not used therefore we save instructions by not retrieving lane_id
3064 /// from the corresponding special registers.  The 4th parameter, which
3065 /// represents the version of the algorithm being used, is set to 0 to
3066 /// signify full warp reduction.
3067 ///
3068 /// In this version, 'ShuffleReduceFn' behaves, per element, as follows:
3069 ///
3070 /// #reduce_elem refers to an element in the local lane's data structure
3071 /// #remote_elem is retrieved from a remote lane
3072 /// remote_elem = shuffle_down(reduce_elem, offset, WARPSIZE);
3073 /// reduce_elem = reduce_elem REDUCE_OP remote_elem;
3074 ///
3075 /// Contiguous Partial Warp Reduction
3076 ///
3077 /// This reduce algorithm is used within a warp where only the first
3078 /// 'n' (n <= WARPSIZE) lanes are active.  It is typically used when the
3079 /// number of OpenMP threads in a parallel region is not a multiple of
3080 /// WARPSIZE.  The algorithm is implemented in the runtime as follows:
3081 ///
3082 /// void
3083 /// contiguous_partial_reduce(void *reduce_data,
3084 ///                           kmp_ShuffleReductFctPtr ShuffleReduceFn,
3085 ///                           int size, int lane_id) {
3086 ///   int curr_size;
3087 ///   int offset;
3088 ///   curr_size = size;
3089 ///   mask = curr_size/2;
3090 ///   while (offset>0) {
3091 ///     ShuffleReduceFn(reduce_data, lane_id, offset, 1);
3092 ///     curr_size = (curr_size+1)/2;
3093 ///     offset = curr_size/2;
3094 ///   }
3095 /// }
3096 ///
3097 /// In this version, 'ShuffleReduceFn' behaves, per element, as follows:
3098 ///
3099 /// remote_elem = shuffle_down(reduce_elem, offset, WARPSIZE);
3100 /// if (lane_id < offset)
3101 ///     reduce_elem = reduce_elem REDUCE_OP remote_elem
3102 /// else
3103 ///     reduce_elem = remote_elem
3104 ///
3105 /// This algorithm assumes that the data to be reduced are located in a
3106 /// contiguous subset of lanes starting from the first.  When there is
3107 /// an odd number of active lanes, the data in the last lane is not
3108 /// aggregated with any other lane's dat but is instead copied over.
3109 ///
3110 /// Dispersed Partial Warp Reduction
3111 ///
3112 /// This algorithm is used within a warp when any discontiguous subset of
3113 /// lanes are active.  It is used to implement the reduction operation
3114 /// across lanes in an OpenMP simd region or in a nested parallel region.
3115 ///
3116 /// void
3117 /// dispersed_partial_reduce(void *reduce_data,
3118 ///                          kmp_ShuffleReductFctPtr ShuffleReduceFn) {
3119 ///   int size, remote_id;
3120 ///   int logical_lane_id = number_of_active_lanes_before_me() * 2;
3121 ///   do {
3122 ///       remote_id = next_active_lane_id_right_after_me();
3123 ///       # the above function returns 0 of no active lane
3124 ///       # is present right after the current lane.
3125 ///       size = number_of_active_lanes_in_this_warp();
3126 ///       logical_lane_id /= 2;
3127 ///       ShuffleReduceFn(reduce_data, logical_lane_id,
3128 ///                       remote_id-1-threadIdx.x, 2);
3129 ///   } while (logical_lane_id % 2 == 0 && size > 1);
3130 /// }
3131 ///
3132 /// There is no assumption made about the initial state of the reduction.
3133 /// Any number of lanes (>=1) could be active at any position.  The reduction
3134 /// result is returned in the first active lane.
3135 ///
3136 /// In this version, 'ShuffleReduceFn' behaves, per element, as follows:
3137 ///
3138 /// remote_elem = shuffle_down(reduce_elem, offset, WARPSIZE);
3139 /// if (lane_id % 2 == 0 && offset > 0)
3140 ///     reduce_elem = reduce_elem REDUCE_OP remote_elem
3141 /// else
3142 ///     reduce_elem = remote_elem
3143 ///
3144 ///
3145 /// Intra-Team Reduction
3146 ///
3147 /// This function, as implemented in the runtime call
3148 /// '__kmpc_nvptx_parallel_reduce_nowait_v2', aggregates data across OpenMP
3149 /// threads in a team.  It first reduces within a warp using the
3150 /// aforementioned algorithms.  We then proceed to gather all such
3151 /// reduced values at the first warp.
3152 ///
3153 /// The runtime makes use of the function 'InterWarpCpyFn', which copies
3154 /// data from each of the "warp master" (zeroth lane of each warp, where
3155 /// warp-reduced data is held) to the zeroth warp.  This step reduces (in
3156 /// a mathematical sense) the problem of reduction across warp masters in
3157 /// a block to the problem of warp reduction.
3158 ///
3159 ///
3160 /// Inter-Team Reduction
3161 ///
3162 /// Once a team has reduced its data to a single value, it is stored in
3163 /// a global scratchpad array.  Since each team has a distinct slot, this
3164 /// can be done without locking.
3165 ///
3166 /// The last team to write to the scratchpad array proceeds to reduce the
3167 /// scratchpad array.  One or more workers in the last team use the helper
3168 /// 'loadAndReduceDataFn' to load and reduce values from the array, i.e.,
3169 /// the k'th worker reduces every k'th element.
3170 ///
3171 /// Finally, a call is made to '__kmpc_nvptx_parallel_reduce_nowait_v2' to
3172 /// reduce across workers and compute a globally reduced value.
3173 ///
3174 void CGOpenMPRuntimeGPU::emitReduction(
3175     CodeGenFunction &CGF, SourceLocation Loc, ArrayRef<const Expr *> Privates,
3176     ArrayRef<const Expr *> LHSExprs, ArrayRef<const Expr *> RHSExprs,
3177     ArrayRef<const Expr *> ReductionOps, ReductionOptionsTy Options) {
3178   if (!CGF.HaveInsertPoint())
3179     return;
3180 
3181   bool ParallelReduction = isOpenMPParallelDirective(Options.ReductionKind);
3182 #ifndef NDEBUG
3183   bool TeamsReduction = isOpenMPTeamsDirective(Options.ReductionKind);
3184 #endif
3185 
3186   if (Options.SimpleReduction) {
3187     assert(!TeamsReduction && !ParallelReduction &&
3188            "Invalid reduction selection in emitReduction.");
3189     CGOpenMPRuntime::emitReduction(CGF, Loc, Privates, LHSExprs, RHSExprs,
3190                                    ReductionOps, Options);
3191     return;
3192   }
3193 
3194   assert((TeamsReduction || ParallelReduction) &&
3195          "Invalid reduction selection in emitReduction.");
3196 
3197   // Build res = __kmpc_reduce{_nowait}(<gtid>, <n>, sizeof(RedList),
3198   // RedList, shuffle_reduce_func, interwarp_copy_func);
3199   // or
3200   // Build res = __kmpc_reduce_teams_nowait_simple(<loc>, <gtid>, <lck>);
3201   llvm::Value *RTLoc = emitUpdateLocation(CGF, Loc);
3202   llvm::Value *ThreadId = getThreadID(CGF, Loc);
3203 
3204   llvm::Value *Res;
3205   ASTContext &C = CGM.getContext();
3206   // 1. Build a list of reduction variables.
3207   // void *RedList[<n>] = {<ReductionVars>[0], ..., <ReductionVars>[<n>-1]};
3208   auto Size = RHSExprs.size();
3209   for (const Expr *E : Privates) {
3210     if (E->getType()->isVariablyModifiedType())
3211       // Reserve place for array size.
3212       ++Size;
3213   }
3214   llvm::APInt ArraySize(/*unsigned int numBits=*/32, Size);
3215   QualType ReductionArrayTy =
3216       C.getConstantArrayType(C.VoidPtrTy, ArraySize, nullptr, ArrayType::Normal,
3217                              /*IndexTypeQuals=*/0);
3218   Address ReductionList =
3219       CGF.CreateMemTemp(ReductionArrayTy, ".omp.reduction.red_list");
3220   auto IPriv = Privates.begin();
3221   unsigned Idx = 0;
3222   for (unsigned I = 0, E = RHSExprs.size(); I < E; ++I, ++IPriv, ++Idx) {
3223     Address Elem = CGF.Builder.CreateConstArrayGEP(ReductionList, Idx);
3224     CGF.Builder.CreateStore(
3225         CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
3226             CGF.EmitLValue(RHSExprs[I]).getPointer(CGF), CGF.VoidPtrTy),
3227         Elem);
3228     if ((*IPriv)->getType()->isVariablyModifiedType()) {
3229       // Store array size.
3230       ++Idx;
3231       Elem = CGF.Builder.CreateConstArrayGEP(ReductionList, Idx);
3232       llvm::Value *Size = CGF.Builder.CreateIntCast(
3233           CGF.getVLASize(
3234                  CGF.getContext().getAsVariableArrayType((*IPriv)->getType()))
3235               .NumElts,
3236           CGF.SizeTy, /*isSigned=*/false);
3237       CGF.Builder.CreateStore(CGF.Builder.CreateIntToPtr(Size, CGF.VoidPtrTy),
3238                               Elem);
3239     }
3240   }
3241 
3242   llvm::Value *RL = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
3243       ReductionList.getPointer(), CGF.VoidPtrTy);
3244   llvm::Function *ReductionFn = emitReductionFunction(
3245       Loc, CGF.ConvertTypeForMem(ReductionArrayTy)->getPointerTo(), Privates,
3246       LHSExprs, RHSExprs, ReductionOps);
3247   llvm::Value *ReductionArrayTySize = CGF.getTypeSize(ReductionArrayTy);
3248   llvm::Function *ShuffleAndReduceFn = emitShuffleAndReduceFunction(
3249       CGM, Privates, ReductionArrayTy, ReductionFn, Loc);
3250   llvm::Value *InterWarpCopyFn =
3251       emitInterWarpCopyFunction(CGM, Privates, ReductionArrayTy, Loc);
3252 
3253   if (ParallelReduction) {
3254     llvm::Value *Args[] = {RTLoc,
3255                            ThreadId,
3256                            CGF.Builder.getInt32(RHSExprs.size()),
3257                            ReductionArrayTySize,
3258                            RL,
3259                            ShuffleAndReduceFn,
3260                            InterWarpCopyFn};
3261 
3262     Res = CGF.EmitRuntimeCall(
3263         OMPBuilder.getOrCreateRuntimeFunction(
3264             CGM.getModule(), OMPRTL___kmpc_nvptx_parallel_reduce_nowait_v2),
3265         Args);
3266   } else {
3267     assert(TeamsReduction && "expected teams reduction.");
3268     llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *> VarFieldMap;
3269     llvm::SmallVector<const ValueDecl *, 4> PrivatesReductions(Privates.size());
3270     int Cnt = 0;
3271     for (const Expr *DRE : Privates) {
3272       PrivatesReductions[Cnt] = cast<DeclRefExpr>(DRE)->getDecl();
3273       ++Cnt;
3274     }
3275     const RecordDecl *TeamReductionRec = ::buildRecordForGlobalizedVars(
3276         CGM.getContext(), PrivatesReductions, llvm::None, VarFieldMap,
3277         C.getLangOpts().OpenMPCUDAReductionBufNum);
3278     TeamsReductions.push_back(TeamReductionRec);
3279     if (!KernelTeamsReductionPtr) {
3280       KernelTeamsReductionPtr = new llvm::GlobalVariable(
3281           CGM.getModule(), CGM.VoidPtrTy, /*isConstant=*/true,
3282           llvm::GlobalValue::InternalLinkage, nullptr,
3283           "_openmp_teams_reductions_buffer_$_$ptr");
3284     }
3285     llvm::Value *GlobalBufferPtr = CGF.EmitLoadOfScalar(
3286         Address(KernelTeamsReductionPtr, CGM.getPointerAlign()),
3287         /*Volatile=*/false, C.getPointerType(C.VoidPtrTy), Loc);
3288     llvm::Value *GlobalToBufferCpyFn = ::emitListToGlobalCopyFunction(
3289         CGM, Privates, ReductionArrayTy, Loc, TeamReductionRec, VarFieldMap);
3290     llvm::Value *GlobalToBufferRedFn = ::emitListToGlobalReduceFunction(
3291         CGM, Privates, ReductionArrayTy, Loc, TeamReductionRec, VarFieldMap,
3292         ReductionFn);
3293     llvm::Value *BufferToGlobalCpyFn = ::emitGlobalToListCopyFunction(
3294         CGM, Privates, ReductionArrayTy, Loc, TeamReductionRec, VarFieldMap);
3295     llvm::Value *BufferToGlobalRedFn = ::emitGlobalToListReduceFunction(
3296         CGM, Privates, ReductionArrayTy, Loc, TeamReductionRec, VarFieldMap,
3297         ReductionFn);
3298 
3299     llvm::Value *Args[] = {
3300         RTLoc,
3301         ThreadId,
3302         GlobalBufferPtr,
3303         CGF.Builder.getInt32(C.getLangOpts().OpenMPCUDAReductionBufNum),
3304         RL,
3305         ShuffleAndReduceFn,
3306         InterWarpCopyFn,
3307         GlobalToBufferCpyFn,
3308         GlobalToBufferRedFn,
3309         BufferToGlobalCpyFn,
3310         BufferToGlobalRedFn};
3311 
3312     Res = CGF.EmitRuntimeCall(
3313         OMPBuilder.getOrCreateRuntimeFunction(
3314             CGM.getModule(), OMPRTL___kmpc_nvptx_teams_reduce_nowait_v2),
3315         Args);
3316   }
3317 
3318   // 5. Build if (res == 1)
3319   llvm::BasicBlock *ExitBB = CGF.createBasicBlock(".omp.reduction.done");
3320   llvm::BasicBlock *ThenBB = CGF.createBasicBlock(".omp.reduction.then");
3321   llvm::Value *Cond = CGF.Builder.CreateICmpEQ(
3322       Res, llvm::ConstantInt::get(CGM.Int32Ty, /*V=*/1));
3323   CGF.Builder.CreateCondBr(Cond, ThenBB, ExitBB);
3324 
3325   // 6. Build then branch: where we have reduced values in the master
3326   //    thread in each team.
3327   //    __kmpc_end_reduce{_nowait}(<gtid>);
3328   //    break;
3329   CGF.EmitBlock(ThenBB);
3330 
3331   // Add emission of __kmpc_end_reduce{_nowait}(<gtid>);
3332   auto &&CodeGen = [Privates, LHSExprs, RHSExprs, ReductionOps,
3333                     this](CodeGenFunction &CGF, PrePostActionTy &Action) {
3334     auto IPriv = Privates.begin();
3335     auto ILHS = LHSExprs.begin();
3336     auto IRHS = RHSExprs.begin();
3337     for (const Expr *E : ReductionOps) {
3338       emitSingleReductionCombiner(CGF, E, *IPriv, cast<DeclRefExpr>(*ILHS),
3339                                   cast<DeclRefExpr>(*IRHS));
3340       ++IPriv;
3341       ++ILHS;
3342       ++IRHS;
3343     }
3344   };
3345   llvm::Value *EndArgs[] = {ThreadId};
3346   RegionCodeGenTy RCG(CodeGen);
3347   NVPTXActionTy Action(
3348       nullptr, llvm::None,
3349       OMPBuilder.getOrCreateRuntimeFunction(
3350           CGM.getModule(), OMPRTL___kmpc_nvptx_end_reduce_nowait),
3351       EndArgs);
3352   RCG.setAction(Action);
3353   RCG(CGF);
3354   // There is no need to emit line number for unconditional branch.
3355   (void)ApplyDebugLocation::CreateEmpty(CGF);
3356   CGF.EmitBlock(ExitBB, /*IsFinished=*/true);
3357 }
3358 
3359 const VarDecl *
3360 CGOpenMPRuntimeGPU::translateParameter(const FieldDecl *FD,
3361                                        const VarDecl *NativeParam) const {
3362   if (!NativeParam->getType()->isReferenceType())
3363     return NativeParam;
3364   QualType ArgType = NativeParam->getType();
3365   QualifierCollector QC;
3366   const Type *NonQualTy = QC.strip(ArgType);
3367   QualType PointeeTy = cast<ReferenceType>(NonQualTy)->getPointeeType();
3368   if (const auto *Attr = FD->getAttr<OMPCaptureKindAttr>()) {
3369     if (Attr->getCaptureKind() == OMPC_map) {
3370       PointeeTy = CGM.getContext().getAddrSpaceQualType(PointeeTy,
3371                                                         LangAS::opencl_global);
3372     }
3373   }
3374   ArgType = CGM.getContext().getPointerType(PointeeTy);
3375   QC.addRestrict();
3376   enum { NVPTX_local_addr = 5 };
3377   QC.addAddressSpace(getLangASFromTargetAS(NVPTX_local_addr));
3378   ArgType = QC.apply(CGM.getContext(), ArgType);
3379   if (isa<ImplicitParamDecl>(NativeParam))
3380     return ImplicitParamDecl::Create(
3381         CGM.getContext(), /*DC=*/nullptr, NativeParam->getLocation(),
3382         NativeParam->getIdentifier(), ArgType, ImplicitParamDecl::Other);
3383   return ParmVarDecl::Create(
3384       CGM.getContext(),
3385       const_cast<DeclContext *>(NativeParam->getDeclContext()),
3386       NativeParam->getBeginLoc(), NativeParam->getLocation(),
3387       NativeParam->getIdentifier(), ArgType,
3388       /*TInfo=*/nullptr, SC_None, /*DefArg=*/nullptr);
3389 }
3390 
3391 Address
3392 CGOpenMPRuntimeGPU::getParameterAddress(CodeGenFunction &CGF,
3393                                           const VarDecl *NativeParam,
3394                                           const VarDecl *TargetParam) const {
3395   assert(NativeParam != TargetParam &&
3396          NativeParam->getType()->isReferenceType() &&
3397          "Native arg must not be the same as target arg.");
3398   Address LocalAddr = CGF.GetAddrOfLocalVar(TargetParam);
3399   QualType NativeParamType = NativeParam->getType();
3400   QualifierCollector QC;
3401   const Type *NonQualTy = QC.strip(NativeParamType);
3402   QualType NativePointeeTy = cast<ReferenceType>(NonQualTy)->getPointeeType();
3403   unsigned NativePointeeAddrSpace =
3404       CGF.getContext().getTargetAddressSpace(NativePointeeTy);
3405   QualType TargetTy = TargetParam->getType();
3406   llvm::Value *TargetAddr = CGF.EmitLoadOfScalar(
3407       LocalAddr, /*Volatile=*/false, TargetTy, SourceLocation());
3408   // First cast to generic.
3409   TargetAddr = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
3410       TargetAddr, llvm::PointerType::getWithSamePointeeType(
3411           cast<llvm::PointerType>(TargetAddr->getType()), /*AddrSpace=*/0));
3412   // Cast from generic to native address space.
3413   TargetAddr = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
3414       TargetAddr, llvm::PointerType::getWithSamePointeeType(
3415           cast<llvm::PointerType>(TargetAddr->getType()),
3416                                   NativePointeeAddrSpace));
3417   Address NativeParamAddr = CGF.CreateMemTemp(NativeParamType);
3418   CGF.EmitStoreOfScalar(TargetAddr, NativeParamAddr, /*Volatile=*/false,
3419                         NativeParamType);
3420   return NativeParamAddr;
3421 }
3422 
3423 void CGOpenMPRuntimeGPU::emitOutlinedFunctionCall(
3424     CodeGenFunction &CGF, SourceLocation Loc, llvm::FunctionCallee OutlinedFn,
3425     ArrayRef<llvm::Value *> Args) const {
3426   SmallVector<llvm::Value *, 4> TargetArgs;
3427   TargetArgs.reserve(Args.size());
3428   auto *FnType = OutlinedFn.getFunctionType();
3429   for (unsigned I = 0, E = Args.size(); I < E; ++I) {
3430     if (FnType->isVarArg() && FnType->getNumParams() <= I) {
3431       TargetArgs.append(std::next(Args.begin(), I), Args.end());
3432       break;
3433     }
3434     llvm::Type *TargetType = FnType->getParamType(I);
3435     llvm::Value *NativeArg = Args[I];
3436     if (!TargetType->isPointerTy()) {
3437       TargetArgs.emplace_back(NativeArg);
3438       continue;
3439     }
3440     llvm::Value *TargetArg = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
3441         NativeArg, llvm::PointerType::getWithSamePointeeType(
3442             cast<llvm::PointerType>(NativeArg->getType()), /*AddrSpace*/ 0));
3443     TargetArgs.emplace_back(
3444         CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(TargetArg, TargetType));
3445   }
3446   CGOpenMPRuntime::emitOutlinedFunctionCall(CGF, Loc, OutlinedFn, TargetArgs);
3447 }
3448 
3449 /// Emit function which wraps the outline parallel region
3450 /// and controls the arguments which are passed to this function.
3451 /// The wrapper ensures that the outlined function is called
3452 /// with the correct arguments when data is shared.
3453 llvm::Function *CGOpenMPRuntimeGPU::createParallelDataSharingWrapper(
3454     llvm::Function *OutlinedParallelFn, const OMPExecutableDirective &D) {
3455   ASTContext &Ctx = CGM.getContext();
3456   const auto &CS = *D.getCapturedStmt(OMPD_parallel);
3457 
3458   // Create a function that takes as argument the source thread.
3459   FunctionArgList WrapperArgs;
3460   QualType Int16QTy =
3461       Ctx.getIntTypeForBitwidth(/*DestWidth=*/16, /*Signed=*/false);
3462   QualType Int32QTy =
3463       Ctx.getIntTypeForBitwidth(/*DestWidth=*/32, /*Signed=*/false);
3464   ImplicitParamDecl ParallelLevelArg(Ctx, /*DC=*/nullptr, D.getBeginLoc(),
3465                                      /*Id=*/nullptr, Int16QTy,
3466                                      ImplicitParamDecl::Other);
3467   ImplicitParamDecl WrapperArg(Ctx, /*DC=*/nullptr, D.getBeginLoc(),
3468                                /*Id=*/nullptr, Int32QTy,
3469                                ImplicitParamDecl::Other);
3470   WrapperArgs.emplace_back(&ParallelLevelArg);
3471   WrapperArgs.emplace_back(&WrapperArg);
3472 
3473   const CGFunctionInfo &CGFI =
3474       CGM.getTypes().arrangeBuiltinFunctionDeclaration(Ctx.VoidTy, WrapperArgs);
3475 
3476   auto *Fn = llvm::Function::Create(
3477       CGM.getTypes().GetFunctionType(CGFI), llvm::GlobalValue::InternalLinkage,
3478       Twine(OutlinedParallelFn->getName(), "_wrapper"), &CGM.getModule());
3479 
3480   // Ensure we do not inline the function. This is trivially true for the ones
3481   // passed to __kmpc_fork_call but the ones calles in serialized regions
3482   // could be inlined. This is not a perfect but it is closer to the invariant
3483   // we want, namely, every data environment starts with a new function.
3484   // TODO: We should pass the if condition to the runtime function and do the
3485   //       handling there. Much cleaner code.
3486   Fn->addFnAttr(llvm::Attribute::NoInline);
3487 
3488   CGM.SetInternalFunctionAttributes(GlobalDecl(), Fn, CGFI);
3489   Fn->setLinkage(llvm::GlobalValue::InternalLinkage);
3490   Fn->setDoesNotRecurse();
3491 
3492   CodeGenFunction CGF(CGM, /*suppressNewContext=*/true);
3493   CGF.StartFunction(GlobalDecl(), Ctx.VoidTy, Fn, CGFI, WrapperArgs,
3494                     D.getBeginLoc(), D.getBeginLoc());
3495 
3496   const auto *RD = CS.getCapturedRecordDecl();
3497   auto CurField = RD->field_begin();
3498 
3499   Address ZeroAddr = CGF.CreateDefaultAlignTempAlloca(CGF.Int32Ty,
3500                                                       /*Name=*/".zero.addr");
3501   CGF.Builder.CreateStore(CGF.Builder.getInt32(/*C*/ 0), ZeroAddr);
3502   // Get the array of arguments.
3503   SmallVector<llvm::Value *, 8> Args;
3504 
3505   Args.emplace_back(CGF.GetAddrOfLocalVar(&WrapperArg).getPointer());
3506   Args.emplace_back(ZeroAddr.getPointer());
3507 
3508   CGBuilderTy &Bld = CGF.Builder;
3509   auto CI = CS.capture_begin();
3510 
3511   // Use global memory for data sharing.
3512   // Handle passing of global args to workers.
3513   Address GlobalArgs =
3514       CGF.CreateDefaultAlignTempAlloca(CGF.VoidPtrPtrTy, "global_args");
3515   llvm::Value *GlobalArgsPtr = GlobalArgs.getPointer();
3516   llvm::Value *DataSharingArgs[] = {GlobalArgsPtr};
3517   CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
3518                           CGM.getModule(), OMPRTL___kmpc_get_shared_variables),
3519                       DataSharingArgs);
3520 
3521   // Retrieve the shared variables from the list of references returned
3522   // by the runtime. Pass the variables to the outlined function.
3523   Address SharedArgListAddress = Address::invalid();
3524   if (CS.capture_size() > 0 ||
3525       isOpenMPLoopBoundSharingDirective(D.getDirectiveKind())) {
3526     SharedArgListAddress = CGF.EmitLoadOfPointer(
3527         GlobalArgs, CGF.getContext()
3528                         .getPointerType(CGF.getContext().getPointerType(
3529                             CGF.getContext().VoidPtrTy))
3530                         .castAs<PointerType>());
3531   }
3532   unsigned Idx = 0;
3533   if (isOpenMPLoopBoundSharingDirective(D.getDirectiveKind())) {
3534     Address Src = Bld.CreateConstInBoundsGEP(SharedArgListAddress, Idx);
3535     Address TypedAddress = Bld.CreatePointerBitCastOrAddrSpaceCast(
3536         Src, CGF.SizeTy->getPointerTo());
3537     llvm::Value *LB = CGF.EmitLoadOfScalar(
3538         TypedAddress,
3539         /*Volatile=*/false,
3540         CGF.getContext().getPointerType(CGF.getContext().getSizeType()),
3541         cast<OMPLoopDirective>(D).getLowerBoundVariable()->getExprLoc());
3542     Args.emplace_back(LB);
3543     ++Idx;
3544     Src = Bld.CreateConstInBoundsGEP(SharedArgListAddress, Idx);
3545     TypedAddress = Bld.CreatePointerBitCastOrAddrSpaceCast(
3546         Src, CGF.SizeTy->getPointerTo());
3547     llvm::Value *UB = CGF.EmitLoadOfScalar(
3548         TypedAddress,
3549         /*Volatile=*/false,
3550         CGF.getContext().getPointerType(CGF.getContext().getSizeType()),
3551         cast<OMPLoopDirective>(D).getUpperBoundVariable()->getExprLoc());
3552     Args.emplace_back(UB);
3553     ++Idx;
3554   }
3555   if (CS.capture_size() > 0) {
3556     ASTContext &CGFContext = CGF.getContext();
3557     for (unsigned I = 0, E = CS.capture_size(); I < E; ++I, ++CI, ++CurField) {
3558       QualType ElemTy = CurField->getType();
3559       Address Src = Bld.CreateConstInBoundsGEP(SharedArgListAddress, I + Idx);
3560       Address TypedAddress = Bld.CreatePointerBitCastOrAddrSpaceCast(
3561           Src, CGF.ConvertTypeForMem(CGFContext.getPointerType(ElemTy)));
3562       llvm::Value *Arg = CGF.EmitLoadOfScalar(TypedAddress,
3563                                               /*Volatile=*/false,
3564                                               CGFContext.getPointerType(ElemTy),
3565                                               CI->getLocation());
3566       if (CI->capturesVariableByCopy() &&
3567           !CI->getCapturedVar()->getType()->isAnyPointerType()) {
3568         Arg = castValueToType(CGF, Arg, ElemTy, CGFContext.getUIntPtrType(),
3569                               CI->getLocation());
3570       }
3571       Args.emplace_back(Arg);
3572     }
3573   }
3574 
3575   emitOutlinedFunctionCall(CGF, D.getBeginLoc(), OutlinedParallelFn, Args);
3576   CGF.FinishFunction();
3577   return Fn;
3578 }
3579 
3580 void CGOpenMPRuntimeGPU::emitFunctionProlog(CodeGenFunction &CGF,
3581                                               const Decl *D) {
3582   if (getDataSharingMode(CGM) != CGOpenMPRuntimeGPU::Generic)
3583     return;
3584 
3585   assert(D && "Expected function or captured|block decl.");
3586   assert(FunctionGlobalizedDecls.count(CGF.CurFn) == 0 &&
3587          "Function is registered already.");
3588   assert((!TeamAndReductions.first || TeamAndReductions.first == D) &&
3589          "Team is set but not processed.");
3590   const Stmt *Body = nullptr;
3591   bool NeedToDelayGlobalization = false;
3592   if (const auto *FD = dyn_cast<FunctionDecl>(D)) {
3593     Body = FD->getBody();
3594   } else if (const auto *BD = dyn_cast<BlockDecl>(D)) {
3595     Body = BD->getBody();
3596   } else if (const auto *CD = dyn_cast<CapturedDecl>(D)) {
3597     Body = CD->getBody();
3598     NeedToDelayGlobalization = CGF.CapturedStmtInfo->getKind() == CR_OpenMP;
3599     if (NeedToDelayGlobalization &&
3600         getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD)
3601       return;
3602   }
3603   if (!Body)
3604     return;
3605   CheckVarsEscapingDeclContext VarChecker(CGF, TeamAndReductions.second);
3606   VarChecker.Visit(Body);
3607   const RecordDecl *GlobalizedVarsRecord =
3608       VarChecker.getGlobalizedRecord(IsInTTDRegion);
3609   TeamAndReductions.first = nullptr;
3610   TeamAndReductions.second.clear();
3611   ArrayRef<const ValueDecl *> EscapedVariableLengthDecls =
3612       VarChecker.getEscapedVariableLengthDecls();
3613   if (!GlobalizedVarsRecord && EscapedVariableLengthDecls.empty())
3614     return;
3615   auto I = FunctionGlobalizedDecls.try_emplace(CGF.CurFn).first;
3616   I->getSecond().MappedParams =
3617       std::make_unique<CodeGenFunction::OMPMapVars>();
3618   I->getSecond().EscapedParameters.insert(
3619       VarChecker.getEscapedParameters().begin(),
3620       VarChecker.getEscapedParameters().end());
3621   I->getSecond().EscapedVariableLengthDecls.append(
3622       EscapedVariableLengthDecls.begin(), EscapedVariableLengthDecls.end());
3623   DeclToAddrMapTy &Data = I->getSecond().LocalVarData;
3624   for (const ValueDecl *VD : VarChecker.getEscapedDecls()) {
3625     assert(VD->isCanonicalDecl() && "Expected canonical declaration");
3626     Data.insert(std::make_pair(VD, MappedVarData()));
3627   }
3628   if (!IsInTTDRegion && !NeedToDelayGlobalization && !IsInParallelRegion) {
3629     CheckVarsEscapingDeclContext VarChecker(CGF, llvm::None);
3630     VarChecker.Visit(Body);
3631     I->getSecond().SecondaryLocalVarData.emplace();
3632     DeclToAddrMapTy &Data = I->getSecond().SecondaryLocalVarData.getValue();
3633     for (const ValueDecl *VD : VarChecker.getEscapedDecls()) {
3634       assert(VD->isCanonicalDecl() && "Expected canonical declaration");
3635       Data.insert(std::make_pair(VD, MappedVarData()));
3636     }
3637   }
3638   if (!NeedToDelayGlobalization) {
3639     emitGenericVarsProlog(CGF, D->getBeginLoc(), /*WithSPMDCheck=*/true);
3640     struct GlobalizationScope final : EHScopeStack::Cleanup {
3641       GlobalizationScope() = default;
3642 
3643       void Emit(CodeGenFunction &CGF, Flags flags) override {
3644         static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime())
3645             .emitGenericVarsEpilog(CGF, /*WithSPMDCheck=*/true);
3646       }
3647     };
3648     CGF.EHStack.pushCleanup<GlobalizationScope>(NormalAndEHCleanup);
3649   }
3650 }
3651 
3652 Address CGOpenMPRuntimeGPU::getAddressOfLocalVariable(CodeGenFunction &CGF,
3653                                                         const VarDecl *VD) {
3654   if (VD && VD->hasAttr<OMPAllocateDeclAttr>()) {
3655     const auto *A = VD->getAttr<OMPAllocateDeclAttr>();
3656     auto AS = LangAS::Default;
3657     switch (A->getAllocatorType()) {
3658       // Use the default allocator here as by default local vars are
3659       // threadlocal.
3660     case OMPAllocateDeclAttr::OMPNullMemAlloc:
3661     case OMPAllocateDeclAttr::OMPDefaultMemAlloc:
3662     case OMPAllocateDeclAttr::OMPThreadMemAlloc:
3663     case OMPAllocateDeclAttr::OMPHighBWMemAlloc:
3664     case OMPAllocateDeclAttr::OMPLowLatMemAlloc:
3665       // Follow the user decision - use default allocation.
3666       return Address::invalid();
3667     case OMPAllocateDeclAttr::OMPUserDefinedMemAlloc:
3668       // TODO: implement aupport for user-defined allocators.
3669       return Address::invalid();
3670     case OMPAllocateDeclAttr::OMPConstMemAlloc:
3671       AS = LangAS::cuda_constant;
3672       break;
3673     case OMPAllocateDeclAttr::OMPPTeamMemAlloc:
3674       AS = LangAS::cuda_shared;
3675       break;
3676     case OMPAllocateDeclAttr::OMPLargeCapMemAlloc:
3677     case OMPAllocateDeclAttr::OMPCGroupMemAlloc:
3678       break;
3679     }
3680     llvm::Type *VarTy = CGF.ConvertTypeForMem(VD->getType());
3681     auto *GV = new llvm::GlobalVariable(
3682         CGM.getModule(), VarTy, /*isConstant=*/false,
3683         llvm::GlobalValue::InternalLinkage, llvm::Constant::getNullValue(VarTy),
3684         VD->getName(),
3685         /*InsertBefore=*/nullptr, llvm::GlobalValue::NotThreadLocal,
3686         CGM.getContext().getTargetAddressSpace(AS));
3687     CharUnits Align = CGM.getContext().getDeclAlign(VD);
3688     GV->setAlignment(Align.getAsAlign());
3689     return Address(
3690         CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
3691             GV, VarTy->getPointerTo(CGM.getContext().getTargetAddressSpace(
3692                     VD->getType().getAddressSpace()))),
3693         Align);
3694   }
3695 
3696   if (getDataSharingMode(CGM) != CGOpenMPRuntimeGPU::Generic)
3697     return Address::invalid();
3698 
3699   VD = VD->getCanonicalDecl();
3700   auto I = FunctionGlobalizedDecls.find(CGF.CurFn);
3701   if (I == FunctionGlobalizedDecls.end())
3702     return Address::invalid();
3703   auto VDI = I->getSecond().LocalVarData.find(VD);
3704   if (VDI != I->getSecond().LocalVarData.end())
3705     return VDI->second.PrivateAddr;
3706   if (VD->hasAttrs()) {
3707     for (specific_attr_iterator<OMPReferencedVarAttr> IT(VD->attr_begin()),
3708          E(VD->attr_end());
3709          IT != E; ++IT) {
3710       auto VDI = I->getSecond().LocalVarData.find(
3711           cast<VarDecl>(cast<DeclRefExpr>(IT->getRef())->getDecl())
3712               ->getCanonicalDecl());
3713       if (VDI != I->getSecond().LocalVarData.end())
3714         return VDI->second.PrivateAddr;
3715     }
3716   }
3717 
3718   return Address::invalid();
3719 }
3720 
3721 void CGOpenMPRuntimeGPU::functionFinished(CodeGenFunction &CGF) {
3722   FunctionGlobalizedDecls.erase(CGF.CurFn);
3723   CGOpenMPRuntime::functionFinished(CGF);
3724 }
3725 
3726 void CGOpenMPRuntimeGPU::getDefaultDistScheduleAndChunk(
3727     CodeGenFunction &CGF, const OMPLoopDirective &S,
3728     OpenMPDistScheduleClauseKind &ScheduleKind,
3729     llvm::Value *&Chunk) const {
3730   auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
3731   if (getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD) {
3732     ScheduleKind = OMPC_DIST_SCHEDULE_static;
3733     Chunk = CGF.EmitScalarConversion(
3734         RT.getGPUNumThreads(CGF),
3735         CGF.getContext().getIntTypeForBitwidth(32, /*Signed=*/0),
3736         S.getIterationVariable()->getType(), S.getBeginLoc());
3737     return;
3738   }
3739   CGOpenMPRuntime::getDefaultDistScheduleAndChunk(
3740       CGF, S, ScheduleKind, Chunk);
3741 }
3742 
3743 void CGOpenMPRuntimeGPU::getDefaultScheduleAndChunk(
3744     CodeGenFunction &CGF, const OMPLoopDirective &S,
3745     OpenMPScheduleClauseKind &ScheduleKind,
3746     const Expr *&ChunkExpr) const {
3747   ScheduleKind = OMPC_SCHEDULE_static;
3748   // Chunk size is 1 in this case.
3749   llvm::APInt ChunkSize(32, 1);
3750   ChunkExpr = IntegerLiteral::Create(CGF.getContext(), ChunkSize,
3751       CGF.getContext().getIntTypeForBitwidth(32, /*Signed=*/0),
3752       SourceLocation());
3753 }
3754 
3755 void CGOpenMPRuntimeGPU::adjustTargetSpecificDataForLambdas(
3756     CodeGenFunction &CGF, const OMPExecutableDirective &D) const {
3757   assert(isOpenMPTargetExecutionDirective(D.getDirectiveKind()) &&
3758          " Expected target-based directive.");
3759   const CapturedStmt *CS = D.getCapturedStmt(OMPD_target);
3760   for (const CapturedStmt::Capture &C : CS->captures()) {
3761     // Capture variables captured by reference in lambdas for target-based
3762     // directives.
3763     if (!C.capturesVariable())
3764       continue;
3765     const VarDecl *VD = C.getCapturedVar();
3766     const auto *RD = VD->getType()
3767                          .getCanonicalType()
3768                          .getNonReferenceType()
3769                          ->getAsCXXRecordDecl();
3770     if (!RD || !RD->isLambda())
3771       continue;
3772     Address VDAddr = CGF.GetAddrOfLocalVar(VD);
3773     LValue VDLVal;
3774     if (VD->getType().getCanonicalType()->isReferenceType())
3775       VDLVal = CGF.EmitLoadOfReferenceLValue(VDAddr, VD->getType());
3776     else
3777       VDLVal = CGF.MakeAddrLValue(
3778           VDAddr, VD->getType().getCanonicalType().getNonReferenceType());
3779     llvm::DenseMap<const VarDecl *, FieldDecl *> Captures;
3780     FieldDecl *ThisCapture = nullptr;
3781     RD->getCaptureFields(Captures, ThisCapture);
3782     if (ThisCapture && CGF.CapturedStmtInfo->isCXXThisExprCaptured()) {
3783       LValue ThisLVal =
3784           CGF.EmitLValueForFieldInitialization(VDLVal, ThisCapture);
3785       llvm::Value *CXXThis = CGF.LoadCXXThis();
3786       CGF.EmitStoreOfScalar(CXXThis, ThisLVal);
3787     }
3788     for (const LambdaCapture &LC : RD->captures()) {
3789       if (LC.getCaptureKind() != LCK_ByRef)
3790         continue;
3791       const VarDecl *VD = LC.getCapturedVar();
3792       if (!CS->capturesVariable(VD))
3793         continue;
3794       auto It = Captures.find(VD);
3795       assert(It != Captures.end() && "Found lambda capture without field.");
3796       LValue VarLVal = CGF.EmitLValueForFieldInitialization(VDLVal, It->second);
3797       Address VDAddr = CGF.GetAddrOfLocalVar(VD);
3798       if (VD->getType().getCanonicalType()->isReferenceType())
3799         VDAddr = CGF.EmitLoadOfReferenceLValue(VDAddr,
3800                                                VD->getType().getCanonicalType())
3801                      .getAddress(CGF);
3802       CGF.EmitStoreOfScalar(VDAddr.getPointer(), VarLVal);
3803     }
3804   }
3805 }
3806 
3807 bool CGOpenMPRuntimeGPU::hasAllocateAttributeForGlobalVar(const VarDecl *VD,
3808                                                             LangAS &AS) {
3809   if (!VD || !VD->hasAttr<OMPAllocateDeclAttr>())
3810     return false;
3811   const auto *A = VD->getAttr<OMPAllocateDeclAttr>();
3812   switch(A->getAllocatorType()) {
3813   case OMPAllocateDeclAttr::OMPNullMemAlloc:
3814   case OMPAllocateDeclAttr::OMPDefaultMemAlloc:
3815   // Not supported, fallback to the default mem space.
3816   case OMPAllocateDeclAttr::OMPThreadMemAlloc:
3817   case OMPAllocateDeclAttr::OMPLargeCapMemAlloc:
3818   case OMPAllocateDeclAttr::OMPCGroupMemAlloc:
3819   case OMPAllocateDeclAttr::OMPHighBWMemAlloc:
3820   case OMPAllocateDeclAttr::OMPLowLatMemAlloc:
3821     AS = LangAS::Default;
3822     return true;
3823   case OMPAllocateDeclAttr::OMPConstMemAlloc:
3824     AS = LangAS::cuda_constant;
3825     return true;
3826   case OMPAllocateDeclAttr::OMPPTeamMemAlloc:
3827     AS = LangAS::cuda_shared;
3828     return true;
3829   case OMPAllocateDeclAttr::OMPUserDefinedMemAlloc:
3830     llvm_unreachable("Expected predefined allocator for the variables with the "
3831                      "static storage.");
3832   }
3833   return false;
3834 }
3835 
3836 // Get current CudaArch and ignore any unknown values
3837 static CudaArch getCudaArch(CodeGenModule &CGM) {
3838   if (!CGM.getTarget().hasFeature("ptx"))
3839     return CudaArch::UNKNOWN;
3840   for (const auto &Feature : CGM.getTarget().getTargetOpts().FeatureMap) {
3841     if (Feature.getValue()) {
3842       CudaArch Arch = StringToCudaArch(Feature.getKey());
3843       if (Arch != CudaArch::UNKNOWN)
3844         return Arch;
3845     }
3846   }
3847   return CudaArch::UNKNOWN;
3848 }
3849 
3850 /// Check to see if target architecture supports unified addressing which is
3851 /// a restriction for OpenMP requires clause "unified_shared_memory".
3852 void CGOpenMPRuntimeGPU::processRequiresDirective(
3853     const OMPRequiresDecl *D) {
3854   for (const OMPClause *Clause : D->clauselists()) {
3855     if (Clause->getClauseKind() == OMPC_unified_shared_memory) {
3856       CudaArch Arch = getCudaArch(CGM);
3857       switch (Arch) {
3858       case CudaArch::SM_20:
3859       case CudaArch::SM_21:
3860       case CudaArch::SM_30:
3861       case CudaArch::SM_32:
3862       case CudaArch::SM_35:
3863       case CudaArch::SM_37:
3864       case CudaArch::SM_50:
3865       case CudaArch::SM_52:
3866       case CudaArch::SM_53: {
3867         SmallString<256> Buffer;
3868         llvm::raw_svector_ostream Out(Buffer);
3869         Out << "Target architecture " << CudaArchToString(Arch)
3870             << " does not support unified addressing";
3871         CGM.Error(Clause->getBeginLoc(), Out.str());
3872         return;
3873       }
3874       case CudaArch::SM_60:
3875       case CudaArch::SM_61:
3876       case CudaArch::SM_62:
3877       case CudaArch::SM_70:
3878       case CudaArch::SM_72:
3879       case CudaArch::SM_75:
3880       case CudaArch::SM_80:
3881       case CudaArch::SM_86:
3882       case CudaArch::GFX600:
3883       case CudaArch::GFX601:
3884       case CudaArch::GFX602:
3885       case CudaArch::GFX700:
3886       case CudaArch::GFX701:
3887       case CudaArch::GFX702:
3888       case CudaArch::GFX703:
3889       case CudaArch::GFX704:
3890       case CudaArch::GFX705:
3891       case CudaArch::GFX801:
3892       case CudaArch::GFX802:
3893       case CudaArch::GFX803:
3894       case CudaArch::GFX805:
3895       case CudaArch::GFX810:
3896       case CudaArch::GFX900:
3897       case CudaArch::GFX902:
3898       case CudaArch::GFX904:
3899       case CudaArch::GFX906:
3900       case CudaArch::GFX908:
3901       case CudaArch::GFX909:
3902       case CudaArch::GFX90a:
3903       case CudaArch::GFX90c:
3904       case CudaArch::GFX1010:
3905       case CudaArch::GFX1011:
3906       case CudaArch::GFX1012:
3907       case CudaArch::GFX1013:
3908       case CudaArch::GFX1030:
3909       case CudaArch::GFX1031:
3910       case CudaArch::GFX1032:
3911       case CudaArch::GFX1033:
3912       case CudaArch::GFX1034:
3913       case CudaArch::GFX1035:
3914       case CudaArch::Generic:
3915       case CudaArch::UNUSED:
3916       case CudaArch::UNKNOWN:
3917         break;
3918       case CudaArch::LAST:
3919         llvm_unreachable("Unexpected Cuda arch.");
3920       }
3921     }
3922   }
3923   CGOpenMPRuntime::processRequiresDirective(D);
3924 }
3925 
3926 void CGOpenMPRuntimeGPU::clear() {
3927 
3928   if (!TeamsReductions.empty()) {
3929     ASTContext &C = CGM.getContext();
3930     RecordDecl *StaticRD = C.buildImplicitRecord(
3931         "_openmp_teams_reduction_type_$_", RecordDecl::TagKind::TTK_Union);
3932     StaticRD->startDefinition();
3933     for (const RecordDecl *TeamReductionRec : TeamsReductions) {
3934       QualType RecTy = C.getRecordType(TeamReductionRec);
3935       auto *Field = FieldDecl::Create(
3936           C, StaticRD, SourceLocation(), SourceLocation(), nullptr, RecTy,
3937           C.getTrivialTypeSourceInfo(RecTy, SourceLocation()),
3938           /*BW=*/nullptr, /*Mutable=*/false,
3939           /*InitStyle=*/ICIS_NoInit);
3940       Field->setAccess(AS_public);
3941       StaticRD->addDecl(Field);
3942     }
3943     StaticRD->completeDefinition();
3944     QualType StaticTy = C.getRecordType(StaticRD);
3945     llvm::Type *LLVMReductionsBufferTy =
3946         CGM.getTypes().ConvertTypeForMem(StaticTy);
3947     // FIXME: nvlink does not handle weak linkage correctly (object with the
3948     // different size are reported as erroneous).
3949     // Restore CommonLinkage as soon as nvlink is fixed.
3950     auto *GV = new llvm::GlobalVariable(
3951         CGM.getModule(), LLVMReductionsBufferTy,
3952         /*isConstant=*/false, llvm::GlobalValue::InternalLinkage,
3953         llvm::Constant::getNullValue(LLVMReductionsBufferTy),
3954         "_openmp_teams_reductions_buffer_$_");
3955     KernelTeamsReductionPtr->setInitializer(
3956         llvm::ConstantExpr::getPointerBitCastOrAddrSpaceCast(GV,
3957                                                              CGM.VoidPtrTy));
3958   }
3959   CGOpenMPRuntime::clear();
3960 }
3961 
3962 llvm::Value *CGOpenMPRuntimeGPU::getGPUNumThreads(CodeGenFunction &CGF) {
3963   CGBuilderTy &Bld = CGF.Builder;
3964   llvm::Module *M = &CGF.CGM.getModule();
3965   const char *LocSize = "__kmpc_get_hardware_num_threads_in_block";
3966   llvm::Function *F = M->getFunction(LocSize);
3967   if (!F) {
3968     F = llvm::Function::Create(
3969         llvm::FunctionType::get(CGF.Int32Ty, llvm::None, false),
3970         llvm::GlobalVariable::ExternalLinkage, LocSize, &CGF.CGM.getModule());
3971   }
3972   return Bld.CreateCall(F, llvm::None, "nvptx_num_threads");
3973 }
3974 
3975 llvm::Value *CGOpenMPRuntimeGPU::getGPUThreadID(CodeGenFunction &CGF) {
3976   ArrayRef<llvm::Value *> Args{};
3977   return CGF.EmitRuntimeCall(
3978       OMPBuilder.getOrCreateRuntimeFunction(
3979           CGM.getModule(), OMPRTL___kmpc_get_hardware_thread_id_in_block),
3980       Args);
3981 }
3982 
3983 llvm::Value *CGOpenMPRuntimeGPU::getGPUWarpSize(CodeGenFunction &CGF) {
3984   ArrayRef<llvm::Value *> Args{};
3985   return CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
3986                                  CGM.getModule(), OMPRTL___kmpc_get_warp_size),
3987                              Args);
3988 }
3989