OMPIRBuilder.cpp

Go to the documentation of this file.

//===- OpenMPIRBuilder.cpp - Builder for LLVM-IR for OpenMP directives ----===//
//
// 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
//
//===----------------------------------------------------------------------===//
/// \file
///
/// This file implements the OpenMPIRBuilder class, which is used as a
/// convenient way to create LLVM instructions for OpenMP directives.
///
//===----------------------------------------------------------------------===//
 
#include "llvm/Frontend/OpenMP/OMPIRBuilder.h"
#include "llvm/ADT/SmallSet.h"
#include "llvm/ADT/StringExtras.h"
#include "llvm/ADT/StringRef.h"
#include "llvm/Analysis/AssumptionCache.h"
#include "llvm/Analysis/CodeMetrics.h"
#include "llvm/Analysis/LoopInfo.h"
#include "llvm/Analysis/OptimizationRemarkEmitter.h"
#include "llvm/Analysis/ScalarEvolution.h"
#include "llvm/Analysis/TargetLibraryInfo.h"
#include "llvm/Bitcode/BitcodeReader.h"
#include "llvm/Frontend/Offloading/Utility.h"
#include "llvm/Frontend/OpenMP/OMPGridValues.h"
#include "llvm/IR/Attributes.h"
#include "llvm/IR/BasicBlock.h"
#include "llvm/IR/CFG.h"
#include "llvm/IR/CallingConv.h"
#include "llvm/IR/Constant.h"
#include "llvm/IR/Constants.h"
#include "llvm/IR/DebugInfoMetadata.h"
#include "llvm/IR/DerivedTypes.h"
#include "llvm/IR/Function.h"
#include "llvm/IR/GlobalVariable.h"
#include "llvm/IR/IRBuilder.h"
#include "llvm/IR/LLVMContext.h"
#include "llvm/IR/MDBuilder.h"
#include "llvm/IR/Metadata.h"
#include "llvm/IR/PassManager.h"
#include "llvm/IR/PassInstrumentation.h"
#include "llvm/IR/ReplaceConstant.h"
#include "llvm/IR/Value.h"
#include "llvm/MC/TargetRegistry.h"
#include "llvm/Support/CommandLine.h"
#include "llvm/Support/ErrorHandling.h"
#include "llvm/Support/FileSystem.h"
#include "llvm/Target/TargetMachine.h"
#include "llvm/Target/TargetOptions.h"
#include "llvm/Transforms/Utils/BasicBlockUtils.h"
#include "llvm/Transforms/Utils/Cloning.h"
#include "llvm/Transforms/Utils/CodeExtractor.h"
#include "llvm/Transforms/Utils/LoopPeel.h"
#include "llvm/Transforms/Utils/UnrollLoop.h"
 
#include <cstdint>
#include <optional>
#include <stack>
 
#define DEBUG_TYPE "openmp-ir-builder"
 
using namespace llvm;
using namespace omp;
 
static cl::opt<bool>
    OptimisticAttributes("openmp-ir-builder-optimistic-attributes", cl::Hidden,
                         cl::desc("Use optimistic attributes describing "
                                  "'as-if' properties of runtime calls."),
                         cl::init(false));
 
static cl::opt<double> UnrollThresholdFactor(
    "openmp-ir-builder-unroll-threshold-factor", cl::Hidden,
    cl::desc("Factor for the unroll threshold to account for code "
             "simplifications still taking place"),
    cl::init(1.5));
 
#ifndef NDEBUG
/// Return whether IP1 and IP2 are ambiguous, i.e. that inserting instructions
/// at position IP1 may change the meaning of IP2 or vice-versa. This is because
/// an InsertPoint stores the instruction before something is inserted. For
/// instance, if both point to the same instruction, two IRBuilders alternating
/// creating instruction will cause the instructions to be interleaved.
static bool isConflictIP(IRBuilder<>::InsertPoint IP1,
                         IRBuilder<>::InsertPoint IP2) {
  if (!IP1.isSet() || !IP2.isSet())
    return false;
  return IP1.getBlock() == IP2.getBlock() && IP1.getPoint() == IP2.getPoint();
}
 
static bool isValidWorkshareLoopScheduleType(OMPScheduleType SchedType) {
  // Valid ordered/unordered and base algorithm combinations.
  switch (SchedType & ~OMPScheduleType::MonotonicityMask) {
  case OMPScheduleType::UnorderedStaticChunked:
  case OMPScheduleType::UnorderedStatic:
  case OMPScheduleType::UnorderedDynamicChunked:
  case OMPScheduleType::UnorderedGuidedChunked:
  case OMPScheduleType::UnorderedRuntime:
  case OMPScheduleType::UnorderedAuto:
  case OMPScheduleType::UnorderedTrapezoidal:
  case OMPScheduleType::UnorderedGreedy:
  case OMPScheduleType::UnorderedBalanced:
  case OMPScheduleType::UnorderedGuidedIterativeChunked:
  case OMPScheduleType::UnorderedGuidedAnalyticalChunked:
  case OMPScheduleType::UnorderedSteal:
  case OMPScheduleType::UnorderedStaticBalancedChunked:
  case OMPScheduleType::UnorderedGuidedSimd:
  case OMPScheduleType::UnorderedRuntimeSimd:
  case OMPScheduleType::OrderedStaticChunked:
  case OMPScheduleType::OrderedStatic:
  case OMPScheduleType::OrderedDynamicChunked:
  case OMPScheduleType::OrderedGuidedChunked:
  case OMPScheduleType::OrderedRuntime:
  case OMPScheduleType::OrderedAuto:
  case OMPScheduleType::OrderdTrapezoidal:
  case OMPScheduleType::NomergeUnorderedStaticChunked:
  case OMPScheduleType::NomergeUnorderedStatic:
  case OMPScheduleType::NomergeUnorderedDynamicChunked:
  case OMPScheduleType::NomergeUnorderedGuidedChunked:
  case OMPScheduleType::NomergeUnorderedRuntime:
  case OMPScheduleType::NomergeUnorderedAuto:
  case OMPScheduleType::NomergeUnorderedTrapezoidal:
  case OMPScheduleType::NomergeUnorderedGreedy:
  case OMPScheduleType::NomergeUnorderedBalanced:
  case OMPScheduleType::NomergeUnorderedGuidedIterativeChunked:
  case OMPScheduleType::NomergeUnorderedGuidedAnalyticalChunked:
  case OMPScheduleType::NomergeUnorderedSteal:
  case OMPScheduleType::NomergeOrderedStaticChunked:
  case OMPScheduleType::NomergeOrderedStatic:
  case OMPScheduleType::NomergeOrderedDynamicChunked:
  case OMPScheduleType::NomergeOrderedGuidedChunked:
  case OMPScheduleType::NomergeOrderedRuntime:
  case OMPScheduleType::NomergeOrderedAuto:
  case OMPScheduleType::NomergeOrderedTrapezoidal:
    break;
  default:
    return false;
  }
 
  // Must not set both monotonicity modifiers at the same time.
  OMPScheduleType MonotonicityFlags =
      SchedType & OMPScheduleType::MonotonicityMask;
  if (MonotonicityFlags == OMPScheduleType::MonotonicityMask)
    return false;
 
  return true;
}
#endif
 
static const omp::GV &getGridValue(const Triple &T, Function *Kernel) {
  if (T.isAMDGPU()) {
    StringRef Features =
        Kernel->getFnAttribute("target-features").getValueAsString();
    if (Features.count("+wavefrontsize64"))
      return omp::getAMDGPUGridValues<64>();
    return omp::getAMDGPUGridValues<32>();
  }
  if (T.isNVPTX())
    return omp::NVPTXGridValues;
  llvm_unreachable("No grid value available for this architecture!");
}
 
/// Determine which scheduling algorithm to use, determined from schedule clause
/// arguments.
static OMPScheduleType
getOpenMPBaseScheduleType(llvm::omp::ScheduleKind ClauseKind, bool HasChunks,
                          bool HasSimdModifier) {
  // Currently, the default schedule it static.
  switch (ClauseKind) {
  case OMP_SCHEDULE_Default:
  case OMP_SCHEDULE_Static:
    return HasChunks ? OMPScheduleType::BaseStaticChunked
                     : OMPScheduleType::BaseStatic;
  case OMP_SCHEDULE_Dynamic:
    return OMPScheduleType::BaseDynamicChunked;
  case OMP_SCHEDULE_Guided:
    return HasSimdModifier ? OMPScheduleType::BaseGuidedSimd
                           : OMPScheduleType::BaseGuidedChunked;
  case OMP_SCHEDULE_Auto:
    return llvm::omp::OMPScheduleType::BaseAuto;
  case OMP_SCHEDULE_Runtime:
    return HasSimdModifier ? OMPScheduleType::BaseRuntimeSimd
                           : OMPScheduleType::BaseRuntime;
  }
  llvm_unreachable("unhandled schedule clause argument");
}
 
/// Adds ordering modifier flags to schedule type.
static OMPScheduleType
getOpenMPOrderingScheduleType(OMPScheduleType BaseScheduleType,
                              bool HasOrderedClause) {
  assert((BaseScheduleType & OMPScheduleType::ModifierMask) ==
             OMPScheduleType::None &&
         "Must not have ordering nor monotonicity flags already set");
 
  OMPScheduleType OrderingModifier = HasOrderedClause
                                         ? OMPScheduleType::ModifierOrdered
                                         : OMPScheduleType::ModifierUnordered;
  OMPScheduleType OrderingScheduleType = BaseScheduleType | OrderingModifier;
 
  // Unsupported combinations
  if (OrderingScheduleType ==
      (OMPScheduleType::BaseGuidedSimd | OMPScheduleType::ModifierOrdered))
    return OMPScheduleType::OrderedGuidedChunked;
  else if (OrderingScheduleType == (OMPScheduleType::BaseRuntimeSimd |
                                    OMPScheduleType::ModifierOrdered))
    return OMPScheduleType::OrderedRuntime;
 
  return OrderingScheduleType;
}
 
/// Adds monotonicity modifier flags to schedule type.
static OMPScheduleType
getOpenMPMonotonicityScheduleType(OMPScheduleType ScheduleType,
                                  bool HasSimdModifier, bool HasMonotonic,
                                  bool HasNonmonotonic, bool HasOrderedClause) {
  assert((ScheduleType & OMPScheduleType::MonotonicityMask) ==
             OMPScheduleType::None &&
         "Must not have monotonicity flags already set");
  assert((!HasMonotonic || !HasNonmonotonic) &&
         "Monotonic and Nonmonotonic are contradicting each other");
 
  if (HasMonotonic) {
    return ScheduleType | OMPScheduleType::ModifierMonotonic;
  } else if (HasNonmonotonic) {
    return ScheduleType | OMPScheduleType::ModifierNonmonotonic;
  } else {
    // OpenMP 5.1, 2.11.4 Worksharing-Loop Construct, Description.
    // If the static schedule kind is specified or if the ordered clause is
    // specified, and if the nonmonotonic modifier is not specified, the
    // effect is as if the monotonic modifier is specified. Otherwise, unless
    // the monotonic modifier is specified, the effect is as if the
    // nonmonotonic modifier is specified.
    OMPScheduleType BaseScheduleType =
        ScheduleType & ~OMPScheduleType::ModifierMask;
    if ((BaseScheduleType == OMPScheduleType::BaseStatic) ||
        (BaseScheduleType == OMPScheduleType::BaseStaticChunked) ||
        HasOrderedClause) {
      // The monotonic is used by default in openmp runtime library, so no need
      // to set it.
      return ScheduleType;
    } else {
      return ScheduleType | OMPScheduleType::ModifierNonmonotonic;
    }
  }
}
 
/// Determine the schedule type using schedule and ordering clause arguments.
static OMPScheduleType
computeOpenMPScheduleType(ScheduleKind ClauseKind, bool HasChunks,
                          bool HasSimdModifier, bool HasMonotonicModifier,
                          bool HasNonmonotonicModifier, bool HasOrderedClause) {
  OMPScheduleType BaseSchedule =
      getOpenMPBaseScheduleType(ClauseKind, HasChunks, HasSimdModifier);
  OMPScheduleType OrderedSchedule =
      getOpenMPOrderingScheduleType(BaseSchedule, HasOrderedClause);
  OMPScheduleType Result = getOpenMPMonotonicityScheduleType(
      OrderedSchedule, HasSimdModifier, HasMonotonicModifier,
      HasNonmonotonicModifier, HasOrderedClause);
 
  assert(isValidWorkshareLoopScheduleType(Result));
  return Result;
}
 
/// Make \p Source branch to \p Target.
///
/// Handles two situations:
/// * \p Source already has an unconditional branch.
/// * \p Source is a degenerate block (no terminator because the BB is
///             the current head of the IR construction).
static void redirectTo(BasicBlock *Source, BasicBlock *Target, DebugLoc DL) {
  if (Instruction *Term = Source->getTerminator()) {
    auto *Br = cast<BranchInst>(Term);
    assert(!Br->isConditional() &&
           "BB's terminator must be an unconditional branch (or degenerate)");
    BasicBlock *Succ = Br->getSuccessor(0);
    Succ->removePredecessor(Source, /*KeepOneInputPHIs=*/true);
    Br->setSuccessor(0, Target);
    return;
  }
 
  auto *NewBr = BranchInst::Create(Target, Source);
  NewBr->setDebugLoc(DL);
}
 
void llvm::spliceBB(IRBuilderBase::InsertPoint IP, BasicBlock *New,
                    bool CreateBranch) {
  assert(New->getFirstInsertionPt() == New->begin() &&
         "Target BB must not have PHI nodes");
 
  // Move instructions to new block.
  BasicBlock *Old = IP.getBlock();
  New->splice(New->begin(), Old, IP.getPoint(), Old->end());
 
  if (CreateBranch)
    BranchInst::Create(New, Old);
}
 
void llvm::spliceBB(IRBuilder<> &Builder, BasicBlock *New, bool CreateBranch) {
  DebugLoc DebugLoc = Builder.getCurrentDebugLocation();
  BasicBlock *Old = Builder.GetInsertBlock();
 
  spliceBB(Builder.saveIP(), New, CreateBranch);
  if (CreateBranch)
    Builder.SetInsertPoint(Old->getTerminator());
  else
    Builder.SetInsertPoint(Old);
 
  // SetInsertPoint also updates the Builder's debug location, but we want to
  // keep the one the Builder was configured to use.
  Builder.SetCurrentDebugLocation(DebugLoc);
}
 
BasicBlock *llvm::splitBB(IRBuilderBase::InsertPoint IP, bool CreateBranch,
                          llvm::Twine Name) {
  BasicBlock *Old = IP.getBlock();
  BasicBlock *New = BasicBlock::Create(
      Old->getContext(), Name.isTriviallyEmpty() ? Old->getName() : Name,
      Old->getParent(), Old->getNextNode());
  spliceBB(IP, New, CreateBranch);
  New->replaceSuccessorsPhiUsesWith(Old, New);
  return New;
}
 
BasicBlock *llvm::splitBB(IRBuilderBase &Builder, bool CreateBranch,
                          llvm::Twine Name) {
  DebugLoc DebugLoc = Builder.getCurrentDebugLocation();
  BasicBlock *New = splitBB(Builder.saveIP(), CreateBranch, Name);
  if (CreateBranch)
    Builder.SetInsertPoint(Builder.GetInsertBlock()->getTerminator());
  else
    Builder.SetInsertPoint(Builder.GetInsertBlock());
  // SetInsertPoint also updates the Builder's debug location, but we want to
  // keep the one the Builder was configured to use.
  Builder.SetCurrentDebugLocation(DebugLoc);
  return New;
}
 
BasicBlock *llvm::splitBB(IRBuilder<> &Builder, bool CreateBranch,
                          llvm::Twine Name) {
  DebugLoc DebugLoc = Builder.getCurrentDebugLocation();
  BasicBlock *New = splitBB(Builder.saveIP(), CreateBranch, Name);
  if (CreateBranch)
    Builder.SetInsertPoint(Builder.GetInsertBlock()->getTerminator());
  else
    Builder.SetInsertPoint(Builder.GetInsertBlock());
  // SetInsertPoint also updates the Builder's debug location, but we want to
  // keep the one the Builder was configured to use.
  Builder.SetCurrentDebugLocation(DebugLoc);
  return New;
}
 
BasicBlock *llvm::splitBBWithSuffix(IRBuilderBase &Builder, bool CreateBranch,
                                    llvm::Twine Suffix) {
  BasicBlock *Old = Builder.GetInsertBlock();
  return splitBB(Builder, CreateBranch, Old->getName() + Suffix);
}
 
// This function creates a fake integer value and a fake use for the integer
// value. It returns the fake value created. This is useful in modeling the
// extra arguments to the outlined functions.
Value *createFakeIntVal(IRBuilderBase &Builder,
                        OpenMPIRBuilder::InsertPointTy OuterAllocaIP,
                        llvm::SmallVectorImpl<Instruction *> &ToBeDeleted,
                        OpenMPIRBuilder::InsertPointTy InnerAllocaIP,
                        const Twine &Name = "", bool AsPtr = true) {
  Builder.restoreIP(OuterAllocaIP);
  Instruction *FakeVal;
  AllocaInst *FakeValAddr =
      Builder.CreateAlloca(Builder.getInt32Ty(), nullptr, Name + ".addr");
  ToBeDeleted.push_back(FakeValAddr);
 
  if (AsPtr) {
    FakeVal = FakeValAddr;
  } else {
    FakeVal =
        Builder.CreateLoad(Builder.getInt32Ty(), FakeValAddr, Name + ".val");
    ToBeDeleted.push_back(FakeVal);
  }
 
  // Generate a fake use of this value
  Builder.restoreIP(InnerAllocaIP);
  Instruction *UseFakeVal;
  if (AsPtr) {
    UseFakeVal =
        Builder.CreateLoad(Builder.getInt32Ty(), FakeVal, Name + ".use");
  } else {
    UseFakeVal =
        cast<BinaryOperator>(Builder.CreateAdd(FakeVal, Builder.getInt32(10)));
  }
  ToBeDeleted.push_back(UseFakeVal);
  return FakeVal;
}
 
//===----------------------------------------------------------------------===//
// OpenMPIRBuilderConfig
//===----------------------------------------------------------------------===//
 
namespace {
LLVM_ENABLE_BITMASK_ENUMS_IN_NAMESPACE();
/// Values for bit flags for marking which requires clauses have been used.
enum OpenMPOffloadingRequiresDirFlags {
  /// flag undefined.
  OMP_REQ_UNDEFINED = 0x000,
  /// no requires directive present.
  OMP_REQ_NONE = 0x001,
  /// reverse_offload clause.
  OMP_REQ_REVERSE_OFFLOAD = 0x002,
  /// unified_address clause.
  OMP_REQ_UNIFIED_ADDRESS = 0x004,
  /// unified_shared_memory clause.
  OMP_REQ_UNIFIED_SHARED_MEMORY = 0x008,
  /// dynamic_allocators clause.
  OMP_REQ_DYNAMIC_ALLOCATORS = 0x010,
  LLVM_MARK_AS_BITMASK_ENUM(/*LargestValue=*/OMP_REQ_DYNAMIC_ALLOCATORS)
};
 
} // anonymous namespace
 
OpenMPIRBuilderConfig::OpenMPIRBuilderConfig()
    : RequiresFlags(OMP_REQ_UNDEFINED) {}
 
OpenMPIRBuilderConfig::OpenMPIRBuilderConfig(
    bool IsTargetDevice, bool IsGPU, bool OpenMPOffloadMandatory,
    bool HasRequiresReverseOffload, bool HasRequiresUnifiedAddress,
    bool HasRequiresUnifiedSharedMemory, bool HasRequiresDynamicAllocators)
    : IsTargetDevice(IsTargetDevice), IsGPU(IsGPU),
      OpenMPOffloadMandatory(OpenMPOffloadMandatory),
      RequiresFlags(OMP_REQ_UNDEFINED) {
  if (HasRequiresReverseOffload)
    RequiresFlags |= OMP_REQ_REVERSE_OFFLOAD;
  if (HasRequiresUnifiedAddress)
    RequiresFlags |= OMP_REQ_UNIFIED_ADDRESS;
  if (HasRequiresUnifiedSharedMemory)
    RequiresFlags |= OMP_REQ_UNIFIED_SHARED_MEMORY;
  if (HasRequiresDynamicAllocators)
    RequiresFlags |= OMP_REQ_DYNAMIC_ALLOCATORS;
}
 
bool OpenMPIRBuilderConfig::hasRequiresReverseOffload() const {
  return RequiresFlags & OMP_REQ_REVERSE_OFFLOAD;
}
 
bool OpenMPIRBuilderConfig::hasRequiresUnifiedAddress() const {
  return RequiresFlags & OMP_REQ_UNIFIED_ADDRESS;
}
 
bool OpenMPIRBuilderConfig::hasRequiresUnifiedSharedMemory() const {
  return RequiresFlags & OMP_REQ_UNIFIED_SHARED_MEMORY;
}
 
bool OpenMPIRBuilderConfig::hasRequiresDynamicAllocators() const {
  return RequiresFlags & OMP_REQ_DYNAMIC_ALLOCATORS;
}
 
int64_t OpenMPIRBuilderConfig::getRequiresFlags() const {
  return hasRequiresFlags() ? RequiresFlags
                            : static_cast<int64_t>(OMP_REQ_NONE);
}
 
void OpenMPIRBuilderConfig::setHasRequiresReverseOffload(bool Value) {
  if (Value)
    RequiresFlags |= OMP_REQ_REVERSE_OFFLOAD;
  else
    RequiresFlags &= ~OMP_REQ_REVERSE_OFFLOAD;
}
 
void OpenMPIRBuilderConfig::setHasRequiresUnifiedAddress(bool Value) {
  if (Value)
    RequiresFlags |= OMP_REQ_UNIFIED_ADDRESS;
  else
    RequiresFlags &= ~OMP_REQ_UNIFIED_ADDRESS;
}
 
void OpenMPIRBuilderConfig::setHasRequiresUnifiedSharedMemory(bool Value) {
  if (Value)
    RequiresFlags |= OMP_REQ_UNIFIED_SHARED_MEMORY;
  else
    RequiresFlags &= ~OMP_REQ_UNIFIED_SHARED_MEMORY;
}
 
void OpenMPIRBuilderConfig::setHasRequiresDynamicAllocators(bool Value) {
  if (Value)
    RequiresFlags |= OMP_REQ_DYNAMIC_ALLOCATORS;
  else
    RequiresFlags &= ~OMP_REQ_DYNAMIC_ALLOCATORS;
}
 
//===----------------------------------------------------------------------===//
// OpenMPIRBuilder
//===----------------------------------------------------------------------===//
 
void OpenMPIRBuilder::getKernelArgsVector(TargetKernelArgs &KernelArgs,
                                          IRBuilderBase &Builder,
                                          SmallVector<Value *> &ArgsVector) {
  Value *Version = Builder.getInt32(OMP_KERNEL_ARG_VERSION);
  Value *PointerNum = Builder.getInt32(KernelArgs.NumTargetItems);
  auto Int32Ty = Type::getInt32Ty(Builder.getContext());
  constexpr const size_t MaxDim = 3;
  Value *ZeroArray = Constant::getNullValue(ArrayType::get(Int32Ty, MaxDim));
  Value *Flags = Builder.getInt64(KernelArgs.HasNoWait);
 
  assert(!KernelArgs.NumTeams.empty());
 
  Value *NumTeams3D =
      Builder.CreateInsertValue(ZeroArray, KernelArgs.NumTeams[0], {0});
  for (unsigned I = 1; I < std::min(KernelArgs.NumTeams.size(), MaxDim); ++I)
    NumTeams3D =
        Builder.CreateInsertValue(NumTeams3D, KernelArgs.NumTeams[I], {I});
  Value *NumThreads3D =
      Builder.CreateInsertValue(ZeroArray, KernelArgs.NumThreads, {0});
 
  ArgsVector = {Version,
                PointerNum,
                KernelArgs.RTArgs.BasePointersArray,
                KernelArgs.RTArgs.PointersArray,
                KernelArgs.RTArgs.SizesArray,
                KernelArgs.RTArgs.MapTypesArray,
                KernelArgs.RTArgs.MapNamesArray,
                KernelArgs.RTArgs.MappersArray,
                KernelArgs.NumIterations,
                Flags,
                NumTeams3D,
                NumThreads3D,
                KernelArgs.DynCGGroupMem};
}
 
void OpenMPIRBuilder::addAttributes(omp::RuntimeFunction FnID, Function &Fn) {
  LLVMContext &Ctx = Fn.getContext();
 
  // Get the function's current attributes.
  auto Attrs = Fn.getAttributes();
  auto FnAttrs = Attrs.getFnAttrs();
  auto RetAttrs = Attrs.getRetAttrs();
  SmallVector<AttributeSet, 4> ArgAttrs;
  for (size_t ArgNo = 0; ArgNo < Fn.arg_size(); ++ArgNo)
    ArgAttrs.emplace_back(Attrs.getParamAttrs(ArgNo));
 
  // Add AS to FnAS while taking special care with integer extensions.
  auto addAttrSet = [&](AttributeSet &FnAS, const AttributeSet &AS,
                        bool Param = true) -> void {
    bool HasSignExt = AS.hasAttribute(Attribute::SExt);
    bool HasZeroExt = AS.hasAttribute(Attribute::ZExt);
    if (HasSignExt || HasZeroExt) {
      assert(AS.getNumAttributes() == 1 &&
             "Currently not handling extension attr combined with others.");
      if (Param) {
        if (auto AK = TargetLibraryInfo::getExtAttrForI32Param(T, HasSignExt))
          FnAS = FnAS.addAttribute(Ctx, AK);
      } else if (auto AK =
                     TargetLibraryInfo::getExtAttrForI32Return(T, HasSignExt))
        FnAS = FnAS.addAttribute(Ctx, AK);
    } else {
      FnAS = FnAS.addAttributes(Ctx, AS);
    }
  };
 
#define OMP_ATTRS_SET(VarName, AttrSet) AttributeSet VarName = AttrSet;
#include "llvm/Frontend/OpenMP/OMPKinds.def"
 
  // Add attributes to the function declaration.
  switch (FnID) {
#define OMP_RTL_ATTRS(Enum, FnAttrSet, RetAttrSet, ArgAttrSets)                \
  case Enum:                                                                   \
    FnAttrs = FnAttrs.addAttributes(Ctx, FnAttrSet);                           \
    addAttrSet(RetAttrs, RetAttrSet, /*Param*/ false);                         \
    for (size_t ArgNo = 0; ArgNo < ArgAttrSets.size(); ++ArgNo)                \
      addAttrSet(ArgAttrs[ArgNo], ArgAttrSets[ArgNo]);                         \
    Fn.setAttributes(AttributeList::get(Ctx, FnAttrs, RetAttrs, ArgAttrs));    \
    break;
#include "llvm/Frontend/OpenMP/OMPKinds.def"
  default:
    // Attributes are optional.
    break;
  }
}
 
