TheRock/compiler/amd-llvm/clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp

//===---- CGOpenMPRuntimeGPU.cpp - Interface to OpenMP GPU Runtimes ----===//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
//
// This provides a generalized class for OpenMP runtime code generation
// specialized by GPU targets NVPTX, AMDGCN and SPIR-V.
//
//===----------------------------------------------------------------------===//

#include "CGOpenMPRuntimeGPU.h"
#include "CGDebugInfo.h"
#include "CodeGenFunction.h"
#include "clang/AST/Attr.h"
#include "clang/AST/DeclOpenMP.h"
#include "clang/AST/OpenMPClause.h"
#include "clang/AST/StmtOpenMP.h"
#include "clang/AST/StmtVisitor.h"
#include "clang/Basic/Cuda.h"
#include "llvm/ADT/SmallPtrSet.h"
#include "llvm/ADT/StringExtras.h"
#include "llvm/Frontend/OpenMP/OMPDeviceConstants.h"
#include "llvm/Frontend/OpenMP/OMPGridValues.h"
#include "llvm/IR/IntrinsicsAMDGPU.h"
#include "llvm/IR/Metadata.h"
#include "llvm/Support/Debug.h"

using namespace clang;
using namespace CodeGen;
using namespace llvm::omp;

namespace {
/// Pre(post)-action for different OpenMP constructs specialized for NVPTX.
class NVPTXActionTy final : public PrePostActionTy {
  llvm::FunctionCallee EnterCallee = nullptr;
  ArrayRef<llvm::Value *> EnterArgs;
  llvm::FunctionCallee ExitCallee = nullptr;
  ArrayRef<llvm::Value *> ExitArgs;
  bool Conditional = false;
  llvm::BasicBlock *ContBlock = nullptr;

public:
  NVPTXActionTy(llvm::FunctionCallee EnterCallee,
                ArrayRef<llvm::Value *> EnterArgs,
                llvm::FunctionCallee ExitCallee,
                ArrayRef<llvm::Value *> ExitArgs, bool Conditional = false)
      : EnterCallee(EnterCallee), EnterArgs(EnterArgs), ExitCallee(ExitCallee),
        ExitArgs(ExitArgs), Conditional(Conditional) {}
  void Enter(CodeGenFunction &CGF) override {
    llvm::Value *EnterRes = CGF.EmitRuntimeCall(EnterCallee, EnterArgs);
    if (Conditional) {
      llvm::Value *CallBool = CGF.Builder.CreateIsNotNull(EnterRes);
      auto *ThenBlock = CGF.createBasicBlock("omp_if.then");
      ContBlock = CGF.createBasicBlock("omp_if.end");
      // Generate the branch (If-stmt)
      CGF.Builder.CreateCondBr(CallBool, ThenBlock, ContBlock);
      CGF.EmitBlock(ThenBlock);
    }
  }
  void Done(CodeGenFunction &CGF) {
    // Emit the rest of blocks/branches
    CGF.EmitBranch(ContBlock);
    CGF.EmitBlock(ContBlock, true);
  }
  void Exit(CodeGenFunction &CGF) override {
    CGF.EmitRuntimeCall(ExitCallee, ExitArgs);
  }
};

/// A class to track the execution mode when codegening directives within
/// a target region. The appropriate mode (SPMD|NON-SPMD) is set on entry
/// to the target region and used by containing directives such as 'parallel'
/// to emit optimized code.
class ExecutionRuntimeModesRAII {
private:
  CGOpenMPRuntimeGPU::ExecutionMode SavedExecMode =
      CGOpenMPRuntimeGPU::EM_Unknown;
  CGOpenMPRuntimeGPU::ExecutionMode &ExecMode;

public:
  ExecutionRuntimeModesRAII(CGOpenMPRuntimeGPU::ExecutionMode &ExecMode,
                            CGOpenMPRuntimeGPU::ExecutionMode EntryMode)
      : ExecMode(ExecMode) {
    SavedExecMode = ExecMode;
    ExecMode = EntryMode;
  }
  ~ExecutionRuntimeModesRAII() { ExecMode = SavedExecMode; }
};

static const ValueDecl *getPrivateItem(const Expr *RefExpr) {
  RefExpr = RefExpr->IgnoreParens();
  if (const auto *ASE = dyn_cast<ArraySubscriptExpr>(RefExpr)) {
    const Expr *Base = ASE->getBase()->IgnoreParenImpCasts();
    while (const auto *TempASE = dyn_cast<ArraySubscriptExpr>(Base))
      Base = TempASE->getBase()->IgnoreParenImpCasts();
    RefExpr = Base;
  } else if (auto *OASE = dyn_cast<ArraySectionExpr>(RefExpr)) {
    const Expr *Base = OASE->getBase()->IgnoreParenImpCasts();
    while (const auto *TempOASE = dyn_cast<ArraySectionExpr>(Base))
      Base = TempOASE->getBase()->IgnoreParenImpCasts();
    while (const auto *TempASE = dyn_cast<ArraySubscriptExpr>(Base))
      Base = TempASE->getBase()->IgnoreParenImpCasts();
    RefExpr = Base;
  }
  RefExpr = RefExpr->IgnoreParenImpCasts();
  if (const auto *DE = dyn_cast<DeclRefExpr>(RefExpr))
    return cast<ValueDecl>(DE->getDecl()->getCanonicalDecl());
  const auto *ME = cast<MemberExpr>(RefExpr);
  return cast<ValueDecl>(ME->getMemberDecl()->getCanonicalDecl());
}

static RecordDecl *buildRecordForGlobalizedVars(
    ASTContext &C, ArrayRef<const ValueDecl *> EscapedDecls,
    ArrayRef<const ValueDecl *> EscapedDeclsForTeams,
    llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>
        &MappedDeclsFields,
    int BufSize) {
  using VarsDataTy = std::pair<CharUnits /*Align*/, const ValueDecl *>;
  if (EscapedDecls.empty() && EscapedDeclsForTeams.empty())
    return nullptr;
  SmallVector<VarsDataTy, 4> GlobalizedVars;
  for (const ValueDecl *D : EscapedDecls)
    GlobalizedVars.emplace_back(C.getDeclAlign(D), D);
  for (const ValueDecl *D : EscapedDeclsForTeams)
    GlobalizedVars.emplace_back(C.getDeclAlign(D), D);

  // Build struct _globalized_locals_ty {
  //         /*  globalized vars  */[WarSize] align (decl_align)
  //         /*  globalized vars  */ for EscapedDeclsForTeams
  //       };
  RecordDecl *GlobalizedRD = C.buildImplicitRecord("_globalized_locals_ty");
  GlobalizedRD->startDefinition();
  llvm::SmallPtrSet<const ValueDecl *, 16> SingleEscaped(llvm::from_range,
                                                         EscapedDeclsForTeams);
  for (const auto &Pair : GlobalizedVars) {
    const ValueDecl *VD = Pair.second;
    QualType Type = VD->getType();
    if (Type->isLValueReferenceType())
      Type = C.getPointerType(Type.getNonReferenceType());
    else
      Type = Type.getNonReferenceType();
    SourceLocation Loc = VD->getLocation();
    FieldDecl *Field;
    if (SingleEscaped.count(VD)) {
      Field = FieldDecl::Create(
          C, GlobalizedRD, Loc, Loc, VD->getIdentifier(), Type,
          C.getTrivialTypeSourceInfo(Type, SourceLocation()),
          /*BW=*/nullptr, /*Mutable=*/false,
          /*InitStyle=*/ICIS_NoInit);
      Field->setAccess(AS_public);
      if (VD->hasAttrs()) {
        for (specific_attr_iterator<AlignedAttr> I(VD->getAttrs().begin()),
             E(VD->getAttrs().end());
             I != E; ++I)
          Field->addAttr(*I);
      }
    } else {
      if (BufSize > 1) {
        llvm::APInt ArraySize(32, BufSize);
        Type = C.getConstantArrayType(Type, ArraySize, nullptr,
                                      ArraySizeModifier::Normal, 0);
      }
      Field = FieldDecl::Create(
          C, GlobalizedRD, Loc, Loc, VD->getIdentifier(), Type,
          C.getTrivialTypeSourceInfo(Type, SourceLocation()),
          /*BW=*/nullptr, /*Mutable=*/false,
          /*InitStyle=*/ICIS_NoInit);
      Field->setAccess(AS_public);
      llvm::APInt Align(32, Pair.first.getQuantity());
      Field->addAttr(AlignedAttr::CreateImplicit(
          C, /*IsAlignmentExpr=*/true,
          IntegerLiteral::Create(C, Align,
                                 C.getIntTypeForBitwidth(32, /*Signed=*/0),
                                 SourceLocation()),
          {}, AlignedAttr::GNU_aligned));
    }
    GlobalizedRD->addDecl(Field);
    MappedDeclsFields.try_emplace(VD, Field);
  }
  GlobalizedRD->completeDefinition();
  return GlobalizedRD;
}

/// Get the list of variables that can escape their declaration context.
class CheckVarsEscapingDeclContext final
    : public ConstStmtVisitor<CheckVarsEscapingDeclContext> {
  CodeGenFunction &CGF;
  llvm::SetVector<const ValueDecl *> EscapedDecls;
  llvm::SetVector<const ValueDecl *> EscapedVariableLengthDecls;
  llvm::SetVector<const ValueDecl *> DelayedVariableLengthDecls;
  llvm::SmallPtrSet<const Decl *, 4> EscapedParameters;
  RecordDecl *GlobalizedRD = nullptr;
  llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *> MappedDeclsFields;
  bool AllEscaped = false;
  bool IsForCombinedParallelRegion = false;

  void markAsEscaped(const ValueDecl *VD) {
    // Do not globalize declare target variables.
    if (!isa<VarDecl>(VD) ||
        OMPDeclareTargetDeclAttr::isDeclareTargetDeclaration(VD))
      return;
    VD = cast<ValueDecl>(VD->getCanonicalDecl());
    // Use user-specified allocation.
    if (VD->hasAttrs() && VD->hasAttr<OMPAllocateDeclAttr>())
      return;
    // Variables captured by value must be globalized.
    bool IsCaptured = false;
    if (auto *CSI = CGF.CapturedStmtInfo) {
      if (const FieldDecl *FD = CSI->lookup(cast<VarDecl>(VD))) {
        // Check if need to capture the variable that was already captured by
        // value in the outer region.
        IsCaptured = true;
        if (!IsForCombinedParallelRegion) {
          if (!FD->hasAttrs())
            return;
          const auto *Attr = FD->getAttr<OMPCaptureKindAttr>();
          if (!Attr)
            return;
          if (((Attr->getCaptureKind() != OMPC_map) &&
               !isOpenMPPrivate(Attr->getCaptureKind())) ||
              ((Attr->getCaptureKind() == OMPC_map) &&
               !FD->getType()->isAnyPointerType()))
            return;
        }
        if (!FD->getType()->isReferenceType()) {
          assert(!VD->getType()->isVariablyModifiedType() &&
                 "Parameter captured by value with variably modified type");
          EscapedParameters.insert(VD);
        } else if (!IsForCombinedParallelRegion) {
          return;
        }
      }
    }
    if ((!CGF.CapturedStmtInfo ||
         (IsForCombinedParallelRegion && CGF.CapturedStmtInfo)) &&
        VD->getType()->isReferenceType())
      // Do not globalize variables with reference type.
      return;
    if (VD->getType()->isVariablyModifiedType()) {
      // If not captured at the target region level then mark the escaped
      // variable as delayed.
      if (IsCaptured)
        EscapedVariableLengthDecls.insert(VD);
      else
        DelayedVariableLengthDecls.insert(VD);
    } else
      EscapedDecls.insert(VD);
  }

  void VisitValueDecl(const ValueDecl *VD) {
    if (VD->getType()->isLValueReferenceType())
      markAsEscaped(VD);
    if (const auto *VarD = dyn_cast<VarDecl>(VD)) {
      if (!isa<ParmVarDecl>(VarD) && VarD->hasInit()) {
        const bool SavedAllEscaped = AllEscaped;
        AllEscaped = VD->getType()->isLValueReferenceType();
        Visit(VarD->getInit());
        AllEscaped = SavedAllEscaped;
      }
    }
  }
  void VisitOpenMPCapturedStmt(const CapturedStmt *S,
                               ArrayRef<OMPClause *> Clauses,
                               bool IsCombinedParallelRegion) {
    if (!S)
      return;
    for (const CapturedStmt::Capture &C : S->captures()) {
      if (C.capturesVariable() && !C.capturesVariableByCopy()) {
        const ValueDecl *VD = C.getCapturedVar();
        bool SavedIsForCombinedParallelRegion = IsForCombinedParallelRegion;
        if (IsCombinedParallelRegion) {
          // Check if the variable is privatized in the combined construct and
          // those private copies must be shared in the inner parallel
          // directive.
          IsForCombinedParallelRegion = false;
          for (const OMPClause *C : Clauses) {
            if (!isOpenMPPrivate(C->getClauseKind()) ||
                C->getClauseKind() == OMPC_reduction ||
                C->getClauseKind() == OMPC_linear ||
                C->getClauseKind() == OMPC_private)
              continue;
            ArrayRef<const Expr *> Vars;
            if (const auto *PC = dyn_cast<OMPFirstprivateClause>(C))
              Vars = PC->getVarRefs();
            else if (const auto *PC = dyn_cast<OMPLastprivateClause>(C))
              Vars = PC->getVarRefs();
            else
              llvm_unreachable("Unexpected clause.");
            for (const auto *E : Vars) {
              const Decl *D =
                  cast<DeclRefExpr>(E)->getDecl()->getCanonicalDecl();
              if (D == VD->getCanonicalDecl()) {
                IsForCombinedParallelRegion = true;
                break;
              }
            }
            if (IsForCombinedParallelRegion)
              break;
          }
        }
        markAsEscaped(VD);
        if (isa<OMPCapturedExprDecl>(VD))
          VisitValueDecl(VD);
        IsForCombinedParallelRegion = SavedIsForCombinedParallelRegion;
      }
    }
  }

  void buildRecordForGlobalizedVars(bool IsInTTDRegion) {
    assert(!GlobalizedRD &&
           "Record for globalized variables is built already.");
    ArrayRef<const ValueDecl *> EscapedDeclsForParallel, EscapedDeclsForTeams;
    unsigned WarpSize = CGF.getTarget().getGridValue().GV_Warp_Size;
    if (IsInTTDRegion)
      EscapedDeclsForTeams = EscapedDecls.getArrayRef();
    else
      EscapedDeclsForParallel = EscapedDecls.getArrayRef();
    GlobalizedRD = ::buildRecordForGlobalizedVars(
        CGF.getContext(), EscapedDeclsForParallel, EscapedDeclsForTeams,
        MappedDeclsFields, WarpSize);
  }

public:
  CheckVarsEscapingDeclContext(CodeGenFunction &CGF,
                               ArrayRef<const ValueDecl *> TeamsReductions)
      : CGF(CGF), EscapedDecls(llvm::from_range, TeamsReductions) {}
  ~CheckVarsEscapingDeclContext() = default;
  void VisitDeclStmt(const DeclStmt *S) {
    if (!S)
      return;
    for (const Decl *D : S->decls())
      if (const auto *VD = dyn_cast_or_null<ValueDecl>(D))
        VisitValueDecl(VD);
  }
  void VisitOMPExecutableDirective(const OMPExecutableDirective *D) {
    if (!D)
      return;
    if (!D->hasAssociatedStmt())
      return;
    if (const auto *S =
            dyn_cast_or_null<CapturedStmt>(D->getAssociatedStmt())) {
      // Do not analyze directives that do not actually require capturing,
      // like `omp for` or `omp simd` directives.
      llvm::SmallVector<OpenMPDirectiveKind, 4> CaptureRegions;
      getOpenMPCaptureRegions(CaptureRegions, D->getDirectiveKind());
      if (CaptureRegions.size() == 1 && CaptureRegions.back() == OMPD_unknown) {
        VisitStmt(S->getCapturedStmt());
        return;
      }
      VisitOpenMPCapturedStmt(
          S, D->clauses(),
          CaptureRegions.back() == OMPD_parallel &&
              isOpenMPDistributeDirective(D->getDirectiveKind()));
    }
  }
  void VisitCapturedStmt(const CapturedStmt *S) {
    if (!S)
      return;
    for (const CapturedStmt::Capture &C : S->captures()) {
      if (C.capturesVariable() && !C.capturesVariableByCopy()) {
        const ValueDecl *VD = C.getCapturedVar();
        markAsEscaped(VD);
        if (isa<OMPCapturedExprDecl>(VD))
          VisitValueDecl(VD);
      }
    }
  }
  void VisitLambdaExpr(const LambdaExpr *E) {
    if (!E)
      return;
    for (const LambdaCapture &C : E->captures()) {
      if (C.capturesVariable()) {
        if (C.getCaptureKind() == LCK_ByRef) {
          const ValueDecl *VD = C.getCapturedVar();
          markAsEscaped(VD);
          if (E->isInitCapture(&C) || isa<OMPCapturedExprDecl>(VD))
            VisitValueDecl(VD);
        }
      }
    }
  }
  void VisitBlockExpr(const BlockExpr *E) {
    if (!E)
      return;
    for (const BlockDecl::Capture &C : E->getBlockDecl()->captures()) {
      if (C.isByRef()) {
        const VarDecl *VD = C.getVariable();
        markAsEscaped(VD);
        if (isa<OMPCapturedExprDecl>(VD) || VD->isInitCapture())
          VisitValueDecl(VD);
      }
    }
  }
  void VisitCallExpr(const CallExpr *E) {
    if (!E)
      return;
    for (const Expr *Arg : E->arguments()) {
      if (!Arg)
        continue;
      if (Arg->isLValue()) {
        const bool SavedAllEscaped = AllEscaped;
        AllEscaped = true;
        Visit(Arg);
        AllEscaped = SavedAllEscaped;
      } else {
        Visit(Arg);
      }
    }
    Visit(E->getCallee());
  }
  void VisitDeclRefExpr(const DeclRefExpr *E) {
    if (!E)
      return;
    const ValueDecl *VD = E->getDecl();
    if (AllEscaped)
      markAsEscaped(VD);
    if (isa<OMPCapturedExprDecl>(VD))
      VisitValueDecl(VD);
    else if (VD->isInitCapture())
      VisitValueDecl(VD);
  }
  void VisitUnaryOperator(const UnaryOperator *E) {
    if (!E)
      return;
    if (E->getOpcode() == UO_AddrOf) {
      const bool SavedAllEscaped = AllEscaped;
      AllEscaped = true;
      Visit(E->getSubExpr());
      AllEscaped = SavedAllEscaped;
    } else {
      Visit(E->getSubExpr());
    }
  }
  void VisitImplicitCastExpr(const ImplicitCastExpr *E) {
    if (!E)
      return;
    if (E->getCastKind() == CK_ArrayToPointerDecay) {
      const bool SavedAllEscaped = AllEscaped;
      AllEscaped = true;
      Visit(E->getSubExpr());
      AllEscaped = SavedAllEscaped;
    } else {
      Visit(E->getSubExpr());
    }
  }
  void VisitExpr(const Expr *E) {
    if (!E)
      return;
    bool SavedAllEscaped = AllEscaped;
    if (!E->isLValue())
      AllEscaped = false;
    for (const Stmt *Child : E->children())
      if (Child)
        Visit(Child);
    AllEscaped = SavedAllEscaped;
  }
  void VisitStmt(const Stmt *S) {
    if (!S)
      return;
    for (const Stmt *Child : S->children())
      if (Child)
        Visit(Child);
  }

  /// Returns the record that handles all the escaped local variables and used
  /// instead of their original storage.
  const RecordDecl *getGlobalizedRecord(bool IsInTTDRegion) {
    if (!GlobalizedRD)
      buildRecordForGlobalizedVars(IsInTTDRegion);
    return GlobalizedRD;
  }

  /// Returns the field in the globalized record for the escaped variable.
  const FieldDecl *getFieldForGlobalizedVar(const ValueDecl *VD) const {
    assert(GlobalizedRD &&
           "Record for globalized variables must be generated already.");
    return MappedDeclsFields.lookup(VD);
  }

  /// Returns the list of the escaped local variables/parameters.
  ArrayRef<const ValueDecl *> getEscapedDecls() const {
    return EscapedDecls.getArrayRef();
  }

  /// Checks if the escaped local variable is actually a parameter passed by
  /// value.
  const llvm::SmallPtrSetImpl<const Decl *> &getEscapedParameters() const {
    return EscapedParameters;
  }

  /// Returns the list of the escaped variables with the variably modified
  /// types.
  ArrayRef<const ValueDecl *> getEscapedVariableLengthDecls() const {
    return EscapedVariableLengthDecls.getArrayRef();
  }

  /// Returns the list of the delayed variables with the variably modified
  /// types.
  ArrayRef<const ValueDecl *> getDelayedVariableLengthDecls() const {
    return DelayedVariableLengthDecls.getArrayRef();
  }
};
} // anonymous namespace

