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 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 95 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 NVPTXRegisterInfo::getCalleeSavedRegs(const MachineFunction *) const { 103 static const MCPhysReg CalleeSavedRegs[] = { 0 }; 104 return CalleeSavedRegs; 105 } 106 107 BitVector NVPTXRegisterInfo::getReservedRegs(const MachineFunction &MF) const { 108 BitVector Reserved(getNumRegs()); 109 return Reserved; 110 } 111 112 void 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(NVPTX::VRFrame, false); 126 MI.getOperand(FIOperandNum + 1).ChangeToImmediate(Offset); 127 } 128 129 Register NVPTXRegisterInfo::getFrameRegister(const MachineFunction &MF) const { 130 return NVPTX::VRFrame; 131 } 132