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