CGOpenMPRuntimeGPU::ExecutionMode
CGOpenMPRuntimeGPU::getExecutionMode() const {
  return CurrentExecutionMode;
}

CGOpenMPRuntimeGPU::DataSharingMode
CGOpenMPRuntimeGPU::getDataSharingMode() const {
  return CurrentDataSharingMode;
}

/// Check for inner (nested) SPMD construct, if any
static bool hasNestedSPMDDirective(ASTContext &Ctx,
                                   const OMPExecutableDirective &D) {
  const auto *CS = D.getInnermostCapturedStmt();
  const auto *Body =
      CS->getCapturedStmt()->IgnoreContainers(/*IgnoreCaptured=*/true);
  const Stmt *ChildStmt = CGOpenMPRuntime::getSingleCompoundChild(Ctx, Body);

  if (const auto *NestedDir =
          dyn_cast_or_null<OMPExecutableDirective>(ChildStmt)) {
    OpenMPDirectiveKind DKind = NestedDir->getDirectiveKind();
    switch (D.getDirectiveKind()) {
    case OMPD_target:
      if (isOpenMPParallelDirective(DKind))
        return true;
      if (DKind == OMPD_teams) {
        Body = NestedDir->getInnermostCapturedStmt()->IgnoreContainers(
            /*IgnoreCaptured=*/true);
        if (!Body)
          return false;
        ChildStmt = CGOpenMPRuntime::getSingleCompoundChild(Ctx, Body);
        if (const auto *NND =
                dyn_cast_or_null<OMPExecutableDirective>(ChildStmt)) {
          DKind = NND->getDirectiveKind();
          if (isOpenMPParallelDirective(DKind))
            return true;
        }
      }
      return false;
    case OMPD_target_teams:
      return isOpenMPParallelDirective(DKind) || (DKind == OMPD_loop);
    case OMPD_target_simd:
    case OMPD_target_parallel:
    case OMPD_target_parallel_for:
    case OMPD_target_parallel_for_simd:
    case OMPD_target_teams_distribute:
    case OMPD_target_teams_distribute_simd:
    case OMPD_target_teams_distribute_parallel_for:
    case OMPD_target_teams_distribute_parallel_for_simd:
    case OMPD_parallel:
    case OMPD_for:
    case OMPD_parallel_for:
    case OMPD_parallel_master:
    case OMPD_parallel_sections:
    case OMPD_for_simd:
    case OMPD_parallel_for_simd:
    case OMPD_cancel:
    case OMPD_cancellation_point:
    case OMPD_ordered:
    case OMPD_threadprivate:
    case OMPD_allocate:
    case OMPD_task:
    case OMPD_simd:
    case OMPD_sections:
    case OMPD_section:
    case OMPD_single:
    case OMPD_master:
    case OMPD_critical:
    case OMPD_taskyield:
    case OMPD_barrier:
    case OMPD_taskwait:
    case OMPD_taskgroup:
    case OMPD_atomic:
    case OMPD_flush:
    case OMPD_depobj:
    case OMPD_scan:
    case OMPD_teams:
    case OMPD_target_data:
    case OMPD_target_exit_data:
    case OMPD_target_enter_data:
    case OMPD_distribute:
    case OMPD_distribute_simd:
    case OMPD_distribute_parallel_for:
    case OMPD_distribute_parallel_for_simd:
    case OMPD_teams_distribute:
    case OMPD_teams_distribute_simd:
    case OMPD_teams_distribute_parallel_for:
    case OMPD_teams_distribute_parallel_for_simd:
    case OMPD_target_update:
    case OMPD_declare_simd:
    case OMPD_declare_variant:
    case OMPD_begin_declare_variant:
    case OMPD_end_declare_variant:
    case OMPD_declare_target:
    case OMPD_end_declare_target:
    case OMPD_declare_reduction:
    case OMPD_declare_mapper:
    case OMPD_taskloop:
    case OMPD_taskloop_simd:
    case OMPD_master_taskloop:
    case OMPD_master_taskloop_simd:
    case OMPD_parallel_master_taskloop:
    case OMPD_parallel_master_taskloop_simd:
    case OMPD_requires:
    case OMPD_unknown:
    default:
      llvm_unreachable("Unexpected directive.");
    }
  }

  return false;
}

static bool supportsSPMDExecutionMode(CodeGenModule &CGM,
                                      const OMPExecutableDirective &D) {
  ASTContext &Ctx = CGM.getContext();
  OpenMPDirectiveKind DirectiveKind = D.getDirectiveKind();
  switch (DirectiveKind) {
  case OMPD_target:
  case OMPD_target_teams:
    return hasNestedSPMDDirective(Ctx, D);
  case OMPD_target_parallel_loop:
  case OMPD_target_parallel:
  case OMPD_target_parallel_for:
  case OMPD_target_parallel_for_simd:
  case OMPD_target_teams_distribute_parallel_for:
  case OMPD_target_teams_distribute_parallel_for_simd:
  case OMPD_target_simd:
  case OMPD_target_teams_distribute_simd:
    return true;
  case OMPD_target_teams_distribute:
    return false;
  case OMPD_target_teams_loop:
    // Whether this is true or not depends on how the directive will
    // eventually be emitted.
    if (auto *TTLD = dyn_cast<OMPTargetTeamsGenericLoopDirective>(&D))
      return TTLD->canBeParallelFor();
    return false;
  case OMPD_parallel:
  case OMPD_for:
  case OMPD_parallel_for:
  case OMPD_parallel_master:
  case OMPD_parallel_sections:
  case OMPD_for_simd:
  case OMPD_parallel_for_simd:
  case OMPD_cancel:
  case OMPD_cancellation_point:
  case OMPD_ordered:
  case OMPD_threadprivate:
  case OMPD_allocate:
  case OMPD_task:
  case OMPD_simd:
  case OMPD_sections:
  case OMPD_section:
  case OMPD_single:
  case OMPD_master:
  case OMPD_critical:
  case OMPD_taskyield:
  case OMPD_barrier:
  case OMPD_taskwait:
  case OMPD_taskgroup:
  case OMPD_atomic:
  case OMPD_flush:
  case OMPD_depobj:
  case OMPD_scan:
  case OMPD_teams:
  case OMPD_target_data:
  case OMPD_target_exit_data:
  case OMPD_target_enter_data:
  case OMPD_distribute:
  case OMPD_distribute_simd:
  case OMPD_distribute_parallel_for:
  case OMPD_distribute_parallel_for_simd:
  case OMPD_teams_distribute:
  case OMPD_teams_distribute_simd:
  case OMPD_teams_distribute_parallel_for:
  case OMPD_teams_distribute_parallel_for_simd:
  case OMPD_target_update:
  case OMPD_declare_simd:
  case OMPD_declare_variant:
  case OMPD_begin_declare_variant:
  case OMPD_end_declare_variant:
  case OMPD_declare_target:
  case OMPD_end_declare_target:
  case OMPD_declare_reduction:
  case OMPD_declare_mapper:
  case OMPD_taskloop:
  case OMPD_taskloop_simd:
  case OMPD_master_taskloop:
  case OMPD_master_taskloop_simd:
  case OMPD_parallel_master_taskloop:
  case OMPD_parallel_master_taskloop_simd:
  case OMPD_requires:
  case OMPD_unknown:
  default:
    break;
  }
  llvm_unreachable(
      "Unknown programming model for OpenMP directive on NVPTX target.");
}

// Create a unique global variable to indicate the flat-work-group-size
// for this region. Values are [1..1024].
static void setPropertyWorkGroupSize(CodeGenModule &CGM, StringRef Name,
                                     int WGSize) {
  auto *GVMode = new llvm::GlobalVariable(
      CGM.getModule(), CGM.Int16Ty,
      /*isConstant=*/true, llvm::GlobalValue::WeakAnyLinkage,
      llvm::ConstantInt::get(CGM.Int16Ty, WGSize), Twine(Name, "_wg_size"));

  CGM.addCompilerUsedGlobal(GVMode);
}

// Create a unique global variable to indicate if the kernel is multi-device.
static void setMultiDeviceStatus(CodeGenModule &CGM, StringRef Name,
                                 int IsMultiDevice) {
  auto *GVMode = new llvm::GlobalVariable(
      CGM.getModule(), CGM.Int8Ty,
      /*isConstant=*/true, llvm::GlobalValue::WeakAnyLinkage,
      llvm::ConstantInt::get(CGM.Int8Ty, IsMultiDevice),
      Twine(Name, "_multi_device"));

  CGM.addCompilerUsedGlobal(GVMode);
}

// Compute the correct number of threads in a team
// to accommodate for a master thread.
// Keep aligned with amdgpu plugin code located in function getLaunchVals
static int ComputeGenericWorkgroupSize(CodeGenModule &CGM, int WorkgroupSize) {
  assert(WorkgroupSize >= 0);
  int MaxWorkGroupSz = CGM.getTarget().getGridValue().GV_Max_WG_Size;
  int WorkgroupSizeWithMaster = -1;

  // Add master thread in additional warp for GENERIC mode
  // Only one additional thread is started, not an entire warp

  if (WorkgroupSize >= MaxWorkGroupSz)
    // Do not exceed max number of threads: sacrifice last warp for
    // the thread master
    WorkgroupSizeWithMaster =
        MaxWorkGroupSz - CGM.getTarget().getGridValue().GV_Warp_Size + 1;
  else if ((unsigned int)WorkgroupSize <
           CGM.getTarget().getGridValue().GV_Warp_Size)
    // Cap threadsPerGroup at WarpSize level as we need a master
    WorkgroupSizeWithMaster = CGM.getTarget().getGridValue().GV_Warp_Size + 1;
  else
    WorkgroupSizeWithMaster =
        CGM.getTarget().getGridValue().GV_Warp_Size *
            (WorkgroupSize / CGM.getTarget().getGridValue().GV_Warp_Size) +
        1;
  return WorkgroupSizeWithMaster;
}

void CGOpenMPRuntimeGPU::GenerateMetaData(CodeGenModule &CGM,
                                          const OMPExecutableDirective &D,
                                          llvm::Function *&OutlinedFn,
                                          bool IsGeneric) {
  if (!CGM.getTriple().isAMDGCN())
    return;

  int FlatAttr = 0;
  bool flatAttrEmitted = false;
  unsigned compileTimeThreadLimit =
      CGM.getTarget().getGridValue().GV_Default_WG_Size;
  bool isXteamRedKernel = CGM.isXteamRedKernel(D);
  bool isBigJumpLoopKernel = CGM.isBigJumpLoopKernel(D);
  bool isNoLoopKernel = CGM.isNoLoopKernel(D);
  // If constant ThreadLimit(), set reqd_work_group_size metadata
  if (isOpenMPTeamsDirective(D.getDirectiveKind()) ||
      isOpenMPParallelDirective(D.getDirectiveKind()) || isXteamRedKernel ||
      isBigJumpLoopKernel || isNoLoopKernel) {
    // Call the work group size calculation based on kernel type.
    if (isXteamRedKernel)
      compileTimeThreadLimit = CGM.getXteamRedBlockSize(D);
    else if (isBigJumpLoopKernel)
      compileTimeThreadLimit = CGM.getBigJumpLoopBlockSize(D);
    else if (isNoLoopKernel)
      compileTimeThreadLimit = CGM.getNoLoopBlockSize(D);
    else
      compileTimeThreadLimit = CGM.getWorkGroupSizeSPMDHelper(D);

    // Add kernel metadata if ThreadLimit Clause is compile time constant > 0
    if (compileTimeThreadLimit > 0) {
      if (IsGeneric)
        compileTimeThreadLimit =
            ComputeGenericWorkgroupSize(CGM, compileTimeThreadLimit);
      FlatAttr = compileTimeThreadLimit;
      OutlinedFn->addFnAttr("amdgpu-flat-work-group-size",
                            "1," + llvm::utostr(compileTimeThreadLimit));
      flatAttrEmitted = true;
    } // end   > 0
  }   // end of amdgcn teams or parallel directive

  // emit amdgpu-flat-work-group-size if not emitted already.
  if (!flatAttrEmitted) {
    // When outermost construct does not have teams or parallel
    // workgroup size is still based on mode
    int GenericModeWorkgroupSize = compileTimeThreadLimit;
    if (IsGeneric)
      GenericModeWorkgroupSize =
          ComputeGenericWorkgroupSize(CGM, compileTimeThreadLimit);
    FlatAttr = GenericModeWorkgroupSize;
    OutlinedFn->addFnAttr("amdgpu-flat-work-group-size",
                          "1," + llvm::utostr(GenericModeWorkgroupSize));
  }
  // Emit a kernel descriptor for runtime.
  setPropertyWorkGroupSize(CGM, OutlinedFn->getName(), FlatAttr);

  // Emit multi-device flag for this kernel.
  setMultiDeviceStatus(CGM, OutlinedFn->getName(), CGM.isMultiDeviceKernel(D));
}

void CGOpenMPRuntimeGPU::emitNonSPMDKernel(const OMPExecutableDirective &D,
                                             StringRef ParentName,
                                             llvm::Function *&OutlinedFn,
                                             llvm::Constant *&OutlinedFnID,
                                             bool IsOffloadEntry,
                                             const RegionCodeGenTy &CodeGen) {
  ExecutionRuntimeModesRAII ModeRAII(CurrentExecutionMode, EM_NonSPMD);
  EntryFunctionState EST;
  WrapperFunctionsMap.clear();

  [[maybe_unused]] bool IsBareKernel = D.getSingleClause<OMPXBareClause>();
  assert(!IsBareKernel && "bare kernel should not be at generic mode");

  // Emit target region as a standalone region.
  class NVPTXPrePostActionTy : public PrePostActionTy {
    CGOpenMPRuntimeGPU::EntryFunctionState &EST;
    const OMPExecutableDirective &D;

  public:
    NVPTXPrePostActionTy(CGOpenMPRuntimeGPU::EntryFunctionState &EST,
                         const OMPExecutableDirective &D)
        : EST(EST), D(D) {}
    void Enter(CodeGenFunction &CGF) override {
      auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
      RT.emitKernelInit(D, CGF, EST, /* IsSPMD */ false);
      // Skip target region initialization.
      RT.setLocThreadIdInsertPt(CGF, /*AtCurrentPoint=*/true);
    }
    void Exit(CodeGenFunction &CGF) override {
      auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
      RT.clearLocThreadIdInsertPt(CGF);
      RT.emitKernelDeinit(CGF, EST, /* IsSPMD */ false);
    }
  } Action(EST, D);
  CodeGen.setAction(Action);
  IsInTTDRegion = true;
  emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,
                                   IsOffloadEntry, CodeGen);
  IsInTTDRegion = false;
  GenerateMetaData(CGM, D, OutlinedFn, /*Generic*/ true);
}

void CGOpenMPRuntimeGPU::emitKernelInit(const OMPExecutableDirective &D,
                                        CodeGenFunction &CGF,
                                        EntryFunctionState &EST, bool IsSPMD) {
  llvm::OpenMPIRBuilder::TargetKernelDefaultAttrs Attrs;
  Attrs.ExecFlags =
      IsSPMD ? llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_SPMD
             : llvm::omp::OMPTgtExecModeFlags::OMP_TGT_EXEC_MODE_GENERIC;
  computeMinAndMaxThreadsAndTeams(D, CGF, Attrs);

  CGBuilderTy &Bld = CGF.Builder;
  Bld.restoreIP(OMPBuilder.createTargetInit(Bld, Attrs));
  if (!IsSPMD)
    emitGenericVarsProlog(CGF, EST.Loc);
}

void CGOpenMPRuntimeGPU::emitKernelDeinit(CodeGenFunction &CGF,
                                          EntryFunctionState &EST,
                                          bool IsSPMD) {
  if (!IsSPMD)
    emitGenericVarsEpilog(CGF);

  // This is temporary until we remove the fixed sized buffer.
  ASTContext &C = CGM.getContext();
  RecordDecl *StaticRD = C.buildImplicitRecord(
      "_openmp_teams_reduction_type_$_", RecordDecl::TagKind::Union);
  StaticRD->startDefinition();
  for (const RecordDecl *TeamReductionRec : TeamsReductions) {
    CanQualType RecTy = C.getCanonicalTagType(TeamReductionRec);
    auto *Field = FieldDecl::Create(
        C, StaticRD, SourceLocation(), SourceLocation(), nullptr, RecTy,
        C.getTrivialTypeSourceInfo(RecTy, SourceLocation()),
        /*BW=*/nullptr, /*Mutable=*/false,
        /*InitStyle=*/ICIS_NoInit);
    Field->setAccess(AS_public);
    StaticRD->addDecl(Field);
  }
  StaticRD->completeDefinition();
  CanQualType StaticTy = C.getCanonicalTagType(StaticRD);
  llvm::Type *LLVMReductionsBufferTy =
      CGM.getTypes().ConvertTypeForMem(StaticTy);
  const auto &DL = CGM.getModule().getDataLayout();
  uint64_t ReductionDataSize =
      TeamsReductions.empty()
          ? 0
          : DL.getTypeAllocSize(LLVMReductionsBufferTy).getFixedValue();
  CGBuilderTy &Bld = CGF.Builder;
  OMPBuilder.createTargetDeinit(Bld, ReductionDataSize,
                                C.getLangOpts().OpenMPCUDAReductionBufNum);
  TeamsReductions.clear();
}

void CGOpenMPRuntimeGPU::emitSPMDKernel(const OMPExecutableDirective &D,
                                          StringRef ParentName,
                                          llvm::Function *&OutlinedFn,
                                          llvm::Constant *&OutlinedFnID,
                                          bool IsOffloadEntry,
                                          const RegionCodeGenTy &CodeGen) {
  ExecutionRuntimeModesRAII ModeRAII(CurrentExecutionMode, EM_SPMD);
  EntryFunctionState EST;

  bool IsBareKernel = D.getSingleClause<OMPXBareClause>();

  // Emit target region as a standalone region.
  class NVPTXPrePostActionTy : public PrePostActionTy {
    CGOpenMPRuntimeGPU &RT;
    CGOpenMPRuntimeGPU::EntryFunctionState &EST;
    bool IsBareKernel;
    DataSharingMode Mode;
    const OMPExecutableDirective &D;

  public:
    NVPTXPrePostActionTy(CGOpenMPRuntimeGPU &RT,
                         CGOpenMPRuntimeGPU::EntryFunctionState &EST,
                         bool IsBareKernel, const OMPExecutableDirective &D)
        : RT(RT), EST(EST), IsBareKernel(IsBareKernel),
          Mode(RT.CurrentDataSharingMode), D(D) {}
    void Enter(CodeGenFunction &CGF) override {
      if (IsBareKernel) {
        RT.CurrentDataSharingMode = DataSharingMode::DS_CUDA;
        return;
      }
      RT.emitKernelInit(D, CGF, EST, /* IsSPMD */ true);
      // Skip target region initialization.
      RT.setLocThreadIdInsertPt(CGF, /*AtCurrentPoint=*/true);
    }
    void Exit(CodeGenFunction &CGF) override {
      if (IsBareKernel) {
        RT.CurrentDataSharingMode = Mode;
        return;
      }
      RT.clearLocThreadIdInsertPt(CGF);
      RT.emitKernelDeinit(CGF, EST, /* IsSPMD */ true);
    }
  } Action(*this, EST, IsBareKernel, D);
  CodeGen.setAction(Action);
  IsInTTDRegion = true;
  emitTargetOutlinedFunctionHelper(D, ParentName, OutlinedFn, OutlinedFnID,
                                   IsOffloadEntry, CodeGen);
  IsInTTDRegion = false;

  GenerateMetaData(CGM, D, OutlinedFn, /*SPMD*/ false);
}