FunctionCallee
OpenMPIRBuilder::getOrCreateRuntimeFunction(Module &M, RuntimeFunction FnID) {
  FunctionType *FnTy = nullptr;
  Function *Fn = nullptr;
 
  // Try to find the declation in the module first.
  switch (FnID) {
#define OMP_RTL(Enum, Str, IsVarArg, ReturnType, ...)                          \
  case Enum:                                                                   \
    FnTy = FunctionType::get(ReturnType, ArrayRef<Type *>{__VA_ARGS__},        \
                             IsVarArg);                                        \
    Fn = M.getFunction(Str);                                                   \
    break;
#include "llvm/Frontend/OpenMP/OMPKinds.def"
  }
 
  if (!Fn) {
    // Create a new declaration if we need one.
    switch (FnID) {
#define OMP_RTL(Enum, Str, ...)                                                \
  case Enum:                                                                   \
    Fn = Function::Create(FnTy, GlobalValue::ExternalLinkage, Str, M);         \
    break;
#include "llvm/Frontend/OpenMP/OMPKinds.def"
    }
 
    // Add information if the runtime function takes a callback function
    if (FnID == OMPRTL___kmpc_fork_call || FnID == OMPRTL___kmpc_fork_teams) {
      if (!Fn->hasMetadata(LLVMContext::MD_callback)) {
        LLVMContext &Ctx = Fn->getContext();
        MDBuilder MDB(Ctx);
        // Annotate the callback behavior of the runtime function:
        //  - The callback callee is argument number 2 (microtask).
        //  - The first two arguments of the callback callee are unknown (-1).
        //  - All variadic arguments to the runtime function are passed to the
        //    callback callee.
        Fn->addMetadata(
            LLVMContext::MD_callback,
            *MDNode::get(Ctx, {MDB.createCallbackEncoding(
                                  2, {-1, -1}, /* VarArgsArePassed */ true)}));
      }
    }
 
    LLVM_DEBUG(dbgs() << "Created OpenMP runtime function " << Fn->getName()
                      << " with type " << *Fn->getFunctionType() << "\n");
    addAttributes(FnID, *Fn);
 
  } else {
    LLVM_DEBUG(dbgs() << "Found OpenMP runtime function " << Fn->getName()
                      << " with type " << *Fn->getFunctionType() << "\n");
  }
 
  assert(Fn && "Failed to create OpenMP runtime function");
 
  return {FnTy, Fn};
}
 
Function *OpenMPIRBuilder::getOrCreateRuntimeFunctionPtr(RuntimeFunction FnID) {
  FunctionCallee RTLFn = getOrCreateRuntimeFunction(M, FnID);
  auto *Fn = dyn_cast<llvm::Function>(RTLFn.getCallee());
  assert(Fn && "Failed to create OpenMP runtime function pointer");
  return Fn;
}
 
void OpenMPIRBuilder::initialize() { initializeTypes(M); }
 
static void raiseUserConstantDataAllocasToEntryBlock(IRBuilderBase &Builder,
                                                     Function *Function) {
  BasicBlock &EntryBlock = Function->getEntryBlock();
  Instruction *MoveLocInst = EntryBlock.getFirstNonPHI();
 
  // Loop over blocks looking for constant allocas, skipping the entry block
  // as any allocas there are already in the desired location.
  for (auto Block = std::next(Function->begin(), 1); Block != Function->end();
       Block++) {
    for (auto Inst = Block->getReverseIterator()->begin();
         Inst != Block->getReverseIterator()->end();) {
      if (auto *AllocaInst = dyn_cast_if_present<llvm::AllocaInst>(Inst)) {
        Inst++;
        if (!isa<ConstantData>(AllocaInst->getArraySize()))
          continue;
        AllocaInst->moveBeforePreserving(MoveLocInst);
      } else {
        Inst++;
      }
    }
  }
}
 
void OpenMPIRBuilder::finalize(Function *Fn) {
  SmallPtrSet<BasicBlock *, 32> ParallelRegionBlockSet;
  SmallVector<BasicBlock *, 32> Blocks;
  SmallVector<OutlineInfo, 16> DeferredOutlines;
  for (OutlineInfo &OI : OutlineInfos) {
    // Skip functions that have not finalized yet; may happen with nested
    // function generation.
    if (Fn && OI.getFunction() != Fn) {
      DeferredOutlines.push_back(OI);
      continue;
    }
 
    ParallelRegionBlockSet.clear();
    Blocks.clear();
    OI.collectBlocks(ParallelRegionBlockSet, Blocks);
 
    Function *OuterFn = OI.getFunction();
    CodeExtractorAnalysisCache CEAC(*OuterFn);
    // If we generate code for the target device, we need to allocate
    // struct for aggregate params in the device default alloca address space.
    // OpenMP runtime requires that the params of the extracted functions are
    // passed as zero address space pointers. This flag ensures that
    // CodeExtractor generates correct code for extracted functions
    // which are used by OpenMP runtime.
    bool ArgsInZeroAddressSpace = Config.isTargetDevice();
    CodeExtractor Extractor(Blocks, /* DominatorTree */ nullptr,
                            /* AggregateArgs */ true,
                            /* BlockFrequencyInfo */ nullptr,
                            /* BranchProbabilityInfo */ nullptr,
                            /* AssumptionCache */ nullptr,
                            /* AllowVarArgs */ true,
                            /* AllowAlloca */ true,
                            /* AllocaBlock*/ OI.OuterAllocaBB,
                            /* Suffix */ ".omp_par", ArgsInZeroAddressSpace);
 
    LLVM_DEBUG(dbgs() << "Before     outlining: " << *OuterFn << "\n");
    LLVM_DEBUG(dbgs() << "Entry " << OI.EntryBB->getName()
                      << " Exit: " << OI.ExitBB->getName() << "\n");
    assert(Extractor.isEligible() &&
           "Expected OpenMP outlining to be possible!");
 
    for (auto *V : OI.ExcludeArgsFromAggregate)
      Extractor.excludeArgFromAggregate(V);
 
    Function *OutlinedFn = Extractor.extractCodeRegion(CEAC);
 
    // Forward target-cpu, target-features attributes to the outlined function.
    auto TargetCpuAttr = OuterFn->getFnAttribute("target-cpu");
    if (TargetCpuAttr.isStringAttribute())
      OutlinedFn->addFnAttr(TargetCpuAttr);
 
    auto TargetFeaturesAttr = OuterFn->getFnAttribute("target-features");
    if (TargetFeaturesAttr.isStringAttribute())
      OutlinedFn->addFnAttr(TargetFeaturesAttr);
 
    LLVM_DEBUG(dbgs() << "After      outlining: " << *OuterFn << "\n");
    LLVM_DEBUG(dbgs() << "   Outlined function: " << *OutlinedFn << "\n");
    assert(OutlinedFn->getReturnType()->isVoidTy() &&
           "OpenMP outlined functions should not return a value!");
 
    // For compability with the clang CG we move the outlined function after the
    // one with the parallel region.
    OutlinedFn->removeFromParent();
    M.getFunctionList().insertAfter(OuterFn->getIterator(), OutlinedFn);
 
    // Remove the artificial entry introduced by the extractor right away, we
    // made our own entry block after all.
    {
      BasicBlock &ArtificialEntry = OutlinedFn->getEntryBlock();
      assert(ArtificialEntry.getUniqueSuccessor() == OI.EntryBB);
      assert(OI.EntryBB->getUniquePredecessor() == &ArtificialEntry);
      // Move instructions from the to-be-deleted ArtificialEntry to the entry
      // basic block of the parallel region. CodeExtractor generates
      // instructions to unwrap the aggregate argument and may sink
      // allocas/bitcasts for values that are solely used in the outlined region
      // and do not escape.
      assert(!ArtificialEntry.empty() &&
             "Expected instructions to add in the outlined region entry");
      for (BasicBlock::reverse_iterator It = ArtificialEntry.rbegin(),
                                        End = ArtificialEntry.rend();
           It != End;) {
        Instruction &I = *It;
        It++;
 
        if (I.isTerminator())
          continue;
 
        I.moveBeforePreserving(*OI.EntryBB, OI.EntryBB->getFirstInsertionPt());
      }
 
      OI.EntryBB->moveBefore(&ArtificialEntry);
      ArtificialEntry.eraseFromParent();
    }
    assert(&OutlinedFn->getEntryBlock() == OI.EntryBB);
    assert(OutlinedFn && OutlinedFn->getNumUses() == 1);
 
    // Run a user callback, e.g. to add attributes.
    if (OI.PostOutlineCB)
      OI.PostOutlineCB(*OutlinedFn);
  }
 
  // Remove work items that have been completed.
  OutlineInfos = std::move(DeferredOutlines);
 
  // The createTarget functions embeds user written code into
  // the target region which may inject allocas which need to
  // be moved to the entry block of our target or risk malformed
  // optimisations by later passes, this is only relevant for
  // the device pass which appears to be a little more delicate
  // when it comes to optimisations (however, we do not block on
  // that here, it's up to the inserter to the list to do so).
  // This notbaly has to occur after the OutlinedInfo candidates
  // have been extracted so we have an end product that will not
  // be implicitly adversely affected by any raises unless
  // intentionally appended to the list.
  // NOTE: This only does so for ConstantData, it could be extended
  // to ConstantExpr's with further effort, however, they should
  // largely be folded when they get here. Extending it to runtime
  // defined/read+writeable allocation sizes would be non-trivial
  // (need to factor in movement of any stores to variables the
  // allocation size depends on, as well as the usual loads,
  // otherwise it'll yield the wrong result after movement) and
  // likely be more suitable as an LLVM optimisation pass.
  for (Function *F : ConstantAllocaRaiseCandidates)
    raiseUserConstantDataAllocasToEntryBlock(Builder, F);
 
  EmitMetadataErrorReportFunctionTy &&ErrorReportFn =
      [](EmitMetadataErrorKind Kind,
         const TargetRegionEntryInfo &EntryInfo) -> void {
    errs() << "Error of kind: " << Kind
           << " when emitting offload entries and metadata during "
              "OMPIRBuilder finalization \n";
  };
 
  if (!OffloadInfoManager.empty())
    createOffloadEntriesAndInfoMetadata(ErrorReportFn);
 
  if (Config.EmitLLVMUsedMetaInfo.value_or(false)) {
    std::vector<WeakTrackingVH> LLVMCompilerUsed = {
        M.getGlobalVariable("__openmp_nvptx_data_transfer_temporary_storage")};
    emitUsed("llvm.compiler.used", LLVMCompilerUsed);
  }
}
 
OpenMPIRBuilder::~OpenMPIRBuilder() {
  assert(OutlineInfos.empty() && "There must be no outstanding outlinings");
}
 
GlobalValue *OpenMPIRBuilder::createGlobalFlag(unsigned Value, StringRef Name) {
  IntegerType *I32Ty = Type::getInt32Ty(M.getContext());
  auto *GV =
      new GlobalVariable(M, I32Ty,
                         /* isConstant = */ true, GlobalValue::WeakODRLinkage,
                         ConstantInt::get(I32Ty, Value), Name);
  GV->setVisibility(GlobalValue::HiddenVisibility);
 
  return GV;
}
 
Constant *OpenMPIRBuilder::getOrCreateIdent(Constant *SrcLocStr,
                                            uint32_t SrcLocStrSize,
                                            IdentFlag LocFlags,
                                            unsigned Reserve2Flags) {
  // Enable "C-mode".
  LocFlags |= OMP_IDENT_FLAG_KMPC;
 
  Constant *&Ident =
      IdentMap[{SrcLocStr, uint64_t(LocFlags) << 31 | Reserve2Flags}];
  if (!Ident) {
    Constant *I32Null = ConstantInt::getNullValue(Int32);
    Constant *IdentData[] = {I32Null,
                             ConstantInt::get(Int32, uint32_t(LocFlags)),
                             ConstantInt::get(Int32, Reserve2Flags),
                             ConstantInt::get(Int32, SrcLocStrSize), SrcLocStr};
    Constant *Initializer =
        ConstantStruct::get(OpenMPIRBuilder::Ident, IdentData);
 
    // Look for existing encoding of the location + flags, not needed but
    // minimizes the difference to the existing solution while we transition.
    for (GlobalVariable &GV : M.globals())
      if (GV.getValueType() == OpenMPIRBuilder::Ident && GV.hasInitializer())
        if (GV.getInitializer() == Initializer)
          Ident = &GV;
 
    if (!Ident) {
      auto *GV = new GlobalVariable(
          M, OpenMPIRBuilder::Ident,
          /* isConstant = */ true, GlobalValue::PrivateLinkage, Initializer, "",
          nullptr, GlobalValue::NotThreadLocal,
          M.getDataLayout().getDefaultGlobalsAddressSpace());
      GV->setUnnamedAddr(GlobalValue::UnnamedAddr::Global);
      GV->setAlignment(Align(8));
      Ident = GV;
    }
  }
 
  return ConstantExpr::getPointerBitCastOrAddrSpaceCast(Ident, IdentPtr);
}
 
Constant *OpenMPIRBuilder::getOrCreateSrcLocStr(StringRef LocStr,
                                                uint32_t &SrcLocStrSize) {
  SrcLocStrSize = LocStr.size();
  Constant *&SrcLocStr = SrcLocStrMap[LocStr];
  if (!SrcLocStr) {
    Constant *Initializer =
        ConstantDataArray::getString(M.getContext(), LocStr);
 
    // Look for existing encoding of the location, not needed but minimizes the
    // difference to the existing solution while we transition.
    for (GlobalVariable &GV : M.globals())
      if (GV.isConstant() && GV.hasInitializer() &&
          GV.getInitializer() == Initializer)
        return SrcLocStr = ConstantExpr::getPointerCast(&GV, Int8Ptr);
 
    SrcLocStr = Builder.CreateGlobalStringPtr(LocStr, /* Name */ "",
                                              /* AddressSpace */ 0, &M);
  }
  return SrcLocStr;
}
 
Constant *OpenMPIRBuilder::getOrCreateSrcLocStr(StringRef FunctionName,
                                                StringRef FileName,
                                                unsigned Line, unsigned Column,
                                                uint32_t &SrcLocStrSize) {
  SmallString<128> Buffer;
  Buffer.push_back(';');
  Buffer.append(FileName);
  Buffer.push_back(';');
  Buffer.append(FunctionName);
  Buffer.push_back(';');
  Buffer.append(std::to_string(Line));
  Buffer.push_back(';');
  Buffer.append(std::to_string(Column));
  Buffer.push_back(';');
  Buffer.push_back(';');
  return getOrCreateSrcLocStr(Buffer.str(), SrcLocStrSize);
}
 
Constant *
OpenMPIRBuilder::getOrCreateDefaultSrcLocStr(uint32_t &SrcLocStrSize) {
  StringRef UnknownLoc = ";unknown;unknown;0;0;;";
  return getOrCreateSrcLocStr(UnknownLoc, SrcLocStrSize);
}
 
Constant *OpenMPIRBuilder::getOrCreateSrcLocStr(DebugLoc DL,
                                                uint32_t &SrcLocStrSize,
                                                Function *F) {
  DILocation *DIL = DL.get();
  if (!DIL)
    return getOrCreateDefaultSrcLocStr(SrcLocStrSize);
  StringRef FileName = M.getName();
  if (DIFile *DIF = DIL->getFile())
    if (std::optional<StringRef> Source = DIF->getSource())
      FileName = *Source;
  StringRef Function = DIL->getScope()->getSubprogram()->getName();
  if (Function.empty() && F)
    Function = F->getName();
  return getOrCreateSrcLocStr(Function, FileName, DIL->getLine(),
                              DIL->getColumn(), SrcLocStrSize);
}
 
Constant *OpenMPIRBuilder::getOrCreateSrcLocStr(const LocationDescription &Loc,
                                                uint32_t &SrcLocStrSize) {
  return getOrCreateSrcLocStr(Loc.DL, SrcLocStrSize,
                              Loc.IP.getBlock()->getParent());
}
 
Value *OpenMPIRBuilder::getOrCreateThreadID(Value *Ident) {
  return Builder.CreateCall(
      getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_global_thread_num), Ident,
      "omp_global_thread_num");
}
 
OpenMPIRBuilder::InsertPointTy
OpenMPIRBuilder::createBarrier(const LocationDescription &Loc, Directive Kind,
                               bool ForceSimpleCall, bool CheckCancelFlag) {
  if (!updateToLocation(Loc))
    return Loc.IP;
 
  // Build call __kmpc_cancel_barrier(loc, thread_id) or
  //            __kmpc_barrier(loc, thread_id);
 
  IdentFlag BarrierLocFlags;
  switch (Kind) {
  case OMPD_for:
    BarrierLocFlags = OMP_IDENT_FLAG_BARRIER_IMPL_FOR;
    break;
  case OMPD_sections:
    BarrierLocFlags = OMP_IDENT_FLAG_BARRIER_IMPL_SECTIONS;
    break;
  case OMPD_single:
    BarrierLocFlags = OMP_IDENT_FLAG_BARRIER_IMPL_SINGLE;
    break;
  case OMPD_barrier:
    BarrierLocFlags = OMP_IDENT_FLAG_BARRIER_EXPL;
    break;
  default:
    BarrierLocFlags = OMP_IDENT_FLAG_BARRIER_IMPL;
    break;
  }
 
  uint32_t SrcLocStrSize;
  Constant *SrcLocStr = getOrCreateSrcLocStr(Loc, SrcLocStrSize);
  Value *Args[] = {
      getOrCreateIdent(SrcLocStr, SrcLocStrSize, BarrierLocFlags),
      getOrCreateThreadID(getOrCreateIdent(SrcLocStr, SrcLocStrSize))};
 
  // If we are in a cancellable parallel region, barriers are cancellation
  // points.
  // TODO: Check why we would force simple calls or to ignore the cancel flag.
  bool UseCancelBarrier =
      !ForceSimpleCall && isLastFinalizationInfoCancellable(OMPD_parallel);
 
  Value *Result =
      Builder.CreateCall(getOrCreateRuntimeFunctionPtr(
                             UseCancelBarrier ? OMPRTL___kmpc_cancel_barrier
                                              : OMPRTL___kmpc_barrier),
                         Args);
 
  if (UseCancelBarrier && CheckCancelFlag)
    emitCancelationCheckImpl(Result, OMPD_parallel);
 
  return Builder.saveIP();
}
 
OpenMPIRBuilder::InsertPointTy
OpenMPIRBuilder::createCancel(const LocationDescription &Loc,
                              Value *IfCondition,
                              omp::Directive CanceledDirective) {
  if (!updateToLocation(Loc))
    return Loc.IP;
 
  // LLVM utilities like blocks with terminators.
  auto *UI = Builder.CreateUnreachable();
 
  Instruction *ThenTI = UI, *ElseTI = nullptr;
  if (IfCondition)
    SplitBlockAndInsertIfThenElse(IfCondition, UI, &ThenTI, &ElseTI);
  Builder.SetInsertPoint(ThenTI);
 
  Value *CancelKind = nullptr;
  switch (CanceledDirective) {
#define OMP_CANCEL_KIND(Enum, Str, DirectiveEnum, Value)                       \
  case DirectiveEnum:                                                          \
    CancelKind = Builder.getInt32(Value);                                      \
    break;
#include "llvm/Frontend/OpenMP/OMPKinds.def"
  default:
    llvm_unreachable("Unknown cancel kind!");
  }
 
  uint32_t SrcLocStrSize;
  Constant *SrcLocStr = getOrCreateSrcLocStr(Loc, SrcLocStrSize);
  Value *Ident = getOrCreateIdent(SrcLocStr, SrcLocStrSize);
  Value *Args[] = {Ident, getOrCreateThreadID(Ident), CancelKind};
  Value *Result = Builder.CreateCall(
      getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_cancel), Args);
  auto ExitCB = [this, CanceledDirective, Loc](InsertPointTy IP) {
    if (CanceledDirective == OMPD_parallel) {
      IRBuilder<>::InsertPointGuard IPG(Builder);
      Builder.restoreIP(IP);
      createBarrier(LocationDescription(Builder.saveIP(), Loc.DL),
                    omp::Directive::OMPD_unknown, /* ForceSimpleCall */ false,
                    /* CheckCancelFlag */ false);
    }
  };
 
  // The actual cancel logic is shared with others, e.g., cancel_barriers.
  emitCancelationCheckImpl(Result, CanceledDirective, ExitCB);
 
  // Update the insertion point and remove the terminator we introduced.
  Builder.SetInsertPoint(UI->getParent());
  UI->eraseFromParent();
 
  return Builder.saveIP();
}
 
OpenMPIRBuilder::InsertPointTy OpenMPIRBuilder::emitTargetKernel(
    const LocationDescription &Loc, InsertPointTy AllocaIP, Value *&Return,
    Value *Ident, Value *DeviceID, Value *NumTeams, Value *NumThreads,
    Value *HostPtr, ArrayRef<Value *> KernelArgs) {
  if (!updateToLocation(Loc))
    return Loc.IP;
 
  Builder.restoreIP(AllocaIP);
  auto *KernelArgsPtr =
      Builder.CreateAlloca(OpenMPIRBuilder::KernelArgs, nullptr, "kernel_args");
  Builder.restoreIP(Loc.IP);
 
  for (unsigned I = 0, Size = KernelArgs.size(); I != Size; ++I) {
    llvm::Value *Arg =
        Builder.CreateStructGEP(OpenMPIRBuilder::KernelArgs, KernelArgsPtr, I);
    Builder.CreateAlignedStore(
        KernelArgs[I], Arg,
        M.getDataLayout().getPrefTypeAlign(KernelArgs[I]->getType()));
  }
 
  SmallVector<Value *> OffloadingArgs{Ident,      DeviceID, NumTeams,
                                      NumThreads, HostPtr,  KernelArgsPtr};
 
  Return = Builder.CreateCall(
      getOrCreateRuntimeFunction(M, OMPRTL___tgt_target_kernel),
      OffloadingArgs);
 
  return Builder.saveIP();
}
 
OpenMPIRBuilder::InsertPointTy OpenMPIRBuilder::emitKernelLaunch(
    const LocationDescription &Loc, Function *OutlinedFn, Value *OutlinedFnID,
    EmitFallbackCallbackTy emitTargetCallFallbackCB, TargetKernelArgs &Args,
    Value *DeviceID, Value *RTLoc, InsertPointTy AllocaIP) {
 
  if (!updateToLocation(Loc))
    return Loc.IP;
 
  Builder.restoreIP(Loc.IP);
  // On top of the arrays that were filled up, the target offloading call
  // takes as arguments the device id as well as the host pointer. The host
  // pointer is used by the runtime library to identify the current target
  // region, so it only has to be unique and not necessarily point to
  // anything. It could be the pointer to the outlined function that
  // implements the target region, but we aren't using that so that the
  // compiler doesn't need to keep that, and could therefore inline the host
  // function if proven worthwhile during optimization.
 
  // From this point on, we need to have an ID of the target region defined.
  assert(OutlinedFnID && "Invalid outlined function ID!");
  (void)OutlinedFnID;
 
  // Return value of the runtime offloading call.
  Value *Return = nullptr;
 
  // Arguments for the target kernel.
  SmallVector<Value *> ArgsVector;
  getKernelArgsVector(Args, Builder, ArgsVector);
 
  // The target region is an outlined function launched by the runtime
  // via calls to __tgt_target_kernel().
  //
  // Note that on the host and CPU targets, the runtime implementation of
  // these calls simply call the outlined function without forking threads.
  // The outlined functions themselves have runtime calls to
  // __kmpc_fork_teams() and __kmpc_fork() for this purpose, codegen'd by
  // the compiler in emitTeamsCall() and emitParallelCall().
  //
  // In contrast, on the NVPTX target, the implementation of
  // __tgt_target_teams() launches a GPU kernel with the requested number
  // of teams and threads so no additional calls to the runtime are required.
  // Check the error code and execute the host version if required.
  Builder.restoreIP(emitTargetKernel(Builder, AllocaIP, Return, RTLoc, DeviceID,
                                     Args.NumTeams.front(), Args.NumThreads,
                                     OutlinedFnID, ArgsVector));
 
  BasicBlock *OffloadFailedBlock =
      BasicBlock::Create(Builder.getContext(), "omp_offload.failed");
  BasicBlock *OffloadContBlock =
      BasicBlock::Create(Builder.getContext(), "omp_offload.cont");
  Value *Failed = Builder.CreateIsNotNull(Return);
  Builder.CreateCondBr(Failed, OffloadFailedBlock, OffloadContBlock);
 
  auto CurFn = Builder.GetInsertBlock()->getParent();
  emitBlock(OffloadFailedBlock, CurFn);
  Builder.restoreIP(emitTargetCallFallbackCB(Builder.saveIP()));
  emitBranch(OffloadContBlock);
  emitBlock(OffloadContBlock, CurFn, /*IsFinished=*/true);
  return Builder.saveIP();
}
 
