brintos / llvm-project-archived public Read only

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-exception6//7//===----------------------------------------------------------------------===//8//9// This provides a generalized class for OpenMP runtime code generation10// specialized by GPU targets NVPTX, AMDGCN and SPIR-V.11//12//===----------------------------------------------------------------------===//13 14#include "CGOpenMPRuntimeGPU.h"15#include "CGDebugInfo.h"16#include "CodeGenFunction.h"17#include "clang/AST/Attr.h"18#include "clang/AST/DeclOpenMP.h"19#include "clang/AST/OpenMPClause.h"20#include "clang/AST/StmtOpenMP.h"21#include "clang/AST/StmtVisitor.h"22#include "clang/Basic/Cuda.h"23#include "llvm/ADT/SmallPtrSet.h"24#include "llvm/Frontend/OpenMP/OMPDeviceConstants.h"25#include "llvm/Frontend/OpenMP/OMPGridValues.h"26 27using namespace clang;28using namespace CodeGen;29using namespace llvm::omp;30 31namespace {32/// Pre(post)-action for different OpenMP constructs specialized for NVPTX.33class NVPTXActionTy final : public PrePostActionTy {34  llvm::FunctionCallee EnterCallee = nullptr;35  ArrayRef<llvm::Value *> EnterArgs;36  llvm::FunctionCallee ExitCallee = nullptr;37  ArrayRef<llvm::Value *> ExitArgs;38  bool Conditional = false;39  llvm::BasicBlock *ContBlock = nullptr;40 41public:42  NVPTXActionTy(llvm::FunctionCallee EnterCallee,43                ArrayRef<llvm::Value *> EnterArgs,44                llvm::FunctionCallee ExitCallee,45                ArrayRef<llvm::Value *> ExitArgs, bool Conditional = false)46      : EnterCallee(EnterCallee), EnterArgs(EnterArgs), ExitCallee(ExitCallee),47        ExitArgs(ExitArgs), Conditional(Conditional) {}48  void Enter(CodeGenFunction &CGF) override {49    llvm::Value *EnterRes = CGF.EmitRuntimeCall(EnterCallee, EnterArgs);50    if (Conditional) {51      llvm::Value *CallBool = CGF.Builder.CreateIsNotNull(EnterRes);52      auto *ThenBlock = CGF.createBasicBlock("omp_if.then");53      ContBlock = CGF.createBasicBlock("omp_if.end");54      // Generate the branch (If-stmt)55      CGF.Builder.CreateCondBr(CallBool, ThenBlock, ContBlock);56      CGF.EmitBlock(ThenBlock);57    }58  }59  void Done(CodeGenFunction &CGF) {60    // Emit the rest of blocks/branches61    CGF.EmitBranch(ContBlock);62    CGF.EmitBlock(ContBlock, true);63  }64  void Exit(CodeGenFunction &CGF) override {65    CGF.EmitRuntimeCall(ExitCallee, ExitArgs);66  }67};68 69/// A class to track the execution mode when codegening directives within70/// a target region. The appropriate mode (SPMD|NON-SPMD) is set on entry71/// to the target region and used by containing directives such as 'parallel'72/// to emit optimized code.73class ExecutionRuntimeModesRAII {74private:75  CGOpenMPRuntimeGPU::ExecutionMode SavedExecMode =76      CGOpenMPRuntimeGPU::EM_Unknown;77  CGOpenMPRuntimeGPU::ExecutionMode &ExecMode;78 79public:80  ExecutionRuntimeModesRAII(CGOpenMPRuntimeGPU::ExecutionMode &ExecMode,81                            CGOpenMPRuntimeGPU::ExecutionMode EntryMode)82      : ExecMode(ExecMode) {83    SavedExecMode = ExecMode;84    ExecMode = EntryMode;85  }86  ~ExecutionRuntimeModesRAII() { ExecMode = SavedExecMode; }87};88 89static const ValueDecl *getPrivateItem(const Expr *RefExpr) {90  RefExpr = RefExpr->IgnoreParens();91  if (const auto *ASE = dyn_cast<ArraySubscriptExpr>(RefExpr)) {92    const Expr *Base = ASE->getBase()->IgnoreParenImpCasts();93    while (const auto *TempASE = dyn_cast<ArraySubscriptExpr>(Base))94      Base = TempASE->getBase()->IgnoreParenImpCasts();95    RefExpr = Base;96  } else if (auto *OASE = dyn_cast<ArraySectionExpr>(RefExpr)) {97    const Expr *Base = OASE->getBase()->IgnoreParenImpCasts();98    while (const auto *TempOASE = dyn_cast<ArraySectionExpr>(Base))99      Base = TempOASE->getBase()->IgnoreParenImpCasts();100    while (const auto *TempASE = dyn_cast<ArraySubscriptExpr>(Base))101      Base = TempASE->getBase()->IgnoreParenImpCasts();102    RefExpr = Base;103  }104  RefExpr = RefExpr->IgnoreParenImpCasts();105  if (const auto *DE = dyn_cast<DeclRefExpr>(RefExpr))106    return cast<ValueDecl>(DE->getDecl()->getCanonicalDecl());107  const auto *ME = cast<MemberExpr>(RefExpr);108  return cast<ValueDecl>(ME->getMemberDecl()->getCanonicalDecl());109}110 111static RecordDecl *buildRecordForGlobalizedVars(112    ASTContext &C, ArrayRef<const ValueDecl *> EscapedDecls,113    ArrayRef<const ValueDecl *> EscapedDeclsForTeams,114    llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>115        &MappedDeclsFields,116    int BufSize) {117  using VarsDataTy = std::pair<CharUnits /*Align*/, const ValueDecl *>;118  if (EscapedDecls.empty() && EscapedDeclsForTeams.empty())119    return nullptr;120  SmallVector<VarsDataTy, 4> GlobalizedVars;121  for (const ValueDecl *D : EscapedDecls)122    GlobalizedVars.emplace_back(C.getDeclAlign(D), D);123  for (const ValueDecl *D : EscapedDeclsForTeams)124    GlobalizedVars.emplace_back(C.getDeclAlign(D), D);125 126  // Build struct _globalized_locals_ty {127  //         /*  globalized vars  */[WarSize] align (decl_align)128  //         /*  globalized vars  */ for EscapedDeclsForTeams129  //       };130  RecordDecl *GlobalizedRD = C.buildImplicitRecord("_globalized_locals_ty");131  GlobalizedRD->startDefinition();132  llvm::SmallPtrSet<const ValueDecl *, 16> SingleEscaped(llvm::from_range,133                                                         EscapedDeclsForTeams);134  for (const auto &Pair : GlobalizedVars) {135    const ValueDecl *VD = Pair.second;136    QualType Type = VD->getType();137    if (Type->isLValueReferenceType())138      Type = C.getPointerType(Type.getNonReferenceType());139    else140      Type = Type.getNonReferenceType();141    SourceLocation Loc = VD->getLocation();142    FieldDecl *Field;143    if (SingleEscaped.count(VD)) {144      Field = FieldDecl::Create(145          C, GlobalizedRD, Loc, Loc, VD->getIdentifier(), Type,146          C.getTrivialTypeSourceInfo(Type, SourceLocation()),147          /*BW=*/nullptr, /*Mutable=*/false,148          /*InitStyle=*/ICIS_NoInit);149      Field->setAccess(AS_public);150      if (VD->hasAttrs()) {151        for (specific_attr_iterator<AlignedAttr> I(VD->getAttrs().begin()),152             E(VD->getAttrs().end());153             I != E; ++I)154          Field->addAttr(*I);155      }156    } else {157      if (BufSize > 1) {158        llvm::APInt ArraySize(32, BufSize);159        Type = C.getConstantArrayType(Type, ArraySize, nullptr,160                                      ArraySizeModifier::Normal, 0);161      }162      Field = FieldDecl::Create(163          C, GlobalizedRD, Loc, Loc, VD->getIdentifier(), Type,164          C.getTrivialTypeSourceInfo(Type, SourceLocation()),165          /*BW=*/nullptr, /*Mutable=*/false,166          /*InitStyle=*/ICIS_NoInit);167      Field->setAccess(AS_public);168      llvm::APInt Align(32, Pair.first.getQuantity());169      Field->addAttr(AlignedAttr::CreateImplicit(170          C, /*IsAlignmentExpr=*/true,171          IntegerLiteral::Create(C, Align,172                                 C.getIntTypeForBitwidth(32, /*Signed=*/0),173                                 SourceLocation()),174          {}, AlignedAttr::GNU_aligned));175    }176    GlobalizedRD->addDecl(Field);177    MappedDeclsFields.try_emplace(VD, Field);178  }179  GlobalizedRD->completeDefinition();180  return GlobalizedRD;181}182 183/// Get the list of variables that can escape their declaration context.184class CheckVarsEscapingDeclContext final185    : public ConstStmtVisitor<CheckVarsEscapingDeclContext> {186  CodeGenFunction &CGF;187  llvm::SetVector<const ValueDecl *> EscapedDecls;188  llvm::SetVector<const ValueDecl *> EscapedVariableLengthDecls;189  llvm::SetVector<const ValueDecl *> DelayedVariableLengthDecls;190  llvm::SmallPtrSet<const Decl *, 4> EscapedParameters;191  RecordDecl *GlobalizedRD = nullptr;192  llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *> MappedDeclsFields;193  bool AllEscaped = false;194  bool IsForCombinedParallelRegion = false;195 196  void markAsEscaped(const ValueDecl *VD) {197    // Do not globalize declare target variables.198    if (!isa<VarDecl>(VD) ||199        OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD))200      return;201    VD = cast<ValueDecl>(VD->getCanonicalDecl());202    // Use user-specified allocation.203    if (VD->hasAttrs() && VD->hasAttr<OMPAllocateDeclAttr>())204      return;205    // Variables captured by value must be globalized.206    bool IsCaptured = false;207    if (auto *CSI = CGF.CapturedStmtInfo) {208      if (const FieldDecl *FD = CSI->lookup(cast<VarDecl>(VD))) {209        // Check if need to capture the variable that was already captured by210        // value in the outer region.211        IsCaptured = true;212        if (!IsForCombinedParallelRegion) {213          if (!FD->hasAttrs())214            return;215          const auto *Attr = FD->getAttr<OMPCaptureKindAttr>();216          if (!Attr)217            return;218          if (((Attr->getCaptureKind() != OMPC_map) &&219               !isOpenMPPrivate(Attr->getCaptureKind())) ||220              ((Attr->getCaptureKind() == OMPC_map) &&221               !FD->getType()->isAnyPointerType()))222            return;223        }224        if (!FD->getType()->isReferenceType()) {225          assert(!VD->getType()->isVariablyModifiedType() &&226                 "Parameter captured by value with variably modified type");227          EscapedParameters.insert(VD);228        } else if (!IsForCombinedParallelRegion) {229          return;230        }231      }232    }233    if ((!CGF.CapturedStmtInfo ||234         (IsForCombinedParallelRegion && CGF.CapturedStmtInfo)) &&235        VD->getType()->isReferenceType())236      // Do not globalize variables with reference type.237      return;238    if (VD->getType()->isVariablyModifiedType()) {239      // If not captured at the target region level then mark the escaped240      // variable as delayed.241      if (IsCaptured)242        EscapedVariableLengthDecls.insert(VD);243      else244        DelayedVariableLengthDecls.insert(VD);245    } else246      EscapedDecls.insert(VD);247  }248 249  void VisitValueDecl(const ValueDecl *VD) {250    if (VD->getType()->isLValueReferenceType())251      markAsEscaped(VD);252    if (const auto *VarD = dyn_cast<VarDecl>(VD)) {253      if (!isa<ParmVarDecl>(VarD) && VarD->hasInit()) {254        const bool SavedAllEscaped = AllEscaped;255        AllEscaped = VD->getType()->isLValueReferenceType();256        Visit(VarD->getInit());257        AllEscaped = SavedAllEscaped;258      }259    }260  }261  void VisitOpenMPCapturedStmt(const CapturedStmt *S,262                               ArrayRef<OMPClause *> Clauses,263                               bool IsCombinedParallelRegion) {264    if (!S)265      return;266    for (const CapturedStmt::Capture &C : S->captures()) {267      if (C.capturesVariable() && !C.capturesVariableByCopy()) {268        const ValueDecl *VD = C.getCapturedVar();269        bool SavedIsForCombinedParallelRegion = IsForCombinedParallelRegion;270        if (IsCombinedParallelRegion) {271          // Check if the variable is privatized in the combined construct and272          // those private copies must be shared in the inner parallel273          // directive.274          IsForCombinedParallelRegion = false;275          for (const OMPClause *C : Clauses) {276            if (!isOpenMPPrivate(C->getClauseKind()) ||277                C->getClauseKind() == OMPC_reduction ||278                C->getClauseKind() == OMPC_linear ||279                C->getClauseKind() == OMPC_private)280              continue;281            ArrayRef<const Expr *> Vars;282            if (const auto *PC = dyn_cast<OMPFirstprivateClause>(C))283              Vars = PC->getVarRefs();284            else if (const auto *PC = dyn_cast<OMPLastprivateClause>(C))285              Vars = PC->getVarRefs();286            else287              llvm_unreachable("Unexpected clause.");288            for (const auto *E : Vars) {289              const Decl *D =290                  cast<DeclRefExpr>(E)->getDecl()->getCanonicalDecl();291              if (D == VD->getCanonicalDecl()) {292                IsForCombinedParallelRegion = true;293                break;294              }295            }296            if (IsForCombinedParallelRegion)297              break;298          }299        }300        markAsEscaped(VD);301        if (isa<OMPCapturedExprDecl>(VD))302          VisitValueDecl(VD);303        IsForCombinedParallelRegion = SavedIsForCombinedParallelRegion;304      }305    }306  }307 308  void buildRecordForGlobalizedVars(bool IsInTTDRegion) {309    assert(!GlobalizedRD &&310           "Record for globalized variables is built already.");311    ArrayRef<const ValueDecl *> EscapedDeclsForParallel, EscapedDeclsForTeams;312    unsigned WarpSize = CGF.getTarget().getGridValue().GV_Warp_Size;313    if (IsInTTDRegion)314      EscapedDeclsForTeams = EscapedDecls.getArrayRef();315    else316      EscapedDeclsForParallel = EscapedDecls.getArrayRef();317    GlobalizedRD = ::buildRecordForGlobalizedVars(318        CGF.getContext(), EscapedDeclsForParallel, EscapedDeclsForTeams,319        MappedDeclsFields, WarpSize);320  }321 322public:323  CheckVarsEscapingDeclContext(CodeGenFunction &CGF,324                               ArrayRef<const ValueDecl *> TeamsReductions)325      : CGF(CGF), EscapedDecls(llvm::from_range, TeamsReductions) {}326  ~CheckVarsEscapingDeclContext() = default;327  void VisitDeclStmt(const DeclStmt *S) {328    if (!S)329      return;330    for (const Decl *D : S->decls())331      if (const auto *VD = dyn_cast_or_null<ValueDecl>(D))332        VisitValueDecl(VD);333  }334  void VisitOMPExecutableDirective(const OMPExecutableDirective *D) {335    if (!D)336      return;337    if (!D->hasAssociatedStmt())338      return;339    if (const auto *S =340            dyn_cast_or_null<CapturedStmt>(D->getAssociatedStmt())) {341      // Do not analyze directives that do not actually require capturing,342      // like `omp for` or `omp simd` directives.343      llvm::SmallVector<OpenMPDirectiveKind, 4> CaptureRegions;344      getOpenMPCaptureRegions(CaptureRegions, D->getDirectiveKind());345      if (CaptureRegions.size() == 1 && CaptureRegions.back() == OMPD_unknown) {346        VisitStmt(S->getCapturedStmt());347        return;348      }349      VisitOpenMPCapturedStmt(350          S, D->clauses(),351          CaptureRegions.back() == OMPD_parallel &&352              isOpenMPDistributeDirective(D->getDirectiveKind()));353    }354  }355  void VisitCapturedStmt(const CapturedStmt *S) {356    if (!S)357      return;358    for (const CapturedStmt::Capture &C : S->captures()) {359      if (C.capturesVariable() && !C.capturesVariableByCopy()) {360        const ValueDecl *VD = C.getCapturedVar();361        markAsEscaped(VD);362        if (isa<OMPCapturedExprDecl>(VD))363          VisitValueDecl(VD);364      }365    }366  }367  void VisitLambdaExpr(const LambdaExpr *E) {368    if (!E)369      return;370    for (const LambdaCapture &C : E->captures()) {371      if (C.capturesVariable()) {372        if (C.getCaptureKind() == LCK_ByRef) {373          const ValueDecl *VD = C.getCapturedVar();374          markAsEscaped(VD);375          if (E->isInitCapture(&C) || isa<OMPCapturedExprDecl>(VD))376            VisitValueDecl(VD);377        }378      }379    }380  }381  void VisitBlockExpr(const BlockExpr *E) {382    if (!E)383      return;384    for (const BlockDecl::Capture &C : E->getBlockDecl()->captures()) {385      if (C.isByRef()) {386        const VarDecl *VD = C.getVariable();387        markAsEscaped(VD);388        if (isa<OMPCapturedExprDecl>(VD) || VD->isInitCapture())389          VisitValueDecl(VD);390      }391    }392  }393  void VisitCallExpr(const CallExpr *E) {394    if (!E)395      return;396    for (const Expr *Arg : E->arguments()) {397      if (!Arg)398        continue;399      if (Arg->isLValue()) {400        const bool SavedAllEscaped = AllEscaped;401        AllEscaped = true;402        Visit(Arg);403        AllEscaped = SavedAllEscaped;404      } else {405        Visit(Arg);406      }407    }408    Visit(E->getCallee());409  }410  void VisitDeclRefExpr(const DeclRefExpr *E) {411    if (!E)412      return;413    const ValueDecl *VD = E->getDecl();414    if (AllEscaped)415      markAsEscaped(VD);416    if (isa<OMPCapturedExprDecl>(VD))417      VisitValueDecl(VD);418    else if (VD->isInitCapture())419      VisitValueDecl(VD);420  }421  void VisitUnaryOperator(const UnaryOperator *E) {422    if (!E)423      return;424    if (E->getOpcode() == UO_AddrOf) {425      const bool SavedAllEscaped = AllEscaped;426      AllEscaped = true;427      Visit(E->getSubExpr());428      AllEscaped = SavedAllEscaped;429    } else {430      Visit(E->getSubExpr());431    }432  }433  void VisitImplicitCastExpr(const ImplicitCastExpr *E) {434    if (!E)435      return;436    if (E->getCastKind() == CK_ArrayToPointerDecay) {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 VisitExpr(const Expr *E) {446    if (!E)447      return;448    bool SavedAllEscaped = AllEscaped;449    if (!E->isLValue())450      AllEscaped = false;451    for (const Stmt *Child : E->children())452      if (Child)453        Visit(Child);454    AllEscaped = SavedAllEscaped;455  }456  void VisitStmt(const Stmt *S) {457    if (!S)458      return;459    for (const Stmt *Child : S->children())460      if (Child)461        Visit(Child);462  }463 464  /// Returns the record that handles all the escaped local variables and used465  /// instead of their original storage.466  const RecordDecl *getGlobalizedRecord(bool IsInTTDRegion) {467    if (!GlobalizedRD)468      buildRecordForGlobalizedVars(IsInTTDRegion);469    return GlobalizedRD;470  }471 472  /// Returns the field in the globalized record for the escaped variable.473  const FieldDecl *getFieldForGlobalizedVar(const ValueDecl *VD) const {474    assert(GlobalizedRD &&475           "Record for globalized variables must be generated already.");476    return MappedDeclsFields.lookup(VD);477  }478 479  /// Returns the list of the escaped local variables/parameters.480  ArrayRef<const ValueDecl *> getEscapedDecls() const {481    return EscapedDecls.getArrayRef();482  }483 484  /// Checks if the escaped local variable is actually a parameter passed by485  /// value.486  const llvm::SmallPtrSetImpl<const Decl *> &getEscapedParameters() const {487    return EscapedParameters;488  }489 490  /// Returns the list of the escaped variables with the variably modified491  /// types.492  ArrayRef<const ValueDecl *> getEscapedVariableLengthDecls() const {493    return EscapedVariableLengthDecls.getArrayRef();494  }495 496  /// Returns the list of the delayed variables with the variably modified497  /// types.498  ArrayRef<const ValueDecl *> getDelayedVariableLengthDecls() const {499    return DelayedVariableLengthDecls.getArrayRef();500  }501};502} // anonymous namespace503 504CGOpenMPRuntimeGPU::ExecutionMode505CGOpenMPRuntimeGPU::getExecutionMode() const {506  return CurrentExecutionMode;507}508 509CGOpenMPRuntimeGPU::DataSharingMode510CGOpenMPRuntimeGPU::getDataSharingMode() const {511  return CurrentDataSharingMode;512}513 514/// Check for inner (nested) SPMD construct, if any515static bool hasNestedSPMDDirective(ASTContext &Ctx,516                                   const OMPExecutableDirective &D) {517  const auto *CS = D.getInnermostCapturedStmt();518  const auto *Body =519      CS->getCapturedStmt()->IgnoreContainers(/*IgnoreCaptured=*/true);520  const Stmt *ChildStmt = CGOpenMPRuntime::getSingleCompoundChild(Ctx, Body);521 522  if (const auto *NestedDir =523          dyn_cast_or_null<OMPExecutableDirective>(ChildStmt)) {524    OpenMPDirectiveKind DKind = NestedDir->getDirectiveKind();525    switch (D.getDirectiveKind()) {526    case OMPD_target:527      if (isOpenMPParallelDirective(DKind))528        return true;529      if (DKind == OMPD_teams) {530        Body = NestedDir->getInnermostCapturedStmt()->IgnoreContainers(531            /*IgnoreCaptured=*/true);532        if (!Body)533          return false;534        ChildStmt = CGOpenMPRuntime::getSingleCompoundChild(Ctx, Body);535        if (const auto *NND =536                dyn_cast_or_null<OMPExecutableDirective>(ChildStmt)) {537          DKind = NND->getDirectiveKind();538          if (isOpenMPParallelDirective(DKind))539            return true;540        }541      }542      return false;543    case OMPD_target_teams:544      return isOpenMPParallelDirective(DKind);545    case OMPD_target_simd:546    case OMPD_target_parallel:547    case OMPD_target_parallel_for:548    case OMPD_target_parallel_for_simd:549    case OMPD_target_teams_distribute:550    case OMPD_target_teams_distribute_simd:551    case OMPD_target_teams_distribute_parallel_for:552    case OMPD_target_teams_distribute_parallel_for_simd:553    case OMPD_parallel:554    case OMPD_for:555    case OMPD_parallel_for:556    case OMPD_parallel_master:557    case OMPD_parallel_sections:558    case OMPD_for_simd:559    case OMPD_parallel_for_simd:560    case OMPD_cancel:561    case OMPD_cancellation_point:562    case OMPD_ordered:563    case OMPD_threadprivate:564    case OMPD_allocate:565    case OMPD_task:566    case OMPD_simd:567    case OMPD_sections:568    case OMPD_section:569    case OMPD_single:570    case OMPD_master:571    case OMPD_critical:572    case OMPD_taskyield:573    case OMPD_barrier:574    case OMPD_taskwait:575    case OMPD_taskgroup:576    case OMPD_atomic:577    case OMPD_flush:578    case OMPD_depobj:579    case OMPD_scan:580    case OMPD_teams:581    case OMPD_target_data:582    case OMPD_target_exit_data:583    case OMPD_target_enter_data:584    case OMPD_distribute:585    case OMPD_distribute_simd:586    case OMPD_distribute_parallel_for:587    case OMPD_distribute_parallel_for_simd:588    case OMPD_teams_distribute:589    case OMPD_teams_distribute_simd:590    case OMPD_teams_distribute_parallel_for:591    case OMPD_teams_distribute_parallel_for_simd:592    case OMPD_target_update:593    case OMPD_declare_simd:594    case OMPD_declare_variant:595    case OMPD_begin_declare_variant:596    case OMPD_end_declare_variant:597    case OMPD_declare_target:598    case OMPD_end_declare_target:599    case OMPD_declare_reduction:600    case OMPD_declare_mapper:601    case OMPD_taskloop:602    case OMPD_taskloop_simd:603    case OMPD_master_taskloop:604    case OMPD_master_taskloop_simd:605    case OMPD_parallel_master_taskloop:606    case OMPD_parallel_master_taskloop_simd:607    case OMPD_requires:608    case OMPD_unknown:609    default:610      llvm_unreachable("Unexpected directive.");611    }612  }613 614  return false;615}616 617static bool supportsSPMDExecutionMode(ASTContext &Ctx,618                                      const OMPExecutableDirective &D) {619  OpenMPDirectiveKind DirectiveKind = D.getDirectiveKind();620  switch (DirectiveKind) {621  case OMPD_target:622  case OMPD_target_teams:623    return hasNestedSPMDDirective(Ctx, D);624  case OMPD_target_parallel_loop:625  case OMPD_target_parallel:626  case OMPD_target_parallel_for:627  case OMPD_target_parallel_for_simd:628  case OMPD_target_teams_distribute_parallel_for:629  case OMPD_target_teams_distribute_parallel_for_simd:630  case OMPD_target_simd:631  case OMPD_target_teams_distribute_simd:632    return true;633  case OMPD_target_teams_distribute:634    return false;635  case OMPD_target_teams_loop:636    // Whether this is true or not depends on how the directive will637    // eventually be emitted.638    if (auto *TTLD = dyn_cast<OMPTargetTeamsGenericLoopDirective>(&D))639      return TTLD->canBeParallelFor();640    return false;641  case OMPD_parallel:642  case OMPD_for:643  case OMPD_parallel_for:644  case OMPD_parallel_master:645  case OMPD_parallel_sections:646  case OMPD_for_simd:647  case OMPD_parallel_for_simd:648  case OMPD_cancel:649  case OMPD_cancellation_point:650  case OMPD_ordered:651  case OMPD_threadprivate:652  case OMPD_allocate:653  case OMPD_task:654  case OMPD_simd:655  case OMPD_sections:656  case OMPD_section:657  case OMPD_single:658  case OMPD_master:659  case OMPD_critical:660  case OMPD_taskyield:661  case OMPD_barrier:662  case OMPD_taskwait:663  case OMPD_taskgroup:664  case OMPD_atomic:665  case OMPD_flush:666  case OMPD_depobj:667  case OMPD_scan:668  case OMPD_teams:669  case OMPD_target_data:670  case OMPD_target_exit_data:671  case OMPD_target_enter_data:672  case OMPD_distribute:673  case OMPD_distribute_simd:674  case OMPD_distribute_parallel_for:675  case OMPD_distribute_parallel_for_simd:676  case OMPD_teams_distribute:677  case OMPD_teams_distribute_simd:678  case OMPD_teams_distribute_parallel_for:679  case OMPD_teams_distribute_parallel_for_simd:680  case OMPD_target_update:681  case OMPD_declare_simd:682  case OMPD_declare_variant:683  case OMPD_begin_declare_variant:684  case OMPD_end_declare_variant:685  case OMPD_declare_target:686  case OMPD_end_declare_target:687  case OMPD_declare_reduction:688  case OMPD_declare_mapper:689  case OMPD_taskloop:690  case OMPD_taskloop_simd:691  case OMPD_master_taskloop:692  case OMPD_master_taskloop_simd:693  case OMPD_parallel_master_taskloop:694  case OMPD_parallel_master_taskloop_simd:695  case OMPD_requires:696  case OMPD_unknown:697  default:698    break;699  }700  llvm_unreachable(701      "Unknown programming model for OpenMP directive on NVPTX target.");702}703 704void CGOpenMPRuntimeGPU::emitNonSPMDKernel(const OMPExecutableDirective &D,705                                             StringRef ParentName,706                                             llvm::Function *&OutlinedFn,707                                             llvm::Constant *&OutlinedFnID,708                                             bool IsOffloadEntry,709                                             const RegionCodeGenTy &CodeGen) {710  ExecutionRuntimeModesRAII ModeRAII(CurrentExecutionMode, EM_NonSPMD);711  EntryFunctionState EST;712  WrapperFunctionsMap.clear();713 714  [[maybe_unused]] bool IsBareKernel = D.getSingleClause<OMPXBareClause>();715  assert(!IsBareKernel && "bare kernel should not be at generic mode");716 717  // Emit target region as a standalone region.718  class NVPTXPrePostActionTy : public PrePostActionTy {719    CGOpenMPRuntimeGPU::EntryFunctionState &EST;720    const OMPExecutableDirective &D;721 722  public:723    NVPTXPrePostActionTy(CGOpenMPRuntimeGPU::EntryFunctionState &EST,724                         const OMPExecutableDirective &D)725        : EST(EST), D(D) {}726    void Enter(CodeGenFunction &CGF) override {727      auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());728      RT.emitKernelInit(D, CGF, EST, /* IsSPMD */ false);729      // Skip target region initialization.730      RT.setLocThreadIdInsertPt(CGF, /*AtCurrentPoint=*/true);731    }732    void Exit(CodeGenFunction &CGF) override {733      auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());734      RT.clearLocThreadIdInsertPt(CGF);735      RT.emitKernelDeinit(CGF, EST, /* IsSPMD */ false);736    }737  } Action(EST, D);738  CodeGen.setAction(Action);739  IsInTTDRegion = true;740  emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,741                                   IsOffloadEntry, CodeGen);742  IsInTTDRegion = false;743}744 745void CGOpenMPRuntimeGPU::emitKernelInit(const OMPExecutableDirective &D,746                                        CodeGenFunction &CGF,747                                        EntryFunctionState &EST, bool IsSPMD) {748  llvm::OpenMPIRBuilder::TargetKernelDefaultAttrs Attrs;749  Attrs.ExecFlags =750      IsSPMD ? llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_SPMD751             : llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_GENERIC;752  computeMinAndMaxThreadsAndTeams(D, CGF, Attrs);753 754  CGBuilderTy &Bld = CGF.Builder;755  Bld.restoreIP(OMPBuilder.createTargetInit(Bld, Attrs));756  if (!IsSPMD)757    emitGenericVarsProlog(CGF, EST.Loc);758}759 760void CGOpenMPRuntimeGPU::emitKernelDeinit(CodeGenFunction &CGF,761                                          EntryFunctionState &EST,762                                          bool IsSPMD) {763  if (!IsSPMD)764    emitGenericVarsEpilog(CGF);765 766  // This is temporary until we remove the fixed sized buffer.767  ASTContext &C = CGM.getContext();768  RecordDecl *StaticRD = C.buildImplicitRecord(769      "_openmp_teams_reduction_type_$_", RecordDecl::TagKind::Union);770  StaticRD->startDefinition();771  for (const RecordDecl *TeamReductionRec : TeamsReductions) {772    CanQualType RecTy = C.getCanonicalTagType(TeamReductionRec);773    auto *Field = FieldDecl::Create(774        C, StaticRD, SourceLocation(), SourceLocation(), nullptr, RecTy,775        C.getTrivialTypeSourceInfo(RecTy, SourceLocation()),776        /*BW=*/nullptr, /*Mutable=*/false,777        /*InitStyle=*/ICIS_NoInit);778    Field->setAccess(AS_public);779    StaticRD->addDecl(Field);780  }781  StaticRD->completeDefinition();782  CanQualType StaticTy = C.getCanonicalTagType(StaticRD);783  llvm::Type *LLVMReductionsBufferTy =784      CGM.getTypes().ConvertTypeForMem(StaticTy);785  const auto &DL = CGM.getModule().getDataLayout();786  uint64_t ReductionDataSize =787      TeamsReductions.empty()788          ? 0789          : DL.getTypeAllocSize(LLVMReductionsBufferTy).getFixedValue();790  CGBuilderTy &Bld = CGF.Builder;791  OMPBuilder.createTargetDeinit(Bld, ReductionDataSize,792                                C.getLangOpts().OpenMPCUDAReductionBufNum);793  TeamsReductions.clear();794}795 796void CGOpenMPRuntimeGPU::emitSPMDKernel(const OMPExecutableDirective &D,797                                          StringRef ParentName,798                                          llvm::Function *&OutlinedFn,799                                          llvm::Constant *&OutlinedFnID,800                                          bool IsOffloadEntry,801                                          const RegionCodeGenTy &CodeGen) {802  ExecutionRuntimeModesRAII ModeRAII(CurrentExecutionMode, EM_SPMD);803  EntryFunctionState EST;804 805  bool IsBareKernel = D.getSingleClause<OMPXBareClause>();806 807  // Emit target region as a standalone region.808  class NVPTXPrePostActionTy : public PrePostActionTy {809    CGOpenMPRuntimeGPU &RT;810    CGOpenMPRuntimeGPU::EntryFunctionState &EST;811    bool IsBareKernel;812    DataSharingMode Mode;813    const OMPExecutableDirective &D;814 815  public:816    NVPTXPrePostActionTy(CGOpenMPRuntimeGPU &RT,817                         CGOpenMPRuntimeGPU::EntryFunctionState &EST,818                         bool IsBareKernel, const OMPExecutableDirective &D)819        : RT(RT), EST(EST), IsBareKernel(IsBareKernel),820          Mode(RT.CurrentDataSharingMode), D(D) {}821    void Enter(CodeGenFunction &CGF) override {822      if (IsBareKernel) {823        RT.CurrentDataSharingMode = DataSharingMode::DS_CUDA;824        return;825      }826      RT.emitKernelInit(D, CGF, EST, /* IsSPMD */ true);827      // Skip target region initialization.828      RT.setLocThreadIdInsertPt(CGF, /*AtCurrentPoint=*/true);829    }830    void Exit(CodeGenFunction &CGF) override {831      if (IsBareKernel) {832        RT.CurrentDataSharingMode = Mode;833        return;834      }835      RT.clearLocThreadIdInsertPt(CGF);836      RT.emitKernelDeinit(CGF, EST, /* IsSPMD */ true);837    }838  } Action(*this, EST, IsBareKernel, D);839  CodeGen.setAction(Action);840  IsInTTDRegion = true;841  emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,842                                   IsOffloadEntry, CodeGen);843  IsInTTDRegion = false;844}845 846void CGOpenMPRuntimeGPU::emitTargetOutlinedFunction(847    const OMPExecutableDirective &D, StringRef ParentName,848    llvm::Function *&OutlinedFn, llvm::Constant *&OutlinedFnID,849    bool IsOffloadEntry, const RegionCodeGenTy &CodeGen) {850  if (!IsOffloadEntry) // Nothing to do.851    return;852 853  assert(!ParentName.empty() && "Invalid target region parent name!");854 855  bool Mode = supportsSPMDExecutionMode(CGM.getContext(), D);856  bool IsBareKernel = D.getSingleClause<OMPXBareClause>();857  if (Mode || IsBareKernel)858    emitSPMDKernel(D, ParentName, OutlinedFn, OutlinedFnID, IsOffloadEntry,859                   CodeGen);860  else861    emitNonSPMDKernel(D, ParentName, OutlinedFn, OutlinedFnID, IsOffloadEntry,862                      CodeGen);863}864 865CGOpenMPRuntimeGPU::CGOpenMPRuntimeGPU(CodeGenModule &CGM)866    : CGOpenMPRuntime(CGM) {867  llvm::OpenMPIRBuilderConfig Config(868      CGM.getLangOpts().OpenMPIsTargetDevice, isGPU(),869      CGM.getLangOpts().OpenMPOffloadMandatory,870      /*HasRequiresReverseOffload*/ false, /*HasRequiresUnifiedAddress*/ false,871      hasRequiresUnifiedSharedMemory(), /*HasRequiresDynamicAllocators*/ false);872  Config.setDefaultTargetAS(873      CGM.getContext().getTargetInfo().getTargetAddressSpace(LangAS::Default));874  Config.setRuntimeCC(CGM.getRuntimeCC());875 876  OMPBuilder.setConfig(Config);877 878  if (!CGM.getLangOpts().OpenMPIsTargetDevice)879    llvm_unreachable("OpenMP can only handle device code.");880 881  if (CGM.getLangOpts().OpenMPCUDAMode)882    CurrentDataSharingMode = CGOpenMPRuntimeGPU::DS_CUDA;883 884  llvm::OpenMPIRBuilder &OMPBuilder = getOMPBuilder();885  if (CGM.getLangOpts().NoGPULib || CGM.getLangOpts().OMPHostIRFile.empty())886    return;887 888  OMPBuilder.createGlobalFlag(CGM.getLangOpts().OpenMPTargetDebug,889                              "__omp_rtl_debug_kind");890  OMPBuilder.createGlobalFlag(CGM.getLangOpts().OpenMPTeamSubscription,891                              "__omp_rtl_assume_teams_oversubscription");892  OMPBuilder.createGlobalFlag(CGM.getLangOpts().OpenMPThreadSubscription,893                              "__omp_rtl_assume_threads_oversubscription");894  OMPBuilder.createGlobalFlag(CGM.getLangOpts().OpenMPNoThreadState,895                              "__omp_rtl_assume_no_thread_state");896  OMPBuilder.createGlobalFlag(CGM.getLangOpts().OpenMPNoNestedParallelism,897                              "__omp_rtl_assume_no_nested_parallelism");898}899 900void CGOpenMPRuntimeGPU::emitProcBindClause(CodeGenFunction &CGF,901                                              ProcBindKind ProcBind,902                                              SourceLocation Loc) {903  // Nothing to do.904}905 906llvm::Value *CGOpenMPRuntimeGPU::emitMessageClause(CodeGenFunction &CGF,907                                                   const Expr *Message,908                                                   SourceLocation Loc) {909  CGM.getDiags().Report(Loc, diag::warn_omp_gpu_unsupported_clause)910      << getOpenMPClauseName(OMPC_message);911  return nullptr;912}913 914llvm::Value *915CGOpenMPRuntimeGPU::emitSeverityClause(OpenMPSeverityClauseKind Severity,916                                       SourceLocation Loc) {917  CGM.getDiags().Report(Loc, diag::warn_omp_gpu_unsupported_clause)918      << getOpenMPClauseName(OMPC_severity);919  return nullptr;920}921 922void CGOpenMPRuntimeGPU::emitNumThreadsClause(923    CodeGenFunction &CGF, llvm::Value *NumThreads, SourceLocation Loc,924    OpenMPNumThreadsClauseModifier Modifier, OpenMPSeverityClauseKind Severity,925    SourceLocation SeverityLoc, const Expr *Message,926    SourceLocation MessageLoc) {927  if (Modifier == OMPC_NUMTHREADS_strict) {928    CGM.getDiags().Report(Loc,929                          diag::warn_omp_gpu_unsupported_modifier_for_clause)930        << "strict" << getOpenMPClauseName(OMPC_num_threads);931    return;932  }933 934  // Nothing to do.935}936 937void CGOpenMPRuntimeGPU::emitNumTeamsClause(CodeGenFunction &CGF,938                                              const Expr *NumTeams,939                                              const Expr *ThreadLimit,940                                              SourceLocation Loc) {}941 942llvm::Function *CGOpenMPRuntimeGPU::emitParallelOutlinedFunction(943    CodeGenFunction &CGF, const OMPExecutableDirective &D,944    const VarDecl *ThreadIDVar, OpenMPDirectiveKind InnermostKind,945    const RegionCodeGenTy &CodeGen) {946  // Emit target region as a standalone region.947  bool PrevIsInTTDRegion = IsInTTDRegion;948  IsInTTDRegion = false;949  auto *OutlinedFun =950      cast<llvm::Function>(CGOpenMPRuntime::emitParallelOutlinedFunction(951          CGF, D, ThreadIDVar, InnermostKind, CodeGen));952  IsInTTDRegion = PrevIsInTTDRegion;953  if (getExecutionMode() != CGOpenMPRuntimeGPU::EM_SPMD) {954    llvm::Function *WrapperFun =955        createParallelDataSharingWrapper(OutlinedFun, D);956    WrapperFunctionsMap[OutlinedFun] = WrapperFun;957  }958 959  return OutlinedFun;960}961 962/// Get list of lastprivate variables from the teams distribute ... or963/// teams {distribute ...} directives.964static void965getDistributeLastprivateVars(ASTContext &Ctx, const OMPExecutableDirective &D,966                             llvm::SmallVectorImpl<const ValueDecl *> &Vars) {967  assert(isOpenMPTeamsDirective(D.getDirectiveKind()) &&968         "expected teams directive.");969  const OMPExecutableDirective *Dir = &D;970  if (!isOpenMPDistributeDirective(D.getDirectiveKind())) {971    if (const Stmt *S = CGOpenMPRuntime::getSingleCompoundChild(972            Ctx,973            D.getInnermostCapturedStmt()->getCapturedStmt()->IgnoreContainers(974                /*IgnoreCaptured=*/true))) {975      Dir = dyn_cast_or_null<OMPExecutableDirective>(S);976      if (Dir && !isOpenMPDistributeDirective(Dir->getDirectiveKind()))977        Dir = nullptr;978    }979  }980  if (!Dir)981    return;982  for (const auto *C : Dir->getClausesOfKind<OMPLastprivateClause>()) {983    for (const Expr *E : C->getVarRefs())984      Vars.push_back(getPrivateItem(E));985  }986}987 988/// Get list of reduction variables from the teams ... directives.989static void990getTeamsReductionVars(ASTContext &Ctx, const OMPExecutableDirective &D,991                      llvm::SmallVectorImpl<const ValueDecl *> &Vars) {992  assert(isOpenMPTeamsDirective(D.getDirectiveKind()) &&993         "expected teams directive.");994  for (const auto *C : D.getClausesOfKind<OMPReductionClause>()) {995    for (const Expr *E : C->privates())996      Vars.push_back(getPrivateItem(E));997  }998}999 1000llvm::Function *CGOpenMPRuntimeGPU::emitTeamsOutlinedFunction(1001    CodeGenFunction &CGF, const OMPExecutableDirective &D,1002    const VarDecl *ThreadIDVar, OpenMPDirectiveKind InnermostKind,1003    const RegionCodeGenTy &CodeGen) {1004  SourceLocation Loc = D.getBeginLoc();1005 1006  const RecordDecl *GlobalizedRD = nullptr;1007  llvm::SmallVector<const ValueDecl *, 4> LastPrivatesReductions;1008  llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *> MappedDeclsFields;1009  unsigned WarpSize = CGM.getTarget().getGridValue().GV_Warp_Size;1010  // Globalize team reductions variable unconditionally in all modes.1011  if (getExecutionMode() != CGOpenMPRuntimeGPU::EM_SPMD)1012    getTeamsReductionVars(CGM.getContext(), D, LastPrivatesReductions);1013  if (getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD) {1014    getDistributeLastprivateVars(CGM.getContext(), D, LastPrivatesReductions);1015    if (!LastPrivatesReductions.empty()) {1016      GlobalizedRD = ::buildRecordForGlobalizedVars(1017          CGM.getContext(), {}, LastPrivatesReductions, MappedDeclsFields,1018          WarpSize);1019    }1020  } else if (!LastPrivatesReductions.empty()) {1021    assert(!TeamAndReductions.first &&1022           "Previous team declaration is not expected.");1023    TeamAndReductions.first = D.getCapturedStmt(OMPD_teams)->getCapturedDecl();1024    std::swap(TeamAndReductions.second, LastPrivatesReductions);1025  }1026 1027  // Emit target region as a standalone region.1028  class NVPTXPrePostActionTy : public PrePostActionTy {1029    SourceLocation &Loc;1030    const RecordDecl *GlobalizedRD;1031    llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>1032        &MappedDeclsFields;1033 1034  public:1035    NVPTXPrePostActionTy(1036        SourceLocation &Loc, const RecordDecl *GlobalizedRD,1037        llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>1038            &MappedDeclsFields)1039        : Loc(Loc), GlobalizedRD(GlobalizedRD),1040          MappedDeclsFields(MappedDeclsFields) {}1041    void Enter(CodeGenFunction &CGF) override {1042      auto &Rt =1043          static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());1044      if (GlobalizedRD) {1045        auto I = Rt.FunctionGlobalizedDecls.try_emplace(CGF.CurFn).first;1046        I->getSecond().MappedParams =1047            std::make_unique<CodeGenFunction::OMPMapVars>();1048        DeclToAddrMapTy &Data = I->getSecond().LocalVarData;1049        for (const auto &Pair : MappedDeclsFields) {1050          assert(Pair.getFirst()->isCanonicalDecl() &&1051                 "Expected canonical declaration");1052          Data.try_emplace(Pair.getFirst());1053        }1054      }1055      Rt.emitGenericVarsProlog(CGF, Loc);1056    }1057    void Exit(CodeGenFunction &CGF) override {1058      static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime())1059          .emitGenericVarsEpilog(CGF);1060    }1061  } Action(Loc, GlobalizedRD, MappedDeclsFields);1062  CodeGen.setAction(Action);1063  llvm::Function *OutlinedFun = CGOpenMPRuntime::emitTeamsOutlinedFunction(1064      CGF, D, ThreadIDVar, InnermostKind, CodeGen);1065 1066  return OutlinedFun;1067}1068 1069void CGOpenMPRuntimeGPU::emitGenericVarsProlog(CodeGenFunction &CGF,1070                                               SourceLocation Loc) {1071  if (getDataSharingMode() != CGOpenMPRuntimeGPU::DS_Generic)1072    return;1073 1074  CGBuilderTy &Bld = CGF.Builder;1075 1076  const auto I = FunctionGlobalizedDecls.find(CGF.CurFn);1077  if (I == FunctionGlobalizedDecls.end())1078    return;1079 1080  for (auto &Rec : I->getSecond().LocalVarData) {1081    const auto *VD = cast<VarDecl>(Rec.first);1082    bool EscapedParam = I->getSecond().EscapedParameters.count(Rec.first);1083    QualType VarTy = VD->getType();1084 1085    // Get the local allocation of a firstprivate variable before sharing1086    llvm::Value *ParValue;1087    if (EscapedParam) {1088      LValue ParLVal =1089          CGF.MakeAddrLValue(CGF.GetAddrOfLocalVar(VD), VD->getType());1090      ParValue = CGF.EmitLoadOfScalar(ParLVal, Loc);1091    }1092 1093    // Allocate space for the variable to be globalized1094    llvm::Value *AllocArgs[] = {CGF.getTypeSize(VD->getType())};1095    llvm::CallBase *VoidPtr =1096        CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(1097                                CGM.getModule(), OMPRTL___kmpc_alloc_shared),1098                            AllocArgs, VD->getName());1099    // FIXME: We should use the variables actual alignment as an argument.1100    VoidPtr->addRetAttr(llvm::Attribute::get(1101        CGM.getLLVMContext(), llvm::Attribute::Alignment,1102        CGM.getContext().getTargetInfo().getNewAlign() / 8));1103 1104    // Cast the void pointer and get the address of the globalized variable.1105    llvm::Value *CastedVoidPtr = Bld.CreatePointerBitCastOrAddrSpaceCast(1106        VoidPtr, Bld.getPtrTy(0), VD->getName() + "_on_stack");1107    LValue VarAddr =1108        CGF.MakeNaturalAlignPointeeRawAddrLValue(CastedVoidPtr, VarTy);1109    Rec.second.PrivateAddr = VarAddr.getAddress();1110    Rec.second.GlobalizedVal = VoidPtr;1111 1112    // Assign the local allocation to the newly globalized location.1113    if (EscapedParam) {1114      CGF.EmitStoreOfScalar(ParValue, VarAddr);1115      I->getSecond().MappedParams->setVarAddr(CGF, VD, VarAddr.getAddress());1116    }1117    if (auto *DI = CGF.getDebugInfo())1118      VoidPtr->setDebugLoc(DI->SourceLocToDebugLoc(VD->getLocation()));1119  }1120 1121  for (const auto *ValueD : I->getSecond().EscapedVariableLengthDecls) {1122    const auto *VD = cast<VarDecl>(ValueD);1123    std::pair<llvm::Value *, llvm::Value *> AddrSizePair =1124        getKmpcAllocShared(CGF, VD);1125    I->getSecond().EscapedVariableLengthDeclsAddrs.emplace_back(AddrSizePair);1126    LValue Base = CGF.MakeAddrLValue(AddrSizePair.first, VD->getType(),1127                                     CGM.getContext().getDeclAlign(VD),1128                                     AlignmentSource::Decl);1129    I->getSecond().MappedParams->setVarAddr(CGF, VD, Base.getAddress());1130  }1131  I->getSecond().MappedParams->apply(CGF);1132}1133 1134bool CGOpenMPRuntimeGPU::isDelayedVariableLengthDecl(CodeGenFunction &CGF,1135                                                     const VarDecl *VD) const {1136  const auto I = FunctionGlobalizedDecls.find(CGF.CurFn);1137  if (I == FunctionGlobalizedDecls.end())1138    return false;1139 1140  // Check variable declaration is delayed:1141  return llvm::is_contained(I->getSecond().DelayedVariableLengthDecls, VD);1142}1143 1144std::pair<llvm::Value *, llvm::Value *>1145CGOpenMPRuntimeGPU::getKmpcAllocShared(CodeGenFunction &CGF,1146                                       const VarDecl *VD) {1147  CGBuilderTy &Bld = CGF.Builder;1148 1149  // Compute size and alignment.1150  llvm::Value *Size = CGF.getTypeSize(VD->getType());1151  CharUnits Align = CGM.getContext().getDeclAlign(VD);1152  Size = Bld.CreateNUWAdd(1153      Size, llvm::ConstantInt::get(CGF.SizeTy, Align.getQuantity() - 1));1154  llvm::Value *AlignVal =1155      llvm::ConstantInt::get(CGF.SizeTy, Align.getQuantity());1156  Size = Bld.CreateUDiv(Size, AlignVal);1157  Size = Bld.CreateNUWMul(Size, AlignVal);1158 1159  // Allocate space for this VLA object to be globalized.1160  llvm::Value *AllocArgs[] = {Size};1161  llvm::CallBase *VoidPtr =1162      CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(1163                              CGM.getModule(), OMPRTL___kmpc_alloc_shared),1164                          AllocArgs, VD->getName());1165  VoidPtr->addRetAttr(llvm::Attribute::get(1166      CGM.getLLVMContext(), llvm::Attribute::Alignment, Align.getQuantity()));1167 1168  return std::make_pair(VoidPtr, Size);1169}1170 1171void CGOpenMPRuntimeGPU::getKmpcFreeShared(1172    CodeGenFunction &CGF,1173    const std::pair<llvm::Value *, llvm::Value *> &AddrSizePair) {1174  // Deallocate the memory for each globalized VLA object1175  CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(1176                          CGM.getModule(), OMPRTL___kmpc_free_shared),1177                      {AddrSizePair.first, AddrSizePair.second});1178}1179 1180void CGOpenMPRuntimeGPU::emitGenericVarsEpilog(CodeGenFunction &CGF) {1181  if (getDataSharingMode() != CGOpenMPRuntimeGPU::DS_Generic)1182    return;1183 1184  const auto I = FunctionGlobalizedDecls.find(CGF.CurFn);1185  if (I != FunctionGlobalizedDecls.end()) {1186    // Deallocate the memory for each globalized VLA object that was1187    // globalized in the prolog (i.e. emitGenericVarsProlog).1188    for (const auto &AddrSizePair :1189         llvm::reverse(I->getSecond().EscapedVariableLengthDeclsAddrs)) {1190      CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(1191                              CGM.getModule(), OMPRTL___kmpc_free_shared),1192                          {AddrSizePair.first, AddrSizePair.second});1193    }1194    // Deallocate the memory for each globalized value1195    for (auto &Rec : llvm::reverse(I->getSecond().LocalVarData)) {1196      const auto *VD = cast<VarDecl>(Rec.first);1197      I->getSecond().MappedParams->restore(CGF);1198 1199      llvm::Value *FreeArgs[] = {Rec.second.GlobalizedVal,1200                                 CGF.getTypeSize(VD->getType())};1201      CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(1202                              CGM.getModule(), OMPRTL___kmpc_free_shared),1203                          FreeArgs);1204    }1205  }1206}1207 1208void CGOpenMPRuntimeGPU::emitTeamsCall(CodeGenFunction &CGF,1209                                         const OMPExecutableDirective &D,1210                                         SourceLocation Loc,1211                                         llvm::Function *OutlinedFn,1212                                         ArrayRef<llvm::Value *> CapturedVars) {1213  if (!CGF.HaveInsertPoint())1214    return;1215 1216  bool IsBareKernel = D.getSingleClause<OMPXBareClause>();1217 1218  RawAddress ZeroAddr = CGF.CreateDefaultAlignTempAlloca(CGF.Int32Ty,1219                                                         /*Name=*/".zero.addr");1220  CGF.Builder.CreateStore(CGF.Builder.getInt32(/*C*/ 0), ZeroAddr);1221  llvm::SmallVector<llvm::Value *, 16> OutlinedFnArgs;1222  // We don't emit any thread id function call in bare kernel, but because the1223  // outlined function has a pointer argument, we emit a nullptr here.1224  if (IsBareKernel)1225    OutlinedFnArgs.push_back(llvm::ConstantPointerNull::get(CGM.VoidPtrTy));1226  else1227    OutlinedFnArgs.push_back(emitThreadIDAddress(CGF, Loc).emitRawPointer(CGF));1228  OutlinedFnArgs.push_back(ZeroAddr.getPointer());1229  OutlinedFnArgs.append(CapturedVars.begin(), CapturedVars.end());1230  emitOutlinedFunctionCall(CGF, Loc, OutlinedFn, OutlinedFnArgs);1231}1232 1233void CGOpenMPRuntimeGPU::emitParallelCall(1234    CodeGenFunction &CGF, SourceLocation Loc, llvm::Function *OutlinedFn,1235    ArrayRef<llvm::Value *> CapturedVars, const Expr *IfCond,1236    llvm::Value *NumThreads, OpenMPNumThreadsClauseModifier NumThreadsModifier,1237    OpenMPSeverityClauseKind Severity, const Expr *Message) {1238  if (!CGF.HaveInsertPoint())1239    return;1240 1241  auto &&ParallelGen = [this, Loc, OutlinedFn, CapturedVars, IfCond,1242                        NumThreads](CodeGenFunction &CGF,1243                                    PrePostActionTy &Action) {1244    CGBuilderTy &Bld = CGF.Builder;1245    llvm::Value *NumThreadsVal = NumThreads;1246    llvm::Function *WFn = WrapperFunctionsMap[OutlinedFn];1247    llvm::PointerType *FnPtrTy = llvm::PointerType::get(1248        CGF.getLLVMContext(), CGM.getDataLayout().getProgramAddressSpace());1249 1250    llvm::Value *ID = llvm::ConstantPointerNull::get(FnPtrTy);1251    if (WFn)1252      ID = Bld.CreateBitOrPointerCast(WFn, FnPtrTy);1253 1254    llvm::Value *FnPtr = Bld.CreateBitOrPointerCast(OutlinedFn, FnPtrTy);1255 1256    // Create a private scope that will globalize the arguments1257    // passed from the outside of the target region.1258    // TODO: Is that needed?1259    CodeGenFunction::OMPPrivateScope PrivateArgScope(CGF);1260 1261    Address CapturedVarsAddrs = CGF.CreateDefaultAlignTempAlloca(1262        llvm::ArrayType::get(CGM.VoidPtrTy, CapturedVars.size()),1263        "captured_vars_addrs");1264    // There's something to share.1265    if (!CapturedVars.empty()) {1266      // Prepare for parallel region. Indicate the outlined function.1267      ASTContext &Ctx = CGF.getContext();1268      unsigned Idx = 0;1269      for (llvm::Value *V : CapturedVars) {1270        Address Dst = Bld.CreateConstArrayGEP(CapturedVarsAddrs, Idx);1271        llvm::Value *PtrV;1272        if (V->getType()->isIntegerTy())1273          PtrV = Bld.CreateIntToPtr(V, CGF.VoidPtrTy);1274        else1275          PtrV = Bld.CreatePointerBitCastOrAddrSpaceCast(V, CGF.VoidPtrTy);1276        CGF.EmitStoreOfScalar(PtrV, Dst, /*Volatile=*/false,1277                              Ctx.getPointerType(Ctx.VoidPtrTy));1278        ++Idx;1279      }1280    }1281 1282    llvm::Value *IfCondVal = nullptr;1283    if (IfCond)1284      IfCondVal = Bld.CreateIntCast(CGF.EvaluateExprAsBool(IfCond), CGF.Int32Ty,1285                                    /* isSigned */ false);1286    else1287      IfCondVal = llvm::ConstantInt::get(CGF.Int32Ty, 1);1288 1289    if (!NumThreadsVal)1290      NumThreadsVal = llvm::ConstantInt::get(CGF.Int32Ty, -1);1291    else1292      NumThreadsVal = Bld.CreateZExtOrTrunc(NumThreadsVal, CGF.Int32Ty);1293 1294    assert(IfCondVal && "Expected a value");1295    llvm::Value *RTLoc = emitUpdateLocation(CGF, Loc);1296    llvm::Value *Args[] = {1297        RTLoc,1298        getThreadID(CGF, Loc),1299        IfCondVal,1300        NumThreadsVal,1301        llvm::ConstantInt::get(CGF.Int32Ty, -1),1302        FnPtr,1303        ID,1304        Bld.CreateBitOrPointerCast(CapturedVarsAddrs.emitRawPointer(CGF),1305                                   CGF.VoidPtrPtrTy),1306        llvm::ConstantInt::get(CGM.SizeTy, CapturedVars.size())};1307    CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(1308                            CGM.getModule(), OMPRTL___kmpc_parallel_51),1309                        Args);1310  };1311 1312  RegionCodeGenTy RCG(ParallelGen);1313  RCG(CGF);1314}1315 1316void CGOpenMPRuntimeGPU::syncCTAThreads(CodeGenFunction &CGF) {1317  // Always emit simple barriers!1318  if (!CGF.HaveInsertPoint())1319    return;1320  // Build call __kmpc_barrier_simple_spmd(nullptr, 0);1321  // This function does not use parameters, so we can emit just default values.1322  llvm::Value *Args[] = {1323      llvm::ConstantPointerNull::get(1324          cast<llvm::PointerType>(getIdentTyPointerTy())),1325      llvm::ConstantInt::get(CGF.Int32Ty, /*V=*/0, /*isSigned=*/true)};1326  CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(1327                          CGM.getModule(), OMPRTL___kmpc_barrier_simple_spmd),1328                      Args);1329}1330 1331void CGOpenMPRuntimeGPU::emitBarrierCall(CodeGenFunction &CGF,1332                                           SourceLocation Loc,1333                                           OpenMPDirectiveKind Kind, bool,1334                                           bool) {1335  // Always emit simple barriers!1336  if (!CGF.HaveInsertPoint())1337    return;1338  // Build call __kmpc_cancel_barrier(loc, thread_id);1339  unsigned Flags = getDefaultFlagsForBarriers(Kind);1340  llvm::Value *Args[] = {emitUpdateLocation(CGF, Loc, Flags),1341                         getThreadID(CGF, Loc)};1342 1343  CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(1344                          CGM.getModule(), OMPRTL___kmpc_barrier),1345                      Args);1346}1347 1348void CGOpenMPRuntimeGPU::emitCriticalRegion(1349    CodeGenFunction &CGF, StringRef CriticalName,1350    const RegionCodeGenTy &CriticalOpGen, SourceLocation Loc,1351    const Expr *Hint) {1352  llvm::BasicBlock *LoopBB = CGF.createBasicBlock("omp.critical.loop");1353  llvm::BasicBlock *TestBB = CGF.createBasicBlock("omp.critical.test");1354  llvm::BasicBlock *SyncBB = CGF.createBasicBlock("omp.critical.sync");1355  llvm::BasicBlock *BodyBB = CGF.createBasicBlock("omp.critical.body");1356  llvm::BasicBlock *ExitBB = CGF.createBasicBlock("omp.critical.exit");1357 1358  auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());1359 1360  // Get the mask of active threads in the warp.1361  llvm::Value *Mask = CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(1362      CGM.getModule(), OMPRTL___kmpc_warp_active_thread_mask));1363  // Fetch team-local id of the thread.1364  llvm::Value *ThreadID = RT.getGPUThreadID(CGF);1365 1366  // Get the width of the team.1367  llvm::Value *TeamWidth = RT.getGPUNumThreads(CGF);1368 1369  // Initialize the counter variable for the loop.1370  QualType Int32Ty =1371      CGF.getContext().getIntTypeForBitwidth(/*DestWidth=*/32, /*Signed=*/0);1372  Address Counter = CGF.CreateMemTemp(Int32Ty, "critical_counter");1373  LValue CounterLVal = CGF.MakeAddrLValue(Counter, Int32Ty);1374  CGF.EmitStoreOfScalar(llvm::Constant::getNullValue(CGM.Int32Ty), CounterLVal,1375                        /*isInit=*/true);1376 1377  // Block checks if loop counter exceeds upper bound.1378  CGF.EmitBlock(LoopBB);1379  llvm::Value *CounterVal = CGF.EmitLoadOfScalar(CounterLVal, Loc);1380  llvm::Value *CmpLoopBound = CGF.Builder.CreateICmpSLT(CounterVal, TeamWidth);1381  CGF.Builder.CreateCondBr(CmpLoopBound, TestBB, ExitBB);1382 1383  // Block tests which single thread should execute region, and which threads1384  // should go straight to synchronisation point.1385  CGF.EmitBlock(TestBB);1386  CounterVal = CGF.EmitLoadOfScalar(CounterLVal, Loc);1387  llvm::Value *CmpThreadToCounter =1388      CGF.Builder.CreateICmpEQ(ThreadID, CounterVal);1389  CGF.Builder.CreateCondBr(CmpThreadToCounter, BodyBB, SyncBB);1390 1391  // Block emits the body of the critical region.1392  CGF.EmitBlock(BodyBB);1393 1394  // Output the critical statement.1395  CGOpenMPRuntime::emitCriticalRegion(CGF, CriticalName, CriticalOpGen, Loc,1396                                      Hint);1397 1398  // After the body surrounded by the critical region, the single executing1399  // thread will jump to the synchronisation point.1400  // Block waits for all threads in current team to finish then increments the1401  // counter variable and returns to the loop.1402  CGF.EmitBlock(SyncBB);1403  // Reconverge active threads in the warp.1404  (void)CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(1405                                CGM.getModule(), OMPRTL___kmpc_syncwarp),1406                            Mask);1407 1408  llvm::Value *IncCounterVal =1409      CGF.Builder.CreateNSWAdd(CounterVal, CGF.Builder.getInt32(1));1410  CGF.EmitStoreOfScalar(IncCounterVal, CounterLVal);1411  CGF.EmitBranch(LoopBB);1412 1413  // Block that is reached when  all threads in the team complete the region.1414  CGF.EmitBlock(ExitBB, /*IsFinished=*/true);1415}1416 1417/// Cast value to the specified type.1418static llvm::Value *castValueToType(CodeGenFunction &CGF, llvm::Value *Val,1419                                    QualType ValTy, QualType CastTy,1420                                    SourceLocation Loc) {1421  assert(!CGF.getContext().getTypeSizeInChars(CastTy).isZero() &&1422         "Cast type must sized.");1423  assert(!CGF.getContext().getTypeSizeInChars(ValTy).isZero() &&1424         "Val type must sized.");1425  llvm::Type *LLVMCastTy = CGF.ConvertTypeForMem(CastTy);1426  if (ValTy == CastTy)1427    return Val;1428  if (CGF.getContext().getTypeSizeInChars(ValTy) ==1429      CGF.getContext().getTypeSizeInChars(CastTy))1430    return CGF.Builder.CreateBitCast(Val, LLVMCastTy);1431  if (CastTy->isIntegerType() && ValTy->isIntegerType())1432    return CGF.Builder.CreateIntCast(Val, LLVMCastTy,1433                                     CastTy->hasSignedIntegerRepresentation());1434  Address CastItem = CGF.CreateMemTemp(CastTy);1435  Address ValCastItem = CastItem.withElementType(Val->getType());1436  CGF.EmitStoreOfScalar(Val, ValCastItem, /*Volatile=*/false, ValTy,1437                        LValueBaseInfo(AlignmentSource::Type),1438                        TBAAAccessInfo());1439  return CGF.EmitLoadOfScalar(CastItem, /*Volatile=*/false, CastTy, Loc,1440                              LValueBaseInfo(AlignmentSource::Type),1441                              TBAAAccessInfo());1442}1443 1444///1445/// Design of OpenMP reductions on the GPU1446///1447/// Consider a typical OpenMP program with one or more reduction1448/// clauses:1449///1450/// float foo;1451/// double bar;1452/// #pragma omp target teams distribute parallel for \1453///             reduction(+:foo) reduction(*:bar)1454/// for (int i = 0; i < N; i++) {1455///   foo += A[i]; bar *= B[i];1456/// }1457///1458/// where 'foo' and 'bar' are reduced across all OpenMP threads in1459/// all teams.  In our OpenMP implementation on the NVPTX device an1460/// OpenMP team is mapped to a CUDA threadblock and OpenMP threads1461/// within a team are mapped to CUDA threads within a threadblock.1462/// Our goal is to efficiently aggregate values across all OpenMP1463/// threads such that:1464///1465///   - the compiler and runtime are logically concise, and1466///   - the reduction is performed efficiently in a hierarchical1467///     manner as follows: within OpenMP threads in the same warp,1468///     across warps in a threadblock, and finally across teams on1469///     the NVPTX device.1470///1471/// Introduction to Decoupling1472///1473/// We would like to decouple the compiler and the runtime so that the1474/// latter is ignorant of the reduction variables (number, data types)1475/// and the reduction operators.  This allows a simpler interface1476/// and implementation while still attaining good performance.1477///1478/// Pseudocode for the aforementioned OpenMP program generated by the1479/// compiler is as follows:1480///1481/// 1. Create private copies of reduction variables on each OpenMP1482///    thread: 'foo_private', 'bar_private'1483/// 2. Each OpenMP thread reduces the chunk of 'A' and 'B' assigned1484///    to it and writes the result in 'foo_private' and 'bar_private'1485///    respectively.1486/// 3. Call the OpenMP runtime on the GPU to reduce within a team1487///    and store the result on the team master:1488///1489///     __kmpc_nvptx_parallel_reduce_nowait_v2(...,1490///        reduceData, shuffleReduceFn, interWarpCpyFn)1491///1492///     where:1493///       struct ReduceData {1494///         double *foo;1495///         double *bar;1496///       } reduceData1497///       reduceData.foo = &foo_private1498///       reduceData.bar = &bar_private1499///1500///     'shuffleReduceFn' and 'interWarpCpyFn' are pointers to two1501///     auxiliary functions generated by the compiler that operate on1502///     variables of type 'ReduceData'.  They aid the runtime perform1503///     algorithmic steps in a data agnostic manner.1504///1505///     'shuffleReduceFn' is a pointer to a function that reduces data1506///     of type 'ReduceData' across two OpenMP threads (lanes) in the1507///     same warp.  It takes the following arguments as input:1508///1509///     a. variable of type 'ReduceData' on the calling lane,1510///     b. its lane_id,1511///     c. an offset relative to the current lane_id to generate a1512///        remote_lane_id.  The remote lane contains the second1513///        variable of type 'ReduceData' that is to be reduced.1514///     d. an algorithm version parameter determining which reduction1515///        algorithm to use.1516///1517///     'shuffleReduceFn' retrieves data from the remote lane using1518///     efficient GPU shuffle intrinsics and reduces, using the1519///     algorithm specified by the 4th parameter, the two operands1520///     element-wise.  The result is written to the first operand.1521///1522///     Different reduction algorithms are implemented in different1523///     runtime functions, all calling 'shuffleReduceFn' to perform1524///     the essential reduction step.  Therefore, based on the 4th1525///     parameter, this function behaves slightly differently to1526///     cooperate with the runtime to ensure correctness under1527///     different circumstances.1528///1529///     'InterWarpCpyFn' is a pointer to a function that transfers1530///     reduced variables across warps.  It tunnels, through CUDA1531///     shared memory, the thread-private data of type 'ReduceData'1532///     from lane 0 of each warp to a lane in the first warp.1533/// 4. Call the OpenMP runtime on the GPU to reduce across teams.1534///    The last team writes the global reduced value to memory.1535///1536///     ret = __kmpc_nvptx_teams_reduce_nowait(...,1537///             reduceData, shuffleReduceFn, interWarpCpyFn,1538///             scratchpadCopyFn, loadAndReduceFn)1539///1540///     'scratchpadCopyFn' is a helper that stores reduced1541///     data from the team master to a scratchpad array in1542///     global memory.1543///1544///     'loadAndReduceFn' is a helper that loads data from1545///     the scratchpad array and reduces it with the input1546///     operand.1547///1548///     These compiler generated functions hide address1549///     calculation and alignment information from the runtime.1550/// 5. if ret == 1:1551///     The team master of the last team stores the reduced1552///     result to the globals in memory.1553///     foo += reduceData.foo; bar *= reduceData.bar1554///1555///1556/// Warp Reduction Algorithms1557///1558/// On the warp level, we have three algorithms implemented in the1559/// OpenMP runtime depending on the number of active lanes:1560///1561/// Full Warp Reduction1562///1563/// The reduce algorithm within a warp where all lanes are active1564/// is implemented in the runtime as follows:1565///1566/// full_warp_reduce(void *reduce_data,1567///                  kmp_ShuffleReductFctPtr ShuffleReduceFn) {1568///   for (int offset = WARPSIZE/2; offset > 0; offset /= 2)1569///     ShuffleReduceFn(reduce_data, 0, offset, 0);1570/// }1571///1572/// The algorithm completes in log(2, WARPSIZE) steps.1573///1574/// 'ShuffleReduceFn' is used here with lane_id set to 0 because it is1575/// not used therefore we save instructions by not retrieving lane_id1576/// from the corresponding special registers.  The 4th parameter, which1577/// represents the version of the algorithm being used, is set to 0 to1578/// signify full warp reduction.1579///1580/// In this version, 'ShuffleReduceFn' behaves, per element, as follows:1581///1582/// #reduce_elem refers to an element in the local lane's data structure1583/// #remote_elem is retrieved from a remote lane1584/// remote_elem = shuffle_down(reduce_elem, offset, WARPSIZE);1585/// reduce_elem = reduce_elem REDUCE_OP remote_elem;1586///1587/// Contiguous Partial Warp Reduction1588///1589/// This reduce algorithm is used within a warp where only the first1590/// 'n' (n <= WARPSIZE) lanes are active.  It is typically used when the1591/// number of OpenMP threads in a parallel region is not a multiple of1592/// WARPSIZE.  The algorithm is implemented in the runtime as follows:1593///1594/// void1595/// contiguous_partial_reduce(void *reduce_data,1596///                           kmp_ShuffleReductFctPtr ShuffleReduceFn,1597///                           int size, int lane_id) {1598///   int curr_size;1599///   int offset;1600///   curr_size = size;1601///   mask = curr_size/2;1602///   while (offset>0) {1603///     ShuffleReduceFn(reduce_data, lane_id, offset, 1);1604///     curr_size = (curr_size+1)/2;1605///     offset = curr_size/2;1606///   }1607/// }1608///1609/// In this version, 'ShuffleReduceFn' behaves, per element, as follows:1610///1611/// remote_elem = shuffle_down(reduce_elem, offset, WARPSIZE);1612/// if (lane_id < offset)1613///     reduce_elem = reduce_elem REDUCE_OP remote_elem1614/// else1615///     reduce_elem = remote_elem1616///1617/// This algorithm assumes that the data to be reduced are located in a1618/// contiguous subset of lanes starting from the first.  When there is1619/// an odd number of active lanes, the data in the last lane is not1620/// aggregated with any other lane's dat but is instead copied over.1621///1622/// Dispersed Partial Warp Reduction1623///1624/// This algorithm is used within a warp when any discontiguous subset of1625/// lanes are active.  It is used to implement the reduction operation1626/// across lanes in an OpenMP simd region or in a nested parallel region.1627///1628/// void1629/// dispersed_partial_reduce(void *reduce_data,1630///                          kmp_ShuffleReductFctPtr ShuffleReduceFn) {1631///   int size, remote_id;1632///   int logical_lane_id = number_of_active_lanes_before_me() * 2;1633///   do {1634///       remote_id = next_active_lane_id_right_after_me();1635///       # the above function returns 0 of no active lane1636///       # is present right after the current lane.1637///       size = number_of_active_lanes_in_this_warp();1638///       logical_lane_id /= 2;1639///       ShuffleReduceFn(reduce_data, logical_lane_id,1640///                       remote_id-1-threadIdx.x, 2);1641///   } while (logical_lane_id % 2 == 0 && size > 1);1642/// }1643///1644/// There is no assumption made about the initial state of the reduction.1645/// Any number of lanes (>=1) could be active at any position.  The reduction1646/// result is returned in the first active lane.1647///1648/// In this version, 'ShuffleReduceFn' behaves, per element, as follows:1649///1650/// remote_elem = shuffle_down(reduce_elem, offset, WARPSIZE);1651/// if (lane_id % 2 == 0 && offset > 0)1652///     reduce_elem = reduce_elem REDUCE_OP remote_elem1653/// else1654///     reduce_elem = remote_elem1655///1656///1657/// Intra-Team Reduction1658///1659/// This function, as implemented in the runtime call1660/// '__kmpc_nvptx_parallel_reduce_nowait_v2', aggregates data across OpenMP1661/// threads in a team.  It first reduces within a warp using the1662/// aforementioned algorithms.  We then proceed to gather all such1663/// reduced values at the first warp.1664///1665/// The runtime makes use of the function 'InterWarpCpyFn', which copies1666/// data from each of the "warp master" (zeroth lane of each warp, where1667/// warp-reduced data is held) to the zeroth warp.  This step reduces (in1668/// a mathematical sense) the problem of reduction across warp masters in1669/// a block to the problem of warp reduction.1670///1671///1672/// Inter-Team Reduction1673///1674/// Once a team has reduced its data to a single value, it is stored in1675/// a global scratchpad array.  Since each team has a distinct slot, this1676/// can be done without locking.1677///1678/// The last team to write to the scratchpad array proceeds to reduce the1679/// scratchpad array.  One or more workers in the last team use the helper1680/// 'loadAndReduceDataFn' to load and reduce values from the array, i.e.,1681/// the k'th worker reduces every k'th element.1682///1683/// Finally, a call is made to '__kmpc_nvptx_parallel_reduce_nowait_v2' to1684/// reduce across workers and compute a globally reduced value.1685///1686void CGOpenMPRuntimeGPU::emitReduction(1687    CodeGenFunction &CGF, SourceLocation Loc, ArrayRef<const Expr *> Privates,1688    ArrayRef<const Expr *> LHSExprs, ArrayRef<const Expr *> RHSExprs,1689    ArrayRef<const Expr *> ReductionOps, ReductionOptionsTy Options) {1690  if (!CGF.HaveInsertPoint())1691    return;1692 1693  bool ParallelReduction = isOpenMPParallelDirective(Options.ReductionKind);1694  bool TeamsReduction = isOpenMPTeamsDirective(Options.ReductionKind);1695 1696  ASTContext &C = CGM.getContext();1697 1698  if (Options.SimpleReduction) {1699    assert(!TeamsReduction && !ParallelReduction &&1700           "Invalid reduction selection in emitReduction.");1701    (void)ParallelReduction;1702    CGOpenMPRuntime::emitReduction(CGF, Loc, Privates, LHSExprs, RHSExprs,1703                                   ReductionOps, Options);1704    return;1705  }1706 1707  llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *> VarFieldMap;1708  llvm::SmallVector<const ValueDecl *, 4> PrivatesReductions(Privates.size());1709  int Cnt = 0;1710  for (const Expr *DRE : Privates) {1711    PrivatesReductions[Cnt] = cast<DeclRefExpr>(DRE)->getDecl();1712    ++Cnt;1713  }1714  const RecordDecl *ReductionRec = ::buildRecordForGlobalizedVars(1715      CGM.getContext(), PrivatesReductions, {}, VarFieldMap, 1);1716 1717  if (TeamsReduction)1718    TeamsReductions.push_back(ReductionRec);1719 1720  // Source location for the ident struct1721  llvm::Value *RTLoc = emitUpdateLocation(CGF, Loc);1722 1723  using InsertPointTy = llvm::OpenMPIRBuilder::InsertPointTy;1724  InsertPointTy AllocaIP(CGF.AllocaInsertPt->getParent(),1725                         CGF.AllocaInsertPt->getIterator());1726  InsertPointTy CodeGenIP(CGF.Builder.GetInsertBlock(),1727                          CGF.Builder.GetInsertPoint());1728  llvm::OpenMPIRBuilder::LocationDescription OmpLoc(1729      CodeGenIP, CGF.SourceLocToDebugLoc(Loc));1730  llvm::SmallVector<llvm::OpenMPIRBuilder::ReductionInfo, 2> ReductionInfos;1731 1732  CodeGenFunction::OMPPrivateScope Scope(CGF);1733  unsigned Idx = 0;1734  for (const Expr *Private : Privates) {1735    llvm::Type *ElementType;1736    llvm::Value *Variable;1737    llvm::Value *PrivateVariable;1738    llvm::OpenMPIRBuilder::ReductionGenAtomicCBTy AtomicReductionGen = nullptr;1739    ElementType = CGF.ConvertTypeForMem(Private->getType());1740    const auto *RHSVar =1741        cast<VarDecl>(cast<DeclRefExpr>(RHSExprs[Idx])->getDecl());1742    PrivateVariable = CGF.GetAddrOfLocalVar(RHSVar).emitRawPointer(CGF);1743    const auto *LHSVar =1744        cast<VarDecl>(cast<DeclRefExpr>(LHSExprs[Idx])->getDecl());1745    Variable = CGF.GetAddrOfLocalVar(LHSVar).emitRawPointer(CGF);1746    llvm::OpenMPIRBuilder::EvalKind EvalKind;1747    switch (CGF.getEvaluationKind(Private->getType())) {1748    case TEK_Scalar:1749      EvalKind = llvm::OpenMPIRBuilder::EvalKind::Scalar;1750      break;1751    case TEK_Complex:1752      EvalKind = llvm::OpenMPIRBuilder::EvalKind::Complex;1753      break;1754    case TEK_Aggregate:1755      EvalKind = llvm::OpenMPIRBuilder::EvalKind::Aggregate;1756      break;1757    }1758    auto ReductionGen = [&](InsertPointTy CodeGenIP, unsigned I,1759                            llvm::Value **LHSPtr, llvm::Value **RHSPtr,1760                            llvm::Function *NewFunc) {1761      CGF.Builder.restoreIP(CodeGenIP);1762      auto *CurFn = CGF.CurFn;1763      CGF.CurFn = NewFunc;1764 1765      *LHSPtr = CGF.GetAddrOfLocalVar(1766                       cast<VarDecl>(cast<DeclRefExpr>(LHSExprs[I])->getDecl()))1767                    .emitRawPointer(CGF);1768      *RHSPtr = CGF.GetAddrOfLocalVar(1769                       cast<VarDecl>(cast<DeclRefExpr>(RHSExprs[I])->getDecl()))1770                    .emitRawPointer(CGF);1771 1772      emitSingleReductionCombiner(CGF, ReductionOps[I], Privates[I],1773                                  cast<DeclRefExpr>(LHSExprs[I]),1774                                  cast<DeclRefExpr>(RHSExprs[I]));1775 1776      CGF.CurFn = CurFn;1777 1778      return InsertPointTy(CGF.Builder.GetInsertBlock(),1779                           CGF.Builder.GetInsertPoint());1780    };1781    ReductionInfos.emplace_back(llvm::OpenMPIRBuilder::ReductionInfo(1782        ElementType, Variable, PrivateVariable, EvalKind,1783        /*ReductionGen=*/nullptr, ReductionGen, AtomicReductionGen,1784        /*DataPtrPtrGen=*/nullptr));1785    Idx++;1786  }1787 1788  llvm::OpenMPIRBuilder::InsertPointTy AfterIP =1789      cantFail(OMPBuilder.createReductionsGPU(1790          OmpLoc, AllocaIP, CodeGenIP, ReductionInfos, /*IsByRef=*/{}, false,1791          TeamsReduction, llvm::OpenMPIRBuilder::ReductionGenCBKind::Clang,1792          CGF.getTarget().getGridValue(),1793          C.getLangOpts().OpenMPCUDAReductionBufNum, RTLoc));1794  CGF.Builder.restoreIP(AfterIP);1795}1796 1797const VarDecl *1798CGOpenMPRuntimeGPU::translateParameter(const FieldDecl *FD,1799                                       const VarDecl *NativeParam) const {1800  if (!NativeParam->getType()->isReferenceType())1801    return NativeParam;1802  QualType ArgType = NativeParam->getType();1803  QualifierCollector QC;1804  const Type *NonQualTy = QC.strip(ArgType);1805  QualType PointeeTy = cast<ReferenceType>(NonQualTy)->getPointeeType();1806  if (const auto *Attr = FD->getAttr<OMPCaptureKindAttr>()) {1807    if (Attr->getCaptureKind() == OMPC_map) {1808      PointeeTy = CGM.getContext().getAddrSpaceQualType(PointeeTy,1809                                                        LangAS::opencl_global);1810    }1811  }1812  ArgType = CGM.getContext().getPointerType(PointeeTy);1813  QC.addRestrict();1814  enum { NVPTX_local_addr = 5 };1815  QC.addAddressSpace(getLangASFromTargetAS(NVPTX_local_addr));1816  ArgType = QC.apply(CGM.getContext(), ArgType);1817  if (isa<ImplicitParamDecl>(NativeParam))1818    return ImplicitParamDecl::Create(1819        CGM.getContext(), /*DC=*/nullptr, NativeParam->getLocation(),1820        NativeParam->getIdentifier(), ArgType, ImplicitParamKind::Other);1821  return ParmVarDecl::Create(1822      CGM.getContext(),1823      const_cast<DeclContext *>(NativeParam->getDeclContext()),1824      NativeParam->getBeginLoc(), NativeParam->getLocation(),1825      NativeParam->getIdentifier(), ArgType,1826      /*TInfo=*/nullptr, SC_None, /*DefArg=*/nullptr);1827}1828 1829Address1830CGOpenMPRuntimeGPU::getParameterAddress(CodeGenFunction &CGF,1831                                          const VarDecl *NativeParam,1832                                          const VarDecl *TargetParam) const {1833  assert(NativeParam != TargetParam &&1834         NativeParam->getType()->isReferenceType() &&1835         "Native arg must not be the same as target arg.");1836  Address LocalAddr = CGF.GetAddrOfLocalVar(TargetParam);1837  QualType NativeParamType = NativeParam->getType();1838  QualifierCollector QC;1839  const Type *NonQualTy = QC.strip(NativeParamType);1840  QualType NativePointeeTy = cast<ReferenceType>(NonQualTy)->getPointeeType();1841  unsigned NativePointeeAddrSpace =1842      CGF.getTypes().getTargetAddressSpace(NativePointeeTy);1843  QualType TargetTy = TargetParam->getType();1844  llvm::Value *TargetAddr = CGF.EmitLoadOfScalar(LocalAddr, /*Volatile=*/false,1845                                                 TargetTy, SourceLocation());1846  // Cast to native address space.1847  TargetAddr = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(1848      TargetAddr,1849      llvm::PointerType::get(CGF.getLLVMContext(), NativePointeeAddrSpace));1850  Address NativeParamAddr = CGF.CreateMemTemp(NativeParamType);1851  CGF.EmitStoreOfScalar(TargetAddr, NativeParamAddr, /*Volatile=*/false,1852                        NativeParamType);1853  return NativeParamAddr;1854}1855 1856void CGOpenMPRuntimeGPU::emitOutlinedFunctionCall(1857    CodeGenFunction &CGF, SourceLocation Loc, llvm::FunctionCallee OutlinedFn,1858    ArrayRef<llvm::Value *> Args) const {1859  SmallVector<llvm::Value *, 4> TargetArgs;1860  TargetArgs.reserve(Args.size());1861  auto *FnType = OutlinedFn.getFunctionType();1862  for (unsigned I = 0, E = Args.size(); I < E; ++I) {1863    if (FnType->isVarArg() && FnType->getNumParams() <= I) {1864      TargetArgs.append(std::next(Args.begin(), I), Args.end());1865      break;1866    }1867    llvm::Type *TargetType = FnType->getParamType(I);1868    llvm::Value *NativeArg = Args[I];1869    if (!TargetType->isPointerTy()) {1870      TargetArgs.emplace_back(NativeArg);1871      continue;1872    }1873    TargetArgs.emplace_back(1874        CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(NativeArg, TargetType));1875  }1876  CGOpenMPRuntime::emitOutlinedFunctionCall(CGF, Loc, OutlinedFn, TargetArgs);1877}1878 1879/// Emit function which wraps the outline parallel region1880/// and controls the arguments which are passed to this function.1881/// The wrapper ensures that the outlined function is called1882/// with the correct arguments when data is shared.1883llvm::Function *CGOpenMPRuntimeGPU::createParallelDataSharingWrapper(1884    llvm::Function *OutlinedParallelFn, const OMPExecutableDirective &D) {1885  ASTContext &Ctx = CGM.getContext();1886  const auto &CS = *D.getCapturedStmt(OMPD_parallel);1887 1888  // Create a function that takes as argument the source thread.1889  FunctionArgList WrapperArgs;1890  QualType Int16QTy =1891      Ctx.getIntTypeForBitwidth(/*DestWidth=*/16, /*Signed=*/false);1892  QualType Int32QTy =1893      Ctx.getIntTypeForBitwidth(/*DestWidth=*/32, /*Signed=*/false);1894  ImplicitParamDecl ParallelLevelArg(Ctx, /*DC=*/nullptr, D.getBeginLoc(),1895                                     /*Id=*/nullptr, Int16QTy,1896                                     ImplicitParamKind::Other);1897  ImplicitParamDecl WrapperArg(Ctx, /*DC=*/nullptr, D.getBeginLoc(),1898                               /*Id=*/nullptr, Int32QTy,1899                               ImplicitParamKind::Other);1900  WrapperArgs.emplace_back(&ParallelLevelArg);1901  WrapperArgs.emplace_back(&WrapperArg);1902 1903  const CGFunctionInfo &CGFI =1904      CGM.getTypes().arrangeBuiltinFunctionDeclaration(Ctx.VoidTy, WrapperArgs);1905 1906  auto *Fn = llvm::Function::Create(1907      CGM.getTypes().GetFunctionType(CGFI), llvm::GlobalValue::InternalLinkage,1908      Twine(OutlinedParallelFn->getName(), "_wrapper"), &CGM.getModule());1909 1910  // Ensure we do not inline the function. This is trivially true for the ones1911  // passed to __kmpc_fork_call but the ones calles in serialized regions1912  // could be inlined. This is not a perfect but it is closer to the invariant1913  // we want, namely, every data environment starts with a new function.1914  // TODO: We should pass the if condition to the runtime function and do the1915  //       handling there. Much cleaner code.1916  Fn->addFnAttr(llvm::Attribute::NoInline);1917 1918  CGM.SetInternalFunctionAttributes(GlobalDecl(), Fn, CGFI);1919  Fn->setLinkage(llvm::GlobalValue::InternalLinkage);1920  Fn->setDoesNotRecurse();1921 1922  CodeGenFunction CGF(CGM, /*suppressNewContext=*/true);1923  CGF.StartFunction(GlobalDecl(), Ctx.VoidTy, Fn, CGFI, WrapperArgs,1924                    D.getBeginLoc(), D.getBeginLoc());1925 1926  const auto *RD = CS.getCapturedRecordDecl();1927  auto CurField = RD->field_begin();1928 1929  Address ZeroAddr = CGF.CreateDefaultAlignTempAlloca(CGF.Int32Ty,1930                                                      /*Name=*/".zero.addr");1931  CGF.Builder.CreateStore(CGF.Builder.getInt32(/*C*/ 0), ZeroAddr);1932  // Get the array of arguments.1933  SmallVector<llvm::Value *, 8> Args;1934 1935  Args.emplace_back(CGF.GetAddrOfLocalVar(&WrapperArg).emitRawPointer(CGF));1936  Args.emplace_back(ZeroAddr.emitRawPointer(CGF));1937 1938  CGBuilderTy &Bld = CGF.Builder;1939  auto CI = CS.capture_begin();1940 1941  // Use global memory for data sharing.1942  // Handle passing of global args to workers.1943  RawAddress GlobalArgs =1944      CGF.CreateDefaultAlignTempAlloca(CGF.VoidPtrPtrTy, "global_args");1945  llvm::Value *GlobalArgsPtr = GlobalArgs.getPointer();1946  llvm::Value *DataSharingArgs[] = {GlobalArgsPtr};1947  CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(1948                          CGM.getModule(), OMPRTL___kmpc_get_shared_variables),1949                      DataSharingArgs);1950 1951  // Retrieve the shared variables from the list of references returned1952  // by the runtime. Pass the variables to the outlined function.1953  Address SharedArgListAddress = Address::invalid();1954  if (CS.capture_size() > 0 ||1955      isOpenMPLoopBoundSharingDirective(D.getDirectiveKind())) {1956    SharedArgListAddress = CGF.EmitLoadOfPointer(1957        GlobalArgs, CGF.getContext()1958                        .getPointerType(CGF.getContext().VoidPtrTy)1959                        .castAs<PointerType>());1960  }1961  unsigned Idx = 0;1962  if (isOpenMPLoopBoundSharingDirective(D.getDirectiveKind())) {1963    Address Src = Bld.CreateConstInBoundsGEP(SharedArgListAddress, Idx);1964    Address TypedAddress = Bld.CreatePointerBitCastOrAddrSpaceCast(1965        Src, Bld.getPtrTy(0), CGF.SizeTy);1966    llvm::Value *LB = CGF.EmitLoadOfScalar(1967        TypedAddress,1968        /*Volatile=*/false,1969        CGF.getContext().getPointerType(CGF.getContext().getSizeType()),1970        cast<OMPLoopDirective>(D).getLowerBoundVariable()->getExprLoc());1971    Args.emplace_back(LB);1972    ++Idx;1973    Src = Bld.CreateConstInBoundsGEP(SharedArgListAddress, Idx);1974    TypedAddress = Bld.CreatePointerBitCastOrAddrSpaceCast(Src, Bld.getPtrTy(0),1975                                                           CGF.SizeTy);1976    llvm::Value *UB = CGF.EmitLoadOfScalar(1977        TypedAddress,1978        /*Volatile=*/false,1979        CGF.getContext().getPointerType(CGF.getContext().getSizeType()),1980        cast<OMPLoopDirective>(D).getUpperBoundVariable()->getExprLoc());1981    Args.emplace_back(UB);1982    ++Idx;1983  }1984  if (CS.capture_size() > 0) {1985    ASTContext &CGFContext = CGF.getContext();1986    for (unsigned I = 0, E = CS.capture_size(); I < E; ++I, ++CI, ++CurField) {1987      QualType ElemTy = CurField->getType();1988      Address Src = Bld.CreateConstInBoundsGEP(SharedArgListAddress, I + Idx);1989      Address TypedAddress = Bld.CreatePointerBitCastOrAddrSpaceCast(1990          Src, CGF.ConvertTypeForMem(CGFContext.getPointerType(ElemTy)),1991          CGF.ConvertTypeForMem(ElemTy));1992      llvm::Value *Arg = CGF.EmitLoadOfScalar(TypedAddress,1993                                              /*Volatile=*/false,1994                                              CGFContext.getPointerType(ElemTy),1995                                              CI->getLocation());1996      if (CI->capturesVariableByCopy() &&1997          !CI->getCapturedVar()->getType()->isAnyPointerType()) {1998        Arg = castValueToType(CGF, Arg, ElemTy, CGFContext.getUIntPtrType(),1999                              CI->getLocation());2000      }2001      Args.emplace_back(Arg);2002    }2003  }2004 2005  emitOutlinedFunctionCall(CGF, D.getBeginLoc(), OutlinedParallelFn, Args);2006  CGF.FinishFunction();2007  return Fn;2008}2009 2010void CGOpenMPRuntimeGPU::emitFunctionProlog(CodeGenFunction &CGF,2011                                              const Decl *D) {2012  if (getDataSharingMode() != CGOpenMPRuntimeGPU::DS_Generic)2013    return;2014 2015  assert(D && "Expected function or captured|block decl.");2016  assert(FunctionGlobalizedDecls.count(CGF.CurFn) == 0 &&2017         "Function is registered already.");2018  assert((!TeamAndReductions.first || TeamAndReductions.first == D) &&2019         "Team is set but not processed.");2020  const Stmt *Body = nullptr;2021  bool NeedToDelayGlobalization = false;2022  if (const auto *FD = dyn_cast<FunctionDecl>(D)) {2023    Body = FD->getBody();2024  } else if (const auto *BD = dyn_cast<BlockDecl>(D)) {2025    Body = BD->getBody();2026  } else if (const auto *CD = dyn_cast<CapturedDecl>(D)) {2027    Body = CD->getBody();2028    NeedToDelayGlobalization = CGF.CapturedStmtInfo->getKind() == CR_OpenMP;2029    if (NeedToDelayGlobalization &&2030        getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD)2031      return;2032  }2033  if (!Body)2034    return;2035  CheckVarsEscapingDeclContext VarChecker(CGF, TeamAndReductions.second);2036  VarChecker.Visit(Body);2037  const RecordDecl *GlobalizedVarsRecord =2038      VarChecker.getGlobalizedRecord(IsInTTDRegion);2039  TeamAndReductions.first = nullptr;2040  TeamAndReductions.second.clear();2041  ArrayRef<const ValueDecl *> EscapedVariableLengthDecls =2042      VarChecker.getEscapedVariableLengthDecls();2043  ArrayRef<const ValueDecl *> DelayedVariableLengthDecls =2044      VarChecker.getDelayedVariableLengthDecls();2045  if (!GlobalizedVarsRecord && EscapedVariableLengthDecls.empty() &&2046      DelayedVariableLengthDecls.empty())2047    return;2048  auto I = FunctionGlobalizedDecls.try_emplace(CGF.CurFn).first;2049  I->getSecond().MappedParams =2050      std::make_unique<CodeGenFunction::OMPMapVars>();2051  I->getSecond().EscapedParameters.insert(2052      VarChecker.getEscapedParameters().begin(),2053      VarChecker.getEscapedParameters().end());2054  I->getSecond().EscapedVariableLengthDecls.append(2055      EscapedVariableLengthDecls.begin(), EscapedVariableLengthDecls.end());2056  I->getSecond().DelayedVariableLengthDecls.append(2057      DelayedVariableLengthDecls.begin(), DelayedVariableLengthDecls.end());2058  DeclToAddrMapTy &Data = I->getSecond().LocalVarData;2059  for (const ValueDecl *VD : VarChecker.getEscapedDecls()) {2060    assert(VD->isCanonicalDecl() && "Expected canonical declaration");2061    Data.try_emplace(VD);2062  }2063  if (!NeedToDelayGlobalization) {2064    emitGenericVarsProlog(CGF, D->getBeginLoc());2065    struct GlobalizationScope final : EHScopeStack::Cleanup {2066      GlobalizationScope() = default;2067 2068      void Emit(CodeGenFunction &CGF, Flags flags) override {2069        static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime())2070            .emitGenericVarsEpilog(CGF);2071      }2072    };2073    CGF.EHStack.pushCleanup<GlobalizationScope>(NormalAndEHCleanup);2074  }2075}2076 2077Address CGOpenMPRuntimeGPU::getAddressOfLocalVariable(CodeGenFunction &CGF,2078                                                        const VarDecl *VD) {2079  if (VD && VD->hasAttr<OMPAllocateDeclAttr>()) {2080    const auto *A = VD->getAttr<OMPAllocateDeclAttr>();2081    auto AS = LangAS::Default;2082    switch (A->getAllocatorType()) {2083    case OMPAllocateDeclAttr::OMPNullMemAlloc:2084    case OMPAllocateDeclAttr::OMPDefaultMemAlloc:2085    case OMPAllocateDeclAttr::OMPHighBWMemAlloc:2086    case OMPAllocateDeclAttr::OMPLowLatMemAlloc:2087      break;2088    case OMPAllocateDeclAttr::OMPThreadMemAlloc:2089      return Address::invalid();2090    case OMPAllocateDeclAttr::OMPUserDefinedMemAlloc:2091      // TODO: implement aupport for user-defined allocators.2092      return Address::invalid();2093    case OMPAllocateDeclAttr::OMPConstMemAlloc:2094      AS = LangAS::cuda_constant;2095      break;2096    case OMPAllocateDeclAttr::OMPPTeamMemAlloc:2097      AS = LangAS::cuda_shared;2098      break;2099    case OMPAllocateDeclAttr::OMPLargeCapMemAlloc:2100    case OMPAllocateDeclAttr::OMPCGroupMemAlloc:2101      break;2102    }2103    llvm::Type *VarTy = CGF.ConvertTypeForMem(VD->getType());2104    auto *GV = new llvm::GlobalVariable(2105        CGM.getModule(), VarTy, /*isConstant=*/false,2106        llvm::GlobalValue::InternalLinkage, llvm::PoisonValue::get(VarTy),2107        VD->getName(),2108        /*InsertBefore=*/nullptr, llvm::GlobalValue::NotThreadLocal,2109        CGM.getContext().getTargetAddressSpace(AS));2110    CharUnits Align = CGM.getContext().getDeclAlign(VD);2111    GV->setAlignment(Align.getAsAlign());2112    return Address(2113        CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(2114            GV, CGF.Builder.getPtrTy(CGM.getContext().getTargetAddressSpace(2115                    VD->getType().getAddressSpace()))),2116        VarTy, Align);2117  }2118 2119  if (getDataSharingMode() != CGOpenMPRuntimeGPU::DS_Generic)2120    return Address::invalid();2121 2122  VD = VD->getCanonicalDecl();2123  auto I = FunctionGlobalizedDecls.find(CGF.CurFn);2124  if (I == FunctionGlobalizedDecls.end())2125    return Address::invalid();2126  auto VDI = I->getSecond().LocalVarData.find(VD);2127  if (VDI != I->getSecond().LocalVarData.end())2128    return VDI->second.PrivateAddr;2129  if (VD->hasAttrs()) {2130    for (specific_attr_iterator<OMPReferencedVarAttr> IT(VD->attr_begin()),2131         E(VD->attr_end());2132         IT != E; ++IT) {2133      auto VDI = I->getSecond().LocalVarData.find(2134          cast<VarDecl>(cast<DeclRefExpr>(IT->getRef())->getDecl())2135              ->getCanonicalDecl());2136      if (VDI != I->getSecond().LocalVarData.end())2137        return VDI->second.PrivateAddr;2138    }2139  }2140 2141  return Address::invalid();2142}2143 2144void CGOpenMPRuntimeGPU::functionFinished(CodeGenFunction &CGF) {2145  FunctionGlobalizedDecls.erase(CGF.CurFn);2146  CGOpenMPRuntime::functionFinished(CGF);2147}2148 2149void CGOpenMPRuntimeGPU::getDefaultDistScheduleAndChunk(2150    CodeGenFunction &CGF, const OMPLoopDirective &S,2151    OpenMPDistScheduleClauseKind &ScheduleKind,2152    llvm::Value *&Chunk) const {2153  auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());2154  if (getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD) {2155    ScheduleKind = OMPC_DIST_SCHEDULE_static;2156    Chunk = CGF.EmitScalarConversion(2157        RT.getGPUNumThreads(CGF),2158        CGF.getContext().getIntTypeForBitwidth(32, /*Signed=*/0),2159        S.getIterationVariable()->getType(), S.getBeginLoc());2160    return;2161  }2162  CGOpenMPRuntime::getDefaultDistScheduleAndChunk(2163      CGF, S, ScheduleKind, Chunk);2164}2165 2166void CGOpenMPRuntimeGPU::getDefaultScheduleAndChunk(2167    CodeGenFunction &CGF, const OMPLoopDirective &S,2168    OpenMPScheduleClauseKind &ScheduleKind,2169    const Expr *&ChunkExpr) const {2170  ScheduleKind = OMPC_SCHEDULE_static;2171  // Chunk size is 1 in this case.2172  llvm::APInt ChunkSize(32, 1);2173  ChunkExpr = IntegerLiteral::Create(CGF.getContext(), ChunkSize,2174      CGF.getContext().getIntTypeForBitwidth(32, /*Signed=*/0),2175      SourceLocation());2176}2177 2178void CGOpenMPRuntimeGPU::adjustTargetSpecificDataForLambdas(2179    CodeGenFunction &CGF, const OMPExecutableDirective &D) const {2180  assert(isOpenMPTargetExecutionDirective(D.getDirectiveKind()) &&2181         " Expected target-based directive.");2182  const CapturedStmt *CS = D.getCapturedStmt(OMPD_target);2183  for (const CapturedStmt::Capture &C : CS->captures()) {2184    // Capture variables captured by reference in lambdas for target-based2185    // directives.2186    if (!C.capturesVariable())2187      continue;2188    const VarDecl *VD = C.getCapturedVar();2189    const auto *RD = VD->getType()2190                         .getCanonicalType()2191                         .getNonReferenceType()2192                         ->getAsCXXRecordDecl();2193    if (!RD || !RD->isLambda())2194      continue;2195    Address VDAddr = CGF.GetAddrOfLocalVar(VD);2196    LValue VDLVal;2197    if (VD->getType().getCanonicalType()->isReferenceType())2198      VDLVal = CGF.EmitLoadOfReferenceLValue(VDAddr, VD->getType());2199    else2200      VDLVal = CGF.MakeAddrLValue(2201          VDAddr, VD->getType().getCanonicalType().getNonReferenceType());2202    llvm::DenseMap<const ValueDecl *, FieldDecl *> Captures;2203    FieldDecl *ThisCapture = nullptr;2204    RD->getCaptureFields(Captures, ThisCapture);2205    if (ThisCapture && CGF.CapturedStmtInfo->isCXXThisExprCaptured()) {2206      LValue ThisLVal =2207          CGF.EmitLValueForFieldInitialization(VDLVal, ThisCapture);2208      llvm::Value *CXXThis = CGF.LoadCXXThis();2209      CGF.EmitStoreOfScalar(CXXThis, ThisLVal);2210    }2211    for (const LambdaCapture &LC : RD->captures()) {2212      if (LC.getCaptureKind() != LCK_ByRef)2213        continue;2214      const ValueDecl *VD = LC.getCapturedVar();2215      // FIXME: For now VD is always a VarDecl because OpenMP does not support2216      //  capturing structured bindings in lambdas yet.2217      if (!CS->capturesVariable(cast<VarDecl>(VD)))2218        continue;2219      auto It = Captures.find(VD);2220      assert(It != Captures.end() && "Found lambda capture without field.");2221      LValue VarLVal = CGF.EmitLValueForFieldInitialization(VDLVal, It->second);2222      Address VDAddr = CGF.GetAddrOfLocalVar(cast<VarDecl>(VD));2223      if (VD->getType().getCanonicalType()->isReferenceType())2224        VDAddr = CGF.EmitLoadOfReferenceLValue(VDAddr,2225                                               VD->getType().getCanonicalType())2226                     .getAddress();2227      CGF.EmitStoreOfScalar(VDAddr.emitRawPointer(CGF), VarLVal);2228    }2229  }2230}2231 2232bool CGOpenMPRuntimeGPU::hasAllocateAttributeForGlobalVar(const VarDecl *VD,2233                                                            LangAS &AS) {2234  if (!VD || !VD->hasAttr<OMPAllocateDeclAttr>())2235    return false;2236  const auto *A = VD->getAttr<OMPAllocateDeclAttr>();2237  switch(A->getAllocatorType()) {2238  case OMPAllocateDeclAttr::OMPNullMemAlloc:2239  case OMPAllocateDeclAttr::OMPDefaultMemAlloc:2240  // Not supported, fallback to the default mem space.2241  case OMPAllocateDeclAttr::OMPLargeCapMemAlloc:2242  case OMPAllocateDeclAttr::OMPCGroupMemAlloc:2243  case OMPAllocateDeclAttr::OMPHighBWMemAlloc:2244  case OMPAllocateDeclAttr::OMPLowLatMemAlloc:2245  case OMPAllocateDeclAttr::OMPThreadMemAlloc:2246    AS = LangAS::Default;2247    return true;2248  case OMPAllocateDeclAttr::OMPConstMemAlloc:2249    AS = LangAS::cuda_constant;2250    return true;2251  case OMPAllocateDeclAttr::OMPPTeamMemAlloc:2252    AS = LangAS::cuda_shared;2253    return true;2254  case OMPAllocateDeclAttr::OMPUserDefinedMemAlloc:2255    llvm_unreachable("Expected predefined allocator for the variables with the "2256                     "static storage.");2257  }2258  return false;2259}2260 2261// Get current OffloadArch and ignore any unknown values2262static OffloadArch getOffloadArch(CodeGenModule &CGM) {2263  if (!CGM.getTarget().hasFeature("ptx"))2264    return OffloadArch::UNKNOWN;2265  for (const auto &Feature : CGM.getTarget().getTargetOpts().FeatureMap) {2266    if (Feature.getValue()) {2267      OffloadArch Arch = StringToOffloadArch(Feature.getKey());2268      if (Arch != OffloadArch::UNKNOWN)2269        return Arch;2270    }2271  }2272  return OffloadArch::UNKNOWN;2273}2274 2275/// Check to see if target architecture supports unified addressing which is2276/// a restriction for OpenMP requires clause "unified_shared_memory".2277void CGOpenMPRuntimeGPU::processRequiresDirective(const OMPRequiresDecl *D) {2278  for (const OMPClause *Clause : D->clauselists()) {2279    if (Clause->getClauseKind() == OMPC_unified_shared_memory) {2280      OffloadArch Arch = getOffloadArch(CGM);2281      switch (Arch) {2282      case OffloadArch::SM_20:2283      case OffloadArch::SM_21:2284      case OffloadArch::SM_30:2285      case OffloadArch::SM_32_:2286      case OffloadArch::SM_35:2287      case OffloadArch::SM_37:2288      case OffloadArch::SM_50:2289      case OffloadArch::SM_52:2290      case OffloadArch::SM_53: {2291        SmallString<256> Buffer;2292        llvm::raw_svector_ostream Out(Buffer);2293        Out << "Target architecture " << OffloadArchToString(Arch)2294            << " does not support unified addressing";2295        CGM.Error(Clause->getBeginLoc(), Out.str());2296        return;2297      }2298      case OffloadArch::SM_60:2299      case OffloadArch::SM_61:2300      case OffloadArch::SM_62:2301      case OffloadArch::SM_70:2302      case OffloadArch::SM_72:2303      case OffloadArch::SM_75:2304      case OffloadArch::SM_80:2305      case OffloadArch::SM_86:2306      case OffloadArch::SM_87:2307      case OffloadArch::SM_89:2308      case OffloadArch::SM_90:2309      case OffloadArch::SM_90a:2310      case OffloadArch::SM_100:2311      case OffloadArch::SM_100a:2312      case OffloadArch::SM_101:2313      case OffloadArch::SM_101a:2314      case OffloadArch::SM_103:2315      case OffloadArch::SM_103a:2316      case OffloadArch::SM_120:2317      case OffloadArch::SM_120a:2318      case OffloadArch::SM_121:2319      case OffloadArch::SM_121a:2320      case OffloadArch::GFX600:2321      case OffloadArch::GFX601:2322      case OffloadArch::GFX602:2323      case OffloadArch::GFX700:2324      case OffloadArch::GFX701:2325      case OffloadArch::GFX702:2326      case OffloadArch::GFX703:2327      case OffloadArch::GFX704:2328      case OffloadArch::GFX705:2329      case OffloadArch::GFX801:2330      case OffloadArch::GFX802:2331      case OffloadArch::GFX803:2332      case OffloadArch::GFX805:2333      case OffloadArch::GFX810:2334      case OffloadArch::GFX9_GENERIC:2335      case OffloadArch::GFX900:2336      case OffloadArch::GFX902:2337      case OffloadArch::GFX904:2338      case OffloadArch::GFX906:2339      case OffloadArch::GFX908:2340      case OffloadArch::GFX909:2341      case OffloadArch::GFX90a:2342      case OffloadArch::GFX90c:2343      case OffloadArch::GFX9_4_GENERIC:2344      case OffloadArch::GFX942:2345      case OffloadArch::GFX950:2346      case OffloadArch::GFX10_1_GENERIC:2347      case OffloadArch::GFX1010:2348      case OffloadArch::GFX1011:2349      case OffloadArch::GFX1012:2350      case OffloadArch::GFX1013:2351      case OffloadArch::GFX10_3_GENERIC:2352      case OffloadArch::GFX1030:2353      case OffloadArch::GFX1031:2354      case OffloadArch::GFX1032:2355      case OffloadArch::GFX1033:2356      case OffloadArch::GFX1034:2357      case OffloadArch::GFX1035:2358      case OffloadArch::GFX1036:2359      case OffloadArch::GFX11_GENERIC:2360      case OffloadArch::GFX1100:2361      case OffloadArch::GFX1101:2362      case OffloadArch::GFX1102:2363      case OffloadArch::GFX1103:2364      case OffloadArch::GFX1150:2365      case OffloadArch::GFX1151:2366      case OffloadArch::GFX1152:2367      case OffloadArch::GFX1153:2368      case OffloadArch::GFX12_GENERIC:2369      case OffloadArch::GFX1200:2370      case OffloadArch::GFX1201:2371      case OffloadArch::GFX1250:2372      case OffloadArch::GFX1251:2373      case OffloadArch::AMDGCNSPIRV:2374      case OffloadArch::Generic:2375      case OffloadArch::GRANITERAPIDS:2376      case OffloadArch::BMG_G21:2377      case OffloadArch::UNUSED:2378      case OffloadArch::UNKNOWN:2379        break;2380      case OffloadArch::LAST:2381        llvm_unreachable("Unexpected GPU arch.");2382      }2383    }2384  }2385  CGOpenMPRuntime::processRequiresDirective(D);2386}2387 2388llvm::Value *CGOpenMPRuntimeGPU::getGPUNumThreads(CodeGenFunction &CGF) {2389  CGBuilderTy &Bld = CGF.Builder;2390  llvm::Module *M = &CGF.CGM.getModule();2391  const char *LocSize = "__kmpc_get_hardware_num_threads_in_block";2392  llvm::Function *F = M->getFunction(LocSize);2393  if (!F) {2394    F = llvm::Function::Create(llvm::FunctionType::get(CGF.Int32Ty, {}, false),2395                               llvm::GlobalVariable::ExternalLinkage, LocSize,2396                               &CGF.CGM.getModule());2397  }2398  return Bld.CreateCall(F, {}, "nvptx_num_threads");2399}2400 2401llvm::Value *CGOpenMPRuntimeGPU::getGPUThreadID(CodeGenFunction &CGF) {2402  ArrayRef<llvm::Value *> Args{};2403  return CGF.EmitRuntimeCall(2404      OMPBuilder.getOrCreateRuntimeFunction(2405          CGM.getModule(), OMPRTL___kmpc_get_hardware_thread_id_in_block),2406      Args);2407}2408