// Create a unique global variable to indicate the execution mode of this target
// region. The execution mode is either 'generic', or 'spmd' depending on the
// target directive. This variable is picked up by the offload library to setup
// the device appropriately before kernel launch. If the execution mode is
// 'generic', the runtime reserves one warp for the master, otherwise, all
// warps participate in parallel work.
static void setPropertyExecutionMode(CodeGenModule &CGM, StringRef Name,
                                     OMPTgtExecModeFlags Mode) {
  auto *GVMode = new llvm::GlobalVariable(
      CGM.getModule(), CGM.Int8Ty, /*isConstant=*/true,
      llvm::GlobalValue::WeakAnyLinkage,
      llvm::ConstantInt::get(CGM.Int8Ty, Mode), Twine(Name, "_exec_mode"));
  CGM.addCompilerUsedGlobal(GVMode);
}

// Create a global variable to indicate whether fast reduction is enabled for
// this file. This variable is read by the runtime while determining the launch
// bounds.
static void setIsFastReduction(CodeGenModule &CGM) {
  auto *GVFastReduction = new llvm::GlobalVariable(
      CGM.getModule(), CGM.Int8Ty, /*isConstant=*/true,
      llvm::GlobalValue::WeakAnyLinkage,
      llvm::ConstantInt::get(CGM.Int8Ty,
                             CGM.getLangOpts().OpenMPTargetFastReduction),
      Twine("__omp_plugin_enable_fast_reduction"));
  CGM.addCompilerUsedGlobal(GVFastReduction);
}

static OMPTgtExecModeFlags
computeExecutionMode(bool Mode, const Stmt *DirectiveStmt, CodeGenModule &CGM) {
  if (!Mode)
    return OMP_TGT_EXEC_MODE_GENERIC;
  if (DirectiveStmt) {
    const Stmt *KernelForStmt = CGM.getSingleForStmt(DirectiveStmt);
    if (KernelForStmt) {
      if (CGM.isNoLoopKernel(KernelForStmt))
        return OMP_TGT_EXEC_MODE_SPMD_NO_LOOP;
      if (CGM.isBigJumpLoopKernel(KernelForStmt))
        return OMP_TGT_EXEC_MODE_SPMD_BIG_JUMP_LOOP;
      if (CGM.isXteamRedKernel(KernelForStmt))
        return OMP_TGT_EXEC_MODE_XTEAM_RED;
    }
  }
  return OMP_TGT_EXEC_MODE_SPMD;
}

void CGOpenMPRuntimeGPU::emitTargetOutlinedFunction(
    const OMPExecutableDirective &D, StringRef ParentName,
    llvm::Function *&OutlinedFn, llvm::Constant *&OutlinedFnID,
    bool IsOffloadEntry, const RegionCodeGenTy &CodeGen) {
  if (!IsOffloadEntry) // Nothing to do.
    return;

  assert(!ParentName.empty() && "Invalid target region parent name!");

  const Stmt *DirectiveStmt = CGM.getOptKernelKey(D);
  bool Mode = supportsSPMDExecutionMode(CGM, D);
  // Used by emitParallelCall
  CGM.setIsSPMDExecutionMode(Mode);
  if (Mode) {
    // For AMDGPU, check if a no-loop or a Xteam reduction kernel should
    // be generated and if so, set metadata that can be used by codegen.
    // This check is done regardless of host or device codegen since the
    // signature of the offloading routine has to match across host and device.
    if (CGM.getTriple().isAMDGCN()) {
      assert(CGM.getLangOpts().OpenMPIsTargetDevice && "Unexpected host path");
      CodeGenModule::NoLoopXteamErr NxStatus = CGM.checkAndSetNoLoopKernel(D);
      DEBUG_WITH_TYPE(NO_LOOP_XTEAM_RED,
                      CGM.emitNxResult("[No-Loop/Big-Jump-Loop]", D, NxStatus));
      if (NxStatus != CodeGenModule::NxSuccess) {
        NxStatus = CGM.checkAndSetXteamRedKernel(D);
        DEBUG_WITH_TYPE(NO_LOOP_XTEAM_RED,
                        CGM.emitNxResult("[Xteam]", D, NxStatus));
      }
    }
  }
  bool IsBareKernel = D.getSingleClause<OMPXBareClause>();
  if (Mode || IsBareKernel)
    emitSPMDKernel(D, ParentName, OutlinedFn, OutlinedFnID, IsOffloadEntry,
                   CodeGen);
  else {
    emitNonSPMDKernel(D, ParentName, OutlinedFn, OutlinedFnID, IsOffloadEntry,
                      CodeGen);
    DEBUG_WITH_TYPE(NO_LOOP_XTEAM_RED,
                    CGM.emitNxResult("[No-Loop/Big-Jump-Loop/Xteam]", D,
                                     CodeGenModule::NxNonSPMD));
  }
  setPropertyExecutionMode(
      CGM, OutlinedFn->getName(),
      IsBareKernel ? OMP_TGT_EXEC_MODE_BARE
                   : computeExecutionMode(Mode, DirectiveStmt, CGM));

  if (Mode && DirectiveStmt)
    CGM.resetOptKernelMetadata(DirectiveStmt);

  // Reset cached mode
  CGM.setIsSPMDExecutionMode(false);
}

CGOpenMPRuntimeGPU::CGOpenMPRuntimeGPU(CodeGenModule &CGM)
    : CGOpenMPRuntime(CGM) {
  llvm::OpenMPIRBuilderConfig Config(
      CGM.getLangOpts().OpenMPIsTargetDevice, isGPU(),
      CGM.getLangOpts().OpenMPOffloadMandatory,
      /*HasRequiresReverseOffload*/ false, /*HasRequiresUnifiedAddress*/ false,
      hasRequiresUnifiedSharedMemory(), /*HasRequiresDynamicAllocators*/ false);
  Config.setDefaultTargetAS(
      CGM.getContext().getTargetInfo().getTargetAddressSpace(LangAS::Default));
  Config.setRuntimeCC(CGM.getRuntimeCC());

  OMPBuilder.setConfig(Config);

  if (!CGM.getLangOpts().OpenMPIsTargetDevice)
    llvm_unreachable("OpenMP can only handle device code.");

  if (CGM.getLangOpts().OpenMPCUDAMode)
    CurrentDataSharingMode = CGOpenMPRuntimeGPU::DS_CUDA;

  // Write a global variable indicating whether fast reduction is enabled.
  // This is done regardless of -nogpulib
  if (!CGM.getLangOpts().OMPHostIRFile.empty())
    setIsFastReduction(CGM);

  llvm::OpenMPIRBuilder &OMPBuilder = getOMPBuilder();
  if (CGM.getLangOpts().NoGPULib || CGM.getLangOpts().OMPHostIRFile.empty())
    return;

  OMPBuilder.createGlobalFlag(CGM.getLangOpts().OpenMPTargetDebug,
                              "__omp_rtl_debug_kind");
  OMPBuilder.createGlobalFlag(CGM.getLangOpts().OpenMPTeamSubscription,
                              "__omp_rtl_assume_teams_oversubscription");
  OMPBuilder.createGlobalFlag(CGM.getLangOpts().OpenMPThreadSubscription,
                              "__omp_rtl_assume_threads_oversubscription");
  OMPBuilder.createGlobalFlag(CGM.getLangOpts().OpenMPNoThreadState,
                              "__omp_rtl_assume_no_thread_state");
  OMPBuilder.createGlobalFlag(CGM.getLangOpts().OpenMPNoNestedParallelism,
                              "__omp_rtl_assume_no_nested_parallelism");
}

void CGOpenMPRuntimeGPU::emitProcBindClause(CodeGenFunction &CGF,
                                              ProcBindKind ProcBind,
                                              SourceLocation Loc) {
  // Nothing to do.
}

llvm::Value *CGOpenMPRuntimeGPU::emitMessageClause(CodeGenFunction &CGF,
                                                   const Expr *Message,
                                                   SourceLocation Loc) {
  CGM.getDiags().Report(Loc, diag::warn_omp_gpu_unsupported_clause)
      << getOpenMPClauseName(OMPC_message);
  return nullptr;
}

llvm::Value *
CGOpenMPRuntimeGPU::emitSeverityClause(OpenMPSeverityClauseKind Severity,
                                       SourceLocation Loc) {
  CGM.getDiags().Report(Loc, diag::warn_omp_gpu_unsupported_clause)
      << getOpenMPClauseName(OMPC_severity);
  return nullptr;
}

void CGOpenMPRuntimeGPU::emitNumThreadsClause(
    CodeGenFunction &CGF, llvm::Value *NumThreads, SourceLocation Loc,
    OpenMPNumThreadsClauseModifier Modifier, OpenMPSeverityClauseKind Severity,
    SourceLocation SeverityLoc, const Expr *Message,
    SourceLocation MessageLoc) {
  if (Modifier == OMPC_NUMTHREADS_strict) {
    CGM.getDiags().Report(Loc,
                          diag::warn_omp_gpu_unsupported_modifier_for_clause)
        << "strict" << getOpenMPClauseName(OMPC_num_threads);
    return;
  }

  // Nothing to do.
}

void CGOpenMPRuntimeGPU::emitNumTeamsClause(CodeGenFunction &CGF,
                                              const Expr *NumTeams,
                                              const Expr *ThreadLimit,
                                              SourceLocation Loc) {}

llvm::Function *CGOpenMPRuntimeGPU::emitParallelOutlinedFunction(
    CodeGenFunction &CGF, const OMPExecutableDirective &D,
    const VarDecl *ThreadIDVar, OpenMPDirectiveKind InnermostKind,
    const RegionCodeGenTy &CodeGen) {
  // Emit target region as a standalone region.
  bool PrevIsInTTDRegion = IsInTTDRegion;
  IsInTTDRegion = false;
  auto *OutlinedFun =
      cast<llvm::Function>(CGOpenMPRuntime::emitParallelOutlinedFunction(
          CGF, D, ThreadIDVar, InnermostKind, CodeGen));
  IsInTTDRegion = PrevIsInTTDRegion;
  if (getExecutionMode() != CGOpenMPRuntimeGPU::EM_SPMD) {
    llvm::Function *WrapperFun =
        createParallelDataSharingWrapper(OutlinedFun, D);
    WrapperFunctionsMap[OutlinedFun] = WrapperFun;
  }

  return OutlinedFun;
}

/// Get list of lastprivate variables from the teams distribute ... or
/// teams {distribute ...} directives.
static void
getDistributeLastprivateVars(ASTContext &Ctx, const OMPExecutableDirective &D,
                             llvm::SmallVectorImpl<const ValueDecl *> &Vars) {
  assert(isOpenMPTeamsDirective(D.getDirectiveKind()) &&
         "expected teams directive.");
  const OMPExecutableDirective *Dir = &D;
  if (!isOpenMPDistributeDirective(D.getDirectiveKind())) {
    if (const Stmt *S = CGOpenMPRuntime::getSingleCompoundChild(
            Ctx,
            D.getInnermostCapturedStmt()->getCapturedStmt()->IgnoreContainers(
                /*IgnoreCaptured=*/true))) {
      Dir = dyn_cast_or_null<OMPExecutableDirective>(S);
      if (Dir && !isOpenMPDistributeDirective(Dir->getDirectiveKind()))
        Dir = nullptr;
    }
  }
  if (!Dir)
    return;
  for (const auto *C : Dir->getClausesOfKind<OMPLastprivateClause>()) {
    for (const Expr *E : C->getVarRefs())
      Vars.push_back(getPrivateItem(E));
  }
}

/// Get list of reduction variables from the teams ... directives.
static void
getTeamsReductionVars(ASTContext &Ctx, const OMPExecutableDirective &D,
                      llvm::SmallVectorImpl<const ValueDecl *> &Vars) {
  assert(isOpenMPTeamsDirective(D.getDirectiveKind()) &&
         "expected teams directive.");
  for (const auto *C : D.getClausesOfKind<OMPReductionClause>()) {
    for (const Expr *E : C->privates())
      Vars.push_back(getPrivateItem(E));
  }
}

llvm::Function *CGOpenMPRuntimeGPU::emitTeamsOutlinedFunction(
    CodeGenFunction &CGF, const OMPExecutableDirective &D,
    const VarDecl *ThreadIDVar, OpenMPDirectiveKind InnermostKind,
    const RegionCodeGenTy &CodeGen) {
  SourceLocation Loc = D.getBeginLoc();

  const RecordDecl *GlobalizedRD = nullptr;
  llvm::SmallVector<const ValueDecl *, 4> LastPrivatesReductions;
  llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *> MappedDeclsFields;
  unsigned WarpSize = CGM.getTarget().getGridValue().GV_Warp_Size;
  // Globalize team reductions variable unconditionally in all modes.
  if (getExecutionMode() != CGOpenMPRuntimeGPU::EM_SPMD)
    getTeamsReductionVars(CGM.getContext(), D, LastPrivatesReductions);
  if (getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD) {
    getDistributeLastprivateVars(CGM.getContext(), D, LastPrivatesReductions);
    if (!LastPrivatesReductions.empty()) {
      GlobalizedRD = ::buildRecordForGlobalizedVars(
          CGM.getContext(), {}, LastPrivatesReductions, MappedDeclsFields,
          WarpSize);
    }
  } else if (!LastPrivatesReductions.empty()) {
    assert(!TeamAndReductions.first &&
           "Previous team declaration is not expected.");
    TeamAndReductions.first = D.getCapturedStmt(OMPD_teams)->getCapturedDecl();
    std::swap(TeamAndReductions.second, LastPrivatesReductions);
  }

  // Emit target region as a standalone region.
  class NVPTXPrePostActionTy : public PrePostActionTy {
    SourceLocation &Loc;
    const RecordDecl *GlobalizedRD;
    llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>
        &MappedDeclsFields;

  public:
    NVPTXPrePostActionTy(
        SourceLocation &Loc, const RecordDecl *GlobalizedRD,
        llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *>
            &MappedDeclsFields)
        : Loc(Loc), GlobalizedRD(GlobalizedRD),
          MappedDeclsFields(MappedDeclsFields) {}
    void Enter(CodeGenFunction &CGF) override {
      auto &Rt =
          static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
      if (GlobalizedRD) {
        auto I = Rt.FunctionGlobalizedDecls.try_emplace(CGF.CurFn).first;
        I->getSecond().MappedParams =
            std::make_unique<CodeGenFunction::OMPMapVars>();
        DeclToAddrMapTy &Data = I->getSecond().LocalVarData;
        for (const auto &Pair : MappedDeclsFields) {
          assert(Pair.getFirst()->isCanonicalDecl() &&
                 "Expected canonical declaration");
          Data.try_emplace(Pair.getFirst());
        }
      }
      Rt.emitGenericVarsProlog(CGF, Loc);
    }
    void Exit(CodeGenFunction &CGF) override {
      static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime())
          .emitGenericVarsEpilog(CGF);
    }
  } Action(Loc, GlobalizedRD, MappedDeclsFields);
  CodeGen.setAction(Action);

  llvm::Function *OutlinedFun = CGOpenMPRuntime::emitTeamsOutlinedFunction(
      CGF, D, ThreadIDVar, InnermostKind, CodeGen);

  return OutlinedFun;
}

void CGOpenMPRuntimeGPU::emitGenericVarsProlog(CodeGenFunction &CGF,
                                               SourceLocation Loc) {
  if (getDataSharingMode() != CGOpenMPRuntimeGPU::DS_Generic)
    return;

  CGBuilderTy &Bld = CGF.Builder;

  const auto I = FunctionGlobalizedDecls.find(CGF.CurFn);
  if (I == FunctionGlobalizedDecls.end())
    return;

  for (auto &Rec : I->getSecond().LocalVarData) {
    const auto *VD = cast<VarDecl>(Rec.first);
    bool EscapedParam = I->getSecond().EscapedParameters.count(Rec.first);
    QualType VarTy = VD->getType();

    // Get the local allocation of a firstprivate variable before sharing
    llvm::Value *ParValue;
    if (EscapedParam) {
      LValue ParLVal =
          CGF.MakeAddrLValue(CGF.GetAddrOfLocalVar(VD), VD->getType());
      ParValue = CGF.EmitLoadOfScalar(ParLVal, Loc);
    }

    // Allocate space for the variable to be globalized
    llvm::Value *AllocArgs[] = {CGF.getTypeSize(VD->getType())};
    llvm::CallBase *VoidPtr =
        CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
                                CGM.getModule(), OMPRTL___kmpc_alloc_shared),
                            AllocArgs, VD->getName());
    // FIXME: We should use the variables actual alignment as an argument.
    VoidPtr->addRetAttr(llvm::Attribute::get(
        CGM.getLLVMContext(), llvm::Attribute::Alignment,
        CGM.getContext().getTargetInfo().getNewAlign() / 8));

    // Cast the void pointer and get the address of the globalized variable.
    llvm::Value *CastedVoidPtr = Bld.CreatePointerBitCastOrAddrSpaceCast(
        VoidPtr, Bld.getPtrTy(0), VD->getName() + "_on_stack");
    LValue VarAddr =
        CGF.MakeNaturalAlignPointeeRawAddrLValue(CastedVoidPtr, VarTy);
    Rec.second.PrivateAddr = VarAddr.getAddress();
    Rec.second.GlobalizedVal = VoidPtr;

    // Assign the local allocation to the newly globalized location.
    if (EscapedParam) {
      CGF.EmitStoreOfScalar(ParValue, VarAddr);
      I->getSecond().MappedParams->setVarAddr(CGF, VD, VarAddr.getAddress());
    }
    if (auto *DI = CGF.getDebugInfo())
      VoidPtr->setDebugLoc(DI->SourceLocToDebugLoc(VD->getLocation()));
  }

  for (const auto *ValueD : I->getSecond().EscapedVariableLengthDecls) {
    const auto *VD = cast<VarDecl>(ValueD);
    std::pair<llvm::Value *, llvm::Value *> AddrSizePair =
        getKmpcAllocShared(CGF, VD);
    I->getSecond().EscapedVariableLengthDeclsAddrs.emplace_back(AddrSizePair);
    LValue Base = CGF.MakeAddrLValue(AddrSizePair.first, VD->getType(),
                                     CGM.getContext().getDeclAlign(VD),
                                     AlignmentSource::Decl);
    I->getSecond().MappedParams->setVarAddr(CGF, VD, Base.getAddress());
  }
  I->getSecond().MappedParams->apply(CGF);
}

bool CGOpenMPRuntimeGPU::isDelayedVariableLengthDecl(CodeGenFunction &CGF,
                                                     const VarDecl *VD) const {
  const auto I = FunctionGlobalizedDecls.find(CGF.CurFn);
  if (I == FunctionGlobalizedDecls.end())
    return false;

  // Check variable declaration is delayed:
  return llvm::is_contained(I->getSecond().DelayedVariableLengthDecls, VD);
}

std::pair<llvm::Value *, llvm::Value *>
CGOpenMPRuntimeGPU::getKmpcAllocShared(CodeGenFunction &CGF,
                                       const VarDecl *VD) {
  CGBuilderTy &Bld = CGF.Builder;

  // Compute size and alignment.
  llvm::Value *Size = CGF.getTypeSize(VD->getType());
  CharUnits Align = CGM.getContext().getDeclAlign(VD);
  Size = Bld.CreateNUWAdd(
      Size, llvm::ConstantInt::get(CGF.SizeTy, Align.getQuantity() - 1));
  llvm::Value *AlignVal =
      llvm::ConstantInt::get(CGF.SizeTy, Align.getQuantity());
  Size = Bld.CreateUDiv(Size, AlignVal);
  Size = Bld.CreateNUWMul(Size, AlignVal);

  // Allocate space for this VLA object to be globalized.
  llvm::Value *AllocArgs[] = {Size};
  llvm::CallBase *VoidPtr =
      CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
                              CGM.getModule(), OMPRTL___kmpc_alloc_shared),
                          AllocArgs, VD->getName());
  VoidPtr->addRetAttr(llvm::Attribute::get(
      CGM.getLLVMContext(), llvm::Attribute::Alignment, Align.getQuantity()));

  return std::make_pair(VoidPtr, Size);
}

void CGOpenMPRuntimeGPU::getKmpcFreeShared(
    CodeGenFunction &CGF,
    const std::pair<llvm::Value *, llvm::Value *> &AddrSizePair) {
  // Deallocate the memory for each globalized VLA object
  CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
                          CGM.getModule(), OMPRTL___kmpc_free_shared),
                      {AddrSizePair.first, AddrSizePair.second});
}