void OpenMPIRBuilder::emitCancelationCheckImpl(Value *CancelFlag,
                                               omp::Directive CanceledDirective,
                                               FinalizeCallbackTy ExitCB) {
  assert(isLastFinalizationInfoCancellable(CanceledDirective) &&
         "Unexpected cancellation!");
 
  // For a cancel barrier we create two new blocks.
  BasicBlock *BB = Builder.GetInsertBlock();
  BasicBlock *NonCancellationBlock;
  if (Builder.GetInsertPoint() == BB->end()) {
    // TODO: This branch will not be needed once we moved to the
    // OpenMPIRBuilder codegen completely.
    NonCancellationBlock = BasicBlock::Create(
        BB->getContext(), BB->getName() + ".cont", BB->getParent());
  } else {
    NonCancellationBlock = SplitBlock(BB, &*Builder.GetInsertPoint());
    BB->getTerminator()->eraseFromParent();
    Builder.SetInsertPoint(BB);
  }
  BasicBlock *CancellationBlock = BasicBlock::Create(
      BB->getContext(), BB->getName() + ".cncl", BB->getParent());
 
  // Jump to them based on the return value.
  Value *Cmp = Builder.CreateIsNull(CancelFlag);
  Builder.CreateCondBr(Cmp, NonCancellationBlock, CancellationBlock,
                       /* TODO weight */ nullptr, nullptr);
 
  // From the cancellation block we finalize all variables and go to the
  // post finalization block that is known to the FiniCB callback.
  Builder.SetInsertPoint(CancellationBlock);
  if (ExitCB)
    ExitCB(Builder.saveIP());
  auto &FI = FinalizationStack.back();
  FI.FiniCB(Builder.saveIP());
 
  // The continuation block is where code generation continues.
  Builder.SetInsertPoint(NonCancellationBlock, NonCancellationBlock->begin());
}
 
// Callback used to create OpenMP runtime calls to support
// omp parallel clause for the device.
// We need to use this callback to replace call to the OutlinedFn in OuterFn
// by the call to the OpenMP DeviceRTL runtime function (kmpc_parallel_51)
static void targetParallelCallback(
    OpenMPIRBuilder *OMPIRBuilder, Function &OutlinedFn, Function *OuterFn,
    BasicBlock *OuterAllocaBB, Value *Ident, Value *IfCondition,
    Value *NumThreads, Instruction *PrivTID, AllocaInst *PrivTIDAddr,
    Value *ThreadID, const SmallVector<Instruction *, 4> &ToBeDeleted) {
  // Add some known attributes.
  IRBuilder<> &Builder = OMPIRBuilder->Builder;
  OutlinedFn.addParamAttr(0, Attribute::NoAlias);
  OutlinedFn.addParamAttr(1, Attribute::NoAlias);
  OutlinedFn.addParamAttr(0, Attribute::NoUndef);
  OutlinedFn.addParamAttr(1, Attribute::NoUndef);
  OutlinedFn.addFnAttr(Attribute::NoUnwind);
 
  assert(OutlinedFn.arg_size() >= 2 &&
         "Expected at least tid and bounded tid as arguments");
  unsigned NumCapturedVars = OutlinedFn.arg_size() - /* tid & bounded tid */ 2;
 
  CallInst *CI = cast<CallInst>(OutlinedFn.user_back());
  assert(CI && "Expected call instruction to outlined function");
  CI->getParent()->setName("omp_parallel");
 
  Builder.SetInsertPoint(CI);
  Type *PtrTy = OMPIRBuilder->VoidPtr;
  Value *NullPtrValue = Constant::getNullValue(PtrTy);
 
  // Add alloca for kernel args
  OpenMPIRBuilder ::InsertPointTy CurrentIP = Builder.saveIP();
  Builder.SetInsertPoint(OuterAllocaBB, OuterAllocaBB->getFirstInsertionPt());
  AllocaInst *ArgsAlloca =
      Builder.CreateAlloca(ArrayType::get(PtrTy, NumCapturedVars));
  Value *Args = ArgsAlloca;
  // Add address space cast if array for storing arguments is not allocated
  // in address space 0
  if (ArgsAlloca->getAddressSpace())
    Args = Builder.CreatePointerCast(ArgsAlloca, PtrTy);
  Builder.restoreIP(CurrentIP);
 
  // Store captured vars which are used by kmpc_parallel_51
  for (unsigned Idx = 0; Idx < NumCapturedVars; Idx++) {
    Value *V = *(CI->arg_begin() + 2 + Idx);
    Value *StoreAddress = Builder.CreateConstInBoundsGEP2_64(
        ArrayType::get(PtrTy, NumCapturedVars), Args, 0, Idx);
    Builder.CreateStore(V, StoreAddress);
  }
 
  Value *Cond =
      IfCondition ? Builder.CreateSExtOrTrunc(IfCondition, OMPIRBuilder->Int32)
                  : Builder.getInt32(1);
 
  // Build kmpc_parallel_51 call
  Value *Parallel51CallArgs[] = {
      /* identifier*/ Ident,
      /* global thread num*/ ThreadID,
      /* if expression */ Cond,
      /* number of threads */ NumThreads ? NumThreads : Builder.getInt32(-1),
      /* Proc bind */ Builder.getInt32(-1),
      /* outlined function */
      Builder.CreateBitCast(&OutlinedFn, OMPIRBuilder->ParallelTaskPtr),
      /* wrapper function */ NullPtrValue,
      /* arguments of the outlined funciton*/ Args,
      /* number of arguments */ Builder.getInt64(NumCapturedVars)};
 
  FunctionCallee RTLFn =
      OMPIRBuilder->getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_parallel_51);
 
  Builder.CreateCall(RTLFn, Parallel51CallArgs);
 
  LLVM_DEBUG(dbgs() << "With kmpc_parallel_51 placed: "
                    << *Builder.GetInsertBlock()->getParent() << "\n");
 
  // Initialize the local TID stack location with the argument value.
  Builder.SetInsertPoint(PrivTID);
  Function::arg_iterator OutlinedAI = OutlinedFn.arg_begin();
  Builder.CreateStore(Builder.CreateLoad(OMPIRBuilder->Int32, OutlinedAI),
                      PrivTIDAddr);
 
  // Remove redundant call to the outlined function.
  CI->eraseFromParent();
 
  for (Instruction *I : ToBeDeleted) {
    I->eraseFromParent();
  }
}
 
// Callback used to create OpenMP runtime calls to support
// omp parallel clause for the host.
// We need to use this callback to replace call to the OutlinedFn in OuterFn
// by the call to the OpenMP host runtime function ( __kmpc_fork_call[_if])
static void
hostParallelCallback(OpenMPIRBuilder *OMPIRBuilder, Function &OutlinedFn,
                     Function *OuterFn, Value *Ident, Value *IfCondition,
                     Instruction *PrivTID, AllocaInst *PrivTIDAddr,
                     const SmallVector<Instruction *, 4> &ToBeDeleted) {
  IRBuilder<> &Builder = OMPIRBuilder->Builder;
  FunctionCallee RTLFn;
  if (IfCondition) {
    RTLFn =
        OMPIRBuilder->getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_fork_call_if);
  } else {
    RTLFn =
        OMPIRBuilder->getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_fork_call);
  }
  if (auto *F = dyn_cast<Function>(RTLFn.getCallee())) {
    if (!F->hasMetadata(LLVMContext::MD_callback)) {
      LLVMContext &Ctx = F->getContext();
      MDBuilder MDB(Ctx);
      // Annotate the callback behavior of the __kmpc_fork_call:
      //  - The callback callee is argument number 2 (microtask).
      //  - The first two arguments of the callback callee are unknown (-1).
      //  - All variadic arguments to the __kmpc_fork_call are passed to the
      //    callback callee.
      F->addMetadata(LLVMContext::MD_callback,
                     *MDNode::get(Ctx, {MDB.createCallbackEncoding(
                                           2, {-1, -1},
                                           /* VarArgsArePassed */ true)}));
    }
  }
  // Add some known attributes.
  OutlinedFn.addParamAttr(0, Attribute::NoAlias);
  OutlinedFn.addParamAttr(1, Attribute::NoAlias);
  OutlinedFn.addFnAttr(Attribute::NoUnwind);
 
  assert(OutlinedFn.arg_size() >= 2 &&
         "Expected at least tid and bounded tid as arguments");
  unsigned NumCapturedVars = OutlinedFn.arg_size() - /* tid & bounded tid */ 2;
 
  CallInst *CI = cast<CallInst>(OutlinedFn.user_back());
  CI->getParent()->setName("omp_parallel");
  Builder.SetInsertPoint(CI);
 
  // Build call __kmpc_fork_call[_if](Ident, n, microtask, var1, .., varn);
  Value *ForkCallArgs[] = {
      Ident, Builder.getInt32(NumCapturedVars),
      Builder.CreateBitCast(&OutlinedFn, OMPIRBuilder->ParallelTaskPtr)};
 
  SmallVector<Value *, 16> RealArgs;
  RealArgs.append(std::begin(ForkCallArgs), std::end(ForkCallArgs));
  if (IfCondition) {
    Value *Cond = Builder.CreateSExtOrTrunc(IfCondition, OMPIRBuilder->Int32);
    RealArgs.push_back(Cond);
  }
  RealArgs.append(CI->arg_begin() + /* tid & bound tid */ 2, CI->arg_end());
 
  // __kmpc_fork_call_if always expects a void ptr as the last argument
  // If there are no arguments, pass a null pointer.
  auto PtrTy = OMPIRBuilder->VoidPtr;
  if (IfCondition && NumCapturedVars == 0) {
    Value *NullPtrValue = Constant::getNullValue(PtrTy);
    RealArgs.push_back(NullPtrValue);
  }
  if (IfCondition && RealArgs.back()->getType() != PtrTy)
    RealArgs.back() = Builder.CreateBitCast(RealArgs.back(), PtrTy);
 
  Builder.CreateCall(RTLFn, RealArgs);
 
  LLVM_DEBUG(dbgs() << "With fork_call placed: "
                    << *Builder.GetInsertBlock()->getParent() << "\n");
 
  // Initialize the local TID stack location with the argument value.
  Builder.SetInsertPoint(PrivTID);
  Function::arg_iterator OutlinedAI = OutlinedFn.arg_begin();
  Builder.CreateStore(Builder.CreateLoad(OMPIRBuilder->Int32, OutlinedAI),
                      PrivTIDAddr);
 
  // Remove redundant call to the outlined function.
  CI->eraseFromParent();
 
  for (Instruction *I : ToBeDeleted) {
    I->eraseFromParent();
  }
}
 
IRBuilder<>::InsertPoint OpenMPIRBuilder::createParallel(
    const LocationDescription &Loc, InsertPointTy OuterAllocaIP,
    BodyGenCallbackTy BodyGenCB, PrivatizeCallbackTy PrivCB,
    FinalizeCallbackTy FiniCB, Value *IfCondition, Value *NumThreads,
    omp::ProcBindKind ProcBind, bool IsCancellable) {
  assert(!isConflictIP(Loc.IP, OuterAllocaIP) && "IPs must not be ambiguous");
 
  if (!updateToLocation(Loc))
    return Loc.IP;
 
  uint32_t SrcLocStrSize;
  Constant *SrcLocStr = getOrCreateSrcLocStr(Loc, SrcLocStrSize);
  Value *Ident = getOrCreateIdent(SrcLocStr, SrcLocStrSize);
  Value *ThreadID = getOrCreateThreadID(Ident);
  // If we generate code for the target device, we need to allocate
  // struct for aggregate params in the device default alloca address space.
  // OpenMP runtime requires that the params of the extracted functions are
  // passed as zero address space pointers. This flag ensures that extracted
  // function arguments are declared in zero address space
  bool ArgsInZeroAddressSpace = Config.isTargetDevice();
 
  // Build call __kmpc_push_num_threads(&Ident, global_tid, num_threads)
  // only if we compile for host side.
  if (NumThreads && !Config.isTargetDevice()) {
    Value *Args[] = {
        Ident, ThreadID,
        Builder.CreateIntCast(NumThreads, Int32, /*isSigned*/ false)};
    Builder.CreateCall(
        getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_push_num_threads), Args);
  }
 
  if (ProcBind != OMP_PROC_BIND_default) {
    // Build call __kmpc_push_proc_bind(&Ident, global_tid, proc_bind)
    Value *Args[] = {
        Ident, ThreadID,
        ConstantInt::get(Int32, unsigned(ProcBind), /*isSigned=*/true)};
    Builder.CreateCall(
        getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_push_proc_bind), Args);
  }
 
  BasicBlock *InsertBB = Builder.GetInsertBlock();
  Function *OuterFn = InsertBB->getParent();
 
  // Save the outer alloca block because the insertion iterator may get
  // invalidated and we still need this later.
  BasicBlock *OuterAllocaBlock = OuterAllocaIP.getBlock();
 
  // Vector to remember instructions we used only during the modeling but which
  // we want to delete at the end.
  SmallVector<Instruction *, 4> ToBeDeleted;
 
  // Change the location to the outer alloca insertion point to create and
  // initialize the allocas we pass into the parallel region.
  InsertPointTy NewOuter(OuterAllocaBlock, OuterAllocaBlock->begin());
  Builder.restoreIP(NewOuter);
  AllocaInst *TIDAddrAlloca = Builder.CreateAlloca(Int32, nullptr, "tid.addr");
  AllocaInst *ZeroAddrAlloca =
      Builder.CreateAlloca(Int32, nullptr, "zero.addr");
  Instruction *TIDAddr = TIDAddrAlloca;
  Instruction *ZeroAddr = ZeroAddrAlloca;
  if (ArgsInZeroAddressSpace && M.getDataLayout().getAllocaAddrSpace() != 0) {
    // Add additional casts to enforce pointers in zero address space
    TIDAddr = new AddrSpaceCastInst(
        TIDAddrAlloca, PointerType ::get(M.getContext(), 0), "tid.addr.ascast");
    TIDAddr->insertAfter(TIDAddrAlloca);
    ToBeDeleted.push_back(TIDAddr);
    ZeroAddr = new AddrSpaceCastInst(ZeroAddrAlloca,
                                     PointerType ::get(M.getContext(), 0),
                                     "zero.addr.ascast");
    ZeroAddr->insertAfter(ZeroAddrAlloca);
    ToBeDeleted.push_back(ZeroAddr);
  }
 
  // We only need TIDAddr and ZeroAddr for modeling purposes to get the
  // associated arguments in the outlined function, so we delete them later.
  ToBeDeleted.push_back(TIDAddrAlloca);
  ToBeDeleted.push_back(ZeroAddrAlloca);
 
  // Create an artificial insertion point that will also ensure the blocks we
  // are about to split are not degenerated.
  auto *UI = new UnreachableInst(Builder.getContext(), InsertBB);
 
  BasicBlock *EntryBB = UI->getParent();
  BasicBlock *PRegEntryBB = EntryBB->splitBasicBlock(UI, "omp.par.entry");
  BasicBlock *PRegBodyBB = PRegEntryBB->splitBasicBlock(UI, "omp.par.region");
  BasicBlock *PRegPreFiniBB =
      PRegBodyBB->splitBasicBlock(UI, "omp.par.pre_finalize");
  BasicBlock *PRegExitBB = PRegPreFiniBB->splitBasicBlock(UI, "omp.par.exit");
 
  auto FiniCBWrapper = [&](InsertPointTy IP) {
    // Hide "open-ended" blocks from the given FiniCB by setting the right jump
    // target to the region exit block.
    if (IP.getBlock()->end() == IP.getPoint()) {
      IRBuilder<>::InsertPointGuard IPG(Builder);
      Builder.restoreIP(IP);
      Instruction *I = Builder.CreateBr(PRegExitBB);
      IP = InsertPointTy(I->getParent(), I->getIterator());
    }
    assert(IP.getBlock()->getTerminator()->getNumSuccessors() == 1 &&
           IP.getBlock()->getTerminator()->getSuccessor(0) == PRegExitBB &&
           "Unexpected insertion point for finalization call!");
    return FiniCB(IP);
  };
 
  FinalizationStack.push_back({FiniCBWrapper, OMPD_parallel, IsCancellable});
 
  // Generate the privatization allocas in the block that will become the entry
  // of the outlined function.
  Builder.SetInsertPoint(PRegEntryBB->getTerminator());
  InsertPointTy InnerAllocaIP = Builder.saveIP();
 
  AllocaInst *PrivTIDAddr =
      Builder.CreateAlloca(Int32, nullptr, "tid.addr.local");
  Instruction *PrivTID = Builder.CreateLoad(Int32, PrivTIDAddr, "tid");
 
  // Add some fake uses for OpenMP provided arguments.
  ToBeDeleted.push_back(Builder.CreateLoad(Int32, TIDAddr, "tid.addr.use"));
  Instruction *ZeroAddrUse =
      Builder.CreateLoad(Int32, ZeroAddr, "zero.addr.use");
  ToBeDeleted.push_back(ZeroAddrUse);
 
  // EntryBB
  //   |
  //   V
  // PRegionEntryBB         <- Privatization allocas are placed here.
  //   |
  //   V
  // PRegionBodyBB          <- BodeGen is invoked here.
  //   |
  //   V
  // PRegPreFiniBB          <- The block we will start finalization from.
  //   |
  //   V
  // PRegionExitBB          <- A common exit to simplify block collection.
  //
 
  LLVM_DEBUG(dbgs() << "Before body codegen: " << *OuterFn << "\n");
 
  // Let the caller create the body.
  assert(BodyGenCB && "Expected body generation callback!");
  InsertPointTy CodeGenIP(PRegBodyBB, PRegBodyBB->begin());
  BodyGenCB(InnerAllocaIP, CodeGenIP);
 
  LLVM_DEBUG(dbgs() << "After  body codegen: " << *OuterFn << "\n");
 
  OutlineInfo OI;
  if (Config.isTargetDevice()) {
    // Generate OpenMP target specific runtime call
    OI.PostOutlineCB = [=, ToBeDeletedVec =
                               std::move(ToBeDeleted)](Function &OutlinedFn) {
      targetParallelCallback(this, OutlinedFn, OuterFn, OuterAllocaBlock, Ident,
                             IfCondition, NumThreads, PrivTID, PrivTIDAddr,
                             ThreadID, ToBeDeletedVec);
    };
  } else {
    // Generate OpenMP host runtime call
    OI.PostOutlineCB = [=, ToBeDeletedVec =
                               std::move(ToBeDeleted)](Function &OutlinedFn) {
      hostParallelCallback(this, OutlinedFn, OuterFn, Ident, IfCondition,
                           PrivTID, PrivTIDAddr, ToBeDeletedVec);
    };
  }
 
  OI.OuterAllocaBB = OuterAllocaBlock;
  OI.EntryBB = PRegEntryBB;
  OI.ExitBB = PRegExitBB;
 
  SmallPtrSet<BasicBlock *, 32> ParallelRegionBlockSet;
  SmallVector<BasicBlock *, 32> Blocks;
  OI.collectBlocks(ParallelRegionBlockSet, Blocks);
 
  // Ensure a single exit node for the outlined region by creating one.
  // We might have multiple incoming edges to the exit now due to finalizations,
  // e.g., cancel calls that cause the control flow to leave the region.
  BasicBlock *PRegOutlinedExitBB = PRegExitBB;
  PRegExitBB = SplitBlock(PRegExitBB, &*PRegExitBB->getFirstInsertionPt());
  PRegOutlinedExitBB->setName("omp.par.outlined.exit");
  Blocks.push_back(PRegOutlinedExitBB);
 
  CodeExtractorAnalysisCache CEAC(*OuterFn);
  CodeExtractor Extractor(Blocks, /* DominatorTree */ nullptr,
                          /* AggregateArgs */ false,
                          /* BlockFrequencyInfo */ nullptr,
                          /* BranchProbabilityInfo */ nullptr,
                          /* AssumptionCache */ nullptr,
                          /* AllowVarArgs */ true,
                          /* AllowAlloca */ true,
                          /* AllocationBlock */ OuterAllocaBlock,
                          /* Suffix */ ".omp_par", ArgsInZeroAddressSpace);
 
  // Find inputs to, outputs from the code region.
  BasicBlock *CommonExit = nullptr;
  SetVector<Value *> Inputs, Outputs, SinkingCands, HoistingCands;
  Extractor.findAllocas(CEAC, SinkingCands, HoistingCands, CommonExit);
  Extractor.findInputsOutputs(Inputs, Outputs, SinkingCands);
 
  LLVM_DEBUG(dbgs() << "Before privatization: " << *OuterFn << "\n");
 
  FunctionCallee TIDRTLFn =
      getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_global_thread_num);
 
  auto PrivHelper = [&](Value &V) {
    if (&V == TIDAddr || &V == ZeroAddr) {
      OI.ExcludeArgsFromAggregate.push_back(&V);
      return;
    }
 
    SetVector<Use *> Uses;
    for (Use &U : V.uses())
      if (auto *UserI = dyn_cast<Instruction>(U.getUser()))
        if (ParallelRegionBlockSet.count(UserI->getParent()))
          Uses.insert(&U);
 
    // __kmpc_fork_call expects extra arguments as pointers. If the input
    // already has a pointer type, everything is fine. Otherwise, store the
    // value onto stack and load it back inside the to-be-outlined region. This
    // will ensure only the pointer will be passed to the function.
    // FIXME: if there are more than 15 trailing arguments, they must be
    // additionally packed in a struct.
    Value *Inner = &V;
    if (!V.getType()->isPointerTy()) {
      IRBuilder<>::InsertPointGuard Guard(Builder);
      LLVM_DEBUG(llvm::dbgs() << "Forwarding input as pointer: " << V << "\n");
 
      Builder.restoreIP(OuterAllocaIP);
      Value *Ptr =
          Builder.CreateAlloca(V.getType(), nullptr, V.getName() + ".reloaded");
 
      // Store to stack at end of the block that currently branches to the entry
      // block of the to-be-outlined region.
      Builder.SetInsertPoint(InsertBB,
                             InsertBB->getTerminator()->getIterator());
      Builder.CreateStore(&V, Ptr);
 
      // Load back next to allocations in the to-be-outlined region.
      Builder.restoreIP(InnerAllocaIP);
      Inner = Builder.CreateLoad(V.getType(), Ptr);
    }
 
    Value *ReplacementValue = nullptr;
    CallInst *CI = dyn_cast<CallInst>(&V);
    if (CI && CI->getCalledFunction() == TIDRTLFn.getCallee()) {
      ReplacementValue = PrivTID;
    } else {
      Builder.restoreIP(
          PrivCB(InnerAllocaIP, Builder.saveIP(), V, *Inner, ReplacementValue));
      InnerAllocaIP = {
          InnerAllocaIP.getBlock(),
          InnerAllocaIP.getBlock()->getTerminator()->getIterator()};
 
      assert(ReplacementValue &&
             "Expected copy/create callback to set replacement value!");
      if (ReplacementValue == &V)
        return;
    }
 
    for (Use *UPtr : Uses)
      UPtr->set(ReplacementValue);
  };
 
  // Reset the inner alloca insertion as it will be used for loading the values
  // wrapped into pointers before passing them into the to-be-outlined region.
  // Configure it to insert immediately after the fake use of zero address so
  // that they are available in the generated body and so that the
  // OpenMP-related values (thread ID and zero address pointers) remain leading
  // in the argument list.
  InnerAllocaIP = IRBuilder<>::InsertPoint(
      ZeroAddrUse->getParent(), ZeroAddrUse->getNextNode()->getIterator());
 
  // Reset the outer alloca insertion point to the entry of the relevant block
  // in case it was invalidated.
  OuterAllocaIP = IRBuilder<>::InsertPoint(
      OuterAllocaBlock, OuterAllocaBlock->getFirstInsertionPt());
 
  for (Value *Input : Inputs) {
    LLVM_DEBUG(dbgs() << "Captured input: " << *Input << "\n");
    PrivHelper(*Input);
  }
  LLVM_DEBUG({
    for (Value *Output : Outputs)
      LLVM_DEBUG(dbgs() << "Captured output: " << *Output << "\n");
  });
  assert(Outputs.empty() &&
         "OpenMP outlining should not produce live-out values!");
 
  LLVM_DEBUG(dbgs() << "After  privatization: " << *OuterFn << "\n");
  LLVM_DEBUG({
    for (auto *BB : Blocks)
      dbgs() << " PBR: " << BB->getName() << "\n";
  });
 
  // Adjust the finalization stack, verify the adjustment, and call the
  // finalize function a last time to finalize values between the pre-fini
  // block and the exit block if we left the parallel "the normal way".
  auto FiniInfo = FinalizationStack.pop_back_val();
  (void)FiniInfo;
  assert(FiniInfo.DK == OMPD_parallel &&
         "Unexpected finalization stack state!");
 
  Instruction *PRegPreFiniTI = PRegPreFiniBB->getTerminator();
 
  InsertPointTy PreFiniIP(PRegPreFiniBB, PRegPreFiniTI->getIterator());
  FiniCB(PreFiniIP);
 
  // Register the outlined info.
  addOutlineInfo(std::move(OI));
 
  InsertPointTy AfterIP(UI->getParent(), UI->getParent()->end());
  UI->eraseFromParent();
 
  return AfterIP;
}
 
void OpenMPIRBuilder::emitFlush(const LocationDescription &Loc) {
  // Build call void __kmpc_flush(ident_t *loc)
  uint32_t SrcLocStrSize;
  Constant *SrcLocStr = getOrCreateSrcLocStr(Loc, SrcLocStrSize);
  Value *Args[] = {getOrCreateIdent(SrcLocStr, SrcLocStrSize)};
 
  Builder.CreateCall(getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_flush), Args);
}
 
void OpenMPIRBuilder::createFlush(const LocationDescription &Loc) {
  if (!updateToLocation(Loc))
    return;
  emitFlush(Loc);
}
 
void OpenMPIRBuilder::emitTaskwaitImpl(const LocationDescription &Loc) {
  // Build call kmp_int32 __kmpc_omp_taskwait(ident_t *loc, kmp_int32
  // global_tid);
  uint32_t SrcLocStrSize;
  Constant *SrcLocStr = getOrCreateSrcLocStr(Loc, SrcLocStrSize);
  Value *Ident = getOrCreateIdent(SrcLocStr, SrcLocStrSize);
  Value *Args[] = {Ident, getOrCreateThreadID(Ident)};
 
  // Ignore return result until untied tasks are supported.
  Builder.CreateCall(getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_omp_taskwait),
                     Args);
}
 
void OpenMPIRBuilder::createTaskwait(const LocationDescription &Loc) {
  if (!updateToLocation(Loc))
    return;
  emitTaskwaitImpl(Loc);
}
 
