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" |
18 | #include "llvm/CodeGen/MachineFrameInfo.h" |
19 | #include "llvm/CodeGen/MachineFunction.h" |
20 | #include "llvm/CodeGen/MachineInstrBuilder.h" |
21 | #include "llvm/CodeGen/TargetInstrInfo.h" |
22 | #include "llvm/MC/MachineLocation.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::Float64RegsRegClass) |
33 | return ".f64" ; |
34 | if (RC == &NVPTX::Int64RegsRegClass) |
35 | // We use untyped (.b) integer registers here as NVCC does. |
36 | // Correctness of generated code does not depend on register type, |
37 | // but using .s/.u registers runs into ptxas bug that prevents |
38 | // assembly of otherwise valid PTX into SASS. Despite PTX ISA |
39 | // specifying only argument size for fp16 instructions, ptxas does |
40 | // not allow using .s16 or .u16 arguments for .fp16 |
41 | // instructions. At the same time it allows using .s32/.u32 |
42 | // arguments for .fp16v2 instructions: |
43 | // |
44 | // .reg .b16 rb16 |
45 | // .reg .s16 rs16 |
46 | // add.f16 rb16,rb16,rb16; // OK |
47 | // add.f16 rs16,rs16,rs16; // Arguments mismatch for instruction 'add' |
48 | // but: |
49 | // .reg .b32 rb32 |
50 | // .reg .s32 rs32 |
51 | // add.f16v2 rb32,rb32,rb32; // OK |
52 | // add.f16v2 rs32,rs32,rs32; // OK |
53 | return ".b64" ; |
54 | if (RC == &NVPTX::Int32RegsRegClass) |
55 | return ".b32" ; |
56 | if (RC == &NVPTX::Int16RegsRegClass) |
57 | return ".b16" ; |
58 | if (RC == &NVPTX::Int1RegsRegClass) |
59 | return ".pred" ; |
60 | if (RC == &NVPTX::SpecialRegsRegClass) |
61 | return "!Special!" ; |
62 | return "INTERNAL" ; |
63 | } |
64 | |
65 | std::string getNVPTXRegClassStr(TargetRegisterClass const *RC) { |
66 | if (RC == &NVPTX::Float32RegsRegClass) |
67 | return "%f" ; |
68 | if (RC == &NVPTX::Float64RegsRegClass) |
69 | return "%fd" ; |
70 | if (RC == &NVPTX::Int64RegsRegClass) |
71 | return "%rd" ; |
72 | if (RC == &NVPTX::Int32RegsRegClass) |
73 | return "%r" ; |
74 | if (RC == &NVPTX::Int16RegsRegClass) |
75 | return "%rs" ; |
76 | if (RC == &NVPTX::Int1RegsRegClass) |
77 | return "%p" ; |
78 | if (RC == &NVPTX::SpecialRegsRegClass) |
79 | return "!Special!" ; |
80 | return "INTERNAL" ; |
81 | } |
82 | } |
83 | |
84 | NVPTXRegisterInfo::NVPTXRegisterInfo() |
85 | : NVPTXGenRegisterInfo(0), StrPool(StrAlloc) {} |
86 | |
87 | #define GET_REGINFO_TARGET_DESC |
88 | #include "NVPTXGenRegisterInfo.inc" |
89 | |
90 | /// NVPTX Callee Saved Registers |
91 | const MCPhysReg * |
92 | NVPTXRegisterInfo::getCalleeSavedRegs(const MachineFunction *) const { |
93 | static const MCPhysReg CalleeSavedRegs[] = { 0 }; |
94 | return CalleeSavedRegs; |
95 | } |
96 | |
97 | BitVector NVPTXRegisterInfo::getReservedRegs(const MachineFunction &MF) const { |
98 | BitVector Reserved(getNumRegs()); |
99 | for (unsigned Reg = NVPTX::ENVREG0; Reg <= NVPTX::ENVREG31; ++Reg) { |
100 | markSuperRegs(Reserved, Reg); |
101 | } |
102 | markSuperRegs(Reserved, NVPTX::VRFrame32); |
103 | markSuperRegs(Reserved, NVPTX::VRFrameLocal32); |
104 | markSuperRegs(Reserved, NVPTX::VRFrame64); |
105 | markSuperRegs(Reserved, NVPTX::VRFrameLocal64); |
106 | markSuperRegs(Reserved, NVPTX::VRDepot); |
107 | return Reserved; |
108 | } |
109 | |
110 | bool NVPTXRegisterInfo::eliminateFrameIndex(MachineBasicBlock::iterator II, |
111 | int SPAdj, unsigned FIOperandNum, |
112 | RegScavenger *RS) const { |
113 | assert(SPAdj == 0 && "Unexpected" ); |
114 | |
115 | MachineInstr &MI = *II; |
116 | int FrameIndex = MI.getOperand(i: FIOperandNum).getIndex(); |
117 | |
118 | MachineFunction &MF = *MI.getParent()->getParent(); |
119 | int Offset = MF.getFrameInfo().getObjectOffset(ObjectIdx: FrameIndex) + |
120 | MI.getOperand(i: FIOperandNum + 1).getImm(); |
121 | |
122 | // Using I0 as the frame pointer |
123 | MI.getOperand(i: FIOperandNum).ChangeToRegister(Reg: getFrameRegister(MF), isDef: false); |
124 | MI.getOperand(i: FIOperandNum + 1).ChangeToImmediate(ImmVal: Offset); |
125 | return false; |
126 | } |
127 | |
128 | Register NVPTXRegisterInfo::getFrameRegister(const MachineFunction &MF) const { |
129 | const NVPTXTargetMachine &TM = |
130 | static_cast<const NVPTXTargetMachine &>(MF.getTarget()); |
131 | return TM.is64Bit() ? NVPTX::VRFrame64 : NVPTX::VRFrame32; |
132 | } |
133 | |
134 | Register |
135 | NVPTXRegisterInfo::getFrameLocalRegister(const MachineFunction &MF) const { |
136 | const NVPTXTargetMachine &TM = |
137 | static_cast<const NVPTXTargetMachine &>(MF.getTarget()); |
138 | return TM.is64Bit() ? NVPTX::VRFrameLocal64 : NVPTX::VRFrameLocal32; |
139 | } |
140 | |