void CGOpenMPRuntimeGPU::emitGenericVarsEpilog(CodeGenFunction &CGF) {
  if (getDataSharingMode() != CGOpenMPRuntimeGPU::DS_Generic)
    return;

  const auto I = FunctionGlobalizedDecls.find(CGF.CurFn);
  if (I != FunctionGlobalizedDecls.end()) {
    // Deallocate the memory for each globalized VLA object that was
    // globalized in the prolog (i.e. emitGenericVarsProlog).
    for (const auto &AddrSizePair :
         llvm::reverse(I->getSecond().EscapedVariableLengthDeclsAddrs)) {
      CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
                              CGM.getModule(), OMPRTL___kmpc_free_shared),
                          {AddrSizePair.first, AddrSizePair.second});
    }
    // Deallocate the memory for each globalized value
    for (auto &Rec : llvm::reverse(I->getSecond().LocalVarData)) {
      const auto *VD = cast<VarDecl>(Rec.first);
      I->getSecond().MappedParams->restore(CGF);

      llvm::Value *FreeArgs[] = {Rec.second.GlobalizedVal,
                                 CGF.getTypeSize(VD->getType())};
      CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
                              CGM.getModule(), OMPRTL___kmpc_free_shared),
                          FreeArgs);
    }
  }
}

void CGOpenMPRuntimeGPU::emitTeamsCall(CodeGenFunction &CGF,
                                         const OMPExecutableDirective &D,
                                         SourceLocation Loc,
                                         llvm::Function *OutlinedFn,
                                         ArrayRef<llvm::Value *> CapturedVars) {
  if (!CGF.HaveInsertPoint())
    return;

  bool IsBareKernel = D.getSingleClause<OMPXBareClause>();

  RawAddress ZeroAddr = CGF.CreateDefaultAlignTempAlloca(CGF.Int32Ty,
                                                         /*Name=*/".zero.addr");
  CGF.Builder.CreateStore(CGF.Builder.getInt32(/*C*/ 0), ZeroAddr);
  llvm::SmallVector<llvm::Value *, 16> OutlinedFnArgs;
  // We don't emit any thread id function call in bare kernel, but because the
  // outlined function has a pointer argument, we emit a nullptr here.
  if (IsBareKernel)
    OutlinedFnArgs.push_back(llvm::ConstantPointerNull::get(CGM.VoidPtrTy));
  else
    OutlinedFnArgs.push_back(emitThreadIDAddress(CGF, Loc).emitRawPointer(CGF));
  OutlinedFnArgs.push_back(ZeroAddr.getPointer());

  // If this is a kernel we can run on multiple devices then we need to add
  // the arguments for multi-device targets. This is needed for the case when
  // we emit an outlined teams function which needs to be passed the multi
  // device LB and UB.
  if (CGM.isMultiDeviceKernel(D)) {
    Address LBAddr =
        CGF.GetAddrOfLocalVar(CGM.getMultiDeviceLBArg(D, CGF.CurFn));
    OutlinedFnArgs.push_back(CGF.Builder.CreateLoad(LBAddr));
    Address UBAddr =
        CGF.GetAddrOfLocalVar(CGM.getMultiDeviceUBArg(D, CGF.CurFn));
    OutlinedFnArgs.push_back(CGF.Builder.CreateLoad(UBAddr));
  }

  OutlinedFnArgs.append(CapturedVars.begin(), CapturedVars.end());
  emitOutlinedFunctionCall(CGF, Loc, OutlinedFn, OutlinedFnArgs);
}

void CGOpenMPRuntimeGPU::emitParallelCall(
    CodeGenFunction &CGF, SourceLocation Loc, llvm::Function *OutlinedFn,
    ArrayRef<llvm::Value *> CapturedVars, const Expr *IfCond,
    llvm::Value *NumThreads, OpenMPNumThreadsClauseModifier NumThreadsModifier,
    OpenMPSeverityClauseKind Severity, const Expr *Message) {
  if (!CGF.HaveInsertPoint())
    return;

  auto &&ParallelGen = [this, Loc, OutlinedFn, CapturedVars, IfCond,
                        NumThreads](CodeGenFunction &CGF,
                                    PrePostActionTy &Action) {
    CGBuilderTy &Bld = CGF.Builder;
    llvm::Value *NumThreadsVal = NumThreads;
    llvm::Function *WFn = WrapperFunctionsMap[OutlinedFn];
    llvm::PointerType *FnPtrTy = llvm::PointerType::get(
        CGF.getLLVMContext(), CGM.getDataLayout().getProgramAddressSpace());

    llvm::Value *ID = llvm::ConstantPointerNull::get(FnPtrTy);
    if (WFn)
      ID = Bld.CreateBitOrPointerCast(WFn, FnPtrTy);

    llvm::Value *FnPtr = Bld.CreateBitOrPointerCast(OutlinedFn, FnPtrTy);

    // Create a private scope that will globalize the arguments
    // passed from the outside of the target region.
    // TODO: Is that needed?
    CodeGenFunction::OMPPrivateScope PrivateArgScope(CGF);

    Address CapturedVarsAddrs = CGF.CreateDefaultAlignTempAlloca(
        llvm::ArrayType::get(CGM.VoidPtrTy, CapturedVars.size()),
        "captured_vars_addrs");
    // There's something to share.
    if (!CapturedVars.empty()) {
      // Prepare for parallel region. Indicate the outlined function.
      ASTContext &Ctx = CGF.getContext();
      unsigned Idx = 0;
      for (llvm::Value *V : CapturedVars) {
        Address Dst = Bld.CreateConstArrayGEP(CapturedVarsAddrs, Idx);
        llvm::Value *PtrV;
        if (V->getType()->isIntegerTy())
          PtrV = Bld.CreateIntToPtr(V, CGF.VoidPtrTy);
        else
          PtrV = Bld.CreatePointerBitCastOrAddrSpaceCast(V, CGF.VoidPtrTy);
        CGF.EmitStoreOfScalar(PtrV, Dst, /*Volatile=*/false,
                              Ctx.getPointerType(Ctx.VoidPtrTy));
        ++Idx;
      }
    }

    llvm::Value *IfCondVal = nullptr;
    if (IfCond)
      IfCondVal = Bld.CreateIntCast(CGF.EvaluateExprAsBool(IfCond), CGF.Int32Ty,
                                    /* isSigned */ false);
    else
      IfCondVal = llvm::ConstantInt::get(CGF.Int32Ty, 1);

    if (!NumThreadsVal)
      NumThreadsVal = llvm::ConstantInt::getAllOnesValue(CGF.Int32Ty);
    else
      NumThreadsVal = Bld.CreateZExtOrTrunc(NumThreadsVal, CGF.Int32Ty);

    // No strict prescriptiveness for the number of threads.
    llvm::Value *StrictNumThreadsVal = llvm::ConstantInt::get(CGF.Int32Ty, 0);

    assert(IfCondVal && "Expected a value");
    llvm::Value *RTLoc = emitUpdateLocation(CGF, Loc);
    llvm::Value *Args[] = {
        RTLoc,
        getThreadID(CGF, Loc),
        IfCondVal,
        NumThreadsVal,
        llvm::ConstantInt::getAllOnesValue(CGF.Int32Ty),
        FnPtr,
        ID,
        Bld.CreateBitOrPointerCast(CapturedVarsAddrs.emitRawPointer(CGF),
                                   CGF.VoidPtrPtrTy),
        llvm::ConstantInt::get(CGM.SizeTy, CapturedVars.size()),
        StrictNumThreadsVal};

    CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
                            CGM.getModule(), OMPRTL___kmpc_parallel_60),
                        Args);
  };

  RegionCodeGenTy RCG(ParallelGen);
  RCG(CGF);
}

void CGOpenMPRuntimeGPU::syncCTAThreads(CodeGenFunction &CGF) {
  // Always emit simple barriers!
  if (!CGF.HaveInsertPoint())
    return;
  // Build call __kmpc_barrier_simple_spmd(nullptr, 0);
  // This function does not use parameters, so we can emit just default values.
  llvm::Value *Args[] = {
      llvm::ConstantPointerNull::get(
          cast<llvm::PointerType>(getIdentTyPointerTy())),
      llvm::ConstantInt::get(CGF.Int32Ty, /*V=*/0, /*isSigned=*/true)};
  CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
                          CGM.getModule(), OMPRTL___kmpc_barrier_simple_spmd),
                      Args);
}

void CGOpenMPRuntimeGPU::emitBarrierCall(CodeGenFunction &CGF,
                                           SourceLocation Loc,
                                           OpenMPDirectiveKind Kind, bool,
                                           bool) {
  // Always emit simple barriers!
  if (!CGF.HaveInsertPoint())
    return;
  // Build call __kmpc_cancel_barrier(loc, thread_id);
  unsigned Flags = getDefaultFlagsForBarriers(Kind);
  llvm::Value *Args[] = {emitUpdateLocation(CGF, Loc, Flags),
                         getThreadID(CGF, Loc)};

  CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
                          CGM.getModule(), OMPRTL___kmpc_barrier),
                      Args);
}

void CGOpenMPRuntimeGPU::emitCriticalRegion(
    CodeGenFunction &CGF, StringRef CriticalName,
    const RegionCodeGenTy &CriticalOpGen, SourceLocation Loc,
    const Expr *Hint) {
  llvm::BasicBlock *LoopBB = CGF.createBasicBlock("omp.critical.loop");
  llvm::BasicBlock *TestBB = CGF.createBasicBlock("omp.critical.test");
  llvm::BasicBlock *SyncBB = CGF.createBasicBlock("omp.critical.sync");
  llvm::BasicBlock *BodyBB = CGF.createBasicBlock("omp.critical.body");
  llvm::BasicBlock *ExitBB = CGF.createBasicBlock("omp.critical.exit");

  auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());

  // Get the mask of active threads in the warp.
  llvm::Value *Mask = CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
      CGM.getModule(), OMPRTL___kmpc_warp_active_thread_mask));
  // Fetch team-local id of the thread.
  llvm::Value *ThreadID = RT.getGPUThreadID(CGF);

  // Get the width of the team.
  llvm::Value *TeamWidth = RT.getGPUNumThreads(CGF);

  // Initialize the counter variable for the loop.
  QualType Int32Ty =
      CGF.getContext().getIntTypeForBitwidth(/*DestWidth=*/32, /*Signed=*/0);
  Address Counter = CGF.CreateMemTempWithoutCast(Int32Ty, "critical_counter");
  LValue CounterLVal = CGF.MakeAddrLValue(Counter, Int32Ty);
  CGF.EmitStoreOfScalar(llvm::Constant::getNullValue(CGM.Int32Ty), CounterLVal,
                        /*isInit=*/true);

  // Block checks if loop counter exceeds upper bound.
  CGF.EmitBlock(LoopBB);
  llvm::Value *CounterVal = CGF.EmitLoadOfScalar(CounterLVal, Loc);
  llvm::Value *CmpLoopBound = CGF.Builder.CreateICmpSLT(CounterVal, TeamWidth);
  CGF.Builder.CreateCondBr(CmpLoopBound, TestBB, ExitBB);

  // Block tests which single thread should execute region, and which threads
  // should go straight to synchronisation point.
  CGF.EmitBlock(TestBB);
  CounterVal = CGF.EmitLoadOfScalar(CounterLVal, Loc);
  llvm::Value *CmpThreadToCounter =
      CGF.Builder.CreateICmpEQ(ThreadID, CounterVal);
  CGF.Builder.CreateCondBr(CmpThreadToCounter, BodyBB, SyncBB);

  // Block emits the body of the critical region.
  CGF.EmitBlock(BodyBB);

  // Output the critical statement.
  CGOpenMPRuntime::emitCriticalRegion(CGF, CriticalName, CriticalOpGen, Loc,
                                      Hint);

  // After the body surrounded by the critical region, the single executing
  // thread will jump to the synchronisation point.
  // Block waits for all threads in current team to finish then increments the
  // counter variable and returns to the loop.
  CGF.EmitBlock(SyncBB);
  // Reconverge active threads in the warp.
  (void)CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
                                CGM.getModule(), OMPRTL___kmpc_syncwarp),
                            Mask);

  llvm::Value *IncCounterVal =
      CGF.Builder.CreateNSWAdd(CounterVal, CGF.Builder.getInt32(1));
  CGF.EmitStoreOfScalar(IncCounterVal, CounterLVal);
  CGF.EmitBranch(LoopBB);

  // Block that is reached when  all threads in the team complete the region.
  CGF.EmitBlock(ExitBB, /*IsFinished=*/true);
}

/// Cast value to the specified type.
static llvm::Value *castValueToType(CodeGenFunction &CGF, llvm::Value *Val,
                                    QualType ValTy, QualType CastTy,
                                    SourceLocation Loc) {
  assert(!CGF.getContext().getTypeSizeInChars(CastTy).isZero() &&
         "Cast type must sized.");
  assert(!CGF.getContext().getTypeSizeInChars(ValTy).isZero() &&
         "Val type must sized.");
  llvm::Type *LLVMCastTy = CGF.ConvertTypeForMem(CastTy);
  if (ValTy == CastTy)
    return Val;
  if (CGF.getContext().getTypeSizeInChars(ValTy) ==
      CGF.getContext().getTypeSizeInChars(CastTy))
    return CGF.Builder.CreateBitCast(Val, LLVMCastTy);
  if (CastTy->isIntegerType() && ValTy->isIntegerType())
    return CGF.Builder.CreateIntCast(Val, LLVMCastTy,
                                     CastTy->hasSignedIntegerRepresentation());
  Address CastItem = CGF.CreateMemTempWithoutCast(CastTy);
  Address ValCastItem = CastItem.withElementType(Val->getType());
  CGF.EmitStoreOfScalar(Val, ValCastItem, /*Volatile=*/false, ValTy,
                        LValueBaseInfo(AlignmentSource::Type),
                        TBAAAccessInfo());
  return CGF.EmitLoadOfScalar(CastItem, /*Volatile=*/false, CastTy, Loc,
                              LValueBaseInfo(AlignmentSource::Type),
                              TBAAAccessInfo());
}