void OpenMPIRBuilder::emitTaskyieldImpl(const LocationDescription &Loc) {
  // Build call __kmpc_omp_taskyield(loc, thread_id, 0);
  uint32_t SrcLocStrSize;
  Constant *SrcLocStr = getOrCreateSrcLocStr(Loc, SrcLocStrSize);
  Value *Ident = getOrCreateIdent(SrcLocStr, SrcLocStrSize);
  Constant *I32Null = ConstantInt::getNullValue(Int32);
  Value *Args[] = {Ident, getOrCreateThreadID(Ident), I32Null};
 
  Builder.CreateCall(getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_omp_taskyield),
                     Args);
}
 
void OpenMPIRBuilder::createTaskyield(const LocationDescription &Loc) {
  if (!updateToLocation(Loc))
    return;
  emitTaskyieldImpl(Loc);
}
 
// Processes the dependencies in Dependencies and does the following
// - Allocates space on the stack of an array of DependInfo objects
// - Populates each DependInfo object with relevant information of
//   the corresponding dependence.
// - All code is inserted in the entry block of the current function.
static Value *emitTaskDependencies(
    OpenMPIRBuilder &OMPBuilder,
    SmallVectorImpl<OpenMPIRBuilder::DependData> &Dependencies) {
  // Early return if we have no dependencies to process
  if (Dependencies.empty())
    return nullptr;
 
  // Given a vector of DependData objects, in this function we create an
  // array on the stack that holds kmp_dep_info objects corresponding
  // to each dependency. This is then passed to the OpenMP runtime.
  // For example, if there are 'n' dependencies then the following psedo
  // code is generated. Assume the first dependence is on a variable 'a'
  //
  // \code{c}
  // DepArray = alloc(n x sizeof(kmp_depend_info);
  // idx = 0;
  // DepArray[idx].base_addr = ptrtoint(&a);
  // DepArray[idx].len = 8;
  // DepArray[idx].flags = Dep.DepKind; /*(See OMPContants.h for DepKind)*/
  // ++idx;
  // DepArray[idx].base_addr = ...;
  // \endcode
 
  IRBuilderBase &Builder = OMPBuilder.Builder;
  Type *DependInfo = OMPBuilder.DependInfo;
  Module &M = OMPBuilder.M;
 
  Value *DepArray = nullptr;
  OpenMPIRBuilder::InsertPointTy OldIP = Builder.saveIP();
  Builder.SetInsertPoint(
      OldIP.getBlock()->getParent()->getEntryBlock().getTerminator());
 
  Type *DepArrayTy = ArrayType::get(DependInfo, Dependencies.size());
  DepArray = Builder.CreateAlloca(DepArrayTy, nullptr, ".dep.arr.addr");
 
  for (const auto &[DepIdx, Dep] : enumerate(Dependencies)) {
    Value *Base =
        Builder.CreateConstInBoundsGEP2_64(DepArrayTy, DepArray, 0, DepIdx);
    // Store the pointer to the variable
    Value *Addr = Builder.CreateStructGEP(
        DependInfo, Base,
        static_cast<unsigned int>(RTLDependInfoFields::BaseAddr));
    Value *DepValPtr = Builder.CreatePtrToInt(Dep.DepVal, Builder.getInt64Ty());
    Builder.CreateStore(DepValPtr, Addr);
    // Store the size of the variable
    Value *Size = Builder.CreateStructGEP(
        DependInfo, Base, static_cast<unsigned int>(RTLDependInfoFields::Len));
    Builder.CreateStore(
        Builder.getInt64(M.getDataLayout().getTypeStoreSize(Dep.DepValueType)),
        Size);
    // Store the dependency kind
    Value *Flags = Builder.CreateStructGEP(
        DependInfo, Base,
        static_cast<unsigned int>(RTLDependInfoFields::Flags));
    Builder.CreateStore(
        ConstantInt::get(Builder.getInt8Ty(),
                         static_cast<unsigned int>(Dep.DepKind)),
        Flags);
  }
  Builder.restoreIP(OldIP);
  return DepArray;
}
 
OpenMPIRBuilder::InsertPointTy
OpenMPIRBuilder::createTask(const LocationDescription &Loc,
                            InsertPointTy AllocaIP, BodyGenCallbackTy BodyGenCB,
                            bool Tied, Value *Final, Value *IfCondition,
                            SmallVector<DependData> Dependencies) {
 
  if (!updateToLocation(Loc))
    return InsertPointTy();
 
  uint32_t SrcLocStrSize;
  Constant *SrcLocStr = getOrCreateSrcLocStr(Loc, SrcLocStrSize);
  Value *Ident = getOrCreateIdent(SrcLocStr, SrcLocStrSize);
  // The current basic block is split into four basic blocks. After outlining,
  // they will be mapped as follows:
  // ```
  // def current_fn() {
  //   current_basic_block:
  //     br label %task.exit
  //   task.exit:
  //     ; instructions after task
  // }
  // def outlined_fn() {
  //   task.alloca:
  //     br label %task.body
  //   task.body:
  //     ret void
  // }
  // ```
  BasicBlock *TaskExitBB = splitBB(Builder, /*CreateBranch=*/true, "task.exit");
  BasicBlock *TaskBodyBB = splitBB(Builder, /*CreateBranch=*/true, "task.body");
  BasicBlock *TaskAllocaBB =
      splitBB(Builder, /*CreateBranch=*/true, "task.alloca");
 
  InsertPointTy TaskAllocaIP =
      InsertPointTy(TaskAllocaBB, TaskAllocaBB->begin());
  InsertPointTy TaskBodyIP = InsertPointTy(TaskBodyBB, TaskBodyBB->begin());
  BodyGenCB(TaskAllocaIP, TaskBodyIP);
 
  OutlineInfo OI;
  OI.EntryBB = TaskAllocaBB;
  OI.OuterAllocaBB = AllocaIP.getBlock();
  OI.ExitBB = TaskExitBB;
 
  // Add the thread ID argument.
  SmallVector<Instruction *, 4> ToBeDeleted;
  OI.ExcludeArgsFromAggregate.push_back(createFakeIntVal(
      Builder, AllocaIP, ToBeDeleted, TaskAllocaIP, "global.tid", false));
 
  OI.PostOutlineCB = [this, Ident, Tied, Final, IfCondition, Dependencies,
                      TaskAllocaBB, ToBeDeleted](Function &OutlinedFn) mutable {
    // Replace the Stale CI by appropriate RTL function call.
    assert(OutlinedFn.getNumUses() == 1 &&
           "there must be a single user for the outlined function");
    CallInst *StaleCI = cast<CallInst>(OutlinedFn.user_back());
 
    // HasShareds is true if any variables are captured in the outlined region,
    // false otherwise.
    bool HasShareds = StaleCI->arg_size() > 1;
    Builder.SetInsertPoint(StaleCI);
 
    // Gather the arguments for emitting the runtime call for
    // @__kmpc_omp_task_alloc
    Function *TaskAllocFn =
        getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_omp_task_alloc);
 
    // Arguments - `loc_ref` (Ident) and `gtid` (ThreadID)
    // call.
    Value *ThreadID = getOrCreateThreadID(Ident);
 
    // Argument - `flags`
    // Task is tied iff (Flags & 1) == 1.
    // Task is untied iff (Flags & 1) == 0.
    // Task is final iff (Flags & 2) == 2.
    // Task is not final iff (Flags & 2) == 0.
    // TODO: Handle the other flags.
    Value *Flags = Builder.getInt32(Tied);
    if (Final) {
      Value *FinalFlag =
          Builder.CreateSelect(Final, Builder.getInt32(2), Builder.getInt32(0));
      Flags = Builder.CreateOr(FinalFlag, Flags);
    }
 
    // Argument - `sizeof_kmp_task_t` (TaskSize)
    // Tasksize refers to the size in bytes of kmp_task_t data structure
    // including private vars accessed in task.
    // TODO: add kmp_task_t_with_privates (privates)
    Value *TaskSize = Builder.getInt64(
        divideCeil(M.getDataLayout().getTypeSizeInBits(Task), 8));
 
    // Argument - `sizeof_shareds` (SharedsSize)
    // SharedsSize refers to the shareds array size in the kmp_task_t data
    // structure.
    Value *SharedsSize = Builder.getInt64(0);
    if (HasShareds) {
      AllocaInst *ArgStructAlloca =
          dyn_cast<AllocaInst>(StaleCI->getArgOperand(1));
      assert(ArgStructAlloca &&
             "Unable to find the alloca instruction corresponding to arguments "
             "for extracted function");
      StructType *ArgStructType =
          dyn_cast<StructType>(ArgStructAlloca->getAllocatedType());
      assert(ArgStructType && "Unable to find struct type corresponding to "
                              "arguments for extracted function");
      SharedsSize =
          Builder.getInt64(M.getDataLayout().getTypeStoreSize(ArgStructType));
    }
    // Emit the @__kmpc_omp_task_alloc runtime call
    // The runtime call returns a pointer to an area where the task captured
    // variables must be copied before the task is run (TaskData)
    CallInst *TaskData = Builder.CreateCall(
        TaskAllocFn, {/*loc_ref=*/Ident, /*gtid=*/ThreadID, /*flags=*/Flags,
                      /*sizeof_task=*/TaskSize, /*sizeof_shared=*/SharedsSize,
                      /*task_func=*/&OutlinedFn});
 
    // Copy the arguments for outlined function
    if (HasShareds) {
      Value *Shareds = StaleCI->getArgOperand(1);
      Align Alignment = TaskData->getPointerAlignment(M.getDataLayout());
      Value *TaskShareds = Builder.CreateLoad(VoidPtr, TaskData);
      Builder.CreateMemCpy(TaskShareds, Alignment, Shareds, Alignment,
                           SharedsSize);
    }
 
    Value *DepArray = nullptr;
    if (Dependencies.size()) {
      InsertPointTy OldIP = Builder.saveIP();
      Builder.SetInsertPoint(
          &OldIP.getBlock()->getParent()->getEntryBlock().back());
 
      Type *DepArrayTy = ArrayType::get(DependInfo, Dependencies.size());
      DepArray = Builder.CreateAlloca(DepArrayTy, nullptr, ".dep.arr.addr");
 
      unsigned P = 0;
      for (const DependData &Dep : Dependencies) {
        Value *Base =
            Builder.CreateConstInBoundsGEP2_64(DepArrayTy, DepArray, 0, P);
        // Store the pointer to the variable
        Value *Addr = Builder.CreateStructGEP(
            DependInfo, Base,
            static_cast<unsigned int>(RTLDependInfoFields::BaseAddr));
        Value *DepValPtr =
            Builder.CreatePtrToInt(Dep.DepVal, Builder.getInt64Ty());
        Builder.CreateStore(DepValPtr, Addr);
        // Store the size of the variable
        Value *Size = Builder.CreateStructGEP(
            DependInfo, Base,
            static_cast<unsigned int>(RTLDependInfoFields::Len));
        Builder.CreateStore(Builder.getInt64(M.getDataLayout().getTypeStoreSize(
                                Dep.DepValueType)),
                            Size);
        // Store the dependency kind
        Value *Flags = Builder.CreateStructGEP(
            DependInfo, Base,
            static_cast<unsigned int>(RTLDependInfoFields::Flags));
        Builder.CreateStore(
            ConstantInt::get(Builder.getInt8Ty(),
                             static_cast<unsigned int>(Dep.DepKind)),
            Flags);
        ++P;
      }
 
      Builder.restoreIP(OldIP);
    }
 
    // In the presence of the `if` clause, the following IR is generated:
    //    ...
    //    %data = call @__kmpc_omp_task_alloc(...)
    //    br i1 %if_condition, label %then, label %else
    //  then:
    //    call @__kmpc_omp_task(...)
    //    br label %exit
    //  else:
    //    ;; Wait for resolution of dependencies, if any, before
    //    ;; beginning the task
    //    call @__kmpc_omp_wait_deps(...)
    //    call @__kmpc_omp_task_begin_if0(...)
    //    call @outlined_fn(...)
    //    call @__kmpc_omp_task_complete_if0(...)
    //    br label %exit
    //  exit:
    //    ...
    if (IfCondition) {
      // `SplitBlockAndInsertIfThenElse` requires the block to have a
      // terminator.
      splitBB(Builder, /*CreateBranch=*/true, "if.end");
      Instruction *IfTerminator =
          Builder.GetInsertPoint()->getParent()->getTerminator();
      Instruction *ThenTI = IfTerminator, *ElseTI = nullptr;
      Builder.SetInsertPoint(IfTerminator);
      SplitBlockAndInsertIfThenElse(IfCondition, IfTerminator, &ThenTI,
                                    &ElseTI);
      Builder.SetInsertPoint(ElseTI);
 
      if (Dependencies.size()) {
        Function *TaskWaitFn =
            getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_omp_wait_deps);
        Builder.CreateCall(
            TaskWaitFn,
            {Ident, ThreadID, Builder.getInt32(Dependencies.size()), DepArray,
             ConstantInt::get(Builder.getInt32Ty(), 0),
             ConstantPointerNull::get(PointerType::getUnqual(M.getContext()))});
      }
      Function *TaskBeginFn =
          getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_omp_task_begin_if0);
      Function *TaskCompleteFn =
          getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_omp_task_complete_if0);
      Builder.CreateCall(TaskBeginFn, {Ident, ThreadID, TaskData});
      CallInst *CI = nullptr;
      if (HasShareds)
        CI = Builder.CreateCall(&OutlinedFn, {ThreadID, TaskData});
      else
        CI = Builder.CreateCall(&OutlinedFn, {ThreadID});
      CI->setDebugLoc(StaleCI->getDebugLoc());
      Builder.CreateCall(TaskCompleteFn, {Ident, ThreadID, TaskData});
      Builder.SetInsertPoint(ThenTI);
    }
 
    if (Dependencies.size()) {
      Function *TaskFn =
          getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_omp_task_with_deps);
      Builder.CreateCall(
          TaskFn,
          {Ident, ThreadID, TaskData, Builder.getInt32(Dependencies.size()),
           DepArray, ConstantInt::get(Builder.getInt32Ty(), 0),
           ConstantPointerNull::get(PointerType::getUnqual(M.getContext()))});
 
    } else {
      // Emit the @__kmpc_omp_task runtime call to spawn the task
      Function *TaskFn = getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_omp_task);
      Builder.CreateCall(TaskFn, {Ident, ThreadID, TaskData});
    }
 
    StaleCI->eraseFromParent();
 
    Builder.SetInsertPoint(TaskAllocaBB, TaskAllocaBB->begin());
    if (HasShareds) {
      LoadInst *Shareds = Builder.CreateLoad(VoidPtr, OutlinedFn.getArg(1));
      OutlinedFn.getArg(1)->replaceUsesWithIf(
          Shareds, [Shareds](Use &U) { return U.getUser() != Shareds; });
    }
 
    llvm::for_each(llvm::reverse(ToBeDeleted),
                   [](Instruction *I) { I->eraseFromParent(); });
  };
 
  addOutlineInfo(std::move(OI));
  Builder.SetInsertPoint(TaskExitBB, TaskExitBB->begin());
 
  return Builder.saveIP();
}
 
OpenMPIRBuilder::InsertPointTy
OpenMPIRBuilder::createTaskgroup(const LocationDescription &Loc,
                                 InsertPointTy AllocaIP,
                                 BodyGenCallbackTy BodyGenCB) {
  if (!updateToLocation(Loc))
    return InsertPointTy();
 
  uint32_t SrcLocStrSize;
  Constant *SrcLocStr = getOrCreateSrcLocStr(Loc, SrcLocStrSize);
  Value *Ident = getOrCreateIdent(SrcLocStr, SrcLocStrSize);
  Value *ThreadID = getOrCreateThreadID(Ident);
 
  // Emit the @__kmpc_taskgroup runtime call to start the taskgroup
  Function *TaskgroupFn =
      getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_taskgroup);
  Builder.CreateCall(TaskgroupFn, {Ident, ThreadID});
 
  BasicBlock *TaskgroupExitBB = splitBB(Builder, true, "taskgroup.exit");
  BodyGenCB(AllocaIP, Builder.saveIP());
 
  Builder.SetInsertPoint(TaskgroupExitBB);
  // Emit the @__kmpc_end_taskgroup runtime call to end the taskgroup
  Function *EndTaskgroupFn =
      getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_end_taskgroup);
  Builder.CreateCall(EndTaskgroupFn, {Ident, ThreadID});
 
  return Builder.saveIP();
}
 
OpenMPIRBuilder::InsertPointTy OpenMPIRBuilder::createSections(
    const LocationDescription &Loc, InsertPointTy AllocaIP,
    ArrayRef<StorableBodyGenCallbackTy> SectionCBs, PrivatizeCallbackTy PrivCB,
    FinalizeCallbackTy FiniCB, bool IsCancellable, bool IsNowait) {
  assert(!isConflictIP(AllocaIP, Loc.IP) && "Dedicated IP allocas required");
 
  if (!updateToLocation(Loc))
    return Loc.IP;
 
  auto FiniCBWrapper = [&](InsertPointTy IP) {
    if (IP.getBlock()->end() != IP.getPoint())
      return FiniCB(IP);
    // This must be done otherwise any nested constructs using FinalizeOMPRegion
    // will fail because that function requires the Finalization Basic Block to
    // have a terminator, which is already removed by EmitOMPRegionBody.
    // IP is currently at cancelation block.
    // We need to backtrack to the condition block to fetch
    // the exit block and create a branch from cancelation
    // to exit block.
    IRBuilder<>::InsertPointGuard IPG(Builder);
    Builder.restoreIP(IP);
    auto *CaseBB = IP.getBlock()->getSinglePredecessor();
    auto *CondBB = CaseBB->getSinglePredecessor()->getSinglePredecessor();
    auto *ExitBB = CondBB->getTerminator()->getSuccessor(1);
    Instruction *I = Builder.CreateBr(ExitBB);
    IP = InsertPointTy(I->getParent(), I->getIterator());
    return FiniCB(IP);
  };
 
  FinalizationStack.push_back({FiniCBWrapper, OMPD_sections, IsCancellable});
 
  // Each section is emitted as a switch case
  // Each finalization callback is handled from clang.EmitOMPSectionDirective()
  // -> OMP.createSection() which generates the IR for each section
  // Iterate through all sections and emit a switch construct:
  // switch (IV) {
  //   case 0:
  //     <SectionStmt[0]>;
  //     break;
  // ...
  //   case <NumSection> - 1:
  //     <SectionStmt[<NumSection> - 1]>;
  //     break;
  // }
  // ...
  // section_loop.after:
  // <FiniCB>;
  auto LoopBodyGenCB = [&](InsertPointTy CodeGenIP, Value *IndVar) {
    Builder.restoreIP(CodeGenIP);
    BasicBlock *Continue =
        splitBBWithSuffix(Builder, /*CreateBranch=*/false, ".sections.after");
    Function *CurFn = Continue->getParent();
    SwitchInst *SwitchStmt = Builder.CreateSwitch(IndVar, Continue);
 
    unsigned CaseNumber = 0;
    for (auto SectionCB : SectionCBs) {
      BasicBlock *CaseBB = BasicBlock::Create(
          M.getContext(), "omp_section_loop.body.case", CurFn, Continue);
      SwitchStmt->addCase(Builder.getInt32(CaseNumber), CaseBB);
      Builder.SetInsertPoint(CaseBB);
      BranchInst *CaseEndBr = Builder.CreateBr(Continue);
      SectionCB(InsertPointTy(),
                {CaseEndBr->getParent(), CaseEndBr->getIterator()});
      CaseNumber++;
    }
    // remove the existing terminator from body BB since there can be no
    // terminators after switch/case
  };
  // Loop body ends here
  // LowerBound, UpperBound, and STride for createCanonicalLoop
  Type *I32Ty = Type::getInt32Ty(M.getContext());
  Value *LB = ConstantInt::get(I32Ty, 0);
  Value *UB = ConstantInt::get(I32Ty, SectionCBs.size());
  Value *ST = ConstantInt::get(I32Ty, 1);
  llvm::CanonicalLoopInfo *LoopInfo = createCanonicalLoop(
      Loc, LoopBodyGenCB, LB, UB, ST, true, false, AllocaIP, "section_loop");
  InsertPointTy AfterIP =
      applyStaticWorkshareLoop(Loc.DL, LoopInfo, AllocaIP, !IsNowait);
 
  // Apply the finalization callback in LoopAfterBB
  auto FiniInfo = FinalizationStack.pop_back_val();
  assert(FiniInfo.DK == OMPD_sections &&
         "Unexpected finalization stack state!");
  if (FinalizeCallbackTy &CB = FiniInfo.FiniCB) {
    Builder.restoreIP(AfterIP);
    BasicBlock *FiniBB =
        splitBBWithSuffix(Builder, /*CreateBranch=*/true, "sections.fini");
    CB(Builder.saveIP());
    AfterIP = {FiniBB, FiniBB->begin()};
  }
 
  return AfterIP;
}
 
OpenMPIRBuilder::InsertPointTy
OpenMPIRBuilder::createSection(const LocationDescription &Loc,
                               BodyGenCallbackTy BodyGenCB,
                               FinalizeCallbackTy FiniCB) {
  if (!updateToLocation(Loc))
    return Loc.IP;
 
  auto FiniCBWrapper = [&](InsertPointTy IP) {
    if (IP.getBlock()->end() != IP.getPoint())
      return FiniCB(IP);
    // This must be done otherwise any nested constructs using FinalizeOMPRegion
    // will fail because that function requires the Finalization Basic Block to
    // have a terminator, which is already removed by EmitOMPRegionBody.
    // IP is currently at cancelation block.
    // We need to backtrack to the condition block to fetch
    // the exit block and create a branch from cancelation
    // to exit block.
    IRBuilder<>::InsertPointGuard IPG(Builder);
    Builder.restoreIP(IP);
    auto *CaseBB = Loc.IP.getBlock();
    auto *CondBB = CaseBB->getSinglePredecessor()->getSinglePredecessor();
    auto *ExitBB = CondBB->getTerminator()->getSuccessor(1);
    Instruction *I = Builder.CreateBr(ExitBB);
    IP = InsertPointTy(I->getParent(), I->getIterator());
    return FiniCB(IP);
  };
 
  Directive OMPD = Directive::OMPD_sections;
  // Since we are using Finalization Callback here, HasFinalize
  // and IsCancellable have to be true
  return EmitOMPInlinedRegion(OMPD, nullptr, nullptr, BodyGenCB, FiniCBWrapper,
                              /*Conditional*/ false, /*hasFinalize*/ true,
                              /*IsCancellable*/ true);
}
 
static OpenMPIRBuilder::InsertPointTy getInsertPointAfterInstr(Instruction *I) {
  BasicBlock::iterator IT(I);
  IT++;
  return OpenMPIRBuilder::InsertPointTy(I->getParent(), IT);
}
 
void OpenMPIRBuilder::emitUsed(StringRef Name,
                               std::vector<WeakTrackingVH> &List) {
  if (List.empty())
    return;
 
  // Convert List to what ConstantArray needs.
  SmallVector<Constant *, 8> UsedArray;
  UsedArray.resize(List.size());
  for (unsigned I = 0, E = List.size(); I != E; ++I)
    UsedArray[I] = ConstantExpr::getPointerBitCastOrAddrSpaceCast(
        cast<Constant>(&*List[I]), Builder.getPtrTy());
 
  if (UsedArray.empty())
    return;
  ArrayType *ATy = ArrayType::get(Builder.getPtrTy(), UsedArray.size());
 
  auto *GV = new GlobalVariable(M, ATy, false, GlobalValue::AppendingLinkage,
                                ConstantArray::get(ATy, UsedArray), Name);
 
  GV->setSection("llvm.metadata");
}
 
Value *OpenMPIRBuilder::getGPUThreadID() {
  return Builder.CreateCall(
      getOrCreateRuntimeFunction(M,
                                 OMPRTL___kmpc_get_hardware_thread_id_in_block),
      {});
}
 
Value *OpenMPIRBuilder::getGPUWarpSize() {
  return Builder.CreateCall(
      getOrCreateRuntimeFunction(M, OMPRTL___kmpc_get_warp_size), {});
}
 
Value *OpenMPIRBuilder::getNVPTXWarpID() {
  unsigned LaneIDBits = Log2_32(Config.getGridValue().GV_Warp_Size);
  return Builder.CreateAShr(getGPUThreadID(), LaneIDBits, "nvptx_warp_id");
}
 
Value *OpenMPIRBuilder::getNVPTXLaneID() {
  unsigned LaneIDBits = Log2_32(Config.getGridValue().GV_Warp_Size);
  assert(LaneIDBits < 32 && "Invalid LaneIDBits size in NVPTX device.");
  unsigned LaneIDMask = ~0u >> (32u - LaneIDBits);
  return Builder.CreateAnd(getGPUThreadID(), Builder.getInt32(LaneIDMask),
                           "nvptx_lane_id");
}
 
Value *OpenMPIRBuilder::castValueToType(InsertPointTy AllocaIP, Value *From,
                                        Type *ToType) {
  Type *FromType = From->getType();
  uint64_t FromSize = M.getDataLayout().getTypeStoreSize(FromType);
  uint64_t ToSize = M.getDataLayout().getTypeStoreSize(ToType);
  assert(FromSize > 0 && "From size must be greater than zero");
  assert(ToSize > 0 && "To size must be greater than zero");
  if (FromType == ToType)
    return From;
  if (FromSize == ToSize)
    return Builder.CreateBitCast(From, ToType);
  if (ToType->isIntegerTy() && FromType->isIntegerTy())
    return Builder.CreateIntCast(From, ToType, /*isSigned*/ true);
  InsertPointTy SaveIP = Builder.saveIP();
  Builder.restoreIP(AllocaIP);
  Value *CastItem = Builder.CreateAlloca(ToType);
  Builder.restoreIP(SaveIP);
 
  Value *ValCastItem = Builder.CreatePointerBitCastOrAddrSpaceCast(
      CastItem, FromType->getPointerTo());
  Builder.CreateStore(From, ValCastItem);
  return Builder.CreateLoad(ToType, CastItem);
}
 
