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 "MCTargetDesc/NVPTXInstPrinter.h" 15 #include "NVPTX.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/TargetInstrInfo.h" 21 22 using namespace llvm; 23 24 #define DEBUG_TYPE "nvptx-reg-info" 25 26 namespace llvm { 27 std::string getNVPTXRegClassName(TargetRegisterClass const *RC) { 28 if (RC == &NVPTX::Float32RegsRegClass) 29 return ".f32"; 30 if (RC == &NVPTX::Float64RegsRegClass) 31 return ".f64"; 32 if (RC == &NVPTX::Int128RegsRegClass) 33 return ".b128"; 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::Int128RegsRegClass) 71 return "%rq"; 72 if (RC == &NVPTX::Int64RegsRegClass) 73 return "%rd"; 74 if (RC == &NVPTX::Int32RegsRegClass) 75 return "%r"; 76 if (RC == &NVPTX::Int16RegsRegClass) 77 return "%rs"; 78 if (RC == &NVPTX::Int1RegsRegClass) 79 return "%p"; 80 if (RC == &NVPTX::SpecialRegsRegClass) 81 return "!Special!"; 82 return "INTERNAL"; 83 } 84 } 85 86 NVPTXRegisterInfo::NVPTXRegisterInfo() 87 : NVPTXGenRegisterInfo(0), StrPool(StrAlloc) {} 88 89 #define GET_REGINFO_TARGET_DESC 90 #include "NVPTXGenRegisterInfo.inc" 91 92 /// NVPTX Callee Saved Registers 93 const MCPhysReg * 94 NVPTXRegisterInfo::getCalleeSavedRegs(const MachineFunction *) const { 95 static const MCPhysReg CalleeSavedRegs[] = { 0 }; 96 return CalleeSavedRegs; 97 } 98 99 BitVector NVPTXRegisterInfo::getReservedRegs(const MachineFunction &MF) const { 100 BitVector Reserved(getNumRegs()); 101 for (unsigned Reg = NVPTX::ENVREG0; Reg <= NVPTX::ENVREG31; ++Reg) { 102 markSuperRegs(Reserved, Reg); 103 } 104 markSuperRegs(Reserved, NVPTX::VRFrame32); 105 markSuperRegs(Reserved, NVPTX::VRFrameLocal32); 106 markSuperRegs(Reserved, NVPTX::VRFrame64); 107 markSuperRegs(Reserved, NVPTX::VRFrameLocal64); 108 markSuperRegs(Reserved, NVPTX::VRDepot); 109 return Reserved; 110 } 111 112 bool NVPTXRegisterInfo::eliminateFrameIndex(MachineBasicBlock::iterator II, 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(getFrameRegister(MF), false); 126 MI.getOperand(FIOperandNum + 1).ChangeToImmediate(Offset); 127 return false; 128 } 129 130 Register NVPTXRegisterInfo::getFrameRegister(const MachineFunction &MF) const { 131 const NVPTXTargetMachine &TM = 132 static_cast<const NVPTXTargetMachine &>(MF.getTarget()); 133 return TM.is64Bit() ? NVPTX::VRFrame64 : NVPTX::VRFrame32; 134 } 135 136 Register 137 NVPTXRegisterInfo::getFrameLocalRegister(const MachineFunction &MF) const { 138 const NVPTXTargetMachine &TM = 139 static_cast<const NVPTXTargetMachine &>(MF.getTarget()); 140 return TM.is64Bit() ? NVPTX::VRFrameLocal64 : NVPTX::VRFrameLocal32; 141 } 142 143 void NVPTXRegisterInfo::clearDebugRegisterMap() const { 144 debugRegisterMap.clear(); 145 } 146 147 static uint64_t encodeRegisterForDwarf(std::string registerName) { 148 if (registerName.length() > 8) { 149 // The name is more than 8 characters long, and so won't fit into 64 bits. 150 return 0; 151 } 152 153 // Encode the name string into a DWARF register number using cuda-gdb's 154 // encoding. See cuda_check_dwarf2_reg_ptx_virtual_register in cuda-tdep.c, 155 // https://github.com/NVIDIA/cuda-gdb/blob/e5cf3bddae520ffb326f95b4d98ce5c7474b828b/gdb/cuda/cuda-tdep.c#L353 156 // IE the bytes of the string are concatenated in reverse into a single 157 // number, which is stored in ULEB128, but in practice must be no more than 8 158 // bytes (excluding null terminator, which is not included). 159 uint64_t result = 0; 160 for (unsigned char c : registerName) 161 result = (result << 8) | c; 162 return result; 163 } 164 165 void NVPTXRegisterInfo::addToDebugRegisterMap( 166 uint64_t preEncodedVirtualRegister, std::string registerName) const { 167 uint64_t mapped = encodeRegisterForDwarf(registerName); 168 if (mapped == 0) 169 return; 170 debugRegisterMap.insert({preEncodedVirtualRegister, mapped}); 171 } 172 173 int64_t NVPTXRegisterInfo::getDwarfRegNum(MCRegister RegNum, bool isEH) const { 174 if (RegNum.isPhysical()) { 175 std::string name = NVPTXInstPrinter::getRegisterName(RegNum.id()); 176 // In NVPTXFrameLowering.cpp, we do arrange for %Depot to be accessible from 177 // %SP. Using the %Depot register doesn't provide any debug info in 178 // cuda-gdb, but switching it to %SP does. 179 if (RegNum.id() == NVPTX::VRDepot) 180 name = "%SP"; 181 return encodeRegisterForDwarf(name); 182 } 183 uint64_t lookup = debugRegisterMap.lookup(RegNum.id()); 184 if (lookup) 185 return lookup; 186 return -1; 187 } 188