///
/// Design of OpenMP reductions on the GPU
///
/// Consider a typical OpenMP program with one or more reduction
/// clauses:
///
/// float foo;
/// double bar;
/// #pragma omp target teams distribute parallel for \
///             reduction(+:foo) reduction(*:bar)
/// for (int i = 0; i < N; i++) {
///   foo += A[i]; bar *= B[i];
/// }
///
/// where 'foo' and 'bar' are reduced across all OpenMP threads in
/// all teams.  In our OpenMP implementation on the NVPTX device an
/// OpenMP team is mapped to a CUDA threadblock and OpenMP threads
/// within a team are mapped to CUDA threads within a threadblock.
/// Our goal is to efficiently aggregate values across all OpenMP
/// threads such that:
///
///   - the compiler and runtime are logically concise, and
///   - the reduction is performed efficiently in a hierarchical
///     manner as follows: within OpenMP threads in the same warp,
///     across warps in a threadblock, and finally across teams on
///     the NVPTX device.
///
/// Introduction to Decoupling
///
/// We would like to decouple the compiler and the runtime so that the
/// latter is ignorant of the reduction variables (number, data types)
/// and the reduction operators.  This allows a simpler interface
/// and implementation while still attaining good performance.
///
/// Pseudocode for the aforementioned OpenMP program generated by the
/// compiler is as follows:
///
/// 1. Create private copies of reduction variables on each OpenMP
///    thread: 'foo_private', 'bar_private'
/// 2. Each OpenMP thread reduces the chunk of 'A' and 'B' assigned
///    to it and writes the result in 'foo_private' and 'bar_private'
///    respectively.
/// 3. Call the OpenMP runtime on the GPU to reduce within a team
///    and store the result on the team master:
///
///     __kmpc_nvptx_parallel_reduce_nowait_v2(...,
///        reduceData, shuffleReduceFn, interWarpCpyFn)
///
///     where:
///       struct ReduceData {
///         double *foo;
///         double *bar;
///       } reduceData
///       reduceData.foo = &foo_private
///       reduceData.bar = &bar_private
///
///     'shuffleReduceFn' and 'interWarpCpyFn' are pointers to two
///     auxiliary functions generated by the compiler that operate on
///     variables of type 'ReduceData'.  They aid the runtime perform
///     algorithmic steps in a data agnostic manner.
///
///     'shuffleReduceFn' is a pointer to a function that reduces data
///     of type 'ReduceData' across two OpenMP threads (lanes) in the
///     same warp.  It takes the following arguments as input:
///
///     a. variable of type 'ReduceData' on the calling lane,
///     b. its lane_id,
///     c. an offset relative to the current lane_id to generate a
///        remote_lane_id.  The remote lane contains the second
///        variable of type 'ReduceData' that is to be reduced.
///     d. an algorithm version parameter determining which reduction
///        algorithm to use.
///
///     'shuffleReduceFn' retrieves data from the remote lane using
///     efficient GPU shuffle intrinsics and reduces, using the
///     algorithm specified by the 4th parameter, the two operands
///     element-wise.  The result is written to the first operand.
///
///     Different reduction algorithms are implemented in different
///     runtime functions, all calling 'shuffleReduceFn' to perform
///     the essential reduction step.  Therefore, based on the 4th
///     parameter, this function behaves slightly differently to
///     cooperate with the runtime to ensure correctness under
///     different circumstances.
///
///     'InterWarpCpyFn' is a pointer to a function that transfers
///     reduced variables across warps.  It tunnels, through CUDA
///     shared memory, the thread-private data of type 'ReduceData'
///     from lane 0 of each warp to a lane in the first warp.
/// 4. Call the OpenMP runtime on the GPU to reduce across teams.
///    The last team writes the global reduced value to memory.
///
///     ret = __kmpc_nvptx_teams_reduce_nowait(...,
///             reduceData, shuffleReduceFn, interWarpCpyFn,
///             scratchpadCopyFn, loadAndReduceFn)
///
///     'scratchpadCopyFn' is a helper that stores reduced
///     data from the team master to a scratchpad array in
///     global memory.
///
///     'loadAndReduceFn' is a helper that loads data from
///     the scratchpad array and reduces it with the input
///     operand.
///
///     These compiler generated functions hide address
///     calculation and alignment information from the runtime.
/// 5. if ret == 1:
///     The team master of the last team stores the reduced
///     result to the globals in memory.
///     foo += reduceData.foo; bar *= reduceData.bar
///
///
/// Warp Reduction Algorithms
///
/// On the warp level, we have three algorithms implemented in the
/// OpenMP runtime depending on the number of active lanes:
///
/// Full Warp Reduction
///
/// The reduce algorithm within a warp where all lanes are active
/// is implemented in the runtime as follows:
///
/// full_warp_reduce(void *reduce_data,
///                  kmp_ShuffleReductFctPtr ShuffleReduceFn) {
///   for (int offset = WARPSIZE/2; offset > 0; offset /= 2)
///     ShuffleReduceFn(reduce_data, 0, offset, 0);
/// }
///
/// The algorithm completes in log(2, WARPSIZE) steps.
///
/// 'ShuffleReduceFn' is used here with lane_id set to 0 because it is
/// not used therefore we save instructions by not retrieving lane_id
/// from the corresponding special registers.  The 4th parameter, which
/// represents the version of the algorithm being used, is set to 0 to
/// signify full warp reduction.
///
/// In this version, 'ShuffleReduceFn' behaves, per element, as follows:
///
/// #reduce_elem refers to an element in the local lane's data structure
/// #remote_elem is retrieved from a remote lane
/// remote_elem = shuffle_down(reduce_elem, offset, WARPSIZE);
/// reduce_elem = reduce_elem REDUCE_OP remote_elem;
///
/// Contiguous Partial Warp Reduction
///
/// This reduce algorithm is used within a warp where only the first
/// 'n' (n <= WARPSIZE) lanes are active.  It is typically used when the
/// number of OpenMP threads in a parallel region is not a multiple of
/// WARPSIZE.  The algorithm is implemented in the runtime as follows:
///
/// void
/// contiguous_partial_reduce(void *reduce_data,
///                           kmp_ShuffleReductFctPtr ShuffleReduceFn,
///                           int size, int lane_id) {
///   int curr_size;
///   int offset;
///   curr_size = size;
///   mask = curr_size/2;
///   while (offset>0) {
///     ShuffleReduceFn(reduce_data, lane_id, offset, 1);
///     curr_size = (curr_size+1)/2;
///     offset = curr_size/2;
///   }
/// }
///
/// In this version, 'ShuffleReduceFn' behaves, per element, as follows:
///
/// remote_elem = shuffle_down(reduce_elem, offset, WARPSIZE);
/// if (lane_id < offset)
///     reduce_elem = reduce_elem REDUCE_OP remote_elem
/// else
///     reduce_elem = remote_elem
///
/// This algorithm assumes that the data to be reduced are located in a
/// contiguous subset of lanes starting from the first.  When there is
/// an odd number of active lanes, the data in the last lane is not
/// aggregated with any other lane's dat but is instead copied over.
///
/// Dispersed Partial Warp Reduction
///
/// This algorithm is used within a warp when any discontiguous subset of
/// lanes are active.  It is used to implement the reduction operation
/// across lanes in an OpenMP simd region or in a nested parallel region.
///
/// void
/// dispersed_partial_reduce(void *reduce_data,
///                          kmp_ShuffleReductFctPtr ShuffleReduceFn) {
///   int size, remote_id;
///   int logical_lane_id = number_of_active_lanes_before_me() * 2;
///   do {
///       remote_id = next_active_lane_id_right_after_me();
///       # the above function returns 0 of no active lane
///       # is present right after the current lane.
///       size = number_of_active_lanes_in_this_warp();
///       logical_lane_id /= 2;
///       ShuffleReduceFn(reduce_data, logical_lane_id,
///                       remote_id-1-threadIdx.x, 2);
///   } while (logical_lane_id % 2 == 0 && size > 1);
/// }
///
/// There is no assumption made about the initial state of the reduction.
/// Any number of lanes (>=1) could be active at any position.  The reduction
/// result is returned in the first active lane.
///
/// In this version, 'ShuffleReduceFn' behaves, per element, as follows:
///
/// remote_elem = shuffle_down(reduce_elem, offset, WARPSIZE);
/// if (lane_id % 2 == 0 && offset > 0)
///     reduce_elem = reduce_elem REDUCE_OP remote_elem
/// else
///     reduce_elem = remote_elem
///
///
/// Intra-Team Reduction
///
/// This function, as implemented in the runtime call
/// '__kmpc_nvptx_parallel_reduce_nowait_v2', aggregates data across OpenMP
/// threads in a team.  It first reduces within a warp using the
/// aforementioned algorithms.  We then proceed to gather all such
/// reduced values at the first warp.
///
/// The runtime makes use of the function 'InterWarpCpyFn', which copies
/// data from each of the "warp master" (zeroth lane of each warp, where
/// warp-reduced data is held) to the zeroth warp.  This step reduces (in
/// a mathematical sense) the problem of reduction across warp masters in
/// a block to the problem of warp reduction.
///
///
/// Inter-Team Reduction
///
/// Once a team has reduced its data to a single value, it is stored in
/// a global scratchpad array.  Since each team has a distinct slot, this
/// can be done without locking.
///
/// The last team to write to the scratchpad array proceeds to reduce the
/// scratchpad array.  One or more workers in the last team use the helper
/// 'loadAndReduceDataFn' to load and reduce values from the array, i.e.,
/// the k'th worker reduces every k'th element.
///
/// Finally, a call is made to '__kmpc_nvptx_parallel_reduce_nowait_v2' to
/// reduce across workers and compute a globally reduced value.
///
void CGOpenMPRuntimeGPU::emitReduction(
    CodeGenFunction &CGF, SourceLocation Loc, ArrayRef<const Expr *> Privates,
    ArrayRef<const Expr *> LHSExprs, ArrayRef<const Expr *> RHSExprs,
    ArrayRef<const Expr *> ReductionOps, ReductionOptionsTy Options) {
  if (!CGF.HaveInsertPoint())
    return;

  bool ParallelReduction = isOpenMPParallelDirective(Options.ReductionKind);
  bool TeamsReduction = isOpenMPTeamsDirective(Options.ReductionKind);

  ASTContext &C = CGM.getContext();

  if (Options.SimpleReduction) {
    assert(!TeamsReduction && !ParallelReduction &&
           "Invalid reduction selection in emitReduction.");
    (void)ParallelReduction;
    CGOpenMPRuntime::emitReduction(CGF, Loc, Privates, LHSExprs, RHSExprs,
                                   ReductionOps, Options);
    return;
  }

  llvm::SmallDenseMap<const ValueDecl *, const FieldDecl *> VarFieldMap;
  llvm::SmallVector<const ValueDecl *, 4> PrivatesReductions(Privates.size());
  int Cnt = 0;
  for (const Expr *DRE : Privates) {
    PrivatesReductions[Cnt] = cast<DeclRefExpr>(DRE)->getDecl();
    ++Cnt;
  }
  const RecordDecl *ReductionRec = ::buildRecordForGlobalizedVars(
      CGM.getContext(), PrivatesReductions, {}, VarFieldMap, 1);

  if (!ParallelReduction)
    TeamsReductions.push_back(ReductionRec);

  // Source location for the ident struct
  llvm::Value *RTLoc = emitUpdateLocation(CGF, Loc);

  using InsertPointTy = llvm::OpenMPIRBuilder::InsertPointTy;
  InsertPointTy AllocaIP(CGF.AllocaInsertPt->getParent(),
                         CGF.AllocaInsertPt->getIterator());
  InsertPointTy CodeGenIP(CGF.Builder.GetInsertBlock(),
                          CGF.Builder.GetInsertPoint());
  llvm::OpenMPIRBuilder::LocationDescription OmpLoc(
      CodeGenIP, CGF.SourceLocToDebugLoc(Loc));
  llvm::SmallVector<llvm::OpenMPIRBuilder::ReductionInfo, 2> ReductionInfos;

  CodeGenFunction::OMPPrivateScope Scope(CGF);
  unsigned Idx = 0;
  for (const Expr *Private : Privates) {
    llvm::Type *ElementType;
    llvm::Value *Variable;
    llvm::Value *PrivateVariable;
    llvm::OpenMPIRBuilder::ReductionGenAtomicCBTy AtomicReductionGen = nullptr;
    ElementType = CGF.ConvertTypeForMem(Private->getType());
    const auto *RHSVar =
        cast<VarDecl>(cast<DeclRefExpr>(RHSExprs[Idx])->getDecl());
    PrivateVariable = CGF.GetAddrOfLocalVar(RHSVar).emitRawPointer(CGF);
    const auto *LHSVar =
        cast<VarDecl>(cast<DeclRefExpr>(LHSExprs[Idx])->getDecl());
    Variable = CGF.GetAddrOfLocalVar(LHSVar).emitRawPointer(CGF);
    llvm::OpenMPIRBuilder::EvalKind EvalKind;
    switch (CGF.getEvaluationKind(Private->getType())) {
    case TEK_Scalar:
      EvalKind = llvm::OpenMPIRBuilder::EvalKind::Scalar;
      break;
    case TEK_Complex:
      EvalKind = llvm::OpenMPIRBuilder::EvalKind::Complex;
      break;
    case TEK_Aggregate:
      EvalKind = llvm::OpenMPIRBuilder::EvalKind::Aggregate;
      break;
    }
    auto ReductionGen = [&](InsertPointTy CodeGenIP, unsigned I,
                            llvm::Value **LHSPtr, llvm::Value **RHSPtr,
                            llvm::Function *NewFunc) {
      CGF.Builder.restoreIP(CodeGenIP);
      auto *CurFn = CGF.CurFn;
      CGF.CurFn = NewFunc;

      *LHSPtr = CGF.GetAddrOfLocalVar(
                       cast<VarDecl>(cast<DeclRefExpr>(LHSExprs[I])->getDecl()))
                    .emitRawPointer(CGF);
      *RHSPtr = CGF.GetAddrOfLocalVar(
                       cast<VarDecl>(cast<DeclRefExpr>(RHSExprs[I])->getDecl()))
                    .emitRawPointer(CGF);

      emitSingleReductionCombiner(CGF, ReductionOps[I], Privates[I],
                                  cast<DeclRefExpr>(LHSExprs[I]),
                                  cast<DeclRefExpr>(RHSExprs[I]));

      CGF.CurFn = CurFn;

      return InsertPointTy(CGF.Builder.GetInsertBlock(),
                           CGF.Builder.GetInsertPoint());
    };
    ReductionInfos.emplace_back(llvm::OpenMPIRBuilder::ReductionInfo(
        ElementType, Variable, PrivateVariable, EvalKind,
        /*ReductionGen=*/nullptr, ReductionGen, AtomicReductionGen,
        /*DataPtrPtrGen=*/nullptr));
    Idx++;
  }

  llvm::OpenMPIRBuilder::InsertPointTy AfterIP =
      cantFail(OMPBuilder.createReductionsGPU(
          OmpLoc, AllocaIP, CodeGenIP, ReductionInfos, /*IsByRef=*/{}, false,
          TeamsReduction, llvm::OpenMPIRBuilder::ReductionGenCBKind::Clang,
          CGF.getTarget().getGridValue(),
          C.getLangOpts().OpenMPCUDAReductionBufNum, RTLoc));
  CGF.Builder.restoreIP(AfterIP);
}

const VarDecl *
CGOpenMPRuntimeGPU::translateParameter(const FieldDecl *FD,
                                       const VarDecl *NativeParam) const {
  if (!NativeParam->getType()->isReferenceType())
    return NativeParam;
  QualType ArgType = NativeParam->getType();
  QualifierCollector QC;
  const Type *NonQualTy = QC.strip(ArgType);
  QualType PointeeTy = cast<ReferenceType>(NonQualTy)->getPointeeType();
  if (const auto *Attr = FD->getAttr<OMPCaptureKindAttr>()) {
    if (Attr->getCaptureKind() == OMPC_map) {
      PointeeTy = CGM.getContext().getAddrSpaceQualType(PointeeTy,
                                                        LangAS::opencl_global);
    }
  }
  ArgType = CGM.getContext().getPointerType(PointeeTy);
  QC.addRestrict();
  ArgType = QC.apply(CGM.getContext(), ArgType);
  if (isa<ImplicitParamDecl>(NativeParam))
    return ImplicitParamDecl::Create(
        CGM.getContext(), /*DC=*/nullptr, NativeParam->getLocation(),
        NativeParam->getIdentifier(), ArgType, ImplicitParamKind::Other);
  return ParmVarDecl::Create(
      CGM.getContext(),
      const_cast<DeclContext *>(NativeParam->getDeclContext()),
      NativeParam->getBeginLoc(), NativeParam->getLocation(),
      NativeParam->getIdentifier(), ArgType,
      /*TInfo=*/nullptr, SC_None, /*DefArg=*/nullptr);
}

Address
CGOpenMPRuntimeGPU::getParameterAddress(CodeGenFunction &CGF,
                                          const VarDecl *NativeParam,
                                          const VarDecl *TargetParam) const {
  assert(NativeParam != TargetParam &&
         NativeParam->getType()->isReferenceType() &&
         "Native arg must not be the same as target arg.");
  Address LocalAddr = CGF.GetAddrOfLocalVar(TargetParam);
  QualType NativeParamType = NativeParam->getType();
  QualifierCollector QC;
  const Type *NonQualTy = QC.strip(NativeParamType);
  QualType NativePointeeTy = cast<ReferenceType>(NonQualTy)->getPointeeType();
  unsigned NativePointeeAddrSpace =
      CGF.getTypes().getTargetAddressSpace(NativePointeeTy);
  QualType TargetTy = TargetParam->getType();
  llvm::Value *TargetAddr = CGF.EmitLoadOfScalar(LocalAddr, /*Volatile=*/false,
                                                 TargetTy, SourceLocation());
  // Cast to native address space.
  TargetAddr = CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
      TargetAddr,
      llvm::PointerType::get(CGF.getLLVMContext(), NativePointeeAddrSpace));
  Address NativeParamAddr = CGF.CreateMemTemp(NativeParamType);
  CGF.EmitStoreOfScalar(TargetAddr, NativeParamAddr, /*Volatile=*/false,
                        NativeParamType);
  return NativeParamAddr;
}

void CGOpenMPRuntimeGPU::emitOutlinedFunctionCall(
    CodeGenFunction &CGF, SourceLocation Loc, llvm::FunctionCallee OutlinedFn,
    ArrayRef<llvm::Value *> Args) const {
  SmallVector<llvm::Value *, 4> TargetArgs;
  TargetArgs.reserve(Args.size());
  auto *FnType = OutlinedFn.getFunctionType();
  for (unsigned I = 0, E = Args.size(); I < E; ++I) {
    if (FnType->isVarArg() && FnType->getNumParams() <= I) {
      TargetArgs.append(std::next(Args.begin(), I), Args.end());
      break;
    }
    llvm::Type *TargetType = FnType->getParamType(I);
    llvm::Value *NativeArg = Args[I];
    if (!TargetType->isPointerTy()) {
      TargetArgs.emplace_back(NativeArg);
      continue;
    }
    TargetArgs.emplace_back(
        CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(NativeArg, TargetType));
  }
  CGOpenMPRuntime::emitOutlinedFunctionCall(CGF, Loc, OutlinedFn, TargetArgs);
}

/// Emit function which wraps the outline parallel region
/// and controls the arguments which are passed to this function.
/// The wrapper ensures that the outlined function is called
/// with the correct arguments when data is shared.
llvm::Function *CGOpenMPRuntimeGPU::createParallelDataSharingWrapper(
    llvm::Function *OutlinedParallelFn, const OMPExecutableDirective &D) {
  ASTContext &Ctx = CGM.getContext();
  const auto &CS = *D.getCapturedStmt(OMPD_parallel);

  // Create a function that takes as argument the source thread.
  FunctionArgList WrapperArgs;
  QualType Int16QTy =
      Ctx.getIntTypeForBitwidth(/*DestWidth=*/16, /*Signed=*/false);
  QualType Int32QTy =
      Ctx.getIntTypeForBitwidth(/*DestWidth=*/32, /*Signed=*/false);
  auto *ParallelLevelArg = ImplicitParamDecl::Create(
      Ctx, /*DC=*/nullptr, D.getBeginLoc(),
      /*Id=*/nullptr, Int16QTy, ImplicitParamKind::Other);
  auto *WrapperArg = ImplicitParamDecl::Create(
      Ctx, /*DC=*/nullptr, D.getBeginLoc(),
      /*Id=*/nullptr, Int32QTy, ImplicitParamKind::Other);
  WrapperArgs.emplace_back(ParallelLevelArg);
  WrapperArgs.emplace_back(WrapperArg);

  const CGFunctionInfo &CGFI =
      CGM.getTypes().arrangeBuiltinFunctionDeclaration(Ctx.VoidTy, WrapperArgs);

  auto *Fn = llvm::Function::Create(
      CGM.getTypes().GetFunctionType(CGFI), llvm::GlobalValue::InternalLinkage,
      Twine(OutlinedParallelFn->getName(), "_wrapper"), &CGM.getModule());

  // Ensure we do not inline the function. This is trivially true for the ones
  // passed to __kmpc_fork_call but the ones calles in serialized regions
  // could be inlined. This is not a perfect but it is closer to the invariant
  // we want, namely, every data environment starts with a new function.
  // TODO: We should pass the if condition to the runtime function and do the
  //       handling there. Much cleaner code.
  Fn->addFnAttr(llvm::Attribute::NoInline);

  CGM.SetInternalFunctionAttributes(GlobalDecl(), Fn, CGFI);
  Fn->setLinkage(llvm::GlobalValue::InternalLinkage);

  Fn->setDoesNotRecurse();

  CodeGenFunction CGF(CGM, /*suppressNewContext=*/true);
  CGF.StartFunction(GlobalDecl(), Ctx.VoidTy, Fn, CGFI, WrapperArgs,
                    D.getBeginLoc(), D.getBeginLoc());

  const auto *RD = CS.getCapturedRecordDecl();
  auto CurField = RD->field_begin();

  Address ZeroAddr = CGF.CreateDefaultAlignTempAlloca(CGF.Int32Ty,
                                                      /*Name=*/".zero.addr");
  CGF.Builder.CreateStore(CGF.Builder.getInt32(/*C*/ 0), ZeroAddr);
  // Get the array of arguments.
  SmallVector<llvm::Value *, 8> Args;

  Args.emplace_back(CGF.GetAddrOfLocalVar(WrapperArg).emitRawPointer(CGF));
  Args.emplace_back(ZeroAddr.emitRawPointer(CGF));

  CGBuilderTy &Bld = CGF.Builder;
  auto CI = CS.capture_begin();

  // Use global memory for data sharing.
  // Handle passing of global args to workers.
  RawAddress GlobalArgs =
      CGF.CreateDefaultAlignTempAlloca(CGF.VoidPtrPtrTy, "global_args");
  llvm::Value *GlobalArgsPtr = GlobalArgs.getPointer();
  llvm::Value *DataSharingArgs[] = {GlobalArgsPtr};
  CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
                          CGM.getModule(), OMPRTL___kmpc_get_shared_variables),
                      DataSharingArgs);

  // Retrieve the shared variables from the list of references returned
  // by the runtime. Pass the variables to the outlined function.
  Address SharedArgListAddress = Address::invalid();
  if (CS.capture_size() > 0 ||
      isOpenMPLoopBoundSharingDirective(D.getDirectiveKind())) {
    SharedArgListAddress = CGF.EmitLoadOfPointer(
        GlobalArgs, CGF.getContext()
                        .getPointerType(CGF.getContext().VoidPtrTy)
                        .castAs<PointerType>());
  }
  unsigned Idx = 0;
  if (isOpenMPLoopBoundSharingDirective(D.getDirectiveKind())) {
    Address Src = Bld.CreateConstInBoundsGEP(SharedArgListAddress, Idx);
    Address TypedAddress = Bld.CreatePointerBitCastOrAddrSpaceCast(
        Src, Bld.getPtrTy(0), CGF.SizeTy);
    llvm::Value *LB = CGF.EmitLoadOfScalar(
        TypedAddress,
        /*Volatile=*/false,
        CGF.getContext().getPointerType(CGF.getContext().getSizeType()),
        cast<OMPLoopDirective>(D).getLowerBoundVariable()->getExprLoc());
    Args.emplace_back(LB);
    ++Idx;
    Src = Bld.CreateConstInBoundsGEP(SharedArgListAddress, Idx);
    TypedAddress = Bld.CreatePointerBitCastOrAddrSpaceCast(Src, Bld.getPtrTy(0),
                                                           CGF.SizeTy);
    llvm::Value *UB = CGF.EmitLoadOfScalar(
        TypedAddress,
        /*Volatile=*/false,
        CGF.getContext().getPointerType(CGF.getContext().getSizeType()),
        cast<OMPLoopDirective>(D).getUpperBoundVariable()->getExprLoc());
    Args.emplace_back(UB);
    ++Idx;
  }
  if (CS.capture_size() > 0) {
    ASTContext &CGFContext = CGF.getContext();
    for (unsigned I = 0, E = CS.capture_size(); I < E; ++I, ++CI, ++CurField) {
      QualType ElemTy = CurField->getType();
      Address Src = Bld.CreateConstInBoundsGEP(SharedArgListAddress, I + Idx);
      Address TypedAddress = Bld.CreatePointerBitCastOrAddrSpaceCast(
          Src, CGF.ConvertTypeForMem(CGFContext.getPointerType(ElemTy)),
          CGF.ConvertTypeForMem(ElemTy));
      llvm::Value *Arg = CGF.EmitLoadOfScalar(TypedAddress,
                                              /*Volatile=*/false,
                                              CGFContext.getPointerType(ElemTy),
                                              CI->getLocation());
      if (CI->capturesVariableByCopy() &&
          !CI->getCapturedVar()->getType()->isAnyPointerType()) {
        Arg = castValueToType(CGF, Arg, ElemTy, CGFContext.getUIntPtrType(),
                              CI->getLocation());
      }
      Args.emplace_back(Arg);
    }
  }

  emitOutlinedFunctionCall(CGF, D.getBeginLoc(), OutlinedParallelFn, Args);
  CGF.FinishFunction();
  return Fn;
}