Value *OpenMPIRBuilder::createRuntimeShuffleFunction(InsertPointTy AllocaIP,
                                                     Value *Element,
                                                     Type *ElementType,
                                                     Value *Offset) {
  uint64_t Size = M.getDataLayout().getTypeStoreSize(ElementType);
  assert(Size <= 8 && "Unsupported bitwidth in shuffle instruction");
 
  // Cast all types to 32- or 64-bit values before calling shuffle routines.
  Type *CastTy = Builder.getIntNTy(Size <= 4 ? 32 : 64);
  Value *ElemCast = castValueToType(AllocaIP, Element, CastTy);
  Value *WarpSize =
      Builder.CreateIntCast(getGPUWarpSize(), Builder.getInt16Ty(), true);
  Function *ShuffleFunc = getOrCreateRuntimeFunctionPtr(
      Size <= 4 ? RuntimeFunction::OMPRTL___kmpc_shuffle_int32
                : RuntimeFunction::OMPRTL___kmpc_shuffle_int64);
  Value *WarpSizeCast =
      Builder.CreateIntCast(WarpSize, Builder.getInt16Ty(), /*isSigned=*/true);
  Value *ShuffleCall =
      Builder.CreateCall(ShuffleFunc, {ElemCast, Offset, WarpSizeCast});
  return castValueToType(AllocaIP, ShuffleCall, CastTy);
}
 
void OpenMPIRBuilder::shuffleAndStore(InsertPointTy AllocaIP, Value *SrcAddr,
                                      Value *DstAddr, Type *ElemType,
                                      Value *Offset, Type *ReductionArrayTy) {
  uint64_t Size = M.getDataLayout().getTypeStoreSize(ElemType);
  // Create the loop over the big sized data.
  // ptr = (void*)Elem;
  // ptrEnd = (void*) Elem + 1;
  // Step = 8;
  // while (ptr + Step < ptrEnd)
  //   shuffle((int64_t)*ptr);
  // Step = 4;
  // while (ptr + Step < ptrEnd)
  //   shuffle((int32_t)*ptr);
  // ...
  Type *IndexTy = Builder.getIndexTy(
      M.getDataLayout(), M.getDataLayout().getDefaultGlobalsAddressSpace());
  Value *ElemPtr = DstAddr;
  Value *Ptr = SrcAddr;
  for (unsigned IntSize = 8; IntSize >= 1; IntSize /= 2) {
    if (Size < IntSize)
      continue;
    Type *IntType = Builder.getIntNTy(IntSize * 8);
    Ptr = Builder.CreatePointerBitCastOrAddrSpaceCast(
        Ptr, IntType->getPointerTo(), Ptr->getName() + ".ascast");
    Value *SrcAddrGEP =
        Builder.CreateGEP(ElemType, SrcAddr, {ConstantInt::get(IndexTy, 1)});
    ElemPtr = Builder.CreatePointerBitCastOrAddrSpaceCast(
        ElemPtr, IntType->getPointerTo(), ElemPtr->getName() + ".ascast");
 
    Function *CurFunc = Builder.GetInsertBlock()->getParent();
    if ((Size / IntSize) > 1) {
      Value *PtrEnd = Builder.CreatePointerBitCastOrAddrSpaceCast(
          SrcAddrGEP, Builder.getPtrTy());
      BasicBlock *PreCondBB =
          BasicBlock::Create(M.getContext(), ".shuffle.pre_cond");
      BasicBlock *ThenBB = BasicBlock::Create(M.getContext(), ".shuffle.then");
      BasicBlock *ExitBB = BasicBlock::Create(M.getContext(), ".shuffle.exit");
      BasicBlock *CurrentBB = Builder.GetInsertBlock();
      emitBlock(PreCondBB, CurFunc);
      PHINode *PhiSrc =
          Builder.CreatePHI(Ptr->getType(), /*NumReservedValues=*/2);
      PhiSrc->addIncoming(Ptr, CurrentBB);
      PHINode *PhiDest =
          Builder.CreatePHI(ElemPtr->getType(), /*NumReservedValues=*/2);
      PhiDest->addIncoming(ElemPtr, CurrentBB);
      Ptr = PhiSrc;
      ElemPtr = PhiDest;
      Value *PtrDiff = Builder.CreatePtrDiff(
          Builder.getInt8Ty(), PtrEnd,
          Builder.CreatePointerBitCastOrAddrSpaceCast(Ptr, Builder.getPtrTy()));
      Builder.CreateCondBr(
          Builder.CreateICmpSGT(PtrDiff, Builder.getInt64(IntSize - 1)), ThenBB,
          ExitBB);
      emitBlock(ThenBB, CurFunc);
      Value *Res = createRuntimeShuffleFunction(
          AllocaIP,
          Builder.CreateAlignedLoad(
              IntType, Ptr, M.getDataLayout().getPrefTypeAlign(ElemType)),
          IntType, Offset);
      Builder.CreateAlignedStore(Res, ElemPtr,
                                 M.getDataLayout().getPrefTypeAlign(ElemType));
      Value *LocalPtr =
          Builder.CreateGEP(IntType, Ptr, {ConstantInt::get(IndexTy, 1)});
      Value *LocalElemPtr =
          Builder.CreateGEP(IntType, ElemPtr, {ConstantInt::get(IndexTy, 1)});
      PhiSrc->addIncoming(LocalPtr, ThenBB);
      PhiDest->addIncoming(LocalElemPtr, ThenBB);
      emitBranch(PreCondBB);
      emitBlock(ExitBB, CurFunc);
    } else {
      Value *Res = createRuntimeShuffleFunction(
          AllocaIP, Builder.CreateLoad(IntType, Ptr), IntType, Offset);
      if (ElemType->isIntegerTy() && ElemType->getScalarSizeInBits() <
                                         Res->getType()->getScalarSizeInBits())
        Res = Builder.CreateTrunc(Res, ElemType);
      Builder.CreateStore(Res, ElemPtr);
      Ptr = Builder.CreateGEP(IntType, Ptr, {ConstantInt::get(IndexTy, 1)});
      ElemPtr =
          Builder.CreateGEP(IntType, ElemPtr, {ConstantInt::get(IndexTy, 1)});
    }
    Size = Size % IntSize;
  }
}
 
void OpenMPIRBuilder::emitReductionListCopy(
    InsertPointTy AllocaIP, CopyAction Action, Type *ReductionArrayTy,
    ArrayRef<ReductionInfo> ReductionInfos, Value *SrcBase, Value *DestBase,
    CopyOptionsTy CopyOptions) {
  Type *IndexTy = Builder.getIndexTy(
      M.getDataLayout(), M.getDataLayout().getDefaultGlobalsAddressSpace());
  Value *RemoteLaneOffset = CopyOptions.RemoteLaneOffset;
 
  // Iterates, element-by-element, through the source Reduce list and
  // make a copy.
  for (auto En : enumerate(ReductionInfos)) {
    const ReductionInfo &RI = En.value();
    Value *SrcElementAddr = nullptr;
    Value *DestElementAddr = nullptr;
    Value *DestElementPtrAddr = nullptr;
    // Should we shuffle in an element from a remote lane?
    bool ShuffleInElement = false;
    // Set to true to update the pointer in the dest Reduce list to a
    // newly created element.
    bool UpdateDestListPtr = false;
 
    // Step 1.1: Get the address for the src element in the Reduce list.
    Value *SrcElementPtrAddr = Builder.CreateInBoundsGEP(
        ReductionArrayTy, SrcBase,
        {ConstantInt::get(IndexTy, 0), ConstantInt::get(IndexTy, En.index())});
    SrcElementAddr = Builder.CreateLoad(Builder.getPtrTy(), SrcElementPtrAddr);
 
    // Step 1.2: Create a temporary to store the element in the destination
    // Reduce list.
    DestElementPtrAddr = Builder.CreateInBoundsGEP(
        ReductionArrayTy, DestBase,
        {ConstantInt::get(IndexTy, 0), ConstantInt::get(IndexTy, En.index())});
    switch (Action) {
    case CopyAction::RemoteLaneToThread: {
      InsertPointTy CurIP = Builder.saveIP();
      Builder.restoreIP(AllocaIP);
      AllocaInst *DestAlloca = Builder.CreateAlloca(RI.ElementType, nullptr,
                                                    ".omp.reduction.element");
      DestAlloca->setAlignment(
          M.getDataLayout().getPrefTypeAlign(RI.ElementType));
      DestElementAddr = DestAlloca;
      DestElementAddr =
          Builder.CreateAddrSpaceCast(DestElementAddr, Builder.getPtrTy(),
                                      DestElementAddr->getName() + ".ascast");
      Builder.restoreIP(CurIP);
      ShuffleInElement = true;
      UpdateDestListPtr = true;
      break;
    }
    case CopyAction::ThreadCopy: {
      DestElementAddr =
          Builder.CreateLoad(Builder.getPtrTy(), DestElementPtrAddr);
      break;
    }
    }
 
    // Now that all active lanes have read the element in the
    // Reduce list, shuffle over the value from the remote lane.
    if (ShuffleInElement) {
      shuffleAndStore(AllocaIP, SrcElementAddr, DestElementAddr, RI.ElementType,
                      RemoteLaneOffset, ReductionArrayTy);
    } else {
      switch (RI.EvaluationKind) {
      case EvalKind::Scalar: {
        Value *Elem = Builder.CreateLoad(RI.ElementType, SrcElementAddr);
        // Store the source element value to the dest element address.
        Builder.CreateStore(Elem, DestElementAddr);
        break;
      }
      case EvalKind::Complex: {
        Value *SrcRealPtr = Builder.CreateConstInBoundsGEP2_32(
            RI.ElementType, SrcElementAddr, 0, 0, ".realp");
        Value *SrcReal = Builder.CreateLoad(
            RI.ElementType->getStructElementType(0), SrcRealPtr, ".real");
        Value *SrcImgPtr = Builder.CreateConstInBoundsGEP2_32(
            RI.ElementType, SrcElementAddr, 0, 1, ".imagp");
        Value *SrcImg = Builder.CreateLoad(
            RI.ElementType->getStructElementType(1), SrcImgPtr, ".imag");
 
        Value *DestRealPtr = Builder.CreateConstInBoundsGEP2_32(
            RI.ElementType, DestElementAddr, 0, 0, ".realp");
        Value *DestImgPtr = Builder.CreateConstInBoundsGEP2_32(
            RI.ElementType, DestElementAddr, 0, 1, ".imagp");
        Builder.CreateStore(SrcReal, DestRealPtr);
        Builder.CreateStore(SrcImg, DestImgPtr);
        break;
      }
      case EvalKind::Aggregate: {
        Value *SizeVal = Builder.getInt64(
            M.getDataLayout().getTypeStoreSize(RI.ElementType));
        Builder.CreateMemCpy(
            DestElementAddr, M.getDataLayout().getPrefTypeAlign(RI.ElementType),
            SrcElementAddr, M.getDataLayout().getPrefTypeAlign(RI.ElementType),
            SizeVal, false);
        break;
      }
      };
    }
 
    // Step 3.1: Modify reference in dest Reduce list as needed.
    // Modifying the reference in Reduce list to point to the newly
    // created element.  The element is live in the current function
    // scope and that of functions it invokes (i.e., reduce_function).
    // RemoteReduceData[i] = (void*)&RemoteElem
    if (UpdateDestListPtr) {
      Value *CastDestAddr = Builder.CreatePointerBitCastOrAddrSpaceCast(
          DestElementAddr, Builder.getPtrTy(),
          DestElementAddr->getName() + ".ascast");
      Builder.CreateStore(CastDestAddr, DestElementPtrAddr);
    }
  }
}
 
Function *OpenMPIRBuilder::emitInterWarpCopyFunction(
    const LocationDescription &Loc, ArrayRef<ReductionInfo> ReductionInfos,
    AttributeList FuncAttrs) {
  InsertPointTy SavedIP = Builder.saveIP();
  LLVMContext &Ctx = M.getContext();
  FunctionType *FuncTy = FunctionType::get(
      Builder.getVoidTy(), {Builder.getPtrTy(), Builder.getInt32Ty()},
      /* IsVarArg */ false);
  Function *WcFunc =
      Function::Create(FuncTy, GlobalVariable::InternalLinkage,
                       "_omp_reduction_inter_warp_copy_func", &M);
  WcFunc->setAttributes(FuncAttrs);
  WcFunc->addParamAttr(0, Attribute::NoUndef);
  WcFunc->addParamAttr(1, Attribute::NoUndef);
  BasicBlock *EntryBB = BasicBlock::Create(M.getContext(), "entry", WcFunc);
  Builder.SetInsertPoint(EntryBB);
 
  // ReduceList: thread local Reduce list.
  // At the stage of the computation when this function is called, partially
  // aggregated values reside in the first lane of every active warp.
  Argument *ReduceListArg = WcFunc->getArg(0);
  // NumWarps: number of warps active in the parallel region.  This could
  // be smaller than 32 (max warps in a CTA) for partial block reduction.
  Argument *NumWarpsArg = WcFunc->getArg(1);
 
  // This array is used as a medium to transfer, one reduce element at a time,
  // the data from the first lane of every warp to lanes in the first warp
  // in order to perform the final step of a reduction in a parallel region
  // (reduction across warps).  The array is placed in NVPTX __shared__ memory
  // for reduced latency, as well as to have a distinct copy for concurrently
  // executing target regions.  The array is declared with common linkage so
  // as to be shared across compilation units.
  StringRef TransferMediumName =
      "__openmp_nvptx_data_transfer_temporary_storage";
  GlobalVariable *TransferMedium = M.getGlobalVariable(TransferMediumName);
  unsigned WarpSize = Config.getGridValue().GV_Warp_Size;
  ArrayType *ArrayTy = ArrayType::get(Builder.getInt32Ty(), WarpSize);
  if (!TransferMedium) {
    TransferMedium = new GlobalVariable(
        M, ArrayTy, /*isConstant=*/false, GlobalVariable::WeakAnyLinkage,
        UndefValue::get(ArrayTy), TransferMediumName,
        /*InsertBefore=*/nullptr, GlobalVariable::NotThreadLocal,
        /*AddressSpace=*/3);
  }
 
  // Get the CUDA thread id of the current OpenMP thread on the GPU.
  Value *GPUThreadID = getGPUThreadID();
  // nvptx_lane_id = nvptx_id % warpsize
  Value *LaneID = getNVPTXLaneID();
  // nvptx_warp_id = nvptx_id / warpsize
  Value *WarpID = getNVPTXWarpID();
 
  InsertPointTy AllocaIP =
      InsertPointTy(Builder.GetInsertBlock(),
                    Builder.GetInsertBlock()->getFirstInsertionPt());
  Type *Arg0Type = ReduceListArg->getType();
  Type *Arg1Type = NumWarpsArg->getType();
  Builder.restoreIP(AllocaIP);
  AllocaInst *ReduceListAlloca = Builder.CreateAlloca(
      Arg0Type, nullptr, ReduceListArg->getName() + ".addr");
  AllocaInst *NumWarpsAlloca =
      Builder.CreateAlloca(Arg1Type, nullptr, NumWarpsArg->getName() + ".addr");
  Value *ReduceListAddrCast = Builder.CreatePointerBitCastOrAddrSpaceCast(
      ReduceListAlloca, Arg0Type, ReduceListAlloca->getName() + ".ascast");
  Value *NumWarpsAddrCast = Builder.CreatePointerBitCastOrAddrSpaceCast(
      NumWarpsAlloca, Arg1Type->getPointerTo(),
      NumWarpsAlloca->getName() + ".ascast");
  Builder.CreateStore(ReduceListArg, ReduceListAddrCast);
  Builder.CreateStore(NumWarpsArg, NumWarpsAddrCast);
  AllocaIP = getInsertPointAfterInstr(NumWarpsAlloca);
  InsertPointTy CodeGenIP =
      getInsertPointAfterInstr(&Builder.GetInsertBlock()->back());
  Builder.restoreIP(CodeGenIP);
 
  Value *ReduceList =
      Builder.CreateLoad(Builder.getPtrTy(), ReduceListAddrCast);
 
  for (auto En : enumerate(ReductionInfos)) {
    //
    // Warp master copies reduce element to transfer medium in __shared__
    // memory.
    //
    const ReductionInfo &RI = En.value();
    unsigned RealTySize = M.getDataLayout().getTypeAllocSize(RI.ElementType);
    for (unsigned TySize = 4; TySize > 0 && RealTySize > 0; TySize /= 2) {
      Type *CType = Builder.getIntNTy(TySize * 8);
 
      unsigned NumIters = RealTySize / TySize;
      if (NumIters == 0)
        continue;
      Value *Cnt = nullptr;
      Value *CntAddr = nullptr;
      BasicBlock *PrecondBB = nullptr;
      BasicBlock *ExitBB = nullptr;
      if (NumIters > 1) {
        CodeGenIP = Builder.saveIP();
        Builder.restoreIP(AllocaIP);
        CntAddr =
            Builder.CreateAlloca(Builder.getInt32Ty(), nullptr, ".cnt.addr");
 
        CntAddr = Builder.CreateAddrSpaceCast(CntAddr, Builder.getPtrTy(),
                                              CntAddr->getName() + ".ascast");
        Builder.restoreIP(CodeGenIP);
        Builder.CreateStore(Constant::getNullValue(Builder.getInt32Ty()),
                            CntAddr,
                            /*Volatile=*/false);
        PrecondBB = BasicBlock::Create(Ctx, "precond");
        ExitBB = BasicBlock::Create(Ctx, "exit");
        BasicBlock *BodyBB = BasicBlock::Create(Ctx, "body");
        emitBlock(PrecondBB, Builder.GetInsertBlock()->getParent());
        Cnt = Builder.CreateLoad(Builder.getInt32Ty(), CntAddr,
                                 /*Volatile=*/false);
        Value *Cmp = Builder.CreateICmpULT(
            Cnt, ConstantInt::get(Builder.getInt32Ty(), NumIters));
        Builder.CreateCondBr(Cmp, BodyBB, ExitBB);
        emitBlock(BodyBB, Builder.GetInsertBlock()->getParent());
      }
 
      // kmpc_barrier.
      createBarrier(LocationDescription(Builder.saveIP(), Loc.DL),
                    omp::Directive::OMPD_unknown,
                    /* ForceSimpleCall */ false,
                    /* CheckCancelFlag */ true);
      BasicBlock *ThenBB = BasicBlock::Create(Ctx, "then");
      BasicBlock *ElseBB = BasicBlock::Create(Ctx, "else");
      BasicBlock *MergeBB = BasicBlock::Create(Ctx, "ifcont");
 
      // if (lane_id  == 0)
      Value *IsWarpMaster = Builder.CreateIsNull(LaneID, "warp_master");
      Builder.CreateCondBr(IsWarpMaster, ThenBB, ElseBB);
      emitBlock(ThenBB, Builder.GetInsertBlock()->getParent());
 
      // Reduce element = LocalReduceList[i]
      auto *RedListArrayTy =
          ArrayType::get(Builder.getPtrTy(), ReductionInfos.size());
      Type *IndexTy = Builder.getIndexTy(
          M.getDataLayout(), M.getDataLayout().getDefaultGlobalsAddressSpace());
      Value *ElemPtrPtr =
          Builder.CreateInBoundsGEP(RedListArrayTy, ReduceList,
                                    {ConstantInt::get(IndexTy, 0),
                                     ConstantInt::get(IndexTy, En.index())});
      // elemptr = ((CopyType*)(elemptrptr)) + I
      Value *ElemPtr = Builder.CreateLoad(Builder.getPtrTy(), ElemPtrPtr);
      if (NumIters > 1)
        ElemPtr = Builder.CreateGEP(Builder.getInt32Ty(), ElemPtr, Cnt);
 
      // Get pointer to location in transfer medium.
      // MediumPtr = &medium[warp_id]
      Value *MediumPtr = Builder.CreateInBoundsGEP(
          ArrayTy, TransferMedium, {Builder.getInt64(0), WarpID});
      // elem = *elemptr
      //*MediumPtr = elem
      Value *Elem = Builder.CreateLoad(CType, ElemPtr);
      // Store the source element value to the dest element address.
      Builder.CreateStore(Elem, MediumPtr,
                          /*IsVolatile*/ true);
      Builder.CreateBr(MergeBB);
 
      // else
      emitBlock(ElseBB, Builder.GetInsertBlock()->getParent());
      Builder.CreateBr(MergeBB);
 
      // endif
      emitBlock(MergeBB, Builder.GetInsertBlock()->getParent());
      createBarrier(LocationDescription(Builder.saveIP(), Loc.DL),
                    omp::Directive::OMPD_unknown,
                    /* ForceSimpleCall */ false,
                    /* CheckCancelFlag */ true);
 
      // Warp 0 copies reduce element from transfer medium
      BasicBlock *W0ThenBB = BasicBlock::Create(Ctx, "then");
      BasicBlock *W0ElseBB = BasicBlock::Create(Ctx, "else");
      BasicBlock *W0MergeBB = BasicBlock::Create(Ctx, "ifcont");
 
      Value *NumWarpsVal =
          Builder.CreateLoad(Builder.getInt32Ty(), NumWarpsAddrCast);
      // Up to 32 threads in warp 0 are active.
      Value *IsActiveThread =
          Builder.CreateICmpULT(GPUThreadID, NumWarpsVal, "is_active_thread");
      Builder.CreateCondBr(IsActiveThread, W0ThenBB, W0ElseBB);
 
      emitBlock(W0ThenBB, Builder.GetInsertBlock()->getParent());
 
      // SecMediumPtr = &medium[tid]
      // SrcMediumVal = *SrcMediumPtr
      Value *SrcMediumPtrVal = Builder.CreateInBoundsGEP(
          ArrayTy, TransferMedium, {Builder.getInt64(0), GPUThreadID});
      // TargetElemPtr = (CopyType*)(SrcDataAddr[i]) + I
      Value *TargetElemPtrPtr =
          Builder.CreateInBoundsGEP(RedListArrayTy, ReduceList,
                                    {ConstantInt::get(IndexTy, 0),
                                     ConstantInt::get(IndexTy, En.index())});
      Value *TargetElemPtrVal =
          Builder.CreateLoad(Builder.getPtrTy(), TargetElemPtrPtr);
      Value *TargetElemPtr = TargetElemPtrVal;
      if (NumIters > 1)
        TargetElemPtr =
            Builder.CreateGEP(Builder.getInt32Ty(), TargetElemPtr, Cnt);
 
      // *TargetElemPtr = SrcMediumVal;
      Value *SrcMediumValue =
          Builder.CreateLoad(CType, SrcMediumPtrVal, /*IsVolatile*/ true);
      Builder.CreateStore(SrcMediumValue, TargetElemPtr);
      Builder.CreateBr(W0MergeBB);
 
      emitBlock(W0ElseBB, Builder.GetInsertBlock()->getParent());
      Builder.CreateBr(W0MergeBB);
 
      emitBlock(W0MergeBB, Builder.GetInsertBlock()->getParent());
 
      if (NumIters > 1) {
        Cnt = Builder.CreateNSWAdd(
            Cnt, ConstantInt::get(Builder.getInt32Ty(), /*V=*/1));
        Builder.CreateStore(Cnt, CntAddr, /*Volatile=*/false);
 
        auto *CurFn = Builder.GetInsertBlock()->getParent();
        emitBranch(PrecondBB);
        emitBlock(ExitBB, CurFn);
      }
      RealTySize %= TySize;
    }
  }
 
  Builder.CreateRetVoid();
  Builder.restoreIP(SavedIP);
 
  return WcFunc;
}
 
