LLVM  15.0.0git
NVPTXRegisterInfo.cpp
Go to the documentation of this file.
1 //===- NVPTXRegisterInfo.cpp - NVPTX 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 // This file contains the NVPTX implementation of the TargetRegisterInfo class.
10 //
11 //===----------------------------------------------------------------------===//
12 
13 #include "NVPTXRegisterInfo.h"
14 #include "NVPTX.h"
15 #include "NVPTXSubtarget.h"
16 #include "NVPTXTargetMachine.h"
17 #include "llvm/ADT/BitVector.h"
23 
24 using namespace llvm;
25 
26 #define DEBUG_TYPE "nvptx-reg-info"
27 
28 namespace llvm {
29 std::string getNVPTXRegClassName(TargetRegisterClass const *RC) {
30  if (RC == &NVPTX::Float32RegsRegClass)
31  return ".f32";
32  if (RC == &NVPTX::Float16RegsRegClass)
33  // Ideally fp16 registers should be .f16, but this syntax is only
34  // supported on sm_53+. On the other hand, .b16 registers are
35  // accepted for all supported fp16 instructions on all GPU
36  // variants, so we can use them instead.
37  return ".b16";
38  if (RC == &NVPTX::Float16x2RegsRegClass)
39  return ".b32";
40  if (RC == &NVPTX::Float64RegsRegClass)
41  return ".f64";
42  if (RC == &NVPTX::Int64RegsRegClass)
43  // We use untyped (.b) integer registers here as NVCC does.
44  // Correctness of generated code does not depend on register type,
45  // but using .s/.u registers runs into ptxas bug that prevents
46  // assembly of otherwise valid PTX into SASS. Despite PTX ISA
47  // specifying only argument size for fp16 instructions, ptxas does
48  // not allow using .s16 or .u16 arguments for .fp16
49  // instructions. At the same time it allows using .s32/.u32
50  // arguments for .fp16v2 instructions:
51  //
52  // .reg .b16 rb16
53  // .reg .s16 rs16
54  // add.f16 rb16,rb16,rb16; // OK
55  // add.f16 rs16,rs16,rs16; // Arguments mismatch for instruction 'add'
56  // but:
57  // .reg .b32 rb32
58  // .reg .s32 rs32
59  // add.f16v2 rb32,rb32,rb32; // OK
60  // add.f16v2 rs32,rs32,rs32; // OK
61  return ".b64";
62  if (RC == &NVPTX::Int32RegsRegClass)
63  return ".b32";
64  if (RC == &NVPTX::Int16RegsRegClass)
65  return ".b16";
66  if (RC == &NVPTX::Int1RegsRegClass)
67  return ".pred";
68  if (RC == &NVPTX::SpecialRegsRegClass)
69  return "!Special!";
70  return "INTERNAL";
71 }
72 
73 std::string getNVPTXRegClassStr(TargetRegisterClass const *RC) {
74  if (RC == &NVPTX::Float32RegsRegClass)
75  return "%f";
76  if (RC == &NVPTX::Float16RegsRegClass)
77  return "%h";
78  if (RC == &NVPTX::Float16x2RegsRegClass)
79  return "%hh";
80  if (RC == &NVPTX::Float64RegsRegClass)
81  return "%fd";
82  if (RC == &NVPTX::Int64RegsRegClass)
83  return "%rd";
84  if (RC == &NVPTX::Int32RegsRegClass)
85  return "%r";
86  if (RC == &NVPTX::Int16RegsRegClass)
87  return "%rs";
88  if (RC == &NVPTX::Int1RegsRegClass)
89  return "%p";
90  if (RC == &NVPTX::SpecialRegsRegClass)
91  return "!Special!";
92  return "INTERNAL";
93 }
94 }
95 
97 
98 #define GET_REGINFO_TARGET_DESC
99 #include "NVPTXGenRegisterInfo.inc"
100 
101 /// NVPTX Callee Saved Registers
102 const MCPhysReg *
104  static const MCPhysReg CalleeSavedRegs[] = { 0 };
105  return CalleeSavedRegs;
106 }
107 
109  BitVector Reserved(getNumRegs());
110  for (unsigned Reg = NVPTX::ENVREG0; Reg <= NVPTX::ENVREG31; ++Reg) {
111  markSuperRegs(Reserved, Reg);
112  }
113  markSuperRegs(Reserved, NVPTX::VRFrame32);
114  markSuperRegs(Reserved, NVPTX::VRFrameLocal32);
115  markSuperRegs(Reserved, NVPTX::VRFrame64);
116  markSuperRegs(Reserved, NVPTX::VRFrameLocal64);
117  markSuperRegs(Reserved, NVPTX::VRDepot);
118  return Reserved;
119 }
120 
122  int SPAdj, unsigned FIOperandNum,
123  RegScavenger *RS) const {
124  assert(SPAdj == 0 && "Unexpected");
125 
126  MachineInstr &MI = *II;
127  int FrameIndex = MI.getOperand(FIOperandNum).getIndex();
128 
129  MachineFunction &MF = *MI.getParent()->getParent();
130  int Offset = MF.getFrameInfo().getObjectOffset(FrameIndex) +
131  MI.getOperand(FIOperandNum + 1).getImm();
132 
133  // Using I0 as the frame pointer
134  MI.getOperand(FIOperandNum).ChangeToRegister(getFrameRegister(MF), false);
135  MI.getOperand(FIOperandNum + 1).ChangeToImmediate(Offset);
136 }
137 
139  const NVPTXTargetMachine &TM =
140  static_cast<const NVPTXTargetMachine &>(MF.getTarget());
141  return TM.is64Bit() ? NVPTX::VRFrame64 : NVPTX::VRFrame32;
142 }
143 
144 Register
146  const NVPTXTargetMachine &TM =
147  static_cast<const NVPTXTargetMachine &>(MF.getTarget());
148  return TM.is64Bit() ? NVPTX::VRFrameLocal64 : NVPTX::VRFrameLocal32;
149 }
llvm::NVPTXRegisterInfo::eliminateFrameIndex
void eliminateFrameIndex(MachineBasicBlock::iterator MI, int SPAdj, unsigned FIOperandNum, RegScavenger *RS=nullptr) const override
Definition: NVPTXRegisterInfo.cpp:121
MI
IRTranslator LLVM IR MI
Definition: IRTranslator.cpp:104
llvm
This is an optimization pass for GlobalISel generic memory operations.
Definition: AddressRanges.h:17
llvm::X86Disassembler::Reg
Reg
All possible values of the reg field in the ModR/M byte.
Definition: X86DisassemblerDecoder.h:462
TargetInstrInfo.h
llvm::NVPTXRegisterInfo::getFrameRegister
Register getFrameRegister(const MachineFunction &MF) const override
Definition: NVPTXRegisterInfo.cpp:138
NVPTXRegisterInfo.h
llvm::NVPTXTargetMachine
NVPTXTargetMachine.
Definition: NVPTXTargetMachine.h:25
NVPTX.h
llvm::NVPTXRegisterInfo::getFrameLocalRegister
Register getFrameLocalRegister(const MachineFunction &MF) const
Definition: NVPTXRegisterInfo.cpp:145
llvm::TargetRegisterClass
Definition: TargetRegisterInfo.h:45
BitVector.h
llvm::MachineFrameInfo::getObjectOffset
int64_t getObjectOffset(int ObjectIdx) const
Return the assigned stack offset of the specified object from the incoming stack pointer.
Definition: MachineFrameInfo.h:518
llvm::NVPTXRegisterInfo::NVPTXRegisterInfo
NVPTXRegisterInfo()
Definition: NVPTXRegisterInfo.cpp:96
llvm::BitVector
Definition: BitVector.h:75
llvm::getNVPTXRegClassStr
std::string getNVPTXRegClassStr(TargetRegisterClass const *RC)
Definition: NVPTXRegisterInfo.cpp:73
llvm::MachineInstr
Representation of each machine instruction.
Definition: MachineInstr.h:66
llvm::MCPhysReg
uint16_t MCPhysReg
An unsigned integer type large enough to represent all physical registers, but not necessarily virtua...
Definition: MCRegister.h:21
llvm::RegScavenger
Definition: RegisterScavenging.h:34
assert
assert(ImpDefSCC.getReg()==AMDGPU::SCC &&ImpDefSCC.isDef())
llvm::MachineFunction::getFrameInfo
MachineFrameInfo & getFrameInfo()
getFrameInfo - Return the frame info object for the current function.
Definition: MachineFunction.h:656
llvm::MachineFunction
Definition: MachineFunction.h:241
MachineLocation.h
llvm::NVPTXRegisterInfo::getReservedRegs
BitVector getReservedRegs(const MachineFunction &MF) const override
Definition: NVPTXRegisterInfo.cpp:108
llvm::Register
Wrapper class representing virtual and physical registers.
Definition: Register.h:19
llvm::ISD::FrameIndex
@ FrameIndex
Definition: ISDOpcodes.h:80
llvm::NVPTXRegisterInfo::getCalleeSavedRegs
const MCPhysReg * getCalleeSavedRegs(const MachineFunction *MF) const override
NVPTX Callee Saved Registers.
Definition: NVPTXRegisterInfo.cpp:103
uint16_t
llvm::MachineFunction::getTarget
const LLVMTargetMachine & getTarget() const
getTarget - Return the target machine this machine code is compiled with
Definition: MachineFunction.h:636
MachineFrameInfo.h
NVPTXTargetMachine.h
MachineInstrBuilder.h
NVPTXSubtarget.h
TM
const char LLVMTargetMachineRef TM
Definition: PassBuilderBindings.cpp:47
MachineFunction.h
llvm::MachineInstrBundleIterator< MachineInstr >
llvm::getNVPTXRegClassName
std::string getNVPTXRegClassName(TargetRegisterClass const *RC)
Definition: NVPTXRegisterInfo.cpp:29
NVPTXGenRegisterInfo