void CGOpenMPRuntimeGPU::emitFunctionProlog(CodeGenFunction &CGF,
                                              const Decl *D) {
  if (getDataSharingMode() != CGOpenMPRuntimeGPU::DS_Generic)
    return;

  assert(D && "Expected function or captured|block decl.");
  assert(FunctionGlobalizedDecls.count(CGF.CurFn) == 0 &&
         "Function is registered already.");
  assert((!TeamAndReductions.first || TeamAndReductions.first == D) &&
         "Team is set but not processed.");
  const Stmt *Body = nullptr;
  bool NeedToDelayGlobalization = false;
  if (const auto *FD = dyn_cast<FunctionDecl>(D)) {
    Body = FD->getBody();
  } else if (const auto *BD = dyn_cast<BlockDecl>(D)) {
    Body = BD->getBody();
  } else if (const auto *CD = dyn_cast<CapturedDecl>(D)) {
    Body = CD->getBody();
    NeedToDelayGlobalization = CGF.CapturedStmtInfo->getKind() == CR_OpenMP;
    if (NeedToDelayGlobalization &&
        getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD)
      return;
  }
  if (!Body)
    return;
  CheckVarsEscapingDeclContext VarChecker(CGF, TeamAndReductions.second);
  VarChecker.Visit(Body);
  const RecordDecl *GlobalizedVarsRecord =
      VarChecker.getGlobalizedRecord(IsInTTDRegion);
  TeamAndReductions.first = nullptr;
  TeamAndReductions.second.clear();
  ArrayRef<const ValueDecl *> EscapedVariableLengthDecls =
      VarChecker.getEscapedVariableLengthDecls();
  ArrayRef<const ValueDecl *> DelayedVariableLengthDecls =
      VarChecker.getDelayedVariableLengthDecls();
  if (!GlobalizedVarsRecord && EscapedVariableLengthDecls.empty() &&
      DelayedVariableLengthDecls.empty())
    return;
  auto I = FunctionGlobalizedDecls.try_emplace(CGF.CurFn).first;
  I->getSecond().MappedParams =
      std::make_unique<CodeGenFunction::OMPMapVars>();
  I->getSecond().EscapedParameters.insert(
      VarChecker.getEscapedParameters().begin(),
      VarChecker.getEscapedParameters().end());
  I->getSecond().EscapedVariableLengthDecls.append(
      EscapedVariableLengthDecls.begin(), EscapedVariableLengthDecls.end());
  I->getSecond().DelayedVariableLengthDecls.append(
      DelayedVariableLengthDecls.begin(), DelayedVariableLengthDecls.end());
  DeclToAddrMapTy &Data = I->getSecond().LocalVarData;
  for (const ValueDecl *VD : VarChecker.getEscapedDecls()) {
    assert(VD->isCanonicalDecl() && "Expected canonical declaration");
    Data.try_emplace(VD);
  }
  if (!NeedToDelayGlobalization) {
    emitGenericVarsProlog(CGF, D->getBeginLoc());
    struct GlobalizationScope final : EHScopeStack::Cleanup {
      GlobalizationScope() = default;

      void Emit(CodeGenFunction &CGF, Flags flags) override {
        static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime())
            .emitGenericVarsEpilog(CGF);
      }
    };
    CGF.EHStack.pushCleanup<GlobalizationScope>(NormalAndEHCleanup);
  }
}

Address CGOpenMPRuntimeGPU::getAddressOfLocalVariable(CodeGenFunction &CGF,
                                                        const VarDecl *VD) {
  if (VD && VD->hasAttr<OMPAllocateDeclAttr>()) {
    const auto *A = VD->getAttr<OMPAllocateDeclAttr>();
    auto AS = LangAS::Default;
    switch (A->getAllocatorType()) {
    case OMPAllocateDeclAttr::OMPNullMemAlloc:
    case OMPAllocateDeclAttr::OMPDefaultMemAlloc:
    case OMPAllocateDeclAttr::OMPHighBWMemAlloc:
    case OMPAllocateDeclAttr::OMPLowLatMemAlloc:
      break;
    case OMPAllocateDeclAttr::OMPThreadMemAlloc:
      return Address::invalid();
    case OMPAllocateDeclAttr::OMPUserDefinedMemAlloc:
      // TODO: implement aupport for user-defined allocators.
      return Address::invalid();
    case OMPAllocateDeclAttr::OMPConstMemAlloc:
      AS = LangAS::cuda_constant;
      break;
    case OMPAllocateDeclAttr::OMPPTeamMemAlloc:
      AS = LangAS::cuda_shared;
      break;
    case OMPAllocateDeclAttr::OMPLargeCapMemAlloc:
    case OMPAllocateDeclAttr::OMPCGroupMemAlloc:
      break;
    }
    llvm::Type *VarTy = CGF.ConvertTypeForMem(VD->getType());
    auto *GV = new llvm::GlobalVariable(
        CGM.getModule(), VarTy, /*isConstant=*/false,
        llvm::GlobalValue::InternalLinkage, llvm::PoisonValue::get(VarTy),
        VD->getName(),
        /*InsertBefore=*/nullptr, llvm::GlobalValue::NotThreadLocal,
        CGM.getContext().getTargetAddressSpace(AS));
    CharUnits Align = CGM.getContext().getDeclAlign(VD);
    GV->setAlignment(Align.getAsAlign());
    return Address(
        CGF.Builder.CreatePointerBitCastOrAddrSpaceCast(
            GV, CGF.Builder.getPtrTy(CGM.getContext().getTargetAddressSpace(
                    VD->getType().getAddressSpace()))),
        VarTy, Align);
  }

  if (getDataSharingMode() != CGOpenMPRuntimeGPU::DS_Generic)
    return Address::invalid();

  VD = VD->getCanonicalDecl();
  auto I = FunctionGlobalizedDecls.find(CGF.CurFn);
  if (I == FunctionGlobalizedDecls.end())
    return Address::invalid();
  auto VDI = I->getSecond().LocalVarData.find(VD);
  if (VDI != I->getSecond().LocalVarData.end())
    return VDI->second.PrivateAddr;
  if (VD->hasAttrs()) {
    for (specific_attr_iterator<OMPReferencedVarAttr> IT(VD->attr_begin()),
         E(VD->attr_end());
         IT != E; ++IT) {
      auto VDI = I->getSecond().LocalVarData.find(
          cast<VarDecl>(cast<DeclRefExpr>(IT->getRef())->getDecl())
              ->getCanonicalDecl());
      if (VDI != I->getSecond().LocalVarData.end())
        return VDI->second.PrivateAddr;
    }
  }

  return Address::invalid();
}

void CGOpenMPRuntimeGPU::functionFinished(CodeGenFunction &CGF) {
  FunctionGlobalizedDecls.erase(CGF.CurFn);
  CGOpenMPRuntime::functionFinished(CGF);
}

void CGOpenMPRuntimeGPU::getDefaultDistScheduleAndChunk(
    CodeGenFunction &CGF, const OMPLoopDirective &S,
    OpenMPDistScheduleClauseKind &ScheduleKind,
    llvm::Value *&Chunk) const {
  auto &RT = static_cast<CGOpenMPRuntimeGPU &>(CGF.CGM.getOpenMPRuntime());
  if (getExecutionMode() == CGOpenMPRuntimeGPU::EM_SPMD) {
    ScheduleKind = OMPC_DIST_SCHEDULE_static;
    Chunk = CGF.EmitScalarConversion(
        RT.getGPUNumThreads(CGF),
        CGF.getContext().getIntTypeForBitwidth(32, /*Signed=*/0),
        S.getIterationVariable()->getType(), S.getBeginLoc());
    return;
  }
  CGOpenMPRuntime::getDefaultDistScheduleAndChunk(
      CGF, S, ScheduleKind, Chunk);
}

void CGOpenMPRuntimeGPU::getDefaultScheduleAndChunk(
    CodeGenFunction &CGF, const OMPLoopDirective &S,
    OpenMPScheduleClauseKind &ScheduleKind,
    const Expr *&ChunkExpr) const {
  ScheduleKind = OMPC_SCHEDULE_static;
  // Chunk size is 1 in this case.
  llvm::APInt ChunkSize(32, 1);
  ChunkExpr = IntegerLiteral::Create(CGF.getContext(), ChunkSize,
      CGF.getContext().getIntTypeForBitwidth(32, /*Signed=*/0),
      SourceLocation());
}

void CGOpenMPRuntimeGPU::adjustTargetSpecificDataForLambdas(
    CodeGenFunction &CGF, const OMPExecutableDirective &D) const {
  assert(isOpenMPTargetExecutionDirective(D.getDirectiveKind()) &&
         " Expected target-based directive.");
  const CapturedStmt *CS = D.getCapturedStmt(OMPD_target);
  for (const CapturedStmt::Capture &C : CS->captures()) {
    // Capture variables captured by reference in lambdas for target-based
    // directives.
    if (!C.capturesVariable())
      continue;
    const VarDecl *VD = C.getCapturedVar();
    const auto *RD = VD->getType()
                         .getCanonicalType()
                         .getNonReferenceType()
                         ->getAsCXXRecordDecl();
    if (!RD || !RD->isLambda())
      continue;
    Address VDAddr = CGF.GetAddrOfLocalVar(VD);
    LValue VDLVal;
    if (VD->getType().getCanonicalType()->isReferenceType())
      VDLVal = CGF.EmitLoadOfReferenceLValue(VDAddr, VD->getType());
    else
      VDLVal = CGF.MakeAddrLValue(
          VDAddr, VD->getType().getCanonicalType().getNonReferenceType());
    llvm::DenseMap<const ValueDecl *, FieldDecl *> Captures;
    FieldDecl *ThisCapture = nullptr;
    RD->getCaptureFields(Captures, ThisCapture);
    if (ThisCapture && CGF.CapturedStmtInfo->isCXXThisExprCaptured()) {
      LValue ThisLVal =
          CGF.EmitLValueForFieldInitialization(VDLVal, ThisCapture);
      llvm::Value *CXXThis = CGF.LoadCXXThis();
      CGF.EmitStoreOfScalar(CXXThis, ThisLVal);
    }
    for (const LambdaCapture &LC : RD->captures()) {
      if (LC.getCaptureKind() != LCK_ByRef)
        continue;
      const ValueDecl *VD = LC.getCapturedVar();
      // FIXME: For now VD is always a VarDecl because OpenMP does not support
      //  capturing structured bindings in lambdas yet.
      if (!CS->capturesVariable(cast<VarDecl>(VD)))
        continue;
      auto It = Captures.find(VD);
      assert(It != Captures.end() && "Found lambda capture without field.");
      LValue VarLVal = CGF.EmitLValueForFieldInitialization(VDLVal, It->second);
      Address VDAddr = CGF.GetAddrOfLocalVar(cast<VarDecl>(VD));
      if (VD->getType().getCanonicalType()->isReferenceType())
        VDAddr = CGF.EmitLoadOfReferenceLValue(VDAddr,
                                               VD->getType().getCanonicalType())
                     .getAddress();
      CGF.EmitStoreOfScalar(VDAddr.emitRawPointer(CGF), VarLVal);
    }
  }
}

bool CGOpenMPRuntimeGPU::hasAllocateAttributeForGlobalVar(const VarDecl *VD,
                                                            LangAS &AS) {
  if (!VD || !VD->hasAttr<OMPAllocateDeclAttr>())
    return false;
  const auto *A = VD->getAttr<OMPAllocateDeclAttr>();
  switch(A->getAllocatorType()) {
  case OMPAllocateDeclAttr::OMPNullMemAlloc:
  case OMPAllocateDeclAttr::OMPDefaultMemAlloc:
  // Not supported, fallback to the default mem space.
  case OMPAllocateDeclAttr::OMPLargeCapMemAlloc:
  case OMPAllocateDeclAttr::OMPCGroupMemAlloc:
  case OMPAllocateDeclAttr::OMPHighBWMemAlloc:
  case OMPAllocateDeclAttr::OMPLowLatMemAlloc:
  case OMPAllocateDeclAttr::OMPThreadMemAlloc:
    AS = LangAS::Default;
    return true;
  case OMPAllocateDeclAttr::OMPConstMemAlloc:
    AS = LangAS::cuda_constant;
    return true;
  case OMPAllocateDeclAttr::OMPPTeamMemAlloc:
    AS = LangAS::cuda_shared;
    return true;
  case OMPAllocateDeclAttr::OMPUserDefinedMemAlloc:
    llvm_unreachable("Expected predefined allocator for the variables with the "
                     "static storage.");
  }
  return false;
}

static OffloadArch getOffloadArch(const CodeGenModule &CGM) {
  // FIXME: This should not require parsing
  return StringToOffloadArch(CGM.getTarget().getTargetOpts().CPU);
}

/// Check to see if target architecture supports unified addressing which is
/// a restriction for OpenMP requires clause "unified_shared_memory".
void CGOpenMPRuntimeGPU::processRequiresDirective(const OMPRequiresDecl *D) {
  for (const OMPClause *Clause : D->clauselists()) {
    if (Clause->getClauseKind() == OMPC_unified_shared_memory ||
        Clause->getClauseKind() == OMPC_unified_address) {
      OffloadArch Arch = getOffloadArch(CGM);
      switch (Arch) {
      case OffloadArch::SM_20:
      case OffloadArch::SM_21:
      case OffloadArch::SM_30:
      case OffloadArch::SM_32_:
      case OffloadArch::SM_35:
      case OffloadArch::SM_37:
      case OffloadArch::SM_50:
      case OffloadArch::SM_52:
      case OffloadArch::SM_53: {
        SmallString<256> Buffer;
        llvm::raw_svector_ostream Out(Buffer);
        Out << "Target architecture " << OffloadArchToString(Arch)
            << " does not support unified addressing";
        CGM.Error(Clause->getBeginLoc(), Out.str());
        return;
      }
      case OffloadArch::SM_60:
      case OffloadArch::SM_61:
      case OffloadArch::SM_62:
      case OffloadArch::SM_70:
      case OffloadArch::SM_72:
      case OffloadArch::SM_75:
      case OffloadArch::SM_80:
      case OffloadArch::SM_86:
      case OffloadArch::SM_87:
      case OffloadArch::SM_88:
      case OffloadArch::SM_89:
      case OffloadArch::SM_90:
      case OffloadArch::SM_90a:
      case OffloadArch::SM_100:
      case OffloadArch::SM_100a:
      case OffloadArch::SM_100f:
      case OffloadArch::SM_101:
      case OffloadArch::SM_101a:
      case OffloadArch::SM_101f:
      case OffloadArch::SM_103:
      case OffloadArch::SM_103a:
      case OffloadArch::SM_103f:
      case OffloadArch::SM_110:
      case OffloadArch::SM_110a:
      case OffloadArch::SM_110f:
      case OffloadArch::SM_120:
      case OffloadArch::SM_120a:
      case OffloadArch::SM_120f:
      case OffloadArch::SM_121:
      case OffloadArch::SM_121a:
      case OffloadArch::SM_121f:
      case OffloadArch::GFX600:
      case OffloadArch::GFX601:
      case OffloadArch::GFX602:
      case OffloadArch::GFX700:
      case OffloadArch::GFX701:
      case OffloadArch::GFX702:
      case OffloadArch::GFX703:
      case OffloadArch::GFX704:
      case OffloadArch::GFX705:
      case OffloadArch::GFX801:
      case OffloadArch::GFX802:
      case OffloadArch::GFX803:
      case OffloadArch::GFX805:
      case OffloadArch::GFX810:
      case OffloadArch::GFX9_GENERIC:
      case OffloadArch::GFX900:
      case OffloadArch::GFX902:
      case OffloadArch::GFX904:
      case OffloadArch::GFX906:
      case OffloadArch::GFX908:
      case OffloadArch::GFX909:
      case OffloadArch::GFX90a:
      case OffloadArch::GFX90c:
      case OffloadArch::GFX9_4_GENERIC:
      case OffloadArch::GFX942:
      case OffloadArch::GFX950:
      case OffloadArch::GFX10_1_GENERIC:
      case OffloadArch::GFX1010:
      case OffloadArch::GFX1011:
      case OffloadArch::GFX1012:
      case OffloadArch::GFX1013:
      case OffloadArch::GFX10_3_GENERIC:
      case OffloadArch::GFX1030:
      case OffloadArch::GFX1031:
      case OffloadArch::GFX1032:
      case OffloadArch::GFX1033:
      case OffloadArch::GFX1034:
      case OffloadArch::GFX1035:
      case OffloadArch::GFX1036:
      case OffloadArch::GFX11_GENERIC:
      case OffloadArch::GFX1100:
      case OffloadArch::GFX1101:
      case OffloadArch::GFX1102:
      case OffloadArch::GFX1103:
      case OffloadArch::GFX1150:
      case OffloadArch::GFX1151:
      case OffloadArch::GFX1152:
      case OffloadArch::GFX1153:
      case OffloadArch::GFX1170:
      case OffloadArch::GFX1171:
      case OffloadArch::GFX1172:
      case OffloadArch::GFX12_GENERIC:
      case OffloadArch::GFX1200:
      case OffloadArch::GFX1201:
      case OffloadArch::GFX1250:
      case OffloadArch::GFX1251:
      case OffloadArch::GFX12_5_GENERIC:
      case OffloadArch::GFX1310:
      case OffloadArch::AMDGCNSPIRV:
      case OffloadArch::Generic:
      case OffloadArch::GRANITERAPIDS:
      case OffloadArch::BMG_G21:
      case OffloadArch::Unused:
      case OffloadArch::Unknown:
        break;
      }
    }
  }
  CGOpenMPRuntime::processRequiresDirective(D);
}

llvm::Value *CGOpenMPRuntimeGPU::getGPUNumThreads(CodeGenFunction &CGF) {
  CGBuilderTy &Bld = CGF.Builder;
  llvm::Module *M = &CGF.CGM.getModule();
  const char *LocSize = "__kmpc_get_hardware_num_threads_in_block";
  llvm::Function *F = M->getFunction(LocSize);
  if (!F) {
    F = llvm::Function::Create(llvm::FunctionType::get(CGF.Int32Ty, {}, false),
                               llvm::GlobalVariable::ExternalLinkage, LocSize,
                               &CGF.CGM.getModule());
  }
  return Bld.CreateCall(F, {}, "nvptx_num_threads");
}

llvm::Value *CGOpenMPRuntimeGPU::getGPUThreadID(CodeGenFunction &CGF) {
  ArrayRef<llvm::Value *> Args{};
  return CGF.EmitRuntimeCall(
      OMPBuilder.getOrCreateRuntimeFunction(
          CGM.getModule(), OMPRTL___kmpc_get_hardware_thread_id_in_block),
      Args);
}
llvm::Value *CGOpenMPRuntimeGPU::getGPUBlockID(CodeGenFunction &CGF) {
  CGBuilderTy &Bld = CGF.Builder;
  llvm::Function *F =
      CGF.CGM.getIntrinsic(llvm::Intrinsic::amdgcn_workgroup_id_x);
  return Bld.CreateCall(F, {}, "gpu_block_id");
}

llvm::Value *CGOpenMPRuntimeGPU::getGPUNumBlocks(CodeGenFunction &CGF) {
  return CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
      CGM.getModule(), OMPRTL___kmpc_get_hardware_num_blocks));
}

llvm::Value *CGOpenMPRuntimeGPU::initSpecializedKernel(CodeGenFunction &CGF) {
  return CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
      CGM.getModule(), OMPRTL___kmpc_specialized_kernel_init));
}