Function *OpenMPIRBuilder::emitShuffleAndReduceFunction(
    ArrayRef<ReductionInfo> ReductionInfos, Function *ReduceFn,
    AttributeList FuncAttrs) {
  LLVMContext &Ctx = M.getContext();
  FunctionType *FuncTy =
      FunctionType::get(Builder.getVoidTy(),
                        {Builder.getPtrTy(), Builder.getInt16Ty(),
                         Builder.getInt16Ty(), Builder.getInt16Ty()},
                        /* IsVarArg */ false);
  Function *SarFunc =
      Function::Create(FuncTy, GlobalVariable::InternalLinkage,
                       "_omp_reduction_shuffle_and_reduce_func", &M);
  SarFunc->setAttributes(FuncAttrs);
  SarFunc->addParamAttr(0, Attribute::NoUndef);
  SarFunc->addParamAttr(1, Attribute::NoUndef);
  SarFunc->addParamAttr(2, Attribute::NoUndef);
  SarFunc->addParamAttr(3, Attribute::NoUndef);
  SarFunc->addParamAttr(1, Attribute::SExt);
  SarFunc->addParamAttr(2, Attribute::SExt);
  SarFunc->addParamAttr(3, Attribute::SExt);
  BasicBlock *EntryBB = BasicBlock::Create(M.getContext(), "entry", SarFunc);
  Builder.SetInsertPoint(EntryBB);
 
  // Thread local Reduce list used to host the values of data to be reduced.
  Argument *ReduceListArg = SarFunc->getArg(0);
  // Current lane id; could be logical.
  Argument *LaneIDArg = SarFunc->getArg(1);
  // Offset of the remote source lane relative to the current lane.
  Argument *RemoteLaneOffsetArg = SarFunc->getArg(2);
  // Algorithm version.  This is expected to be known at compile time.
  Argument *AlgoVerArg = SarFunc->getArg(3);
 
  Type *ReduceListArgType = ReduceListArg->getType();
  Type *LaneIDArgType = LaneIDArg->getType();
  Type *LaneIDArgPtrType = LaneIDArg->getType()->getPointerTo();
  Value *ReduceListAlloca = Builder.CreateAlloca(
      ReduceListArgType, nullptr, ReduceListArg->getName() + ".addr");
  Value *LaneIdAlloca = Builder.CreateAlloca(LaneIDArgType, nullptr,
                                             LaneIDArg->getName() + ".addr");
  Value *RemoteLaneOffsetAlloca = Builder.CreateAlloca(
      LaneIDArgType, nullptr, RemoteLaneOffsetArg->getName() + ".addr");
  Value *AlgoVerAlloca = Builder.CreateAlloca(LaneIDArgType, nullptr,
                                              AlgoVerArg->getName() + ".addr");
  ArrayType *RedListArrayTy =
      ArrayType::get(Builder.getPtrTy(), ReductionInfos.size());
 
  // Create a local thread-private variable to host the Reduce list
  // from a remote lane.
  Instruction *RemoteReductionListAlloca = Builder.CreateAlloca(
      RedListArrayTy, nullptr, ".omp.reduction.remote_reduce_list");
 
  Value *ReduceListAddrCast = Builder.CreatePointerBitCastOrAddrSpaceCast(
      ReduceListAlloca, ReduceListArgType,
      ReduceListAlloca->getName() + ".ascast");
  Value *LaneIdAddrCast = Builder.CreatePointerBitCastOrAddrSpaceCast(
      LaneIdAlloca, LaneIDArgPtrType, LaneIdAlloca->getName() + ".ascast");
  Value *RemoteLaneOffsetAddrCast = Builder.CreatePointerBitCastOrAddrSpaceCast(
      RemoteLaneOffsetAlloca, LaneIDArgPtrType,
      RemoteLaneOffsetAlloca->getName() + ".ascast");
  Value *AlgoVerAddrCast = Builder.CreatePointerBitCastOrAddrSpaceCast(
      AlgoVerAlloca, LaneIDArgPtrType, AlgoVerAlloca->getName() + ".ascast");
  Value *RemoteListAddrCast = Builder.CreatePointerBitCastOrAddrSpaceCast(
      RemoteReductionListAlloca, Builder.getPtrTy(),
      RemoteReductionListAlloca->getName() + ".ascast");
 
  Builder.CreateStore(ReduceListArg, ReduceListAddrCast);
  Builder.CreateStore(LaneIDArg, LaneIdAddrCast);
  Builder.CreateStore(RemoteLaneOffsetArg, RemoteLaneOffsetAddrCast);
  Builder.CreateStore(AlgoVerArg, AlgoVerAddrCast);
 
  Value *ReduceList = Builder.CreateLoad(ReduceListArgType, ReduceListAddrCast);
  Value *LaneId = Builder.CreateLoad(LaneIDArgType, LaneIdAddrCast);
  Value *RemoteLaneOffset =
      Builder.CreateLoad(LaneIDArgType, RemoteLaneOffsetAddrCast);
  Value *AlgoVer = Builder.CreateLoad(LaneIDArgType, AlgoVerAddrCast);
 
  InsertPointTy AllocaIP = getInsertPointAfterInstr(RemoteReductionListAlloca);
 
  // This loop iterates through the list of reduce elements and copies,
  // element by element, from a remote lane in the warp to RemoteReduceList,
  // hosted on the thread's stack.
  emitReductionListCopy(
      AllocaIP, CopyAction::RemoteLaneToThread, RedListArrayTy, ReductionInfos,
      ReduceList, RemoteListAddrCast, {RemoteLaneOffset, nullptr, nullptr});
 
  // The actions to be performed on the Remote Reduce list is dependent
  // on the algorithm version.
  //
  //  if (AlgoVer==0) || (AlgoVer==1 && (LaneId < Offset)) || (AlgoVer==2 &&
  //  LaneId % 2 == 0 && Offset > 0):
  //    do the reduction value aggregation
  //
  //  The thread local variable Reduce list is mutated in place to host the
  //  reduced data, which is the aggregated value produced from local and
  //  remote lanes.
  //
  //  Note that AlgoVer is expected to be a constant integer known at compile
  //  time.
  //  When AlgoVer==0, the first conjunction evaluates to true, making
  //    the entire predicate true during compile time.
  //  When AlgoVer==1, the second conjunction has only the second part to be
  //    evaluated during runtime.  Other conjunctions evaluates to false
  //    during compile time.
  //  When AlgoVer==2, the third conjunction has only the second part to be
  //    evaluated during runtime.  Other conjunctions evaluates to false
  //    during compile time.
  Value *CondAlgo0 = Builder.CreateIsNull(AlgoVer);
  Value *Algo1 = Builder.CreateICmpEQ(AlgoVer, Builder.getInt16(1));
  Value *LaneComp = Builder.CreateICmpULT(LaneId, RemoteLaneOffset);
  Value *CondAlgo1 = Builder.CreateAnd(Algo1, LaneComp);
  Value *Algo2 = Builder.CreateICmpEQ(AlgoVer, Builder.getInt16(2));
  Value *LaneIdAnd1 = Builder.CreateAnd(LaneId, Builder.getInt16(1));
  Value *LaneIdComp = Builder.CreateIsNull(LaneIdAnd1);
  Value *Algo2AndLaneIdComp = Builder.CreateAnd(Algo2, LaneIdComp);
  Value *RemoteOffsetComp =
      Builder.CreateICmpSGT(RemoteLaneOffset, Builder.getInt16(0));
  Value *CondAlgo2 = Builder.CreateAnd(Algo2AndLaneIdComp, RemoteOffsetComp);
  Value *CA0OrCA1 = Builder.CreateOr(CondAlgo0, CondAlgo1);
  Value *CondReduce = Builder.CreateOr(CA0OrCA1, CondAlgo2);
 
  BasicBlock *ThenBB = BasicBlock::Create(Ctx, "then");
  BasicBlock *ElseBB = BasicBlock::Create(Ctx, "else");
  BasicBlock *MergeBB = BasicBlock::Create(Ctx, "ifcont");
 
  Builder.CreateCondBr(CondReduce, ThenBB, ElseBB);
  emitBlock(ThenBB, Builder.GetInsertBlock()->getParent());
  Value *LocalReduceListPtr = Builder.CreatePointerBitCastOrAddrSpaceCast(
      ReduceList, Builder.getPtrTy());
  Value *RemoteReduceListPtr = Builder.CreatePointerBitCastOrAddrSpaceCast(
      RemoteListAddrCast, Builder.getPtrTy());
  Builder.CreateCall(ReduceFn, {LocalReduceListPtr, RemoteReduceListPtr})
      ->addFnAttr(Attribute::NoUnwind);
  Builder.CreateBr(MergeBB);
 
  emitBlock(ElseBB, Builder.GetInsertBlock()->getParent());
  Builder.CreateBr(MergeBB);
 
  emitBlock(MergeBB, Builder.GetInsertBlock()->getParent());
 
  // if (AlgoVer==1 && (LaneId >= Offset)) copy Remote Reduce list to local
  // Reduce list.
  Algo1 = Builder.CreateICmpEQ(AlgoVer, Builder.getInt16(1));
  Value *LaneIdGtOffset = Builder.CreateICmpUGE(LaneId, RemoteLaneOffset);
  Value *CondCopy = Builder.CreateAnd(Algo1, LaneIdGtOffset);
 
  BasicBlock *CpyThenBB = BasicBlock::Create(Ctx, "then");
  BasicBlock *CpyElseBB = BasicBlock::Create(Ctx, "else");
  BasicBlock *CpyMergeBB = BasicBlock::Create(Ctx, "ifcont");
  Builder.CreateCondBr(CondCopy, CpyThenBB, CpyElseBB);
 
  emitBlock(CpyThenBB, Builder.GetInsertBlock()->getParent());
  emitReductionListCopy(AllocaIP, CopyAction::ThreadCopy, RedListArrayTy,
                        ReductionInfos, RemoteListAddrCast, ReduceList);
  Builder.CreateBr(CpyMergeBB);
 
  emitBlock(CpyElseBB, Builder.GetInsertBlock()->getParent());
  Builder.CreateBr(CpyMergeBB);
 
  emitBlock(CpyMergeBB, Builder.GetInsertBlock()->getParent());
 
  Builder.CreateRetVoid();
 
  return SarFunc;
}
 
Function *OpenMPIRBuilder::emitListToGlobalCopyFunction(
    ArrayRef<ReductionInfo> ReductionInfos, Type *ReductionsBufferTy,
    AttributeList FuncAttrs) {
  OpenMPIRBuilder::InsertPointTy OldIP = Builder.saveIP();
  LLVMContext &Ctx = M.getContext();
  FunctionType *FuncTy = FunctionType::get(
      Builder.getVoidTy(),
      {Builder.getPtrTy(), Builder.getInt32Ty(), Builder.getPtrTy()},
      /* IsVarArg */ false);
  Function *LtGCFunc =
      Function::Create(FuncTy, GlobalVariable::InternalLinkage,
                       "_omp_reduction_list_to_global_copy_func", &M);
  LtGCFunc->setAttributes(FuncAttrs);
  LtGCFunc->addParamAttr(0, Attribute::NoUndef);
  LtGCFunc->addParamAttr(1, Attribute::NoUndef);
  LtGCFunc->addParamAttr(2, Attribute::NoUndef);
 
  BasicBlock *EntryBlock = BasicBlock::Create(Ctx, "entry", LtGCFunc);
  Builder.SetInsertPoint(EntryBlock);
 
  // Buffer: global reduction buffer.
  Argument *BufferArg = LtGCFunc->getArg(0);
  // Idx: index of the buffer.
  Argument *IdxArg = LtGCFunc->getArg(1);
  // ReduceList: thread local Reduce list.
  Argument *ReduceListArg = LtGCFunc->getArg(2);
 
  Value *BufferArgAlloca = Builder.CreateAlloca(Builder.getPtrTy(), nullptr,
                                                BufferArg->getName() + ".addr");
  Value *IdxArgAlloca = Builder.CreateAlloca(Builder.getInt32Ty(), nullptr,
                                             IdxArg->getName() + ".addr");
  Value *ReduceListArgAlloca = Builder.CreateAlloca(
      Builder.getPtrTy(), nullptr, ReduceListArg->getName() + ".addr");
  Value *BufferArgAddrCast = Builder.CreatePointerBitCastOrAddrSpaceCast(
      BufferArgAlloca, Builder.getPtrTy(),
      BufferArgAlloca->getName() + ".ascast");
  Value *IdxArgAddrCast = Builder.CreatePointerBitCastOrAddrSpaceCast(
      IdxArgAlloca, Builder.getPtrTy(), IdxArgAlloca->getName() + ".ascast");
  Value *ReduceListArgAddrCast = Builder.CreatePointerBitCastOrAddrSpaceCast(
      ReduceListArgAlloca, Builder.getPtrTy(),
      ReduceListArgAlloca->getName() + ".ascast");
 
  Builder.CreateStore(BufferArg, BufferArgAddrCast);
  Builder.CreateStore(IdxArg, IdxArgAddrCast);
  Builder.CreateStore(ReduceListArg, ReduceListArgAddrCast);
 
  Value *LocalReduceList =
      Builder.CreateLoad(Builder.getPtrTy(), ReduceListArgAddrCast);
  Value *BufferArgVal =
      Builder.CreateLoad(Builder.getPtrTy(), BufferArgAddrCast);
  Value *Idxs[] = {Builder.CreateLoad(Builder.getInt32Ty(), IdxArgAddrCast)};
  Type *IndexTy = Builder.getIndexTy(
      M.getDataLayout(), M.getDataLayout().getDefaultGlobalsAddressSpace());
  for (auto En : enumerate(ReductionInfos)) {
    const ReductionInfo &RI = En.value();
    auto *RedListArrayTy =
        ArrayType::get(Builder.getPtrTy(), ReductionInfos.size());
    // Reduce element = LocalReduceList[i]
    Value *ElemPtrPtr = Builder.CreateInBoundsGEP(
        RedListArrayTy, LocalReduceList,
        {ConstantInt::get(IndexTy, 0), ConstantInt::get(IndexTy, En.index())});
    // elemptr = ((CopyType*)(elemptrptr)) + I
    Value *ElemPtr = Builder.CreateLoad(Builder.getPtrTy(), ElemPtrPtr);
 
    // Global = Buffer.VD[Idx];
    Value *BufferVD =
        Builder.CreateInBoundsGEP(ReductionsBufferTy, BufferArgVal, Idxs);
    Value *GlobVal = Builder.CreateConstInBoundsGEP2_32(
        ReductionsBufferTy, BufferVD, 0, En.index());
 
    switch (RI.EvaluationKind) {
    case EvalKind::Scalar: {
      Value *TargetElement = Builder.CreateLoad(RI.ElementType, ElemPtr);
      Builder.CreateStore(TargetElement, GlobVal);
      break;
    }
    case EvalKind::Complex: {
      Value *SrcRealPtr = Builder.CreateConstInBoundsGEP2_32(
          RI.ElementType, ElemPtr, 0, 0, ".realp");
      Value *SrcReal = Builder.CreateLoad(
          RI.ElementType->getStructElementType(0), SrcRealPtr, ".real");
      Value *SrcImgPtr = Builder.CreateConstInBoundsGEP2_32(
          RI.ElementType, ElemPtr, 0, 1, ".imagp");
      Value *SrcImg = Builder.CreateLoad(
          RI.ElementType->getStructElementType(1), SrcImgPtr, ".imag");
 
      Value *DestRealPtr = Builder.CreateConstInBoundsGEP2_32(
          RI.ElementType, GlobVal, 0, 0, ".realp");
      Value *DestImgPtr = Builder.CreateConstInBoundsGEP2_32(
          RI.ElementType, GlobVal, 0, 1, ".imagp");
      Builder.CreateStore(SrcReal, DestRealPtr);
      Builder.CreateStore(SrcImg, DestImgPtr);
      break;
    }
    case EvalKind::Aggregate: {
      Value *SizeVal =
          Builder.getInt64(M.getDataLayout().getTypeStoreSize(RI.ElementType));
      Builder.CreateMemCpy(
          GlobVal, M.getDataLayout().getPrefTypeAlign(RI.ElementType), ElemPtr,
          M.getDataLayout().getPrefTypeAlign(RI.ElementType), SizeVal, false);
      break;
    }
    }
  }
 
  Builder.CreateRetVoid();
  Builder.restoreIP(OldIP);
  return LtGCFunc;
}
 
Function *OpenMPIRBuilder::emitListToGlobalReduceFunction(
    ArrayRef<ReductionInfo> ReductionInfos, Function *ReduceFn,
    Type *ReductionsBufferTy, AttributeList FuncAttrs) {
  OpenMPIRBuilder::InsertPointTy OldIP = Builder.saveIP();
  LLVMContext &Ctx = M.getContext();
  FunctionType *FuncTy = FunctionType::get(
      Builder.getVoidTy(),
      {Builder.getPtrTy(), Builder.getInt32Ty(), Builder.getPtrTy()},
      /* IsVarArg */ false);
  Function *LtGRFunc =
      Function::Create(FuncTy, GlobalVariable::InternalLinkage,
                       "_omp_reduction_list_to_global_reduce_func", &M);
  LtGRFunc->setAttributes(FuncAttrs);
  LtGRFunc->addParamAttr(0, Attribute::NoUndef);
  LtGRFunc->addParamAttr(1, Attribute::NoUndef);
  LtGRFunc->addParamAttr(2, Attribute::NoUndef);
 
  BasicBlock *EntryBlock = BasicBlock::Create(Ctx, "entry", LtGRFunc);
  Builder.SetInsertPoint(EntryBlock);
 
  // Buffer: global reduction buffer.
  Argument *BufferArg = LtGRFunc->getArg(0);
  // Idx: index of the buffer.
  Argument *IdxArg = LtGRFunc->getArg(1);
  // ReduceList: thread local Reduce list.
  Argument *ReduceListArg = LtGRFunc->getArg(2);
 
  Value *BufferArgAlloca = Builder.CreateAlloca(Builder.getPtrTy(), nullptr,
                                                BufferArg->getName() + ".addr");
  Value *IdxArgAlloca = Builder.CreateAlloca(Builder.getInt32Ty(), nullptr,
                                             IdxArg->getName() + ".addr");
  Value *ReduceListArgAlloca = Builder.CreateAlloca(
      Builder.getPtrTy(), nullptr, ReduceListArg->getName() + ".addr");
  auto *RedListArrayTy =
      ArrayType::get(Builder.getPtrTy(), ReductionInfos.size());
 
  // 1. Build a list of reduction variables.
  // void *RedList[<n>] = {<ReductionVars>[0], ..., <ReductionVars>[<n>-1]};
  Value *LocalReduceList =
      Builder.CreateAlloca(RedListArrayTy, nullptr, ".omp.reduction.red_list");
 
  Value *BufferArgAddrCast = Builder.CreatePointerBitCastOrAddrSpaceCast(
      BufferArgAlloca, Builder.getPtrTy(),
      BufferArgAlloca->getName() + ".ascast");
  Value *IdxArgAddrCast = Builder.CreatePointerBitCastOrAddrSpaceCast(
      IdxArgAlloca, Builder.getPtrTy(), IdxArgAlloca->getName() + ".ascast");
  Value *ReduceListArgAddrCast = Builder.CreatePointerBitCastOrAddrSpaceCast(
      ReduceListArgAlloca, Builder.getPtrTy(),
      ReduceListArgAlloca->getName() + ".ascast");
  Value *LocalReduceListAddrCast = Builder.CreatePointerBitCastOrAddrSpaceCast(
      LocalReduceList, Builder.getPtrTy(),
      LocalReduceList->getName() + ".ascast");
 
  Builder.CreateStore(BufferArg, BufferArgAddrCast);
  Builder.CreateStore(IdxArg, IdxArgAddrCast);
  Builder.CreateStore(ReduceListArg, ReduceListArgAddrCast);
 
  Value *BufferVal = Builder.CreateLoad(Builder.getPtrTy(), BufferArgAddrCast);
  Value *Idxs[] = {Builder.CreateLoad(Builder.getInt32Ty(), IdxArgAddrCast)};
  Type *IndexTy = Builder.getIndexTy(
      M.getDataLayout(), M.getDataLayout().getDefaultGlobalsAddressSpace());
  for (auto En : enumerate(ReductionInfos)) {
    Value *TargetElementPtrPtr = Builder.CreateInBoundsGEP(
        RedListArrayTy, LocalReduceListAddrCast,
        {ConstantInt::get(IndexTy, 0), ConstantInt::get(IndexTy, En.index())});
    Value *BufferVD =
        Builder.CreateInBoundsGEP(ReductionsBufferTy, BufferVal, Idxs);
    // Global = Buffer.VD[Idx];
    Value *GlobValPtr = Builder.CreateConstInBoundsGEP2_32(
        ReductionsBufferTy, BufferVD, 0, En.index());
    Builder.CreateStore(GlobValPtr, TargetElementPtrPtr);
  }
 
  // Call reduce_function(GlobalReduceList, ReduceList)
  Value *ReduceList =
      Builder.CreateLoad(Builder.getPtrTy(), ReduceListArgAddrCast);
  Builder.CreateCall(ReduceFn, {LocalReduceListAddrCast, ReduceList})
      ->addFnAttr(Attribute::NoUnwind);
  Builder.CreateRetVoid();
  Builder.restoreIP(OldIP);
  return LtGRFunc;
}
 
Function *OpenMPIRBuilder::emitGlobalToListCopyFunction(
    ArrayRef<ReductionInfo> ReductionInfos, Type *ReductionsBufferTy,
    AttributeList FuncAttrs) {
  OpenMPIRBuilder::InsertPointTy OldIP = Builder.saveIP();
  LLVMContext &Ctx = M.getContext();
  FunctionType *FuncTy = FunctionType::get(
      Builder.getVoidTy(),
      {Builder.getPtrTy(), Builder.getInt32Ty(), Builder.getPtrTy()},
      /* IsVarArg */ false);
  Function *LtGCFunc =
      Function::Create(FuncTy, GlobalVariable::InternalLinkage,
                       "_omp_reduction_global_to_list_copy_func", &M);
  LtGCFunc->setAttributes(FuncAttrs);
  LtGCFunc->addParamAttr(0, Attribute::NoUndef);
  LtGCFunc->addParamAttr(1, Attribute::NoUndef);
  LtGCFunc->addParamAttr(2, Attribute::NoUndef);
 
  BasicBlock *EntryBlock = BasicBlock::Create(Ctx, "entry", LtGCFunc);
  Builder.SetInsertPoint(EntryBlock);
 
  // Buffer: global reduction buffer.
  Argument *BufferArg = LtGCFunc->getArg(0);
  // Idx: index of the buffer.
  Argument *IdxArg = LtGCFunc->getArg(1);
  // ReduceList: thread local Reduce list.
  Argument *ReduceListArg = LtGCFunc->getArg(2);
 
  Value *BufferArgAlloca = Builder.CreateAlloca(Builder.getPtrTy(), nullptr,
                                                BufferArg->getName() + ".addr");
  Value *IdxArgAlloca = Builder.CreateAlloca(Builder.getInt32Ty(), nullptr,
                                             IdxArg->getName() + ".addr");
  Value *ReduceListArgAlloca = Builder.CreateAlloca(
      Builder.getPtrTy(), nullptr, ReduceListArg->getName() + ".addr");
  Value *BufferArgAddrCast = Builder.CreatePointerBitCastOrAddrSpaceCast(
      BufferArgAlloca, Builder.getPtrTy(),
      BufferArgAlloca->getName() + ".ascast");
  Value *IdxArgAddrCast = Builder.CreatePointerBitCastOrAddrSpaceCast(
      IdxArgAlloca, Builder.getPtrTy(), IdxArgAlloca->getName() + ".ascast");
  Value *ReduceListArgAddrCast = Builder.CreatePointerBitCastOrAddrSpaceCast(
      ReduceListArgAlloca, Builder.getPtrTy(),
      ReduceListArgAlloca->getName() + ".ascast");
  Builder.CreateStore(BufferArg, BufferArgAddrCast);
  Builder.CreateStore(IdxArg, IdxArgAddrCast);
  Builder.CreateStore(ReduceListArg, ReduceListArgAddrCast);
 
  Value *LocalReduceList =
      Builder.CreateLoad(Builder.getPtrTy(), ReduceListArgAddrCast);
  Value *BufferVal = Builder.CreateLoad(Builder.getPtrTy(), BufferArgAddrCast);
  Value *Idxs[] = {Builder.CreateLoad(Builder.getInt32Ty(), IdxArgAddrCast)};
  Type *IndexTy = Builder.getIndexTy(
      M.getDataLayout(), M.getDataLayout().getDefaultGlobalsAddressSpace());
  for (auto En : enumerate(ReductionInfos)) {
    const OpenMPIRBuilder::ReductionInfo &RI = En.value();
    auto *RedListArrayTy =
        ArrayType::get(Builder.getPtrTy(), ReductionInfos.size());
    // Reduce element = LocalReduceList[i]
    Value *ElemPtrPtr = Builder.CreateInBoundsGEP(
        RedListArrayTy, LocalReduceList,
        {ConstantInt::get(IndexTy, 0), ConstantInt::get(IndexTy, En.index())});
    // elemptr = ((CopyType*)(elemptrptr)) + I
    Value *ElemPtr = Builder.CreateLoad(Builder.getPtrTy(), ElemPtrPtr);
    // Global = Buffer.VD[Idx];
    Value *BufferVD =
        Builder.CreateInBoundsGEP(ReductionsBufferTy, BufferVal, Idxs);
    Value *GlobValPtr = Builder.CreateConstInBoundsGEP2_32(
        ReductionsBufferTy, BufferVD, 0, En.index());
 
    switch (RI.EvaluationKind) {
    case EvalKind::Scalar: {
      Value *TargetElement = Builder.CreateLoad(RI.ElementType, GlobValPtr);
      Builder.CreateStore(TargetElement, ElemPtr);
      break;
    }
    case EvalKind::Complex: {
      Value *SrcRealPtr = Builder.CreateConstInBoundsGEP2_32(
          RI.ElementType, GlobValPtr, 0, 0, ".realp");
      Value *SrcReal = Builder.CreateLoad(
          RI.ElementType->getStructElementType(0), SrcRealPtr, ".real");
      Value *SrcImgPtr = Builder.CreateConstInBoundsGEP2_32(
          RI.ElementType, GlobValPtr, 0, 1, ".imagp");
      Value *SrcImg = Builder.CreateLoad(
          RI.ElementType->getStructElementType(1), SrcImgPtr, ".imag");
 
      Value *DestRealPtr = Builder.CreateConstInBoundsGEP2_32(
          RI.ElementType, ElemPtr, 0, 0, ".realp");
      Value *DestImgPtr = Builder.CreateConstInBoundsGEP2_32(
          RI.ElementType, ElemPtr, 0, 1, ".imagp");
      Builder.CreateStore(SrcReal, DestRealPtr);
      Builder.CreateStore(SrcImg, DestImgPtr);
      break;
    }
    case EvalKind::Aggregate: {
      Value *SizeVal =
          Builder.getInt64(M.getDataLayout().getTypeStoreSize(RI.ElementType));
      Builder.CreateMemCpy(
          ElemPtr, M.getDataLayout().getPrefTypeAlign(RI.ElementType),
          GlobValPtr, M.getDataLayout().getPrefTypeAlign(RI.ElementType),
          SizeVal, false);
      break;
    }
    }
  }
 
  Builder.CreateRetVoid();
  Builder.restoreIP(OldIP);
  return LtGCFunc;
}
 
