1 //===- NVPTXRegisterInfo.cpp - NVPTX Register Information -----------------===//
2 //
3 //                     The LLVM Compiler Infrastructure
4 //
5 // This file is distributed under the University of Illinois Open Source
6 // License. See LICENSE.TXT for details.
7 //
8 //===----------------------------------------------------------------------===//
9 //
10 // This file contains the NVPTX implementation of the TargetRegisterInfo class.
11 //
12 //===----------------------------------------------------------------------===//
13 
14 #include "NVPTXRegisterInfo.h"
15 #include "NVPTX.h"
16 #include "NVPTXSubtarget.h"
17 #include "llvm/ADT/BitVector.h"
18 #include "llvm/CodeGen/MachineFrameInfo.h"
19 #include "llvm/CodeGen/MachineFunction.h"
20 #include "llvm/CodeGen/MachineInstrBuilder.h"
21 #include "llvm/MC/MachineLocation.h"
22 #include "llvm/Target/TargetInstrInfo.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   }
33   if (RC == &NVPTX::Float64RegsRegClass) {
34     return ".f64";
35   } else if (RC == &NVPTX::Int64RegsRegClass) {
36     // We use untyped (.b) integer registers here as NVCC does.
37     // Correctness of generated code does not depend on register type,
38     // but using .s/.u registers runs into ptxas bug that prevents
39     // assembly of otherwise valid PTX into SASS. Despite PTX ISA
40     // specifying only argument size for fp16 instructions, ptxas does
41     // not allow using .s16 or .u16 arguments for .fp16
42     // instructions. At the same time it allows using .s32/.u32
43     // arguments for .fp16v2 instructions:
44     //
45     //   .reg .b16 rb16
46     //   .reg .s16 rs16
47     //   add.f16 rb16,rb16,rb16; // OK
48     //   add.f16 rs16,rs16,rs16; // Arguments mismatch for instruction 'add'
49     // but:
50     //   .reg .b32 rb32
51     //   .reg .s32 rs32
52     //   add.f16v2 rb32,rb32,rb32; // OK
53     //   add.f16v2 rs32,rs32,rs32; // OK
54     return ".b64";
55   } else if (RC == &NVPTX::Int32RegsRegClass) {
56     return ".b32";
57   } else if (RC == &NVPTX::Int16RegsRegClass) {
58     return ".b16";
59   } else if (RC == &NVPTX::Int1RegsRegClass) {
60     return ".pred";
61   } else if (RC == &NVPTX::SpecialRegsRegClass) {
62     return "!Special!";
63   } else {
64     return "INTERNAL";
65   }
66   return "";
67 }
68 
69 std::string getNVPTXRegClassStr(TargetRegisterClass const *RC) {
70   if (RC == &NVPTX::Float32RegsRegClass) {
71     return "%f";
72   }
73   if (RC == &NVPTX::Float64RegsRegClass) {
74     return "%fd";
75   } else if (RC == &NVPTX::Int64RegsRegClass) {
76     return "%rd";
77   } else if (RC == &NVPTX::Int32RegsRegClass) {
78     return "%r";
79   } else if (RC == &NVPTX::Int16RegsRegClass) {
80     return "%rs";
81   } else if (RC == &NVPTX::Int1RegsRegClass) {
82     return "%p";
83   } else if (RC == &NVPTX::SpecialRegsRegClass) {
84     return "!Special!";
85   } else {
86     return "INTERNAL";
87   }
88   return "";
89 }
90 }
91 
92 NVPTXRegisterInfo::NVPTXRegisterInfo() : NVPTXGenRegisterInfo(0) {}
93 
94 #define GET_REGINFO_TARGET_DESC
95 #include "NVPTXGenRegisterInfo.inc"
96 
97 /// NVPTX Callee Saved Registers
98 const MCPhysReg *
99 NVPTXRegisterInfo::getCalleeSavedRegs(const MachineFunction *) const {
100   static const MCPhysReg CalleeSavedRegs[] = { 0 };
101   return CalleeSavedRegs;
102 }
103 
104 BitVector NVPTXRegisterInfo::getReservedRegs(const MachineFunction &MF) const {
105   BitVector Reserved(getNumRegs());
106   return Reserved;
107 }
108 
109 void NVPTXRegisterInfo::eliminateFrameIndex(MachineBasicBlock::iterator II,
110                                             int SPAdj, unsigned FIOperandNum,
111                                             RegScavenger *RS) const {
112   assert(SPAdj == 0 && "Unexpected");
113 
114   MachineInstr &MI = *II;
115   int FrameIndex = MI.getOperand(FIOperandNum).getIndex();
116 
117   MachineFunction &MF = *MI.getParent()->getParent();
118   int Offset = MF.getFrameInfo().getObjectOffset(FrameIndex) +
119                MI.getOperand(FIOperandNum + 1).getImm();
120 
121   // Using I0 as the frame pointer
122   MI.getOperand(FIOperandNum).ChangeToRegister(NVPTX::VRFrame, false);
123   MI.getOperand(FIOperandNum + 1).ChangeToImmediate(Offset);
124 }
125 
126 unsigned NVPTXRegisterInfo::getFrameRegister(const MachineFunction &MF) const {
127   return NVPTX::VRFrame;
128 }
129