std::pair<llvm::Value *, llvm::Value *>
CGOpenMPRuntimeGPU::getXteamRedFunctionPtrs(
    CodeGenFunction &CGF, llvm::Type *RedVarType,
    CodeGenModule::XteamRedOpKind Opcode) {
  if (RedVarType->isIntegerTy()) {
    if (RedVarType->getPrimitiveSizeInBits() == 16) {
      switch (Opcode) {
      case CodeGenModule::XR_OP_unknown:
        llvm_unreachable("Xteam reduction opcode cannot be unknown");
      case CodeGenModule::XR_OP_add:
        return std::make_pair(
            OMPBuilder
                .getOrCreateRuntimeFunction(CGM.getModule(),
                                            OMPRTL___kmpc_rfun_sum_s)
                .getCallee(),
            OMPBuilder
                .getOrCreateRuntimeFunction(CGM.getModule(),
                                            OMPRTL___kmpc_rfun_sum_lds_s)
                .getCallee());
      case CodeGenModule::XR_OP_min:
        return std::make_pair(
            OMPBuilder
                .getOrCreateRuntimeFunction(CGM.getModule(),
                                            OMPRTL___kmpc_rfun_min_s)
                .getCallee(),
            OMPBuilder
                .getOrCreateRuntimeFunction(CGM.getModule(),
                                            OMPRTL___kmpc_rfun_min_lds_s)
                .getCallee());
      case CodeGenModule::XR_OP_max:
        return std::make_pair(
            OMPBuilder
                .getOrCreateRuntimeFunction(CGM.getModule(),
                                            OMPRTL___kmpc_rfun_max_s)
                .getCallee(),
            OMPBuilder
                .getOrCreateRuntimeFunction(CGM.getModule(),
                                            OMPRTL___kmpc_rfun_max_lds_s)
                .getCallee());
      }
    }
    if (RedVarType->getPrimitiveSizeInBits() == 32) {
      switch (Opcode) {
      case CodeGenModule::XR_OP_unknown:
        llvm_unreachable("Xteam reduction opcode cannot be unknown");
      case CodeGenModule::XR_OP_add:
        return std::make_pair(
            OMPBuilder
                .getOrCreateRuntimeFunction(CGM.getModule(),
                                            OMPRTL___kmpc_rfun_sum_i)
                .getCallee(),
            OMPBuilder
                .getOrCreateRuntimeFunction(CGM.getModule(),
                                            OMPRTL___kmpc_rfun_sum_lds_i)
                .getCallee());
      case CodeGenModule::XR_OP_min:
        return std::make_pair(
            OMPBuilder
                .getOrCreateRuntimeFunction(CGM.getModule(),
                                            OMPRTL___kmpc_rfun_min_i)
                .getCallee(),
            OMPBuilder
                .getOrCreateRuntimeFunction(CGM.getModule(),
                                            OMPRTL___kmpc_rfun_min_lds_i)
                .getCallee());
      case CodeGenModule::XR_OP_max:
        return std::make_pair(
            OMPBuilder
                .getOrCreateRuntimeFunction(CGM.getModule(),
                                            OMPRTL___kmpc_rfun_max_i)
                .getCallee(),
            OMPBuilder
                .getOrCreateRuntimeFunction(CGM.getModule(),
                                            OMPRTL___kmpc_rfun_max_lds_i)
                .getCallee());
      }
    }
    if (RedVarType->getPrimitiveSizeInBits() == 64) {
      switch (Opcode) {
      case CodeGenModule::XR_OP_unknown:
        llvm_unreachable("Xteam reduction opcode cannot be unknown");
      case CodeGenModule::XR_OP_add:
        return std::make_pair(
            OMPBuilder
                .getOrCreateRuntimeFunction(CGM.getModule(),
                                            OMPRTL___kmpc_rfun_sum_l)
                .getCallee(),
            OMPBuilder
                .getOrCreateRuntimeFunction(CGM.getModule(),
                                            OMPRTL___kmpc_rfun_sum_lds_l)
                .getCallee());
      case CodeGenModule::XR_OP_min:
        return std::make_pair(
            OMPBuilder
                .getOrCreateRuntimeFunction(CGM.getModule(),
                                            OMPRTL___kmpc_rfun_min_l)
                .getCallee(),
            OMPBuilder
                .getOrCreateRuntimeFunction(CGM.getModule(),
                                            OMPRTL___kmpc_rfun_min_lds_l)
                .getCallee());
      case CodeGenModule::XR_OP_max:
        return std::make_pair(
            OMPBuilder
                .getOrCreateRuntimeFunction(CGM.getModule(),
                                            OMPRTL___kmpc_rfun_max_l)
                .getCallee(),
            OMPBuilder
                .getOrCreateRuntimeFunction(CGM.getModule(),
                                            OMPRTL___kmpc_rfun_max_lds_l)
                .getCallee());
      }
    }
  }

  if (RedVarType->isFloatTy()) {
    switch (Opcode) {
    case CodeGenModule::XR_OP_unknown:
      llvm_unreachable("Xteam reduction opcode cannot be unknown");
    case CodeGenModule::XR_OP_add:
      return std::make_pair(
          OMPBuilder
              .getOrCreateRuntimeFunction(CGM.getModule(),
                                          OMPRTL___kmpc_rfun_sum_f)
              .getCallee(),
          OMPBuilder
              .getOrCreateRuntimeFunction(CGM.getModule(),
                                          OMPRTL___kmpc_rfun_sum_lds_f)
              .getCallee());
    case CodeGenModule::XR_OP_min:
      return std::make_pair(
          OMPBuilder
              .getOrCreateRuntimeFunction(CGM.getModule(),
                                          OMPRTL___kmpc_rfun_min_f)
              .getCallee(),
          OMPBuilder
              .getOrCreateRuntimeFunction(CGM.getModule(),
                                          OMPRTL___kmpc_rfun_min_lds_f)
              .getCallee());
    case CodeGenModule::XR_OP_max:
      return std::make_pair(
          OMPBuilder
              .getOrCreateRuntimeFunction(CGM.getModule(),
                                          OMPRTL___kmpc_rfun_max_f)
              .getCallee(),
          OMPBuilder
              .getOrCreateRuntimeFunction(CGM.getModule(),
                                          OMPRTL___kmpc_rfun_max_lds_f)
              .getCallee());
    }
  }

  if (RedVarType->isDoubleTy()) {
    switch (Opcode) {
    case CodeGenModule::XR_OP_unknown:
      llvm_unreachable("Xteam reduction opcode cannot be unknown");
    case CodeGenModule::XR_OP_add:
      return std::make_pair(
          OMPBuilder
              .getOrCreateRuntimeFunction(CGM.getModule(),
                                          OMPRTL___kmpc_rfun_sum_d)
              .getCallee(),
          OMPBuilder
              .getOrCreateRuntimeFunction(CGM.getModule(),
                                          OMPRTL___kmpc_rfun_sum_lds_d)
              .getCallee());
    case CodeGenModule::XR_OP_min:
      return std::make_pair(
          OMPBuilder
              .getOrCreateRuntimeFunction(CGM.getModule(),
                                          OMPRTL___kmpc_rfun_min_d)
              .getCallee(),
          OMPBuilder
              .getOrCreateRuntimeFunction(CGM.getModule(),
                                          OMPRTL___kmpc_rfun_min_lds_d)
              .getCallee());
    case CodeGenModule::XR_OP_max:
      return std::make_pair(
          OMPBuilder
              .getOrCreateRuntimeFunction(CGM.getModule(),
                                          OMPRTL___kmpc_rfun_max_d)
              .getCallee(),
          OMPBuilder
              .getOrCreateRuntimeFunction(CGM.getModule(),
                                          OMPRTL___kmpc_rfun_max_lds_d)
              .getCallee());
    }
  }

  if (RedVarType->isHalfTy()) {
    switch (Opcode) {
    case CodeGenModule::XR_OP_unknown:
      llvm_unreachable("Xteam reduction opcode cannot be unknown");
    case CodeGenModule::XR_OP_add:
      return std::make_pair(
          OMPBuilder
              .getOrCreateRuntimeFunction(CGM.getModule(),
                                          OMPRTL___kmpc_rfun_sum_h)
              .getCallee(),
          OMPBuilder
              .getOrCreateRuntimeFunction(CGM.getModule(),
                                          OMPRTL___kmpc_rfun_sum_lds_h)
              .getCallee());
    case CodeGenModule::XR_OP_min:
      return std::make_pair(
          OMPBuilder
              .getOrCreateRuntimeFunction(CGM.getModule(),
                                          OMPRTL___kmpc_rfun_min_h)
              .getCallee(),
          OMPBuilder
              .getOrCreateRuntimeFunction(CGM.getModule(),
                                          OMPRTL___kmpc_rfun_min_lds_h)
              .getCallee());
    case CodeGenModule::XR_OP_max:
      return std::make_pair(
          OMPBuilder
              .getOrCreateRuntimeFunction(CGM.getModule(),
                                          OMPRTL___kmpc_rfun_max_h)
              .getCallee(),
          OMPBuilder
              .getOrCreateRuntimeFunction(CGM.getModule(),
                                          OMPRTL___kmpc_rfun_max_lds_h)
              .getCallee());
    }
  }

  if (RedVarType->isBFloatTy()) {
    switch (Opcode) {
    case CodeGenModule::XR_OP_unknown:
      llvm_unreachable("Xteam reduction opcode cannot be unknown");
    case CodeGenModule::XR_OP_add:
      return std::make_pair(
          OMPBuilder
              .getOrCreateRuntimeFunction(CGM.getModule(),
                                          OMPRTL___kmpc_rfun_sum_bf)
              .getCallee(),
          OMPBuilder
              .getOrCreateRuntimeFunction(CGM.getModule(),
                                          OMPRTL___kmpc_rfun_sum_lds_bf)
              .getCallee());
    case CodeGenModule::XR_OP_min:
      return std::make_pair(
          OMPBuilder
              .getOrCreateRuntimeFunction(CGM.getModule(),
                                          OMPRTL___kmpc_rfun_min_bf)
              .getCallee(),
          OMPBuilder
              .getOrCreateRuntimeFunction(CGM.getModule(),
                                          OMPRTL___kmpc_rfun_min_lds_bf)
              .getCallee());
    case CodeGenModule::XR_OP_max:
      return std::make_pair(
          OMPBuilder
              .getOrCreateRuntimeFunction(CGM.getModule(),
                                          OMPRTL___kmpc_rfun_max_bf)
              .getCallee(),
          OMPBuilder
              .getOrCreateRuntimeFunction(CGM.getModule(),
                                          OMPRTL___kmpc_rfun_max_lds_bf)
              .getCallee());
    }
  }
  llvm_unreachable("No support for other types currently.");
}

llvm::Value *CGOpenMPRuntimeGPU::getXteamRedOperation(
    CodeGenFunction &CGF, llvm::Value *Val, llvm::Value *OrigVarPtr,
    llvm::Value *DTeamVals, llvm::Value *DTeamsDonePtr,
    llvm::Value *ThreadStartIndex, llvm::Value *NumTeams, int BlockSize,
    CodeGenModule::XteamRedOpKind Opcode, bool IsFast) {
  // TODO handle more types
  llvm::Type *RedVarType = Val->getType();
  assert((RedVarType->isFloatTy() || RedVarType->isDoubleTy() ||
          RedVarType->isHalfTy() || RedVarType->isBFloatTy() ||
          (RedVarType->isIntegerTy() &&
           (RedVarType->getPrimitiveSizeInBits() == 16 ||
            RedVarType->getPrimitiveSizeInBits() == 32 ||
            RedVarType->getPrimitiveSizeInBits() == 64))) &&
         "Unhandled type");
  assert((Opcode == CodeGenModule::XR_OP_add ||
          Opcode == CodeGenModule::XR_OP_min ||
          Opcode == CodeGenModule::XR_OP_max) &&
         "Unexpected Xteam reduction operator");
  std::pair<llvm::Value *, llvm::Value *> RfunPair =
      getXteamRedFunctionPtrs(CGF, RedVarType, Opcode);
  // The initial value (referred to as the sentinel value) of the local
  // reduction variable depends on the opcode.
  llvm::Value *SentinelVal = CGF.getXteamRedSentinel(RedVarType, Opcode);

  llvm::Value *Args[] = {
      Val,
      OrigVarPtr,
      DTeamVals,
      DTeamsDonePtr,
      RfunPair.first,
      RfunPair.second,
      SentinelVal,
      ThreadStartIndex,
      NumTeams,
      CGF.CGM.getLangOpts().OpenMPTargetMultiDevice
          ? llvm::ConstantInt::get(CGF.CGM.Int32Ty,
                                   0) /* __MEMORY_SCOPE_SYSTEM */
          : llvm::ConstantInt::get(CGF.CGM.Int32Ty,
                                   1) /* __MEMORY_SCOPE_DEVICE */};

  unsigned WarpSize = CGF.getTarget().getGridValue().GV_Warp_Size;
  assert(WarpSize == 32 || WarpSize == 64);

  assert(BlockSize > 0 && BlockSize <= llvm::omp::xteam_red::MaxBlockSize &&
         "XTeam Reduction blocksize outside expected range");
  assert(((BlockSize & (BlockSize - 1)) == 0) &&
         "XTeam Reduction blocksize must be a power of two");

  // Prior analysis ensures that Xteam min/max reduction is not initiated if
  // fast reduction is requested by the user.
  if (IsFast)
    assert(Opcode == CodeGenModule::XR_OP_add &&
           "Fast reduction is not enabled for min and max");

  if (RedVarType->isIntegerTy()) {
    if (RedVarType->getPrimitiveSizeInBits() == 16) {
      if (WarpSize == 32) {
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(
                CGM.getModule(), IsFast ? OMPRTL___kmpc_xteamr_s_32x32_fast_sum
                                        : OMPRTL___kmpc_xteamr_s_32x32),
            Args);
      } else {
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(
                CGM.getModule(), IsFast ? OMPRTL___kmpc_xteamr_s_16x64_fast_sum
                                        : OMPRTL___kmpc_xteamr_s_16x64),
            Args);
      }
    }
    if (RedVarType->getPrimitiveSizeInBits() == 32) {
      if (WarpSize == 32) {
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(
                CGM.getModule(), IsFast ? OMPRTL___kmpc_xteamr_i_32x32_fast_sum
                                        : OMPRTL___kmpc_xteamr_i_32x32),
            Args);
      } else {
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(
                CGM.getModule(), IsFast ? OMPRTL___kmpc_xteamr_i_16x64_fast_sum
                                        : OMPRTL___kmpc_xteamr_i_16x64),
            Args);
      }
    }
    if (RedVarType->getPrimitiveSizeInBits() == 64) {
      if (WarpSize == 32) {
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(
                CGM.getModule(), IsFast ? OMPRTL___kmpc_xteamr_l_32x32_fast_sum
                                        : OMPRTL___kmpc_xteamr_l_32x32),
            Args);
      } else {
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(
                CGM.getModule(), IsFast ? OMPRTL___kmpc_xteamr_l_16x64_fast_sum
                                        : OMPRTL___kmpc_xteamr_l_16x64),
            Args);
      }
    }
  }
  if (RedVarType->isFloatTy()) {
    if (WarpSize == 32) {
      return CGF.EmitRuntimeCall(
          OMPBuilder.getOrCreateRuntimeFunction(
              CGM.getModule(), IsFast ? OMPRTL___kmpc_xteamr_f_32x32_fast_sum
                                      : OMPRTL___kmpc_xteamr_f_32x32),
          Args);
    } else {
      return CGF.EmitRuntimeCall(
          OMPBuilder.getOrCreateRuntimeFunction(
              CGM.getModule(), IsFast ? OMPRTL___kmpc_xteamr_f_16x64_fast_sum
                                      : OMPRTL___kmpc_xteamr_f_16x64),
          Args);
    }
  }
  if (RedVarType->isDoubleTy()) {
    if (WarpSize == 32) {
      return CGF.EmitRuntimeCall(
          OMPBuilder.getOrCreateRuntimeFunction(
              CGM.getModule(), IsFast ? OMPRTL___kmpc_xteamr_d_32x32_fast_sum
                                      : OMPRTL___kmpc_xteamr_d_32x32),
          Args);
    } else {
      return CGF.EmitRuntimeCall(
          OMPBuilder.getOrCreateRuntimeFunction(
              CGM.getModule(), IsFast ? OMPRTL___kmpc_xteamr_d_16x64_fast_sum
                                      : OMPRTL___kmpc_xteamr_d_16x64),
          Args);
    }
  }
  if (RedVarType->isHalfTy()) {
    if (WarpSize == 32) {
      return CGF.EmitRuntimeCall(
          OMPBuilder.getOrCreateRuntimeFunction(
              CGM.getModule(), IsFast ? OMPRTL___kmpc_xteamr_h_32x32_fast_sum
                                      : OMPRTL___kmpc_xteamr_h_32x32),
          Args);
    } else {
      return CGF.EmitRuntimeCall(
          OMPBuilder.getOrCreateRuntimeFunction(
              CGM.getModule(), IsFast ? OMPRTL___kmpc_xteamr_h_16x64_fast_sum
                                      : OMPRTL___kmpc_xteamr_h_16x64),
          Args);
    }
  }
  if (RedVarType->isBFloatTy()) {
    if (WarpSize == 32) {
      return CGF.EmitRuntimeCall(
          OMPBuilder.getOrCreateRuntimeFunction(
              CGM.getModule(), IsFast ? OMPRTL___kmpc_xteamr_bf_32x32_fast_sum
                                      : OMPRTL___kmpc_xteamr_bf_32x32),
          Args);
    } else {
      return CGF.EmitRuntimeCall(
          OMPBuilder.getOrCreateRuntimeFunction(
              CGM.getModule(), IsFast ? OMPRTL___kmpc_xteamr_bf_16x64_fast_sum
                                      : OMPRTL___kmpc_xteamr_bf_16x64),
          Args);
    }
  }
  llvm_unreachable("No support for other types currently.");
}

llvm::Value *CGOpenMPRuntimeGPU::getXteamScanSum(
    CodeGenFunction &CGF, llvm::Value *Val, llvm::Value *SumPtr,
    llvm::Value *DTeamVals, llvm::Value *DTeamsDonePtr,
    llvm::Value *DScanStorage, llvm::Value *ThreadStartIndex,
    llvm::Value *NumTeams, int BlockSize, bool IsFast) {
  // TODO handle more types
  llvm::Type *SumType = Val->getType();
  assert(
      (SumType->isFloatTy() || SumType->isDoubleTy() ||
       (SumType->isIntegerTy() && (SumType->getPrimitiveSizeInBits() == 32 ||
                                   SumType->getPrimitiveSizeInBits() == 64))) &&
      "Unhandled type");

  llvm::Type *Int32Ty = llvm::Type::getInt32Ty(CGM.getLLVMContext());
  llvm::Type *Int64Ty = llvm::Type::getInt64Ty(CGM.getLLVMContext());

  std::pair<llvm::Value *, llvm::Value *> RfunPair =
      getXteamRedFunctionPtrs(CGF, SumType, CodeGenModule::XR_OP_add);
  llvm::Value *ZeroVal = (SumType->isFloatTy() || SumType->isDoubleTy())
                             ? llvm::ConstantFP::getZero(SumType)
                         : SumType->getPrimitiveSizeInBits() == 32
                             ? llvm::ConstantInt::get(Int32Ty, 0)
                             : llvm::ConstantInt::get(Int64Ty, 0);

  // TODO: The argument 'SumPtr' is useless for Xteam Scan. Plan to get rid of
  // it in the future from both here and the DeviceRTL implementation.
  llvm::Value *Args[] = {Val,
                         DScanStorage,
                         SumPtr,
                         DTeamVals,
                         DTeamsDonePtr,
                         RfunPair.first,
                         RfunPair.second,
                         ZeroVal,
                         ThreadStartIndex,
                         NumTeams};

  unsigned WarpSize = CGF.getTarget().getGridValue().GV_Warp_Size;
  assert(WarpSize == 32 || WarpSize == 64);

  assert(BlockSize > 0 && BlockSize <= llvm::omp::xteam_red::MaxBlockSize &&
         "XTeam Reduction blocksize outside expected range");
  assert(((BlockSize & (BlockSize - 1)) == 0) &&
         "XTeam Reduction blocksize must be a power of two");

  if (SumType->isIntegerTy()) {
    if (SumType->getPrimitiveSizeInBits() == 64) {
      if (WarpSize == 64) {
        if (BlockSize == 1024)
          return CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_xteams_l_16x64),
              Args);
        else if (BlockSize == 512)
          return CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_xteams_l_8x64),
              Args);
        else if (BlockSize == 256)
          return CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_xteams_l_4x64),
              Args);
        else
          llvm_unreachable("Block size unsupported.");
      } else if (WarpSize == 32) {
        if (BlockSize == 1024)
          return CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_xteams_l_32x32),
              Args);
        else if (BlockSize == 512)
          return CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_xteams_l_16x32),
              Args);
        else if (BlockSize == 256)
          return CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_xteams_l_8x32),
              Args);
        else
          llvm_unreachable("Block size unsupported.");
      } else
        llvm_unreachable("Warp size should be 32 or 64.");
    } else if (SumType->getPrimitiveSizeInBits() == 32) {
      if (WarpSize == 64) {
        if (BlockSize == 1024)
          return CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_xteams_i_16x64),
              Args);
        else if (BlockSize == 512)
          return CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_xteams_i_8x64),
              Args);
        else if (BlockSize == 256)
          return CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_xteams_i_4x64),
              Args);
        else
          llvm_unreachable("Block size unsupported.");
      } else if (WarpSize == 32) {
        if (BlockSize == 1024)
          return CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_xteams_i_32x32),
              Args);
        else if (BlockSize == 512)
          return CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_xteams_i_16x32),
              Args);
        else if (BlockSize == 256)
          return CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_xteams_i_8x32),
              Args);
        else
          llvm_unreachable("Block size unsupported.");
      } else
        llvm_unreachable("Warp size should be 32 or 64.");
    }
  }
  if (SumType->isDoubleTy()) {
    if (WarpSize == 64) {
      if (BlockSize == 1024)
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(),
                                                  OMPRTL___kmpc_xteams_d_16x64),
            Args);
      else if (BlockSize == 512)
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(),
                                                  OMPRTL___kmpc_xteams_d_8x64),
            Args);
      else if (BlockSize == 256)
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(),
                                                  OMPRTL___kmpc_xteams_d_4x64),
            Args);
      else
        llvm_unreachable("Block size unsupported.");
    } else if (WarpSize == 32) {
      if (BlockSize == 1024)
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(),
                                                  OMPRTL___kmpc_xteams_d_32x32),
            Args);
      else if (BlockSize == 512)
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(),
                                                  OMPRTL___kmpc_xteams_d_16x32),
            Args);
      else if (BlockSize == 256)
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(),
                                                  OMPRTL___kmpc_xteams_d_8x32),
            Args);
      else
        llvm_unreachable("Block size unsupported.");
    } else
      llvm_unreachable("Warp size should be 32 or 64.");
  }
  if (SumType->isFloatTy()) {
    // FIXME: The Xteam Scan Implementation exhibits unpredictable behavior for
    // 'float' datatype when number of elements to be scanned goes beyond 1
    // million. This issue requires further debugging.
    if (WarpSize == 64) {
      if (BlockSize == 1024)
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(),
                                                  OMPRTL___kmpc_xteams_f_16x64),
            Args);
      else if (BlockSize == 512)
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(),
                                                  OMPRTL___kmpc_xteams_f_8x64),
            Args);
      else if (BlockSize == 256)
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(),
                                                  OMPRTL___kmpc_xteams_f_4x64),
            Args);
      else
        llvm_unreachable("BBlock size unsupported.");
    } else if (WarpSize == 32) {
      if (BlockSize == 1024)
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(),
                                                  OMPRTL___kmpc_xteams_f_32x32),
            Args);
      else if (BlockSize == 512)
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(),
                                                  OMPRTL___kmpc_xteams_f_16x32),
            Args);
      else if (BlockSize == 256)
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(),
                                                  OMPRTL___kmpc_xteams_f_8x32),
            Args);
      else
        llvm_unreachable("Block size unsupported.");
    } else
      llvm_unreachable("Warp size should be 32 or 64.");
  }
  llvm_unreachable("No support for other types currently.");
}

