Home
last modified time | relevance | path

Searched refs:AMDGPU (Results 1 – 25 of 211) sorted by relevance

123456789

/freebsd/contrib/llvm-project/llvm/lib/Target/AMDGPU/
H A DSIRegisterInfo.cpp96 Register TmpVGPR = AMDGPU::NoRegister;
102 Register SavedExecReg = AMDGPU::NoRegister;
136 ExecReg = AMDGPU::EXEC_LO; in SGPRSpillBuilder()
137 MovOpc = AMDGPU::S_MOV_B32; in SGPRSpillBuilder()
138 NotOpc = AMDGPU::S_NOT_B32; in SGPRSpillBuilder()
140 ExecReg = AMDGPU::EXEC; in SGPRSpillBuilder()
141 MovOpc = AMDGPU::S_MOV_B64; in SGPRSpillBuilder()
142 NotOpc = AMDGPU::S_NOT_B64; in SGPRSpillBuilder()
145 assert(SuperReg != AMDGPU::M0 && "m0 should never spill"); in SGPRSpillBuilder()
146 assert(SuperReg != AMDGPU::EXEC_LO && SuperReg != AMDGPU::EXEC_HI && in SGPRSpillBuilder()
[all …]
H A DAMDGPUResourceUsageAnalysis.cpp29 using namespace llvm::AMDGPU;
86 if (AMDGPU::getAMDHSACodeObjectVersion(*MF.getFunction().getParent()) >= in runOnMachineFunction()
87 AMDGPU::AMDHSA_COV5 || in runOnMachineFunction()
113 if (AMDGPU::getAMDHSACodeObjectVersion(*MF.getFunction().getParent()) >= in run()
114 AMDGPU::AMDHSA_COV5 || in run()
140 Info.UsesFlatScratch = MRI.isPhysRegUsed(AMDGPU::FLAT_SCR_LO) || in analyzeResourceUsage()
141 MRI.isPhysRegUsed(AMDGPU::FLAT_SCR_HI) || in analyzeResourceUsage()
152 (!hasAnyNonFlatUseOfReg(MRI, *TII, AMDGPU::FLAT_SCR) && in analyzeResourceUsage()
153 !hasAnyNonFlatUseOfReg(MRI, *TII, AMDGPU::FLAT_SCR_LO) && in analyzeResourceUsage()
154 !hasAnyNonFlatUseOfReg(MRI, *TII, AMDGPU::FLAT_SCR_HI))) { in analyzeResourceUsage()
[all …]
H A DSIInstrInfo.cpp43 namespace llvm::AMDGPU { namespace
64 : AMDGPUGenInstrInfo(AMDGPU::ADJCALLSTACKUP, AMDGPU::ADJCALLSTACKDOWN), in SIInstrInfo()
83 AMDGPU::OpName OpName) { in nodesHaveSameOperandValue()
87 int Op0Idx = AMDGPU::getNamedOperandIdx(Opc0, OpName); in nodesHaveSameOperandValue()
88 int Op1Idx = AMDGPU::getNamedOperandIdx(Opc1, OpName); in nodesHaveSameOperandValue()
154 const MachineOperand *Dst = getNamedOperand(MI, AMDGPU::OpName::sdst); in resultDependsOnExec()
165 case AMDGPU::S_AND_SAVEEXEC_B32: in resultDependsOnExec()
166 case AMDGPU::S_AND_SAVEEXEC_B64: in resultDependsOnExec()
168 case AMDGPU::S_AND_B32: in resultDependsOnExec()
169 case AMDGPU::S_AND_B64: in resultDependsOnExec()
[all …]
H A DSILoadStoreOptimizer.cpp168 AddrOp->getReg() != AMDGPU::SGPR_NULL) in hasMergeableAddress()
232 AMDGPU::OpName OpName, Register DestReg) const;
235 AMDGPU::OpName OpName) const;
335 return AMDGPU::getMUBUFElements(Opc); in getOpcodeWidth()
339 TII.getNamedOperand(MI, AMDGPU::OpName::dmask)->getImm(); in getOpcodeWidth()
343 return AMDGPU::getMTBUFElements(Opc); in getOpcodeWidth()
347 case AMDGPU::S_BUFFER_LOAD_DWORD_IMM: in getOpcodeWidth()
348 case AMDGPU::S_BUFFER_LOAD_DWORD_SGPR_IMM: in getOpcodeWidth()
349 case AMDGPU::S_LOAD_DWORD_IMM: in getOpcodeWidth()
350 case AMDGPU::GLOBAL_LOAD_DWORD: in getOpcodeWidth()
[all …]
H A DGCNSubtarget.cpp102 } else if (!hasFeature(AMDGPU::FeatureWavefrontSize32) && in initializeSubtargetDependencies()
103 !hasFeature(AMDGPU::FeatureWavefrontSize64)) { in initializeSubtargetDependencies()
107 ToggleFeature(AMDGPU::FeatureWavefrontSize32); in initializeSubtargetDependencies()
123 ToggleFeature(AMDGPU::FeatureFlatForGlobal); in initializeSubtargetDependencies()
129 ToggleFeature(AMDGPU::FeatureFlatForGlobal); in initializeSubtargetDependencies()
144 if (AMDGPU::isGFX10Plus(*this) && in initializeSubtargetDependencies()
145 !getFeatureBits().test(AMDGPU::FeatureCuMode)) in initializeSubtargetDependencies()
163 if (hasFeature(AMDGPU::FeatureWavefrontSize32) && in checkSubtargetFeatures()
164 hasFeature(AMDGPU::FeatureWavefrontSize64)) { in checkSubtargetFeatures()
182 MaxWavesPerEU = AMDGPU::IsaInfo::getMaxWavesPerEU(this); in GCNSubtarget()
[all …]
H A DAMDGPURegisterBankInfo.cpp124 if (Opc == AMDGPU::G_ANYEXT || Opc == AMDGPU::G_ZEXT || in applyBank()
125 Opc == AMDGPU::G_SEXT) { in applyBank()
132 if (SrcBank == &AMDGPU::VCCRegBank) { in applyBank()
136 assert(NewBank == &AMDGPU::VGPRRegBank); in applyBank()
142 auto True = B.buildConstant(S32, Opc == AMDGPU::G_SEXT ? -1 : 1); in applyBank()
156 if (Opc == AMDGPU::G_TRUNC) { in applyBank()
159 assert(DstBank != &AMDGPU::VCCRegBank); in applyBank()
174 assert(NewBank == &AMDGPU::VGPRRegBank && in applyBank()
176 assert((MI.getOpcode() != AMDGPU::G_TRUNC && in applyBank()
177 MI.getOpcode() != AMDGPU::G_ANYEXT) && in applyBank()
[all …]
H A DAMDGPUCombinerHelper.cpp30 case AMDGPU::G_FADD: in fnegFoldsIntoMI()
31 case AMDGPU::G_FSUB: in fnegFoldsIntoMI()
32 case AMDGPU::G_FMUL: in fnegFoldsIntoMI()
33 case AMDGPU::G_FMA: in fnegFoldsIntoMI()
34 case AMDGPU::G_FMAD: in fnegFoldsIntoMI()
35 case AMDGPU::G_FMINNUM: in fnegFoldsIntoMI()
36 case AMDGPU::G_FMAXNUM: in fnegFoldsIntoMI()
37 case AMDGPU::G_FMINNUM_IEEE: in fnegFoldsIntoMI()
38 case AMDGPU::G_FMAXNUM_IEEE: in fnegFoldsIntoMI()
39 case AMDGPU::G_FMINIMUM: in fnegFoldsIntoMI()
[all …]
H A DSIOptimizeExecMasking.cpp117 case AMDGPU::COPY: in isCopyFromExec()
118 case AMDGPU::S_MOV_B64: in isCopyFromExec()
119 case AMDGPU::S_MOV_B64_term: in isCopyFromExec()
120 case AMDGPU::S_MOV_B32: in isCopyFromExec()
121 case AMDGPU::S_MOV_B32_term: { in isCopyFromExec()
128 return AMDGPU::NoRegister; in isCopyFromExec()
134 case AMDGPU::COPY: in isCopyToExec()
135 case AMDGPU::S_MOV_B64: in isCopyToExec()
136 case AMDGPU::S_MOV_B32: { in isCopyToExec()
142 case AMDGPU::S_MOV_B64_term: in isCopyToExec()
[all …]
H A DGCNDPPCombine.cpp73 bool hasNoImmOrEqual(MachineInstr &MI, AMDGPU::OpName OpndName, int64_t Value,
129 if (AMDGPU::isTrue16Inst(Op)) in isShrinkable()
131 if (const auto *SDst = TII->getNamedOperand(MI, AMDGPU::OpName::sdst)) { in isShrinkable()
141 if (!hasNoImmOrEqual(MI, AMDGPU::OpName::src0_modifiers, 0, Mask) || in isShrinkable()
142 !hasNoImmOrEqual(MI, AMDGPU::OpName::src1_modifiers, 0, Mask) || in isShrinkable()
143 !hasNoImmOrEqual(MI, AMDGPU::OpName::clamp, 0) || in isShrinkable()
144 !hasNoImmOrEqual(MI, AMDGPU::OpName::omod, 0) || in isShrinkable()
145 !hasNoImmOrEqual(MI, AMDGPU::OpName::byte_sel, 0)) { in isShrinkable()
153 int DPP32 = AMDGPU::getDPPOp32(Op); in getDPPOp()
156 int E32 = AMDGPU::getVOPe32(Op); in getDPPOp()
[all …]
H A DSIPeepholeSDWA.cpp90 using namespace AMDGPU::SDWA;
365 if (TII->getNamedOperand(*MI, AMDGPU::OpName::src0) == SrcOp) { in getSrcMods()
366 if (auto *Mod = TII->getNamedOperand(*MI, AMDGPU::OpName::src0_modifiers)) { in getSrcMods()
369 } else if (TII->getNamedOperand(*MI, AMDGPU::OpName::src1) == SrcOp) { in getSrcMods()
370 if (auto *Mod = TII->getNamedOperand(*MI, AMDGPU::OpName::src1_modifiers)) { in getSrcMods()
427 case AMDGPU::V_CVT_F32_FP8_sdwa: in convertToSDWA()
428 case AMDGPU::V_CVT_F32_BF8_sdwa: in convertToSDWA()
429 case AMDGPU::V_CVT_PK_F32_FP8_sdwa: in convertToSDWA()
430 case AMDGPU::V_CVT_PK_F32_BF8_sdwa: in convertToSDWA()
433 case AMDGPU::V_CNDMASK_B32_sdwa: in convertToSDWA()
[all …]
H A DSIShrinkInstructions.cpp100 int Src0Idx = AMDGPU::getNamedOperandIdx(MI.getOpcode(), AMDGPU::OpName::src0); in foldImmediates()
159 if (AMDGPU::VGPR_32RegClass.contains(Reg) && in shouldShrinkTrue16()
160 !AMDGPU::VGPR_32_Lo128RegClass.contains(Reg)) in shouldShrinkTrue16()
163 if (AMDGPU::VGPR_16RegClass.contains(Reg) && in shouldShrinkTrue16()
164 !AMDGPU::VGPR_16_Lo128RegClass.contains(Reg)) in shouldShrinkTrue16()
218 return AMDGPU::V_NOT_B32_e32; in canModifyToInlineImmOp32()
223 return Scalar ? AMDGPU::S_BREV_B32 : AMDGPU::V_BFREV_B32_e32; in canModifyToInlineImmOp32()
262 int SOPKOpc = AMDGPU::getSOPKOp(MI.getOpcode()); in shrinkScalarCompare()
268 if (SOPKOpc == AMDGPU::S_CMPK_EQ_U32 || SOPKOpc == AMDGPU::S_CMPK_LG_U32) { in shrinkScalarCompare()
272 SOPKOpc = (SOPKOpc == AMDGPU::S_CMPK_EQ_U32) ? in shrinkScalarCompare()
[all …]
H A DAMDGPUInstructionSelector.cpp72 return Def->getOpcode() == AMDGPU::G_AMDGPU_WAVE_ADDRESS in getWaveAddress()
91 return MRI.getVRegDef(Reg)->getOpcode() != AMDGPU::G_TRUNC && in isVCC()
96 return RB->getID() == AMDGPU::VCCRegBankID; in isVCC()
103 MI.addOperand(*MF, MachineOperand::CreateReg(AMDGPU::EXEC, false, true)); in constrainCopyLikeIntrin()
134 if (SrcReg == AMDGPU::SCC) { in selectCOPY()
154 STI.isWave64() ? AMDGPU::S_MOV_B64 : AMDGPU::S_MOV_B32; in selectCOPY()
164 if (AMDGPU::getRegBitWidth(SrcRC->getID()) == 16) { in selectCOPY()
167 BuildMI(*BB, &I, DL, TII.get(AMDGPU::V_AND_B16_t16_e64), MaskedReg) in selectCOPY()
173 BuildMI(*BB, &I, DL, TII.get(AMDGPU::V_CMP_NE_U16_t16_e64), DstReg) in selectCOPY()
181 unsigned AndOpc = IsSGPR ? AMDGPU::S_AND_B32 : AMDGPU::V_AND_B32_e32; in selectCOPY()
[all …]
H A DSIFoldOperands.cpp44 unsigned DefSubReg = AMDGPU::NoSubRegister;
51 unsigned DefSubReg = AMDGPU::NoSubRegister) in FoldableDef()
67 unsigned DefSubReg = AMDGPU::NoSubRegister) in FoldableDef()
127 if (DefSubReg != AMDGPU::NoSubRegister) in isOperandLegal()
135 if (DefSubReg != AMDGPU::NoSubRegister) in isOperandLegal()
188 case AMDGPU::S_ADD_I32: { in convertToVALUOp()
190 return UseVOP3 ? AMDGPU::V_ADD_U32_e64 : AMDGPU::V_ADD_U32_e32; in convertToVALUOp()
191 return UseVOP3 ? AMDGPU::V_ADD_CO_U32_e64 : AMDGPU::V_ADD_CO_U32_e32; in convertToVALUOp()
193 case AMDGPU::S_OR_B32: in convertToVALUOp()
194 return UseVOP3 ? AMDGPU::V_OR_B32_e64 : AMDGPU::V_OR_B32_e32; in convertToVALUOp()
[all …]
H A DGCNHazardRecognizer.cpp64 MaxLookAhead = MF.getRegInfo().isPhysRegUsed(AMDGPU::AGPR0) ? 19 : 5; in GCNHazardRecognizer()
81 return Opcode == AMDGPU::V_DIV_FMAS_F32_e64 || Opcode == AMDGPU::V_DIV_FMAS_F64_e64; in isDivFMas()
85 return Opcode == AMDGPU::S_GETREG_B32; in isSGetReg()
90 case AMDGPU::S_SETREG_B32: in isSSetReg()
91 case AMDGPU::S_SETREG_B32_mode: in isSSetReg()
92 case AMDGPU::S_SETREG_IMM32_B32: in isSSetReg()
93 case AMDGPU::S_SETREG_IMM32_B32_mode: in isSSetReg()
100 return Opcode == AMDGPU::V_READLANE_B32 || Opcode == AMDGPU::V_WRITELANE_B32; in isRWLane()
104 return Opcode == AMDGPU::S_RFE_B64; in isRFE()
109 case AMDGPU::S_MOVRELS_B32: in isSMovRel()
[all …]
H A DSILateBranchLowering.cpp86 bool HasColorExports = AMDGPU::getHasColorExport(F); in generateEndPgm()
87 bool HasDepthExports = AMDGPU::getHasDepthExport(F); in generateEndPgm()
91 bool MustExport = !AMDGPU::isGFX10Plus(TII->getSubtarget()); in generateEndPgm()
98 ? AMDGPU::Exp::ET_NULL in generateEndPgm()
99 : (HasColorExports ? AMDGPU::Exp::ET_MRT0 : AMDGPU::Exp::ET_MRTZ); in generateEndPgm()
100 BuildMI(MBB, I, DL, TII->get(AMDGPU::EXP_DONE)) in generateEndPgm()
102 .addReg(AMDGPU::VGPR0, RegState::Undef) in generateEndPgm()
103 .addReg(AMDGPU::VGPR0, RegState::Undef) in generateEndPgm()
104 .addReg(AMDGPU::VGPR0, RegState::Undef) in generateEndPgm()
105 .addReg(AMDGPU::VGPR0, RegState::Undef) in generateEndPgm()
[all …]
H A DAMDGPUArgumentUsageInfo.cpp95 &AMDGPU::SGPR_128RegClass, LLT::fixed_vector(4, 32)); in getPreloadedValue()
99 &AMDGPU::SGPR_64RegClass, in getPreloadedValue()
103 &AMDGPU::SGPR_32RegClass, LLT::scalar(32)); in getPreloadedValue()
106 &AMDGPU::SGPR_32RegClass, LLT::scalar(32)); in getPreloadedValue()
109 &AMDGPU::SGPR_32RegClass, LLT::scalar(32)); in getPreloadedValue()
112 &AMDGPU::SGPR_32RegClass, LLT::scalar(32)); in getPreloadedValue()
116 &AMDGPU::SGPR_32RegClass, LLT::scalar(32)); in getPreloadedValue()
119 &AMDGPU::SGPR_32RegClass, LLT::scalar(32)}; in getPreloadedValue()
122 &AMDGPU::SGPR_64RegClass, in getPreloadedValue()
126 &AMDGPU::SGPR_64RegClass, in getPreloadedValue()
[all …]
H A DAMDGPUWaitSGPRHazards.cpp63 case AMDGPU::M0: in sgprNumber()
64 case AMDGPU::EXEC: in sgprNumber()
65 case AMDGPU::EXEC_LO: in sgprNumber()
66 case AMDGPU::EXEC_HI: in sgprNumber()
67 case AMDGPU::SGPR_NULL: in sgprNumber()
68 case AMDGPU::SGPR_NULL64: in sgprNumber()
80 return Reg == AMDGPU::VCC || Reg == AMDGPU::VCC_LO || Reg == AMDGPU::VCC_HI; in isVCC()
97 if (I->getOpcode() != AMDGPU::S_GETPC_B64) in updateGetPCBundle()
102 assert(NewMI->getOpcode() == AMDGPU::S_WAITCNT_DEPCTR && in updateGetPCBundle()
164 BuildMI(MBB, MI, MI->getDebugLoc(), TII->get(AMDGPU::DS_NOP)); in insertHazardCull()
[all …]
H A DGCNCreateVOPD.cpp63 AMDGPU::getVOPDEncodingFamily(SII->getSubtarget()); in doReplace()
64 int NewOpcode = AMDGPU::getVOPDFull(AMDGPU::getVOPDOpcode(Opc1, CI.IsVOPD3), in doReplace()
65 AMDGPU::getVOPDOpcode(Opc2, CI.IsVOPD3), in doReplace()
74 namespace VOPD = AMDGPU::VOPD; in doReplace()
77 AMDGPU::getVOPDInstInfo(FirstMI->getDesc(), SecondMI->getDesc()); in doReplace()
84 const AMDGPU::OpName Mods[2][3] = { in doReplace()
85 {AMDGPU::OpName::src0X_modifiers, AMDGPU::OpName::vsrc1X_modifiers, in doReplace()
86 AMDGPU::OpName::vsrc2X_modifiers}, in doReplace()
87 {AMDGPU::OpName::src0Y_modifiers, AMDGPU::OpName::vsrc1Y_modifiers, in doReplace()
88 AMDGPU::OpName::vsrc2Y_modifiers}}; in doReplace()
[all …]
/freebsd/contrib/llvm-project/clang/lib/CodeGen/TargetBuiltins/
H A DAMDGPU.cpp303 case AMDGPU::BI__builtin_amdgcn_div_scale: in EmitAMDGPUBuiltinExpr()
304 case AMDGPU::BI__builtin_amdgcn_div_scalef: { in EmitAMDGPUBuiltinExpr()
328 case AMDGPU::BI__builtin_amdgcn_div_fmas: in EmitAMDGPUBuiltinExpr()
329 case AMDGPU::BI__builtin_amdgcn_div_fmasf: { in EmitAMDGPUBuiltinExpr()
341 case AMDGPU::BI__builtin_amdgcn_ds_swizzle: in EmitAMDGPUBuiltinExpr()
344 case AMDGPU::BI__builtin_amdgcn_mov_dpp8: in EmitAMDGPUBuiltinExpr()
345 case AMDGPU::BI__builtin_amdgcn_mov_dpp: in EmitAMDGPUBuiltinExpr()
346 case AMDGPU::BI__builtin_amdgcn_update_dpp: { in EmitAMDGPUBuiltinExpr()
359 CGM.getIntrinsic(BuiltinID == AMDGPU::BI__builtin_amdgcn_mov_dpp8 in EmitAMDGPUBuiltinExpr()
365 bool InsertOld = BuiltinID == AMDGPU::BI__builtin_amdgcn_mov_dpp; in EmitAMDGPUBuiltinExpr()
[all …]
/freebsd/contrib/llvm-project/llvm/lib/Target/AMDGPU/MCTargetDesc/
H A DAMDGPUMCCodeEmitter.cpp155 STI.hasFeature(AMDGPU::FeatureInv2PiInlineImm)) in getLit16Encoding()
212 STI.hasFeature(AMDGPU::FeatureInv2PiInlineImm)) in getLit32Encoding()
253 STI.hasFeature(AMDGPU::FeatureInv2PiInlineImm)) in getLit64Encoding()
259 return STI.hasFeature(AMDGPU::Feature64BitLiterals) && Lo_32(Val) ? 254 in getLit64Encoding()
263 return STI.hasFeature(AMDGPU::Feature64BitLiterals) && in getLit64Encoding()
275 return (STI.hasFeature(AMDGPU::Feature64BitLiterals) && in getLitEncoding()
276 OpInfo.OperandType == AMDGPU::OPERAND_REG_IMM_INT64) in getLitEncoding()
289 case AMDGPU::OPERAND_REG_IMM_INT32: in getLitEncoding()
290 case AMDGPU::OPERAND_REG_IMM_FP32: in getLitEncoding()
291 case AMDGPU::OPERAND_REG_INLINE_C_INT32: in getLitEncoding()
[all …]
H A DAMDGPUInstPrinter.cpp25 using namespace llvm::AMDGPU;
85 if (STI.hasFeature(AMDGPU::Feature64BitLiterals) && Lo_32(Imm)) in printFP64ImmOperand()
108 if (AMDGPU::isGFX12(STI) && IsVBuffer) in printOffset()
125 AMDGPU::isGFX12(STI); in printFlatOffset()
128 O << formatDec(SignExtend32(Imm, AMDGPU::getNumFlatOffsetBits(STI))); in printFlatOffset()
156 if (AMDGPU::isGFX12Plus(STI)) { in printCPol()
167 O << ((AMDGPU::isGFX940(STI) && in printCPol()
171 O << (AMDGPU::isGFX940(STI) ? " nt" : " slc"); in printCPol()
172 if ((Imm & CPol::DLC) && AMDGPU::isGFX10Plus(STI)) in printCPol()
174 if ((Imm & CPol::SCC) && AMDGPU::isGFX90A(STI)) in printCPol()
[all …]
/freebsd/contrib/llvm-project/llvm/lib/Target/AMDGPU/MCA/
H A DAMDGPUCustomBehaviour.cpp27 case AMDGPU::S_WAITCNT: in postProcessInstruction()
28 case AMDGPU::S_WAITCNT_soft: in postProcessInstruction()
29 case AMDGPU::S_WAITCNT_EXPCNT: in postProcessInstruction()
30 case AMDGPU::S_WAITCNT_LGKMCNT: in postProcessInstruction()
31 case AMDGPU::S_WAITCNT_VMCNT: in postProcessInstruction()
32 case AMDGPU::S_WAITCNT_VSCNT: in postProcessInstruction()
33 case AMDGPU::S_WAITCNT_VSCNT_soft: in postProcessInstruction()
34 case AMDGPU::S_WAITCNT_EXPCNT_gfx10: in postProcessInstruction()
35 case AMDGPU::S_WAITCNT_LGKMCNT_gfx10: in postProcessInstruction()
36 case AMDGPU::S_WAITCNT_VMCNT_gfx10: in postProcessInstruction()
[all …]
/freebsd/contrib/llvm-project/llvm/lib/Target/AMDGPU/Disassembler/
H A DAMDGPUDisassembler.cpp44 (isGFX10Plus() ? AMDGPU::EncValues::SGPR_MAX_GFX10 \
45 : AMDGPU::EncValues::SGPR_MAX_SI)
58 CodeObjectVersion(AMDGPU::getDefaultAMDHSACodeObjectVersion()) { in AMDGPUDisassembler()
60 if (!STI.hasFeature(AMDGPU::FeatureGCN3Encoding) && !isGFX10Plus()) in AMDGPUDisassembler()
63 for (auto [Symbol, Code] : AMDGPU::UCVersion::getGFXVersions()) in AMDGPUDisassembler()
72 CodeObjectVersion = AMDGPU::getAMDHSACodeObjectVersion(Version); in setABIVersion()
84 AMDGPU::OpName Name) { in insertNamedMCOperand()
85 int OpIdx = AMDGPU::getNamedOperandIdx(MI.getOpcode(), Name); in insertNamedMCOperand()
157 Inst, DAsm->createRegOperand(AMDGPU::RegClass##RegClassID, Imm)); \
191 return decodeSrcOp(Inst, 10, OpWidth, Imm, Imm | AMDGPU::EncValues::IS_VGPR, in decodeAV10()
[all …]
/freebsd/contrib/llvm-project/llvm/lib/Target/AMDGPU/AsmParser/
H A DAMDGPUAsmParser.cpp49 using namespace llvm::AMDGPU;
287 return isRegOrImmWithInputMods(AMDGPU::VS_32RegClassID, MVT::i16); in isRegOrImmWithInt16InputMods()
292 IsFake16 ? AMDGPU::VS_32RegClassID : AMDGPU::VS_16RegClassID, MVT::i16); in isRegOrImmWithIntT16InputMods()
296 return isRegOrImmWithInputMods(AMDGPU::VS_32RegClassID, MVT::i32); in isRegOrImmWithInt32InputMods()
300 return isRegOrInline(AMDGPU::VS_32RegClassID, MVT::i16); in isRegOrInlineImmWithInt16InputMods()
305 IsFake16 ? AMDGPU::VS_32RegClassID : AMDGPU::VS_16RegClassID, MVT::i16); in isRegOrInlineImmWithIntT16InputMods()
309 return isRegOrInline(AMDGPU::VS_32RegClassID, MVT::i32); in isRegOrInlineImmWithInt32InputMods()
313 return isRegOrImmWithInputMods(AMDGPU::VS_64RegClassID, MVT::i64); in isRegOrImmWithInt64InputMods()
317 return isRegOrImmWithInputMods(AMDGPU::VS_32RegClassID, MVT::f16); in isRegOrImmWithFP16InputMods()
322 IsFake16 ? AMDGPU::VS_32RegClassID : AMDGPU::VS_16RegClassID, MVT::f16); in isRegOrImmWithFPT16InputMods()
[all …]
/freebsd/contrib/llvm-project/llvm/lib/Target/AMDGPU/Utils/
H A DAMDGPUBaseInfo.cpp37 llvm::cl::init(llvm::AMDGPU::AMDHSA_COV6),
188 namespace AMDGPU { namespace
251 return AMDGPU::ImplicitArg::MULTIGRID_SYNC_ARG_OFFSET; in getMultigridSyncArgImplicitArgPosition()
264 return AMDGPU::ImplicitArg::HOSTCALL_PTR_OFFSET; in getHostcallImplicitArgPosition()
275 return AMDGPU::ImplicitArg::DEFAULT_QUEUE_OFFSET; in getDefaultQueueImplicitArgPosition()
286 return AMDGPU::ImplicitArg::COMPLETION_ACTION_OFFSET; in getCompletionActionImplicitArgPosition()
596 if (ST.hasFeature(AMDGPU::FeatureGFX1250Insts)) in getVOPDEncodingFamily()
598 if (ST.hasFeature(AMDGPU::FeatureGFX12Insts)) in getVOPDEncodingFamily()
600 if (ST.hasFeature(AMDGPU::FeatureGFX11Insts)) in getVOPDEncodingFamily()
607 Opc = IsConvertibleToBitOp ? (unsigned)AMDGPU::V_BITOP3_B32_e64 : Opc; in getCanBeVOPD()
[all …]

123456789