LLVM  9.0.0svn
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 "llvm/ADT/BitVector.h"
22 
23 using namespace llvm;
24 
25 #define DEBUG_TYPE "nvptx-reg-info"
26 
27 namespace llvm {
28 std::string getNVPTXRegClassName(TargetRegisterClass const *RC) {
29  if (RC == &NVPTX::Float32RegsRegClass)
30  return ".f32";
31  if (RC == &NVPTX::Float16RegsRegClass)
32  // Ideally fp16 registers should be .f16, but this syntax is only
33  // supported on sm_53+. On the other hand, .b16 registers are
34  // accepted for all supported fp16 instructions on all GPU
35  // variants, so we can use them instead.
36  return ".b16";
37  if (RC == &NVPTX::Float16x2RegsRegClass)
38  return ".b32";
39  if (RC == &NVPTX::Float64RegsRegClass)
40  return ".f64";
41  if (RC == &NVPTX::Int64RegsRegClass)
42  // We use untyped (.b) integer registers here as NVCC does.
43  // Correctness of generated code does not depend on register type,
44  // but using .s/.u registers runs into ptxas bug that prevents
45  // assembly of otherwise valid PTX into SASS. Despite PTX ISA
46  // specifying only argument size for fp16 instructions, ptxas does
47  // not allow using .s16 or .u16 arguments for .fp16
48  // instructions. At the same time it allows using .s32/.u32
49  // arguments for .fp16v2 instructions:
50  //
51  // .reg .b16 rb16
52  // .reg .s16 rs16
53  // add.f16 rb16,rb16,rb16; // OK
54  // add.f16 rs16,rs16,rs16; // Arguments mismatch for instruction 'add'
55  // but:
56  // .reg .b32 rb32
57  // .reg .s32 rs32
58  // add.f16v2 rb32,rb32,rb32; // OK
59  // add.f16v2 rs32,rs32,rs32; // OK
60  return ".b64";
61  if (RC == &NVPTX::Int32RegsRegClass)
62  return ".b32";
63  if (RC == &NVPTX::Int16RegsRegClass)
64  return ".b16";
65  if (RC == &NVPTX::Int1RegsRegClass)
66  return ".pred";
67  if (RC == &NVPTX::SpecialRegsRegClass)
68  return "!Special!";
69  return "INTERNAL";
70 }
71 
72 std::string getNVPTXRegClassStr(TargetRegisterClass const *RC) {
73  if (RC == &NVPTX::Float32RegsRegClass)
74  return "%f";
75  if (RC == &NVPTX::Float16RegsRegClass)
76  return "%h";
77  if (RC == &NVPTX::Float16x2RegsRegClass)
78  return "%hh";
79  if (RC == &NVPTX::Float64RegsRegClass)
80  return "%fd";
81  if (RC == &NVPTX::Int64RegsRegClass)
82  return "%rd";
83  if (RC == &NVPTX::Int32RegsRegClass)
84  return "%r";
85  if (RC == &NVPTX::Int16RegsRegClass)
86  return "%rs";
87  if (RC == &NVPTX::Int1RegsRegClass)
88  return "%p";
89  if (RC == &NVPTX::SpecialRegsRegClass)
90  return "!Special!";
91  return "INTERNAL";
92 }
93 }
94 
96 
97 #define GET_REGINFO_TARGET_DESC
98 #include "NVPTXGenRegisterInfo.inc"
99 
100 /// NVPTX Callee Saved Registers
101 const MCPhysReg *
103  static const MCPhysReg CalleeSavedRegs[] = { 0 };
104  return CalleeSavedRegs;
105 }
106 
108  BitVector Reserved(getNumRegs());
109  return Reserved;
110 }
111 
113  int SPAdj, unsigned FIOperandNum,
114  RegScavenger *RS) const {
115  assert(SPAdj == 0 && "Unexpected");
116 
117  MachineInstr &MI = *II;
118  int FrameIndex = MI.getOperand(FIOperandNum).getIndex();
119 
120  MachineFunction &MF = *MI.getParent()->getParent();
121  int Offset = MF.getFrameInfo().getObjectOffset(FrameIndex) +
122  MI.getOperand(FIOperandNum + 1).getImm();
123 
124  // Using I0 as the frame pointer
125  MI.getOperand(FIOperandNum).ChangeToRegister(NVPTX::VRFrame, false);
126  MI.getOperand(FIOperandNum + 1).ChangeToImmediate(Offset);
127 }
128 
130  return NVPTX::VRFrame;
131 }
This class represents lattice values for constants.
Definition: AllocatorList.h:23
void ChangeToRegister(unsigned Reg, bool isDef, bool isImp=false, bool isKill=false, bool isDead=false, bool isUndef=false, bool isDebug=false)
ChangeToRegister - Replace this operand with a new register operand of the specified value...
const MCPhysReg * getCalleeSavedRegs(const MachineFunction *MF) const override
NVPTX Callee Saved Registers.
void eliminateFrameIndex(MachineBasicBlock::iterator MI, int SPAdj, unsigned FIOperandNum, RegScavenger *RS=nullptr) const override
int64_t getObjectOffset(int ObjectIdx) const
Return the assigned stack offset of the specified object from the incoming stack pointer.
uint16_t MCPhysReg
An unsigned integer type large enough to represent all physical registers, but not necessarily virtua...
void ChangeToImmediate(int64_t ImmVal)
ChangeToImmediate - Replace this operand with a new immediate operand of the specified value...
MachineFrameInfo & getFrameInfo()
getFrameInfo - Return the frame info object for the current function.
BitVector getReservedRegs(const MachineFunction &MF) const override
int64_t getImm() const
const MachineBasicBlock * getParent() const
Definition: MachineInstr.h:255
Representation of each machine instruction.
Definition: MachineInstr.h:63
const MachineFunction * getParent() const
Return the MachineFunction containing this basic block.
unsigned getFrameRegister(const MachineFunction &MF) const override
std::string getNVPTXRegClassStr(TargetRegisterClass const *RC)
std::string getNVPTXRegClassName(TargetRegisterClass const *RC)
assert(ImpDefSCC.getReg()==AMDGPU::SCC &&ImpDefSCC.isDef())
IRTranslator LLVM IR MI
const MachineOperand & getOperand(unsigned i) const
Definition: MachineInstr.h:415