llvm::Value *CGOpenMPRuntimeGPU::getXteamScanPhaseTwo(
    CodeGenFunction &CGF, llvm::Value *Val, llvm::Value *SegmentSize,
    llvm::Value *DTeamVals, llvm::Value *DScanStorage,
    llvm::Value *DSegmentVals, llvm::Value *ThreadStartIndex, int BlockSize,
    bool IsInclusiveScan) {
  // TODO handle more types
  llvm::Type *SumType = Val->getType();
  assert(
      (SumType->isFloatTy() || SumType->isDoubleTy() ||
       (SumType->isIntegerTy() && (SumType->getPrimitiveSizeInBits() == 32 ||
                                   SumType->getPrimitiveSizeInBits() == 64))) &&
      "Unhandled type");

  llvm::Type *Int32Ty = llvm::Type::getInt32Ty(CGM.getLLVMContext());
  llvm::Type *Int64Ty = llvm::Type::getInt64Ty(CGM.getLLVMContext());

  std::pair<llvm::Value *, llvm::Value *> RfunPair =
      getXteamRedFunctionPtrs(CGF, SumType, CodeGenModule::XR_OP_add);
  llvm::Value *ZeroVal = (SumType->isFloatTy() || SumType->isDoubleTy())
                             ? llvm::ConstantFP::getZero(SumType)
                         : SumType->getPrimitiveSizeInBits() == 32
                             ? llvm::ConstantInt::get(Int32Ty, 0)
                             : llvm::ConstantInt::get(Int64Ty, 0);

  llvm::Value *IsInclusiveScanVal =
      llvm::ConstantInt::get(Int32Ty, IsInclusiveScan);
  llvm::Value *Args[] = {DScanStorage,     SegmentSize,       DTeamVals,
                         DSegmentVals,     RfunPair.first,    ZeroVal,
                         ThreadStartIndex, IsInclusiveScanVal};

  unsigned WarpSize = CGF.getTarget().getGridValue().GV_Warp_Size;
  assert(WarpSize == 32 || WarpSize == 64);

  assert(BlockSize > 0 && BlockSize <= llvm::omp::xteam_red::MaxBlockSize &&
         "XTeam Reduction blocksize outside expected range");
  assert(((BlockSize & (BlockSize - 1)) == 0) &&
         "XTeam Reduction blocksize must be a power of two");

  if (SumType->isIntegerTy()) {
    if (SumType->getPrimitiveSizeInBits() == 64) {
      if (WarpSize == 64) {
        if (BlockSize == 1024)
          return CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_xteams_phase2_l_16x64),
              Args);
        else if (BlockSize == 512)
          return CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_xteams_phase2_l_8x64),
              Args);
        else if (BlockSize == 256)
          return CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_xteams_phase2_l_4x64),
              Args);
        else
          llvm_unreachable("Block size unsupported.");
      } else if (WarpSize == 32) {
        if (BlockSize == 1024)
          return CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_xteams_phase2_l_32x32),
              Args);
        else if (BlockSize == 512)
          return CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_xteams_phase2_l_16x32),
              Args);
        else if (BlockSize == 256)
          return CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_xteams_phase2_l_8x32),
              Args);
        else
          llvm_unreachable("Block size unsupported.");
      } else
        llvm_unreachable("Warp size should be 32 or 64.");
    } else if (SumType->getPrimitiveSizeInBits() == 32) {
      if (WarpSize == 64) {
        if (BlockSize == 1024)
          return CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_xteams_phase2_i_16x64),
              Args);
        else if (BlockSize == 512)
          return CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_xteams_phase2_i_8x64),
              Args);
        else if (BlockSize == 256)
          return CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_xteams_phase2_i_4x64),
              Args);
        else
          llvm_unreachable("Block size unsupported.");
      } else if (WarpSize == 32) {
        if (BlockSize == 1024)
          return CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_xteams_phase2_i_32x32),
              Args);
        else if (BlockSize == 512)
          return CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_xteams_phase2_i_16x32),
              Args);
        else if (BlockSize == 256)
          return CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_xteams_phase2_i_8x32),
              Args);
        else
          llvm_unreachable("Block size unsupported.");
      } else
        llvm_unreachable("Warp size should be 32 or 64.");
    }
  }
  if (SumType->isDoubleTy()) {
    if (WarpSize == 64) {
      if (BlockSize == 1024)
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(
                CGM.getModule(), OMPRTL___kmpc_xteams_phase2_d_16x64),
            Args);
      else if (BlockSize == 512)
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(
                CGM.getModule(), OMPRTL___kmpc_xteams_phase2_d_8x64),
            Args);
      else if (BlockSize == 256)
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(
                CGM.getModule(), OMPRTL___kmpc_xteams_phase2_d_4x64),
            Args);
      else
        llvm_unreachable("Block size unsupported.");
    } else if (WarpSize == 32) {
      if (BlockSize == 1024)
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(
                CGM.getModule(), OMPRTL___kmpc_xteams_phase2_d_32x32),
            Args);
      else if (BlockSize == 512)
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(
                CGM.getModule(), OMPRTL___kmpc_xteams_phase2_d_16x32),
            Args);
      else if (BlockSize == 256)
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(
                CGM.getModule(), OMPRTL___kmpc_xteams_phase2_d_8x32),
            Args);
      else
        llvm_unreachable("Block size unsupported.");
    } else
      llvm_unreachable("Warp size should be 32 or 64.");
  }
  if (SumType->isFloatTy()) {
    if (WarpSize == 64) {
      if (BlockSize == 1024)
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(
                CGM.getModule(), OMPRTL___kmpc_xteams_phase2_f_16x64),
            Args);
      else if (BlockSize == 512)
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(
                CGM.getModule(), OMPRTL___kmpc_xteams_phase2_f_8x64),
            Args);
      else if (BlockSize == 256)
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(
                CGM.getModule(), OMPRTL___kmpc_xteams_phase2_f_4x64),
            Args);
      else
        llvm_unreachable("BBlock size unsupported.");
    } else if (WarpSize == 32) {
      if (BlockSize == 1024)
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(
                CGM.getModule(), OMPRTL___kmpc_xteams_phase2_f_32x32),
            Args);
      else if (BlockSize == 512)
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(
                CGM.getModule(), OMPRTL___kmpc_xteams_phase2_f_16x32),
            Args);
      else if (BlockSize == 256)
        return CGF.EmitRuntimeCall(
            OMPBuilder.getOrCreateRuntimeFunction(
                CGM.getModule(), OMPRTL___kmpc_xteams_phase2_f_8x32),
            Args);
      else
        llvm_unreachable("Block size unsupported.");
    } else
      llvm_unreachable("Warp size should be 32 or 64.");
  }
  llvm_unreachable("No support for other types currently.");
}

bool CGOpenMPRuntimeGPU::needsHintsForFastFPAtomics() {
  return getOffloadArch(CGM) == OffloadArch::GFX90a;
}

bool CGOpenMPRuntimeGPU::supportFastFPAtomics() {
  OffloadArch Arch = getOffloadArch(CGM);
  switch (Arch) {
  case OffloadArch::GFX90a:
  case OffloadArch::GFX942:
    return true;
  default:
    break;
  }
  return false;
}

std::pair<bool, RValue>
CGOpenMPRuntimeGPU::emitFastFPAtomicCall(CodeGenFunction &CGF, LValue X,
                                         RValue Update, BinaryOperatorKind BO,
                                         bool IsXBinopExpr) {
  CGBuilderTy &Bld = CGF.Builder;
  llvm::AtomicRMWInst::BinOp Kind = llvm::AtomicRMWInst::FAdd;
  switch (BO) {
  case BO_Sub:
    Kind = llvm::AtomicRMWInst::FSub;
    break;
  case BO_Add:
    Kind = llvm::AtomicRMWInst::FAdd;
    break;
  case BO_LT:
    Kind = IsXBinopExpr ? llvm::AtomicRMWInst::FMax : llvm::AtomicRMWInst::FMin;
    break;
  case BO_GT:
    Kind = IsXBinopExpr ? llvm::AtomicRMWInst::FMin : llvm::AtomicRMWInst::FMax;
    break;
  default:
    // remaining operations are not supported yet
    return std::make_pair(false, RValue::get(nullptr));
  }

  llvm::Value *UpdateVal = Update.getScalarVal();

  // The scope of the atomic, currently set to 'agent'. By default, if this
  // scope is not specified the scope will be 'system' scope.
  llvm::SyncScope::ID SSID =
      CGM.getLLVMContext().getOrInsertSyncScopeID("agent");
  llvm::AtomicRMWInst *CallInst = Bld.CreateAtomicRMW(
      Kind, X.getAddress(), UpdateVal, llvm::AtomicOrdering::Monotonic, SSID);

  // The following settings are used to get the atomicrmw instruction to
  // be closer in spirit to the previous use of the intrinsic.
  // Setting of amdgpu.no.fine.grained.memory property
  llvm::MDTuple *EmptyMD = llvm::MDNode::get(CGM.getLLVMContext(), {});
  CallInst->setMetadata("amdgpu.no.fine.grained.memory", EmptyMD);

  // Setting of amdgpu.ignore.denormal.mode
  if (Kind == llvm::AtomicRMWInst::FAdd && UpdateVal->getType()->isFloatTy())
    CallInst->setMetadata("amdgpu.ignore.denormal.mode", EmptyMD);

  // Note: breaks fp_atomics test so volatile cannot be used
  // CallInst->setVolatile(true);

  return std::make_pair(true, RValue::get(CallInst));
}

void CGOpenMPRuntimeGPU::emitFlush(CodeGenFunction &CGF, ArrayRef<const Expr *>,
                                   SourceLocation Loc,
                                   llvm::AtomicOrdering AO) {
  if (CGF.CGM.getLangOpts().OpenMPIRBuilder) {
    OMPBuilder.createFlush(CGF.Builder);
  } else {
    if (!CGF.HaveInsertPoint())
      return;
    // Build call void __kmpc_flush(ident_t *loc) and variants
    //__kmpc_flush_acquire, __kmpc_flush_release, __kmpc_flush_acqrel
    if (AO == llvm::AtomicOrdering::NotAtomic ||
        AO == llvm::AtomicOrdering::SequentiallyConsistent)
      CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
                              CGM.getModule(), OMPRTL___kmpc_flush),
                          emitUpdateLocation(CGF, Loc));
    else
      switch (AO) {
      case llvm::AtomicOrdering::Acquire:
        CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
                                CGM.getModule(), OMPRTL___kmpc_flush_acquire),
                            emitUpdateLocation(CGF, Loc));
        return;
      case llvm::AtomicOrdering::Release:
        CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
                                CGM.getModule(), OMPRTL___kmpc_flush_release),
                            emitUpdateLocation(CGF, Loc));
        return;
      case llvm::AtomicOrdering::AcquireRelease:
        CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
                                CGM.getModule(), OMPRTL___kmpc_flush_acqrel),
                            emitUpdateLocation(CGF, Loc));
        return;
      default:
        llvm_unreachable("Unexpected atomic ordering for flush directive.");
      }
  }
}

std::pair<bool, RValue>
CGOpenMPRuntimeGPU::emitAtomicCASLoop(CodeGenFunction &CGF, LValue X,
                                      RValue Update, BinaryOperatorKind BO) {
  ASTContext &Context = CGF.getContext();
  SmallVector<llvm::Value *> CASLoopArgs;
  CASLoopArgs.reserve(2);
  CASLoopArgs.push_back(X.getPointer(CGF));
  CASLoopArgs.push_back(Update.getScalarVal());
  llvm::Value *CallInst = nullptr;
  switch (BO) {
  case BO_LT: { // unavailable for both float, double, and integer types (32 and
                // 64 bits)
    if (Update.getScalarVal()->getType()->isIntegerTy() &&
        !(Context.getTypeSize(X.getType()) == 32 ||
          Context.getTypeSize(X.getType()) == 64))
      llvm_unreachable("Atomic Min types available for CAS loop conversion is "
                       "double, float, int (32 and 64 bits)");

    if (Update.getScalarVal()->getType()->isDoubleTy())
      CallInst = CGF.EmitRuntimeCall(
          OMPBuilder.getOrCreateRuntimeFunction(
              CGM.getModule(), OMPRTL___kmpc_atomicCASLoopMin_double),
          CASLoopArgs);
    else if (Update.getScalarVal()->getType()->isFloatTy())
      CallInst = CGF.EmitRuntimeCall(
          OMPBuilder.getOrCreateRuntimeFunction(
              CGM.getModule(), OMPRTL___kmpc_atomicCASLoopMin_float),
          CASLoopArgs);

    else if (Update.getScalarVal()->getType()->isIntegerTy()) {
      if (Context.getTypeSize(X.getType()) == 32) {
        if (X.getType()->hasSignedIntegerRepresentation()) {
          CallInst = CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_atomicCASLoopMin_int32_t),
              CASLoopArgs);
        } else {
          const llvm::StringRef FunNameStr = "__kmpc_atomicCASLoopMin_uint32_t";
          CallInst = CGF.EmitRuntimeCall(
              OMPBuilder.unsignedGetOrCreateAtomicCASRuntimeFunction(
                  CGM.getModule(), FunNameStr,
                  /*RetType=*/CGF.Builder.getVoidTy(),
                  X.getPointer(CGF)->getType(),
                  Update.getScalarVal()->getType()),
              CASLoopArgs);
        }
      } else if (Context.getTypeSize(X.getType()) == 64) {
        if (X.getType()->hasSignedIntegerRepresentation()) {
          CallInst = CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_atomicCASLoopMin_int64_t),
              CASLoopArgs);
        } else {
          const llvm::StringRef FunNameStr = "__kmpc_atomicCASLoopMin_uint64_t";
          CallInst = CGF.EmitRuntimeCall(
              OMPBuilder.unsignedGetOrCreateAtomicCASRuntimeFunction(
                  CGM.getModule(), FunNameStr,
                  /*RetType=*/CGF.Builder.getVoidTy(),
                  X.getPointer(CGF)->getType(),
                  Update.getScalarVal()->getType()),
              CASLoopArgs);
        }
      }
    }
    // other types (e.g., int8_t) are handled by backend directly
    return std::make_pair(true, RValue::get(CallInst));
  }
  case BO_GT: { // unavailable for both float, double, and integer types (32 and
                // 664 bits)
    if (Update.getScalarVal()->getType()->isIntegerTy() &&
        !(Context.getTypeSize(X.getType()) == 32 ||
          Context.getTypeSize(X.getType()) == 64))
      llvm_unreachable("Atomic Max types available for CAS loop conversion is "
                       "double, float, int (32 and 64 bits)");

    if (Update.getScalarVal()->getType()->isDoubleTy())
      CallInst = CGF.EmitRuntimeCall(
          OMPBuilder.getOrCreateRuntimeFunction(
              CGM.getModule(), OMPRTL___kmpc_atomicCASLoopMax_double),
          CASLoopArgs);
    else if (Update.getScalarVal()->getType()->isFloatTy())
      CallInst = CGF.EmitRuntimeCall(
          OMPBuilder.getOrCreateRuntimeFunction(
              CGM.getModule(), OMPRTL___kmpc_atomicCASLoopMax_float),
          CASLoopArgs);

    else if (Update.getScalarVal()->getType()->isIntegerTy()) {
      if (Context.getTypeSize(X.getType()) == 32) {
        if (X.getType()->hasSignedIntegerRepresentation()) {
          CallInst = CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_atomicCASLoopMax_int32_t),
              CASLoopArgs);
        } else {
          const llvm::StringRef FunNameStr = "__kmpc_atomicCASLoopMax_uint32_t";
          CallInst = CGF.EmitRuntimeCall(
              OMPBuilder.unsignedGetOrCreateAtomicCASRuntimeFunction(
                  CGM.getModule(), FunNameStr,
                  /*RetType=*/CGF.Builder.getVoidTy(),
                  X.getPointer(CGF)->getType(),
                  Update.getScalarVal()->getType()),
              CASLoopArgs);
        }
      } else if (Context.getTypeSize(X.getType()) == 64) {
        if (X.getType()->hasSignedIntegerRepresentation()) {
          CallInst = CGF.EmitRuntimeCall(
              OMPBuilder.getOrCreateRuntimeFunction(
                  CGM.getModule(), OMPRTL___kmpc_atomicCASLoopMax_int64_t),
              CASLoopArgs);
        } else {
          const llvm::StringRef FunNameStr = "__kmpc_atomicCASLoopMax_uint64_t";
          CallInst = CGF.EmitRuntimeCall(
              OMPBuilder.unsignedGetOrCreateAtomicCASRuntimeFunction(
                  CGM.getModule(), FunNameStr,
                  /*RetType=*/CGF.Builder.getVoidTy(),
                  X.getPointer(CGF)->getType(),
                  Update.getScalarVal()->getType()),
              CASLoopArgs);
        }
      }
    }
    return std::make_pair(true, RValue::get(CallInst));
  }
  case BO_Add:
  case BO_Sub:
  case BO_And:
  case BO_Or:
  case BO_Xor:
    llvm_unreachable("Atomic operation must be generated via clang atomic "
                     "support and not via OpenMP runtime");
    break;
  default:
    llvm_unreachable(
        "Operation is not supported by kmpc_atomicCASLoop functions");
    break;
  }
  return std::make_pair(false, RValue::get(nullptr));
}