/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp

Bug Summary

File:	llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp
Warning:	line 141, column 27 The result of the left shift is undefined due to shifting by '64', which is greater or equal to the width of type 'long long'

Annotated Source Code

Press '?' to see keyboard shortcuts

Show analyzer invocation

clang -cc1 -cc1 -triple x86_64-pc-linux-gnu -analyze -disable-free -disable-llvm-verifier -discard-value-names -main-file-name SIRegisterInfo.cpp -analyzer-store=region -analyzer-opt-analyze-nested-blocks -analyzer-checker=core -analyzer-checker=apiModeling -analyzer-checker=unix -analyzer-checker=deadcode -analyzer-checker=cplusplus -analyzer-checker=security.insecureAPI.UncheckedReturn -analyzer-checker=security.insecureAPI.getpw -analyzer-checker=security.insecureAPI.gets -analyzer-checker=security.insecureAPI.mktemp -analyzer-checker=security.insecureAPI.mkstemp -analyzer-checker=security.insecureAPI.vfork -analyzer-checker=nullability.NullPassedToNonnull -analyzer-checker=nullability.NullReturnedFromNonnull -analyzer-output plist -w -setup-static-analyzer -analyzer-config-compatibility-mode=true -mrelocation-model pic -pic-level 2 -mframe-pointer=none -fmath-errno -fno-rounding-math -mconstructor-aliases -munwind-tables -target-cpu x86-64 -tune-cpu generic -debugger-tuning=gdb -ffunction-sections -fdata-sections -fcoverage-compilation-dir=/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/build-llvm/lib/Target/AMDGPU -resource-dir /usr/lib/llvm-13/lib/clang/13.0.0 -D _DEBUG -D _GNU_SOURCE -D __STDC_CONSTANT_MACROS -D __STDC_FORMAT_MACROS -D __STDC_LIMIT_MACROS -I /build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/build-llvm/lib/Target/AMDGPU -I /build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU -I /build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/build-llvm/include -I /build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/include -D NDEBUG -U NDEBUG -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/10/../../../../include/c++/10 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/10/../../../../include/x86_64-linux-gnu/c++/10 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/10/../../../../include/c++/10/backward -internal-isystem /usr/lib/llvm-13/lib/clang/13.0.0/include -internal-isystem /usr/local/include -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/10/../../../../x86_64-linux-gnu/include -internal-externc-isystem /usr/include/x86_64-linux-gnu -internal-externc-isystem /include -internal-externc-isystem /usr/include -O2 -Wno-unused-parameter -Wwrite-strings -Wno-missing-field-initializers -Wno-long-long -Wno-maybe-uninitialized -Wno-class-memaccess -Wno-redundant-move -Wno-pessimizing-move -Wno-noexcept-type -Wno-comment -std=c++14 -fdeprecated-macro -fdebug-compilation-dir=/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/build-llvm/lib/Target/AMDGPU -fdebug-prefix-map=/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c=. -ferror-limit 19 -fvisibility hidden -fvisibility-inlines-hidden -stack-protector 2 -fgnuc-version=4.2.1 -vectorize-loops -vectorize-slp -analyzer-output=html -analyzer-config stable-report-filename=true -faddrsig -D__GCC_HAVE_DWARF2_CFI_ASM=1 -o /tmp/scan-build-2021-07-26-235520-9401-1 -x c++ /build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp

/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp

→

1//===-- SIRegisterInfo.cpp - SI Register Information ---------------------===//
2//
3// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4// See https://llvm.org/LICENSE.txt for license information.
5// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6//
7//===----------------------------------------------------------------------===//
8//
9/// \file
10/// SI implementation of the TargetRegisterInfo class.
11//
12//===----------------------------------------------------------------------===//

14#include "SIRegisterInfo.h"
15#include "AMDGPU.h"
16#include "AMDGPURegisterBankInfo.h"
17#include "GCNSubtarget.h"
18#include "MCTargetDesc/AMDGPUInstPrinter.h"
19#include "MCTargetDesc/AMDGPUMCTargetDesc.h"
20#include "SIMachineFunctionInfo.h"
21#include "llvm/CodeGen/LiveIntervals.h"
22#include "llvm/CodeGen/MachineDominators.h"
23#include "llvm/CodeGen/RegisterScavenging.h"

25using namespace llvm;

27#define GET_REGINFO_TARGET_DESC
28#include "AMDGPUGenRegisterInfo.inc"

30static cl::opt<bool> EnableSpillSGPRToVGPR(
"amdgpu-spill-sgpr-to-vgpr",
cl::desc("Enable spilling VGPRs to SGPRs"),
cl::ReallyHidden,
cl::init(true));

36std::array<std::vector<int16_t>, 16> SIRegisterInfo::RegSplitParts;
37std::array<std::array<uint16_t, 32>, 9> SIRegisterInfo::SubRegFromChannelTable;

39// Map numbers of DWORDs to indexes in SubRegFromChannelTable.
40// Valid indexes are shifted 1, such that a 0 mapping means unsupported.
41// e.g. for 8 DWORDs (256-bit), SubRegFromChannelTableWidthMap[8] = 8,
42//      meaning index 7 in SubRegFromChannelTable.
43static const std::array<unsigned, 17> SubRegFromChannelTableWidthMap = {
  0, 1, 2, 3, 4, 5, 6, 7, 8, 0, 0, 0, 0, 0, 0, 0, 9};

