/Users/buildslave/jenkins/workspace/clang-stage2-coverage-R/llvm/lib/Target/NVPTX/NVPTXRegisterInfo.cpp
Line | Count | Source |
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" |
17 | | #include "llvm/CodeGen/MachineFrameInfo.h" |
18 | | #include "llvm/CodeGen/MachineFunction.h" |
19 | | #include "llvm/CodeGen/MachineInstrBuilder.h" |
20 | | #include "llvm/CodeGen/TargetInstrInfo.h" |
21 | | #include "llvm/MC/MachineLocation.h" |
22 | | |
23 | | using namespace llvm; |
24 | | |
25 | | #define DEBUG_TYPE "nvptx-reg-info" |
26 | | |
27 | | namespace llvm { |
28 | 21.9k | std::string getNVPTXRegClassName(TargetRegisterClass const *RC) { |
29 | 21.9k | if (RC == &NVPTX::Float32RegsRegClass) |
30 | 1.68k | return ".f32"; |
31 | 20.2k | if (RC == &NVPTX::Float16RegsRegClass) |
32 | 1.68k | // Ideally fp16 registers should be .f16, but this syntax is only |
33 | 1.68k | // supported on sm_53+. On the other hand, .b16 registers are |
34 | 1.68k | // accepted for all supported fp16 instructions on all GPU |
35 | 1.68k | // variants, so we can use them instead. |
36 | 1.68k | return ".b16"; |
37 | 18.5k | if (RC == &NVPTX::Float16x2RegsRegClass) |
38 | 1.68k | return ".b32"; |
39 | 16.8k | if (RC == &NVPTX::Float64RegsRegClass) |
40 | 1.68k | return ".f64"; |
41 | 15.1k | if (RC == &NVPTX::Int64RegsRegClass) |
42 | 1.68k | // We use untyped (.b) integer registers here as NVCC does. |
43 | 1.68k | // Correctness of generated code does not depend on register type, |
44 | 1.68k | // but using .s/.u registers runs into ptxas bug that prevents |
45 | 1.68k | // assembly of otherwise valid PTX into SASS. Despite PTX ISA |
46 | 1.68k | // specifying only argument size for fp16 instructions, ptxas does |
47 | 1.68k | // not allow using .s16 or .u16 arguments for .fp16 |
48 | 1.68k | // instructions. At the same time it allows using .s32/.u32 |
49 | 1.68k | // arguments for .fp16v2 instructions: |
50 | 1.68k | // |
51 | 1.68k | // .reg .b16 rb16 |
52 | 1.68k | // .reg .s16 rs16 |
53 | 1.68k | // add.f16 rb16,rb16,rb16; // OK |
54 | 1.68k | // add.f16 rs16,rs16,rs16; // Arguments mismatch for instruction 'add' |
55 | 1.68k | // but: |
56 | 1.68k | // .reg .b32 rb32 |
57 | 1.68k | // .reg .s32 rs32 |
58 | 1.68k | // add.f16v2 rb32,rb32,rb32; // OK |
59 | 1.68k | // add.f16v2 rs32,rs32,rs32; // OK |
60 | 1.68k | return ".b64"; |
61 | 13.4k | if (RC == &NVPTX::Int32RegsRegClass) |
62 | 1.68k | return ".b32"; |
63 | 11.7k | if (RC == &NVPTX::Int16RegsRegClass) |
64 | 1.68k | return ".b16"; |
65 | 10.1k | if (RC == &NVPTX::Int1RegsRegClass) |
66 | 1.68k | return ".pred"; |
67 | 8.42k | if (RC == &NVPTX::SpecialRegsRegClass) |
68 | 1.68k | return "!Special!"; |
69 | 6.74k | return "INTERNAL"; |
70 | 6.74k | } |
71 | | |
72 | 21.9k | std::string getNVPTXRegClassStr(TargetRegisterClass const *RC) { |
73 | 21.9k | if (RC == &NVPTX::Float32RegsRegClass) |
74 | 1.68k | return "%f"; |
75 | 20.2k | if (RC == &NVPTX::Float16RegsRegClass) |
76 | 1.68k | return "%h"; |
77 | 18.5k | if (RC == &NVPTX::Float16x2RegsRegClass) |
78 | 1.68k | return "%hh"; |
79 | 16.8k | if (RC == &NVPTX::Float64RegsRegClass) |
80 | 1.68k | return "%fd"; |
81 | 15.1k | if (RC == &NVPTX::Int64RegsRegClass) |
82 | 1.68k | return "%rd"; |
83 | 13.4k | if (RC == &NVPTX::Int32RegsRegClass) |
84 | 1.68k | return "%r"; |
85 | 11.7k | if (RC == &NVPTX::Int16RegsRegClass) |
86 | 1.68k | return "%rs"; |
87 | 10.1k | if (RC == &NVPTX::Int1RegsRegClass) |
88 | 1.68k | return "%p"; |
89 | 8.42k | if (RC == &NVPTX::SpecialRegsRegClass) |
90 | 1.68k | return "!Special!"; |
91 | 6.74k | return "INTERNAL"; |
92 | 6.74k | } |
93 | | } |
94 | | |
95 | 455 | NVPTXRegisterInfo::NVPTXRegisterInfo() : NVPTXGenRegisterInfo(0) {} |
96 | | |
97 | | #define GET_REGINFO_TARGET_DESC |
98 | | #include "NVPTXGenRegisterInfo.inc" |
99 | | |
100 | | /// NVPTX Callee Saved Registers |
101 | | const MCPhysReg * |
102 | 1.07k | NVPTXRegisterInfo::getCalleeSavedRegs(const MachineFunction *) const { |
103 | 1.07k | static const MCPhysReg CalleeSavedRegs[] = { 0 }; |
104 | 1.07k | return CalleeSavedRegs; |
105 | 1.07k | } |
106 | | |
107 | 3.40k | BitVector NVPTXRegisterInfo::getReservedRegs(const MachineFunction &MF) const { |
108 | 3.40k | BitVector Reserved(getNumRegs()); |
109 | 3.40k | return Reserved; |
110 | 3.40k | } |
111 | | |
112 | | void NVPTXRegisterInfo::eliminateFrameIndex(MachineBasicBlock::iterator II, |
113 | | int SPAdj, unsigned FIOperandNum, |
114 | 98 | RegScavenger *RS) const { |
115 | 98 | assert(SPAdj == 0 && "Unexpected"); |
116 | 98 | |
117 | 98 | MachineInstr &MI = *II; |
118 | 98 | int FrameIndex = MI.getOperand(FIOperandNum).getIndex(); |
119 | 98 | |
120 | 98 | MachineFunction &MF = *MI.getParent()->getParent(); |
121 | 98 | int Offset = MF.getFrameInfo().getObjectOffset(FrameIndex) + |
122 | 98 | MI.getOperand(FIOperandNum + 1).getImm(); |
123 | 98 | |
124 | 98 | // Using I0 as the frame pointer |
125 | 98 | MI.getOperand(FIOperandNum).ChangeToRegister(NVPTX::VRFrame, false); |
126 | 98 | MI.getOperand(FIOperandNum + 1).ChangeToImmediate(Offset); |
127 | 98 | } |
128 | | |
129 | 15 | Register NVPTXRegisterInfo::getFrameRegister(const MachineFunction &MF) const { |
130 | 15 | return NVPTX::VRFrame; |
131 | 15 | } |