Function *OpenMPIRBuilder::emitGlobalToListReduceFunction(
    ArrayRef<ReductionInfo> ReductionInfos, Function *ReduceFn,
    Type *ReductionsBufferTy, AttributeList FuncAttrs) {
  OpenMPIRBuilder::InsertPointTy OldIP = Builder.saveIP();
  LLVMContext &Ctx = M.getContext();
  auto *FuncTy = FunctionType::get(
      Builder.getVoidTy(),
      {Builder.getPtrTy(), Builder.getInt32Ty(), Builder.getPtrTy()},
      /* IsVarArg */ false);
  Function *LtGRFunc =
      Function::Create(FuncTy, GlobalVariable::InternalLinkage,
                       "_omp_reduction_global_to_list_reduce_func", &M);
  LtGRFunc->setAttributes(FuncAttrs);
  LtGRFunc->addParamAttr(0, Attribute::NoUndef);
  LtGRFunc->addParamAttr(1, Attribute::NoUndef);
  LtGRFunc->addParamAttr(2, Attribute::NoUndef);
 
  BasicBlock *EntryBlock = BasicBlock::Create(Ctx, "entry", LtGRFunc);
  Builder.SetInsertPoint(EntryBlock);
 
  // Buffer: global reduction buffer.
  Argument *BufferArg = LtGRFunc->getArg(0);
  // Idx: index of the buffer.
  Argument *IdxArg = LtGRFunc->getArg(1);
  // ReduceList: thread local Reduce list.
  Argument *ReduceListArg = LtGRFunc->getArg(2);
 
  Value *BufferArgAlloca = Builder.CreateAlloca(Builder.getPtrTy(), nullptr,
                                                BufferArg->getName() + ".addr");
  Value *IdxArgAlloca = Builder.CreateAlloca(Builder.getInt32Ty(), nullptr,
                                             IdxArg->getName() + ".addr");
  Value *ReduceListArgAlloca = Builder.CreateAlloca(
      Builder.getPtrTy(), nullptr, ReduceListArg->getName() + ".addr");
  ArrayType *RedListArrayTy =
      ArrayType::get(Builder.getPtrTy(), ReductionInfos.size());
 
  // 1. Build a list of reduction variables.
  // void *RedList[<n>] = {<ReductionVars>[0], ..., <ReductionVars>[<n>-1]};
  Value *LocalReduceList =
      Builder.CreateAlloca(RedListArrayTy, nullptr, ".omp.reduction.red_list");
 
  Value *BufferArgAddrCast = Builder.CreatePointerBitCastOrAddrSpaceCast(
      BufferArgAlloca, Builder.getPtrTy(),
      BufferArgAlloca->getName() + ".ascast");
  Value *IdxArgAddrCast = Builder.CreatePointerBitCastOrAddrSpaceCast(
      IdxArgAlloca, Builder.getPtrTy(), IdxArgAlloca->getName() + ".ascast");
  Value *ReduceListArgAddrCast = Builder.CreatePointerBitCastOrAddrSpaceCast(
      ReduceListArgAlloca, Builder.getPtrTy(),
      ReduceListArgAlloca->getName() + ".ascast");
  Value *ReductionList = Builder.CreatePointerBitCastOrAddrSpaceCast(
      LocalReduceList, Builder.getPtrTy(),
      LocalReduceList->getName() + ".ascast");
 
  Builder.CreateStore(BufferArg, BufferArgAddrCast);
  Builder.CreateStore(IdxArg, IdxArgAddrCast);
  Builder.CreateStore(ReduceListArg, ReduceListArgAddrCast);
 
  Value *BufferVal = Builder.CreateLoad(Builder.getPtrTy(), BufferArgAddrCast);
  Value *Idxs[] = {Builder.CreateLoad(Builder.getInt32Ty(), IdxArgAddrCast)};
  Type *IndexTy = Builder.getIndexTy(
      M.getDataLayout(), M.getDataLayout().getDefaultGlobalsAddressSpace());
  for (auto En : enumerate(ReductionInfos)) {
    Value *TargetElementPtrPtr = Builder.CreateInBoundsGEP(
        RedListArrayTy, ReductionList,
        {ConstantInt::get(IndexTy, 0), ConstantInt::get(IndexTy, En.index())});
    // Global = Buffer.VD[Idx];
    Value *BufferVD =
        Builder.CreateInBoundsGEP(ReductionsBufferTy, BufferVal, Idxs);
    Value *GlobValPtr = Builder.CreateConstInBoundsGEP2_32(
        ReductionsBufferTy, BufferVD, 0, En.index());
    Builder.CreateStore(GlobValPtr, TargetElementPtrPtr);
  }
 
  // Call reduce_function(ReduceList, GlobalReduceList)
  Value *ReduceList =
      Builder.CreateLoad(Builder.getPtrTy(), ReduceListArgAddrCast);
  Builder.CreateCall(ReduceFn, {ReduceList, ReductionList})
      ->addFnAttr(Attribute::NoUnwind);
  Builder.CreateRetVoid();
  Builder.restoreIP(OldIP);
  return LtGRFunc;
}
 
std::string OpenMPIRBuilder::getReductionFuncName(StringRef Name) const {
  std::string Suffix =
      createPlatformSpecificName({"omp", "reduction", "reduction_func"});
  return (Name + Suffix).str();
}
 
Function *OpenMPIRBuilder::createReductionFunction(
    StringRef ReducerName, ArrayRef<ReductionInfo> ReductionInfos,
    ReductionGenCBKind ReductionGenCBKind, AttributeList FuncAttrs) {
  auto *FuncTy = FunctionType::get(Builder.getVoidTy(),
                                   {Builder.getPtrTy(), Builder.getPtrTy()},
                                   /* IsVarArg */ false);
  std::string Name = getReductionFuncName(ReducerName);
  Function *ReductionFunc =
      Function::Create(FuncTy, GlobalVariable::InternalLinkage, Name, &M);
  ReductionFunc->setAttributes(FuncAttrs);
  ReductionFunc->addParamAttr(0, Attribute::NoUndef);
  ReductionFunc->addParamAttr(1, Attribute::NoUndef);
  BasicBlock *EntryBB =
      BasicBlock::Create(M.getContext(), "entry", ReductionFunc);
  Builder.SetInsertPoint(EntryBB);
 
  // Need to alloca memory here and deal with the pointers before getting
  // LHS/RHS pointers out
  Value *LHSArrayPtr = nullptr;
  Value *RHSArrayPtr = nullptr;
  Argument *Arg0 = ReductionFunc->getArg(0);
  Argument *Arg1 = ReductionFunc->getArg(1);
  Type *Arg0Type = Arg0->getType();
  Type *Arg1Type = Arg1->getType();
 
  Value *LHSAlloca =
      Builder.CreateAlloca(Arg0Type, nullptr, Arg0->getName() + ".addr");
  Value *RHSAlloca =
      Builder.CreateAlloca(Arg1Type, nullptr, Arg1->getName() + ".addr");
  Value *LHSAddrCast = Builder.CreatePointerBitCastOrAddrSpaceCast(
      LHSAlloca, Arg0Type, LHSAlloca->getName() + ".ascast");
  Value *RHSAddrCast = Builder.CreatePointerBitCastOrAddrSpaceCast(
      RHSAlloca, Arg1Type, RHSAlloca->getName() + ".ascast");
  Builder.CreateStore(Arg0, LHSAddrCast);
  Builder.CreateStore(Arg1, RHSAddrCast);
  LHSArrayPtr = Builder.CreateLoad(Arg0Type, LHSAddrCast);
  RHSArrayPtr = Builder.CreateLoad(Arg1Type, RHSAddrCast);
 
  Type *RedArrayTy = ArrayType::get(Builder.getPtrTy(), ReductionInfos.size());
  Type *IndexTy = Builder.getIndexTy(
      M.getDataLayout(), M.getDataLayout().getDefaultGlobalsAddressSpace());
  SmallVector<Value *> LHSPtrs, RHSPtrs;
  for (auto En : enumerate(ReductionInfos)) {
    const ReductionInfo &RI = En.value();
    Value *RHSI8PtrPtr = Builder.CreateInBoundsGEP(
        RedArrayTy, RHSArrayPtr,
        {ConstantInt::get(IndexTy, 0), ConstantInt::get(IndexTy, En.index())});
    Value *RHSI8Ptr = Builder.CreateLoad(Builder.getPtrTy(), RHSI8PtrPtr);
    Value *RHSPtr = Builder.CreatePointerBitCastOrAddrSpaceCast(
        RHSI8Ptr, RI.PrivateVariable->getType(),
        RHSI8Ptr->getName() + ".ascast");
 
    Value *LHSI8PtrPtr = Builder.CreateInBoundsGEP(
        RedArrayTy, LHSArrayPtr,
        {ConstantInt::get(IndexTy, 0), ConstantInt::get(IndexTy, En.index())});
    Value *LHSI8Ptr = Builder.CreateLoad(Builder.getPtrTy(), LHSI8PtrPtr);
    Value *LHSPtr = Builder.CreatePointerBitCastOrAddrSpaceCast(
        LHSI8Ptr, RI.Variable->getType(), LHSI8Ptr->getName() + ".ascast");
 
    if (ReductionGenCBKind == ReductionGenCBKind::Clang) {
      LHSPtrs.emplace_back(LHSPtr);
      RHSPtrs.emplace_back(RHSPtr);
    } else {
      Value *LHS = Builder.CreateLoad(RI.ElementType, LHSPtr);
      Value *RHS = Builder.CreateLoad(RI.ElementType, RHSPtr);
      Value *Reduced;
      RI.ReductionGen(Builder.saveIP(), LHS, RHS, Reduced);
      if (!Builder.GetInsertBlock())
        return ReductionFunc;
      Builder.CreateStore(Reduced, LHSPtr);
    }
  }
 
  if (ReductionGenCBKind == ReductionGenCBKind::Clang)
    for (auto En : enumerate(ReductionInfos)) {
      unsigned Index = En.index();
      const ReductionInfo &RI = En.value();
      Value *LHSFixupPtr, *RHSFixupPtr;
      Builder.restoreIP(RI.ReductionGenClang(
          Builder.saveIP(), Index, &LHSFixupPtr, &RHSFixupPtr, ReductionFunc));
 
      // Fix the CallBack code genereated to use the correct Values for the LHS
      // and RHS
      LHSFixupPtr->replaceUsesWithIf(
          LHSPtrs[Index], [ReductionFunc](const Use &U) {
            return cast<Instruction>(U.getUser())->getParent()->getParent() ==
                   ReductionFunc;
          });
      RHSFixupPtr->replaceUsesWithIf(
          RHSPtrs[Index], [ReductionFunc](const Use &U) {
            return cast<Instruction>(U.getUser())->getParent()->getParent() ==
                   ReductionFunc;
          });
    }
 
  Builder.CreateRetVoid();
  return ReductionFunc;
}
 
static void
checkReductionInfos(ArrayRef<OpenMPIRBuilder::ReductionInfo> ReductionInfos,
                    bool IsGPU) {
  for (const OpenMPIRBuilder::ReductionInfo &RI : ReductionInfos) {
    (void)RI;
    assert(RI.Variable && "expected non-null variable");
    assert(RI.PrivateVariable && "expected non-null private variable");
    assert((RI.ReductionGen || RI.ReductionGenClang) &&
           "expected non-null reduction generator callback");
    if (!IsGPU) {
      assert(
          RI.Variable->getType() == RI.PrivateVariable->getType() &&
          "expected variables and their private equivalents to have the same "
          "type");
    }
    assert(RI.Variable->getType()->isPointerTy() &&
           "expected variables to be pointers");
  }
}
 
OpenMPIRBuilder::InsertPointTy OpenMPIRBuilder::createReductionsGPU(
    const LocationDescription &Loc, InsertPointTy AllocaIP,
    InsertPointTy CodeGenIP, ArrayRef<ReductionInfo> ReductionInfos,
    bool IsNoWait, bool IsTeamsReduction, bool HasDistribute,
    ReductionGenCBKind ReductionGenCBKind, std::optional<omp::GV> GridValue,
    unsigned ReductionBufNum, Value *SrcLocInfo) {
  if (!updateToLocation(Loc))
    return InsertPointTy();
  Builder.restoreIP(CodeGenIP);
  checkReductionInfos(ReductionInfos, /*IsGPU*/ true);
  LLVMContext &Ctx = M.getContext();
 
  // Source location for the ident struct
  if (!SrcLocInfo) {
    uint32_t SrcLocStrSize;
    Constant *SrcLocStr = getOrCreateSrcLocStr(Loc, SrcLocStrSize);
    SrcLocInfo = getOrCreateIdent(SrcLocStr, SrcLocStrSize);
  }
 
  if (ReductionInfos.size() == 0)
    return Builder.saveIP();
 
  Function *CurFunc = Builder.GetInsertBlock()->getParent();
  AttributeList FuncAttrs;
  AttrBuilder AttrBldr(Ctx);
  for (auto Attr : CurFunc->getAttributes().getFnAttrs())
    AttrBldr.addAttribute(Attr);
  AttrBldr.removeAttribute(Attribute::OptimizeNone);
  FuncAttrs = FuncAttrs.addFnAttributes(Ctx, AttrBldr);
 
  Function *ReductionFunc = nullptr;
  CodeGenIP = Builder.saveIP();
  ReductionFunc =
      createReductionFunction(Builder.GetInsertBlock()->getParent()->getName(),
                              ReductionInfos, ReductionGenCBKind, FuncAttrs);
  Builder.restoreIP(CodeGenIP);
 
  // Set the grid value in the config needed for lowering later on
  if (GridValue.has_value())
    Config.setGridValue(GridValue.value());
  else
    Config.setGridValue(getGridValue(T, ReductionFunc));
 
  // Build res = __kmpc_reduce{_nowait}(<gtid>, <n>, sizeof(RedList),
  // RedList, shuffle_reduce_func, interwarp_copy_func);
  // or
  // Build res = __kmpc_reduce_teams_nowait_simple(<loc>, <gtid>, <lck>);
  Value *Res;
 
  // 1. Build a list of reduction variables.
  // void *RedList[<n>] = {<ReductionVars>[0], ..., <ReductionVars>[<n>-1]};
  auto Size = ReductionInfos.size();
  Type *PtrTy = PointerType::getUnqual(Ctx);
  Type *RedArrayTy = ArrayType::get(PtrTy, Size);
  CodeGenIP = Builder.saveIP();
  Builder.restoreIP(AllocaIP);
  Value *ReductionListAlloca =
      Builder.CreateAlloca(RedArrayTy, nullptr, ".omp.reduction.red_list");
  Value *ReductionList = Builder.CreatePointerBitCastOrAddrSpaceCast(
      ReductionListAlloca, PtrTy, ReductionListAlloca->getName() + ".ascast");
  Builder.restoreIP(CodeGenIP);
  Type *IndexTy = Builder.getIndexTy(
      M.getDataLayout(), M.getDataLayout().getDefaultGlobalsAddressSpace());
  for (auto En : enumerate(ReductionInfos)) {
    const ReductionInfo &RI = En.value();
    Value *ElemPtr = Builder.CreateInBoundsGEP(
        RedArrayTy, ReductionList,
        {ConstantInt::get(IndexTy, 0), ConstantInt::get(IndexTy, En.index())});
    Value *CastElem =
        Builder.CreatePointerBitCastOrAddrSpaceCast(RI.PrivateVariable, PtrTy);
    Builder.CreateStore(CastElem, ElemPtr);
  }
  CodeGenIP = Builder.saveIP();
  Function *SarFunc =
      emitShuffleAndReduceFunction(ReductionInfos, ReductionFunc, FuncAttrs);
  Function *WcFunc = emitInterWarpCopyFunction(Loc, ReductionInfos, FuncAttrs);
  Builder.restoreIP(CodeGenIP);
 
  Value *RL = Builder.CreatePointerBitCastOrAddrSpaceCast(ReductionList, PtrTy);
 
  unsigned MaxDataSize = 0;
  SmallVector<Type *> ReductionTypeArgs;
  for (auto En : enumerate(ReductionInfos)) {
    auto Size = M.getDataLayout().getTypeStoreSize(En.value().ElementType);
    if (Size > MaxDataSize)
      MaxDataSize = Size;
    ReductionTypeArgs.emplace_back(En.value().ElementType);
  }
  Value *ReductionDataSize =
      Builder.getInt64(MaxDataSize * ReductionInfos.size());
  if (!IsTeamsReduction) {
    Value *SarFuncCast =
        Builder.CreatePointerBitCastOrAddrSpaceCast(SarFunc, PtrTy);
    Value *WcFuncCast =
        Builder.CreatePointerBitCastOrAddrSpaceCast(WcFunc, PtrTy);
    Value *Args[] = {SrcLocInfo, ReductionDataSize, RL, SarFuncCast,
                     WcFuncCast};
    Function *Pv2Ptr = getOrCreateRuntimeFunctionPtr(
        RuntimeFunction::OMPRTL___kmpc_nvptx_parallel_reduce_nowait_v2);
    Res = Builder.CreateCall(Pv2Ptr, Args);
  } else {
    CodeGenIP = Builder.saveIP();
    StructType *ReductionsBufferTy = StructType::create(
        Ctx, ReductionTypeArgs, "struct._globalized_locals_ty");
    Function *RedFixedBuferFn = getOrCreateRuntimeFunctionPtr(
        RuntimeFunction::OMPRTL___kmpc_reduction_get_fixed_buffer);
    Function *LtGCFunc = emitListToGlobalCopyFunction(
        ReductionInfos, ReductionsBufferTy, FuncAttrs);
    Function *LtGRFunc = emitListToGlobalReduceFunction(
        ReductionInfos, ReductionFunc, ReductionsBufferTy, FuncAttrs);
    Function *GtLCFunc = emitGlobalToListCopyFunction(
        ReductionInfos, ReductionsBufferTy, FuncAttrs);
    Function *GtLRFunc = emitGlobalToListReduceFunction(
        ReductionInfos, ReductionFunc, ReductionsBufferTy, FuncAttrs);
    Builder.restoreIP(CodeGenIP);
 
    Value *KernelTeamsReductionPtr = Builder.CreateCall(
        RedFixedBuferFn, {}, "_openmp_teams_reductions_buffer_$_$ptr");
 
    Value *Args3[] = {SrcLocInfo,
                      KernelTeamsReductionPtr,
                      Builder.getInt32(ReductionBufNum),
                      ReductionDataSize,
                      RL,
                      SarFunc,
                      WcFunc,
                      LtGCFunc,
                      LtGRFunc,
                      GtLCFunc,
                      GtLRFunc};
 
    Function *TeamsReduceFn = getOrCreateRuntimeFunctionPtr(
        RuntimeFunction::OMPRTL___kmpc_nvptx_teams_reduce_nowait_v2);
    Res = Builder.CreateCall(TeamsReduceFn, Args3);
  }
 
  // 5. Build if (res == 1)
  BasicBlock *ExitBB = BasicBlock::Create(Ctx, ".omp.reduction.done");
  BasicBlock *ThenBB = BasicBlock::Create(Ctx, ".omp.reduction.then");
  Value *Cond = Builder.CreateICmpEQ(Res, Builder.getInt32(1));
  Builder.CreateCondBr(Cond, ThenBB, ExitBB);
 
  // 6. Build then branch: where we have reduced values in the master
  //    thread in each team.
  //    __kmpc_end_reduce{_nowait}(<gtid>);
  //    break;
  emitBlock(ThenBB, CurFunc);
 
  // Add emission of __kmpc_end_reduce{_nowait}(<gtid>);
  for (auto En : enumerate(ReductionInfos)) {
    const ReductionInfo &RI = En.value();
    Value *LHS = RI.Variable;
    Value *RHS =
        Builder.CreatePointerBitCastOrAddrSpaceCast(RI.PrivateVariable, PtrTy);
 
    if (ReductionGenCBKind == ReductionGenCBKind::Clang) {
      Value *LHSPtr, *RHSPtr;
      Builder.restoreIP(RI.ReductionGenClang(Builder.saveIP(), En.index(),
                                             &LHSPtr, &RHSPtr, CurFunc));
 
      // Fix the CallBack code genereated to use the correct Values for the LHS
      // and RHS
      LHSPtr->replaceUsesWithIf(LHS, [ReductionFunc](const Use &U) {
        return cast<Instruction>(U.getUser())->getParent()->getParent() ==
               ReductionFunc;
      });
      RHSPtr->replaceUsesWithIf(RHS, [ReductionFunc](const Use &U) {
        return cast<Instruction>(U.getUser())->getParent()->getParent() ==
               ReductionFunc;
      });
    } else {
      assert(false && "Unhandled ReductionGenCBKind");
    }
  }
  emitBlock(ExitBB, CurFunc);
 
  Config.setEmitLLVMUsed();
 
  return Builder.saveIP();
}
 
static Function *getFreshReductionFunc(Module &M) {
  Type *VoidTy = Type::getVoidTy(M.getContext());
  Type *Int8PtrTy = PointerType::getUnqual(M.getContext());
  auto *FuncTy =
      FunctionType::get(VoidTy, {Int8PtrTy, Int8PtrTy}, /* IsVarArg */ false);
  return Function::Create(FuncTy, GlobalVariable::InternalLinkage,
                          ".omp.reduction.func", &M);
}
 
OpenMPIRBuilder::InsertPointTy
OpenMPIRBuilder::createReductions(const LocationDescription &Loc,
                                  InsertPointTy AllocaIP,
                                  ArrayRef<ReductionInfo> ReductionInfos,
                                  ArrayRef<bool> IsByRef, bool IsNoWait) {
  assert(ReductionInfos.size() == IsByRef.size());
  for (const ReductionInfo &RI : ReductionInfos) {
    (void)RI;
    assert(RI.Variable && "expected non-null variable");
    assert(RI.PrivateVariable && "expected non-null private variable");
    assert(RI.ReductionGen && "expected non-null reduction generator callback");
    assert(RI.Variable->getType() == RI.PrivateVariable->getType() &&
           "expected variables and their private equivalents to have the same "
           "type");
    assert(RI.Variable->getType()->isPointerTy() &&
           "expected variables to be pointers");
  }
 
  if (!updateToLocation(Loc))
    return InsertPointTy();
 
  BasicBlock *InsertBlock = Loc.IP.getBlock();
  BasicBlock *ContinuationBlock =
      InsertBlock->splitBasicBlock(Loc.IP.getPoint(), "reduce.finalize");
  InsertBlock->getTerminator()->eraseFromParent();
 
  // Create and populate array of type-erased pointers to private reduction
  // values.
  unsigned NumReductions = ReductionInfos.size();
  Type *RedArrayTy = ArrayType::get(Builder.getPtrTy(), NumReductions);
  Builder.SetInsertPoint(AllocaIP.getBlock()->getTerminator());
  Value *RedArray = Builder.CreateAlloca(RedArrayTy, nullptr, "red.array");
 
  Builder.SetInsertPoint(InsertBlock, InsertBlock->end());
 
  for (auto En : enumerate(ReductionInfos)) {
    unsigned Index = En.index();
    const ReductionInfo &RI = En.value();
    Value *RedArrayElemPtr = Builder.CreateConstInBoundsGEP2_64(
        RedArrayTy, RedArray, 0, Index, "red.array.elem." + Twine(Index));
    Builder.CreateStore(RI.PrivateVariable, RedArrayElemPtr);
  }
 
  // Emit a call to the runtime function that orchestrates the reduction.
  // Declare the reduction function in the process.
  Function *Func = Builder.GetInsertBlock()->getParent();
  Module *Module = Func->getParent();
  uint32_t SrcLocStrSize;
  Constant *SrcLocStr = getOrCreateSrcLocStr(Loc, SrcLocStrSize);
  bool CanGenerateAtomic = all_of(ReductionInfos, [](const ReductionInfo &RI) {
    return RI.AtomicReductionGen;
  });
  Value *Ident = getOrCreateIdent(SrcLocStr, SrcLocStrSize,
                                  CanGenerateAtomic
                                      ? IdentFlag::OMP_IDENT_FLAG_ATOMIC_REDUCE
                                      : IdentFlag(0));
  Value *ThreadId = getOrCreateThreadID(Ident);
  Constant *NumVariables = Builder.getInt32(NumReductions);
  const DataLayout &DL = Module->getDataLayout();
  unsigned RedArrayByteSize = DL.getTypeStoreSize(RedArrayTy);
  Constant *RedArraySize = Builder.getInt64(RedArrayByteSize);
  Function *ReductionFunc = getFreshReductionFunc(*Module);
  Value *Lock = getOMPCriticalRegionLock(".reduction");
  Function *ReduceFunc = getOrCreateRuntimeFunctionPtr(
      IsNoWait ? RuntimeFunction::OMPRTL___kmpc_reduce_nowait
               : RuntimeFunction::OMPRTL___kmpc_reduce);
  CallInst *ReduceCall =
      Builder.CreateCall(ReduceFunc,
                         {Ident, ThreadId, NumVariables, RedArraySize, RedArray,
                          ReductionFunc, Lock},
                         "reduce");
 
  // Create final reduction entry blocks for the atomic and non-atomic case.
  // Emit IR that dispatches control flow to one of the blocks based on the
  // reduction supporting the atomic mode.
  BasicBlock *NonAtomicRedBlock =
      BasicBlock::Create(Module->getContext(), "reduce.switch.nonatomic", Func);
  BasicBlock *AtomicRedBlock =
      BasicBlock::Create(Module->getContext(), "reduce.switch.atomic", Func);
  SwitchInst *Switch =
      Builder.CreateSwitch(ReduceCall, ContinuationBlock, /* NumCases */ 2);
  Switch->addCase(Builder.getInt32(1), NonAtomicRedBlock);
  Switch->addCase(Builder.getInt32(2), AtomicRedBlock);
 
  // Populate the non-atomic reduction using the elementwise reduction function.
  // This loads the elements from the global and private variables and reduces
  // them before storing back the result to the global variable.
  Builder.SetInsertPoint(NonAtomicRedBlock);
  for (auto En : enumerate(ReductionInfos)) {
    const ReductionInfo &RI = En.value();
    Type *ValueType = RI.ElementType;
    // We have one less load for by-ref case because that load is now inside of
    // the reduction region
    Value *RedValue = nullptr;
    if (!IsByRef[En.index()]) {
      RedValue = Builder.CreateLoad(ValueType, RI.Variable,
                                    "red.value." + Twine(En.index()));
    }
    Value *PrivateRedValue =
        Builder.CreateLoad(ValueType, RI.PrivateVariable,
                           "red.private.value." + Twine(En.index()));
    Value *Reduced;
    if (IsByRef[En.index()]) {
      Builder.restoreIP(RI.ReductionGen(Builder.saveIP(), RI.Variable,
                                        PrivateRedValue, Reduced));
    } else {
      Builder.restoreIP(RI.ReductionGen(Builder.saveIP(), RedValue,
                                        PrivateRedValue, Reduced));
    }
    if (!Builder.GetInsertBlock())
      return InsertPointTy();
    // for by-ref case, the load is inside of the reduction region
    if (!IsByRef[En.index()])
      Builder.CreateStore(Reduced, RI.Variable);
  }
  Function *EndReduceFunc = getOrCreateRuntimeFunctionPtr(
      IsNoWait ? RuntimeFunction::OMPRTL___kmpc_end_reduce_nowait
               : RuntimeFunction::OMPRTL___kmpc_end_reduce);
  Builder.CreateCall(EndReduceFunc, {Ident, ThreadId, Lock});
  Builder.CreateBr(ContinuationBlock);
 
  // Populate the atomic reduction using the atomic elementwise reduction
  // function. There are no loads/stores here because they will be happening
  // inside the atomic elementwise reduction.
  Builder.SetInsertPoint(AtomicRedBlock);
  if (CanGenerateAtomic && llvm::none_of(IsByRef, [](bool P) { return P; })) {
    for (const ReductionInfo &RI : ReductionInfos) {
      Builder.restoreIP(RI.AtomicReductionGen(Builder.saveIP(), RI.ElementType,
                                              RI.Variable, RI.PrivateVariable));
      if (!Builder.GetInsertBlock())
        return InsertPointTy();
    }
    Builder.CreateBr(ContinuationBlock);
  } else {
    Builder.CreateUnreachable();
  }
 
  // Populate the outlined reduction function using the elementwise reduction
  // function. Partial values are extracted from the type-erased array of
  // pointers to private variables.
  BasicBlock *ReductionFuncBlock =
      BasicBlock::Create(Module->getContext(), "", ReductionFunc);
  Builder.SetInsertPoint(ReductionFuncBlock);
  Value *LHSArrayPtr = ReductionFunc->getArg(0);
  Value *RHSArrayPtr = ReductionFunc->getArg(1);
 
  for (auto En : enumerate(ReductionInfos)) {
    const ReductionInfo &RI = En.value();
    Value *LHSI8PtrPtr = Builder.CreateConstInBoundsGEP2_64(
        RedArrayTy, LHSArrayPtr, 0, En.index());
    Value *LHSI8Ptr = Builder.CreateLoad(Builder.getPtrTy(), LHSI8PtrPtr);
    Value *LHSPtr = Builder.CreateBitCast(LHSI8Ptr, RI.Variable->getType());
    Value *LHS = Builder.CreateLoad(RI.ElementType, LHSPtr);
    Value *RHSI8PtrPtr = Builder.CreateConstInBoundsGEP2_64(
        RedArrayTy, RHSArrayPtr, 0, En.index());
    Value *RHSI8Ptr = Builder.CreateLoad(Builder.getPtrTy(), RHSI8PtrPtr);
    Value *RHSPtr =
        Builder.CreateBitCast(RHSI8Ptr, RI.PrivateVariable->getType());
    Value *RHS = Builder.CreateLoad(RI.ElementType, RHSPtr);
    Value *Reduced;
    Builder.restoreIP(RI.ReductionGen(Builder.saveIP(), LHS, RHS, Reduced));
    if (!Builder.GetInsertBlock())
      return InsertPointTy();
    // store is inside of the reduction region when using by-ref
    if (!IsByRef[En.index()])
      Builder.CreateStore(Reduced, LHSPtr);
  }
  Builder.CreateRetVoid();
 
  Builder.SetInsertPoint(ContinuationBlock);
  return Builder.saveIP();
}
 