46namespace llvm {

48// A temporary struct to spill SGPRs.
49// This is mostly to spill SGPRs to memory. Spilling SGPRs into VGPR lanes emits
50// just v_writelane and v_readlane.
51//
52// When spilling to memory, the SGPRs are written into VGPR lanes and the VGPR
53// is saved to scratch (or the other way around for loads).
54// For this, a VGPR is required where the needed lanes can be clobbered. The
55// RegScavenger can provide a VGPR where currently active lanes can be
56// clobbered, but we still need to save inactive lanes.
57// The high-level steps are:
58// - Try to scavenge SGPR(s) to save exec
59// - Try to scavenge VGPR
60// - Save needed, all or inactive lanes of a TmpVGPR
61// - Spill/Restore SGPRs using TmpVGPR
62// - Restore TmpVGPR
63//
64// To save all lanes of TmpVGPR, exec needs to be saved and modified. If we
65// cannot scavenge temporary SGPRs to save exec, we use the following code:
66// buffer_store_dword TmpVGPR ; only if active lanes need to be saved
67// s_not exec, exec
68// buffer_store_dword TmpVGPR ; save inactive lanes
69// s_not exec, exec
70struct SGPRSpillBuilder {
struct PerVGPRData {
  unsigned PerVGPR;
  unsigned NumVGPRs;
  int64_t VGPRLanes;
};

// The SGPR to save
Register SuperReg;
MachineBasicBlock::iterator MI;
ArrayRef<int16_t> SplitParts;
unsigned NumSubRegs;
bool IsKill;
const DebugLoc &DL;

/* When spilling to stack */
// The SGPRs are written into this VGPR, which is then written to scratch
// (or vice versa for loads).
Register TmpVGPR = AMDGPU::NoRegister;
// Temporary spill slot to save TmpVGPR to.
int TmpVGPRIndex = 0;
// If TmpVGPR is live before the spill or if it is scavenged.
bool TmpVGPRLive = false;
// Scavenged SGPR to save EXEC.
Register SavedExecReg = AMDGPU::NoRegister;
// Stack index to write the SGPRs to.
int Index;
unsigned EltSize = 4;

RegScavenger *RS;
MachineBasicBlock &MBB;
MachineFunction &MF;
SIMachineFunctionInfo &MFI;
const SIInstrInfo &TII;
const SIRegisterInfo &TRI;
bool IsWave32;
Register ExecReg;
unsigned MovOpc;
unsigned NotOpc;

SGPRSpillBuilder(const SIRegisterInfo &TRI, const SIInstrInfo &TII,
                 bool IsWave32, MachineBasicBlock::iterator MI, int Index,
                 RegScavenger *RS)
    : SuperReg(MI->getOperand(0).getReg()), MI(MI),
      IsKill(MI->getOperand(0).isKill()), DL(MI->getDebugLoc()), Index(Index),
      RS(RS), MBB(*MI->getParent()), MF(*MBB.getParent()),
      MFI(*MF.getInfo<SIMachineFunctionInfo>()), TII(TII), TRI(TRI),
      IsWave32(IsWave32) {
  const TargetRegisterClass *RC = TRI.getPhysRegClass(SuperReg);
  SplitParts = TRI.getRegSplitParts(RC, EltSize);
  NumSubRegs = SplitParts.empty() ? 1 : SplitParts.size();

  if (IsWave32) {
    ExecReg = AMDGPU::EXEC_LO;
    MovOpc = AMDGPU::S_MOV_B32;
    NotOpc = AMDGPU::S_NOT_B32;
  } else {
    ExecReg = AMDGPU::EXEC;
    MovOpc = AMDGPU::S_MOV_B64;
    NotOpc = AMDGPU::S_NOT_B64;
  }

  assert(SuperReg != AMDGPU::M0 && "m0 should never spill")(static_cast <bool> (SuperReg != AMDGPU::M0 && "m0 should never spill"
) ? void (0) : __assert_fail ("SuperReg != AMDGPU::M0 && \"m0 should never spill\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 132, __extension__ __PRETTY_FUNCTION__));
  assert(SuperReg != AMDGPU::EXEC_LO && SuperReg != AMDGPU::EXEC_HI &&(static_cast <bool> (SuperReg != AMDGPU::EXEC_LO &&
 SuperReg != AMDGPU::EXEC_HI && SuperReg != AMDGPU::EXEC
 && "exec should never spill") ? void (0) : __assert_fail
 ("SuperReg != AMDGPU::EXEC_LO && SuperReg != AMDGPU::EXEC_HI && SuperReg != AMDGPU::EXEC && \"exec should never spill\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 134, __extension__ __PRETTY_FUNCTION__))
         SuperReg != AMDGPU::EXEC && "exec should never spill")(static_cast <bool> (SuperReg != AMDGPU::EXEC_LO &&
 SuperReg != AMDGPU::EXEC_HI && SuperReg != AMDGPU::EXEC
 && "exec should never spill") ? void (0) : __assert_fail
 ("SuperReg != AMDGPU::EXEC_LO && SuperReg != AMDGPU::EXEC_HI && SuperReg != AMDGPU::EXEC && \"exec should never spill\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 134, __extension__ __PRETTY_FUNCTION__));
}

PerVGPRData getPerVGPRData() {
  PerVGPRData Data;
  Data.PerVGPR = IsWave3215.1
Field 'IsWave32' is false
1
Field 'IsWave32' is false
 ? 32 : 64;
16
←
'?' condition is false→
17
←
The value 64 is assigned to 'Data.PerVGPR'→
  Data.NumVGPRs = (NumSubRegs + (Data.PerVGPR - 1)) / Data.PerVGPR;
  Data.VGPRLanes = (1LL << std::min(Data.PerVGPR, NumSubRegs)) - 1LL;
18
←
Passing value via 1st parameter '__a'→
19
←
Calling 'min<unsigned int>'→
23
←
Returning from 'min<unsigned int>'→
24
←
The result of the left shift is undefined due to shifting by '64', which is greater or equal to the width of type 'long long'
  return Data;
}

// Tries to scavenge SGPRs to save EXEC and a VGPR. Uses v0 if no VGPR is
// free.
// Writes these instructions if an SGPR can be scavenged:
// s_mov_b64 s[6:7], exec   ; Save exec
// s_mov_b64 exec, 3        ; Wanted lanemask
// buffer_store_dword v1    ; Write scavenged VGPR to emergency slot
//
// Writes these instructions if no SGPR can be scavenged:
// buffer_store_dword v0    ; Only if no free VGPR was found
// s_not_b64 exec, exec
// buffer_store_dword v0    ; Save inactive lanes
//                          ; exec stays inverted, it is flipped back in
//                          ; restore.
void prepare() {
  // Scavenged temporary VGPR to use. It must be scavenged once for any number
  // of spilled subregs.
  // FIXME: The liveness analysis is limited and does not tell if a register
  // is in use in lanes that are currently inactive. We can never be sure if
  // a register as actually in use in another lane, so we need to save all
  // used lanes of the chosen VGPR.
  assert(RS && "Cannot spill SGPR to memory without RegScavenger")(static_cast <bool> (RS && "Cannot spill SGPR to memory without RegScavenger"
) ? void (0) : __assert_fail ("RS && \"Cannot spill SGPR to memory without RegScavenger\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 165, __extension__ __PRETTY_FUNCTION__));
8
←
Assuming field 'RS' is non-null→
9
←
'?' condition is true→
  TmpVGPR = RS->scavengeRegister(&AMDGPU::VGPR_32RegClass, MI, 0, false);

  // Reserve temporary stack slot
  TmpVGPRIndex = MFI.getScavengeFI(MF.getFrameInfo(), TRI);
  if (TmpVGPR) {
10
←
Assuming the condition is true→
11
←
Taking true branch→
    // Found a register that is dead in the currently active lanes, we only
    // need to spill inactive lanes.
    TmpVGPRLive = false;
  } else {
    // Pick v0 because it doesn't make a difference.
    TmpVGPR = AMDGPU::VGPR0;
    TmpVGPRLive = true;
  }

  // Try to scavenge SGPRs to save exec
  assert(!SavedExecReg && "Exec is already saved, refuse to save again")(static_cast <bool> (!SavedExecReg && "Exec is already saved, refuse to save again"
) ? void (0) : __assert_fail ("!SavedExecReg && \"Exec is already saved, refuse to save again\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 181, __extension__ __PRETTY_FUNCTION__));
12
←
Assuming the condition is true→
13
←
'?' condition is true→
  const TargetRegisterClass &RC =
      IsWave3213.1
Field 'IsWave32' is false
1
Field 'IsWave32' is false
 ? AMDGPU::SGPR_32RegClass : AMDGPU::SGPR_64RegClass;
14
←
'?' condition is false→
  RS->setRegUsed(SuperReg);
  SavedExecReg = RS->scavengeRegister(&RC, MI, 0, false);

  int64_t VGPRLanes = getPerVGPRData().VGPRLanes;
15
←
Calling 'SGPRSpillBuilder::getPerVGPRData'→

  if (SavedExecReg) {
    RS->setRegUsed(SavedExecReg);
    // Set exec to needed lanes
    BuildMI(MBB, MI, DL, TII.get(MovOpc), SavedExecReg).addReg(ExecReg);
    auto I = BuildMI(MBB, MI, DL, TII.get(MovOpc), ExecReg).addImm(VGPRLanes);
    if (!TmpVGPRLive)
      I.addReg(TmpVGPR, RegState::ImplicitDefine);
    // Spill needed lanes
    TRI.buildVGPRSpillLoadStore(*this, TmpVGPRIndex, 0, /*IsLoad*/ false);
  } else {
    // Spill active lanes
    if (TmpVGPRLive)
      TRI.buildVGPRSpillLoadStore(*this, TmpVGPRIndex, 0, /*IsLoad*/ false,
                                  /*IsKill*/ false);
    // Spill inactive lanes
    auto I = BuildMI(MBB, MI, DL, TII.get(NotOpc), ExecReg).addReg(ExecReg);
    if (!TmpVGPRLive)
      I.addReg(TmpVGPR, RegState::ImplicitDefine);
    TRI.buildVGPRSpillLoadStore(*this, TmpVGPRIndex, 0, /*IsLoad*/ false);
  }
}

// Writes these instructions if an SGPR can be scavenged:
// buffer_load_dword v1     ; Write scavenged VGPR to emergency slot
// s_waitcnt vmcnt(0)       ; If a free VGPR was found
// s_mov_b64 exec, s[6:7]   ; Save exec
//
// Writes these instructions if no SGPR can be scavenged:
// buffer_load_dword v0     ; Restore inactive lanes
// s_waitcnt vmcnt(0)       ; If a free VGPR was found
// s_not_b64 exec, exec
// buffer_load_dword v0     ; Only if no free VGPR was found
void restore() {
  if (SavedExecReg) {
    // Restore used lanes
    TRI.buildVGPRSpillLoadStore(*this, TmpVGPRIndex, 0, /*IsLoad*/ true,
                                /*IsKill*/ false);
    // Restore exec
    auto I = BuildMI(MBB, MI, DL, TII.get(MovOpc), ExecReg)
                 .addReg(SavedExecReg, RegState::Kill);
    // Add an implicit use of the load so it is not dead.
    // FIXME This inserts an unnecessary waitcnt
    if (!TmpVGPRLive) {
      I.addReg(TmpVGPR, RegState::ImplicitKill);
    }
  } else {
    // Restore inactive lanes
    TRI.buildVGPRSpillLoadStore(*this, TmpVGPRIndex, 0, /*IsLoad*/ true,
                                /*IsKill*/ false);
    auto I = BuildMI(MBB, MI, DL, TII.get(NotOpc), ExecReg).addReg(ExecReg);
    if (!TmpVGPRLive) {
      I.addReg(TmpVGPR, RegState::ImplicitKill);
    }
    // Restore active lanes
    if (TmpVGPRLive)
      TRI.buildVGPRSpillLoadStore(*this, TmpVGPRIndex, 0, /*IsLoad*/ true);
  }
}

// Write TmpVGPR to memory or read TmpVGPR from memory.
// Either using a single buffer_load/store if exec is set to the needed mask
// or using
// buffer_load
// s_not exec, exec
// buffer_load
// s_not exec, exec
void readWriteTmpVGPR(unsigned Offset, bool IsLoad) {
  if (SavedExecReg) {
    // Spill needed lanes
    TRI.buildVGPRSpillLoadStore(*this, Index, Offset, IsLoad);
  } else {
    // Spill active lanes
    TRI.buildVGPRSpillLoadStore(*this, Index, Offset, IsLoad,
                                /*IsKill*/ false);
    // Spill inactive lanes
    BuildMI(MBB, MI, DL, TII.get(NotOpc), ExecReg).addReg(ExecReg);
    TRI.buildVGPRSpillLoadStore(*this, Index, Offset, IsLoad);
    BuildMI(MBB, MI, DL, TII.get(NotOpc), ExecReg).addReg(ExecReg);
  }
}
269};

271} // namespace llvm

273SIRegisterInfo::SIRegisterInfo(const GCNSubtarget &ST)
  : AMDGPUGenRegisterInfo(AMDGPU::PC_REG, ST.getAMDGPUDwarfFlavour()), ST(ST),
    SpillSGPRToVGPR(EnableSpillSGPRToVGPR), isWave32(ST.isWave32()) {

assert(getSubRegIndexLaneMask(AMDGPU::sub0).getAsInteger() == 3 &&(static_cast <bool> (getSubRegIndexLaneMask(AMDGPU::sub0
).getAsInteger() == 3 && getSubRegIndexLaneMask(AMDGPU
::sub31).getAsInteger() == (3ULL << 62) && (getSubRegIndexLaneMask
(AMDGPU::lo16) | getSubRegIndexLaneMask(AMDGPU::hi16)).getAsInteger
() == getSubRegIndexLaneMask(AMDGPU::sub0).getAsInteger() &&
 "getNumCoveredRegs() will not work with generated subreg masks!"
) ? void (0) : __assert_fail ("getSubRegIndexLaneMask(AMDGPU::sub0).getAsInteger() == 3 && getSubRegIndexLaneMask(AMDGPU::sub31).getAsInteger() == (3ULL << 62) && (getSubRegIndexLaneMask(AMDGPU::lo16) | getSubRegIndexLaneMask(AMDGPU::hi16)).getAsInteger() == getSubRegIndexLaneMask(AMDGPU::sub0).getAsInteger() && \"getNumCoveredRegs() will not work with generated subreg masks!\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 282, __extension__ __PRETTY_FUNCTION__))
       getSubRegIndexLaneMask(AMDGPU::sub31).getAsInteger() == (3ULL << 62) &&(static_cast <bool> (getSubRegIndexLaneMask(AMDGPU::sub0
).getAsInteger() == 3 && getSubRegIndexLaneMask(AMDGPU
::sub31).getAsInteger() == (3ULL << 62) && (getSubRegIndexLaneMask
(AMDGPU::lo16) | getSubRegIndexLaneMask(AMDGPU::hi16)).getAsInteger
() == getSubRegIndexLaneMask(AMDGPU::sub0).getAsInteger() &&
 "getNumCoveredRegs() will not work with generated subreg masks!"
) ? void (0) : __assert_fail ("getSubRegIndexLaneMask(AMDGPU::sub0).getAsInteger() == 3 && getSubRegIndexLaneMask(AMDGPU::sub31).getAsInteger() == (3ULL << 62) && (getSubRegIndexLaneMask(AMDGPU::lo16) | getSubRegIndexLaneMask(AMDGPU::hi16)).getAsInteger() == getSubRegIndexLaneMask(AMDGPU::sub0).getAsInteger() && \"getNumCoveredRegs() will not work with generated subreg masks!\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 282, __extension__ __PRETTY_FUNCTION__))
       (getSubRegIndexLaneMask(AMDGPU::lo16) |(static_cast <bool> (getSubRegIndexLaneMask(AMDGPU::sub0
).getAsInteger() == 3 && getSubRegIndexLaneMask(AMDGPU
::sub31).getAsInteger() == (3ULL << 62) && (getSubRegIndexLaneMask
(AMDGPU::lo16) | getSubRegIndexLaneMask(AMDGPU::hi16)).getAsInteger
() == getSubRegIndexLaneMask(AMDGPU::sub0).getAsInteger() &&
 "getNumCoveredRegs() will not work with generated subreg masks!"
) ? void (0) : __assert_fail ("getSubRegIndexLaneMask(AMDGPU::sub0).getAsInteger() == 3 && getSubRegIndexLaneMask(AMDGPU::sub31).getAsInteger() == (3ULL << 62) && (getSubRegIndexLaneMask(AMDGPU::lo16) | getSubRegIndexLaneMask(AMDGPU::hi16)).getAsInteger() == getSubRegIndexLaneMask(AMDGPU::sub0).getAsInteger() && \"getNumCoveredRegs() will not work with generated subreg masks!\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 282, __extension__ __PRETTY_FUNCTION__))
        getSubRegIndexLaneMask(AMDGPU::hi16)).getAsInteger() ==(static_cast <bool> (getSubRegIndexLaneMask(AMDGPU::sub0
).getAsInteger() == 3 && getSubRegIndexLaneMask(AMDGPU
::sub31).getAsInteger() == (3ULL << 62) && (getSubRegIndexLaneMask
(AMDGPU::lo16) | getSubRegIndexLaneMask(AMDGPU::hi16)).getAsInteger
() == getSubRegIndexLaneMask(AMDGPU::sub0).getAsInteger() &&
 "getNumCoveredRegs() will not work with generated subreg masks!"
) ? void (0) : __assert_fail ("getSubRegIndexLaneMask(AMDGPU::sub0).getAsInteger() == 3 && getSubRegIndexLaneMask(AMDGPU::sub31).getAsInteger() == (3ULL << 62) && (getSubRegIndexLaneMask(AMDGPU::lo16) | getSubRegIndexLaneMask(AMDGPU::hi16)).getAsInteger() == getSubRegIndexLaneMask(AMDGPU::sub0).getAsInteger() && \"getNumCoveredRegs() will not work with generated subreg masks!\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 282, __extension__ __PRETTY_FUNCTION__))
         getSubRegIndexLaneMask(AMDGPU::sub0).getAsInteger() &&(static_cast <bool> (getSubRegIndexLaneMask(AMDGPU::sub0
).getAsInteger() == 3 && getSubRegIndexLaneMask(AMDGPU
::sub31).getAsInteger() == (3ULL << 62) && (getSubRegIndexLaneMask
(AMDGPU::lo16) | getSubRegIndexLaneMask(AMDGPU::hi16)).getAsInteger
() == getSubRegIndexLaneMask(AMDGPU::sub0).getAsInteger() &&
 "getNumCoveredRegs() will not work with generated subreg masks!"
) ? void (0) : __assert_fail ("getSubRegIndexLaneMask(AMDGPU::sub0).getAsInteger() == 3 && getSubRegIndexLaneMask(AMDGPU::sub31).getAsInteger() == (3ULL << 62) && (getSubRegIndexLaneMask(AMDGPU::lo16) | getSubRegIndexLaneMask(AMDGPU::hi16)).getAsInteger() == getSubRegIndexLaneMask(AMDGPU::sub0).getAsInteger() && \"getNumCoveredRegs() will not work with generated subreg masks!\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 282, __extension__ __PRETTY_FUNCTION__))
       "getNumCoveredRegs() will not work with generated subreg masks!")(static_cast <bool> (getSubRegIndexLaneMask(AMDGPU::sub0
).getAsInteger() == 3 && getSubRegIndexLaneMask(AMDGPU
::sub31).getAsInteger() == (3ULL << 62) && (getSubRegIndexLaneMask
(AMDGPU::lo16) | getSubRegIndexLaneMask(AMDGPU::hi16)).getAsInteger
() == getSubRegIndexLaneMask(AMDGPU::sub0).getAsInteger() &&
 "getNumCoveredRegs() will not work with generated subreg masks!"
) ? void (0) : __assert_fail ("getSubRegIndexLaneMask(AMDGPU::sub0).getAsInteger() == 3 && getSubRegIndexLaneMask(AMDGPU::sub31).getAsInteger() == (3ULL << 62) && (getSubRegIndexLaneMask(AMDGPU::lo16) | getSubRegIndexLaneMask(AMDGPU::hi16)).getAsInteger() == getSubRegIndexLaneMask(AMDGPU::sub0).getAsInteger() && \"getNumCoveredRegs() will not work with generated subreg masks!\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 282, __extension__ __PRETTY_FUNCTION__));

RegPressureIgnoredUnits.resize(getNumRegUnits());
RegPressureIgnoredUnits.set(
    *MCRegUnitIterator(MCRegister::from(AMDGPU::M0), this));
for (auto Reg : AMDGPU::VGPR_HI16RegClass)
  RegPressureIgnoredUnits.set(*MCRegUnitIterator(Reg, this));

// HACK: Until this is fully tablegen'd.
static llvm::once_flag InitializeRegSplitPartsFlag;

static auto InitializeRegSplitPartsOnce = [this]() {
  for (unsigned Idx = 1, E = getNumSubRegIndices() - 1; Idx < E; ++Idx) {
    unsigned Size = getSubRegIdxSize(Idx);
    if (Size & 31)
      continue;
    std::vector<int16_t> &Vec = RegSplitParts[Size / 32 - 1];
    unsigned Pos = getSubRegIdxOffset(Idx);
    if (Pos % Size)
      continue;
    Pos /= Size;
    if (Vec.empty()) {
      unsigned MaxNumParts = 1024 / Size; // Maximum register is 1024 bits.
      Vec.resize(MaxNumParts);
    }
    Vec[Pos] = Idx;
  }
};

static llvm::once_flag InitializeSubRegFromChannelTableFlag;

static auto InitializeSubRegFromChannelTableOnce = [this]() {
  for (auto &Row : SubRegFromChannelTable)
    Row.fill(AMDGPU::NoSubRegister);
  for (uint16_t Idx = 1; Idx < getNumSubRegIndices(); ++Idx) {
    unsigned Width = AMDGPUSubRegIdxRanges[Idx].Size / 32;
    unsigned Offset = AMDGPUSubRegIdxRanges[Idx].Offset / 32;
    assert(Width < SubRegFromChannelTableWidthMap.size())(static_cast <bool> (Width < SubRegFromChannelTableWidthMap
.size()) ? void (0) : __assert_fail ("Width < SubRegFromChannelTableWidthMap.size()"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 319, __extension__ __PRETTY_FUNCTION__));
    Width = SubRegFromChannelTableWidthMap[Width];
    if (Width == 0)
      continue;
    unsigned TableIdx = Width - 1;
    assert(TableIdx < SubRegFromChannelTable.size())(static_cast <bool> (TableIdx < SubRegFromChannelTable
.size()) ? void (0) : __assert_fail ("TableIdx < SubRegFromChannelTable.size()"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 324, __extension__ __PRETTY_FUNCTION__));
    assert(Offset < SubRegFromChannelTable[TableIdx].size())(static_cast <bool> (Offset < SubRegFromChannelTable
[TableIdx].size()) ? void (0) : __assert_fail ("Offset < SubRegFromChannelTable[TableIdx].size()"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 325, __extension__ __PRETTY_FUNCTION__));
    SubRegFromChannelTable[TableIdx][Offset] = Idx;
  }
};

llvm::call_once(InitializeRegSplitPartsFlag, InitializeRegSplitPartsOnce);
llvm::call_once(InitializeSubRegFromChannelTableFlag,
                InitializeSubRegFromChannelTableOnce);
333}

335void SIRegisterInfo::reserveRegisterTuples(BitVector &Reserved,
                                         MCRegister Reg) const {
MCRegAliasIterator R(Reg, this, true);

for (; R.isValid(); ++R)
  Reserved.set(*R);
341}

343// Forced to be here by one .inc
344const MCPhysReg *SIRegisterInfo::getCalleeSavedRegs(
const MachineFunction *MF) const {
CallingConv::ID CC = MF->getFunction().getCallingConv();
switch (CC) {
case CallingConv::C:
case CallingConv::Fast:
case CallingConv::Cold:
case CallingConv::AMDGPU_Gfx:
  return MF->getSubtarget<GCNSubtarget>().hasGFX90AInsts()
      ? CSR_AMDGPU_HighRegs_With_AGPRs_SaveList
      : CSR_AMDGPU_HighRegs_SaveList;
default: {
  // Dummy to not crash RegisterClassInfo.
  static const MCPhysReg NoCalleeSavedReg = AMDGPU::NoRegister;
  return &NoCalleeSavedReg;
}
}
361}

363const MCPhysReg *
364SIRegisterInfo::getCalleeSavedRegsViaCopy(const MachineFunction *MF) const {
return nullptr;
366}

368const uint32_t *SIRegisterInfo::getCallPreservedMask(const MachineFunction &MF,
                                                   CallingConv::ID CC) const {
switch (CC) {
case CallingConv::C:
case CallingConv::Fast:
case CallingConv::Cold:
case CallingConv::AMDGPU_Gfx:
  return MF.getSubtarget<GCNSubtarget>().hasGFX90AInsts()
      ? CSR_AMDGPU_HighRegs_With_AGPRs_RegMask
      : CSR_AMDGPU_HighRegs_RegMask;
default:
  return nullptr;
}
381}

383const uint32_t *SIRegisterInfo::getNoPreservedMask() const {
return CSR_AMDGPU_NoRegs_RegMask;
385}

387Register SIRegisterInfo::getFrameRegister(const MachineFunction &MF) const {
const SIFrameLowering *TFI =
    MF.getSubtarget<GCNSubtarget>().getFrameLowering();
const SIMachineFunctionInfo *FuncInfo = MF.getInfo<SIMachineFunctionInfo>();
// During ISel lowering we always reserve the stack pointer in entry
// functions, but never actually want to reference it when accessing our own
// frame. If we need a frame pointer we use it, but otherwise we can just use
// an immediate "0" which we represent by returning NoRegister.
if (FuncInfo->isEntryFunction()) {
  return TFI->hasFP(MF) ? FuncInfo->getFrameOffsetReg() : Register();
}
return TFI->hasFP(MF) ? FuncInfo->getFrameOffsetReg()
                      : FuncInfo->getStackPtrOffsetReg();
400}

402bool SIRegisterInfo::hasBasePointer(const MachineFunction &MF) const {
// When we need stack realignment, we can't reference off of the
// stack pointer, so we reserve a base pointer.
const MachineFrameInfo &MFI = MF.getFrameInfo();
return MFI.getNumFixedObjects() && shouldRealignStack(MF);
407}

409Register SIRegisterInfo::getBaseRegister() const { return AMDGPU::SGPR34; }

411const uint32_t *SIRegisterInfo::getAllVGPRRegMask() const {
return CSR_AMDGPU_AllVGPRs_RegMask;
413}

415const uint32_t *SIRegisterInfo::getAllAGPRRegMask() const {
return CSR_AMDGPU_AllAGPRs_RegMask;
417}

419const uint32_t *SIRegisterInfo::getAllVectorRegMask() const {
return CSR_AMDGPU_AllVectorRegs_RegMask;
421}

423const uint32_t *SIRegisterInfo::getAllAllocatableSRegMask() const {
return CSR_AMDGPU_AllAllocatableSRegs_RegMask;
425}

427unsigned SIRegisterInfo::getSubRegFromChannel(unsigned Channel,
                                            unsigned NumRegs) {
assert(NumRegs < SubRegFromChannelTableWidthMap.size())(static_cast <bool> (NumRegs < SubRegFromChannelTableWidthMap
.size()) ? void (0) : __assert_fail ("NumRegs < SubRegFromChannelTableWidthMap.size()"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 429, __extension__ __PRETTY_FUNCTION__));
unsigned NumRegIndex = SubRegFromChannelTableWidthMap[NumRegs];
assert(NumRegIndex && "Not implemented")(static_cast <bool> (NumRegIndex && "Not implemented"
) ? void (0) : __assert_fail ("NumRegIndex && \"Not implemented\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 431, __extension__ __PRETTY_FUNCTION__));
assert(Channel < SubRegFromChannelTable[NumRegIndex - 1].size())(static_cast <bool> (Channel < SubRegFromChannelTable
[NumRegIndex - 1].size()) ? void (0) : __assert_fail ("Channel < SubRegFromChannelTable[NumRegIndex - 1].size()"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 432, __extension__ __PRETTY_FUNCTION__));
return SubRegFromChannelTable[NumRegIndex - 1][Channel];
434}

436MCRegister SIRegisterInfo::reservedPrivateSegmentBufferReg(
const MachineFunction &MF) const {
unsigned BaseIdx = alignDown(ST.getMaxNumSGPRs(MF), 4) - 4;
MCRegister BaseReg(AMDGPU::SGPR_32RegClass.getRegister(BaseIdx));
return getMatchingSuperReg(BaseReg, AMDGPU::sub0, &AMDGPU::SGPR_128RegClass);
441}

443BitVector SIRegisterInfo::getReservedRegs(const MachineFunction &MF) const {
BitVector Reserved(getNumRegs());
Reserved.set(AMDGPU::MODE);

// EXEC_LO and EXEC_HI could be allocated and used as regular register, but
// this seems likely to result in bugs, so I'm marking them as reserved.
reserveRegisterTuples(Reserved, AMDGPU::EXEC);
reserveRegisterTuples(Reserved, AMDGPU::FLAT_SCR);

// M0 has to be reserved so that llvm accepts it as a live-in into a block.
reserveRegisterTuples(Reserved, AMDGPU::M0);

// Reserve src_vccz, src_execz, src_scc.
reserveRegisterTuples(Reserved, AMDGPU::SRC_VCCZ);
reserveRegisterTuples(Reserved, AMDGPU::SRC_EXECZ);
reserveRegisterTuples(Reserved, AMDGPU::SRC_SCC);

// Reserve the memory aperture registers.
reserveRegisterTuples(Reserved, AMDGPU::SRC_SHARED_BASE);
reserveRegisterTuples(Reserved, AMDGPU::SRC_SHARED_LIMIT);
reserveRegisterTuples(Reserved, AMDGPU::SRC_PRIVATE_BASE);
reserveRegisterTuples(Reserved, AMDGPU::SRC_PRIVATE_LIMIT);

// Reserve src_pops_exiting_wave_id - support is not implemented in Codegen.
reserveRegisterTuples(Reserved, AMDGPU::SRC_POPS_EXITING_WAVE_ID);

// Reserve xnack_mask registers - support is not implemented in Codegen.
reserveRegisterTuples(Reserved, AMDGPU::XNACK_MASK);

// Reserve lds_direct register - support is not implemented in Codegen.
reserveRegisterTuples(Reserved, AMDGPU::LDS_DIRECT);

// Reserve Trap Handler registers - support is not implemented in Codegen.
reserveRegisterTuples(Reserved, AMDGPU::TBA);
reserveRegisterTuples(Reserved, AMDGPU::TMA);
reserveRegisterTuples(Reserved, AMDGPU::TTMP0_TTMP1);
reserveRegisterTuples(Reserved, AMDGPU::TTMP2_TTMP3);
reserveRegisterTuples(Reserved, AMDGPU::TTMP4_TTMP5);
reserveRegisterTuples(Reserved, AMDGPU::TTMP6_TTMP7);
reserveRegisterTuples(Reserved, AMDGPU::TTMP8_TTMP9);
reserveRegisterTuples(Reserved, AMDGPU::TTMP10_TTMP11);
reserveRegisterTuples(Reserved, AMDGPU::TTMP12_TTMP13);
reserveRegisterTuples(Reserved, AMDGPU::TTMP14_TTMP15);

// Reserve null register - it shall never be allocated
reserveRegisterTuples(Reserved, AMDGPU::SGPR_NULL);

// Disallow vcc_hi allocation in wave32. It may be allocated but most likely
// will result in bugs.
if (isWave32) {
  Reserved.set(AMDGPU::VCC);
  Reserved.set(AMDGPU::VCC_HI);
}

unsigned MaxNumSGPRs = ST.getMaxNumSGPRs(MF);
unsigned TotalNumSGPRs = AMDGPU::SGPR_32RegClass.getNumRegs();
for (unsigned i = MaxNumSGPRs; i < TotalNumSGPRs; ++i) {
  unsigned Reg = AMDGPU::SGPR_32RegClass.getRegister(i);
  reserveRegisterTuples(Reserved, Reg);
}

unsigned MaxNumVGPRs = ST.getMaxNumVGPRs(MF);
// TODO: In an entry function without calls and AGPRs used it is possible
//       to use the whole register budget for VGPRs. Even more it shall
//       be possible to estimate maximum AGPR/VGPR pressure and split
//       register file accordingly.
if (ST.hasGFX90AInsts())
  MaxNumVGPRs /= 2;
unsigned TotalNumVGPRs = AMDGPU::VGPR_32RegClass.getNumRegs();
for (unsigned i = MaxNumVGPRs; i < TotalNumVGPRs; ++i) {
  unsigned Reg = AMDGPU::VGPR_32RegClass.getRegister(i);
  reserveRegisterTuples(Reserved, Reg);
  Reg = AMDGPU::AGPR_32RegClass.getRegister(i);
  reserveRegisterTuples(Reserved, Reg);
}

for (auto Reg : AMDGPU::SReg_32RegClass) {
  Reserved.set(getSubReg(Reg, AMDGPU::hi16));
  Register Low = getSubReg(Reg, AMDGPU::lo16);
  // This is to prevent BB vcc liveness errors.
  if (!AMDGPU::SGPR_LO16RegClass.contains(Low))
    Reserved.set(Low);
}

for (auto Reg : AMDGPU::AGPR_32RegClass) {
  Reserved.set(getSubReg(Reg, AMDGPU::hi16));
}

// Reserve all the rest AGPRs if there are no instructions to use it.
if (!ST.hasMAIInsts()) {
  for (unsigned i = 0; i < MaxNumVGPRs; ++i) {
    unsigned Reg = AMDGPU::AGPR_32RegClass.getRegister(i);
    reserveRegisterTuples(Reserved, Reg);
  }
}

const SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();

Register ScratchRSrcReg = MFI->getScratchRSrcReg();
if (ScratchRSrcReg != AMDGPU::NoRegister) {
  // Reserve 4 SGPRs for the scratch buffer resource descriptor in case we need
  // to spill.
  // TODO: May need to reserve a VGPR if doing LDS spilling.
  reserveRegisterTuples(Reserved, ScratchRSrcReg);
}

// We have to assume the SP is needed in case there are calls in the function,
// which is detected after the function is lowered. If we aren't really going
// to need SP, don't bother reserving it.
MCRegister StackPtrReg = MFI->getStackPtrOffsetReg();

if (StackPtrReg) {
  reserveRegisterTuples(Reserved, StackPtrReg);
  assert(!isSubRegister(ScratchRSrcReg, StackPtrReg))(static_cast <bool> (!isSubRegister(ScratchRSrcReg, StackPtrReg
)) ? void (0) : __assert_fail ("!isSubRegister(ScratchRSrcReg, StackPtrReg)"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 556, __extension__ __PRETTY_FUNCTION__));
}

MCRegister FrameReg = MFI->getFrameOffsetReg();
if (FrameReg) {
  reserveRegisterTuples(Reserved, FrameReg);
  assert(!isSubRegister(ScratchRSrcReg, FrameReg))(static_cast <bool> (!isSubRegister(ScratchRSrcReg, FrameReg
)) ? void (0) : __assert_fail ("!isSubRegister(ScratchRSrcReg, FrameReg)"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 562, __extension__ __PRETTY_FUNCTION__));
}

if (hasBasePointer(MF)) {
  MCRegister BasePtrReg = getBaseRegister();
  reserveRegisterTuples(Reserved, BasePtrReg);
  assert(!isSubRegister(ScratchRSrcReg, BasePtrReg))(static_cast <bool> (!isSubRegister(ScratchRSrcReg, BasePtrReg
)) ? void (0) : __assert_fail ("!isSubRegister(ScratchRSrcReg, BasePtrReg)"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 568, __extension__ __PRETTY_FUNCTION__));
}

for (auto Reg : MFI->WWMReservedRegs) {
  reserveRegisterTuples(Reserved, Reg.first);
}

// Reserve VGPRs used for SGPR spilling.
// Note we treat freezeReservedRegs unusually because we run register
// allocation in two phases. It's OK to re-freeze with new registers for the
// second run.
579#if 0
for (auto &SpilledFI : MFI->sgpr_spill_vgprs()) {
  for (auto &SpilledVGPR : SpilledFI.second)
    reserveRegisterTuples(Reserved, SpilledVGPR.VGPR);
}
584#endif

// FIXME: Stop using reserved registers for this.
for (MCPhysReg Reg : MFI->getAGPRSpillVGPRs())
  reserveRegisterTuples(Reserved, Reg);

for (MCPhysReg Reg : MFI->getVGPRSpillAGPRs())
  reserveRegisterTuples(Reserved, Reg);

for (auto SSpill : MFI->getSGPRSpillVGPRs())
  reserveRegisterTuples(Reserved, SSpill.VGPR);

return Reserved;
597}

599bool SIRegisterInfo::shouldRealignStack(const MachineFunction &MF) const {
const SIMachineFunctionInfo *Info = MF.getInfo<SIMachineFunctionInfo>();
// On entry, the base address is 0, so it can't possibly need any more
// alignment.

// FIXME: Should be able to specify the entry frame alignment per calling
// convention instead.
if (Info->isEntryFunction())
  return false;

return TargetRegisterInfo::shouldRealignStack(MF);
610}

612bool SIRegisterInfo::requiresRegisterScavenging(const MachineFunction &Fn) const {
const SIMachineFunctionInfo *Info = Fn.getInfo<SIMachineFunctionInfo>();
if (Info->isEntryFunction()) {
  const MachineFrameInfo &MFI = Fn.getFrameInfo();
  return MFI.hasStackObjects() || MFI.hasCalls();
}

// May need scavenger for dealing with callee saved registers.
return true;
621}

623bool SIRegisterInfo::requiresFrameIndexScavenging(
const MachineFunction &MF) const {
// Do not use frame virtual registers. They used to be used for SGPRs, but
// once we reach PrologEpilogInserter, we can no longer spill SGPRs. If the
// scavenger fails, we can increment/decrement the necessary SGPRs to avoid a
// spill.
return false;
630}

632bool SIRegisterInfo::requiresFrameIndexReplacementScavenging(
const MachineFunction &MF) const {
const MachineFrameInfo &MFI = MF.getFrameInfo();
return MFI.hasStackObjects();
636}

638bool SIRegisterInfo::requiresVirtualBaseRegisters(
const MachineFunction &) const {
// There are no special dedicated stack or frame pointers.
return true;
642}

644int64_t SIRegisterInfo::getScratchInstrOffset(const MachineInstr *MI) const {
assert(SIInstrInfo::isMUBUF(*MI) || SIInstrInfo::isFLATScratch(*MI))(static_cast <bool> (SIInstrInfo::isMUBUF(*MI) || SIInstrInfo
::isFLATScratch(*MI)) ? void (0) : __assert_fail ("SIInstrInfo::isMUBUF(*MI) || SIInstrInfo::isFLATScratch(*MI)"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 645, __extension__ __PRETTY_FUNCTION__));

int OffIdx = AMDGPU::getNamedOperandIdx(MI->getOpcode(),
                                        AMDGPU::OpName::offset);
return MI->getOperand(OffIdx).getImm();
650}

652int64_t SIRegisterInfo::getFrameIndexInstrOffset(const MachineInstr *MI,
                                               int Idx) const {
if (!SIInstrInfo::isMUBUF(*MI) && !SIInstrInfo::isFLATScratch(*MI))
  return 0;

assert((Idx == AMDGPU::getNamedOperandIdx(MI->getOpcode(),(static_cast <bool> ((Idx == AMDGPU::getNamedOperandIdx
(MI->getOpcode(), AMDGPU::OpName::vaddr) || (Idx == AMDGPU
::getNamedOperandIdx(MI->getOpcode(), AMDGPU::OpName::saddr
))) && "Should never see frame index on non-address operand"
) ? void (0) : __assert_fail ("(Idx == AMDGPU::getNamedOperandIdx(MI->getOpcode(), AMDGPU::OpName::vaddr) || (Idx == AMDGPU::getNamedOperandIdx(MI->getOpcode(), AMDGPU::OpName::saddr))) && \"Should never see frame index on non-address operand\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 661, __extension__ __PRETTY_FUNCTION__))
                                          AMDGPU::OpName::vaddr) ||(static_cast <bool> ((Idx == AMDGPU::getNamedOperandIdx
(MI->getOpcode(), AMDGPU::OpName::vaddr) || (Idx == AMDGPU
::getNamedOperandIdx(MI->getOpcode(), AMDGPU::OpName::saddr
))) && "Should never see frame index on non-address operand"
) ? void (0) : __assert_fail ("(Idx == AMDGPU::getNamedOperandIdx(MI->getOpcode(), AMDGPU::OpName::vaddr) || (Idx == AMDGPU::getNamedOperandIdx(MI->getOpcode(), AMDGPU::OpName::saddr))) && \"Should never see frame index on non-address operand\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 661, __extension__ __PRETTY_FUNCTION__))
       (Idx == AMDGPU::getNamedOperandIdx(MI->getOpcode(),(static_cast <bool> ((Idx == AMDGPU::getNamedOperandIdx
(MI->getOpcode(), AMDGPU::OpName::vaddr) || (Idx == AMDGPU
::getNamedOperandIdx(MI->getOpcode(), AMDGPU::OpName::saddr
))) && "Should never see frame index on non-address operand"
) ? void (0) : __assert_fail ("(Idx == AMDGPU::getNamedOperandIdx(MI->getOpcode(), AMDGPU::OpName::vaddr) || (Idx == AMDGPU::getNamedOperandIdx(MI->getOpcode(), AMDGPU::OpName::saddr))) && \"Should never see frame index on non-address operand\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 661, __extension__ __PRETTY_FUNCTION__))
                                          AMDGPU::OpName::saddr))) &&(static_cast <bool> ((Idx == AMDGPU::getNamedOperandIdx
(MI->getOpcode(), AMDGPU::OpName::vaddr) || (Idx == AMDGPU
::getNamedOperandIdx(MI->getOpcode(), AMDGPU::OpName::saddr
))) && "Should never see frame index on non-address operand"
) ? void (0) : __assert_fail ("(Idx == AMDGPU::getNamedOperandIdx(MI->getOpcode(), AMDGPU::OpName::vaddr) || (Idx == AMDGPU::getNamedOperandIdx(MI->getOpcode(), AMDGPU::OpName::saddr))) && \"Should never see frame index on non-address operand\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 661, __extension__ __PRETTY_FUNCTION__))
       "Should never see frame index on non-address operand")(static_cast <bool> ((Idx == AMDGPU::getNamedOperandIdx
(MI->getOpcode(), AMDGPU::OpName::vaddr) || (Idx == AMDGPU
::getNamedOperandIdx(MI->getOpcode(), AMDGPU::OpName::saddr
))) && "Should never see frame index on non-address operand"
) ? void (0) : __assert_fail ("(Idx == AMDGPU::getNamedOperandIdx(MI->getOpcode(), AMDGPU::OpName::vaddr) || (Idx == AMDGPU::getNamedOperandIdx(MI->getOpcode(), AMDGPU::OpName::saddr))) && \"Should never see frame index on non-address operand\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 661, __extension__ __PRETTY_FUNCTION__));

return getScratchInstrOffset(MI);
664}

666bool SIRegisterInfo::needsFrameBaseReg(MachineInstr *MI, int64_t Offset) const {
if (!SIInstrInfo::isMUBUF(*MI) && !SIInstrInfo::isFLATScratch(*MI))
  return false;

int64_t FullOffset = Offset + getScratchInstrOffset(MI);

if (SIInstrInfo::isMUBUF(*MI))
  return !SIInstrInfo::isLegalMUBUFImmOffset(FullOffset);

const SIInstrInfo *TII = ST.getInstrInfo();
return !TII->isLegalFLATOffset(FullOffset, AMDGPUAS::PRIVATE_ADDRESS,
                               SIInstrFlags::FlatScratch);
678}

680Register SIRegisterInfo::materializeFrameBaseRegister(MachineBasicBlock *MBB,
                                                    int FrameIdx,
                                                    int64_t Offset) const {
MachineBasicBlock::iterator Ins = MBB->begin();
DebugLoc DL; // Defaults to "unknown"

if (Ins != MBB->end())
  DL = Ins->getDebugLoc();

MachineFunction *MF = MBB->getParent();
const SIInstrInfo *TII = ST.getInstrInfo();
MachineRegisterInfo &MRI = MF->getRegInfo();
unsigned MovOpc = ST.enableFlatScratch() ? AMDGPU::S_MOV_B32
                                         : AMDGPU::V_MOV_B32_e32;

Register BaseReg = MRI.createVirtualRegister(
    ST.enableFlatScratch() ? &AMDGPU::SReg_32_XEXEC_HIRegClass
                           : &AMDGPU::VGPR_32RegClass);

if (Offset == 0) {
  BuildMI(*MBB, Ins, DL, TII->get(MovOpc), BaseReg)
    .addFrameIndex(FrameIdx);
  return BaseReg;
}

Register OffsetReg = MRI.createVirtualRegister(&AMDGPU::SReg_32_XM0RegClass);

Register FIReg = MRI.createVirtualRegister(
    ST.enableFlatScratch() ? &AMDGPU::SReg_32_XM0RegClass
                           : &AMDGPU::VGPR_32RegClass);

BuildMI(*MBB, Ins, DL, TII->get(AMDGPU::S_MOV_B32), OffsetReg)
  .addImm(Offset);
BuildMI(*MBB, Ins, DL, TII->get(MovOpc), FIReg)
  .addFrameIndex(FrameIdx);

if (ST.enableFlatScratch() ) {
  BuildMI(*MBB, Ins, DL, TII->get(AMDGPU::S_ADD_I32), BaseReg)
      .addReg(OffsetReg, RegState::Kill)
      .addReg(FIReg);
  return BaseReg;
}

TII->getAddNoCarry(*MBB, Ins, DL, BaseReg)
  .addReg(OffsetReg, RegState::Kill)
  .addReg(FIReg)
  .addImm(0); // clamp bit

return BaseReg;
729}

731void SIRegisterInfo::resolveFrameIndex(MachineInstr &MI, Register BaseReg,
                                     int64_t Offset) const {
const SIInstrInfo *TII = ST.getInstrInfo();
bool IsFlat = TII->isFLATScratch(MI);

736#ifndef NDEBUG
// FIXME: Is it possible to be storing a frame index to itself?
bool SeenFI = false;
for (const MachineOperand &MO: MI.operands()) {
  if (MO.isFI()) {
    if (SeenFI)
      llvm_unreachable("should not see multiple frame indices")::llvm::llvm_unreachable_internal("should not see multiple frame indices"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 742);

    SeenFI = true;
  }
}
747#endif

MachineOperand *FIOp =
    TII->getNamedOperand(MI, IsFlat ? AMDGPU::OpName::saddr
                                    : AMDGPU::OpName::vaddr);

MachineOperand *OffsetOp = TII->getNamedOperand(MI, AMDGPU::OpName::offset);
int64_t NewOffset = OffsetOp->getImm() + Offset;

assert(FIOp && FIOp->isFI() && "frame index must be address operand")(static_cast <bool> (FIOp && FIOp->isFI() &&
 "frame index must be address operand") ? void (0) : __assert_fail
 ("FIOp && FIOp->isFI() && \"frame index must be address operand\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 756, __extension__ __PRETTY_FUNCTION__));
assert(TII->isMUBUF(MI) || TII->isFLATScratch(MI))(static_cast <bool> (TII->isMUBUF(MI) || TII->isFLATScratch
(MI)) ? void (0) : __assert_fail ("TII->isMUBUF(MI) || TII->isFLATScratch(MI)"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 757, __extension__ __PRETTY_FUNCTION__));

if (IsFlat) {
  assert(TII->isLegalFLATOffset(NewOffset, AMDGPUAS::PRIVATE_ADDRESS,(static_cast <bool> (TII->isLegalFLATOffset(NewOffset
, AMDGPUAS::PRIVATE_ADDRESS, SIInstrFlags::FlatScratch) &&
 "offset should be legal") ? void (0) : __assert_fail ("TII->isLegalFLATOffset(NewOffset, AMDGPUAS::PRIVATE_ADDRESS, SIInstrFlags::FlatScratch) && \"offset should be legal\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 762, __extension__ __PRETTY_FUNCTION__))
                                SIInstrFlags::FlatScratch) &&(static_cast <bool> (TII->isLegalFLATOffset(NewOffset
, AMDGPUAS::PRIVATE_ADDRESS, SIInstrFlags::FlatScratch) &&
 "offset should be legal") ? void (0) : __assert_fail ("TII->isLegalFLATOffset(NewOffset, AMDGPUAS::PRIVATE_ADDRESS, SIInstrFlags::FlatScratch) && \"offset should be legal\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 762, __extension__ __PRETTY_FUNCTION__))
         "offset should be legal")(static_cast <bool> (TII->isLegalFLATOffset(NewOffset
, AMDGPUAS::PRIVATE_ADDRESS, SIInstrFlags::FlatScratch) &&
 "offset should be legal") ? void (0) : __assert_fail ("TII->isLegalFLATOffset(NewOffset, AMDGPUAS::PRIVATE_ADDRESS, SIInstrFlags::FlatScratch) && \"offset should be legal\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 762, __extension__ __PRETTY_FUNCTION__));
  FIOp->ChangeToRegister(BaseReg, false);
  OffsetOp->setImm(NewOffset);
  return;
}

768#ifndef NDEBUG
MachineOperand *SOffset = TII->getNamedOperand(MI, AMDGPU::OpName::soffset);
assert(SOffset->isImm() && SOffset->getImm() == 0)(static_cast <bool> (SOffset->isImm() && SOffset
->getImm() == 0) ? void (0) : __assert_fail ("SOffset->isImm() && SOffset->getImm() == 0"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 770, __extension__ __PRETTY_FUNCTION__));
771#endif

assert(SIInstrInfo::isLegalMUBUFImmOffset(NewOffset) &&(static_cast <bool> (SIInstrInfo::isLegalMUBUFImmOffset
(NewOffset) && "offset should be legal") ? void (0) :
 __assert_fail ("SIInstrInfo::isLegalMUBUFImmOffset(NewOffset) && \"offset should be legal\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 774, __extension__ __PRETTY_FUNCTION__))
       "offset should be legal")(static_cast <bool> (SIInstrInfo::isLegalMUBUFImmOffset
(NewOffset) && "offset should be legal") ? void (0) :
 __assert_fail ("SIInstrInfo::isLegalMUBUFImmOffset(NewOffset) && \"offset should be legal\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 774, __extension__ __PRETTY_FUNCTION__));

FIOp->ChangeToRegister(BaseReg, false);
OffsetOp->setImm(NewOffset);
778}

780bool SIRegisterInfo::isFrameOffsetLegal(const MachineInstr *MI,
                                      Register BaseReg,
                                      int64_t Offset) const {
if (!SIInstrInfo::isMUBUF(*MI) && !SIInstrInfo::isFLATScratch(*MI))
  return false;

int64_t NewOffset = Offset + getScratchInstrOffset(MI);

if (SIInstrInfo::isMUBUF(*MI))
  return SIInstrInfo::isLegalMUBUFImmOffset(NewOffset);

const SIInstrInfo *TII = ST.getInstrInfo();
return TII->isLegalFLATOffset(NewOffset, AMDGPUAS::PRIVATE_ADDRESS,
                              SIInstrFlags::FlatScratch);
794}

796const TargetRegisterClass *SIRegisterInfo::getPointerRegClass(
const MachineFunction &MF, unsigned Kind) const {
// This is inaccurate. It depends on the instruction and address space. The
// only place where we should hit this is for dealing with frame indexes /
// private accesses, so this is correct in that case.
return &AMDGPU::VGPR_32RegClass;
802}

804static unsigned getNumSubRegsForSpillOp(unsigned Op) {

switch (Op) {
case AMDGPU::SI_SPILL_S1024_SAVE:
case AMDGPU::SI_SPILL_S1024_RESTORE:
case AMDGPU::SI_SPILL_V1024_SAVE:
case AMDGPU::SI_SPILL_V1024_RESTORE:
case AMDGPU::SI_SPILL_A1024_SAVE:
case AMDGPU::SI_SPILL_A1024_RESTORE:
  return 32;
case AMDGPU::SI_SPILL_S512_SAVE:
case AMDGPU::SI_SPILL_S512_RESTORE:
case AMDGPU::SI_SPILL_V512_SAVE:
case AMDGPU::SI_SPILL_V512_RESTORE:
case AMDGPU::SI_SPILL_A512_SAVE:
case AMDGPU::SI_SPILL_A512_RESTORE:
  return 16;
case AMDGPU::SI_SPILL_S256_SAVE:
case AMDGPU::SI_SPILL_S256_RESTORE:
case AMDGPU::SI_SPILL_V256_SAVE:
case AMDGPU::SI_SPILL_V256_RESTORE:
case AMDGPU::SI_SPILL_A256_SAVE:
case AMDGPU::SI_SPILL_A256_RESTORE:
  return 8;
case AMDGPU::SI_SPILL_S224_SAVE:
case AMDGPU::SI_SPILL_S224_RESTORE:
case AMDGPU::SI_SPILL_V224_SAVE:
case AMDGPU::SI_SPILL_V224_RESTORE:
case AMDGPU::SI_SPILL_A224_SAVE:
case AMDGPU::SI_SPILL_A224_RESTORE:
  return 7;
case AMDGPU::SI_SPILL_S192_SAVE:
case AMDGPU::SI_SPILL_S192_RESTORE:
case AMDGPU::SI_SPILL_V192_SAVE:
case AMDGPU::SI_SPILL_V192_RESTORE:
case AMDGPU::SI_SPILL_A192_SAVE:
case AMDGPU::SI_SPILL_A192_RESTORE:
  return 6;
case AMDGPU::SI_SPILL_S160_SAVE:
case AMDGPU::SI_SPILL_S160_RESTORE:
case AMDGPU::SI_SPILL_V160_SAVE:
case AMDGPU::SI_SPILL_V160_RESTORE:
case AMDGPU::SI_SPILL_A160_SAVE:
case AMDGPU::SI_SPILL_A160_RESTORE:
  return 5;
case AMDGPU::SI_SPILL_S128_SAVE:
case AMDGPU::SI_SPILL_S128_RESTORE:
case AMDGPU::SI_SPILL_V128_SAVE:
case AMDGPU::SI_SPILL_V128_RESTORE:
case AMDGPU::SI_SPILL_A128_SAVE:
case AMDGPU::SI_SPILL_A128_RESTORE:
  return 4;
case AMDGPU::SI_SPILL_S96_SAVE:
case AMDGPU::SI_SPILL_S96_RESTORE:
case AMDGPU::SI_SPILL_V96_SAVE:
case AMDGPU::SI_SPILL_V96_RESTORE:
case AMDGPU::SI_SPILL_A96_SAVE:
case AMDGPU::SI_SPILL_A96_RESTORE:
  return 3;
case AMDGPU::SI_SPILL_S64_SAVE:
case AMDGPU::SI_SPILL_S64_RESTORE:
case AMDGPU::SI_SPILL_V64_SAVE:
case AMDGPU::SI_SPILL_V64_RESTORE:
case AMDGPU::SI_SPILL_A64_SAVE:
case AMDGPU::SI_SPILL_A64_RESTORE:
  return 2;
case AMDGPU::SI_SPILL_S32_SAVE:
case AMDGPU::SI_SPILL_S32_RESTORE:
case AMDGPU::SI_SPILL_V32_SAVE:
case AMDGPU::SI_SPILL_V32_RESTORE:
case AMDGPU::SI_SPILL_A32_SAVE:
case AMDGPU::SI_SPILL_A32_RESTORE:
  return 1;
default: llvm_unreachable("Invalid spill opcode")::llvm::llvm_unreachable_internal("Invalid spill opcode", "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 877);
}
879}

881static int getOffsetMUBUFStore(unsigned Opc) {
switch (Opc) {
case AMDGPU::BUFFER_STORE_DWORD_OFFEN:
  return AMDGPU::BUFFER_STORE_DWORD_OFFSET;
case AMDGPU::BUFFER_STORE_BYTE_OFFEN:
  return AMDGPU::BUFFER_STORE_BYTE_OFFSET;
case AMDGPU::BUFFER_STORE_SHORT_OFFEN:
  return AMDGPU::BUFFER_STORE_SHORT_OFFSET;
case AMDGPU::BUFFER_STORE_DWORDX2_OFFEN:
  return AMDGPU::BUFFER_STORE_DWORDX2_OFFSET;
case AMDGPU::BUFFER_STORE_DWORDX4_OFFEN:
  return AMDGPU::BUFFER_STORE_DWORDX4_OFFSET;
case AMDGPU::BUFFER_STORE_SHORT_D16_HI_OFFEN:
  return AMDGPU::BUFFER_STORE_SHORT_D16_HI_OFFSET;
case AMDGPU::BUFFER_STORE_BYTE_D16_HI_OFFEN:
  return AMDGPU::BUFFER_STORE_BYTE_D16_HI_OFFSET;
default:
  return -1;
}
900}

902static int getOffsetMUBUFLoad(unsigned Opc) {
switch (Opc) {
case AMDGPU::BUFFER_LOAD_DWORD_OFFEN:
  return AMDGPU::BUFFER_LOAD_DWORD_OFFSET;
case AMDGPU::BUFFER_LOAD_UBYTE_OFFEN:
  return AMDGPU::BUFFER_LOAD_UBYTE_OFFSET;
case AMDGPU::BUFFER_LOAD_SBYTE_OFFEN:
  return AMDGPU::BUFFER_LOAD_SBYTE_OFFSET;
case AMDGPU::BUFFER_LOAD_USHORT_OFFEN:
  return AMDGPU::BUFFER_LOAD_USHORT_OFFSET;
case AMDGPU::BUFFER_LOAD_SSHORT_OFFEN:
  return AMDGPU::BUFFER_LOAD_SSHORT_OFFSET;
case AMDGPU::BUFFER_LOAD_DWORDX2_OFFEN:
  return AMDGPU::BUFFER_LOAD_DWORDX2_OFFSET;
case AMDGPU::BUFFER_LOAD_DWORDX4_OFFEN:
  return AMDGPU::BUFFER_LOAD_DWORDX4_OFFSET;
case AMDGPU::BUFFER_LOAD_UBYTE_D16_OFFEN:
  return AMDGPU::BUFFER_LOAD_UBYTE_D16_OFFSET;
case AMDGPU::BUFFER_LOAD_UBYTE_D16_HI_OFFEN:
  return AMDGPU::BUFFER_LOAD_UBYTE_D16_HI_OFFSET;
case AMDGPU::BUFFER_LOAD_SBYTE_D16_OFFEN:
  return AMDGPU::BUFFER_LOAD_SBYTE_D16_OFFSET;
case AMDGPU::BUFFER_LOAD_SBYTE_D16_HI_OFFEN:
  return AMDGPU::BUFFER_LOAD_SBYTE_D16_HI_OFFSET;
case AMDGPU::BUFFER_LOAD_SHORT_D16_OFFEN:
  return AMDGPU::BUFFER_LOAD_SHORT_D16_OFFSET;
case AMDGPU::BUFFER_LOAD_SHORT_D16_HI_OFFEN:
  return AMDGPU::BUFFER_LOAD_SHORT_D16_HI_OFFSET;
default:
  return -1;
}
933}

935static MachineInstrBuilder spillVGPRtoAGPR(const GCNSubtarget &ST,
                                         MachineBasicBlock &MBB,
                                         MachineBasicBlock::iterator MI,
                                         int Index, unsigned Lane,
                                         unsigned ValueReg, bool IsKill) {
MachineFunction *MF = MBB.getParent();
SIMachineFunctionInfo *MFI = MF->getInfo<SIMachineFunctionInfo>();
const SIInstrInfo *TII = ST.getInstrInfo();

MCPhysReg Reg = MFI->getVGPRToAGPRSpill(Index, Lane);

if (Reg == AMDGPU::NoRegister)
  return MachineInstrBuilder();

bool IsStore = MI->mayStore();
MachineRegisterInfo &MRI = MF->getRegInfo();
auto *TRI = static_cast<const SIRegisterInfo*>(MRI.getTargetRegisterInfo());

unsigned Dst = IsStore ? Reg : ValueReg;
unsigned Src = IsStore ? ValueReg : Reg;
unsigned Opc = (IsStore ^ TRI->isVGPR(MRI, Reg)) ? AMDGPU::V_ACCVGPR_WRITE_B32_e64
                                                 : AMDGPU::V_ACCVGPR_READ_B32_e64;

auto MIB = BuildMI(MBB, MI, MI->getDebugLoc(), TII->get(Opc), Dst)
               .addReg(Src, getKillRegState(IsKill));
MIB->setAsmPrinterFlag(MachineInstr::ReloadReuse);
return MIB;
962}

964// This differs from buildSpillLoadStore by only scavenging a VGPR. It does not
965// need to handle the case where an SGPR may need to be spilled while spilling.
966static bool buildMUBUFOffsetLoadStore(const GCNSubtarget &ST,
                                    MachineFrameInfo &MFI,
                                    MachineBasicBlock::iterator MI,
                                    int Index,
                                    int64_t Offset) {
const SIInstrInfo *TII = ST.getInstrInfo();
MachineBasicBlock *MBB = MI->getParent();
const DebugLoc &DL = MI->getDebugLoc();
bool IsStore = MI->mayStore();

unsigned Opc = MI->getOpcode();
int LoadStoreOp = IsStore ?
  getOffsetMUBUFStore(Opc) : getOffsetMUBUFLoad(Opc);
if (LoadStoreOp == -1)
  return false;

const MachineOperand *Reg = TII->getNamedOperand(*MI, AMDGPU::OpName::vdata);
if (spillVGPRtoAGPR(ST, *MBB, MI, Index, 0, Reg->getReg(), false).getInstr())
  return true;

MachineInstrBuilder NewMI =
    BuildMI(*MBB, MI, DL, TII->get(LoadStoreOp))
        .add(*Reg)
        .add(*TII->getNamedOperand(*MI, AMDGPU::OpName::srsrc))
        .add(*TII->getNamedOperand(*MI, AMDGPU::OpName::soffset))
        .addImm(Offset)
        .addImm(0) // cpol
        .addImm(0) // tfe
        .addImm(0) // swz
        .cloneMemRefs(*MI);

const MachineOperand *VDataIn = TII->getNamedOperand(*MI,
                                                     AMDGPU::OpName::vdata_in);
if (VDataIn)
  NewMI.add(*VDataIn);
return true;
1002}

1004static unsigned getFlatScratchSpillOpcode(const SIInstrInfo *TII,
                                        unsigned LoadStoreOp,
                                        unsigned EltSize) {
bool IsStore = TII->get(LoadStoreOp).mayStore();
bool UseST =
  AMDGPU::getNamedOperandIdx(LoadStoreOp, AMDGPU::OpName::vaddr) < 0 &&
  AMDGPU::getNamedOperandIdx(LoadStoreOp, AMDGPU::OpName::saddr) < 0;

switch (EltSize) {
case 4:
  LoadStoreOp = IsStore ? AMDGPU::SCRATCH_STORE_DWORD_SADDR
                        : AMDGPU::SCRATCH_LOAD_DWORD_SADDR;
  break;
case 8:
  LoadStoreOp = IsStore ? AMDGPU::SCRATCH_STORE_DWORDX2_SADDR
                        : AMDGPU::SCRATCH_LOAD_DWORDX2_SADDR;
  break;
case 12:
  LoadStoreOp = IsStore ? AMDGPU::SCRATCH_STORE_DWORDX3_SADDR
                        : AMDGPU::SCRATCH_LOAD_DWORDX3_SADDR;
  break;
case 16:
  LoadStoreOp = IsStore ? AMDGPU::SCRATCH_STORE_DWORDX4_SADDR
                        : AMDGPU::SCRATCH_LOAD_DWORDX4_SADDR;
  break;
default:
  llvm_unreachable("Unexpected spill load/store size!")::llvm::llvm_unreachable_internal("Unexpected spill load/store size!"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1030);
}

if (UseST)
  LoadStoreOp = AMDGPU::getFlatScratchInstSTfromSS(LoadStoreOp);

return LoadStoreOp;
1037}

1039void SIRegisterInfo::buildSpillLoadStore(
  MachineBasicBlock &MBB, MachineBasicBlock::iterator MI,
  unsigned LoadStoreOp, int Index, Register ValueReg, bool IsKill,
  MCRegister ScratchOffsetReg, int64_t InstOffset, MachineMemOperand *MMO,
  RegScavenger *RS, LivePhysRegs *LiveRegs) const {
assert((!RS || !LiveRegs) && "Only RS or LiveRegs can be set but not both")(static_cast <bool> ((!RS || !LiveRegs) && "Only RS or LiveRegs can be set but not both"
) ? void (0) : __assert_fail ("(!RS || !LiveRegs) && \"Only RS or LiveRegs can be set but not both\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1044, __extension__ __PRETTY_FUNCTION__));

MachineFunction *MF = MBB.getParent();
const SIInstrInfo *TII = ST.getInstrInfo();
const MachineFrameInfo &MFI = MF->getFrameInfo();
const SIMachineFunctionInfo *FuncInfo = MF->getInfo<SIMachineFunctionInfo>();

const MCInstrDesc *Desc = &TII->get(LoadStoreOp);
const DebugLoc &DL = MI != MBB.end() ? MI->getDebugLoc() : DebugLoc();
bool IsStore = Desc->mayStore();
bool IsFlat = TII->isFLATScratch(LoadStoreOp);

bool Scavenged = false;
MCRegister SOffset = ScratchOffsetReg;

const TargetRegisterClass *RC = getRegClassForReg(MF->getRegInfo(), ValueReg);
// On gfx90a+ AGPR is a regular VGPR acceptable for loads and stores.
const bool IsAGPR = !ST.hasGFX90AInsts() && hasAGPRs(RC);
const unsigned RegWidth = AMDGPU::getRegBitWidth(RC->getID()) / 8;

// Always use 4 byte operations for AGPRs because we need to scavenge
// a temporary VGPR.
unsigned EltSize = (IsFlat && !IsAGPR) ? std::min(RegWidth, 16u) : 4u;
unsigned NumSubRegs = RegWidth / EltSize;
unsigned Size = NumSubRegs * EltSize;
unsigned RemSize = RegWidth - Size;
unsigned NumRemSubRegs = RemSize ? 1 : 0;
int64_t Offset = InstOffset + MFI.getObjectOffset(Index);
int64_t MaxOffset = Offset + Size + RemSize - EltSize;
int64_t ScratchOffsetRegDelta = 0;

if (IsFlat && EltSize > 4) {
  LoadStoreOp = getFlatScratchSpillOpcode(TII, LoadStoreOp, EltSize);
  Desc = &TII->get(LoadStoreOp);
}

Align Alignment = MFI.getObjectAlign(Index);
const MachinePointerInfo &BasePtrInfo = MMO->getPointerInfo();

assert((IsFlat || ((Offset % EltSize) == 0)) &&(static_cast <bool> ((IsFlat || ((Offset % EltSize) == 0
)) && "unexpected VGPR spill offset") ? void (0) : __assert_fail
 ("(IsFlat || ((Offset % EltSize) == 0)) && \"unexpected VGPR spill offset\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1084, __extension__ __PRETTY_FUNCTION__))
       "unexpected VGPR spill offset")(static_cast <bool> ((IsFlat || ((Offset % EltSize) == 0
)) && "unexpected VGPR spill offset") ? void (0) : __assert_fail
 ("(IsFlat || ((Offset % EltSize) == 0)) && \"unexpected VGPR spill offset\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1084, __extension__ __PRETTY_FUNCTION__));

bool IsOffsetLegal =
    IsFlat ? TII->isLegalFLATOffset(MaxOffset, AMDGPUAS::PRIVATE_ADDRESS,
                                    SIInstrFlags::FlatScratch)
           : SIInstrInfo::isLegalMUBUFImmOffset(MaxOffset);
if (!IsOffsetLegal || (IsFlat && !SOffset && !ST.hasFlatScratchSTMode())) {
  SOffset = MCRegister();

  // We currently only support spilling VGPRs to EltSize boundaries, meaning
  // we can simplify the adjustment of Offset here to just scale with
  // WavefrontSize.
  if (!IsFlat)
    Offset *= ST.getWavefrontSize();

  // We don't have access to the register scavenger if this function is called
  // during  PEI::scavengeFrameVirtualRegs() so use LiveRegs in this case.
  if (RS) {
    SOffset = RS->scavengeRegister(&AMDGPU::SGPR_32RegClass, MI, 0, false);
  } else if (LiveRegs) {
    for (MCRegister Reg : AMDGPU::SGPR_32RegClass) {
      if (LiveRegs->available(MF->getRegInfo(), Reg)) {
        SOffset = Reg;
        break;
      }
    }
  }

  if (!SOffset) {
    // There are no free SGPRs, and since we are in the process of spilling
    // VGPRs too.  Since we need a VGPR in order to spill SGPRs (this is true
    // on SI/CI and on VI it is true until we implement spilling using scalar
    // stores), we have no way to free up an SGPR.  Our solution here is to
    // add the offset directly to the ScratchOffset or StackPtrOffset
    // register, and then subtract the offset after the spill to return the
    // register to it's original value.
    if (!ScratchOffsetReg)
      ScratchOffsetReg = FuncInfo->getStackPtrOffsetReg();
    SOffset = ScratchOffsetReg;
    ScratchOffsetRegDelta = Offset;
  } else {
    Scavenged = true;
  }

  if (!SOffset)
    report_fatal_error("could not scavenge SGPR to spill in entry function");

  if (ScratchOffsetReg == AMDGPU::NoRegister) {
    BuildMI(MBB, MI, DL, TII->get(AMDGPU::S_MOV_B32), SOffset).addImm(Offset);
  } else {
    BuildMI(MBB, MI, DL, TII->get(AMDGPU::S_ADD_I32), SOffset)
        .addReg(ScratchOffsetReg)
        .addImm(Offset);
  }

  Offset = 0;
}

if (IsFlat && SOffset == AMDGPU::NoRegister) {
  assert(AMDGPU::getNamedOperandIdx(LoadStoreOp, AMDGPU::OpName::vaddr) < 0(static_cast <bool> (AMDGPU::getNamedOperandIdx(LoadStoreOp
, AMDGPU::OpName::vaddr) < 0 && "Unexpected vaddr for flat scratch with a FI operand"
) ? void (0) : __assert_fail ("AMDGPU::getNamedOperandIdx(LoadStoreOp, AMDGPU::OpName::vaddr) < 0 && \"Unexpected vaddr for flat scratch with a FI operand\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1144, __extension__ __PRETTY_FUNCTION__))
         && "Unexpected vaddr for flat scratch with a FI operand")(static_cast <bool> (AMDGPU::getNamedOperandIdx(LoadStoreOp
, AMDGPU::OpName::vaddr) < 0 && "Unexpected vaddr for flat scratch with a FI operand"
) ? void (0) : __assert_fail ("AMDGPU::getNamedOperandIdx(LoadStoreOp, AMDGPU::OpName::vaddr) < 0 && \"Unexpected vaddr for flat scratch with a FI operand\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1144, __extension__ __PRETTY_FUNCTION__));

  assert(ST.hasFlatScratchSTMode())(static_cast <bool> (ST.hasFlatScratchSTMode()) ? void (
0) : __assert_fail ("ST.hasFlatScratchSTMode()", "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1146, __extension__ __PRETTY_FUNCTION__));
  LoadStoreOp = AMDGPU::getFlatScratchInstSTfromSS(LoadStoreOp);
  Desc = &TII->get(LoadStoreOp);
}

Register TmpReg;

for (unsigned i = 0, e = NumSubRegs + NumRemSubRegs, RegOffset = 0; i != e;
     ++i, RegOffset += EltSize) {
  if (i == NumSubRegs) {
    EltSize = RemSize;
    LoadStoreOp = getFlatScratchSpillOpcode(TII, LoadStoreOp, EltSize);
  }
  Desc = &TII->get(LoadStoreOp);

  unsigned NumRegs = EltSize / 4;
  Register SubReg = e == 1
          ? ValueReg
          : Register(getSubReg(ValueReg,
                               getSubRegFromChannel(RegOffset / 4, NumRegs)));

  unsigned SOffsetRegState = 0;
  unsigned SrcDstRegState = getDefRegState(!IsStore);
  if (i + 1 == e) {
    SOffsetRegState |= getKillRegState(Scavenged);
    // The last implicit use carries the "Kill" flag.
    SrcDstRegState |= getKillRegState(IsKill);
  }

  // Make sure the whole register is defined if there are undef components by
  // adding an implicit def of the super-reg on the first instruction.
  bool NeedSuperRegDef = e > 1 && IsStore && i == 0;
  bool NeedSuperRegImpOperand = e > 1;

  unsigned Lane = RegOffset / 4;
  unsigned LaneE = (RegOffset + EltSize) / 4;
  for ( ; Lane != LaneE; ++Lane) {
    bool IsSubReg = e > 1 || EltSize > 4;
    Register Sub = IsSubReg
           ? Register(getSubReg(ValueReg, getSubRegFromChannel(Lane)))
           : ValueReg;
    auto MIB = spillVGPRtoAGPR(ST, MBB, MI, Index, Lane, Sub, IsKill);
    if (!MIB.getInstr())
      break;
    if (NeedSuperRegDef || (IsSubReg && IsStore && Lane == 0)) {
      MIB.addReg(ValueReg, RegState::ImplicitDefine);
      NeedSuperRegDef = false;
    }
    if (IsSubReg || NeedSuperRegImpOperand) {
      NeedSuperRegImpOperand = true;
      unsigned State = SrcDstRegState;
      if (Lane + 1 != LaneE)
        State &= ~RegState::Kill;
      MIB.addReg(ValueReg, RegState::Implicit | State);
    }
  }

  if (Lane == LaneE) // Fully spilled into AGPRs.
    continue;

  // Offset in bytes from the beginning of the ValueReg to its portion we
  // still need to spill. It may differ from RegOffset if a portion of
  // current SubReg has been already spilled into AGPRs by the loop above.
  unsigned RemRegOffset = Lane * 4;
  unsigned RemEltSize = EltSize - (RemRegOffset - RegOffset);
  if (RemEltSize != EltSize) { // Partially spilled to AGPRs
    assert(IsFlat && EltSize > 4)(static_cast <bool> (IsFlat && EltSize > 4) ?
 void (0) : __assert_fail ("IsFlat && EltSize > 4"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1212, __extension__ __PRETTY_FUNCTION__));

    unsigned NumRegs = RemEltSize / 4;
    SubReg = Register(getSubReg(ValueReg,
                      getSubRegFromChannel(RemRegOffset / 4, NumRegs)));
    unsigned Opc = getFlatScratchSpillOpcode(TII, LoadStoreOp, RemEltSize);
    Desc = &TII->get(Opc);
  }

  unsigned FinalReg = SubReg;

  if (IsAGPR) {
    assert(EltSize == 4)(static_cast <bool> (EltSize == 4) ? void (0) : __assert_fail
 ("EltSize == 4", "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1224, __extension__ __PRETTY_FUNCTION__));

    if (!TmpReg) {
      assert(RS && "Needs to have RegScavenger to spill an AGPR!")(static_cast <bool> (RS && "Needs to have RegScavenger to spill an AGPR!"
) ? void (0) : __assert_fail ("RS && \"Needs to have RegScavenger to spill an AGPR!\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1227, __extension__ __PRETTY_FUNCTION__));
      // FIXME: change to scavengeRegisterBackwards()
      TmpReg = RS->scavengeRegister(&AMDGPU::VGPR_32RegClass, MI, 0);
      RS->setRegUsed(TmpReg);
    }
    if (IsStore) {
      auto AccRead = BuildMI(MBB, MI, DL,
                             TII->get(AMDGPU::V_ACCVGPR_READ_B32_e64), TmpReg)
                         .addReg(SubReg, getKillRegState(IsKill));
      if (NeedSuperRegDef)
        AccRead.addReg(ValueReg, RegState::ImplicitDefine);
      AccRead->setAsmPrinterFlag(MachineInstr::ReloadReuse);
    }
    SubReg = TmpReg;
  }

  MachinePointerInfo PInfo = BasePtrInfo.getWithOffset(RemRegOffset);
  MachineMemOperand *NewMMO =
      MF->getMachineMemOperand(PInfo, MMO->getFlags(), RemEltSize,
                               commonAlignment(Alignment, RemRegOffset));

  auto MIB =
      BuildMI(MBB, MI, DL, *Desc)
          .addReg(SubReg, getDefRegState(!IsStore) | getKillRegState(IsKill));
  if (!IsFlat)
    MIB.addReg(FuncInfo->getScratchRSrcReg());

  if (SOffset == AMDGPU::NoRegister) {
    if (!IsFlat)
      MIB.addImm(0);
  } else {
    MIB.addReg(SOffset, SOffsetRegState);
  }
  MIB.addImm(Offset + RemRegOffset)
     .addImm(0); // cpol
  if (!IsFlat)
    MIB.addImm(0)  // tfe
       .addImm(0); // swz
  MIB.addMemOperand(NewMMO);

  if (!IsAGPR && NeedSuperRegDef)
    MIB.addReg(ValueReg, RegState::ImplicitDefine);

  if (!IsStore && TmpReg != AMDGPU::NoRegister) {
    MIB = BuildMI(MBB, MI, DL, TII->get(AMDGPU::V_ACCVGPR_WRITE_B32_e64),
                  FinalReg)
              .addReg(TmpReg, RegState::Kill);
    MIB->setAsmPrinterFlag(MachineInstr::ReloadReuse);
  }

  if (NeedSuperRegImpOperand)
    MIB.addReg(ValueReg, RegState::Implicit | SrcDstRegState);
}

if (ScratchOffsetRegDelta != 0) {
  // Subtract the offset we added to the ScratchOffset register.
  BuildMI(MBB, MI, DL, TII->get(AMDGPU::S_ADD_I32), SOffset)
      .addReg(SOffset)
      .addImm(-ScratchOffsetRegDelta);
}
1287}

1289void SIRegisterInfo::buildVGPRSpillLoadStore(SGPRSpillBuilder &SB, int Index,
                                           int Offset, bool IsLoad,
                                           bool IsKill) const {
// Load/store VGPR
MachineFrameInfo &FrameInfo = SB.MF.getFrameInfo();
assert(FrameInfo.getStackID(Index) != TargetStackID::SGPRSpill)(static_cast <bool> (FrameInfo.getStackID(Index) != TargetStackID
::SGPRSpill) ? void (0) : __assert_fail ("FrameInfo.getStackID(Index) != TargetStackID::SGPRSpill"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1294, __extension__ __PRETTY_FUNCTION__));

Register FrameReg =
    FrameInfo.isFixedObjectIndex(Index) && hasBasePointer(SB.MF)
        ? getBaseRegister()
        : getFrameRegister(SB.MF);

Align Alignment = FrameInfo.getObjectAlign(Index);
MachinePointerInfo PtrInfo = MachinePointerInfo::getFixedStack(SB.MF, Index);
MachineMemOperand *MMO = SB.MF.getMachineMemOperand(
    PtrInfo, IsLoad ? MachineMemOperand::MOLoad : MachineMemOperand::MOStore,
    SB.EltSize, Alignment);

if (IsLoad) {
  unsigned Opc = ST.enableFlatScratch() ? AMDGPU::SCRATCH_LOAD_DWORD_SADDR
                                        : AMDGPU::BUFFER_LOAD_DWORD_OFFSET;
  buildSpillLoadStore(SB.MBB, SB.MI, Opc, Index, SB.TmpVGPR, false, FrameReg,
                      Offset * SB.EltSize, MMO, SB.RS);
} else {
  unsigned Opc = ST.enableFlatScratch() ? AMDGPU::SCRATCH_STORE_DWORD_SADDR
                                        : AMDGPU::BUFFER_STORE_DWORD_OFFSET;
  buildSpillLoadStore(SB.MBB, SB.MI, Opc, Index, SB.TmpVGPR, IsKill, FrameReg,
                      Offset * SB.EltSize, MMO, SB.RS);
  // This only ever adds one VGPR spill
  SB.MFI.addToSpilledVGPRs(1);
}
1320}

1322bool SIRegisterInfo::spillSGPR(MachineBasicBlock::iterator MI,
                             int Index,
                             RegScavenger *RS,
                             LiveIntervals *LIS,
                             bool OnlyToVGPR) const {
SGPRSpillBuilder SB(*this, *ST.getInstrInfo(), isWave32, MI, Index, RS);

ArrayRef<SIMachineFunctionInfo::SpilledReg> VGPRSpills =
    SB.MFI.getSGPRToVGPRSpills(Index);
bool SpillToVGPR = !VGPRSpills.empty();
if (OnlyToVGPR && !SpillToVGPR)
  return false;

assert(SpillToVGPR || (SB.SuperReg != SB.MFI.getStackPtrOffsetReg() &&(static_cast <bool> (SpillToVGPR || (SB.SuperReg != SB.
MFI.getStackPtrOffsetReg() && SB.SuperReg != SB.MFI.getFrameOffsetReg
())) ? void (0) : __assert_fail ("SpillToVGPR || (SB.SuperReg != SB.MFI.getStackPtrOffsetReg() && SB.SuperReg != SB.MFI.getFrameOffsetReg())"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1336, __extension__ __PRETTY_FUNCTION__))
                       SB.SuperReg != SB.MFI.getFrameOffsetReg()))(static_cast <bool> (SpillToVGPR || (SB.SuperReg != SB.
MFI.getStackPtrOffsetReg() && SB.SuperReg != SB.MFI.getFrameOffsetReg
())) ? void (0) : __assert_fail ("SpillToVGPR || (SB.SuperReg != SB.MFI.getStackPtrOffsetReg() && SB.SuperReg != SB.MFI.getFrameOffsetReg())"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1336, __extension__ __PRETTY_FUNCTION__));

if (SpillToVGPR) {
  for (unsigned i = 0, e = SB.NumSubRegs; i < e; ++i) {
    Register SubReg =
        SB.NumSubRegs == 1
            ? SB.SuperReg
            : Register(getSubReg(SB.SuperReg, SB.SplitParts[i]));
    SIMachineFunctionInfo::SpilledReg Spill = VGPRSpills[i];

    bool UseKill = SB.IsKill && i == SB.NumSubRegs - 1;

    // Mark the "old value of vgpr" input undef only if this is the first sgpr
    // spill to this specific vgpr in the first basic block.
    auto MIB = BuildMI(SB.MBB, MI, SB.DL, SB.TII.get(AMDGPU::V_WRITELANE_B32),
                       Spill.VGPR)
                   .addReg(SubReg, getKillRegState(UseKill))
                   .addImm(Spill.Lane)
                   .addReg(Spill.VGPR);
    if (LIS) {
      if (i == 0)
        LIS->ReplaceMachineInstrInMaps(*MI, *MIB);
      else
        LIS->InsertMachineInstrInMaps(*MIB);
    }

    if (i == 0 && SB.NumSubRegs > 1) {
      // We may be spilling a super-register which is only partially defined,
      // and need to ensure later spills think the value is defined.
      MIB.addReg(SB.SuperReg, RegState::ImplicitDefine);
    }

    if (SB.NumSubRegs > 1)
      MIB.addReg(SB.SuperReg, getKillRegState(UseKill) | RegState::Implicit);

    // FIXME: Since this spills to another register instead of an actual
    // frame index, we should delete the frame index when all references to
    // it are fixed.
  }
} else {
  SB.prepare();

  // SubReg carries the "Kill" flag when SubReg == SB.SuperReg.
  unsigned SubKillState = getKillRegState((SB.NumSubRegs == 1) && SB.IsKill);

  // Per VGPR helper data
  auto PVD = SB.getPerVGPRData();

  for (unsigned Offset = 0; Offset < PVD.NumVGPRs; ++Offset) {
    unsigned TmpVGPRFlags = RegState::Undef;

    // Write sub registers into the VGPR
    for (unsigned i = Offset * PVD.PerVGPR,
                  e = std::min((Offset + 1) * PVD.PerVGPR, SB.NumSubRegs);
         i < e; ++i) {
      Register SubReg =
          SB.NumSubRegs == 1
              ? SB.SuperReg
              : Register(getSubReg(SB.SuperReg, SB.SplitParts[i]));

      MachineInstrBuilder WriteLane =
          BuildMI(SB.MBB, MI, SB.DL, SB.TII.get(AMDGPU::V_WRITELANE_B32),
                  SB.TmpVGPR)
              .addReg(SubReg, SubKillState)
              .addImm(i % PVD.PerVGPR)
              .addReg(SB.TmpVGPR, TmpVGPRFlags);
      TmpVGPRFlags = 0;

      if (LIS) {
        if (i == 0)
          LIS->ReplaceMachineInstrInMaps(*MI, *WriteLane);
        else
          LIS->InsertMachineInstrInMaps(*WriteLane);
      }

      // There could be undef components of a spilled super register.
      // TODO: Can we detect this and skip the spill?
      if (SB.NumSubRegs > 1) {
        // The last implicit use of the SB.SuperReg carries the "Kill" flag.
        unsigned SuperKillState = 0;
        if (i + 1 == SB.NumSubRegs)
          SuperKillState |= getKillRegState(SB.IsKill);
        WriteLane.addReg(SB.SuperReg, RegState::Implicit | SuperKillState);
      }
    }

    // Write out VGPR
    SB.readWriteTmpVGPR(Offset, /*IsLoad*/ false);
  }

  SB.restore();
}

MI->eraseFromParent();
SB.MFI.addToSpilledSGPRs(SB.NumSubRegs);

if (LIS)
  LIS->removeAllRegUnitsForPhysReg(SB.SuperReg);

return true;
1436}

1438bool SIRegisterInfo::restoreSGPR(MachineBasicBlock::iterator MI,
                               int Index,
                               RegScavenger *RS,
                               LiveIntervals *LIS,
                               bool OnlyToVGPR) const {
SGPRSpillBuilder SB(*this, *ST.getInstrInfo(), isWave32, MI, Index, RS);

ArrayRef<SIMachineFunctionInfo::SpilledReg> VGPRSpills =
    SB.MFI.getSGPRToVGPRSpills(Index);
bool SpillToVGPR = !VGPRSpills.empty();
5
←
Assuming the condition is false→
if (OnlyToVGPR5.1
'OnlyToVGPR' is false
1
'OnlyToVGPR' is false
 && !SpillToVGPR)
  return false;

if (SpillToVGPR5.2
'SpillToVGPR' is false
2
'SpillToVGPR' is false
) {
6
←
Taking false branch→
  for (unsigned i = 0, e = SB.NumSubRegs; i < e; ++i) {
    Register SubReg =
        SB.NumSubRegs == 1
            ? SB.SuperReg
            : Register(getSubReg(SB.SuperReg, SB.SplitParts[i]));

    SIMachineFunctionInfo::SpilledReg Spill = VGPRSpills[i];
    auto MIB =
        BuildMI(SB.MBB, MI, SB.DL, SB.TII.get(AMDGPU::V_READLANE_B32), SubReg)
            .addReg(Spill.VGPR)
            .addImm(Spill.Lane);
    if (SB.NumSubRegs > 1 && i == 0)
      MIB.addReg(SB.SuperReg, RegState::ImplicitDefine);
    if (LIS) {
      if (i == e - 1)
        LIS->ReplaceMachineInstrInMaps(*MI, *MIB);
      else
        LIS->InsertMachineInstrInMaps(*MIB);
    }

  }
} else {
  SB.prepare();
7
←
Calling 'SGPRSpillBuilder::prepare'→

  // Per VGPR helper data
  auto PVD = SB.getPerVGPRData();

  for (unsigned Offset = 0; Offset < PVD.NumVGPRs; ++Offset) {
    // Load in VGPR data
    SB.readWriteTmpVGPR(Offset, /*IsLoad*/ true);

    // Unpack lanes
    for (unsigned i = Offset * PVD.PerVGPR,
                  e = std::min((Offset + 1) * PVD.PerVGPR, SB.NumSubRegs);
         i < e; ++i) {
      Register SubReg =
          SB.NumSubRegs == 1
              ? SB.SuperReg
              : Register(getSubReg(SB.SuperReg, SB.SplitParts[i]));

      bool LastSubReg = (i + 1 == e);
      auto MIB = BuildMI(SB.MBB, MI, SB.DL,
                         SB.TII.get(AMDGPU::V_READLANE_B32), SubReg)
                     .addReg(SB.TmpVGPR, getKillRegState(LastSubReg))
                     .addImm(i);
      if (SB.NumSubRegs > 1 && i == 0)
        MIB.addReg(SB.SuperReg, RegState::ImplicitDefine);
      if (LIS) {
        if (i == e - 1)
          LIS->ReplaceMachineInstrInMaps(*MI, *MIB);
        else
          LIS->InsertMachineInstrInMaps(*MIB);
      }
    }
  }

  SB.restore();
}

MI->eraseFromParent();

if (LIS)
  LIS->removeAllRegUnitsForPhysReg(SB.SuperReg);

return true;
1517}

1519/// Special case of eliminateFrameIndex. Returns true if the SGPR was spilled to
1520/// a VGPR and the stack slot can be safely eliminated when all other users are
1521/// handled.
1522bool SIRegisterInfo::eliminateSGPRToVGPRSpillFrameIndex(
MachineBasicBlock::iterator MI,
int FI,
RegScavenger *RS,
LiveIntervals *LIS) const {
switch (MI->getOpcode()) {
case AMDGPU::SI_SPILL_S1024_SAVE:
case AMDGPU::SI_SPILL_S512_SAVE:
case AMDGPU::SI_SPILL_S256_SAVE:
case AMDGPU::SI_SPILL_S224_SAVE:
case AMDGPU::SI_SPILL_S192_SAVE:
case AMDGPU::SI_SPILL_S160_SAVE:
case AMDGPU::SI_SPILL_S128_SAVE:
case AMDGPU::SI_SPILL_S96_SAVE:
case AMDGPU::SI_SPILL_S64_SAVE:
case AMDGPU::SI_SPILL_S32_SAVE:
  return spillSGPR(MI, FI, RS, LIS, true);
case AMDGPU::SI_SPILL_S1024_RESTORE:
case AMDGPU::SI_SPILL_S512_RESTORE:
case AMDGPU::SI_SPILL_S256_RESTORE:
case AMDGPU::SI_SPILL_S224_RESTORE:
case AMDGPU::SI_SPILL_S192_RESTORE:
case AMDGPU::SI_SPILL_S160_RESTORE:
case AMDGPU::SI_SPILL_S128_RESTORE:
case AMDGPU::SI_SPILL_S96_RESTORE:
case AMDGPU::SI_SPILL_S64_RESTORE:
case AMDGPU::SI_SPILL_S32_RESTORE:
  return restoreSGPR(MI, FI, RS, LIS, true);
default:
  llvm_unreachable("not an SGPR spill instruction")::llvm::llvm_unreachable_internal("not an SGPR spill instruction"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1551);
}
1553}

1555void SIRegisterInfo::eliminateFrameIndex(MachineBasicBlock::iterator MI,
                                      int SPAdj, unsigned FIOperandNum,
                                      RegScavenger *RS) const {
MachineFunction *MF = MI->getParent()->getParent();
MachineBasicBlock *MBB = MI->getParent();
SIMachineFunctionInfo *MFI = MF->getInfo<SIMachineFunctionInfo>();
MachineFrameInfo &FrameInfo = MF->getFrameInfo();
const SIInstrInfo *TII = ST.getInstrInfo();
DebugLoc DL = MI->getDebugLoc();

assert(SPAdj == 0 && "unhandled SP adjustment in call sequence?")(static_cast <bool> (SPAdj == 0 && "unhandled SP adjustment in call sequence?"
) ? void (0) : __assert_fail ("SPAdj == 0 && \"unhandled SP adjustment in call sequence?\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1565, __extension__ __PRETTY_FUNCTION__));
1
Assuming 'SPAdj' is equal to 0→
2
←
'?' condition is true→

MachineOperand &FIOp = MI->getOperand(FIOperandNum);
int Index = MI->getOperand(FIOperandNum).getIndex();

Register FrameReg = FrameInfo.isFixedObjectIndex(Index) && hasBasePointer(*MF)
                        ? getBaseRegister()
                        : getFrameRegister(*MF);

switch (MI->getOpcode()) {
3
←
Control jumps to 'case SI_SPILL_S32_RESTORE:'  at line 1600→
  // SGPR register spill
  case AMDGPU::SI_SPILL_S1024_SAVE:
  case AMDGPU::SI_SPILL_S512_SAVE:
  case AMDGPU::SI_SPILL_S256_SAVE:
  case AMDGPU::SI_SPILL_S224_SAVE:
  case AMDGPU::SI_SPILL_S192_SAVE:
  case AMDGPU::SI_SPILL_S160_SAVE:
  case AMDGPU::SI_SPILL_S128_SAVE:
  case AMDGPU::SI_SPILL_S96_SAVE:
  case AMDGPU::SI_SPILL_S64_SAVE:
  case AMDGPU::SI_SPILL_S32_SAVE: {
    spillSGPR(MI, Index, RS);
    break;
  }

  // SGPR register restore
  case AMDGPU::SI_SPILL_S1024_RESTORE:
  case AMDGPU::SI_SPILL_S512_RESTORE:
  case AMDGPU::SI_SPILL_S256_RESTORE:
  case AMDGPU::SI_SPILL_S224_RESTORE:
  case AMDGPU::SI_SPILL_S192_RESTORE:
  case AMDGPU::SI_SPILL_S160_RESTORE:
  case AMDGPU::SI_SPILL_S128_RESTORE:
  case AMDGPU::SI_SPILL_S96_RESTORE:
  case AMDGPU::SI_SPILL_S64_RESTORE:
  case AMDGPU::SI_SPILL_S32_RESTORE: {
    restoreSGPR(MI, Index, RS);
4
←
Calling 'SIRegisterInfo::restoreSGPR'→
    break;
  }

  // VGPR register spill
  case AMDGPU::SI_SPILL_V1024_SAVE:
  case AMDGPU::SI_SPILL_V512_SAVE:
  case AMDGPU::SI_SPILL_V256_SAVE:
  case AMDGPU::SI_SPILL_V224_SAVE:
  case AMDGPU::SI_SPILL_V192_SAVE:
  case AMDGPU::SI_SPILL_V160_SAVE:
  case AMDGPU::SI_SPILL_V128_SAVE:
  case AMDGPU::SI_SPILL_V96_SAVE:
  case AMDGPU::SI_SPILL_V64_SAVE:
  case AMDGPU::SI_SPILL_V32_SAVE:
  case AMDGPU::SI_SPILL_A1024_SAVE:
  case AMDGPU::SI_SPILL_A512_SAVE:
  case AMDGPU::SI_SPILL_A256_SAVE:
  case AMDGPU::SI_SPILL_A224_SAVE:
  case AMDGPU::SI_SPILL_A192_SAVE:
  case AMDGPU::SI_SPILL_A160_SAVE:
  case AMDGPU::SI_SPILL_A128_SAVE:
  case AMDGPU::SI_SPILL_A96_SAVE:
  case AMDGPU::SI_SPILL_A64_SAVE:
  case AMDGPU::SI_SPILL_A32_SAVE: {
    const MachineOperand *VData = TII->getNamedOperand(*MI,
                                                       AMDGPU::OpName::vdata);
    assert(TII->getNamedOperand(*MI, AMDGPU::OpName::soffset)->getReg() ==(static_cast <bool> (TII->getNamedOperand(*MI, AMDGPU
::OpName::soffset)->getReg() == MFI->getStackPtrOffsetReg
()) ? void (0) : __assert_fail ("TII->getNamedOperand(*MI, AMDGPU::OpName::soffset)->getReg() == MFI->getStackPtrOffsetReg()"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1629, __extension__ __PRETTY_FUNCTION__))
           MFI->getStackPtrOffsetReg())(static_cast <bool> (TII->getNamedOperand(*MI, AMDGPU
::OpName::soffset)->getReg() == MFI->getStackPtrOffsetReg
()) ? void (0) : __assert_fail ("TII->getNamedOperand(*MI, AMDGPU::OpName::soffset)->getReg() == MFI->getStackPtrOffsetReg()"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1629, __extension__ __PRETTY_FUNCTION__));

    unsigned Opc = ST.enableFlatScratch() ? AMDGPU::SCRATCH_STORE_DWORD_SADDR
                                          : AMDGPU::BUFFER_STORE_DWORD_OFFSET;
    auto *MBB = MI->getParent();
    buildSpillLoadStore(
        *MBB, MI, Opc, Index, VData->getReg(), VData->isKill(), FrameReg,
        TII->getNamedOperand(*MI, AMDGPU::OpName::offset)->getImm(),
        *MI->memoperands_begin(), RS);
    MFI->addToSpilledVGPRs(getNumSubRegsForSpillOp(MI->getOpcode()));
    MI->eraseFromParent();
    break;
  }
  case AMDGPU::SI_SPILL_V32_RESTORE:
  case AMDGPU::SI_SPILL_V64_RESTORE:
  case AMDGPU::SI_SPILL_V96_RESTORE:
  case AMDGPU::SI_SPILL_V128_RESTORE:
  case AMDGPU::SI_SPILL_V160_RESTORE:
  case AMDGPU::SI_SPILL_V192_RESTORE:
  case AMDGPU::SI_SPILL_V224_RESTORE:
  case AMDGPU::SI_SPILL_V256_RESTORE:
  case AMDGPU::SI_SPILL_V512_RESTORE:
  case AMDGPU::SI_SPILL_V1024_RESTORE:
  case AMDGPU::SI_SPILL_A32_RESTORE:
  case AMDGPU::SI_SPILL_A64_RESTORE:
  case AMDGPU::SI_SPILL_A96_RESTORE:
  case AMDGPU::SI_SPILL_A128_RESTORE:
  case AMDGPU::SI_SPILL_A160_RESTORE:
  case AMDGPU::SI_SPILL_A192_RESTORE:
  case AMDGPU::SI_SPILL_A224_RESTORE:
  case AMDGPU::SI_SPILL_A256_RESTORE:
  case AMDGPU::SI_SPILL_A512_RESTORE:
  case AMDGPU::SI_SPILL_A1024_RESTORE: {
    const MachineOperand *VData = TII->getNamedOperand(*MI,
                                                       AMDGPU::OpName::vdata);
    assert(TII->getNamedOperand(*MI, AMDGPU::OpName::soffset)->getReg() ==(static_cast <bool> (TII->getNamedOperand(*MI, AMDGPU
::OpName::soffset)->getReg() == MFI->getStackPtrOffsetReg
()) ? void (0) : __assert_fail ("TII->getNamedOperand(*MI, AMDGPU::OpName::soffset)->getReg() == MFI->getStackPtrOffsetReg()"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1665, __extension__ __PRETTY_FUNCTION__))
           MFI->getStackPtrOffsetReg())(static_cast <bool> (TII->getNamedOperand(*MI, AMDGPU
::OpName::soffset)->getReg() == MFI->getStackPtrOffsetReg
()) ? void (0) : __assert_fail ("TII->getNamedOperand(*MI, AMDGPU::OpName::soffset)->getReg() == MFI->getStackPtrOffsetReg()"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1665, __extension__ __PRETTY_FUNCTION__));

    unsigned Opc = ST.enableFlatScratch() ? AMDGPU::SCRATCH_LOAD_DWORD_SADDR
                                          : AMDGPU::BUFFER_LOAD_DWORD_OFFSET;
    auto *MBB = MI->getParent();
    buildSpillLoadStore(
        *MBB, MI, Opc, Index, VData->getReg(), VData->isKill(), FrameReg,
        TII->getNamedOperand(*MI, AMDGPU::OpName::offset)->getImm(),
        *MI->memoperands_begin(), RS);
    MI->eraseFromParent();
    break;
  }

  default: {
    // Other access to frame index
    const DebugLoc &DL = MI->getDebugLoc();

    int64_t Offset = FrameInfo.getObjectOffset(Index);
    if (ST.enableFlatScratch()) {
      if (TII->isFLATScratch(*MI)) {
        assert((int16_t)FIOperandNum ==(static_cast <bool> ((int16_t)FIOperandNum == AMDGPU::getNamedOperandIdx
(MI->getOpcode(), AMDGPU::OpName::saddr)) ? void (0) : __assert_fail
 ("(int16_t)FIOperandNum == AMDGPU::getNamedOperandIdx(MI->getOpcode(), AMDGPU::OpName::saddr)"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1687, __extension__ __PRETTY_FUNCTION__))
               AMDGPU::getNamedOperandIdx(MI->getOpcode(),(static_cast <bool> ((int16_t)FIOperandNum == AMDGPU::getNamedOperandIdx
(MI->getOpcode(), AMDGPU::OpName::saddr)) ? void (0) : __assert_fail
 ("(int16_t)FIOperandNum == AMDGPU::getNamedOperandIdx(MI->getOpcode(), AMDGPU::OpName::saddr)"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1687, __extension__ __PRETTY_FUNCTION__))
                                          AMDGPU::OpName::saddr))(static_cast <bool> ((int16_t)FIOperandNum == AMDGPU::getNamedOperandIdx
(MI->getOpcode(), AMDGPU::OpName::saddr)) ? void (0) : __assert_fail
 ("(int16_t)FIOperandNum == AMDGPU::getNamedOperandIdx(MI->getOpcode(), AMDGPU::OpName::saddr)"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1687, __extension__ __PRETTY_FUNCTION__));

        // The offset is always swizzled, just replace it
        if (FrameReg)
          FIOp.ChangeToRegister(FrameReg, false);

        if (!Offset)
          return;

        MachineOperand *OffsetOp =
          TII->getNamedOperand(*MI, AMDGPU::OpName::offset);
        int64_t NewOffset = Offset + OffsetOp->getImm();
        if (TII->isLegalFLATOffset(NewOffset, AMDGPUAS::PRIVATE_ADDRESS,
                                   SIInstrFlags::FlatScratch)) {
          OffsetOp->setImm(NewOffset);
          if (FrameReg)
            return;
          Offset = 0;
        }

        assert(!TII->getNamedOperand(*MI, AMDGPU::OpName::vaddr) &&(static_cast <bool> (!TII->getNamedOperand(*MI, AMDGPU
::OpName::vaddr) && "Unexpected vaddr for flat scratch with a FI operand"
) ? void (0) : __assert_fail ("!TII->getNamedOperand(*MI, AMDGPU::OpName::vaddr) && \"Unexpected vaddr for flat scratch with a FI operand\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1708, __extension__ __PRETTY_FUNCTION__))
               "Unexpected vaddr for flat scratch with a FI operand")(static_cast <bool> (!TII->getNamedOperand(*MI, AMDGPU
::OpName::vaddr) && "Unexpected vaddr for flat scratch with a FI operand"
) ? void (0) : __assert_fail ("!TII->getNamedOperand(*MI, AMDGPU::OpName::vaddr) && \"Unexpected vaddr for flat scratch with a FI operand\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1708, __extension__ __PRETTY_FUNCTION__));

        // On GFX10 we have ST mode to use no registers for an address.
        // Otherwise we need to materialize 0 into an SGPR.
        if (!Offset && ST.hasFlatScratchSTMode()) {
          unsigned Opc = MI->getOpcode();
          unsigned NewOpc = AMDGPU::getFlatScratchInstSTfromSS(Opc);
          MI->RemoveOperand(
              AMDGPU::getNamedOperandIdx(Opc, AMDGPU::OpName::saddr));
          MI->setDesc(TII->get(NewOpc));
          return;
        }
      }

      if (!FrameReg) {
        FIOp.ChangeToImmediate(Offset);
        if (TII->isImmOperandLegal(*MI, FIOperandNum, FIOp))
          return;
      }

      // We need to use register here. Check if we can use an SGPR or need
      // a VGPR.
      FIOp.ChangeToRegister(AMDGPU::M0, false);
      bool UseSGPR = TII->isOperandLegal(*MI, FIOperandNum, &FIOp);

      if (!Offset && FrameReg && UseSGPR) {
        FIOp.setReg(FrameReg);
        return;
      }

      const TargetRegisterClass *RC = UseSGPR ? &AMDGPU::SReg_32_XM0RegClass
                                              : &AMDGPU::VGPR_32RegClass;

      Register TmpReg = RS->scavengeRegister(RC, MI, 0, !UseSGPR);
      FIOp.setReg(TmpReg);
      FIOp.setIsKill(true);

      if ((!FrameReg || !Offset) && TmpReg) {
        unsigned Opc = UseSGPR ? AMDGPU::S_MOV_B32 : AMDGPU::V_MOV_B32_e32;
        auto MIB = BuildMI(*MBB, MI, DL, TII->get(Opc), TmpReg);
        if (FrameReg)
          MIB.addReg(FrameReg);
        else
          MIB.addImm(Offset);

        return;
      }

      Register TmpSReg =
          UseSGPR ? TmpReg
                  : RS->scavengeRegister(&AMDGPU::SReg_32_XM0RegClass, MI, 0,
                                         !UseSGPR);

      // TODO: for flat scratch another attempt can be made with a VGPR index
      //       if no SGPRs can be scavenged.
      if ((!TmpSReg && !FrameReg) || (!TmpReg && !UseSGPR))
        report_fatal_error("Cannot scavenge register in FI elimination!");

      if (!TmpSReg) {
        // Use frame register and restore it after.
        TmpSReg = FrameReg;
        FIOp.setReg(FrameReg);
        FIOp.setIsKill(false);
      }

      BuildMI(*MBB, MI, DL, TII->get(AMDGPU::S_ADD_I32), TmpSReg)
          .addReg(FrameReg)
          .addImm(Offset);

      if (!UseSGPR)
        BuildMI(*MBB, MI, DL, TII->get(AMDGPU::V_MOV_B32_e32), TmpReg)
          .addReg(TmpSReg, RegState::Kill);

      if (TmpSReg == FrameReg) {
        // Undo frame register modification.
        BuildMI(*MBB, std::next(MI), DL, TII->get(AMDGPU::S_ADD_I32),
                FrameReg)
            .addReg(FrameReg)
            .addImm(-Offset);
      }

      return;
    }

    bool IsMUBUF = TII->isMUBUF(*MI);

    if (!IsMUBUF && !MFI->isEntryFunction()) {
      // Convert to a swizzled stack address by scaling by the wave size.
      //
      // In an entry function/kernel the offset is already swizzled.

      bool IsCopy = MI->getOpcode() == AMDGPU::V_MOV_B32_e32;
      Register ResultReg =
          IsCopy ? MI->getOperand(0).getReg()
                 : RS->scavengeRegister(&AMDGPU::VGPR_32RegClass, MI, 0);

      int64_t Offset = FrameInfo.getObjectOffset(Index);
      if (Offset == 0) {
        // XXX - This never happens because of emergency scavenging slot at 0?
        BuildMI(*MBB, MI, DL, TII->get(AMDGPU::V_LSHRREV_B32_e64), ResultReg)
          .addImm(ST.getWavefrontSizeLog2())
          .addReg(FrameReg);
      } else {
        if (auto MIB = TII->getAddNoCarry(*MBB, MI, DL, ResultReg, *RS)) {
          // Reuse ResultReg in intermediate step.
          Register ScaledReg = ResultReg;

          BuildMI(*MBB, *MIB, DL, TII->get(AMDGPU::V_LSHRREV_B32_e64),
                  ScaledReg)
            .addImm(ST.getWavefrontSizeLog2())
            .addReg(FrameReg);

          const bool IsVOP2 = MIB->getOpcode() == AMDGPU::V_ADD_U32_e32;

          // TODO: Fold if use instruction is another add of a constant.
          if (IsVOP2 || AMDGPU::isInlinableLiteral32(Offset, ST.hasInv2PiInlineImm())) {
            // FIXME: This can fail
            MIB.addImm(Offset);
            MIB.addReg(ScaledReg, RegState::Kill);
            if (!IsVOP2)
              MIB.addImm(0); // clamp bit
          } else {
            assert(MIB->getOpcode() == AMDGPU::V_ADD_CO_U32_e64 &&(static_cast <bool> (MIB->getOpcode() == AMDGPU::V_ADD_CO_U32_e64
 && "Need to reuse carry out register") ? void (0) : __assert_fail
 ("MIB->getOpcode() == AMDGPU::V_ADD_CO_U32_e64 && \"Need to reuse carry out register\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1831, __extension__ __PRETTY_FUNCTION__))
                   "Need to reuse carry out register")(static_cast <bool> (MIB->getOpcode() == AMDGPU::V_ADD_CO_U32_e64
 && "Need to reuse carry out register") ? void (0) : __assert_fail
 ("MIB->getOpcode() == AMDGPU::V_ADD_CO_U32_e64 && \"Need to reuse carry out register\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1831, __extension__ __PRETTY_FUNCTION__));

            // Use scavenged unused carry out as offset register.
            Register ConstOffsetReg;
            if (!isWave32)
              ConstOffsetReg = getSubReg(MIB.getReg(1), AMDGPU::sub0);
            else
              ConstOffsetReg = MIB.getReg(1);

            BuildMI(*MBB, *MIB, DL, TII->get(AMDGPU::S_MOV_B32), ConstOffsetReg)
              .addImm(Offset);
            MIB.addReg(ConstOffsetReg, RegState::Kill);
            MIB.addReg(ScaledReg, RegState::Kill);
            MIB.addImm(0); // clamp bit
          }
        } else {
          // We have to produce a carry out, and there isn't a free SGPR pair
          // for it. We can keep the whole computation on the SALU to avoid
          // clobbering an additional register at the cost of an extra mov.

          // We may have 1 free scratch SGPR even though a carry out is
          // unavailable. Only one additional mov is needed.
          Register TmpScaledReg =
              RS->scavengeRegister(&AMDGPU::SReg_32_XM0RegClass, MI, 0, false);
          Register ScaledReg = TmpScaledReg.isValid() ? TmpScaledReg : FrameReg;

          BuildMI(*MBB, MI, DL, TII->get(AMDGPU::S_LSHR_B32), ScaledReg)
            .addReg(FrameReg)
            .addImm(ST.getWavefrontSizeLog2());
          BuildMI(*MBB, MI, DL, TII->get(AMDGPU::S_ADD_I32), ScaledReg)
              .addReg(ScaledReg, RegState::Kill)
              .addImm(Offset);
          BuildMI(*MBB, MI, DL, TII->get(AMDGPU::COPY), ResultReg)
            .addReg(ScaledReg, RegState::Kill);

          // If there were truly no free SGPRs, we need to undo everything.
          if (!TmpScaledReg.isValid()) {
            BuildMI(*MBB, MI, DL, TII->get(AMDGPU::S_ADD_I32), ScaledReg)
                .addReg(ScaledReg, RegState::Kill)
                .addImm(-Offset);
            BuildMI(*MBB, MI, DL, TII->get(AMDGPU::S_LSHL_B32), ScaledReg)
              .addReg(FrameReg)
              .addImm(ST.getWavefrontSizeLog2());
          }
        }
      }

      // Don't introduce an extra copy if we're just materializing in a mov.
      if (IsCopy)
        MI->eraseFromParent();
      else
        FIOp.ChangeToRegister(ResultReg, false, false, true);
      return;
    }

    if (IsMUBUF) {
      // Disable offen so we don't need a 0 vgpr base.
      assert(static_cast<int>(FIOperandNum) ==(static_cast <bool> (static_cast<int>(FIOperandNum
) == AMDGPU::getNamedOperandIdx(MI->getOpcode(), AMDGPU::OpName
::vaddr)) ? void (0) : __assert_fail ("static_cast<int>(FIOperandNum) == AMDGPU::getNamedOperandIdx(MI->getOpcode(), AMDGPU::OpName::vaddr)"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1890, __extension__ __PRETTY_FUNCTION__))
             AMDGPU::getNamedOperandIdx(MI->getOpcode(),(static_cast <bool> (static_cast<int>(FIOperandNum
) == AMDGPU::getNamedOperandIdx(MI->getOpcode(), AMDGPU::OpName
::vaddr)) ? void (0) : __assert_fail ("static_cast<int>(FIOperandNum) == AMDGPU::getNamedOperandIdx(MI->getOpcode(), AMDGPU::OpName::vaddr)"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1890, __extension__ __PRETTY_FUNCTION__))
                                        AMDGPU::OpName::vaddr))(static_cast <bool> (static_cast<int>(FIOperandNum
) == AMDGPU::getNamedOperandIdx(MI->getOpcode(), AMDGPU::OpName
::vaddr)) ? void (0) : __assert_fail ("static_cast<int>(FIOperandNum) == AMDGPU::getNamedOperandIdx(MI->getOpcode(), AMDGPU::OpName::vaddr)"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1890, __extension__ __PRETTY_FUNCTION__));

      auto &SOffset = *TII->getNamedOperand(*MI, AMDGPU::OpName::soffset);
      assert((SOffset.isImm() && SOffset.getImm() == 0))(static_cast <bool> ((SOffset.isImm() && SOffset
.getImm() == 0)) ? void (0) : __assert_fail ("(SOffset.isImm() && SOffset.getImm() == 0)"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 1893, __extension__ __PRETTY_FUNCTION__));

      if (FrameReg != AMDGPU::NoRegister)
        SOffset.ChangeToRegister(FrameReg, false);

      int64_t Offset = FrameInfo.getObjectOffset(Index);
      int64_t OldImm
        = TII->getNamedOperand(*MI, AMDGPU::OpName::offset)->getImm();
      int64_t NewOffset = OldImm + Offset;

      if (SIInstrInfo::isLegalMUBUFImmOffset(NewOffset) &&
          buildMUBUFOffsetLoadStore(ST, FrameInfo, MI, Index, NewOffset)) {
        MI->eraseFromParent();
        return;
      }
    }

    // If the offset is simply too big, don't convert to a scratch wave offset
    // relative index.

    FIOp.ChangeToImmediate(Offset);
    if (!TII->isImmOperandLegal(*MI, FIOperandNum, FIOp)) {
      Register TmpReg = RS->scavengeRegister(&AMDGPU::VGPR_32RegClass, MI, 0);
      BuildMI(*MBB, MI, DL, TII->get(AMDGPU::V_MOV_B32_e32), TmpReg)
        .addImm(Offset);
      FIOp.ChangeToRegister(TmpReg, false, false, true);
    }
  }
}
1922}

1924StringRef SIRegisterInfo::getRegAsmName(MCRegister Reg) const {
return AMDGPUInstPrinter::getRegisterName(Reg);
1926}

1928static const TargetRegisterClass *
1929getAnyVGPRClassForBitWidth(unsigned BitWidth) {
if (BitWidth <= 64)
  return &AMDGPU::VReg_64RegClass;
if (BitWidth <= 96)
  return &AMDGPU::VReg_96RegClass;
if (BitWidth <= 128)
  return &AMDGPU::VReg_128RegClass;
if (BitWidth <= 160)
  return &AMDGPU::VReg_160RegClass;
if (BitWidth <= 192)
  return &AMDGPU::VReg_192RegClass;
if (BitWidth <= 224)
  return &AMDGPU::VReg_224RegClass;
if (BitWidth <= 256)
  return &AMDGPU::VReg_256RegClass;
if (BitWidth <= 512)
  return &AMDGPU::VReg_512RegClass;
if (BitWidth <= 1024)
  return &AMDGPU::VReg_1024RegClass;

return nullptr;
1950}

1952static const TargetRegisterClass *
1953getAlignedVGPRClassForBitWidth(unsigned BitWidth) {
if (BitWidth <= 64)
  return &AMDGPU::VReg_64_Align2RegClass;
if (BitWidth <= 96)
  return &AMDGPU::VReg_96_Align2RegClass;
if (BitWidth <= 128)
  return &AMDGPU::VReg_128_Align2RegClass;
if (BitWidth <= 160)
  return &AMDGPU::VReg_160_Align2RegClass;
if (BitWidth <= 192)
  return &AMDGPU::VReg_192_Align2RegClass;
if (BitWidth <= 224)
  return &AMDGPU::VReg_224_Align2RegClass;
if (BitWidth <= 256)
  return &AMDGPU::VReg_256_Align2RegClass;
if (BitWidth <= 512)
  return &AMDGPU::VReg_512_Align2RegClass;
if (BitWidth <= 1024)
  return &AMDGPU::VReg_1024_Align2RegClass;

return nullptr;
1974}

1976const TargetRegisterClass *
1977SIRegisterInfo::getVGPRClassForBitWidth(unsigned BitWidth) const {
if (BitWidth == 1)
  return &AMDGPU::VReg_1RegClass;
if (BitWidth <= 16)
  return &AMDGPU::VGPR_LO16RegClass;
if (BitWidth <= 32)
  return &AMDGPU::VGPR_32RegClass;
return ST.needsAlignedVGPRs() ? getAlignedVGPRClassForBitWidth(BitWidth)
                              : getAnyVGPRClassForBitWidth(BitWidth);
1986}

1988static const TargetRegisterClass *
1989getAnyAGPRClassForBitWidth(unsigned BitWidth) {
if (BitWidth <= 64)
  return &AMDGPU::AReg_64RegClass;
if (BitWidth <= 96)
  return &AMDGPU::AReg_96RegClass;
if (BitWidth <= 128)
  return &AMDGPU::AReg_128RegClass;
if (BitWidth <= 160)
  return &AMDGPU::AReg_160RegClass;
if (BitWidth <= 192)
  return &AMDGPU::AReg_192RegClass;
if (BitWidth <= 224)
  return &AMDGPU::AReg_224RegClass;
if (BitWidth <= 256)
  return &AMDGPU::AReg_256RegClass;
if (BitWidth <= 512)
  return &AMDGPU::AReg_512RegClass;
if (BitWidth <= 1024)
  return &AMDGPU::AReg_1024RegClass;

return nullptr;
2010}

2012static const TargetRegisterClass *
2013getAlignedAGPRClassForBitWidth(unsigned BitWidth) {
if (BitWidth <= 64)
  return &AMDGPU::AReg_64_Align2RegClass;
if (BitWidth <= 96)
  return &AMDGPU::AReg_96_Align2RegClass;
if (BitWidth <= 128)
  return &AMDGPU::AReg_128_Align2RegClass;
if (BitWidth <= 160)
  return &AMDGPU::AReg_160_Align2RegClass;
if (BitWidth <= 192)
  return &AMDGPU::AReg_192_Align2RegClass;
if (BitWidth <= 224)
  return &AMDGPU::AReg_224_Align2RegClass;
if (BitWidth <= 256)
  return &AMDGPU::AReg_256_Align2RegClass;
if (BitWidth <= 512)
  return &AMDGPU::AReg_512_Align2RegClass;
if (BitWidth <= 1024)
  return &AMDGPU::AReg_1024_Align2RegClass;

return nullptr;
2034}

2036const TargetRegisterClass *
2037SIRegisterInfo::getAGPRClassForBitWidth(unsigned BitWidth) const {
if (BitWidth <= 16)
  return &AMDGPU::AGPR_LO16RegClass;
if (BitWidth <= 32)
  return &AMDGPU::AGPR_32RegClass;
return ST.needsAlignedVGPRs() ? getAlignedAGPRClassForBitWidth(BitWidth)
                              : getAnyAGPRClassForBitWidth(BitWidth);
2044}

2046const TargetRegisterClass *
2047SIRegisterInfo::getSGPRClassForBitWidth(unsigned BitWidth) {
if (BitWidth <= 16)
  return &AMDGPU::SGPR_LO16RegClass;
if (BitWidth <= 32)
  return &AMDGPU::SReg_32RegClass;
if (BitWidth <= 64)
  return &AMDGPU::SReg_64RegClass;
if (BitWidth <= 96)
  return &AMDGPU::SGPR_96RegClass;
if (BitWidth <= 128)
  return &AMDGPU::SGPR_128RegClass;
if (BitWidth <= 160)
  return &AMDGPU::SGPR_160RegClass;
if (BitWidth <= 192)
  return &AMDGPU::SGPR_192RegClass;
if (BitWidth <= 224)
  return &AMDGPU::SGPR_224RegClass;
if (BitWidth <= 256)
  return &AMDGPU::SGPR_256RegClass;
if (BitWidth <= 512)
  return &AMDGPU::SGPR_512RegClass;
if (BitWidth <= 1024)
  return &AMDGPU::SGPR_1024RegClass;

return nullptr;
2072}

2074// FIXME: This is very slow. It might be worth creating a map from physreg to
2075// register class.
2076const TargetRegisterClass *
2077SIRegisterInfo::getPhysRegClass(MCRegister Reg) const {
static const TargetRegisterClass *const BaseClasses[] = {
  &AMDGPU::VGPR_LO16RegClass,
  &AMDGPU::VGPR_HI16RegClass,
  &AMDGPU::SReg_LO16RegClass,
  &AMDGPU::AGPR_LO16RegClass,
  &AMDGPU::VGPR_32RegClass,
  &AMDGPU::SReg_32RegClass,
  &AMDGPU::AGPR_32RegClass,
  &AMDGPU::AGPR_32RegClass,
  &AMDGPU::VReg_64_Align2RegClass,
  &AMDGPU::VReg_64RegClass,
  &AMDGPU::SReg_64RegClass,
  &AMDGPU::AReg_64_Align2RegClass,
  &AMDGPU::AReg_64RegClass,
  &AMDGPU::VReg_96_Align2RegClass,
  &AMDGPU::VReg_96RegClass,
  &AMDGPU::SReg_96RegClass,
  &AMDGPU::AReg_96_Align2RegClass,
  &AMDGPU::AReg_96RegClass,
  &AMDGPU::VReg_128_Align2RegClass,
  &AMDGPU::VReg_128RegClass,
  &AMDGPU::SReg_128RegClass,
  &AMDGPU::AReg_128_Align2RegClass,
  &AMDGPU::AReg_128RegClass,
  &AMDGPU::VReg_160_Align2RegClass,
  &AMDGPU::VReg_160RegClass,
  &AMDGPU::SReg_160RegClass,
  &AMDGPU::AReg_160_Align2RegClass,
  &AMDGPU::AReg_160RegClass,
  &AMDGPU::VReg_192_Align2RegClass,
  &AMDGPU::VReg_192RegClass,
  &AMDGPU::SReg_192RegClass,
  &AMDGPU::AReg_192_Align2RegClass,
  &AMDGPU::AReg_192RegClass,
  &AMDGPU::VReg_224_Align2RegClass,
  &AMDGPU::VReg_224RegClass,
  &AMDGPU::SReg_224RegClass,
  &AMDGPU::AReg_224_Align2RegClass,
  &AMDGPU::AReg_224RegClass,
  &AMDGPU::VReg_256_Align2RegClass,
  &AMDGPU::VReg_256RegClass,
  &AMDGPU::SReg_256RegClass,
  &AMDGPU::AReg_256_Align2RegClass,
  &AMDGPU::AReg_256RegClass,
  &AMDGPU::VReg_512_Align2RegClass,
  &AMDGPU::VReg_512RegClass,
  &AMDGPU::SReg_512RegClass,
  &AMDGPU::AReg_512_Align2RegClass,
  &AMDGPU::AReg_512RegClass,
  &AMDGPU::SReg_1024RegClass,
  &AMDGPU::VReg_1024_Align2RegClass,
  &AMDGPU::VReg_1024RegClass,
  &AMDGPU::AReg_1024_Align2RegClass,
  &AMDGPU::AReg_1024RegClass,
  &AMDGPU::SCC_CLASSRegClass,
  &AMDGPU::Pseudo_SReg_32RegClass,
  &AMDGPU::Pseudo_SReg_128RegClass,
};

for (const TargetRegisterClass *BaseClass : BaseClasses) {
  if (BaseClass->contains(Reg)) {
    return BaseClass;
  }
}
return nullptr;
2143}

2145bool SIRegisterInfo::isSGPRReg(const MachineRegisterInfo &MRI,
                             Register Reg) const {
const TargetRegisterClass *RC;
if (Reg.isVirtual())
  RC = MRI.getRegClass(Reg);
else
  RC = getPhysRegClass(Reg);
return isSGPRClass(RC);
2153}

2155// TODO: It might be helpful to have some target specific flags in
2156// TargetRegisterClass to mark which classes are VGPRs to make this trivial.
2157bool SIRegisterInfo::hasVGPRs(const TargetRegisterClass *RC) const {
unsigned Size = getRegSizeInBits(*RC);
if (Size == 16) {
  return getCommonSubClass(&AMDGPU::VGPR_LO16RegClass, RC) != nullptr ||
         getCommonSubClass(&AMDGPU::VGPR_HI16RegClass, RC) != nullptr;
}
const TargetRegisterClass *VRC = getVGPRClassForBitWidth(Size);
if (!VRC) {
  assert(Size < 32 && "Invalid register class size")(static_cast <bool> (Size < 32 && "Invalid register class size"
) ? void (0) : __assert_fail ("Size < 32 && \"Invalid register class size\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 2165, __extension__ __PRETTY_FUNCTION__));
  return false;
}
return getCommonSubClass(VRC, RC) != nullptr;
2169}

2171bool SIRegisterInfo::hasAGPRs(const TargetRegisterClass *RC) const {
unsigned Size = getRegSizeInBits(*RC);
if (Size < 16)
  return false;
const TargetRegisterClass *ARC = getAGPRClassForBitWidth(Size);
if (!ARC) {
  assert(getVGPRClassForBitWidth(Size) && "Invalid register class size")(static_cast <bool> (getVGPRClassForBitWidth(Size) &&
 "Invalid register class size") ? void (0) : __assert_fail ("getVGPRClassForBitWidth(Size) && \"Invalid register class size\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 2177, __extension__ __PRETTY_FUNCTION__));
  return false;
}
return getCommonSubClass(ARC, RC) != nullptr;
2181}

2183const TargetRegisterClass *
2184SIRegisterInfo::getEquivalentVGPRClass(const TargetRegisterClass *SRC) const {
unsigned Size = getRegSizeInBits(*SRC);
const TargetRegisterClass *VRC = getVGPRClassForBitWidth(Size);
assert(VRC && "Invalid register class size")(static_cast <bool> (VRC && "Invalid register class size"
) ? void (0) : __assert_fail ("VRC && \"Invalid register class size\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 2187, __extension__ __PRETTY_FUNCTION__));
return VRC;
2189}

2191const TargetRegisterClass *
2192SIRegisterInfo::getEquivalentAGPRClass(const TargetRegisterClass *SRC) const {
unsigned Size = getRegSizeInBits(*SRC);
const TargetRegisterClass *ARC = getAGPRClassForBitWidth(Size);
assert(ARC && "Invalid register class size")(static_cast <bool> (ARC && "Invalid register class size"
) ? void (0) : __assert_fail ("ARC && \"Invalid register class size\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 2195, __extension__ __PRETTY_FUNCTION__));
return ARC;
2197}

2199const TargetRegisterClass *
2200SIRegisterInfo::getEquivalentSGPRClass(const TargetRegisterClass *VRC) const {
unsigned Size = getRegSizeInBits(*VRC);
if (Size == 32)
  return &AMDGPU::SGPR_32RegClass;
const TargetRegisterClass *SRC = getSGPRClassForBitWidth(Size);
assert(SRC && "Invalid register class size")(static_cast <bool> (SRC && "Invalid register class size"
) ? void (0) : __assert_fail ("SRC && \"Invalid register class size\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 2205, __extension__ __PRETTY_FUNCTION__));
return SRC;
2207}

2209const TargetRegisterClass *SIRegisterInfo::getSubRegClass(
                       const TargetRegisterClass *RC, unsigned SubIdx) const {
if (SubIdx == AMDGPU::NoSubRegister)
  return RC;

// We can assume that each lane corresponds to one 32-bit register.
unsigned Size = getNumChannelsFromSubReg(SubIdx) * 32;
if (isSGPRClass(RC)) {
  if (Size == 32)
    RC = &AMDGPU::SGPR_32RegClass;
  else
    RC = getSGPRClassForBitWidth(Size);
} else if (hasAGPRs(RC)) {
  RC = getAGPRClassForBitWidth(Size);
} else {
  RC = getVGPRClassForBitWidth(Size);
}
assert(RC && "Invalid sub-register class size")(static_cast <bool> (RC && "Invalid sub-register class size"
) ? void (0) : __assert_fail ("RC && \"Invalid sub-register class size\""
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 2226, __extension__ __PRETTY_FUNCTION__));
return RC;
2228}

2230const TargetRegisterClass *
2231SIRegisterInfo::getCompatibleSubRegClass(const TargetRegisterClass *SuperRC,
                                       const TargetRegisterClass *SubRC,
                                       unsigned SubIdx) const {
// Ensure this subregister index is aligned in the super register.
const TargetRegisterClass *MatchRC =
    getMatchingSuperRegClass(SuperRC, SubRC, SubIdx);
return MatchRC && MatchRC->hasSubClassEq(SuperRC) ? MatchRC : nullptr;
2238}

2240bool SIRegisterInfo::opCanUseInlineConstant(unsigned OpType) const {
if (OpType >= AMDGPU::OPERAND_REG_INLINE_AC_FIRST &&
    OpType <= AMDGPU::OPERAND_REG_INLINE_AC_LAST)
  return !ST.hasMFMAInlineLiteralBug();

return OpType >= AMDGPU::OPERAND_SRC_FIRST &&
       OpType <= AMDGPU::OPERAND_SRC_LAST;
2247}

2249bool SIRegisterInfo::shouldRewriteCopySrc(
const TargetRegisterClass *DefRC,
unsigned DefSubReg,
const TargetRegisterClass *SrcRC,
unsigned SrcSubReg) const {
// We want to prefer the smallest register class possible, so we don't want to
// stop and rewrite on anything that looks like a subregister
// extract. Operations mostly don't care about the super register class, so we
// only want to stop on the most basic of copies between the same register
// class.
//
// e.g. if we have something like
// %0 = ...
// %1 = ...
// %2 = REG_SEQUENCE %0, sub0, %1, sub1, %2, sub2
// %3 = COPY %2, sub0
//
// We want to look through the COPY to find:
//  => %3 = COPY %0

// Plain copy.
return getCommonSubClass(DefRC, SrcRC) != nullptr;
2271}

2273bool SIRegisterInfo::opCanUseLiteralConstant(unsigned OpType) const {
// TODO: 64-bit operands have extending behavior from 32-bit literal.
return OpType >= AMDGPU::OPERAND_REG_IMM_FIRST &&
       OpType <= AMDGPU::OPERAND_REG_IMM_LAST;
2277}

2279/// Returns a lowest register that is not used at any point in the function.
2280///        If all registers are used, then this function will return
2281///         AMDGPU::NoRegister. If \p ReserveHighestVGPR = true, then return
2282///         highest unused register.
2283MCRegister SIRegisterInfo::findUnusedRegister(const MachineRegisterInfo &MRI,
                                            const TargetRegisterClass *RC,
                                            const MachineFunction &MF,
                                            bool ReserveHighestVGPR) const {
if (ReserveHighestVGPR) {
  for (MCRegister Reg : reverse(*RC))
    if (MRI.isAllocatable(Reg) && !MRI.isPhysRegUsed(Reg))
      return Reg;
} else {
  for (MCRegister Reg : *RC)
    if (MRI.isAllocatable(Reg) && !MRI.isPhysRegUsed(Reg))
      return Reg;
}
return MCRegister();
2297}

2299ArrayRef<int16_t> SIRegisterInfo::getRegSplitParts(const TargetRegisterClass *RC,
                                                 unsigned EltSize) const {
const unsigned RegBitWidth = AMDGPU::getRegBitWidth(*RC->MC);
assert(RegBitWidth >= 32 && RegBitWidth <= 1024)(static_cast <bool> (RegBitWidth >= 32 && RegBitWidth
 <= 1024) ? void (0) : __assert_fail ("RegBitWidth >= 32 && RegBitWidth <= 1024"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 2302, __extension__ __PRETTY_FUNCTION__));

const unsigned RegDWORDs = RegBitWidth / 32;
const unsigned EltDWORDs = EltSize / 4;
assert(RegSplitParts.size() + 1 >= EltDWORDs)(static_cast <bool> (RegSplitParts.size() + 1 >= EltDWORDs
) ? void (0) : __assert_fail ("RegSplitParts.size() + 1 >= EltDWORDs"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 2306, __extension__ __PRETTY_FUNCTION__));

const std::vector<int16_t> &Parts = RegSplitParts[EltDWORDs - 1];
const unsigned NumParts = RegDWORDs / EltDWORDs;

return makeArrayRef(Parts.data(), NumParts);
2312}

2314const TargetRegisterClass*
2315SIRegisterInfo::getRegClassForReg(const MachineRegisterInfo &MRI,
                                Register Reg) const {
return Reg.isVirtual() ? MRI.getRegClass(Reg) : getPhysRegClass(Reg);
2318}

2320bool SIRegisterInfo::isVGPR(const MachineRegisterInfo &MRI,
                          Register Reg) const {
const TargetRegisterClass *RC = getRegClassForReg(MRI, Reg);
// Registers without classes are unaddressable, SGPR-like registers.
return RC && hasVGPRs(RC);
2325}

2327bool SIRegisterInfo::isAGPR(const MachineRegisterInfo &MRI,
                          Register Reg) const {
const TargetRegisterClass *RC = getRegClassForReg(MRI, Reg);

// Registers without classes are unaddressable, SGPR-like registers.
return RC && hasAGPRs(RC);
2333}

2335bool SIRegisterInfo::shouldCoalesce(MachineInstr *MI,
                                  const TargetRegisterClass *SrcRC,
                                  unsigned SubReg,
                                  const TargetRegisterClass *DstRC,
                                  unsigned DstSubReg,
                                  const TargetRegisterClass *NewRC,
                                  LiveIntervals &LIS) const {
unsigned SrcSize = getRegSizeInBits(*SrcRC);
unsigned DstSize = getRegSizeInBits(*DstRC);
unsigned NewSize = getRegSizeInBits(*NewRC);

// Do not increase size of registers beyond dword, we would need to allocate
// adjacent registers and constraint regalloc more than needed.

// Always allow dword coalescing.
if (SrcSize <= 32 || DstSize <= 32)
  return true;

return NewSize <= DstSize || NewSize <= SrcSize;
2354}

2356unsigned SIRegisterInfo::getRegPressureLimit(const TargetRegisterClass *RC,
                                           MachineFunction &MF) const {
const SIMachineFunctionInfo *MFI = MF.getInfo<SIMachineFunctionInfo>();

unsigned Occupancy = ST.getOccupancyWithLocalMemSize(MFI->getLDSSize(),
                                                     MF.getFunction());
switch (RC->getID()) {
default:
  return AMDGPUGenRegisterInfo::getRegPressureLimit(RC, MF);
case AMDGPU::VGPR_32RegClassID:
case AMDGPU::VGPR_LO16RegClassID:
case AMDGPU::VGPR_HI16RegClassID:
  return std::min(ST.getMaxNumVGPRs(Occupancy), ST.getMaxNumVGPRs(MF));
case AMDGPU::SGPR_32RegClassID:
case AMDGPU::SGPR_LO16RegClassID:
  return std::min(ST.getMaxNumSGPRs(Occupancy, true), ST.getMaxNumSGPRs(MF));
}
2373}

2375unsigned SIRegisterInfo::getRegPressureSetLimit(const MachineFunction &MF,
                                              unsigned Idx) const {
if (Idx == AMDGPU::RegisterPressureSets::VGPR_32 ||
    Idx == AMDGPU::RegisterPressureSets::AGPR_32)
  return getRegPressureLimit(&AMDGPU::VGPR_32RegClass,
                             const_cast<MachineFunction &>(MF));

if (Idx == AMDGPU::RegisterPressureSets::SReg_32)
  return getRegPressureLimit(&AMDGPU::SGPR_32RegClass,
                             const_cast<MachineFunction &>(MF));

llvm_unreachable("Unexpected register pressure set!")::llvm::llvm_unreachable_internal("Unexpected register pressure set!"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 2386);
2387}

2389const int *SIRegisterInfo::getRegUnitPressureSets(unsigned RegUnit) const {
static const int Empty[] = { -1 };

if (RegPressureIgnoredUnits[RegUnit])
  return Empty;

return AMDGPUGenRegisterInfo::getRegUnitPressureSets(RegUnit);
2396}

2398MCRegister SIRegisterInfo::getReturnAddressReg(const MachineFunction &MF) const {
// Not a callee saved register.
return AMDGPU::SGPR30_SGPR31;
2401}

2403const TargetRegisterClass *
2404SIRegisterInfo::getRegClassForSizeOnBank(unsigned Size,
                                       const RegisterBank &RB,
                                       const MachineRegisterInfo &MRI) const {
switch (RB.getID()) {
case AMDGPU::VGPRRegBankID:
  return getVGPRClassForBitWidth(std::max(32u, Size));
case AMDGPU::VCCRegBankID:
  assert(Size == 1)(static_cast <bool> (Size == 1) ? void (0) : __assert_fail
 ("Size == 1", "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 2411, __extension__ __PRETTY_FUNCTION__));
  return isWave32 ? &AMDGPU::SReg_32_XM0_XEXECRegClass
                  : &AMDGPU::SReg_64_XEXECRegClass;
case AMDGPU::SGPRRegBankID:
  return getSGPRClassForBitWidth(std::max(32u, Size));
case AMDGPU::AGPRRegBankID:
  return getAGPRClassForBitWidth(std::max(32u, Size));
default:
  llvm_unreachable("unknown register bank")::llvm::llvm_unreachable_internal("unknown register bank", "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 2419);
}
2421}

2423const TargetRegisterClass *
2424SIRegisterInfo::getConstrainedRegClassForOperand(const MachineOperand &MO,
                                       const MachineRegisterInfo &MRI) const {
const RegClassOrRegBank &RCOrRB = MRI.getRegClassOrRegBank(MO.getReg());
if (const RegisterBank *RB = RCOrRB.dyn_cast<const RegisterBank*>())
  return getRegClassForTypeOnBank(MRI.getType(MO.getReg()), *RB, MRI);

const TargetRegisterClass *RC = RCOrRB.get<const TargetRegisterClass*>();
return getAllocatableClass(RC);
2432}

2434MCRegister SIRegisterInfo::getVCC() const {
return isWave32 ? AMDGPU::VCC_LO : AMDGPU::VCC;
2436}

2438const TargetRegisterClass *SIRegisterInfo::getVGPR64Class() const {
// VGPR tuples have an alignment requirement on gfx90a variants.
return ST.needsAlignedVGPRs() ? &AMDGPU::VReg_64_Align2RegClass
                              : &AMDGPU::VReg_64RegClass;
2442}

2444const TargetRegisterClass *
2445SIRegisterInfo::getRegClass(unsigned RCID) const {
switch ((int)RCID) {
case AMDGPU::SReg_1RegClassID:
  return getBoolRC();
case AMDGPU::SReg_1_XEXECRegClassID:
  return isWave32 ? &AMDGPU::SReg_32_XM0_XEXECRegClass
    : &AMDGPU::SReg_64_XEXECRegClass;
case -1:
  return nullptr;
default:
  return AMDGPUGenRegisterInfo::getRegClass(RCID);
}
2457}

2459// Find reaching register definition
2460MachineInstr *SIRegisterInfo::findReachingDef(Register Reg, unsigned SubReg,
                                            MachineInstr &Use,
                                            MachineRegisterInfo &MRI,
                                            LiveIntervals *LIS) const {
auto &MDT = LIS->getAnalysis<MachineDominatorTree>();
SlotIndex UseIdx = LIS->getInstructionIndex(Use);
SlotIndex DefIdx;

if (Reg.isVirtual()) {
  if (!LIS->hasInterval(Reg))
    return nullptr;
  LiveInterval &LI = LIS->getInterval(Reg);
  LaneBitmask SubLanes = SubReg ? getSubRegIndexLaneMask(SubReg)
                                : MRI.getMaxLaneMaskForVReg(Reg);
  VNInfo *V = nullptr;
  if (LI.hasSubRanges()) {
    for (auto &S : LI.subranges()) {
      if ((S.LaneMask & SubLanes) == SubLanes) {
        V = S.getVNInfoAt(UseIdx);
        break;
      }
    }
  } else {
    V = LI.getVNInfoAt(UseIdx);
  }
  if (!V)
    return nullptr;
  DefIdx = V->def;
} else {
  // Find last def.
  for (MCRegUnitIterator Units(Reg.asMCReg(), this); Units.isValid();
       ++Units) {
    LiveRange &LR = LIS->getRegUnit(*Units);
    if (VNInfo *V = LR.getVNInfoAt(UseIdx)) {
      if (!DefIdx.isValid() ||
          MDT.dominates(LIS->getInstructionFromIndex(DefIdx),
                        LIS->getInstructionFromIndex(V->def)))
        DefIdx = V->def;
    } else {
      return nullptr;
    }
  }
}

MachineInstr *Def = LIS->getInstructionFromIndex(DefIdx);

if (!Def || !MDT.dominates(Def, &Use))
  return nullptr;

assert(Def->modifiesRegister(Reg, this))(static_cast <bool> (Def->modifiesRegister(Reg, this
)) ? void (0) : __assert_fail ("Def->modifiesRegister(Reg, this)"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 2509, __extension__ __PRETTY_FUNCTION__));

return Def;
2512}

2514MCPhysReg SIRegisterInfo::get32BitRegister(MCPhysReg Reg) const {
assert(getRegSizeInBits(*getPhysRegClass(Reg)) <= 32)(static_cast <bool> (getRegSizeInBits(*getPhysRegClass(
Reg)) <= 32) ? void (0) : __assert_fail ("getRegSizeInBits(*getPhysRegClass(Reg)) <= 32"
, "/build/llvm-toolchain-snapshot-13~++20210726100616+dead50d4427c/llvm/lib/Target/AMDGPU/SIRegisterInfo.cpp"
, 2515, __extension__ __PRETTY_FUNCTION__));

for (const TargetRegisterClass &RC : { AMDGPU::VGPR_32RegClass,
                                       AMDGPU::SReg_32RegClass,
                                       AMDGPU::AGPR_32RegClass } ) {
  if (MCPhysReg Super = getMatchingSuperReg(Reg, AMDGPU::lo16, &RC))
    return Super;
}
if (MCPhysReg Super = getMatchingSuperReg(Reg, AMDGPU::hi16,
                                          &AMDGPU::VGPR_32RegClass)) {
    return Super;
}

return AMDGPU::NoRegister;
2529}

2531bool SIRegisterInfo::isProperlyAlignedRC(const TargetRegisterClass &RC) const {
if (!ST.needsAlignedVGPRs())
  return true;

if (hasVGPRs(&RC))
  return RC.hasSuperClassEq(getVGPRClassForBitWidth(getRegSizeInBits(RC)));
if (hasAGPRs(&RC))
  return RC.hasSuperClassEq(getAGPRClassForBitWidth(getRegSizeInBits(RC)));

return true;
2541}

2543bool SIRegisterInfo::isConstantPhysReg(MCRegister PhysReg) const {
switch (PhysReg) {
case AMDGPU::SGPR_NULL:
case AMDGPU::SRC_SHARED_BASE:
case AMDGPU::SRC_PRIVATE_BASE:
case AMDGPU::SRC_SHARED_LIMIT:
case AMDGPU::SRC_PRIVATE_LIMIT:
  return true;
default:
  return false;
}
2554}

2556ArrayRef<MCPhysReg>
2557SIRegisterInfo::getAllSGPR128(const MachineFunction &MF) const {
return makeArrayRef(AMDGPU::SGPR_128RegClass.begin(),
                    ST.getMaxNumSGPRs(MF) / 4);
2560}

2562ArrayRef<MCPhysReg>
2563SIRegisterInfo::getAllSGPR64(const MachineFunction &MF) const {
return makeArrayRef(AMDGPU::SGPR_64RegClass.begin(),
                    ST.getMaxNumSGPRs(MF) / 2);
2566}

2568ArrayRef<MCPhysReg>
2569SIRegisterInfo::getAllSGPR32(const MachineFunction &MF) const {
return makeArrayRef(AMDGPU::SGPR_32RegClass.begin(), ST.getMaxNumSGPRs(MF));
2571}

←

/usr/lib/gcc/x86_64-linux-gnu/10/../../../../include/c++/10/bits/stl_algobase.h

1// Core algorithmic facilities -*- C++ -*-

3// Copyright (C) 2001-2020 Free Software Foundation, Inc.
4//
5// This file is part of the GNU ISO C++ Library.  This library is free
6// software; you can redistribute it and/or modify it under the
7// terms of the GNU General Public License as published by the
8// Free Software Foundation; either version 3, or (at your option)
9// any later version.

11// This library is distributed in the hope that it will be useful,
12// but WITHOUT ANY WARRANTY; without even the implied warranty of
13// MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
14// GNU General Public License for more details.

16// Under Section 7 of GPL version 3, you are granted additional
17// permissions described in the GCC Runtime Library Exception, version
18// 3.1, as published by the Free Software Foundation.

20// You should have received a copy of the GNU General Public License and
21// a copy of the GCC Runtime Library Exception along with this program;
22// see the files COPYING3 and COPYING.RUNTIME respectively.  If not, see
23// <http://www.gnu.org/licenses/>.

25/*
*
* Copyright (c) 1994
* Hewlett-Packard Company
*
* Permission to use, copy, modify, distribute and sell this software
* and its documentation for any purpose is hereby granted without fee,
* provided that the above copyright notice appear in all copies and
* that both that copyright notice and this permission notice appear
* in supporting documentation.  Hewlett-Packard Company makes no
* representations about the suitability of this software for any
* purpose.  It is provided "as is" without express or implied warranty.
*
*
* Copyright (c) 1996-1998
* Silicon Graphics Computer Systems, Inc.
*
* Permission to use, copy, modify, distribute and sell this software
* and its documentation for any purpose is hereby granted without fee,
* provided that the above copyright notice appear in all copies and
* that both that copyright notice and this permission notice appear
* in supporting documentation.  Silicon Graphics makes no
* representations about the suitability of this software for any
* purpose.  It is provided "as is" without express or implied warranty.
*/

51/** @file bits/stl_algobase.h
*  This is an internal header file, included by other library headers.
*  Do not attempt to use it directly. @headername{algorithm}
*/

56#ifndef _STL_ALGOBASE_H1
57#define _STL_ALGOBASE_H1 1

59#include <bits/c++config.h>
60#include <bits/functexcept.h>
61#include <bits/cpp_type_traits.h>
62#include <ext/type_traits.h>
63#include <ext/numeric_traits.h>
64#include <bits/stl_pair.h>
65#include <bits/stl_iterator_base_types.h>
66#include <bits/stl_iterator_base_funcs.h>
67#include <bits/stl_iterator.h>
68#include <bits/concept_check.h>
69#include <debug/debug.h>
70#include <bits/move.h> // For std::swap
71#include <bits/predefined_ops.h>
72#if __cplusplus201402L >= 201103L
73# include <type_traits>
74#endif
75#if __cplusplus201402L > 201703L
76# include <compare>
77#endif

79namespace std _GLIBCXX_VISIBILITY(default)__attribute__ ((__visibility__ ("default")))
80{
81_GLIBCXX_BEGIN_NAMESPACE_VERSION

/*
 * A constexpr wrapper for __builtin_memcmp.
 * @param __num The number of elements of type _Tp (not bytes).
 */
template<typename _Tp, typename _Up>
  _GLIBCXX14_CONSTEXPRconstexpr
  inline int
  __memcmp(const _Tp* __first1, const _Up* __first2, size_t __num)
  {
92#if __cplusplus201402L >= 201103L
    static_assert(sizeof(_Tp) == sizeof(_Up), "can be compared with memcmp");
94#endif
95#ifdef __cpp_lib_is_constant_evaluated
    if (std::is_constant_evaluated())
{
 for(; __num > 0; ++__first1, ++__first2, --__num)
   if (*__first1 != *__first2)
     return *__first1 < *__first2 ? -1 : 1;
 return 0;
}
    else
104#endif
return __builtin_memcmp(__first1, __first2, sizeof(_Tp) * __num);
  }

108#if __cplusplus201402L < 201103L
// See http://gcc.gnu.org/ml/libstdc++/2004-08/msg00167.html: in a
// nutshell, we are partially implementing the resolution of DR 187,
// when it's safe, i.e., the value_types are equal.
template<bool _BoolType>
  struct __iter_swap
  {
    template<typename _ForwardIterator1, typename _ForwardIterator2>
static void
iter_swap(_ForwardIterator1 __a, _ForwardIterator2 __b)
{
 typedef typename iterator_traits<_ForwardIterator1>::value_type
   _ValueType1;
 _ValueType1 __tmp = *__a;
 *__a = *__b;
 *__b = __tmp;
}
  };

template<>
  struct __iter_swap<true>
  {
    template<typename _ForwardIterator1, typename _ForwardIterator2>
static void
iter_swap(_ForwardIterator1 __a, _ForwardIterator2 __b)
{
 swap(*__a, *__b);
}
  };
137#endif // C++03

/**
 *  @brief Swaps the contents of two iterators.
 *  @ingroup mutating_algorithms
 *  @param  __a  An iterator.
 *  @param  __b  Another iterator.
 *  @return   Nothing.
 *
 *  This function swaps the values pointed to by two iterators, not the
 *  iterators themselves.
*/
template<typename _ForwardIterator1, typename _ForwardIterator2>
  _GLIBCXX20_CONSTEXPR
  inline void
  iter_swap(_ForwardIterator1 __a, _ForwardIterator2 __b)
  {
    // concept requirements
    __glibcxx_function_requires(_Mutable_ForwardIteratorConcept<
		  _ForwardIterator1>)
    __glibcxx_function_requires(_Mutable_ForwardIteratorConcept<
		  _ForwardIterator2>)

160#if __cplusplus201402L < 201103L
    typedef typename iterator_traits<_ForwardIterator1>::value_type
_ValueType1;
    typedef typename iterator_traits<_ForwardIterator2>::value_type
_ValueType2;

    __glibcxx_function_requires(_ConvertibleConcept<_ValueType1,
		  _ValueType2>)
    __glibcxx_function_requires(_ConvertibleConcept<_ValueType2,
		  _ValueType1>)

    typedef typename iterator_traits<_ForwardIterator1>::reference
_ReferenceType1;
    typedef typename iterator_traits<_ForwardIterator2>::reference
_ReferenceType2;
    std::__iter_swap<__are_same<_ValueType1, _ValueType2>::__value
&& __are_same<_ValueType1&, _ReferenceType1>::__value
&& __are_same<_ValueType2&, _ReferenceType2>::__value>::
iter_swap(__a, __b);
179#else
    // _GLIBCXX_RESOLVE_LIB_DEFECTS
    // 187. iter_swap underspecified
    swap(*__a, *__b);
183#endif
  }

/**
 *  @brief Swap the elements of two sequences.
 *  @ingroup mutating_algorithms
 *  @param  __first1  A forward iterator.
 *  @param  __last1   A forward iterator.
 *  @param  __first2  A forward iterator.
 *  @return   An iterator equal to @p first2+(last1-first1).
 *
 *  Swaps each element in the range @p [first1,last1) with the
 *  corresponding element in the range @p [first2,(last1-first1)).
 *  The ranges must not overlap.
*/
template<typename _ForwardIterator1, typename _ForwardIterator2>
  _GLIBCXX20_CONSTEXPR
  _ForwardIterator2
  swap_ranges(_ForwardIterator1 __first1, _ForwardIterator1 __last1,
_ForwardIterator2 __first2)
  {
    // concept requirements
    __glibcxx_function_requires(_Mutable_ForwardIteratorConcept<
		  _ForwardIterator1>)
    __glibcxx_function_requires(_Mutable_ForwardIteratorConcept<
		  _ForwardIterator2>)
    __glibcxx_requires_valid_range(__first1, __last1);

    for (; __first1 != __last1; ++__first1, (void)++__first2)
std::iter_swap(__first1, __first2);
    return __first2;
  }

/**
 *  @brief This does what you think it does.
 *  @ingroup sorting_algorithms
 *  @param  __a  A thing of arbitrary type.
 *  @param  __b  Another thing of arbitrary type.
 *  @return   The lesser of the parameters.
 *
 *  This is the simple classic generic implementation.  It will work on
 *  temporary expressions, since they are only evaluated once, unlike a
 *  preprocessor macro.
*/
template<typename _Tp>
  _GLIBCXX14_CONSTEXPRconstexpr
  inline const _Tp&
  min(const _Tp& __a, const _Tp& __b)
  {
    // concept requirements
    __glibcxx_function_requires(_LessThanComparableConcept<_Tp>)
    //return __b < __a ? __b : __a;
    if (__b < __a)
20
←
Assuming '__b' is >= '__a'→
21
←
Taking false branch→
return __b;
    return __a;
22
←
Returning the value 64 (reference to 'Data.PerVGPR')→
  }

/**
 *  @brief This does what you think it does.
 *  @ingroup sorting_algorithms
 *  @param  __a  A thing of arbitrary type.
 *  @param  __b  Another thing of arbitrary type.
 *  @return   The greater of the parameters.
 *
 *  This is the simple classic generic implementation.  It will work on
 *  temporary expressions, since they are only evaluated once, unlike a
 *  preprocessor macro.
*/
template<typename _Tp>
  _GLIBCXX14_CONSTEXPRconstexpr
  inline const _Tp&
  max(const _Tp& __a, const _Tp& __b)
  {
    // concept requirements
    __glibcxx_function_requires(_LessThanComparableConcept<_Tp>)
    //return  __a < __b ? __b : __a;
    if (__a < __b)
return __b;
    return __a;
  }

/**
 *  @brief This does what you think it does.
 *  @ingroup sorting_algorithms
 *  @param  __a  A thing of arbitrary type.
 *  @param  __b  Another thing of arbitrary type.
 *  @param  __comp  A @link comparison_functors comparison functor@endlink.
 *  @return   The lesser of the parameters.
 *
 *  This will work on temporary expressions, since they are only evaluated
 *  once, unlike a preprocessor macro.
*/
template<typename _Tp, typename _Compare>
  _GLIBCXX14_CONSTEXPRconstexpr
  inline const _Tp&
  min(const _Tp& __a, const _Tp& __b, _Compare __comp)
  {
    //return __comp(__b, __a) ? __b : __a;
    if (__comp(__b, __a))
return __b;
    return __a;
  }

/**
 *  @brief This does what you think it does.
 *  @ingroup sorting_algorithms
 *  @param  __a  A thing of arbitrary type.
 *  @param  __b  Another thing of arbitrary type.
 *  @param  __comp  A @link comparison_functors comparison functor@endlink.
 *  @return   The greater of the parameters.
 *
 *  This will work on temporary expressions, since they are only evaluated
 *  once, unlike a preprocessor macro.
*/
template<typename _Tp, typename _Compare>
  _GLIBCXX14_CONSTEXPRconstexpr
  inline const _Tp&
  max(const _Tp& __a, const _Tp& __b, _Compare __comp)
  {
    //return __comp(__a, __b) ? __b : __a;
    if (__comp(__a, __b))
return __b;
    return __a;
  }

// Fallback implementation of the function in bits/stl_iterator.h used to
// remove the __normal_iterator wrapper. See copy, fill, ...
template<typename _Iterator>
  _GLIBCXX20_CONSTEXPR
  inline _Iterator
  __niter_base(_Iterator __it)
  _GLIBCXX_NOEXCEPT_IF(std::is_nothrow_copy_constructible<_Iterator>::value)noexcept(std::is_nothrow_copy_constructible<_Iterator>::
value)
  { return __it; }

// Reverse the __niter_base transformation to get a
// __normal_iterator back again (this assumes that __normal_iterator
// is only used to wrap random access iterators, like pointers).
template<typename _From, typename _To>
  _GLIBCXX20_CONSTEXPR
  inline _From
  __niter_wrap(_From __from, _To __res)
  { return __from + (__res - std::__niter_base(__from)); }

// No need to wrap, iterator already has the right type.
template<typename _Iterator>
  _GLIBCXX20_CONSTEXPR
  inline _Iterator
  __niter_wrap(const _Iterator&, _Iterator __res)
  { return __res; }

// All of these auxiliary structs serve two purposes.  (1) Replace
// calls to copy with memmove whenever possible.  (Memmove, not memcpy,
// because the input and output ranges are permitted to overlap.)
// (2) If we're using random access iterators, then write the loop as
// a for loop with an explicit count.

template<bool _IsMove, bool _IsSimple, typename _Category>
  struct __copy_move
  {
    template<typename _II, typename _OI>
_GLIBCXX20_CONSTEXPR
static _OI
__copy_m(_II __first, _II __last, _OI __result)
{
 for (; __first != __last; ++__result, (void)++__first)
   *__result = *__first;
 return __result;
}
  };

353#if __cplusplus201402L >= 201103L
template<typename _Category>
  struct __copy_move<true, false, _Category>
  {
    template<typename _II, typename _OI>
_GLIBCXX20_CONSTEXPR
static _OI
__copy_m(_II __first, _II __last, _OI __result)
{
 for (; __first != __last; ++__result, (void)++__first)
   *__result = std::move(*__first);
 return __result;
}
  };
367#endif

template<>
  struct __copy_move<false, false, random_access_iterator_tag>
  {
    template<typename _II, typename _OI>
_GLIBCXX20_CONSTEXPR
static _OI
__copy_m(_II __first, _II __last, _OI __result)
{
 typedef typename iterator_traits<_II>::difference_type _Distance;
 for(_Distance __n = __last - __first; __n > 0; --__n)
   {
     *__result = *__first;
     ++__first;
     ++__result;
   }
 return __result;
}
  };

388#if __cplusplus201402L >= 201103L
template<>
  struct __copy_move<true, false, random_access_iterator_tag>
  {
    template<typename _II, typename _OI>
_GLIBCXX20_CONSTEXPR
static _OI
__copy_m(_II __first, _II __last, _OI __result)
{
 typedef typename iterator_traits<_II>::difference_type _Distance;
 for(_Distance __n = __last - __first; __n > 0; --__n)
   {
     *__result = std::move(*__first);
     ++__first;
     ++__result;
   }
 return __result;
}
  };
407#endif

template<bool _IsMove>
  struct __copy_move<_IsMove, true, random_access_iterator_tag>
  {
    template<typename _Tp>
_GLIBCXX20_CONSTEXPR
static _Tp*
__copy_m(const _Tp* __first, const _Tp* __last, _Tp* __result)
{
417#if __cplusplus201402L >= 201103L
 using __assignable = conditional<_IsMove,
			   is_move_assignable<_Tp>,
			   is_copy_assignable<_Tp>>;
 // trivial types can have deleted assignment
 static_assert( __assignable::type::value, "type is not assignable" );
423#endif
 const ptrdiff_t _Num = __last - __first;
 if (_Num)
   __builtin_memmove(__result, __first, sizeof(_Tp) * _Num);
 return __result + _Num;
}
  };

// Helpers for streambuf iterators (either istream or ostream).
// NB: avoid including <iosfwd>, relatively large.
template<typename _CharT>
  struct char_traits;

template<typename _CharT, typename _Traits>
  class istreambuf_iterator;

template<typename _CharT, typename _Traits>
  class ostreambuf_iterator;

template<bool _IsMove, typename _CharT>
  typename __gnu_cxx::__enable_if<__is_char<_CharT>::__value,
    ostreambuf_iterator<_CharT, char_traits<_CharT> > >::__type
  __copy_move_a2(_CharT*, _CharT*,
   ostreambuf_iterator<_CharT, char_traits<_CharT> >);

template<bool _IsMove, typename _CharT>
  typename __gnu_cxx::__enable_if<__is_char<_CharT>::__value,
    ostreambuf_iterator<_CharT, char_traits<_CharT> > >::__type
  __copy_move_a2(const _CharT*, const _CharT*,
   ostreambuf_iterator<_CharT, char_traits<_CharT> >);

template<bool _IsMove, typename _CharT>
  typename __gnu_cxx::__enable_if<__is_char<_CharT>::__value,
		    _CharT*>::__type
  __copy_move_a2(istreambuf_iterator<_CharT, char_traits<_CharT> >,
   istreambuf_iterator<_CharT, char_traits<_CharT> >, _CharT*);

template<bool _IsMove, typename _II, typename _OI>
  _GLIBCXX20_CONSTEXPR
  inline _OI
  __copy_move_a2(_II __first, _II __last, _OI __result)
  {
    typedef typename iterator_traits<_II>::iterator_category _Category;
466#ifdef __cpp_lib_is_constant_evaluated
    if (std::is_constant_evaluated())
return std::__copy_move<_IsMove, false, _Category>::
 __copy_m(__first, __last, __result);
470#endif
    return std::__copy_move<_IsMove, __memcpyable<_OI, _II>::__value,
	      _Category>::__copy_m(__first, __last, __result);
  }

475_GLIBCXX_BEGIN_NAMESPACE_CONTAINER

template<typename _Tp, typename _Ref, typename _Ptr>
  struct _Deque_iterator;

480_GLIBCXX_END_NAMESPACE_CONTAINER

template<bool _IsMove,
  typename _Tp, typename _Ref, typename _Ptr, typename _OI>
  _OI
  __copy_move_a1(_GLIBCXX_STD_Cstd::_Deque_iterator<_Tp, _Ref, _Ptr>,
   _GLIBCXX_STD_Cstd::_Deque_iterator<_Tp, _Ref, _Ptr>,
   _OI);

template<bool _IsMove,
  typename _ITp, typename _IRef, typename _IPtr, typename _OTp>
  _GLIBCXX_STD_Cstd::_Deque_iterator<_OTp, _OTp&, _OTp*>
  __copy_move_a1(_GLIBCXX_STD_Cstd::_Deque_iterator<_ITp, _IRef, _IPtr>,
   _GLIBCXX_STD_Cstd::_Deque_iterator<_ITp, _IRef, _IPtr>,
   _GLIBCXX_STD_Cstd::_Deque_iterator<_OTp, _OTp&, _OTp*>);

template<bool _IsMove, typename _II, typename _Tp>
  typename __gnu_cxx::__enable_if<
    __is_random_access_iter<_II>::__value,
    _GLIBCXX_STD_Cstd::_Deque_iterator<_Tp, _Tp&, _Tp*> >::__type
  __copy_move_a1(_II, _II, _GLIBCXX_STD_Cstd::_Deque_iterator<_Tp, _Tp&, _Tp*>);

template<bool _IsMove, typename _II, typename _OI>
  _GLIBCXX20_CONSTEXPR
  inline _OI
  __copy_move_a1(_II __first, _II __last, _OI __result)
  { return std::__copy_move_a2<_IsMove>(__first, __last, __result); }

template<bool _IsMove, typename _II, typename _OI>
  _GLIBCXX20_CONSTEXPR
  inline _OI
  __copy_move_a(_II __first, _II __last, _OI __result)
  {
    return std::__niter_wrap(__result,
std::__copy_move_a1<_IsMove>(std::__niter_base(__first),
			     std::__niter_base(__last),
			     std::__niter_base(__result)));
  }

template<bool _IsMove,
  typename _Ite, typename _Seq, typename _Cat, typename _OI>
  _OI
  __copy_move_a(const ::__gnu_debug::_Safe_iterator<_Ite, _Seq, _Cat>&,
  const ::__gnu_debug::_Safe_iterator<_Ite, _Seq, _Cat>&,
  _OI);

template<bool _IsMove,
  typename _II, typename _Ite, typename _Seq, typename _Cat>
  __gnu_debug::_Safe_iterator<_Ite, _Seq, _Cat>
  __copy_move_a(_II, _II,
  const ::__gnu_debug::_Safe_iterator<_Ite, _Seq, _Cat>&);

template<bool _IsMove,
  typename _IIte, typename _ISeq, typename _ICat,
  typename _OIte, typename _OSeq, typename _OCat>
  ::__gnu_debug::_Safe_iterator<_OIte, _OSeq, _OCat>
  __copy_move_a(const ::__gnu_debug::_Safe_iterator<_IIte, _ISeq, _ICat>&,
  const ::__gnu_debug::_Safe_iterator<_IIte, _ISeq, _ICat>&,
  const ::__gnu_debug::_Safe_iterator<_OIte, _OSeq, _OCat>&);

/**
 *  @brief Copies the range [first,last) into result.
 *  @ingroup mutating_algorithms
 *  @param  __first  An input iterator.
 *  @param  __last   An input iterator.
 *  @param  __result An output iterator.
 *  @return   result + (last - first)
 *
 *  This inline function will boil down to a call to @c memmove whenever
 *  possible.  Failing that, if random access iterators are passed, then the
 *  loop count will be known (and therefore a candidate for compiler
 *  optimizations such as unrolling).  Result may not be contained within
 *  [first,last); the copy_backward function should be used instead.
 *
 *  Note that the end of the output range is permitted to be contained
 *  within [first,last).
*/
template<typename _II, typename _OI>
  _GLIBCXX20_CONSTEXPR
  inline _OI
  copy(_II __first, _II __last, _OI __result)
  {
    // concept requirements
    __glibcxx_function_requires(_InputIteratorConcept<_II>)
    __glibcxx_function_requires(_OutputIteratorConcept<_OI,
   typename iterator_traits<_II>::value_type>)
    __glibcxx_requires_can_increment_range(__first, __last, __result);

    return std::__copy_move_a<__is_move_iterator<_II>::__value>
    (std::__miter_base(__first), std::__miter_base(__last), __result);
  }

572#if __cplusplus201402L >= 201103L
/**
 *  @brief Moves the range [first,last) into result.
 *  @ingroup mutating_algorithms
 *  @param  __first  An input iterator.
 *  @param  __last   An input iterator.
 *  @param  __result An output iterator.
 *  @return   result + (last - first)
 *
 *  This inline function will boil down to a call to @c memmove whenever
 *  possible.  Failing that, if random access iterators are passed, then the
 *  loop count will be known (and therefore a candidate for compiler
 *  optimizations such as unrolling).  Result may not be contained within
 *  [first,last); the move_backward function should be used instead.
 *
 *  Note that the end of the output range is permitted to be contained
 *  within [first,last).
*/
template<typename _II, typename _OI>
  _GLIBCXX20_CONSTEXPR
  inline _OI
  move(_II __first, _II __last, _OI __result)
  {
    // concept requirements
    __glibcxx_function_requires(_InputIteratorConcept<_II>)
    __glibcxx_function_requires(_OutputIteratorConcept<_OI,
   typename iterator_traits<_II>::value_type>)
    __glibcxx_requires_can_increment_range(__first, __last, __result);

    return std::__copy_move_a<true>(std::__miter_base(__first),
		      std::__miter_base(__last), __result);
  }

605#define _GLIBCXX_MOVE3(_Tp, _Up, _Vp)std::move(_Tp, _Up, _Vp) std::move(_Tp, _Up, _Vp)
606#else
607#define _GLIBCXX_MOVE3(_Tp, _Up, _Vp)std::move(_Tp, _Up, _Vp) std::copy(_Tp, _Up, _Vp)
608#endif

template<bool _IsMove, bool _IsSimple, typename _Category>
  struct __copy_move_backward
  {
    template<typename _BI1, typename _BI2>
_GLIBCXX20_CONSTEXPR
static _BI2
__copy_move_b(_BI1 __first, _BI1 __last, _BI2 __result)
{
 while (__first != __last)
   *--__result = *--__last;
 return __result;
}
  };

624#if __cplusplus201402L >= 201103L
template<typename _Category>
  struct __copy_move_backward<true, false, _Category>
  {
    template<typename _BI1, typename _BI2>
_GLIBCXX20_CONSTEXPR
static _BI2
__copy_move_b(_BI1 __first, _BI1 __last, _BI2 __result)
{
 while (__first != __last)
   *--__result = std::move(*--__last);
 return __result;
}
  };
638#endif

template<>
  struct __copy_move_backward<false, false, random_access_iterator_tag>
  {
    template<typename _BI1, typename _BI2>
_GLIBCXX20_CONSTEXPR
static _BI2
__copy_move_b(_BI1 __first, _BI1 __last, _BI2 __result)
{
 typename iterator_traits<_BI1>::difference_type
   __n = __last - __first;
 for (; __n > 0; --__n)
   *--__result = *--__last;
 return __result;
}
  };

656#if __cplusplus201402L >= 201103L
template<>
  struct __copy_move_backward<true, false, random_access_iterator_tag>
  {
    template<typename _BI1, typename _BI2>
_GLIBCXX20_CONSTEXPR
static _BI2
__copy_move_b(_BI1 __first, _BI1 __last, _BI2 __result)
{
 typename iterator_traits<_BI1>::difference_type
   __n = __last - __first;
 for (; __n > 0; --__n)
   *--__result = std::move(*--__last);
 return __result;
}
  };
672#endif

template<bool _IsMove>
  struct __copy_move_backward<_IsMove, true, random_access_iterator_tag>
  {
    template<typename _Tp>
_GLIBCXX20_CONSTEXPR
static _Tp*
__copy_move_b(const _Tp* __first, const _Tp* __last, _Tp* __result)
{
682#if __cplusplus201402L >= 201103L
 using __assignable = conditional<_IsMove,
			   is_move_assignable<_Tp>,
			   is_copy_assignable<_Tp>>;
 // trivial types can have deleted assignment
 static_assert( __assignable::type::value, "type is not assignable" );
688#endif
 const ptrdiff_t _Num = __last - __first;
 if (_Num)
   __builtin_memmove(__result - _Num, __first, sizeof(_Tp) * _Num);
 return __result - _Num;
}
  };

template<bool _IsMove, typename _BI1, typename _BI2>
  _GLIBCXX20_CONSTEXPR
  inline _BI2
  __copy_move_backward_a2(_BI1 __first, _BI1 __last, _BI2 __result)
  {
    typedef typename iterator_traits<_BI1>::iterator_category _Category;
702#ifdef __cpp_lib_is_constant_evaluated
    if (std::is_constant_evaluated())
return std::__copy_move_backward<_IsMove, false, _Category>::
 __copy_move_b(__first, __last, __result);
706#endif
    return std::__copy_move_backward<_IsMove,
		       __memcpyable<_BI2, _BI1>::__value,
		       _Category>::__copy_move_b(__first,
						 __last,
						 __result);
  }

template<bool _IsMove, typename _BI1, typename _BI2>
  _GLIBCXX20_CONSTEXPR
  inline _BI2
  __copy_move_backward_a1(_BI1 __first, _BI1 __last, _BI2 __result)
  { return std::__copy_move_backward_a2<_IsMove>(__first, __last, __result); }

template<bool _IsMove,
  typename _Tp, typename _Ref, typename _Ptr, typename _OI>
  _OI
  __copy_move_backward_a1(_GLIBCXX_STD_Cstd::_Deque_iterator<_Tp, _Ref, _Ptr>,
	    _GLIBCXX_STD_Cstd::_Deque_iterator<_Tp, _Ref, _Ptr>,
	    _OI);

template<bool _IsMove,
  typename _ITp, typename _IRef, typename _IPtr, typename _OTp>
  _GLIBCXX_STD_Cstd::_Deque_iterator<_OTp, _OTp&, _OTp*>
  __copy_move_backward_a1(
	_GLIBCXX_STD_Cstd::_Deque_iterator<_ITp, _IRef, _IPtr>,
	_GLIBCXX_STD_Cstd::_Deque_iterator<_ITp, _IRef, _IPtr>,
	_GLIBCXX_STD_Cstd::_Deque_iterator<_OTp, _OTp&, _OTp*>);

template<bool _IsMove, typename _II, typename _Tp>
  typename __gnu_cxx::__enable_if<
    __is_random_access_iter<_II>::__value,
    _GLIBCXX_STD_Cstd::_Deque_iterator<_Tp, _Tp&, _Tp*> >::__type
  __copy_move_backward_a1(_II, _II,
	    _GLIBCXX_STD_Cstd::_Deque_iterator<_Tp, _Tp&, _Tp*>);

template<bool _IsMove, typename _II, typename _OI>
  _GLIBCXX20_CONSTEXPR
  inline _OI
  __copy_move_backward_a(_II __first, _II __last, _OI __result)
  {
    return std::__niter_wrap(__result,
std::__copy_move_backward_a1<_IsMove>
  (std::__niter_base(__first), std::__niter_base(__last),
   std::__niter_base(__result)));
  }

template<bool _IsMove,
  typename _Ite, typename _Seq, typename _Cat, typename _OI>
  _OI
  __copy_move_backward_a(
const ::__gnu_debug::_Safe_iterator<_Ite, _Seq, _Cat>&,
const ::__gnu_debug::_Safe_iterator<_Ite, _Seq, _Cat>&,
_OI);

template<bool _IsMove,
  typename _II, typename _Ite, typename _Seq, typename _Cat>
  __gnu_debug::_Safe_iterator<_Ite, _Seq, _Cat>
  __copy_move_backward_a(_II, _II,
const ::__gnu_debug::_Safe_iterator<_Ite, _Seq, _Cat>&);

template<bool _IsMove,
  typename _IIte, typename _ISeq, typename _ICat,
  typename _OIte, typename _OSeq, typename _OCat>
  ::__gnu_debug::_Safe_iterator<_OIte, _OSeq, _OCat>
  __copy_move_backward_a(
const ::__gnu_debug::_Safe_iterator<_IIte, _ISeq, _ICat>&,
const ::__gnu_debug::_Safe_iterator<_IIte, _ISeq, _ICat>&,
const ::__gnu_debug::_Safe_iterator<_OIte, _OSeq, _OCat>&);

/**
 *  @brief Copies the range [first,last) into result.
 *  @ingroup mutating_algorithms
 *  @param  __first  A bidirectional iterator.
 *  @param  __last   A bidirectional iterator.
 *  @param  __result A bidirectional iterator.
 *  @return   result - (last - first)
 *
 *  The function has the same effect as copy, but starts at the end of the
 *  range and works its way to the start, returning the start of the result.
 *  This inline function will boil down to a call to @c memmove whenever
 *  possible.  Failing that, if random access iterators are passed, then the
 *  loop count will be known (and therefore a candidate for compiler
 *  optimizations such as unrolling).
 *
 *  Result may not be in the range (first,last].  Use copy instead.  Note
 *  that the start of the output range may overlap [first,last).
*/
template<typename _BI1, typename _BI2>
  _GLIBCXX20_CONSTEXPR
  inline _BI2
  copy_backward(_BI1 __first, _BI1 __last, _BI2 __result)
  {
    // concept requirements
    __glibcxx_function_requires(_BidirectionalIteratorConcept<_BI1>)
    __glibcxx_function_requires(_Mutable_BidirectionalIteratorConcept<_BI2>)
    __glibcxx_function_requires(_ConvertibleConcept<
   typename iterator_traits<_BI1>::value_type,
   typename iterator_traits<_BI2>::value_type>)
    __glibcxx_requires_can_decrement_range(__first, __last, __result);

    return std::__copy_move_backward_a<__is_move_iterator<_BI1>::__value>
    (std::__miter_base(__first), std::__miter_base(__last), __result);
  }

811#if __cplusplus201402L >= 201103L
/**
 *  @brief Moves the range [first,last) into result.
 *  @ingroup mutating_algorithms
 *  @param  __first  A bidirectional iterator.
 *  @param  __last   A bidirectional iterator.
 *  @param  __result A bidirectional iterator.
 *  @return   result - (last - first)
 *
 *  The function has the same effect as move, but starts at the end of the
 *  range and works its way to the start, returning the start of the result.
 *  This inline function will boil down to a call to @c memmove whenever
 *  possible.  Failing that, if random access iterators are passed, then the
 *  loop count will be known (and therefore a candidate for compiler
 *  optimizations such as unrolling).
 *
 *  Result may not be in the range (first,last].  Use move instead.  Note
 *  that the start of the output range may overlap [first,last).
*/
template<typename _BI1, typename _BI2>
  _GLIBCXX20_CONSTEXPR
  inline _BI2
  move_backward(_BI1 __first, _BI1 __last, _BI2 __result)
  {
    // concept requirements
    __glibcxx_function_requires(_BidirectionalIteratorConcept<_BI1>)
    __glibcxx_function_requires(_Mutable_BidirectionalIteratorConcept<_BI2>)
    __glibcxx_function_requires(_ConvertibleConcept<
   typename iterator_traits<_BI1>::value_type,
   typename iterator_traits<_BI2>::value_type>)
    __glibcxx_requires_can_decrement_range(__first, __last, __result);

    return std::__copy_move_backward_a<true>(std::__miter_base(__first),
			       std::__miter_base(__last),
			       __result);
  }

848#define _GLIBCXX_MOVE_BACKWARD3(_Tp, _Up, _Vp)std::move_backward(_Tp, _Up, _Vp) std::move_backward(_Tp, _Up, _Vp)
849#else
850#define _GLIBCXX_MOVE_BACKWARD3(_Tp, _Up, _Vp)std::move_backward(_Tp, _Up, _Vp) std::copy_backward(_Tp, _Up, _Vp)
851#endif

template<typename _ForwardIterator, typename _Tp>
  _GLIBCXX20_CONSTEXPR
  inline typename
  __gnu_cxx::__enable_if<!__is_scalar<_Tp>::__value, void>::__type
  __fill_a1(_ForwardIterator __first, _ForwardIterator __last,
     const _Tp& __value)
  {
    for (; __first != __last; ++__first)
*__first = __value;
  }

template<typename _ForwardIterator, typename _Tp>
  _GLIBCXX20_CONSTEXPR
  inline typename
  __gnu_cxx::__enable_if<__is_scalar<_Tp>::__value, void>::__type
  __fill_a1(_ForwardIterator __first, _ForwardIterator __last,
     const _Tp& __value)
  {
    const _Tp __tmp = __value;
    for (; __first != __last; ++__first)
*__first = __tmp;
  }

// Specialization: for char types we can use memset.
template<typename _Tp>
  _GLIBCXX20_CONSTEXPR
  inline typename
  __gnu_cxx::__enable_if<__is_byte<_Tp>::__value, void>::__type
  __fill_a1(_Tp* __first, _Tp* __last, const _Tp& __c)
  {
    const _Tp __tmp = __c;
884#if __cpp_lib_is_constant_evaluated
    if (std::is_constant_evaluated())
{
 for (; __first != __last; ++__first)
   *__first = __tmp;
 return;
}
891#endif
    if (const size_t __len = __last - __first)
__builtin_memset(__first, static_cast<unsigned char>(__tmp), __len);
  }

template<typename _Ite, typename _Cont, typename _Tp>
  _GLIBCXX20_CONSTEXPR
  inline void
  __fill_a1(::__gnu_cxx::__normal_iterator<_Ite, _Cont> __first,
     ::__gnu_cxx::__normal_iterator<_Ite, _Cont> __last,
     const _Tp& __value)
  { std::__fill_a1(__first.base(), __last.base(), __value); }

template<typename _Tp, typename _VTp>
  void
  __fill_a1(const _GLIBCXX_STD_Cstd::_Deque_iterator<_Tp, _Tp&, _Tp*>&,
     const _GLIBCXX_STD_Cstd::_Deque_iterator<_Tp, _Tp&, _Tp*>&,
     const _VTp&);

template<typename _FIte, typename _Tp>
  _GLIBCXX20_CONSTEXPR
  inline void
  __fill_a(_FIte __first, _FIte __last, const _Tp& __value)
  { std::__fill_a1(__first, __last, __value); }

template<typename _Ite, typename _Seq, typename _Cat, typename _Tp>
  void
  __fill_a(const ::__gnu_debug::_Safe_iterator<_Ite, _Seq, _Cat>&,
    const ::__gnu_debug::_Safe_iterator<_Ite, _Seq, _Cat>&,
    const _Tp&);

/**
 *  @brief Fills the range [first,last) with copies of value.
 *  @ingroup mutating_algorithms
 *  @param  __first  A forward iterator.
 *  @param  __last   A forward iterator.
 *  @param  __value  A reference-to-const of arbitrary type.
 *  @return   Nothing.
 *
 *  This function fills a range with copies of the same value.  For char
 *  types filling contiguous areas of memory, this becomes an inline call
 *  to @c memset or @c wmemset.
*/
template<typename _ForwardIterator, typename _Tp>
  _GLIBCXX20_CONSTEXPR
  inline void
  fill(_ForwardIterator __first, _ForwardIterator __last, const _Tp& __value)
  {
    // concept requirements
    __glibcxx_function_requires(_Mutable_ForwardIteratorConcept<
		  _ForwardIterator>)
    __glibcxx_requires_valid_range(__first, __last);

    std::__fill_a(__first, __last, __value);
  }

// Used by fill_n, generate_n, etc. to convert _Size to an integral type:
inline _GLIBCXX_CONSTEXPRconstexpr int
__size_to_integer(int __n) { return __n; }
inline _GLIBCXX_CONSTEXPRconstexpr unsigned
__size_to_integer(unsigned __n) { return __n; }
inline _GLIBCXX_CONSTEXPRconstexpr long
__size_to_integer(long __n) { return __n; }
inline _GLIBCXX_CONSTEXPRconstexpr unsigned long
__size_to_integer(unsigned long __n) { return __n; }
inline _GLIBCXX_CONSTEXPRconstexpr long long
__size_to_integer(long long __n) { return __n; }
inline _GLIBCXX_CONSTEXPRconstexpr unsigned long long
__size_to_integer(unsigned long long __n) { return __n; }

961#if defined(__GLIBCXX_TYPE_INT_N_0)
inline _GLIBCXX_CONSTEXPRconstexpr __GLIBCXX_TYPE_INT_N_0
__size_to_integer(__GLIBCXX_TYPE_INT_N_0 __n) { return __n; }
inline _GLIBCXX_CONSTEXPRconstexpr unsigned __GLIBCXX_TYPE_INT_N_0
__size_to_integer(unsigned __GLIBCXX_TYPE_INT_N_0 __n) { return __n; }
966#endif
967#if defined(__GLIBCXX_TYPE_INT_N_1)
inline _GLIBCXX_CONSTEXPRconstexpr __GLIBCXX_TYPE_INT_N_1
__size_to_integer(__GLIBCXX_TYPE_INT_N_1 __n) { return __n; }
inline _GLIBCXX_CONSTEXPRconstexpr unsigned __GLIBCXX_TYPE_INT_N_1
__size_to_integer(unsigned __GLIBCXX_TYPE_INT_N_1 __n) { return __n; }
972#endif
973#if defined(__GLIBCXX_TYPE_INT_N_2)
inline _GLIBCXX_CONSTEXPRconstexpr __GLIBCXX_TYPE_INT_N_2
__size_to_integer(__GLIBCXX_TYPE_INT_N_2 __n) { return __n; }
inline _GLIBCXX_CONSTEXPRconstexpr unsigned __GLIBCXX_TYPE_INT_N_2
__size_to_integer(unsigned __GLIBCXX_TYPE_INT_N_2 __n) { return __n; }
978#endif
979#if defined(__GLIBCXX_TYPE_INT_N_3)
inline _GLIBCXX_CONSTEXPRconstexpr unsigned __GLIBCXX_TYPE_INT_N_3
__size_to_integer(__GLIBCXX_TYPE_INT_N_3 __n) { return __n; }
inline _GLIBCXX_CONSTEXPRconstexpr __GLIBCXX_TYPE_INT_N_3
__size_to_integer(unsigned __GLIBCXX_TYPE_INT_N_3 __n) { return __n; }
984#endif

inline _GLIBCXX_CONSTEXPRconstexpr long long
__size_to_integer(float __n) { return __n; }
inline _GLIBCXX_CONSTEXPRconstexpr long long
__size_to_integer(double __n) { return __n; }
inline _GLIBCXX_CONSTEXPRconstexpr long long
__size_to_integer(long double __n) { return __n; }
992#if !defined(__STRICT_ANSI__1) && defined(_GLIBCXX_USE_FLOAT1281) && !defined(__CUDACC__)
inline _GLIBCXX_CONSTEXPRconstexpr long long
__size_to_integer(__float128 __n) { return __n; }
995#endif

template<typename _OutputIterator, typename _Size, typename _Tp>
  _GLIBCXX20_CONSTEXPR
  inline typename
  __gnu_cxx::__enable_if<!__is_scalar<_Tp>::__value, _OutputIterator>::__type
  __fill_n_a1(_OutputIterator __first, _Size __n, const _Tp& __value)
  {
    for (; __n > 0; --__n, (void) ++__first)
*__first = __value;
    return __first;
  }

template<typename _OutputIterator, typename _Size, typename _Tp>
  _GLIBCXX20_CONSTEXPR
  inline typename
  __gnu_cxx::__enable_if<__is_scalar<_Tp>::__value, _OutputIterator>::__type
  __fill_n_a1(_OutputIterator __first, _Size __n, const _Tp& __value)
  {
    const _Tp __tmp = __value;
    for (; __n > 0; --__n, (void) ++__first)
*__first = __tmp;
    return __first;
  }

template<typename _Ite, typename _Seq, typename _Cat, typename _Size,
  typename _Tp>
  ::__gnu_debug::_Safe_iterator<_Ite, _Seq, _Cat>
  __fill_n_a(const ::__gnu_debug::_Safe_iterator<_Ite, _Seq, _Cat>& __first,
      _Size __n, const _Tp& __value,
      std::input_iterator_tag);

template<typename _OutputIterator, typename _Size, typename _Tp>
  _GLIBCXX20_CONSTEXPR
  inline _OutputIterator
  __fill_n_a(_OutputIterator __first, _Size __n, const _Tp& __value,
      std::output_iterator_tag)
  {
1033#if __cplusplus201402L >= 201103L
    static_assert(is_integral<_Size>{}, "fill_n must pass integral size");
1035#endif
    return __fill_n_a1(__first, __n, __value);
  }

template<typename _OutputIterator, typename _Size, typename _Tp>
  _GLIBCXX20_CONSTEXPR
  inline _OutputIterator
  __fill_n_a(_OutputIterator __first, _Size __n, const _Tp& __value,
      std::input_iterator_tag)
  {
1045#if __cplusplus201402L >= 201103L
    static_assert(is_integral<_Size>{}, "fill_n must pass integral size");
1047#endif
    return __fill_n_a1(__first, __n, __value);
  }

template<typename _OutputIterator, typename _Size, typename _Tp>
  _GLIBCXX20_CONSTEXPR
  inline _OutputIterator
  __fill_n_a(_OutputIterator __first, _Size __n, const _Tp& __value,
      std::random_access_iterator_tag)
  {
1057#if __cplusplus201402L >= 201103L
    static_assert(is_integral<_Size>{}, "fill_n must pass integral size");
1059#endif
    if (__n <= 0)
return __first;

    __glibcxx_requires_can_increment(__first, __n);

    std::__fill_a(__first, __first + __n, __value);
    return __first + __n;
  }

/**
 *  @brief Fills the range [first,first+n) with copies of value.
 *  @ingroup mutating_algorithms
 *  @param  __first  An output iterator.
 *  @param  __n      The count of copies to perform.
 *  @param  __value  A reference-to-const of arbitrary type.
 *  @return   The iterator at first+n.
 *
 *  This function fills a range with copies of the same value.  For char
 *  types filling contiguous areas of memory, this becomes an inline call
 *  to @c memset or @c wmemset.
 *
 *  If @p __n is negative, the function does nothing.
*/
// _GLIBCXX_RESOLVE_LIB_DEFECTS
// DR 865. More algorithms that throw away information
// DR 426. search_n(), fill_n(), and generate_n() with negative n
template<typename _OI, typename _Size, typename _Tp>
  _GLIBCXX20_CONSTEXPR
  inline _OI
  fill_n(_OI __first, _Size __n, const _Tp& __value)
  {
    // concept requirements
    __glibcxx_function_requires(_OutputIteratorConcept<_OI, _Tp>)

    return std::__fill_n_a(__first, std::__size_to_integer(__n), __value,
	       std::__iterator_category(__first));
  }

template<bool _BoolType>
  struct __equal
  {
    template<typename _II1, typename _II2>
_GLIBCXX20_CONSTEXPR
static bool
equal(_II1 __first1, _II1 __last1, _II2 __first2)
{
 for (; __first1 != __last1; ++__first1, (void) ++__first2)
   if (!(*__first1 == *__first2))
     return false;
 return true;
}
  };

template<>
  struct __equal<true>
  {
    template<typename _Tp>
_GLIBCXX20_CONSTEXPR
static bool
equal(const _Tp* __first1, const _Tp* __last1, const _Tp* __first2)
{
 if (const size_t __len = (__last1 - __first1))
   return !std::__memcmp(__first1, __first2, __len);
 return true;
}
  };

template<typename _Tp, typename _Ref, typename _Ptr, typename _II>
  typename __gnu_cxx::__enable_if<
    __is_random_access_iter<_II>::__value, bool>::__type
  __equal_aux1(_GLIBCXX_STD_Cstd::_Deque_iterator<_Tp, _Ref, _Ptr>,
 _GLIBCXX_STD_Cstd::_Deque_iterator<_Tp, _Ref, _Ptr>,
 _II);

template<typename _Tp1, typename _Ref1, typename _Ptr1,
  typename _Tp2, typename _Ref2, typename _Ptr2>
  bool
  __equal_aux1(_GLIBCXX_STD_Cstd::_Deque_iterator<_Tp1, _Ref1, _Ptr1>,
 _GLIBCXX_STD_Cstd::_Deque_iterator<_Tp1, _Ref1, _Ptr1>,
 _GLIBCXX_STD_Cstd::_Deque_iterator<_Tp2, _Ref2, _Ptr2>);

template<typename _II, typename _Tp, typename _Ref, typename _Ptr>
  typename __gnu_cxx::__enable_if<
    __is_random_access_iter<_II>::__value, bool>::__type
  __equal_aux1(_II, _II,
_GLIBCXX_STD_Cstd::_Deque_iterator<_Tp, _Ref, _Ptr>);

template<typename _II1, typename _II2>
  _GLIBCXX20_CONSTEXPR
  inline bool
  __equal_aux1(_II1 __first1, _II1 __last1, _II2 __first2)
  {
    typedef typename iterator_traits<_II1>::value_type _ValueType1;
    const bool __simple = ((__is_integer<_ValueType1>::__value
	      || __is_pointer<_ValueType1>::__value)
	     && __memcmpable<_II1, _II2>::__value);
    return std::__equal<__simple>::equal(__first1, __last1, __first2);
  }

template<typename _II1, typename _II2>
  _GLIBCXX20_CONSTEXPR
  inline bool
  __equal_aux(_II1 __first1, _II1 __last1, _II2 __first2)
  {
    return std::__equal_aux1(std::__niter_base(__first1),
	       std::__niter_base(__last1),
	       std::__niter_base(__first2));
  }

template<typename _II1, typename _Seq1, typename _Cat1, typename _II2>
  bool
  __equal_aux(const ::__gnu_debug::_Safe_iterator<_II1, _Seq1, _Cat1>&,
const ::__gnu_debug::_Safe_iterator<_II1, _Seq1, _Cat1>&,
_II2);

template<typename _II1, typename _II2, typename _Seq2, typename _Cat2>
  bool
  __equal_aux(_II1, _II1,
const ::__gnu_debug::_Safe_iterator<_II2, _Seq2, _Cat2>&);

template<typename _II1, typename _Seq1, typename _Cat1,
  typename _II2, typename _Seq2, typename _Cat2>
  bool
  __equal_aux(const ::__gnu_debug::_Safe_iterator<_II1, _Seq1, _Cat1>&,
const ::__gnu_debug::_Safe_iterator<_II1, _Seq1, _Cat1>&,
const ::__gnu_debug::_Safe_iterator<_II2, _Seq2, _Cat2>&);

template<typename, typename>
  struct __lc_rai
  {
    template<typename _II1, typename _II2>
_GLIBCXX20_CONSTEXPR
static _II1
__newlast1(_II1, _II1 __last1, _II2, _II2)
{ return __last1; }

    template<typename _II>
_GLIBCXX20_CONSTEXPR
static bool
__cnd2(_II __first, _II __last)
{ return __first != __last; }
  };

template<>
  struct __lc_rai<random_access_iterator_tag, random_access_iterator_tag>
  {
    template<typename _RAI1, typename _RAI2>
_GLIBCXX20_CONSTEXPR
static _RAI1
__newlast1(_RAI1 __first1, _RAI1 __last1,
   _RAI2 __first2, _RAI2 __last2)
{
 const typename iterator_traits<_RAI1>::difference_type
   __diff1 = __last1 - __first1;
 const typename iterator_traits<_RAI2>::difference_type
   __diff2 = __last2 - __first2;
 return __diff2 < __diff1 ? __first1 + __diff2 : __last1;
}

    template<typename _RAI>
static _GLIBCXX20_CONSTEXPR bool
__cnd2(_RAI, _RAI)
{ return true; }
  };

template<typename _II1, typename _II2, typename _Compare>
  _GLIBCXX20_CONSTEXPR
  bool
  __lexicographical_compare_impl(_II1 __first1, _II1 __last1,
		   _II2 __first2, _II2 __last2,
		   _Compare __comp)
  {
    typedef typename iterator_traits<_II1>::iterator_category _Category1;
    typedef typename iterator_traits<_II2>::iterator_category _Category2;
    typedef std::__lc_rai<_Category1, _Category2> __rai_type;

    __last1 = __rai_type::__newlast1(__first1, __last1, __first2, __last2);
    for (; __first1 != __last1 && __rai_type::__cnd2(__first2, __last2);
  ++__first1, (void)++__first2)
{
 if (__comp(__first1, __first2))
   return true;
 if (__comp(__first2, __first1))
   return false;
}
    return __first1 == __last1 && __first2 != __last2;
  }

template<bool _BoolType>
  struct __lexicographical_compare
  {
    template<typename _II1, typename _II2>
_GLIBCXX20_CONSTEXPR
static bool
__lc(_II1 __first1, _II1 __last1, _II2 __first2, _II2 __last2)
{
 using __gnu_cxx::__ops::__iter_less_iter;
 return std::__lexicographical_compare_impl(__first1, __last1,
				     __first2, __last2,
				     __iter_less_iter());
}
  };

template<>
  struct __lexicographical_compare<true>
  {
    template<typename _Tp, typename _Up>
_GLIBCXX20_CONSTEXPR
static bool
__lc(const _Tp* __first1, const _Tp* __last1,
    const _Up* __first2, const _Up* __last2)
{
 const size_t __len1 = __last1 - __first1;
 const size_t __len2 = __last2 - __first2;
 if (const size_t __len = std::min(__len1, __len2))
   if (int __result = std::__memcmp(__first1, __first2, __len))
     return __result < 0;
 return __len1 < __len2;
}
  };

template<typename _II1, typename _II2>
  _GLIBCXX20_CONSTEXPR
  inline bool
  __lexicographical_compare_aux(_II1 __first1, _II1 __last1,
		  _II2 __first2, _II2 __last2)
  {
    typedef typename iterator_traits<_II1>::value_type _ValueType1;
    typedef typename iterator_traits<_II2>::value_type _ValueType2;
    const bool __simple =
(__is_memcmp_ordered_with<_ValueType1, _ValueType2>::__value
&& __is_pointer<_II1>::__value
&& __is_pointer<_II2>::__value
1293#if __cplusplus201402L > 201703L && __cpp_lib_concepts
// For C++20 iterator_traits<volatile T*>::value_type is non-volatile
// so __is_byte<T> could be true, but we can't use memcmp with
// volatile data.
&& !is_volatile_v<remove_reference_t<iter_reference_t<_II1>>>
&& !is_volatile_v<remove_reference_t<iter_reference_t<_II2>>>
1299#endif
);

    return std::__lexicographical_compare<__simple>::__lc(__first1, __last1,
					    __first2, __last2);
  }

template<typename _ForwardIterator, typename _Tp, typename _Compare>
  _GLIBCXX20_CONSTEXPR
  _ForwardIterator
  __lower_bound(_ForwardIterator __first, _ForwardIterator __last,
  const _Tp& __val, _Compare __comp)
  {
    typedef typename iterator_traits<_ForwardIterator>::difference_type
_DistanceType;

    _DistanceType __len = std::distance(__first, __last);

    while (__len > 0)
{
 _DistanceType __half = __len >> 1;
 _ForwardIterator __middle = __first;
 std::advance(__middle, __half);
 if (__comp(__middle, __val))
   {
     __first = __middle;
     ++__first;
     __len = __len - __half - 1;
   }
 else
   __len = __half;
}
    return __first;
  }

/**
 *  @brief Finds the first position in which @a val could be inserted
 *         without changing the ordering.
 *  @param  __first   An iterator.
 *  @param  __last    Another iterator.
 *  @param  __val     The search term.
 *  @return         An iterator pointing to the first element <em>not less
 *                  than</em> @a val, or end() if every element is less than
 *                  @a val.
 *  @ingroup binary_search_algorithms
*/
template<typename _ForwardIterator, typename _Tp>
  _GLIBCXX20_CONSTEXPR
  inline _ForwardIterator
  lower_bound(_ForwardIterator __first, _ForwardIterator __last,
const _Tp& __val)
  {
    // concept requirements
    __glibcxx_function_requires(_ForwardIteratorConcept<_ForwardIterator>)
    __glibcxx_function_requires(_LessThanOpConcept<
   typename iterator_traits<_ForwardIterator>::value_type, _Tp>)
    __glibcxx_requires_partitioned_lower(__first, __last, __val);

    return std::__lower_bound(__first, __last, __val,
		__gnu_cxx::__ops::__iter_less_val());
  }

/// This is a helper function for the sort routines and for random.tcc.
//  Precondition: __n > 0.
inline _GLIBCXX_CONSTEXPRconstexpr int
__lg(int __n)
{ return (int)sizeof(int) * __CHAR_BIT__8  - 1 - __builtin_clz(__n); }

inline _GLIBCXX_CONSTEXPRconstexpr unsigned
__lg(unsigned __n)
{ return (int)sizeof(int) * __CHAR_BIT__8  - 1 - __builtin_clz(__n); }

inline _GLIBCXX_CONSTEXPRconstexpr long
__lg(long __n)
{ return (int)sizeof(long) * __CHAR_BIT__8 - 1 - __builtin_clzl(__n); }

inline _GLIBCXX_CONSTEXPRconstexpr unsigned long
__lg(unsigned long __n)
{ return (int)sizeof(long) * __CHAR_BIT__8 - 1 - __builtin_clzl(__n); }

inline _GLIBCXX_CONSTEXPRconstexpr long long
__lg(long long __n)
{ return (int)sizeof(long long) * __CHAR_BIT__8 - 1 - __builtin_clzll(__n); }

inline _GLIBCXX_CONSTEXPRconstexpr unsigned long long
__lg(unsigned long long __n)
{ return (int)sizeof(long long) * __CHAR_BIT__8 - 1 - __builtin_clzll(__n); }

1387_GLIBCXX_BEGIN_NAMESPACE_ALGO

/**
 *  @brief Tests a range for element-wise equality.
 *  @ingroup non_mutating_algorithms
 *  @param  __first1  An input iterator.
 *  @param  __last1   An input iterator.
 *  @param  __first2  An input iterator.
 *  @return   A boolean true or false.
 *
 *  This compares the elements of two ranges using @c == and returns true or
 *  false depending on whether all of the corresponding elements of the
 *  ranges are equal.
*/
template<typename _II1, typename _II2>
  _GLIBCXX20_CONSTEXPR
  inline bool
  equal(_II1 __first1, _II1 __last1, _II2 __first2)
  {
    // concept requirements
    __glibcxx_function_requires(_InputIteratorConcept<_II1>)
    __glibcxx_function_requires(_InputIteratorConcept<_II2>)
    __glibcxx_function_requires(_EqualOpConcept<
   typename iterator_traits<_II1>::value_type,
   typename iterator_traits<_II2>::value_type>)
    __glibcxx_requires_can_increment_range(__first1, __last1, __first2);

    return std::__equal_aux(__first1, __last1, __first2);
  }

/**
 *  @brief Tests a range for element-wise equality.
 *  @ingroup non_mutating_algorithms
 *  @param  __first1  An input iterator.
 *  @param  __last1   An input iterator.
 *  @param  __first2  An input iterator.
 *  @param __binary_pred A binary predicate @link functors
 *                  functor@endlink.
 *  @return         A boolean true or false.
 *
 *  This compares the elements of two ranges using the binary_pred
 *  parameter, and returns true or
 *  false depending on whether all of the corresponding elements of the
 *  ranges are equal.
*/
template<typename _IIter1, typename _IIter2, typename _BinaryPredicate>
  _GLIBCXX20_CONSTEXPR
  inline bool
  equal(_IIter1 __first1, _IIter1 __last1,
 _IIter2 __first2, _BinaryPredicate __binary_pred)
  {
    // concept requirements
    __glibcxx_function_requires(_InputIteratorConcept<_IIter1>)
    __glibcxx_function_requires(_InputIteratorConcept<_IIter2>)
    __glibcxx_requires_valid_range(__first1, __last1);

    for (; __first1 != __last1; ++__first1, (void)++__first2)
if (!bool(__binary_pred(*__first1, *__first2)))
 return false;
    return true;
  }

1449#if __cplusplus201402L >= 201103L
// 4-iterator version of std::equal<It1, It2> for use in C++11.
template<typename _II1, typename _II2>
  _GLIBCXX20_CONSTEXPR
  inline bool
  __equal4(_II1 __first1, _II1 __last1, _II2 __first2, _II2 __last2)
  {
    using _RATag = random_access_iterator_tag;
    using _Cat1 = typename iterator_traits<_II1>::iterator_category;
    using _Cat2 = typename iterator_traits<_II2>::iterator_category;
    using _RAIters = __and_<is_same<_Cat1, _RATag>, is_same<_Cat2, _RATag>>;
    if (_RAIters())
{
 auto __d1 = std::distance(__first1, __last1);
 auto __d2 = std::distance(__first2, __last2);
 if (__d1 != __d2)
   return false;
 return _GLIBCXX_STD_Astd::equal(__first1, __last1, __first2);
}

    for (; __first1 != __last1 && __first2 != __last2;
 ++__first1, (void)++__first2)
if (!(*__first1 == *__first2))
 return false;
    return __first1 == __last1 && __first2 == __last2;
  }

// 4-iterator version of std::equal<It1, It2, BinaryPred> for use in C++11.
template<typename _II1, typename _II2, typename _BinaryPredicate>
  _GLIBCXX20_CONSTEXPR
  inline bool
  __equal4(_II1 __first1, _II1 __last1, _II2 __first2, _II2 __last2,
    _BinaryPredicate __binary_pred)
  {
    using _RATag = random_access_iterator_tag;
    using _Cat1 = typename iterator_traits<_II1>::iterator_category;
    using _Cat2 = typename iterator_traits<_II2>::iterator_category;
    using _RAIters = __and_<is_same<_Cat1, _RATag>, is_same<_Cat2, _RATag>>;
    if (_RAIters())
{
 auto __d1 = std::distance(__first1, __last1);
 auto __d2 = std::distance(__first2, __last2);
 if (__d1 != __d2)
   return false;
 return _GLIBCXX_STD_Astd::equal(__first1, __last1, __first2,
		       __binary_pred);
}

    for (; __first1 != __last1 && __first2 != __last2;
 ++__first1, (void)++__first2)
if (!bool(__binary_pred(*__first1, *__first2)))
 return false;
    return __first1 == __last1 && __first2 == __last2;
  }
1503#endif // C++11

1505#if __cplusplus201402L > 201103L

1507#define __cpp_lib_robust_nonmodifying_seq_ops201304 201304

/**
 *  @brief Tests a range for element-wise equality.
 *  @ingroup non_mutating_algorithms
 *  @param  __first1  An input iterator.
 *  @param  __last1   An input iterator.
 *  @param  __first2  An input iterator.
 *  @param  __last2   An input iterator.
 *  @return   A boolean true or false.
 *
 *  This compares the elements of two ranges using @c == and returns true or
 *  false depending on whether all of the corresponding elements of the
 *  ranges are equal.
*/
template<typename _II1, typename _II2>
  _GLIBCXX20_CONSTEXPR
  inline bool
  equal(_II1 __first1, _II1 __last1, _II2 __first2, _II2 __last2)
  {
    // concept requirements
    __glibcxx_function_requires(_InputIteratorConcept<_II1>)
    __glibcxx_function_requires(_InputIteratorConcept<_II2>)
    __glibcxx_function_requires(_EqualOpConcept<
   typename iterator_traits<_II1>::value_type,
   typename iterator_traits<_II2>::value_type>)
    __glibcxx_requires_valid_range(__first1, __last1);
    __glibcxx_requires_valid_range(__first2, __last2);

    return _GLIBCXX_STD_Astd::__equal4(__first1, __last1, __first2, __last2);
  }

/**
 *  @brief Tests a range for element-wise equality.
 *  @ingroup non_mutating_algorithms
 *  @param  __first1  An input iterator.
 *  @param  __last1   An input iterator.
 *  @param  __first2  An input iterator.
 *  @param  __last2   An input iterator.
 *  @param __binary_pred A binary predicate @link functors
 *                  functor@endlink.
 *  @return         A boolean true or false.
 *
 *  This compares the elements of two ranges using the binary_pred
 *  parameter, and returns true or
 *  false depending on whether all of the corresponding elements of the
 *  ranges are equal.
*/
template<typename _IIter1, typename _IIter2, typename _BinaryPredicate>
  _GLIBCXX20_CONSTEXPR
  inline bool
  equal(_IIter1 __first1, _IIter1 __last1,
 _IIter2 __first2, _IIter2 __last2, _BinaryPredicate __binary_pred)
  {
    // concept requirements
    __glibcxx_function_requires(_InputIteratorConcept<_IIter1>)
    __glibcxx_function_requires(_InputIteratorConcept<_IIter2>)
    __glibcxx_requires_valid_range(__first1, __last1);
    __glibcxx_requires_valid_range(__first2, __last2);

    return _GLIBCXX_STD_Astd::__equal4(__first1, __last1, __first2, __last2,
		      __binary_pred);
  }
1570#endif // C++14

/**
 *  @brief Performs @b dictionary comparison on ranges.
 *  @ingroup sorting_algorithms
 *  @param  __first1  An input iterator.
 *  @param  __last1   An input iterator.
 *  @param  __first2  An input iterator.
 *  @param  __last2   An input iterator.
 *  @return   A boolean true or false.
 *
 *  <em>Returns true if the sequence of elements defined by the range
 *  [first1,last1) is lexicographically less than the sequence of elements
 *  defined by the range [first2,last2).  Returns false otherwise.</em>
 *  (Quoted from [25.3.8]/1.)  If the iterators are all character pointers,
 *  then this is an inline call to @c memcmp.
*/
template<typename _II1, typename _II2>
  _GLIBCXX20_CONSTEXPR
  inline bool
  lexicographical_compare(_II1 __first1, _II1 __last1,
	    _II2 __first2, _II2 __last2)
  {
1593#ifdef _GLIBCXX_CONCEPT_CHECKS
    // concept requirements
    typedef typename iterator_traits<_II1>::value_type _ValueType1;
    typedef typename iterator_traits<_II2>::value_type _ValueType2;
1597#endif
    __glibcxx_function_requires(_InputIteratorConcept<_II1>)
    __glibcxx_function_requires(_InputIteratorConcept<_II2>)
    __glibcxx_function_requires(_LessThanOpConcept<_ValueType1, _ValueType2>)
    __glibcxx_function_requires(_LessThanOpConcept<_ValueType2, _ValueType1>)
    __glibcxx_requires_valid_range(__first1, __last1);
    __glibcxx_requires_valid_range(__first2, __last2);

    return std::__lexicographical_compare_aux(std::__niter_base(__first1),
				std::__niter_base(__last1),
				std::__niter_base(__first2),
				std::__niter_base(__last2));
  }

/**
 *  @brief Performs @b dictionary comparison on ranges.
 *  @ingroup sorting_algorithms
 *  @param  __first1  An input iterator.
 *  @param  __last1   An input iterator.
 *  @param  __first2  An input iterator.
 *  @param  __last2   An input iterator.
 *  @param  __comp  A @link comparison_functors comparison functor@endlink.
 *  @return   A boolean true or false.
 *
 *  The same as the four-parameter @c lexicographical_compare, but uses the
 *  comp parameter instead of @c <.
*/
template<typename _II1, typename _II2, typename _Compare>
  _GLIBCXX20_CONSTEXPR
  inline bool
  lexicographical_compare(_II1 __first1, _II1 __last1,
	    _II2 __first2, _II2 __last2, _Compare __comp)
  {
    // concept requirements
    __glibcxx_function_requires(_InputIteratorConcept<_II1>)
    __glibcxx_function_requires(_InputIteratorConcept<_II2>)
    __glibcxx_requires_valid_range(__first1, __last1);
    __glibcxx_requires_valid_range(__first2, __last2);

    return std::__lexicographical_compare_impl
(__first1, __last1, __first2, __last2,
__gnu_cxx::__ops::__iter_comp_iter(__comp));
  }

1641#if __cpp_lib_three_way_comparison
// Iter points to a contiguous range of unsigned narrow character type
// or std::byte, suitable for comparison by memcmp.
template<typename _Iter>
  concept __is_byte_iter = contiguous_iterator<_Iter>
    && __is_memcmp_ordered<iter_value_t<_Iter>>::__value;

// Return a struct with two members, initialized to the smaller of x and y
// (or x if they compare equal) and the result of the comparison x <=> y.
template<typename _Tp>
  constexpr auto
  __min_cmp(_Tp __x, _Tp __y)
  {
    struct _Res {
_Tp _M_min;
decltype(__x <=> __y) _M_cmp;
    };
    auto __c = __x <=> __y;
    if (__c > 0)
return _Res{__y, __c};
    return _Res{__x, __c};
  }

/**
 *  @brief Performs dictionary comparison on ranges.
 *  @ingroup sorting_algorithms
 *  @param  __first1  An input iterator.
 *  @param  __last1   An input iterator.
 *  @param  __first2  An input iterator.
 *  @param  __last2   An input iterator.
 *  @param  __comp  A @link comparison_functors comparison functor@endlink.
 *  @return   The comparison category that `__comp(*__first1, *__first2)`
 *		returns.
*/
template<typename _InputIter1, typename _InputIter2, typename _Comp>
  constexpr auto
  lexicographical_compare_three_way(_InputIter1 __first1,
		      _InputIter1 __last1,
		      _InputIter2 __first2,
		      _InputIter2 __last2,
		      _Comp __comp)
  -> decltype(__comp(*__first1, *__first2))
  {
    // concept requirements
    __glibcxx_function_requires(_InputIteratorConcept<_InputIter1>)
    __glibcxx_function_requires(_InputIteratorConcept<_InputIter2>)
    __glibcxx_requires_valid_range(__first1, __last1);
    __glibcxx_requires_valid_range(__first2, __last2);

1690#if __cpp_lib_is_constant_evaluated
    using _Cat = decltype(__comp(*__first1, *__first2));
    static_assert(same_as<common_comparison_category_t<_Cat>, _Cat>);

    if (!std::is_constant_evaluated())
if constexpr (same_as<_Comp, __detail::_Synth3way>
      || same_as<_Comp, compare_three_way>)
 if constexpr (__is_byte_iter<_InputIter1>)
   if constexpr (__is_byte_iter<_InputIter2>)
     {
const auto [__len, __lencmp]
  = std::__min_cmp(__last1 - __first1, __last2 - __first2);
if (__len)
  {
    const auto __c
      = __builtin_memcmp(&*__first1, &*__first2, __len) <=> 0;
    if (__c != 0)
      return __c;
  }
return __lencmp;
     }
1711#endif // is_constant_evaluated
    while (__first1 != __last1)
{
 if (__first2 == __last2)
   return strong_ordering::greater;
 if (auto __cmp = __comp(*__first1, *__first2); __cmp != 0)
   return __cmp;
 ++__first1;
 ++__first2;
}
    return (__first2 == __last2) <=> true; // See PR 94006
  }

template<typename _InputIter1, typename _InputIter2>
  constexpr auto
  lexicographical_compare_three_way(_InputIter1 __first1,
		      _InputIter1 __last1,
		      _InputIter2 __first2,
		      _InputIter2 __last2)
  {
    return std::lexicographical_compare_three_way(__first1, __last1,
				    __first2, __last2,
				    compare_three_way{});
  }
1735#endif // three_way_comparison

template<typename _InputIterator1, typename _InputIterator2,
  typename _BinaryPredicate>
  _GLIBCXX20_CONSTEXPR
  pair<_InputIterator1, _InputIterator2>
  __mismatch(_InputIterator1 __first1, _InputIterator1 __last1,
      _InputIterator2 __first2, _BinaryPredicate __binary_pred)
  {
    while (__first1 != __last1 && __binary_pred(__first1, __first2))
{
 ++__first1;
 ++__first2;
}
    return pair<_InputIterator1, _InputIterator2>(__first1, __first2);
  }

/**
 *  @brief Finds the places in ranges which don't match.
 *  @ingroup non_mutating_algorithms
 *  @param  __first1  An input iterator.
 *  @param  __last1   An input iterator.
 *  @param  __first2  An input iterator.
 *  @return   A pair of iterators pointing to the first mismatch.
 *
 *  This compares the elements of two ranges using @c == and returns a pair
 *  of iterators.  The first iterator points into the first range, the
 *  second iterator points into the second range, and the elements pointed
 *  to by the iterators are not equal.
*/
template<typename _InputIterator1, typename _InputIterator2>
  _GLIBCXX20_CONSTEXPR
  inline pair<_InputIterator1, _InputIterator2>
  mismatch(_InputIterator1 __first1, _InputIterator1 __last1,
    _InputIterator2 __first2)
  {
    // concept requirements
    __glibcxx_function_requires(_InputIteratorConcept<_InputIterator1>)
    __glibcxx_function_requires(_InputIteratorConcept<_InputIterator2>)
    __glibcxx_function_requires(_EqualOpConcept<
   typename iterator_traits<_InputIterator1>::value_type,
   typename iterator_traits<_InputIterator2>::value_type>)
    __glibcxx_requires_valid_range(__first1, __last1);

    return _GLIBCXX_STD_Astd::__mismatch(__first1, __last1, __first2,
	     __gnu_cxx::__ops::__iter_equal_to_iter());
  }

/**
 *  @brief Finds the places in ranges which don't match.
 *  @ingroup non_mutating_algorithms
 *  @param  __first1  An input iterator.
 *  @param  __last1   An input iterator.
 *  @param  __first2  An input iterator.
 *  @param __binary_pred A binary predicate @link functors
 *         functor@endlink.
 *  @return   A pair of iterators pointing to the first mismatch.
 *
 *  This compares the elements of two ranges using the binary_pred
 *  parameter, and returns a pair
 *  of iterators.  The first iterator points into the first range, the
 *  second iterator points into the second range, and the elements pointed
 *  to by the iterators are not equal.
*/
template<typename _InputIterator1, typename _InputIterator2,
  typename _BinaryPredicate>
  _GLIBCXX20_CONSTEXPR
  inline pair<_InputIterator1, _InputIterator2>
  mismatch(_InputIterator1 __first1, _InputIterator1 __last1,
    _InputIterator2 __first2, _BinaryPredicate __binary_pred)
  {
    // concept requirements
    __glibcxx_function_requires(_InputIteratorConcept<_InputIterator1>)
    __glibcxx_function_requires(_InputIteratorConcept<_InputIterator2>)
    __glibcxx_requires_valid_range(__first1, __last1);

    return _GLIBCXX_STD_Astd::__mismatch(__first1, __last1, __first2,
__gnu_cxx::__ops::__iter_comp_iter(__binary_pred));
  }

1815#if __cplusplus201402L > 201103L

template<typename _InputIterator1, typename _InputIterator2,
  typename _BinaryPredicate>
  _GLIBCXX20_CONSTEXPR
  pair<_InputIterator1, _InputIterator2>
  __mismatch(_InputIterator1 __first1, _InputIterator1 __last1,
      _InputIterator2 __first2, _InputIterator2 __last2,
      _BinaryPredicate __binary_pred)
  {
    while (__first1 != __last1 && __first2 != __last2
    && __binary_pred(__first1, __first2))
{
 ++__first1;
 ++__first2;
}
    return pair<_InputIterator1, _InputIterator2>(__first1, __first2);
  }

/**
 *  @brief Finds the places in ranges which don't match.
 *  @ingroup non_mutating_algorithms
 *  @param  __first1  An input iterator.
 *  @param  __last1   An input iterator.
 *  @param  __first2  An input iterator.
 *  @param  __last2   An input iterator.
 *  @return   A pair of iterators pointing to the first mismatch.
 *
 *  This compares the elements of two ranges using @c == and returns a pair
 *  of iterators.  The first iterator points into the first range, the
 *  second iterator points into the second range, and the elements pointed
 *  to by the iterators are not equal.
*/
template<typename _InputIterator1, typename _InputIterator2>
  _GLIBCXX20_CONSTEXPR
  inline pair<_InputIterator1, _InputIterator2>
  mismatch(_InputIterator1 __first1, _InputIterator1 __last1,
    _InputIterator2 __first2, _InputIterator2 __last2)
  {
    // concept requirements
    __glibcxx_function_requires(_InputIteratorConcept<_InputIterator1>)
    __glibcxx_function_requires(_InputIteratorConcept<_InputIterator2>)
    __glibcxx_function_requires(_EqualOpConcept<
   typename iterator_traits<_InputIterator1>::value_type,
   typename iterator_traits<_InputIterator2>::value_type>)
    __glibcxx_requires_valid_range(__first1, __last1);
    __glibcxx_requires_valid_range(__first2, __last2);

    return _GLIBCXX_STD_Astd::__mismatch(__first1, __last1, __first2, __last2,
	     __gnu_cxx::__ops::__iter_equal_to_iter());
  }

/**
 *  @brief Finds the places in ranges which don't match.
 *  @ingroup non_mutating_algorithms
 *  @param  __first1  An input iterator.
 *  @param  __last1   An input iterator.
 *  @param  __first2  An input iterator.
 *  @param  __last2   An input iterator.
 *  @param __binary_pred A binary predicate @link functors
 *         functor@endlink.
 *  @return   A pair of iterators pointing to the first mismatch.
 *
 *  This compares the elements of two ranges using the binary_pred
 *  parameter, and returns a pair
 *  of iterators.  The first iterator points into the first range, the
 *  second iterator points into the second range, and the elements pointed
 *  to by the iterators are not equal.
*/
template<typename _InputIterator1, typename _InputIterator2,
  typename _BinaryPredicate>
  _GLIBCXX20_CONSTEXPR
  inline pair<_InputIterator1, _InputIterator2>
  mismatch(_InputIterator1 __first1, _InputIterator1 __last1,
    _InputIterator2 __first2, _InputIterator2 __last2,
    _BinaryPredicate __binary_pred)
  {
    // concept requirements
    __glibcxx_function_requires(_InputIteratorConcept<_InputIterator1>)
    __glibcxx_function_requires(_InputIteratorConcept<_InputIterator2>)
    __glibcxx_requires_valid_range(__first1, __last1);
    __glibcxx_requires_valid_range(__first2, __last2);

    return _GLIBCXX_STD_Astd::__mismatch(__first1, __last1, __first2, __last2,
	     __gnu_cxx::__ops::__iter_comp_iter(__binary_pred));
  }
1901#endif

1903_GLIBCXX_END_NAMESPACE_ALGO

/// This is an overload used by find algos for the Input Iterator case.
template<typename _InputIterator, typename _Predicate>
  _GLIBCXX20_CONSTEXPR
  inline _InputIterator
  __find_if(_InputIterator __first, _InputIterator __last,
     _Predicate __pred, input_iterator_tag)
  {
    while (__first != __last && !__pred(__first))
++__first;
    return __first;
  }

/// This is an overload used by find algos for the RAI case.
template<typename _RandomAccessIterator, typename _Predicate>
  _GLIBCXX20_CONSTEXPR
  _RandomAccessIterator
  __find_if(_RandomAccessIterator __first, _RandomAccessIterator __last,
     _Predicate __pred, random_access_iterator_tag)
  {
    typename iterator_traits<_RandomAccessIterator>::difference_type
__trip_count = (__last - __first) >> 2;

    for (; __trip_count > 0; --__trip_count)
{
 if (__pred(__first))
   return __first;
 ++__first;

 if (__pred(__first))
   return __first;
 ++__first;

 if (__pred(__first))
   return __first;
 ++__first;

 if (__pred(__first))
   return __first;
 ++__first;
}

    switch (__last - __first)
{
case 3:
 if (__pred(__first))
   return __first;
 ++__first;
 // FALLTHRU
case 2:
 if (__pred(__first))
   return __first;
 ++__first;
 // FALLTHRU
case 1:
 if (__pred(__first))
   return __first;
 ++__first;
 // FALLTHRU
case 0:
default:
 return __last;
}
  }

template<typename _Iterator, typename _Predicate>
  _GLIBCXX20_CONSTEXPR
  inline _Iterator
  __find_if(_Iterator __first, _Iterator __last, _Predicate __pred)
  {
    return __find_if(__first, __last, __pred,
       std::__iterator_category(__first));
  }

template<typename _InputIterator, typename _Predicate>
  _GLIBCXX20_CONSTEXPR
  typename iterator_traits<_InputIterator>::difference_type
  __count_if(_InputIterator __first, _InputIterator __last, _Predicate __pred)
  {
    typename iterator_traits<_InputIterator>::difference_type __n = 0;
    for (; __first != __last; ++__first)
if (__pred(__first))
 ++__n;
    return __n;
  }

1990#if __cplusplus201402L >= 201103L
template<typename _ForwardIterator1, typename _ForwardIterator2,
  typename _BinaryPredicate>
  _GLIBCXX20_CONSTEXPR
  bool
  __is_permutation(_ForwardIterator1 __first1, _ForwardIterator1 __last1,
     _ForwardIterator2 __first2, _BinaryPredicate __pred)
  {
    // Efficiently compare identical prefixes:  O(N) if sequences
    // have the same elements in the same order.
    for (; __first1 != __last1; ++__first1, (void)++__first2)
if (!__pred(__first1, __first2))
 break;

    if (__first1 == __last1)
return true;

    // Establish __last2 assuming equal ranges by iterating over the
    // rest of the list.
    _ForwardIterator2 __last2 = __first2;
    std::advance(__last2, std::distance(__first1, __last1));
    for (_ForwardIterator1 __scan = __first1; __scan != __last1; ++__scan)
{
 if (__scan != std::__find_if(__first1, __scan,
	  __gnu_cxx::__ops::__iter_comp_iter(__pred, __scan)))
   continue; // We've seen this one before.

 auto __matches
   = std::__count_if(__first2, __last2,
	__gnu_cxx::__ops::__iter_comp_iter(__pred, __scan));
 if (0 == __matches ||
     std::__count_if(__scan, __last1,
	__gnu_cxx::__ops::__iter_comp_iter(__pred, __scan))
     != __matches)
   return false;
}
    return true;
  }

/**
 *  @brief  Checks whether a permutation of the second sequence is equal
 *          to the first sequence.
 *  @ingroup non_mutating_algorithms
 *  @param  __first1  Start of first range.
 *  @param  __last1   End of first range.
 *  @param  __first2  Start of second range.
 *  @return true if there exists a permutation of the elements in the range
 *          [__first2, __first2 + (__last1 - __first1)), beginning with
 *          ForwardIterator2 begin, such that equal(__first1, __last1, begin)
 *          returns true; otherwise, returns false.
*/
template<typename _ForwardIterator1, typename _ForwardIterator2>
  _GLIBCXX20_CONSTEXPR
  inline bool
  is_permutation(_ForwardIterator1 __first1, _ForwardIterator1 __last1,
   _ForwardIterator2 __first2)
  {
    // concept requirements
    __glibcxx_function_requires(_ForwardIteratorConcept<_ForwardIterator1>)
    __glibcxx_function_requires(_ForwardIteratorConcept<_ForwardIterator2>)
    __glibcxx_function_requires(_EqualOpConcept<
typename iterator_traits<_ForwardIterator1>::value_type,
typename iterator_traits<_ForwardIterator2>::value_type>)
    __glibcxx_requires_valid_range(__first1, __last1);

    return std::__is_permutation(__first1, __last1, __first2,
		   __gnu_cxx::__ops::__iter_equal_to_iter());
  }
2058#endif // C++11

2060_GLIBCXX_END_NAMESPACE_VERSION
2061} // namespace std

2063// NB: This file is included within many other C++ includes, as a way
2064// of getting the base algorithms. So, make sure that parallel bits
2065// come in too if requested.
2066#ifdef _GLIBCXX_PARALLEL
2067# include <parallel/algobase.h>
2068#endif

2070#endif