OpenMPIRBuilder::InsertPointTy
OpenMPIRBuilder::createMaster(const LocationDescription &Loc,
                              BodyGenCallbackTy BodyGenCB,
                              FinalizeCallbackTy FiniCB) {
 
  if (!updateToLocation(Loc))
    return Loc.IP;
 
  Directive OMPD = Directive::OMPD_master;
  uint32_t SrcLocStrSize;
  Constant *SrcLocStr = getOrCreateSrcLocStr(Loc, SrcLocStrSize);
  Value *Ident = getOrCreateIdent(SrcLocStr, SrcLocStrSize);
  Value *ThreadId = getOrCreateThreadID(Ident);
  Value *Args[] = {Ident, ThreadId};
 
  Function *EntryRTLFn = getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_master);
  Instruction *EntryCall = Builder.CreateCall(EntryRTLFn, Args);
 
  Function *ExitRTLFn = getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_end_master);
  Instruction *ExitCall = Builder.CreateCall(ExitRTLFn, Args);
 
  return EmitOMPInlinedRegion(OMPD, EntryCall, ExitCall, BodyGenCB, FiniCB,
                              /*Conditional*/ true, /*hasFinalize*/ true);
}
 
OpenMPIRBuilder::InsertPointTy
OpenMPIRBuilder::createMasked(const LocationDescription &Loc,
                              BodyGenCallbackTy BodyGenCB,
                              FinalizeCallbackTy FiniCB, Value *Filter) {
  if (!updateToLocation(Loc))
    return Loc.IP;
 
  Directive OMPD = Directive::OMPD_masked;
  uint32_t SrcLocStrSize;
  Constant *SrcLocStr = getOrCreateSrcLocStr(Loc, SrcLocStrSize);
  Value *Ident = getOrCreateIdent(SrcLocStr, SrcLocStrSize);
  Value *ThreadId = getOrCreateThreadID(Ident);
  Value *Args[] = {Ident, ThreadId, Filter};
  Value *ArgsEnd[] = {Ident, ThreadId};
 
  Function *EntryRTLFn = getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_masked);
  Instruction *EntryCall = Builder.CreateCall(EntryRTLFn, Args);
 
  Function *ExitRTLFn = getOrCreateRuntimeFunctionPtr(OMPRTL___kmpc_end_masked);
  Instruction *ExitCall = Builder.CreateCall(ExitRTLFn, ArgsEnd);
 
  return EmitOMPInlinedRegion(OMPD, EntryCall, ExitCall, BodyGenCB, FiniCB,
                              /*Conditional*/ true, /*hasFinalize*/ true);
}
 
CanonicalLoopInfo *OpenMPIRBuilder::createLoopSkeleton(
    DebugLoc DL, Value *TripCount, Function *F, BasicBlock *PreInsertBefore,
    BasicBlock *PostInsertBefore, const Twine &Name) {
  Module *M = F->getParent();
  LLVMContext &Ctx = M->getContext();
  Type *IndVarTy = TripCount->getType();
 
  // Create the basic block structure.
  BasicBlock *Preheader =
      BasicBlock::Create(Ctx, "omp_" + Name + ".preheader", F, PreInsertBefore);
  BasicBlock *Header =
      BasicBlock::Create(Ctx, "omp_" + Name + ".header", F, PreInsertBefore);
  BasicBlock *Cond =
      BasicBlock::Create(Ctx, "omp_" + Name + ".cond", F, PreInsertBefore);
  BasicBlock *Body =
      BasicBlock::Create(Ctx, "omp_" + Name + ".body", F, PreInsertBefore);
  BasicBlock *Latch =
      BasicBlock::Create(Ctx, "omp_" + Name + ".inc", F, PostInsertBefore);
  BasicBlock *Exit =
      BasicBlock::Create(Ctx, "omp_" + Name + ".exit", F, PostInsertBefore);
  BasicBlock *After =
      BasicBlock::Create(Ctx, "omp_" + Name + ".after", F, PostInsertBefore);
 
  // Use specified DebugLoc for new instructions.
  Builder.SetCurrentDebugLocation(DL);
 
  Builder.SetInsertPoint(Preheader);
  Builder.CreateBr(Header);
 
  Builder.SetInsertPoint(Header);
  PHINode *IndVarPHI = Builder.CreatePHI(IndVarTy, 2, "omp_" + Name + ".iv");
  IndVarPHI->addIncoming(ConstantInt::get(IndVarTy, 0), Preheader);
  Builder.CreateBr(Cond);
 
  Builder.SetInsertPoint(Cond);
  Value *Cmp =
      Builder.CreateICmpULT(IndVarPHI, TripCount, "omp_" + Name + ".cmp");
  Builder.CreateCondBr(Cmp, Body, Exit);
 
  Builder.SetInsertPoint(Body);
  Builder.CreateBr(Latch);
 
  Builder.SetInsertPoint(Latch);
  Value *Next = Builder.CreateAdd(IndVarPHI, ConstantInt::get(IndVarTy, 1),
                                  "omp_" + Name + ".next", /*HasNUW=*/true);
  Builder.CreateBr(Header);
  IndVarPHI->addIncoming(Next, Latch);
 
  Builder.SetInsertPoint(Exit);
  Builder.CreateBr(After);
 
  // Remember and return the canonical control flow.
  LoopInfos.emplace_front();
  CanonicalLoopInfo *CL = &LoopInfos.front();
 
  CL->Header = Header;
  CL->Cond = Cond;
  CL->Latch = Latch;
  CL->Exit = Exit;
 
#ifndef NDEBUG
  CL->assertOK();
#endif
  return CL;
}
 
CanonicalLoopInfo *
OpenMPIRBuilder::createCanonicalLoop(const LocationDescription &Loc,
                                     LoopBodyGenCallbackTy BodyGenCB,
                                     Value *TripCount, const Twine &Name) {
  BasicBlock *BB = Loc.IP.getBlock();
  BasicBlock *NextBB = BB->getNextNode();
 
  CanonicalLoopInfo *CL = createLoopSkeleton(Loc.DL, TripCount, BB->getParent(),
                                             NextBB, NextBB, Name);
  BasicBlock *After = CL->getAfter();
 
  // If location is not set, don't connect the loop.
  if (updateToLocation(Loc)) {
    // Split the loop at the insertion point: Branch to the preheader and move
    // every following instruction to after the loop (the After BB). Also, the
    // new successor is the loop's after block.
    spliceBB(Builder, After, /*CreateBranch=*/false);
    Builder.CreateBr(CL->getPreheader());
  }
 
  // Emit the body content. We do it after connecting the loop to the CFG to
  // avoid that the callback encounters degenerate BBs.
  BodyGenCB(CL->getBodyIP(), CL->getIndVar());
 
#ifndef NDEBUG
  CL->assertOK();
#endif
  return CL;
}
 
CanonicalLoopInfo *OpenMPIRBuilder::createCanonicalLoop(
    const LocationDescription &Loc, LoopBodyGenCallbackTy BodyGenCB,
    Value *Start, Value *Stop, Value *Step, bool IsSigned, bool InclusiveStop,
    InsertPointTy ComputeIP, const Twine &Name) {
 
  // Consider the following difficulties (assuming 8-bit signed integers):
  //  * Adding \p Step to the loop counter which passes \p Stop may overflow:
  //      DO I = 1, 100, 50
  ///  * A \p Step of INT_MIN cannot not be normalized to a positive direction:
  //      DO I = 100, 0, -128
 
  // Start, Stop and Step must be of the same integer type.
  auto *IndVarTy = cast<IntegerType>(Start->getType());
  assert(IndVarTy == Stop->getType() && "Stop type mismatch");
  assert(IndVarTy == Step->getType() && "Step type mismatch");
 
  LocationDescription ComputeLoc =
      ComputeIP.isSet() ? LocationDescription(ComputeIP, Loc.DL) : Loc;
  updateToLocation(ComputeLoc);
 
  ConstantInt *Zero = ConstantInt::get(IndVarTy, 0);
  ConstantInt *One = ConstantInt::get(IndVarTy, 1);
 
  // Like Step, but always positive.
  Value *Incr = Step;
 
  // Distance between Start and Stop; always positive.
  Value *Span;
 
  // Condition whether there are no iterations are executed at all, e.g. because
  // UB < LB.
  Value *ZeroCmp;
 
  if (IsSigned) {
    // Ensure that increment is positive. If not, negate and invert LB and UB.
    Value *IsNeg = Builder.CreateICmpSLT(Step, Zero);
    Incr = Builder.CreateSelect(IsNeg, Builder.CreateNeg(Step), Step);
    Value *LB = Builder.CreateSelect(IsNeg, Stop, Start);
    Value *UB = Builder.CreateSelect(IsNeg, Start, Stop);
    Span = Builder.CreateSub(UB, LB, "", false, true);
    ZeroCmp = Builder.CreateICmp(
        InclusiveStop ? CmpInst::ICMP_SLT : CmpInst::ICMP_SLE, UB, LB);
  } else {
    Span = Builder.CreateSub(Stop, Start, "", true);
    ZeroCmp = Builder.CreateICmp(
        InclusiveStop ? CmpInst::ICMP_ULT : CmpInst::ICMP_ULE, Stop, Start);
  }
 
  Value *CountIfLooping;
  if (InclusiveStop) {
    CountIfLooping = Builder.CreateAdd(Builder.CreateUDiv(Span, Incr), One);
  } else {
    // Avoid incrementing past stop since it could overflow.
    Value *CountIfTwo = Builder.CreateAdd(
        Builder.CreateUDiv(Builder.CreateSub(Span, One), Incr), One);
    Value *OneCmp = Builder.CreateICmp(CmpInst::ICMP_ULE, Span, Incr);
    CountIfLooping = Builder.CreateSelect(OneCmp, One, CountIfTwo);
  }
  Value *TripCount = Builder.CreateSelect(ZeroCmp, Zero, CountIfLooping,
                                          "omp_" + Name + ".tripcount");
 
  auto BodyGen = [=](InsertPointTy CodeGenIP, Value *IV) {
    Builder.restoreIP(CodeGenIP);
    Value *Span = Builder.CreateMul(IV, Step);
    Value *IndVar = Builder.CreateAdd(Span, Start);
    BodyGenCB(Builder.saveIP(), IndVar);
  };
  LocationDescription LoopLoc = ComputeIP.isSet() ? Loc.IP : Builder.saveIP();
  return createCanonicalLoop(LoopLoc, BodyGen, TripCount, Name);
}
 
// Returns an LLVM function to call for initializing loop bounds using OpenMP
// static scheduling depending on `type`. Only i32 and i64 are supported by the
// runtime. Always interpret integers as unsigned similarly to
// CanonicalLoopInfo.
static FunctionCallee getKmpcForStaticInitForType(Type *Ty, Module &M,
                                                  OpenMPIRBuilder &OMPBuilder) {
  unsigned Bitwidth = Ty->getIntegerBitWidth();
  if (Bitwidth == 32)
    return OMPBuilder.getOrCreateRuntimeFunction(
        M, omp::RuntimeFunction::OMPRTL___kmpc_for_static_init_4u);
  if (Bitwidth == 64)
    return OMPBuilder.getOrCreateRuntimeFunction(
        M, omp::RuntimeFunction::OMPRTL___kmpc_for_static_init_8u);
  llvm_unreachable("unknown OpenMP loop iterator bitwidth");
}
 
OpenMPIRBuilder::InsertPointTy
OpenMPIRBuilder::applyStaticWorkshareLoop(DebugLoc DL, CanonicalLoopInfo *CLI,
                                          InsertPointTy AllocaIP,
                                          bool NeedsBarrier) {
  assert(CLI->isValid() && "Requires a valid canonical loop");
  assert(!isConflictIP(AllocaIP, CLI->getPreheaderIP()) &&
         "Require dedicated allocate IP");
 
  // Set up the source location value for OpenMP runtime.
  Builder.restoreIP(CLI->getPreheaderIP());
  Builder.SetCurrentDebugLocation(DL);
 
  uint32_t SrcLocStrSize;
  Constant *SrcLocStr = getOrCreateSrcLocStr(DL, SrcLocStrSize);
  Value *SrcLoc = getOrCreateIdent(SrcLocStr, SrcLocStrSize);
 
  // Declare useful OpenMP runtime functions.
  Value *IV = CLI->getIndVar();
  Type *IVTy = IV->getType();
  FunctionCallee StaticInit = getKmpcForStaticInitForType(IVTy, M, *this);
  FunctionCallee StaticFini =
      getOrCreateRuntimeFunction(M, omp::OMPRTL___kmpc_for_static_fini);
 
  // Allocate space for computed loop bounds as expected by the "init" function.
  Builder.SetInsertPoint(AllocaIP.getBlock()->getFirstNonPHIOrDbgOrAlloca());
 
  Type *I32Type = Type::getInt32Ty(M.getContext());
  Value *PLastIter = Builder.CreateAlloca(I32Type, nullptr, "p.lastiter");
  Value *PLowerBound = Builder.CreateAlloca(IVTy, nullptr, "p.lowerbound");
  Value *PUpperBound = Builder.CreateAlloca(IVTy, nullptr, "p.upperbound");
  Value *PStride = Builder.CreateAlloca(IVTy, nullptr, "p.stride");
 
  // At the end of the preheader, prepare for calling the "init" function by
  // storing the current loop bounds into the allocated space. A canonical loop
  // always iterates from 0 to trip-count with step 1. Note that "init" expects
  // and produces an inclusive upper bound.
  Builder.SetInsertPoint(CLI->getPreheader()->getTerminator());
  Constant *Zero = ConstantInt::get(IVTy, 0);
  Constant *One = ConstantInt::get(IVTy, 1);
  Builder.CreateStore(Zero, PLowerBound);
  Value *UpperBound = Builder.CreateSub(CLI->getTripCount(), One);
  Builder.CreateStore(UpperBound, PUpperBound);
  Builder.CreateStore(One, PStride);
 
  Value *ThreadNum = getOrCreateThreadID(SrcLoc);
 
  Constant *SchedulingType = ConstantInt::get(
      I32Type, static_cast<int>(OMPScheduleType::UnorderedStatic));
 
  // Call the "init" function and update the trip count of the loop with the
  // value it produced.
  Builder.CreateCall(StaticInit,
                     {SrcLoc, ThreadNum, SchedulingType, PLastIter, PLowerBound,
                      PUpperBound, PStride, One, Zero});
  Value *LowerBound = Builder.CreateLoad(IVTy, PLowerBound);
  Value *InclusiveUpperBound = Builder.CreateLoad(IVTy, PUpperBound);
  Value *TripCountMinusOne = Builder.CreateSub(InclusiveUpperBound, LowerBound);
  Value *TripCount = Builder.CreateAdd(TripCountMinusOne, One);
  CLI->setTripCount(TripCount);
 
  // Update all uses of the induction variable except the one in the condition
  // block that compares it with the actual upper bound, and the increment in
  // the latch block.
 
  CLI->mapIndVar([&](Instruction *OldIV) -> Value * {
    Builder.SetInsertPoint(CLI->getBody(),
                           CLI->getBody()->getFirstInsertionPt());
    Builder.SetCurrentDebugLocation(DL);
    return Builder.CreateAdd(OldIV, LowerBound);
  });
 
  // In the "exit" block, call the "fini" function.
  Builder.SetInsertPoint(CLI->getExit(),
                         CLI->getExit()->getTerminator()->getIterator());
  Builder.CreateCall(StaticFini, {SrcLoc, ThreadNum});
 
  // Add the barrier if requested.
  if (NeedsBarrier)
    createBarrier(LocationDescription(Builder.saveIP(), DL),
                  omp::Directive::OMPD_for, /* ForceSimpleCall */ false,
                  /* CheckCancelFlag */ false);
 
  InsertPointTy AfterIP = CLI->getAfterIP();
  CLI->invalidate();
 
  return AfterIP;
}
 
OpenMPIRBuilder::InsertPointTy OpenMPIRBuilder::applyStaticChunkedWorkshareLoop(
    DebugLoc DL, CanonicalLoopInfo *CLI, InsertPointTy AllocaIP,
    bool NeedsBarrier, Value *ChunkSize) {
  assert(CLI->isValid() && "Requires a valid canonical loop");
  assert(ChunkSize && "Chunk size is required");
 
  LLVMContext &Ctx = CLI->getFunction()->getContext();
  Value *IV = CLI->getIndVar();
  Value *OrigTripCount = CLI->getTripCount();
  Type *IVTy = IV->getType();
  assert(IVTy->getIntegerBitWidth() <= 64 &&
         "Max supported tripcount bitwidth is 64 bits");
  Type *InternalIVTy = IVTy->getIntegerBitWidth() <= 32 ? Type::getInt32Ty(Ctx)
                                                        : Type::getInt64Ty(Ctx);
  Type *I32Type = Type::getInt32Ty(M.getContext());
  Constant *Zero = ConstantInt::get(InternalIVTy, 0);
  Constant *One = ConstantInt::get(InternalIVTy, 1);
 
  // Declare useful OpenMP runtime functions.
  FunctionCallee StaticInit =
      getKmpcForStaticInitForType(InternalIVTy, M, *this);
  FunctionCallee StaticFini =
      getOrCreateRuntimeFunction(M, omp::OMPRTL___kmpc_for_static_fini);
 
  // Allocate space for computed loop bounds as expected by the "init" function.
  Builder.restoreIP(AllocaIP);
  Builder.SetCurrentDebugLocation(DL);
  Value *PLastIter = Builder.CreateAlloca(I32Type, nullptr, "p.lastiter");
  Value *PLowerBound =
      Builder.CreateAlloca(InternalIVTy, nullptr, "p.lowerbound");
  Value *PUpperBound =
      Builder.CreateAlloca(InternalIVTy, nullptr, "p.upperbound");
  Value *PStride = Builder.CreateAlloca(InternalIVTy, nullptr, "p.stride");
 
  // Set up the source location value for the OpenMP runtime.
  Builder.restoreIP(CLI->getPreheaderIP());
  Builder.SetCurrentDebugLocation(DL);
 
  // TODO: Detect overflow in ubsan or max-out with current tripcount.
  Value *CastedChunkSize =
      Builder.CreateZExtOrTrunc(ChunkSize, InternalIVTy, "chunksize");
  Value *CastedTripCount =
      Builder.CreateZExt(OrigTripCount, InternalIVTy, "tripcount");
 
  Constant *SchedulingType = ConstantInt::get(
      I32Type, static_cast<int>(OMPScheduleType::UnorderedStaticChunked));
  Builder.CreateStore(Zero, PLowerBound);
  Value *OrigUpperBound = Builder.CreateSub(CastedTripCount, One);
  Builder.CreateStore(OrigUpperBound, PUpperBound);
  Builder.CreateStore(One, PStride);
 
  // Call the "init" function and update the trip count of the loop with the
  // value it produced.
  uint32_t SrcLocStrSize;
  Constant *SrcLocStr = getOrCreateSrcLocStr(DL, SrcLocStrSize);
  Value *SrcLoc = getOrCreateIdent(SrcLocStr, SrcLocStrSize);
  Value *ThreadNum = getOrCreateThreadID(SrcLoc);
  Builder.CreateCall(StaticInit,
                     {/*loc=*/SrcLoc, /*global_tid=*/ThreadNum,
                      /*schedtype=*/SchedulingType, /*plastiter=*/PLastIter,
                      /*plower=*/PLowerBound, /*pupper=*/PUpperBound,
                      /*pstride=*/PStride, /*incr=*/One,
                      /*chunk=*/CastedChunkSize});
 
  // Load values written by the "init" function.
  Value *FirstChunkStart =
      Builder.CreateLoad(InternalIVTy, PLowerBound, "omp_firstchunk.lb");
  Value *FirstChunkStop =
      Builder.CreateLoad(InternalIVTy, PUpperBound, "omp_firstchunk.ub");
  Value *FirstChunkEnd = Builder.CreateAdd(FirstChunkStop, One);
  Value *ChunkRange =
      Builder.CreateSub(FirstChunkEnd, FirstChunkStart, "omp_chunk.range");
  Value *NextChunkStride =
      Builder.CreateLoad(InternalIVTy, PStride, "omp_dispatch.stride");
 
  // Create outer "dispatch" loop for enumerating the chunks.
  BasicBlock *DispatchEnter = splitBB(Builder, true);
  Value *DispatchCounter;
  CanonicalLoopInfo *DispatchCLI = createCanonicalLoop(
      {Builder.saveIP(), DL},
      [&](InsertPointTy BodyIP, Value *Counter) { DispatchCounter = Counter; },
      FirstChunkStart, CastedTripCount, NextChunkStride,
      /*IsSigned=*/false, /*InclusiveStop=*/false, /*ComputeIP=*/{},
      "dispatch");
 
  // Remember the BasicBlocks of the dispatch loop we need, then invalidate to
  // not have to preserve the canonical invariant.
  BasicBlock *DispatchBody = DispatchCLI->getBody();
  BasicBlock *DispatchLatch = DispatchCLI->getLatch();
  BasicBlock *DispatchExit = DispatchCLI->getExit();
  BasicBlock *DispatchAfter = DispatchCLI->getAfter();
  DispatchCLI->invalidate();
 
  // Rewire the original loop to become the chunk loop inside the dispatch loop.
  redirectTo(DispatchAfter, CLI->getAfter(), DL);
  redirectTo(CLI->getExit(), DispatchLatch, DL);
  redirectTo(DispatchBody, DispatchEnter, DL);
 
  // Prepare the prolog of the chunk loop.
  Builder.restoreIP(CLI->getPreheaderIP());
  Builder.SetCurrentDebugLocation(DL);
 
  // Compute the number of iterations of the chunk loop.
  Builder.SetInsertPoint(CLI->getPreheader()->getTerminator());
  Value *ChunkEnd = Builder.CreateAdd(DispatchCounter, ChunkRange);
  Value *IsLastChunk =
      Builder.CreateICmpUGE(ChunkEnd, CastedTripCount, "omp_chunk.is_last");
  Value *CountUntilOrigTripCount =
      Builder.CreateSub(CastedTripCount, DispatchCounter);
  Value *ChunkTripCount = Builder.CreateSelect(
      IsLastChunk, CountUntilOrigTripCount, ChunkRange, "omp_chunk.tripcount");
  Value *BackcastedChunkTC =
      Builder.CreateTrunc(ChunkTripCount, IVTy, "omp_chunk.tripcount.trunc");
  CLI->setTripCount(BackcastedChunkTC);
 
  // Update all uses of the induction variable except the one in the condition
  // block that compares it with the actual upper bound, and the increment in
  // the latch block.
  Value *BackcastedDispatchCounter =
      Builder.CreateTrunc(DispatchCounter, IVTy, "omp_dispatch.iv.trunc");
  CLI->mapIndVar([&](Instruction *) -> Value * {
    Builder.restoreIP(CLI->getBodyIP());
    return Builder.CreateAdd(IV, BackcastedDispatchCounter);
  });
 
  // In the "exit" block, call the "fini" function.
  Builder.SetInsertPoint(DispatchExit, DispatchExit->getFirstInsertionPt());
  Builder.CreateCall(StaticFini, {SrcLoc, ThreadNum});
 
  // Add the barrier if requested.
  if (NeedsBarrier)
    createBarrier(LocationDescription(Builder.saveIP(), DL), OMPD_for,
                  /*ForceSimpleCall=*/false, /*CheckCancelFlag=*/false);
 
#ifndef NDEBUG
  // Even though we currently do not support applying additional methods to it,