32#include "llvm/IR/IntrinsicsSPIRV.h"
36#define DEBUG_TYPE "spirv-isel"
43 std::vector<std::pair<SPIRV::InstructionSet::InstructionSet, uint32_t>>;
47llvm::SPIRV::SelectionControl::SelectionControl
48getSelectionOperandForImm(
int Imm) {
50 return SPIRV::SelectionControl::Flatten;
52 return SPIRV::SelectionControl::DontFlatten;
54 return SPIRV::SelectionControl::None;
58#define GET_GLOBALISEL_PREDICATE_BITSET
59#include "SPIRVGenGlobalISel.inc"
60#undef GET_GLOBALISEL_PREDICATE_BITSET
87#define GET_GLOBALISEL_PREDICATES_DECL
88#include "SPIRVGenGlobalISel.inc"
89#undef GET_GLOBALISEL_PREDICATES_DECL
91#define GET_GLOBALISEL_TEMPORARIES_DECL
92#include "SPIRVGenGlobalISel.inc"
93#undef GET_GLOBALISEL_TEMPORARIES_DECL
117 unsigned BitSetOpcode)
const;
121 unsigned BitSetOpcode)
const;
125 unsigned BitSetOpcode,
bool SwapPrimarySide)
const;
129 unsigned BitSetOpcode,
130 bool SwapPrimarySide)
const;
137 unsigned Opcode)
const;
140 unsigned Opcode)
const;
160 unsigned NegateOpcode = 0)
const;
220 template <
bool Signed>
223 template <
bool Signed>
247 bool IsSigned,
unsigned Opcode)
const;
249 bool IsSigned)
const;
255 bool IsSigned)
const;
294 GL::GLSLExtInst GLInst)
const;
299 GL::GLSLExtInst GLInst)
const;
321 bool selectCounterHandleFromBinding(
Register &ResVReg,
334 bool selectResourceNonUniformIndex(
Register &ResVReg,
346 std::pair<Register, bool>
348 const SPIRVType *ResType =
nullptr)
const;
351 bool isScalarOrVectorIntConstantZero(
Register Reg)
const;
361 SPIRV::StorageClass::StorageClass SC)
const;
368 SPIRV::StorageClass::StorageClass SC,
380 bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
383 bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
388 void decorateUsesAsNonUniform(
Register &NonUniformReg)
const;
392bool sampledTypeIsSignedInteger(
const llvm::Type *HandleType) {
394 if (
TET->getTargetExtName() ==
"spirv.Image") {
397 assert(
TET->getTargetExtName() ==
"spirv.SignedImage");
398 return TET->getTypeParameter(0)->isIntegerTy();
402#define GET_GLOBALISEL_IMPL
403#include "SPIRVGenGlobalISel.inc"
404#undef GET_GLOBALISEL_IMPL
410 TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
413#include
"SPIRVGenGlobalISel.inc"
416#include
"SPIRVGenGlobalISel.inc"
428 InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
432void SPIRVInstructionSelector::resetVRegsType(MachineFunction &MF) {
433 if (HasVRegsReset == &MF)
438 for (
unsigned I = 0,
E =
MRI.getNumVirtRegs();
I !=
E; ++
I) {
440 LLT RegType =
MRI.getType(
Reg);
448 for (
const auto &
MBB : MF) {
449 for (
const auto &
MI :
MBB) {
452 if (
MI.getOpcode() != SPIRV::ASSIGN_TYPE)
456 LLT DstType =
MRI.getType(DstReg);
458 LLT SrcType =
MRI.getType(SrcReg);
459 if (DstType != SrcType)
460 MRI.setType(DstReg,
MRI.getType(SrcReg));
462 const TargetRegisterClass *DstRC =
MRI.getRegClassOrNull(DstReg);
463 const TargetRegisterClass *SrcRC =
MRI.getRegClassOrNull(SrcReg);
464 if (DstRC != SrcRC && SrcRC)
465 MRI.setRegClass(DstReg, SrcRC);
481 case TargetOpcode::G_CONSTANT:
482 case TargetOpcode::G_FCONSTANT:
483 case TargetOpcode::G_IMPLICIT_DEF:
485 case TargetOpcode::G_INTRINSIC:
486 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
487 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
489 Intrinsic::spv_const_composite;
490 case TargetOpcode::G_BUILD_VECTOR:
491 case TargetOpcode::G_SPLAT_VECTOR: {
502 case SPIRV::OpConstantTrue:
503 case SPIRV::OpConstantFalse:
504 case SPIRV::OpConstantI:
505 case SPIRV::OpConstantF:
506 case SPIRV::OpConstantComposite:
507 case SPIRV::OpConstantCompositeContinuedINTEL:
508 case SPIRV::OpConstantSampler:
509 case SPIRV::OpConstantNull:
511 case SPIRV::OpConstantFunctionPointerINTEL:
537 case Intrinsic::spv_all:
538 case Intrinsic::spv_alloca:
539 case Intrinsic::spv_any:
540 case Intrinsic::spv_bitcast:
541 case Intrinsic::spv_const_composite:
542 case Intrinsic::spv_cross:
543 case Intrinsic::spv_degrees:
544 case Intrinsic::spv_distance:
545 case Intrinsic::spv_extractelt:
546 case Intrinsic::spv_extractv:
547 case Intrinsic::spv_faceforward:
548 case Intrinsic::spv_fdot:
549 case Intrinsic::spv_firstbitlow:
550 case Intrinsic::spv_firstbitshigh:
551 case Intrinsic::spv_firstbituhigh:
552 case Intrinsic::spv_frac:
553 case Intrinsic::spv_gep:
554 case Intrinsic::spv_global_offset:
555 case Intrinsic::spv_global_size:
556 case Intrinsic::spv_group_id:
557 case Intrinsic::spv_insertelt:
558 case Intrinsic::spv_insertv:
559 case Intrinsic::spv_isinf:
560 case Intrinsic::spv_isnan:
561 case Intrinsic::spv_lerp:
562 case Intrinsic::spv_length:
563 case Intrinsic::spv_normalize:
564 case Intrinsic::spv_num_subgroups:
565 case Intrinsic::spv_num_workgroups:
566 case Intrinsic::spv_ptrcast:
567 case Intrinsic::spv_radians:
568 case Intrinsic::spv_reflect:
569 case Intrinsic::spv_refract:
570 case Intrinsic::spv_resource_getpointer:
571 case Intrinsic::spv_resource_handlefrombinding:
572 case Intrinsic::spv_resource_handlefromimplicitbinding:
573 case Intrinsic::spv_resource_nonuniformindex:
574 case Intrinsic::spv_resource_sample:
575 case Intrinsic::spv_rsqrt:
576 case Intrinsic::spv_saturate:
577 case Intrinsic::spv_sdot:
578 case Intrinsic::spv_sign:
579 case Intrinsic::spv_smoothstep:
580 case Intrinsic::spv_step:
581 case Intrinsic::spv_subgroup_id:
582 case Intrinsic::spv_subgroup_local_invocation_id:
583 case Intrinsic::spv_subgroup_max_size:
584 case Intrinsic::spv_subgroup_size:
585 case Intrinsic::spv_thread_id:
586 case Intrinsic::spv_thread_id_in_group:
587 case Intrinsic::spv_udot:
588 case Intrinsic::spv_undef:
589 case Intrinsic::spv_value_md:
590 case Intrinsic::spv_workgroup_size:
602 case SPIRV::OpTypeVoid:
603 case SPIRV::OpTypeBool:
604 case SPIRV::OpTypeInt:
605 case SPIRV::OpTypeFloat:
606 case SPIRV::OpTypeVector:
607 case SPIRV::OpTypeMatrix:
608 case SPIRV::OpTypeImage:
609 case SPIRV::OpTypeSampler:
610 case SPIRV::OpTypeSampledImage:
611 case SPIRV::OpTypeArray:
612 case SPIRV::OpTypeRuntimeArray:
613 case SPIRV::OpTypeStruct:
614 case SPIRV::OpTypeOpaque:
615 case SPIRV::OpTypePointer:
616 case SPIRV::OpTypeFunction:
617 case SPIRV::OpTypeEvent:
618 case SPIRV::OpTypeDeviceEvent:
619 case SPIRV::OpTypeReserveId:
620 case SPIRV::OpTypeQueue:
621 case SPIRV::OpTypePipe:
622 case SPIRV::OpTypeForwardPointer:
623 case SPIRV::OpTypePipeStorage:
624 case SPIRV::OpTypeNamedBarrier:
625 case SPIRV::OpTypeAccelerationStructureNV:
626 case SPIRV::OpTypeCooperativeMatrixNV:
627 case SPIRV::OpTypeCooperativeMatrixKHR:
637 if (
MI.getNumDefs() == 0)
640 for (
const auto &MO :
MI.all_defs()) {
642 if (
Reg.isPhysical()) {
646 for (
const auto &
UseMI :
MRI.use_nodbg_instructions(
Reg)) {
647 if (
UseMI.getOpcode() != SPIRV::OpName) {
654 if (
MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE ||
MI.isFakeUse() ||
655 MI.isLifetimeMarker()) {
658 <<
"Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
669 if (
MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
670 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
673 LLVM_DEBUG(
dbgs() <<
"Dead: Intrinsic with no real side effects.\n");
678 if (
MI.mayStore() ||
MI.isCall() ||
679 (
MI.mayLoad() &&
MI.hasOrderedMemoryRef()) ||
MI.isPosition() ||
680 MI.isDebugInstr() ||
MI.isTerminator() ||
MI.isJumpTableDebugInfo()) {
681 LLVM_DEBUG(
dbgs() <<
"Not dead: instruction has side effects.\n");
692 LLVM_DEBUG(
dbgs() <<
"Dead: known opcode with no side effects\n");
699void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &
MI)
const {
701 for (
const auto &MO :
MI.all_defs()) {
705 SmallVector<MachineInstr *, 4> UselessOpNames;
706 for (MachineInstr &
UseMI :
MRI->use_nodbg_instructions(
Reg)) {
708 "There is still a use of the dead function.");
711 for (MachineInstr *OpNameMI : UselessOpNames) {
713 OpNameMI->eraseFromParent();
718void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &
MI)
const {
721 removeOpNamesForDeadMI(
MI);
722 MI.eraseFromParent();
725bool SPIRVInstructionSelector::select(MachineInstr &
I) {
726 resetVRegsType(*
I.getParent()->getParent());
728 assert(
I.getParent() &&
"Instruction should be in a basic block!");
729 assert(
I.getParent()->getParent() &&
"Instruction should be in a function!");
734 removeDeadInstruction(
I);
741 if (Opcode == SPIRV::ASSIGN_TYPE) {
742 Register DstReg =
I.getOperand(0).getReg();
743 Register SrcReg =
I.getOperand(1).getReg();
744 auto *
Def =
MRI->getVRegDef(SrcReg);
746 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
747 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
749 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
750 Register SelectDstReg =
Def->getOperand(0).getReg();
754 Def->removeFromParent();
755 MRI->replaceRegWith(DstReg, SelectDstReg);
757 I.removeFromParent();
759 Res = selectImpl(
I, *CoverageInfo);
761 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
762 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
766 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
773 MRI->setRegClass(SrcReg,
MRI->getRegClass(DstReg));
774 MRI->replaceRegWith(SrcReg, DstReg);
776 I.removeFromParent();
778 }
else if (
I.getNumDefs() == 1) {
785 if (DeadMIs.contains(&
I)) {
789 removeDeadInstruction(
I);
793 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
794 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
800 bool HasDefs =
I.getNumDefs() > 0;
803 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
804 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
805 if (spvSelect(ResVReg, ResType,
I)) {
807 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
810 I.removeFromParent();
818 case TargetOpcode::G_CONSTANT:
819 case TargetOpcode::G_FCONSTANT:
821 case TargetOpcode::G_SADDO:
822 case TargetOpcode::G_SSUBO:
829 MachineInstr &
I)
const {
830 const TargetRegisterClass *DstRC =
MRI->getRegClassOrNull(DestReg);
831 const TargetRegisterClass *SrcRC =
MRI->getRegClassOrNull(SrcReg);
832 if (DstRC != SrcRC && SrcRC)
833 MRI->setRegClass(DestReg, SrcRC);
834 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
835 TII.get(TargetOpcode::COPY))
841bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
843 MachineInstr &
I)
const {
844 const unsigned Opcode =
I.getOpcode();
846 return selectImpl(
I, *CoverageInfo);
848 case TargetOpcode::G_CONSTANT:
849 case TargetOpcode::G_FCONSTANT:
850 return selectConst(ResVReg, ResType,
I);
851 case TargetOpcode::G_GLOBAL_VALUE:
852 return selectGlobalValue(ResVReg,
I);
853 case TargetOpcode::G_IMPLICIT_DEF:
854 return selectOpUndef(ResVReg, ResType,
I);
855 case TargetOpcode::G_FREEZE:
856 return selectFreeze(ResVReg, ResType,
I);
858 case TargetOpcode::G_INTRINSIC:
859 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
860 case TargetOpcode::G_INTRINSIC_CONVERGENT:
861 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
862 return selectIntrinsic(ResVReg, ResType,
I);
863 case TargetOpcode::G_BITREVERSE:
864 return selectBitreverse(ResVReg, ResType,
I);
866 case TargetOpcode::G_BUILD_VECTOR:
867 return selectBuildVector(ResVReg, ResType,
I);
868 case TargetOpcode::G_SPLAT_VECTOR:
869 return selectSplatVector(ResVReg, ResType,
I);
871 case TargetOpcode::G_SHUFFLE_VECTOR: {
872 MachineBasicBlock &BB = *
I.getParent();
873 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
876 .
addUse(
I.getOperand(1).getReg())
877 .
addUse(
I.getOperand(2).getReg());
878 for (
auto V :
I.getOperand(3).getShuffleMask())
882 case TargetOpcode::G_MEMMOVE:
883 case TargetOpcode::G_MEMCPY:
884 case TargetOpcode::G_MEMSET:
885 return selectMemOperation(ResVReg,
I);
887 case TargetOpcode::G_ICMP:
888 return selectICmp(ResVReg, ResType,
I);
889 case TargetOpcode::G_FCMP:
890 return selectFCmp(ResVReg, ResType,
I);
892 case TargetOpcode::G_FRAME_INDEX:
893 return selectFrameIndex(ResVReg, ResType,
I);
895 case TargetOpcode::G_LOAD:
896 return selectLoad(ResVReg, ResType,
I);
897 case TargetOpcode::G_STORE:
898 return selectStore(
I);
900 case TargetOpcode::G_BR:
901 return selectBranch(
I);
902 case TargetOpcode::G_BRCOND:
903 return selectBranchCond(
I);
905 case TargetOpcode::G_PHI:
906 return selectPhi(ResVReg, ResType,
I);
908 case TargetOpcode::G_FPTOSI:
909 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
910 case TargetOpcode::G_FPTOUI:
911 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
913 case TargetOpcode::G_FPTOSI_SAT:
914 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
915 case TargetOpcode::G_FPTOUI_SAT:
916 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
918 case TargetOpcode::G_SITOFP:
919 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
920 case TargetOpcode::G_UITOFP:
921 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
923 case TargetOpcode::G_CTPOP:
924 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitCount);
925 case TargetOpcode::G_SMIN:
926 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
927 case TargetOpcode::G_UMIN:
928 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
930 case TargetOpcode::G_SMAX:
931 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
932 case TargetOpcode::G_UMAX:
933 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
935 case TargetOpcode::G_SCMP:
936 return selectSUCmp(ResVReg, ResType,
I,
true);
937 case TargetOpcode::G_UCMP:
938 return selectSUCmp(ResVReg, ResType,
I,
false);
939 case TargetOpcode::G_LROUND:
940 case TargetOpcode::G_LLROUND: {
942 MRI->createVirtualRegister(
MRI->getRegClass(ResVReg),
"lround");
943 MRI->setRegClass(regForLround, &SPIRV::iIDRegClass);
945 regForLround, *(
I.getParent()->getParent()));
947 I, CL::round, GL::Round);
949 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
955 case TargetOpcode::G_STRICT_FMA:
956 case TargetOpcode::G_FMA: {
959 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
962 .
addUse(
I.getOperand(1).getReg())
963 .
addUse(
I.getOperand(2).getReg())
964 .
addUse(
I.getOperand(3).getReg())
968 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
971 case TargetOpcode::G_STRICT_FLDEXP:
972 return selectExtInst(ResVReg, ResType,
I, CL::ldexp);
974 case TargetOpcode::G_FPOW:
975 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
976 case TargetOpcode::G_FPOWI:
977 return selectExtInst(ResVReg, ResType,
I, CL::pown);
979 case TargetOpcode::G_FEXP:
980 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
981 case TargetOpcode::G_FEXP2:
982 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
983 case TargetOpcode::G_FMODF:
984 return selectModf(ResVReg, ResType,
I);
986 case TargetOpcode::G_FLOG:
987 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
988 case TargetOpcode::G_FLOG2:
989 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
990 case TargetOpcode::G_FLOG10:
991 return selectLog10(ResVReg, ResType,
I);
993 case TargetOpcode::G_FABS:
994 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
995 case TargetOpcode::G_ABS:
996 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
998 case TargetOpcode::G_FMINNUM:
999 case TargetOpcode::G_FMINIMUM:
1000 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1001 case TargetOpcode::G_FMAXNUM:
1002 case TargetOpcode::G_FMAXIMUM:
1003 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1005 case TargetOpcode::G_FCOPYSIGN:
1006 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1008 case TargetOpcode::G_FCEIL:
1009 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1010 case TargetOpcode::G_FFLOOR:
1011 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1013 case TargetOpcode::G_FCOS:
1014 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1015 case TargetOpcode::G_FSIN:
1016 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1017 case TargetOpcode::G_FTAN:
1018 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1019 case TargetOpcode::G_FACOS:
1020 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1021 case TargetOpcode::G_FASIN:
1022 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1023 case TargetOpcode::G_FATAN:
1024 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1025 case TargetOpcode::G_FATAN2:
1026 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1027 case TargetOpcode::G_FCOSH:
1028 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1029 case TargetOpcode::G_FSINH:
1030 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1031 case TargetOpcode::G_FTANH:
1032 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1034 case TargetOpcode::G_STRICT_FSQRT:
1035 case TargetOpcode::G_FSQRT:
1036 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1038 case TargetOpcode::G_CTTZ:
1039 case TargetOpcode::G_CTTZ_ZERO_UNDEF:
1040 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1041 case TargetOpcode::G_CTLZ:
1042 case TargetOpcode::G_CTLZ_ZERO_UNDEF:
1043 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1045 case TargetOpcode::G_INTRINSIC_ROUND:
1046 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1047 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1048 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1049 case TargetOpcode::G_INTRINSIC_TRUNC:
1050 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1051 case TargetOpcode::G_FRINT:
1052 case TargetOpcode::G_FNEARBYINT:
1053 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1055 case TargetOpcode::G_SMULH:
1056 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1057 case TargetOpcode::G_UMULH:
1058 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1060 case TargetOpcode::G_SADDSAT:
1061 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1062 case TargetOpcode::G_UADDSAT:
1063 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1064 case TargetOpcode::G_SSUBSAT:
1065 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1066 case TargetOpcode::G_USUBSAT:
1067 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1069 case TargetOpcode::G_FFREXP:
1070 return selectFrexp(ResVReg, ResType,
I);
1072 case TargetOpcode::G_UADDO:
1073 return selectOverflowArith(ResVReg, ResType,
I,
1074 ResType->
getOpcode() == SPIRV::OpTypeVector
1075 ? SPIRV::OpIAddCarryV
1076 : SPIRV::OpIAddCarryS);
1077 case TargetOpcode::G_USUBO:
1078 return selectOverflowArith(ResVReg, ResType,
I,
1079 ResType->
getOpcode() == SPIRV::OpTypeVector
1080 ? SPIRV::OpISubBorrowV
1081 : SPIRV::OpISubBorrowS);
1082 case TargetOpcode::G_UMULO:
1083 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1084 case TargetOpcode::G_SMULO:
1085 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1087 case TargetOpcode::G_SEXT:
1088 return selectExt(ResVReg, ResType,
I,
true);
1089 case TargetOpcode::G_ANYEXT:
1090 case TargetOpcode::G_ZEXT:
1091 return selectExt(ResVReg, ResType,
I,
false);
1092 case TargetOpcode::G_TRUNC:
1093 return selectTrunc(ResVReg, ResType,
I);
1094 case TargetOpcode::G_FPTRUNC:
1095 case TargetOpcode::G_FPEXT:
1096 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1098 case TargetOpcode::G_PTRTOINT:
1099 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1100 case TargetOpcode::G_INTTOPTR:
1101 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1102 case TargetOpcode::G_BITCAST:
1103 return selectBitcast(ResVReg, ResType,
I);
1104 case TargetOpcode::G_ADDRSPACE_CAST:
1105 return selectAddrSpaceCast(ResVReg, ResType,
I);
1106 case TargetOpcode::G_PTR_ADD: {
1108 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1112 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1113 (*II).getOpcode() == TargetOpcode::COPY ||
1114 (*II).getOpcode() == SPIRV::OpVariable) &&
1117 bool IsGVInit =
false;
1119 UseIt =
MRI->use_instr_begin(
I.getOperand(0).getReg()),
1120 UseEnd =
MRI->use_instr_end();
1121 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1122 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1123 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1124 (*UseIt).getOpcode() == SPIRV::OpVariable) {
1134 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1137 Register NewVReg =
MRI->createGenericVirtualRegister(
MRI->getType(GV));
1138 MRI->setRegClass(NewVReg,
MRI->getRegClass(GV));
1147 "incompatible result and operand types in a bitcast");
1149 MachineInstrBuilder MIB =
1150 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1157 ? SPIRV::OpInBoundsAccessChain
1158 : SPIRV::OpInBoundsPtrAccessChain))
1162 .
addUse(
I.getOperand(2).getReg())
1165 return BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1169 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1171 .
addUse(
I.getOperand(2).getReg())
1179 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1182 .
addImm(
static_cast<uint32_t
>(
1183 SPIRV::Opcode::InBoundsPtrAccessChain))
1186 .
addUse(
I.getOperand(2).getReg());
1190 case TargetOpcode::G_ATOMICRMW_OR:
1191 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1192 case TargetOpcode::G_ATOMICRMW_ADD:
1193 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1194 case TargetOpcode::G_ATOMICRMW_AND:
1195 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1196 case TargetOpcode::G_ATOMICRMW_MAX:
1197 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1198 case TargetOpcode::G_ATOMICRMW_MIN:
1199 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1200 case TargetOpcode::G_ATOMICRMW_SUB:
1201 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1202 case TargetOpcode::G_ATOMICRMW_XOR:
1203 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1204 case TargetOpcode::G_ATOMICRMW_UMAX:
1205 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1206 case TargetOpcode::G_ATOMICRMW_UMIN:
1207 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1208 case TargetOpcode::G_ATOMICRMW_XCHG:
1209 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1210 case TargetOpcode::G_ATOMIC_CMPXCHG:
1211 return selectAtomicCmpXchg(ResVReg, ResType,
I);
1213 case TargetOpcode::G_ATOMICRMW_FADD:
1214 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1215 case TargetOpcode::G_ATOMICRMW_FSUB:
1217 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1218 ResType->
getOpcode() == SPIRV::OpTypeVector
1220 : SPIRV::OpFNegate);
1221 case TargetOpcode::G_ATOMICRMW_FMIN:
1222 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1223 case TargetOpcode::G_ATOMICRMW_FMAX:
1224 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1226 case TargetOpcode::G_FENCE:
1227 return selectFence(
I);
1229 case TargetOpcode::G_STACKSAVE:
1230 return selectStackSave(ResVReg, ResType,
I);
1231 case TargetOpcode::G_STACKRESTORE:
1232 return selectStackRestore(
I);
1234 case TargetOpcode::G_UNMERGE_VALUES:
1240 case TargetOpcode::G_TRAP:
1241 case TargetOpcode::G_UBSANTRAP:
1242 case TargetOpcode::DBG_LABEL:
1244 case TargetOpcode::G_DEBUGTRAP:
1245 return selectDebugTrap(ResVReg, ResType,
I);
1252bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1254 MachineInstr &
I)
const {
1255 unsigned Opcode = SPIRV::OpNop;
1257 return BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
1261bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1264 GL::GLSLExtInst GLInst)
const {
1266 SPIRV::InstructionSet::InstructionSet::GLSL_std_450)) {
1267 std::string DiagMsg;
1268 raw_string_ostream OS(DiagMsg);
1269 I.print(OS,
true,
false,
false,
false);
1270 DiagMsg +=
" is only supported with the GLSL extended instruction set.\n";
1273 return selectExtInst(ResVReg, ResType,
I,
1274 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}});
1277bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1280 CL::OpenCLExtInst CLInst)
const {
1281 return selectExtInst(ResVReg, ResType,
I,
1282 {{SPIRV::InstructionSet::OpenCL_std, CLInst}});
1285bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1288 CL::OpenCLExtInst CLInst,
1289 GL::GLSLExtInst GLInst)
const {
1290 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1291 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1292 return selectExtInst(ResVReg, ResType,
I, ExtInsts);
1295bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1300 for (
const auto &Ex : Insts) {
1301 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1302 uint32_t Opcode = Ex.second;
1305 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1308 .
addImm(
static_cast<uint32_t
>(Set))
1311 const unsigned NumOps =
I.getNumOperands();
1314 I.getOperand(Index).getType() ==
1315 MachineOperand::MachineOperandType::MO_IntrinsicID)
1318 MIB.
add(
I.getOperand(Index));
1324bool SPIRVInstructionSelector::selectExtInstForLRound(
1326 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst)
const {
1327 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1328 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1329 return selectExtInstForLRound(ResVReg, ResType,
I, ExtInsts);
1332bool SPIRVInstructionSelector::selectExtInstForLRound(
1335 for (
const auto &Ex : Insts) {
1336 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1337 uint32_t Opcode = Ex.second;
1340 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1343 .
addImm(
static_cast<uint32_t
>(Set))
1345 const unsigned NumOps =
I.getNumOperands();
1348 I.getOperand(Index).getType() ==
1349 MachineOperand::MachineOperandType::MO_IntrinsicID)
1352 MIB.
add(
I.getOperand(Index));
1360bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1362 MachineInstr &
I)
const {
1363 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1364 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1365 for (
const auto &Ex : ExtInsts) {
1366 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1367 uint32_t Opcode = Ex.second;
1371 MachineIRBuilder MIRBuilder(
I);
1374 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1379 auto MIB =
BuildMI(*It->getParent(), It, It->getDebugLoc(),
1380 TII.get(SPIRV::OpVariable))
1383 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1387 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1390 .
addImm(
static_cast<uint32_t
>(Ex.first))
1392 .
add(
I.getOperand(2))
1397 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1398 .
addDef(
I.getOperand(1).getReg())
1407bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1410 std::vector<Register> Srcs,
1411 unsigned Opcode)
const {
1412 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1421bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1424 unsigned Opcode)
const {
1426 Register SrcReg =
I.getOperand(1).getReg();
1429 MRI->def_instr_begin(SrcReg);
1430 DefIt !=
MRI->def_instr_end(); DefIt = std::next(DefIt)) {
1431 unsigned DefOpCode = DefIt->getOpcode();
1432 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1435 if (
auto *VRD =
getVRegDef(*
MRI, DefIt->getOperand(1).getReg()))
1436 DefOpCode = VRD->getOpcode();
1438 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1439 DefOpCode == TargetOpcode::G_CONSTANT ||
1440 DefOpCode == SPIRV::OpVariable || DefOpCode == SPIRV::OpConstantI) {
1446 uint32_t SpecOpcode = 0;
1448 case SPIRV::OpConvertPtrToU:
1449 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1451 case SPIRV::OpConvertUToPtr:
1452 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1456 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1457 TII.get(SPIRV::OpSpecConstantOp))
1465 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1469bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1471 MachineInstr &
I)
const {
1472 Register OpReg =
I.getOperand(1).getReg();
1476 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1484 if (
MemOp->isVolatile())
1485 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1486 if (
MemOp->isNonTemporal())
1487 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1488 if (
MemOp->getAlign().value())
1489 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1495 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1496 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1500 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1502 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1506 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1510 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1512 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1524 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1526 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1528 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
1532bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
1534 MachineInstr &
I)
const {
1536 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1541 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
1542 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1544 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1546 MRI->createVirtualRegister(
MRI->getRegClass(HandleReg));
1548 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1552 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1553 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
1554 I.getDebugLoc(),
I);
1558 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1562 if (!
I.getNumMemOperands()) {
1563 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1565 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1568 MachineIRBuilder MIRBuilder(
I);
1574bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
1576 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
1577 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1582 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
1583 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1585 MRI->createVirtualRegister(
MRI->getRegClass(HandleReg));
1588 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1592 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1593 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1594 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1595 TII.get(SPIRV::OpImageWrite))
1601 if (sampledTypeIsSignedInteger(LLVMHandleType))
1604 return BMI.constrainAllUses(
TII,
TRI, RBI);
1609 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpStore))
1612 if (!
I.getNumMemOperands()) {
1613 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1615 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1618 MachineIRBuilder MIRBuilder(
I);
1624bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
1626 MachineInstr &
I)
const {
1627 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
1629 "llvm.stacksave intrinsic: this instruction requires the following "
1630 "SPIR-V extension: SPV_INTEL_variable_length_array",
1633 return BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
1639bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
1640 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
1642 "llvm.stackrestore intrinsic: this instruction requires the following "
1643 "SPIR-V extension: SPV_INTEL_variable_length_array",
1645 if (!
I.getOperand(0).isReg())
1648 return BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
1649 .
addUse(
I.getOperand(0).getReg())
1654SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
1655 MachineIRBuilder MIRBuilder(
I);
1656 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1663 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
1667 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
1668 Type *ArrTy = ArrayType::get(ValTy, Num);
1670 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
1673 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
1680 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
1683 .
addImm(SPIRV::StorageClass::UniformConstant)
1685 if (!MIBVar.constrainAllUses(
TII,
TRI, RBI))
1695bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
1698 Register DstReg =
I.getOperand(0).getReg();
1708 "Unable to determine pointee type size for OpCopyMemory");
1709 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
1710 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
1712 "OpCopyMemory requires the size to match the pointee type size");
1713 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemory))
1716 if (
I.getNumMemOperands()) {
1717 MachineIRBuilder MIRBuilder(
I);
1723bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &
I,
1726 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemorySized))
1727 .
addUse(
I.getOperand(0).getReg())
1729 .
addUse(
I.getOperand(2).getReg());
1730 if (
I.getNumMemOperands()) {
1731 MachineIRBuilder MIRBuilder(
I);
1737bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
1738 MachineInstr &
I)
const {
1739 Register SrcReg =
I.getOperand(1).getReg();
1741 if (
I.getOpcode() == TargetOpcode::G_MEMSET) {
1742 Register VarReg = getOrCreateMemSetGlobal(
I);
1745 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
1747 ValTy,
I, SPIRV::StorageClass::UniformConstant);
1749 Result &= selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast);
1752 Result &= selectCopyMemory(
I, SrcReg);
1754 Result &= selectCopyMemorySized(
I, SrcReg);
1756 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
1757 Result &= BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I);
1761bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
1765 unsigned NegateOpcode)
const {
1768 const MachineMemOperand *MemOp = *
I.memoperands_begin();
1771 auto ScopeConstant = buildI32Constant(Scope,
I);
1772 Register ScopeReg = ScopeConstant.first;
1773 Result &= ScopeConstant.second;
1775 Register Ptr =
I.getOperand(1).getReg();
1781 auto MemSemConstant = buildI32Constant(MemSem ,
I);
1782 Register MemSemReg = MemSemConstant.first;
1783 Result &= MemSemConstant.second;
1785 Register ValueReg =
I.getOperand(2).getReg();
1786 if (NegateOpcode != 0) {
1789 Result &= selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode);
1794 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
1804bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
1805 unsigned ArgI =
I.getNumOperands() - 1;
1807 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
1810 if (!SrcType || SrcType->
getOpcode() != SPIRV::OpTypeVector)
1812 "cannot select G_UNMERGE_VALUES with a non-vector argument");
1818 unsigned CurrentIndex = 0;
1819 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
1820 Register ResVReg =
I.getOperand(i).getReg();
1823 LLT ResLLT =
MRI->getType(ResVReg);
1829 ResType = ScalarType;
1835 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
1838 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
1844 for (
unsigned j = 0;
j < NumElements; ++
j) {
1845 MIB.
addImm(CurrentIndex + j);
1847 CurrentIndex += NumElements;
1851 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
1863bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
1866 auto MemSemConstant = buildI32Constant(MemSem,
I);
1867 Register MemSemReg = MemSemConstant.first;
1868 bool Result = MemSemConstant.second;
1870 uint32_t
Scope =
static_cast<uint32_t
>(
1872 auto ScopeConstant = buildI32Constant(Scope,
I);
1873 Register ScopeReg = ScopeConstant.first;
1874 Result &= ScopeConstant.second;
1877 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
1883bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
1886 unsigned Opcode)
const {
1887 Type *ResTy =
nullptr;
1891 "Not enough info to select the arithmetic with overflow instruction");
1894 "with overflow instruction");
1900 MachineIRBuilder MIRBuilder(
I);
1902 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
1903 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
1909 Register ZeroReg = buildZerosVal(ResType,
I);
1912 MRI->setRegClass(StructVReg, &SPIRV::IDRegClass);
1914 if (ResName.
size() > 0)
1919 BuildMI(BB, MIRBuilder.getInsertPt(),
I.getDebugLoc(),
TII.get(Opcode))
1922 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
1923 MIB.
addUse(
I.getOperand(i).getReg());
1928 MRI->setRegClass(HigherVReg, &SPIRV::iIDRegClass);
1929 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
1931 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
1932 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
1940 .
addDef(
I.getOperand(1).getReg())
1947bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
1949 MachineInstr &
I)
const {
1954 Register Ptr =
I.getOperand(2).getReg();
1957 const MachineMemOperand *MemOp = *
I.memoperands_begin();
1960 auto ScopeConstant = buildI32Constant(Scope,
I);
1961 ScopeReg = ScopeConstant.first;
1962 Result &= ScopeConstant.second;
1964 unsigned ScSem =
static_cast<uint32_t
>(
1967 unsigned MemSemEq =
static_cast<uint32_t
>(
getMemSemantics(AO)) | ScSem;
1968 auto MemSemEqConstant = buildI32Constant(MemSemEq,
I);
1969 MemSemEqReg = MemSemEqConstant.first;
1970 Result &= MemSemEqConstant.second;
1972 unsigned MemSemNeq =
static_cast<uint32_t
>(
getMemSemantics(FO)) | ScSem;
1973 if (MemSemEq == MemSemNeq)
1974 MemSemNeqReg = MemSemEqReg;
1976 auto MemSemNeqConstant = buildI32Constant(MemSemEq,
I);
1977 MemSemNeqReg = MemSemNeqConstant.first;
1978 Result &= MemSemNeqConstant.second;
1981 ScopeReg =
I.getOperand(5).getReg();
1982 MemSemEqReg =
I.getOperand(6).getReg();
1983 MemSemNeqReg =
I.getOperand(7).getReg();
1987 Register Val =
I.getOperand(4).getReg();
1992 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2019 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2030 case SPIRV::StorageClass::DeviceOnlyINTEL:
2031 case SPIRV::StorageClass::HostOnlyINTEL:
2040 bool IsGRef =
false;
2041 bool IsAllowedRefs =
2042 llvm::all_of(
MRI->use_instructions(ResVReg), [&IsGRef](
auto const &It) {
2043 unsigned Opcode = It.getOpcode();
2044 if (Opcode == SPIRV::OpConstantComposite ||
2045 Opcode == SPIRV::OpVariable ||
2046 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2047 return IsGRef = true;
2048 return Opcode == SPIRV::OpName;
2050 return IsAllowedRefs && IsGRef;
2053Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2054 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2056 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2060SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2062 uint32_t Opcode)
const {
2063 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2064 TII.get(SPIRV::OpSpecConstantOp))
2072SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2076 Register Tmp =
MRI->createVirtualRegister(&SPIRV::pIDRegClass);
2078 SPIRV::StorageClass::Generic),
2080 MachineFunction *MF =
I.getParent()->getParent();
2082 MachineInstrBuilder MIB = buildSpecConstantOp(
2084 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2094bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2096 MachineInstr &
I)
const {
2100 Register SrcPtr =
I.getOperand(1).getReg();
2104 if (SrcPtrTy->
getOpcode() != SPIRV::OpTypePointer ||
2105 ResType->
getOpcode() != SPIRV::OpTypePointer)
2106 return BuildCOPY(ResVReg, SrcPtr,
I);
2116 unsigned SpecOpcode =
2118 ?
static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric)
2121 ? static_cast<uint32_t>(
SPIRV::Opcode::GenericCastToPtr)
2128 return buildSpecConstantOp(
I, ResVReg, SrcPtr,
2129 getUcharPtrTypeReg(
I, DstSC), SpecOpcode)
2130 .constrainAllUses(
TII,
TRI, RBI);
2132 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2134 buildSpecConstantOp(
2136 getUcharPtrTypeReg(
I, DstSC),
2137 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2138 .constrainAllUses(
TII,
TRI, RBI);
2144 return BuildCOPY(ResVReg, SrcPtr,
I);
2146 if ((SrcSC == SPIRV::StorageClass::Function &&
2147 DstSC == SPIRV::StorageClass::Private) ||
2148 (DstSC == SPIRV::StorageClass::Function &&
2149 SrcSC == SPIRV::StorageClass::Private))
2150 return BuildCOPY(ResVReg, SrcPtr,
I);
2154 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2157 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2178 return selectUnOp(ResVReg, ResType,
I,
2179 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
2181 return selectUnOp(ResVReg, ResType,
I,
2182 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
2184 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2186 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2196 return SPIRV::OpFOrdEqual;
2198 return SPIRV::OpFOrdGreaterThanEqual;
2200 return SPIRV::OpFOrdGreaterThan;
2202 return SPIRV::OpFOrdLessThanEqual;
2204 return SPIRV::OpFOrdLessThan;
2206 return SPIRV::OpFOrdNotEqual;
2208 return SPIRV::OpOrdered;
2210 return SPIRV::OpFUnordEqual;
2212 return SPIRV::OpFUnordGreaterThanEqual;
2214 return SPIRV::OpFUnordGreaterThan;
2216 return SPIRV::OpFUnordLessThanEqual;
2218 return SPIRV::OpFUnordLessThan;
2220 return SPIRV::OpFUnordNotEqual;
2222 return SPIRV::OpUnordered;
2232 return SPIRV::OpIEqual;
2234 return SPIRV::OpINotEqual;
2236 return SPIRV::OpSGreaterThanEqual;
2238 return SPIRV::OpSGreaterThan;
2240 return SPIRV::OpSLessThanEqual;
2242 return SPIRV::OpSLessThan;
2244 return SPIRV::OpUGreaterThanEqual;
2246 return SPIRV::OpUGreaterThan;
2248 return SPIRV::OpULessThanEqual;
2250 return SPIRV::OpULessThan;
2259 return SPIRV::OpPtrEqual;
2261 return SPIRV::OpPtrNotEqual;
2272 return SPIRV::OpLogicalEqual;
2274 return SPIRV::OpLogicalNotEqual;
2308bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
2311 unsigned OpAnyOrAll)
const {
2312 assert(
I.getNumOperands() == 3);
2313 assert(
I.getOperand(2).isReg());
2315 Register InputRegister =
I.getOperand(2).getReg();
2322 bool IsVectorTy = InputType->
getOpcode() == SPIRV::OpTypeVector;
2323 if (IsBoolTy && !IsVectorTy) {
2324 assert(ResVReg ==
I.getOperand(0).getReg());
2325 return BuildCOPY(ResVReg, InputRegister,
I);
2329 unsigned SpirvNotEqualId =
2330 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
2337 IsBoolTy ? InputRegister
2346 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
2366bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
2368 MachineInstr &
I)
const {
2369 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
2372bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
2374 MachineInstr &
I)
const {
2375 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
2379bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
2381 MachineInstr &
I)
const {
2382 assert(
I.getNumOperands() == 4);
2383 assert(
I.getOperand(2).isReg());
2384 assert(
I.getOperand(3).isReg());
2391 "dot product requires a vector of at least 2 components");
2399 return BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpDot))
2402 .
addUse(
I.getOperand(2).getReg())
2403 .
addUse(
I.getOperand(3).getReg())
2407bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
2411 assert(
I.getNumOperands() == 4);
2412 assert(
I.getOperand(2).isReg());
2413 assert(
I.getOperand(3).isReg());
2416 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
2420 .
addUse(
I.getOperand(2).getReg())
2421 .
addUse(
I.getOperand(3).getReg())
2427bool SPIRVInstructionSelector::selectIntegerDotExpansion(
2429 assert(
I.getNumOperands() == 4);
2430 assert(
I.getOperand(2).isReg());
2431 assert(
I.getOperand(3).isReg());
2435 Register Vec0 =
I.getOperand(2).getReg();
2436 Register Vec1 =
I.getOperand(3).getReg();
2449 "dot product requires a vector of at least 2 components");
2463 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2486bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
2488 MachineInstr &
I)
const {
2490 return BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
2493 .
addUse(
I.getOperand(2).getReg())
2497bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
2499 MachineInstr &
I)
const {
2501 return BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
2504 .
addUse(
I.getOperand(2).getReg())
2508template <
bool Signed>
2509bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
2511 MachineInstr &
I)
const {
2512 assert(
I.getNumOperands() == 5);
2513 assert(
I.getOperand(2).isReg());
2514 assert(
I.getOperand(3).isReg());
2515 assert(
I.getOperand(4).isReg());
2518 Register Acc =
I.getOperand(2).getReg();
2522 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
2542template <
bool Signed>
2543bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
2545 assert(
I.getNumOperands() == 5);
2546 assert(
I.getOperand(2).isReg());
2547 assert(
I.getOperand(3).isReg());
2548 assert(
I.getOperand(4).isReg());
2553 Register Acc =
I.getOperand(2).getReg();
2559 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
2563 for (
unsigned i = 0; i < 4; i++) {
2565 Register AElt =
MRI->createVirtualRegister(&SPIRV::IDRegClass);
2576 Register BElt =
MRI->createVirtualRegister(&SPIRV::IDRegClass);
2596 Register MaskMul =
MRI->createVirtualRegister(&SPIRV::IDRegClass);
2608 i < 3 ?
MRI->createVirtualRegister(&SPIRV::IDRegClass) : ResVReg;
2624bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
2626 MachineInstr &
I)
const {
2627 assert(
I.getNumOperands() == 3);
2628 assert(
I.getOperand(2).isReg());
2630 Register VZero = buildZerosValF(ResType,
I);
2631 Register VOne = buildOnesValF(ResType,
I);
2633 return BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
2636 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
2638 .
addUse(
I.getOperand(2).getReg())
2644bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
2646 MachineInstr &
I)
const {
2647 assert(
I.getNumOperands() == 3);
2648 assert(
I.getOperand(2).isReg());
2650 Register InputRegister =
I.getOperand(2).getReg();
2652 auto &
DL =
I.getDebugLoc();
2662 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
2664 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
2666 ?
MRI->createVirtualRegister(&SPIRV::IDRegClass)
2673 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
2678 if (NeedsConversion) {
2679 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
2690bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
2693 unsigned Opcode)
const {
2697 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
2703 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
2704 BMI.addUse(
I.getOperand(J).getReg());
2710bool SPIRVInstructionSelector::selectWaveActiveCountBits(
2716 bool Result = selectWaveOpInst(BallotReg, BallotType,
I,
2717 SPIRV::OpGroupNonUniformBallot);
2721 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
2726 .
addImm(SPIRV::GroupOperation::Reduce)
2733bool SPIRVInstructionSelector::selectWavePrefixBitCount(
2736 assert(
I.getNumOperands() == 3);
2738 auto Op =
I.getOperand(2);
2750 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
2761 Register BallotVReg =
MRI->createVirtualRegister(&SPIRV::IDRegClass);
2772 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
2776 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
2783bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
2786 bool IsUnsigned)
const {
2787 assert(
I.getNumOperands() == 3);
2788 assert(
I.getOperand(2).isReg());
2790 Register InputRegister =
I.getOperand(2).getReg();
2799 auto IntegerOpcodeType =
2800 IsUnsigned ? SPIRV::OpGroupNonUniformUMax : SPIRV::OpGroupNonUniformSMax;
2801 auto Opcode = IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntegerOpcodeType;
2802 return BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
2807 .
addImm(SPIRV::GroupOperation::Reduce)
2808 .
addUse(
I.getOperand(2).getReg())
2812bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
2815 bool IsUnsigned)
const {
2816 assert(
I.getNumOperands() == 3);
2817 assert(
I.getOperand(2).isReg());
2819 Register InputRegister =
I.getOperand(2).getReg();
2828 auto IntegerOpcodeType =
2829 IsUnsigned ? SPIRV::OpGroupNonUniformUMin : SPIRV::OpGroupNonUniformSMin;
2830 auto Opcode = IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntegerOpcodeType;
2831 return BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
2836 .
addImm(SPIRV::GroupOperation::Reduce)
2837 .
addUse(
I.getOperand(2).getReg())
2841bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
2843 MachineInstr &
I)
const {
2844 assert(
I.getNumOperands() == 3);
2845 assert(
I.getOperand(2).isReg());
2847 Register InputRegister =
I.getOperand(2).getReg();
2857 IsFloatTy ? SPIRV::OpGroupNonUniformFAdd : SPIRV::OpGroupNonUniformIAdd;
2858 return BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
2863 .
addImm(SPIRV::GroupOperation::Reduce)
2864 .
addUse(
I.getOperand(2).getReg());
2867bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
2869 MachineInstr &
I)
const {
2871 return BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
2874 .
addUse(
I.getOperand(1).getReg())
2878bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
2880 MachineInstr &
I)
const {
2886 if (!
I.getOperand(0).isReg() || !
I.getOperand(1).isReg())
2888 Register OpReg =
I.getOperand(1).getReg();
2889 if (MachineInstr *Def =
MRI->getVRegDef(OpReg)) {
2890 if (
Def->getOpcode() == TargetOpcode::COPY)
2891 Def =
MRI->getVRegDef(
Def->getOperand(1).getReg());
2893 switch (
Def->getOpcode()) {
2894 case SPIRV::ASSIGN_TYPE:
2895 if (MachineInstr *AssignToDef =
2896 MRI->getVRegDef(
Def->getOperand(1).getReg())) {
2897 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
2898 Reg =
Def->getOperand(2).getReg();
2901 case SPIRV::OpUndef:
2902 Reg =
Def->getOperand(1).getReg();
2905 unsigned DestOpCode;
2907 DestOpCode = SPIRV::OpConstantNull;
2909 DestOpCode = TargetOpcode::COPY;
2912 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
2913 .
addDef(
I.getOperand(0).getReg())
2920bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
2922 MachineInstr &
I)
const {
2924 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
2926 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
2930 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
2935 for (
unsigned i =
I.getNumExplicitDefs();
2936 i <
I.getNumExplicitOperands() && IsConst; ++i)
2940 if (!IsConst &&
N < 2)
2942 "There must be at least two constituent operands in a vector");
2945 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2946 TII.get(IsConst ? SPIRV::OpConstantComposite
2947 : SPIRV::OpCompositeConstruct))
2950 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
2951 MIB.
addUse(
I.getOperand(i).getReg());
2955bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
2957 MachineInstr &
I)
const {
2959 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
2961 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
2967 if (!
I.getOperand(
OpIdx).isReg())
2974 if (!IsConst &&
N < 2)
2976 "There must be at least two constituent operands in a vector");
2979 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2980 TII.get(IsConst ? SPIRV::OpConstantComposite
2981 : SPIRV::OpCompositeConstruct))
2984 for (
unsigned i = 0; i <
N; ++i)
2989bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
2991 MachineInstr &
I)
const {
2996 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
2998 Opcode = SPIRV::OpDemoteToHelperInvocation;
3000 Opcode = SPIRV::OpKill;
3002 if (MachineInstr *NextI =
I.getNextNode()) {
3004 NextI->removeFromParent();
3009 return BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
3013bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
3016 MachineInstr &
I)
const {
3017 Register Cmp0 =
I.getOperand(2).getReg();
3018 Register Cmp1 =
I.getOperand(3).getReg();
3021 "CMP operands should have the same type");
3022 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
3031bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
3033 MachineInstr &
I)
const {
3034 auto Pred =
I.getOperand(1).getPredicate();
3037 Register CmpOperand =
I.getOperand(2).getReg();
3044 return selectCmp(ResVReg, ResType, CmpOpc,
I);
3047std::pair<Register, bool>
3048SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
3054 auto ConstInt = ConstantInt::get(LLVMTy, Val);
3062 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
3065 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
3068 .
addImm(APInt(32, Val).getZExtValue());
3070 GR.
add(ConstInt,
MI);
3075bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
3077 MachineInstr &
I)
const {
3079 return selectCmp(ResVReg, ResType, CmpOp,
I);
3083 MachineInstr &
I)
const {
3086 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3091bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
3097 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
3105 if (
Def->getOpcode() == SPIRV::OpConstantNull)
3108 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
3109 Def->getOpcode() == SPIRV::OpConstantI)
3118 MachineInstr *
Def =
MRI->getVRegDef(
Reg);
3122 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
3123 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
3125 Intrinsic::spv_const_composite)) {
3126 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
3127 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
3128 if (!IsZero(
Def->getOperand(i).getReg()))
3138 MachineInstr &
I)
const {
3142 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3148 MachineInstr &
I)
const {
3152 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3159 MachineInstr &
I)
const {
3163 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3168bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
3170 MachineInstr &
I)
const {
3171 Register SelectFirstArg =
I.getOperand(2).getReg();
3172 Register SelectSecondArg =
I.getOperand(3).getReg();
3181 SPIRV::OpTypeVector;
3188 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
3189 }
else if (IsPtrTy) {
3190 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
3192 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
3196 Opcode = IsScalarBool ? SPIRV::OpSelectSFSCond : SPIRV::OpSelectVFVCond;
3197 }
else if (IsPtrTy) {
3198 Opcode = IsScalarBool ? SPIRV::OpSelectSPSCond : SPIRV::OpSelectVPVCond;
3200 Opcode = IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
3203 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
3206 .
addUse(
I.getOperand(1).getReg())
3212bool SPIRVInstructionSelector::selectSelectDefaultArgs(
Register ResVReg,
3215 bool IsSigned)
const {
3217 Register ZeroReg = buildZerosVal(ResType,
I);
3218 Register OneReg = buildOnesVal(IsSigned, ResType,
I);
3222 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
3223 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
3226 .
addUse(
I.getOperand(1).getReg())
3232bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
3234 MachineInstr &
I,
bool IsSigned,
3235 unsigned Opcode)
const {
3236 Register SrcReg =
I.getOperand(1).getReg();
3242 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
3247 selectSelectDefaultArgs(SrcReg, TmpType,
I,
false);
3249 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
3252bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
3254 MachineInstr &
I,
bool IsSigned)
const {
3255 Register SrcReg =
I.getOperand(1).getReg();
3257 return selectSelectDefaultArgs(ResVReg, ResType,
I, IsSigned);
3260 if (SrcType == ResType)
3261 return BuildCOPY(ResVReg, SrcReg,
I);
3263 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
3264 return selectUnOp(ResVReg, ResType,
I, Opcode);
3267bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
3270 bool IsSigned)
const {
3271 MachineIRBuilder MIRBuilder(
I);
3272 MachineRegisterInfo *
MRI = MIRBuilder.getMRI();
3287 TII.get(IsSigned ? SPIRV::OpSLessThanEqual
3288 : SPIRV::OpULessThanEqual))
3291 .
addUse(
I.getOperand(1).getReg())
3292 .
addUse(
I.getOperand(2).getReg())
3298 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
3301 .
addUse(
I.getOperand(1).getReg())
3302 .
addUse(
I.getOperand(2).getReg())
3310 unsigned SelectOpcode =
3311 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
3316 .
addUse(buildOnesVal(
true, ResType,
I))
3317 .
addUse(buildZerosVal(ResType,
I))
3324 .
addUse(buildOnesVal(
false, ResType,
I))
3328bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
3335 bool IsVectorTy = IntTy->
getOpcode() == SPIRV::OpTypeVector;
3336 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
3338 Register One = buildOnesVal(
false, IntTy,
I);
3354bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
3356 MachineInstr &
I)
const {
3357 Register IntReg =
I.getOperand(1).getReg();
3360 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
3361 if (ArgType == ResType)
3362 return BuildCOPY(ResVReg, IntReg,
I);
3364 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
3365 return selectUnOp(ResVReg, ResType,
I, Opcode);
3368bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
3370 MachineInstr &
I)
const {
3371 unsigned Opcode =
I.getOpcode();
3372 unsigned TpOpcode = ResType->
getOpcode();
3374 if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {
3375 assert(Opcode == TargetOpcode::G_CONSTANT &&
3376 I.getOperand(1).getCImm()->isZero());
3377 MachineBasicBlock &DepMBB =
I.getMF()->front();
3380 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
3387 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
3390bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
3392 MachineInstr &
I)
const {
3393 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
3399bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
3401 MachineInstr &
I)
const {
3403 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
3407 .
addUse(
I.getOperand(3).getReg())
3409 .
addUse(
I.getOperand(2).getReg());
3410 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
3415bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
3417 MachineInstr &
I)
const {
3418 Type *MaybeResTy =
nullptr;
3424 "Expected aggregate type for extractv instruction");
3426 SPIRV::AccessQualifier::ReadWrite,
false);
3430 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3433 .
addUse(
I.getOperand(2).getReg());
3434 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
3439bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
3441 MachineInstr &
I)
const {
3443 return selectInsertVal(ResVReg, ResType,
I);
3445 return BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
3448 .
addUse(
I.getOperand(2).getReg())
3449 .
addUse(
I.getOperand(3).getReg())
3450 .
addUse(
I.getOperand(4).getReg())
3454bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
3456 MachineInstr &
I)
const {
3458 return selectExtractVal(ResVReg, ResType,
I);
3460 return BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
3463 .
addUse(
I.getOperand(2).getReg())
3464 .
addUse(
I.getOperand(3).getReg())
3468bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
3470 MachineInstr &
I)
const {
3471 const bool IsGEPInBounds =
I.getOperand(2).getImm();
3477 ? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
3478 : SPIRV::OpAccessChain)
3479 : (IsGEPInBounds ?
SPIRV::OpInBoundsPtrAccessChain
3480 :
SPIRV::OpPtrAccessChain);
3482 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
3486 .
addUse(
I.getOperand(3).getReg());
3488 (Opcode == SPIRV::OpPtrAccessChain ||
3489 Opcode == SPIRV::OpInBoundsPtrAccessChain ||
3491 "Cannot translate GEP to OpAccessChain. First index must be 0.");
3494 const unsigned StartingIndex =
3495 (Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)
3498 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
3499 Res.addUse(
I.getOperand(i).getReg());
3500 return Res.constrainAllUses(
TII,
TRI, RBI);
3504bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
3507 unsigned Lim =
I.getNumExplicitOperands();
3508 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
3509 Register OpReg =
I.getOperand(i).getReg();
3510 MachineInstr *OpDefine =
MRI->getVRegDef(OpReg);
3512 SmallPtrSet<SPIRVType *, 4> Visited;
3513 if (!OpDefine || !OpType ||
isConstReg(
MRI, OpDefine, Visited) ||
3514 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
3515 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
3522 MachineFunction *MF =
I.getMF();
3534 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3535 TII.get(SPIRV::OpSpecConstantOp))
3538 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
3540 GR.
add(OpDefine, MIB);
3548bool SPIRVInstructionSelector::selectDerivativeInst(
3550 const unsigned DPdOpCode)
const {
3553 errorIfInstrOutsideShader(
I);
3558 Register SrcReg =
I.getOperand(2).getReg();
3563 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
3566 .
addUse(
I.getOperand(2).getReg());
3568 MachineIRBuilder MIRBuilder(
I);
3571 if (componentCount != 1)
3575 const TargetRegisterClass *RegClass = GR.
getRegClass(SrcType);
3576 Register ConvertToVReg =
MRI->createVirtualRegister(RegClass);
3577 Register DpdOpVReg =
MRI->createVirtualRegister(RegClass);
3580 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
3591 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
3599bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
3601 MachineInstr &
I)
const {
3605 case Intrinsic::spv_load:
3606 return selectLoad(ResVReg, ResType,
I);
3607 case Intrinsic::spv_store:
3608 return selectStore(
I);
3609 case Intrinsic::spv_extractv:
3610 return selectExtractVal(ResVReg, ResType,
I);
3611 case Intrinsic::spv_insertv:
3612 return selectInsertVal(ResVReg, ResType,
I);
3613 case Intrinsic::spv_extractelt:
3614 return selectExtractElt(ResVReg, ResType,
I);
3615 case Intrinsic::spv_insertelt:
3616 return selectInsertElt(ResVReg, ResType,
I);
3617 case Intrinsic::spv_gep:
3618 return selectGEP(ResVReg, ResType,
I);
3619 case Intrinsic::spv_bitcast: {
3620 Register OpReg =
I.getOperand(2).getReg();
3625 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
3627 case Intrinsic::spv_unref_global:
3628 case Intrinsic::spv_init_global: {
3629 MachineInstr *
MI =
MRI->getVRegDef(
I.getOperand(1).getReg());
3630 MachineInstr *Init =
I.getNumExplicitOperands() > 2
3631 ?
MRI->getVRegDef(
I.getOperand(2).getReg())
3634 Register GVarVReg =
MI->getOperand(0).getReg();
3635 bool Res = selectGlobalValue(GVarVReg, *
MI, Init);
3639 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
3641 MI->removeFromParent();
3645 case Intrinsic::spv_undef: {
3646 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
3651 case Intrinsic::spv_const_composite: {
3653 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
3659 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
3661 MachineIRBuilder MIR(
I);
3663 MIR, SPIRV::OpConstantComposite, 3,
3664 SPIRV::OpConstantCompositeContinuedINTEL, CompositeArgs, ResVReg,
3666 for (
auto *Instr : Instructions) {
3667 Instr->setDebugLoc(
I.getDebugLoc());
3673 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
3679 case Intrinsic::spv_assign_name: {
3680 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
3681 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
3682 for (
unsigned i =
I.getNumExplicitDefs() + 2;
3683 i <
I.getNumExplicitOperands(); ++i) {
3684 MIB.
addImm(
I.getOperand(i).getImm());
3688 case Intrinsic::spv_switch: {
3689 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
3690 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
3691 if (
I.getOperand(i).isReg())
3692 MIB.
addReg(
I.getOperand(i).getReg());
3693 else if (
I.getOperand(i).isCImm())
3694 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
3695 else if (
I.getOperand(i).isMBB())
3696 MIB.
addMBB(
I.getOperand(i).getMBB());
3702 case Intrinsic::spv_loop_merge: {
3703 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
3704 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
3705 if (
I.getOperand(i).isMBB())
3706 MIB.
addMBB(
I.getOperand(i).getMBB());
3712 case Intrinsic::spv_selection_merge: {
3714 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
3715 assert(
I.getOperand(1).isMBB() &&
3716 "operand 1 to spv_selection_merge must be a basic block");
3717 MIB.
addMBB(
I.getOperand(1).getMBB());
3718 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
3721 case Intrinsic::spv_cmpxchg:
3722 return selectAtomicCmpXchg(ResVReg, ResType,
I);
3723 case Intrinsic::spv_unreachable:
3724 return BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
3726 case Intrinsic::spv_alloca:
3727 return selectFrameIndex(ResVReg, ResType,
I);
3728 case Intrinsic::spv_alloca_array:
3729 return selectAllocaArray(ResVReg, ResType,
I);
3730 case Intrinsic::spv_assume:
3732 return BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
3733 .
addUse(
I.getOperand(1).getReg())
3736 case Intrinsic::spv_expect:
3738 return BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
3741 .
addUse(
I.getOperand(2).getReg())
3742 .
addUse(
I.getOperand(3).getReg())
3745 case Intrinsic::arithmetic_fence:
3748 TII.get(SPIRV::OpArithmeticFenceEXT))
3751 .
addUse(
I.getOperand(2).getReg())
3754 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
3756 case Intrinsic::spv_thread_id:
3762 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
3764 case Intrinsic::spv_thread_id_in_group:
3770 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
3772 case Intrinsic::spv_group_id:
3778 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
3780 case Intrinsic::spv_flattened_thread_id_in_group:
3787 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
3789 case Intrinsic::spv_workgroup_size:
3790 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
3792 case Intrinsic::spv_global_size:
3793 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
3795 case Intrinsic::spv_global_offset:
3796 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
3798 case Intrinsic::spv_num_workgroups:
3799 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
3801 case Intrinsic::spv_subgroup_size:
3802 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
3804 case Intrinsic::spv_num_subgroups:
3805 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
3807 case Intrinsic::spv_subgroup_id:
3808 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
3809 case Intrinsic::spv_subgroup_local_invocation_id:
3810 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
3811 ResVReg, ResType,
I);
3812 case Intrinsic::spv_subgroup_max_size:
3813 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
3815 case Intrinsic::spv_fdot:
3816 return selectFloatDot(ResVReg, ResType,
I);
3817 case Intrinsic::spv_udot:
3818 case Intrinsic::spv_sdot:
3819 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
3821 return selectIntegerDot(ResVReg, ResType,
I,
3822 IID == Intrinsic::spv_sdot);
3823 return selectIntegerDotExpansion(ResVReg, ResType,
I);
3824 case Intrinsic::spv_dot4add_i8packed:
3825 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
3827 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
3828 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
3829 case Intrinsic::spv_dot4add_u8packed:
3830 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
3832 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
3833 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
3834 case Intrinsic::spv_all:
3835 return selectAll(ResVReg, ResType,
I);
3836 case Intrinsic::spv_any:
3837 return selectAny(ResVReg, ResType,
I);
3838 case Intrinsic::spv_cross:
3839 return selectExtInst(ResVReg, ResType,
I, CL::cross, GL::Cross);
3840 case Intrinsic::spv_distance:
3841 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
3842 case Intrinsic::spv_lerp:
3843 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
3844 case Intrinsic::spv_length:
3845 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
3846 case Intrinsic::spv_degrees:
3847 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
3848 case Intrinsic::spv_faceforward:
3849 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
3850 case Intrinsic::spv_frac:
3851 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
3852 case Intrinsic::spv_isinf:
3853 return selectOpIsInf(ResVReg, ResType,
I);
3854 case Intrinsic::spv_isnan:
3855 return selectOpIsNan(ResVReg, ResType,
I);
3856 case Intrinsic::spv_normalize:
3857 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
3858 case Intrinsic::spv_refract:
3859 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
3860 case Intrinsic::spv_reflect:
3861 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
3862 case Intrinsic::spv_rsqrt:
3863 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
3864 case Intrinsic::spv_sign:
3865 return selectSign(ResVReg, ResType,
I);
3866 case Intrinsic::spv_smoothstep:
3867 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
3868 case Intrinsic::spv_firstbituhigh:
3869 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
3870 case Intrinsic::spv_firstbitshigh:
3871 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
3872 case Intrinsic::spv_firstbitlow:
3873 return selectFirstBitLow(ResVReg, ResType,
I);
3874 case Intrinsic::spv_group_memory_barrier_with_group_sync: {
3876 auto MemSemConstant =
3877 buildI32Constant(SPIRV::MemorySemantics::SequentiallyConsistent,
I);
3878 Register MemSemReg = MemSemConstant.first;
3879 Result &= MemSemConstant.second;
3880 auto ScopeConstant = buildI32Constant(SPIRV::Scope::Workgroup,
I);
3881 Register ScopeReg = ScopeConstant.first;
3882 Result &= ScopeConstant.second;
3885 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpControlBarrier))
3891 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
3892 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
3893 SPIRV::StorageClass::StorageClass ResSC =
3897 "Generic storage class");
3899 TII.get(SPIRV::OpGenericCastToPtrExplicit))
3906 case Intrinsic::spv_lifetime_start:
3907 case Intrinsic::spv_lifetime_end: {
3908 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
3909 : SPIRV::OpLifetimeStop;
3910 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
3911 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
3919 case Intrinsic::spv_saturate:
3920 return selectSaturate(ResVReg, ResType,
I);
3921 case Intrinsic::spv_nclamp:
3922 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
3923 case Intrinsic::spv_uclamp:
3924 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
3925 case Intrinsic::spv_sclamp:
3926 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
3927 case Intrinsic::spv_subgroup_prefix_bit_count:
3928 return selectWavePrefixBitCount(ResVReg, ResType,
I);
3929 case Intrinsic::spv_wave_active_countbits:
3930 return selectWaveActiveCountBits(ResVReg, ResType,
I);
3931 case Intrinsic::spv_wave_all:
3932 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
3933 case Intrinsic::spv_wave_any:
3934 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
3935 case Intrinsic::spv_subgroup_ballot:
3936 return selectWaveOpInst(ResVReg, ResType,
I,
3937 SPIRV::OpGroupNonUniformBallot);
3938 case Intrinsic::spv_wave_is_first_lane:
3939 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
3940 case Intrinsic::spv_wave_reduce_umax:
3941 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
3942 case Intrinsic::spv_wave_reduce_max:
3943 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
3944 case Intrinsic::spv_wave_reduce_umin:
3945 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
3946 case Intrinsic::spv_wave_reduce_min:
3947 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
3948 case Intrinsic::spv_wave_reduce_sum:
3949 return selectWaveReduceSum(ResVReg, ResType,
I);
3950 case Intrinsic::spv_wave_readlane:
3951 return selectWaveOpInst(ResVReg, ResType,
I,
3952 SPIRV::OpGroupNonUniformShuffle);
3953 case Intrinsic::spv_step:
3954 return selectExtInst(ResVReg, ResType,
I, CL::step, GL::Step);
3955 case Intrinsic::spv_radians:
3956 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
3960 case Intrinsic::instrprof_increment:
3961 case Intrinsic::instrprof_increment_step:
3962 case Intrinsic::instrprof_value_profile:
3965 case Intrinsic::spv_value_md:
3967 case Intrinsic::spv_resource_handlefrombinding: {
3968 return selectHandleFromBinding(ResVReg, ResType,
I);
3970 case Intrinsic::spv_resource_counterhandlefrombinding:
3971 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
3972 case Intrinsic::spv_resource_updatecounter:
3973 return selectUpdateCounter(ResVReg, ResType,
I);
3974 case Intrinsic::spv_resource_store_typedbuffer: {
3975 return selectImageWriteIntrinsic(
I);
3977 case Intrinsic::spv_resource_load_typedbuffer: {
3978 return selectReadImageIntrinsic(ResVReg, ResType,
I);
3980 case Intrinsic::spv_resource_sample:
3981 case Intrinsic::spv_resource_sample_clamp: {
3982 return selectSampleIntrinsic(ResVReg, ResType,
I);
3984 case Intrinsic::spv_resource_getpointer: {
3985 return selectResourceGetPointer(ResVReg, ResType,
I);
3987 case Intrinsic::spv_pushconstant_getpointer: {
3988 return selectPushConstantGetPointer(ResVReg, ResType,
I);
3990 case Intrinsic::spv_discard: {
3991 return selectDiscard(ResVReg, ResType,
I);
3993 case Intrinsic::spv_resource_nonuniformindex: {
3994 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
3996 case Intrinsic::spv_unpackhalf2x16: {
3997 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
3999 case Intrinsic::spv_packhalf2x16: {
4000 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
4002 case Intrinsic::spv_ddx:
4003 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
4004 case Intrinsic::spv_ddy:
4005 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
4006 case Intrinsic::spv_ddx_coarse:
4007 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
4008 case Intrinsic::spv_ddy_coarse:
4009 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
4010 case Intrinsic::spv_ddx_fine:
4011 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
4012 case Intrinsic::spv_ddy_fine:
4013 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
4014 case Intrinsic::spv_fwidth:
4015 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
4017 std::string DiagMsg;
4018 raw_string_ostream OS(DiagMsg);
4020 DiagMsg =
"Intrinsic selection not implemented: " + DiagMsg;
4027bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
4029 MachineInstr &
I)
const {
4032 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
4039bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
4042 assert(Intr.getIntrinsicID() ==
4043 Intrinsic::spv_resource_counterhandlefrombinding);
4046 Register MainHandleReg = Intr.getOperand(2).getReg();
4048 assert(MainHandleDef->getIntrinsicID() ==
4049 Intrinsic::spv_resource_handlefrombinding);
4053 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(),
MRI);
4054 Register IndexReg = MainHandleDef->getOperand(5).getReg();
4055 std::string CounterName =
4060 MachineIRBuilder MIRBuilder(
I);
4061 Register CounterVarReg = buildPointerToResource(
4063 Binding, ArraySize, IndexReg, CounterName, MIRBuilder);
4065 return BuildCOPY(ResVReg, CounterVarReg,
I);
4068bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
4070 MachineInstr &
I)
const {
4072 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
4074 Register CounterHandleReg = Intr.getOperand(2).getReg();
4075 Register IncrReg = Intr.getOperand(3).getReg();
4083 assert(CounterVarPointeeType &&
4084 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
4085 "Counter variable must be a struct");
4087 SPIRV::StorageClass::StorageBuffer &&
4088 "Counter variable must be in the storage buffer storage class");
4090 "Counter variable must have exactly 1 member in the struct");
4094 "Counter variable struct must have a single i32 member");
4098 MachineIRBuilder MIRBuilder(
I);
4100 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
4103 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
4105 auto Zero = buildI32Constant(0,
I);
4111 if (!
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4112 TII.get(SPIRV::OpAccessChain))
4115 .
addUse(CounterHandleReg)
4123 auto Scope = buildI32Constant(SPIRV::Scope::Device,
I);
4126 auto Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
4127 if (!Semantics.second)
4131 auto Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
4136 if (!
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
4147 return BuildCOPY(ResVReg, AtomicRes,
I);
4155 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
4162bool SPIRVInstructionSelector::selectReadImageIntrinsic(
4171 Register ImageReg =
I.getOperand(2).getReg();
4173 Register NewImageReg =
MRI->createVirtualRegister(
MRI->getRegClass(ImageReg));
4179 Register IdxReg =
I.getOperand(3).getReg();
4181 MachineInstr &Pos =
I;
4183 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
4187bool SPIRVInstructionSelector::selectSampleIntrinsic(
Register &ResVReg,
4189 MachineInstr &
I)
const {
4190 Register ImageReg =
I.getOperand(2).getReg();
4191 Register SamplerReg =
I.getOperand(3).getReg();
4192 Register CoordinateReg =
I.getOperand(4).getReg();
4193 std::optional<Register> OffsetReg;
4194 std::optional<Register> ClampReg;
4196 if (
I.getNumOperands() > 5)
4197 OffsetReg =
I.getOperand(5).getReg();
4198 if (
I.getNumOperands() > 6)
4199 ClampReg =
I.getOperand(6).getReg();
4204 Register NewImageReg =
MRI->createVirtualRegister(
MRI->getRegClass(ImageReg));
4212 MRI->createVirtualRegister(
MRI->getRegClass(SamplerReg));
4213 if (!loadHandleBeforePosition(
4218 MachineIRBuilder MIRBuilder(
I);
4224 bool Succeed =
BuildMI(*
I.getParent(),
I, Loc,
TII.get(SPIRV::OpSampledImage))
4234 BuildMI(*
I.getParent(),
I, Loc,
TII.get(SPIRV::OpImageSampleImplicitLod))
4240 uint32_t ImageOperands = 0;
4241 if (OffsetReg && !isScalarOrVectorIntConstantZero(*OffsetReg)) {
4242 ImageOperands |= 0x8;
4246 ImageOperands |= 0x80;
4249 if (ImageOperands != 0) {
4250 MIB.
addImm(ImageOperands);
4251 if (ImageOperands & 0x8)
4253 if (ImageOperands & 0x80)
4260bool SPIRVInstructionSelector::generateImageReadOrFetch(
4265 "ImageReg is not an image type.");
4267 bool IsSignedInteger =
4272 bool IsFetch = (SampledOp.getImm() == 1);
4275 if (ResultSize == 4) {
4278 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
4284 if (IsSignedInteger)
4289 SPIRVType *ReadType = widenTypeToVec4(ResType, Pos);
4293 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
4298 if (IsSignedInteger)
4300 bool Succeed = BMI.constrainAllUses(
TII,
TRI, RBI);
4304 if (ResultSize == 1) {
4306 TII.get(SPIRV::OpCompositeExtract))
4313 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
4316bool SPIRVInstructionSelector::selectResourceGetPointer(
4318 Register ResourcePtr =
I.getOperand(2).getReg();
4320 if (RegType->
getOpcode() == SPIRV::OpTypeImage) {
4329 MachineIRBuilder MIRBuilder(
I);
4331 Register IndexReg =
I.getOperand(3).getReg();
4334 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4335 TII.get(SPIRV::OpAccessChain))
4344bool SPIRVInstructionSelector::selectPushConstantGetPointer(
4346 MRI->replaceRegWith(ResVReg,
I.getOperand(2).getReg());
4350bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
4352 Register ObjReg =
I.getOperand(2).getReg();
4353 if (!BuildCOPY(ResVReg, ObjReg,
I))
4363 decorateUsesAsNonUniform(ResVReg);
4367void SPIRVInstructionSelector::decorateUsesAsNonUniform(
4370 while (WorkList.
size() > 0) {
4374 bool IsDecorated =
false;
4375 for (MachineInstr &Use :
MRI->use_instructions(CurrentReg)) {
4376 if (
Use.getOpcode() == SPIRV::OpDecorate &&
4377 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
4383 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
4385 if (ResultReg == CurrentReg)
4393 SPIRV::Decoration::NonUniformEXT, {});
4398bool SPIRVInstructionSelector::extractSubvector(
4400 MachineInstr &InsertionPoint)
const {
4402 [[maybe_unused]] uint64_t InputSize =
4405 assert(InputSize > 1 &&
"The input must be a vector.");
4406 assert(ResultSize > 1 &&
"The result must be a vector.");
4407 assert(ResultSize < InputSize &&
4408 "Cannot extract more element than there are in the input.");
4411 const TargetRegisterClass *ScalarRegClass = GR.
getRegClass(ScalarType);
4412 for (uint64_t
I = 0;
I < ResultSize;
I++) {
4413 Register ComponentReg =
MRI->createVirtualRegister(ScalarRegClass);
4416 TII.get(SPIRV::OpCompositeExtract))
4427 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
4429 TII.get(SPIRV::OpCompositeConstruct))
4433 for (
Register ComponentReg : ComponentRegisters)
4434 MIB.
addUse(ComponentReg);
4438bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
4439 MachineInstr &
I)
const {
4446 Register ImageReg =
I.getOperand(1).getReg();
4448 Register NewImageReg =
MRI->createVirtualRegister(
MRI->getRegClass(ImageReg));
4454 Register CoordinateReg =
I.getOperand(2).getReg();
4455 Register DataReg =
I.getOperand(3).getReg();
4458 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4459 TII.get(SPIRV::OpImageWrite))
4466Register SPIRVInstructionSelector::buildPointerToResource(
4467 const SPIRVType *SpirvResType, SPIRV::StorageClass::StorageClass SC,
4468 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
4469 StringRef Name, MachineIRBuilder MIRBuilder)
const {
4471 if (ArraySize == 1) {
4475 "SpirvResType did not have an explicit layout.");
4480 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
4484 VarPointerType, Set,
Binding, Name, MIRBuilder);
4499bool SPIRVInstructionSelector::selectFirstBitSet16(
4501 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
4503 bool Result = selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
4507 selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
4510bool SPIRVInstructionSelector::selectFirstBitSet32(
4512 Register SrcReg,
unsigned BitSetOpcode)
const {
4513 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
4516 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
4522bool SPIRVInstructionSelector::selectFirstBitSet64Overflow(
4524 Register SrcReg,
unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
4531 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
4533 MachineIRBuilder MIRBuilder(
I);
4541 std::vector<Register> PartialRegs;
4544 unsigned CurrentComponent = 0;
4545 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
4551 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4552 TII.get(SPIRV::OpVectorShuffle))
4557 .
addImm(CurrentComponent)
4558 .
addImm(CurrentComponent + 1);
4566 if (!selectFirstBitSet64(SubVecBitSetReg, Vec2ResType,
I, BitSetResult,
4567 BitSetOpcode, SwapPrimarySide))
4570 PartialRegs.push_back(SubVecBitSetReg);
4574 if (CurrentComponent != ComponentCount) {
4580 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
4581 SPIRV::OpVectorExtractDynamic))
4587 if (!selectFirstBitSet64(FinalElemBitSetReg,
BaseType,
I, FinalElemReg,
4588 BitSetOpcode, SwapPrimarySide))
4591 PartialRegs.push_back(FinalElemBitSetReg);
4596 return selectOpWithSrcs(ResVReg, ResType,
I, std::move(PartialRegs),
4597 SPIRV::OpCompositeConstruct);
4600bool SPIRVInstructionSelector::selectFirstBitSet64(
4602 Register SrcReg,
unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
4615 if (ComponentCount > 2) {
4616 return selectFirstBitSet64Overflow(ResVReg, ResType,
I, SrcReg,
4617 BitSetOpcode, SwapPrimarySide);
4621 MachineIRBuilder MIRBuilder(
I);
4623 BaseType, 2 * ComponentCount, MIRBuilder,
false);
4627 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
4633 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
4640 bool IsScalarRes = ResType->
getOpcode() != SPIRV::OpTypeVector;
4643 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntZero},
4644 SPIRV::OpVectorExtractDynamic))
4646 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntOne},
4647 SPIRV::OpVectorExtractDynamic))
4651 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4652 TII.get(SPIRV::OpVectorShuffle))
4660 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
4667 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4668 TII.get(SPIRV::OpVectorShuffle))
4676 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
4697 SelectOp = SPIRV::OpSelectSISCond;
4698 AddOp = SPIRV::OpIAddS;
4706 SelectOp = SPIRV::OpSelectVIVCond;
4707 AddOp = SPIRV::OpIAddV;
4717 if (SwapPrimarySide) {
4718 PrimaryReg = LowReg;
4719 SecondaryReg = HighReg;
4720 PrimaryShiftReg = Reg0;
4721 SecondaryShiftReg = Reg32;
4726 if (!selectOpWithSrcs(BReg, BoolType,
I, {PrimaryReg, NegOneReg},
4732 if (!selectOpWithSrcs(TmpReg, ResType,
I, {
BReg, SecondaryReg, PrimaryReg},
4738 if (!selectOpWithSrcs(ValReg, ResType,
I,
4739 {
BReg, SecondaryShiftReg, PrimaryShiftReg}, SelectOp))
4742 return selectOpWithSrcs(ResVReg, ResType,
I, {ValReg, TmpReg}, AddOp);
4745bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
4748 bool IsSigned)
const {
4750 Register OpReg =
I.getOperand(2).getReg();
4753 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4754 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
4758 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
4760 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
4762 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
4766 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
4770bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
4772 MachineInstr &
I)
const {
4774 Register OpReg =
I.getOperand(2).getReg();
4779 unsigned ExtendOpcode = SPIRV::OpUConvert;
4780 unsigned BitSetOpcode = GL::FindILsb;
4784 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
4786 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
4788 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
4795bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
4797 MachineInstr &
I)
const {
4801 bool Res =
BuildMI(BB,
I,
I.getDebugLoc(),
4802 TII.get(SPIRV::OpVariableLengthArrayINTEL))
4805 .
addUse(
I.getOperand(2).getReg())
4808 unsigned Alignment =
I.getOperand(3).getImm();
4814bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
4816 MachineInstr &
I)
const {
4820 bool Res =
BuildMI(*It->getParent(), It, It->getDebugLoc(),
4821 TII.get(SPIRV::OpVariable))
4824 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
4827 unsigned Alignment =
I.getOperand(2).getImm();
4834bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
4839 const MachineInstr *PrevI =
I.getPrevNode();
4841 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
4842 return BuildMI(
MBB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBranchConditional))
4845 .
addMBB(
I.getOperand(0).getMBB())
4849 .
addMBB(
I.getOperand(0).getMBB())
4853bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
4864 const MachineInstr *NextI =
I.getNextNode();
4866 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
4872 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
4873 return BuildMI(
MBB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBranchConditional))
4874 .
addUse(
I.getOperand(0).getReg())
4875 .
addMBB(
I.getOperand(1).getMBB())
4880bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
4882 MachineInstr &
I)
const {
4883 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpPhi))
4886 const unsigned NumOps =
I.getNumOperands();
4887 for (
unsigned i = 1; i <
NumOps; i += 2) {
4888 MIB.
addUse(
I.getOperand(i + 0).getReg());
4889 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
4897bool SPIRVInstructionSelector::selectGlobalValue(
4898 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
4900 MachineIRBuilder MIRBuilder(
I);
4901 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
4904 std::string GlobalIdent;
4906 unsigned &
ID = UnnamedGlobalIDs[GV];
4908 ID = UnnamedGlobalIDs.size();
4909 GlobalIdent =
"__unnamed_" + Twine(
ID).str();
4936 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
4943 MachineRegisterInfo *
MRI = MIRBuilder.
getMRI();
4946 MRI->setRegClass(FuncVReg, &SPIRV::pIDRegClass);
4947 MachineInstrBuilder MIB1 =
4948 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4951 MachineInstrBuilder MIB2 =
4953 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
4957 GR.
add(ConstVal, MIB2);
4963 MachineInstrBuilder MIB3 =
4964 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4967 GR.
add(ConstVal, MIB3);
4970 assert(NewReg != ResVReg);
4971 return BuildCOPY(ResVReg, NewReg,
I);
4981 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
4990 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
4995 if (
GlobalVar->isExternallyInitialized() &&
4996 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
4997 constexpr unsigned ReadWriteINTEL = 3u;
5000 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
5006bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
5008 MachineInstr &
I)
const {
5010 return selectExtInst(ResVReg, ResType,
I, CL::log10);
5018 MachineIRBuilder MIRBuilder(
I);
5024 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
5027 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
5029 .
add(
I.getOperand(1))
5034 ResType->
getOpcode() == SPIRV::OpTypeFloat);
5037 ResType->
getOpcode() == SPIRV::OpTypeVector
5044 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
5045 ? SPIRV::OpVectorTimesScalar
5055bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
5057 MachineInstr &
I)
const {
5073 MachineIRBuilder MIRBuilder(
I);
5076 ResType, MIRBuilder, SPIRV::StorageClass::Function);
5088 MachineBasicBlock &EntryBB =
I.getMF()->front();
5092 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
5095 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
5101 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
5104 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
5107 .
add(
I.getOperand(
I.getNumExplicitDefs()))
5111 Register IntegralPartReg =
I.getOperand(1).getReg();
5112 if (IntegralPartReg.
isValid()) {
5114 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
5123 assert(
false &&
"GLSL::Modf is deprecated.");
5134bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
5135 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
5136 const SPIRVType *ResType, MachineInstr &
I)
const {
5137 MachineIRBuilder MIRBuilder(
I);
5141 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
5153 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
5157 MachineRegisterInfo *
MRI = MIRBuilder.
getMRI();
5158 Register LoadedRegister =
MRI->createVirtualRegister(&SPIRV::iIDRegClass);
5164 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
5171 assert(
I.getOperand(2).isReg());
5172 const uint32_t ThreadId =
foldImm(
I.getOperand(2),
MRI);
5176 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
5186bool SPIRVInstructionSelector::loadBuiltinInputID(
5187 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
5188 const SPIRVType *ResType, MachineInstr &
I)
const {
5189 MachineIRBuilder MIRBuilder(
I);
5191 ResType, MIRBuilder, SPIRV::StorageClass::Input);
5206 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
5210 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
5219 MachineInstr &
I)
const {
5220 MachineIRBuilder MIRBuilder(
I);
5221 if (
Type->getOpcode() != SPIRV::OpTypeVector)
5225 if (VectorSize == 4)
5233bool SPIRVInstructionSelector::loadHandleBeforePosition(
5235 MachineInstr &Pos)
const {
5238 Intrinsic::spv_resource_handlefrombinding);
5246 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
5247 MachineIRBuilder MIRBuilder(HandleDef);
5249 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
5251 if (IsStructuredBuffer) {
5256 Register VarReg = buildPointerToResource(VarType, SC, Set,
Binding, ArraySize,
5257 IndexReg, Name, MIRBuilder);
5261 uint32_t LoadOpcode =
5262 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
5265 TII.get(LoadOpcode))
5272void SPIRVInstructionSelector::errorIfInstrOutsideShader(
5273 MachineInstr &
I)
const {
5275 std::string DiagMsg;
5276 raw_string_ostream OS(DiagMsg);
5277 I.print(OS,
true,
false,
false,
false);
5278 DiagMsg +=
" is only supported in shaders.\n";
5284InstructionSelector *
5288 return new SPIRVInstructionSelector(TM, Subtarget, RBI);
unsigned const MachineRegisterInfo * MRI
MachineInstrBuilder & UseMI
#define GET_GLOBALISEL_PREDICATES_INIT
#define GET_GLOBALISEL_TEMPORARIES_INIT
assert(UImm &&(UImm !=~static_cast< T >(0)) &&"Invalid immediate!")
This file declares a class to represent arbitrary precision floating point values and provide a varie...
static bool selectUnmergeValues(MachineInstrBuilder &MIB, const ARMBaseInstrInfo &TII, MachineRegisterInfo &MRI, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI)
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
static GCRegistry::Add< CoreCLRGC > E("coreclr", "CoreCLR-compatible GC")
DXIL Resource Implicit Binding
Declares convenience wrapper classes for interpreting MachineInstr instances as specific generic oper...
const HexagonInstrInfo * TII
LLVMTypeRef LLVMIntType(unsigned NumBits)
const size_t AbstractManglingParser< Derived, Alloc >::NumOps
Register const TargetRegisterInfo * TRI
Promote Memory to Register
MachineInstr unsigned OpIdx
uint64_t IntrinsicInst * II
static StringRef getName(Value *V)
static unsigned getFCmpOpcode(CmpInst::Predicate Pred, unsigned Size)
static APFloat getOneFP(const Type *LLVMFloatTy)
static bool isUSMStorageClass(SPIRV::StorageClass::StorageClass SC)
static bool isASCastInGVar(MachineRegisterInfo *MRI, Register ResVReg)
static bool mayApplyGenericSelection(unsigned Opcode)
static APFloat getZeroFP(const Type *LLVMFloatTy)
std::vector< std::pair< SPIRV::InstructionSet::InstructionSet, uint32_t > > ExtInstList
static bool intrinsicHasSideEffects(Intrinsic::ID ID)
static unsigned getBoolCmpOpcode(unsigned PredNum)
static unsigned getICmpOpcode(unsigned PredNum)
static bool isOpcodeWithNoSideEffects(unsigned Opcode)
static void addMemoryOperands(MachineMemOperand *MemOp, MachineInstrBuilder &MIB, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry &GR)
static bool isConstReg(MachineRegisterInfo *MRI, MachineInstr *OpDef, SmallPtrSet< SPIRVType *, 4 > &Visited)
static unsigned getPtrCmpOpcode(unsigned Pred)
bool isDead(const MachineInstr &MI, const MachineRegisterInfo &MRI)
BaseType
A given derived pointer can have multiple base pointers through phi/selects.
static TableGen::Emitter::Opt Y("gen-skeleton-entry", EmitSkeleton, "Generate example skeleton entry")
static TableGen::Emitter::OptClass< SkeletonEmitter > X("gen-skeleton-class", "Generate example skeleton class")
static const fltSemantics & IEEEsingle()
static const fltSemantics & IEEEdouble()
static const fltSemantics & IEEEhalf()
static APFloat getOne(const fltSemantics &Sem, bool Negative=false)
Factory for Positive and Negative One.
static APFloat getZero(const fltSemantics &Sem, bool Negative=false)
Factory for Positive and Negative Zero.
static APInt getAllOnes(unsigned numBits)
Return an APInt of a specified width with all bits set.
uint64_t getZExtValue() const
Get zero extended value.
BlockFrequencyInfo pass uses BlockFrequencyInfoImpl implementation to estimate IR basic block frequen...
Predicate
This enumeration lists the possible predicates for CmpInst subclasses.
@ FCMP_OEQ
0 0 0 1 True if ordered and equal
@ ICMP_SLT
signed less than
@ ICMP_SLE
signed less or equal
@ FCMP_OLT
0 1 0 0 True if ordered and less than
@ FCMP_ULE
1 1 0 1 True if unordered, less than, or equal
@ FCMP_OGT
0 0 1 0 True if ordered and greater than
@ FCMP_OGE
0 0 1 1 True if ordered and greater than or equal
@ ICMP_UGE
unsigned greater or equal
@ ICMP_UGT
unsigned greater than
@ ICMP_SGT
signed greater than
@ FCMP_ULT
1 1 0 0 True if unordered or less than
@ FCMP_ONE
0 1 1 0 True if ordered and operands are unequal
@ FCMP_UEQ
1 0 0 1 True if unordered or equal
@ ICMP_ULT
unsigned less than
@ FCMP_UGT
1 0 1 0 True if unordered or greater than
@ FCMP_OLE
0 1 0 1 True if ordered and less than or equal
@ FCMP_ORD
0 1 1 1 True if ordered (no nans)
@ ICMP_SGE
signed greater or equal
@ FCMP_UNE
1 1 1 0 True if unordered or not equal
@ ICMP_ULE
unsigned less or equal
@ FCMP_UGE
1 0 1 1 True if unordered, greater than, or equal
@ FCMP_UNO
1 0 0 0 True if unordered: isnan(X) | isnan(Y)
static LLVM_ABI Constant * getNullValue(Type *Ty)
Constructor to create a '0' constant of arbitrary type.
LLVMContext & getContext() const
getContext - Return a reference to the LLVMContext associated with this function.
Represents a call to an intrinsic.
Intrinsic::ID getIntrinsicID() const
unsigned getAddressSpace() const
Module * getParent()
Get the module that this global value is contained inside of...
@ InternalLinkage
Rename collisions when linking (static functions).
static LLVM_ABI IntegerType * get(LLVMContext &C, unsigned NumBits)
This static method is the primary way of constructing an IntegerType.
constexpr bool isScalar() const
static constexpr LLT scalar(unsigned SizeInBits)
Get a low-level scalar or aggregate "bag of bits".
constexpr bool isValid() const
constexpr uint16_t getNumElements() const
Returns the number of elements in a vector LLT.
constexpr bool isVector() const
static constexpr LLT pointer(unsigned AddressSpace, unsigned SizeInBits)
Get a low-level pointer in the given address space.
constexpr bool isPointer() const
static constexpr LLT fixed_vector(unsigned NumElements, unsigned ScalarSizeInBits)
Get a low-level fixed-width vector of some number of elements and element width.
int getNumber() const
MachineBasicBlocks are uniquely numbered at the function level, unless they're not in a MachineFuncti...
LLVM_ABI iterator getFirstNonPHI()
Returns a pointer to the first instruction in this block that is not a PHINode instruction.
const MachineFunction * getParent() const
Return the MachineFunction containing this basic block.
MachineInstrBundleIterator< MachineInstr > iterator
MachineRegisterInfo & getRegInfo()
getRegInfo - Return information about the registers currently in use.
Function & getFunction()
Return the LLVM function that this machine code represents.
Helper class to build MachineInstr.
MachineBasicBlock::iterator getInsertPt()
Current insertion point for new instructions.
MachineInstrBuilder buildInstr(unsigned Opcode)
Build and insert <empty> = Opcode <empty>.
MachineFunction & getMF()
Getter for the function we currently build.
MachineRegisterInfo * getMRI()
Getter for MRI.
const MachineInstrBuilder & addUse(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a virtual register use operand.
const MachineInstrBuilder & addReg(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a new virtual register operand.
const MachineInstrBuilder & addImm(int64_t Val) const
Add a new immediate operand.
const MachineInstrBuilder & add(const MachineOperand &MO) const
bool constrainAllUses(const TargetInstrInfo &TII, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI) const
const MachineInstrBuilder & addMBB(MachineBasicBlock *MBB, unsigned TargetFlags=0) const
const MachineInstrBuilder & addDef(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a virtual register definition operand.
const MachineInstrBuilder & setMIFlags(unsigned Flags) const
MachineInstr * getInstr() const
If conversion operators fail, use this method to get the MachineInstr explicitly.
Representation of each machine instruction.
unsigned getOpcode() const
Returns the opcode of this MachineInstr.
const MachineBasicBlock * getParent() const
unsigned getNumOperands() const
Retuns the total number of operands.
LLVM_ABI void setDesc(const MCInstrDesc &TID)
Replace the instruction descriptor (thus opcode) of the current instruction with a new one.
LLVM_ABI unsigned getNumExplicitDefs() const
Returns the number of non-implicit definitions.
LLVM_ABI const MachineFunction * getMF() const
Return the function that contains the basic block that this instruction belongs to.
const DebugLoc & getDebugLoc() const
Returns the debug location id of this MachineInstr.
LLVM_ABI void removeOperand(unsigned OpNo)
Erase an operand from an instruction, leaving it with one fewer operand than it started with.
const MachineOperand & getOperand(unsigned i) const
A description of a memory reference used in the backend.
@ MOVolatile
The memory access is volatile.
@ MONonTemporal
The memory access is non-temporal.
bool isReg() const
isReg - Tests if this is a MO_Register operand.
MachineBasicBlock * getMBB() const
Register getReg() const
getReg - Returns the register number.
MachineRegisterInfo - Keep track of information for virtual and physical registers,...
defusechain_instr_iterator< true, false, false, true > use_instr_iterator
use_instr_iterator/use_instr_begin/use_instr_end - Walk all uses of the specified register,...
defusechain_instr_iterator< false, true, false, true > def_instr_iterator
def_instr_iterator/def_instr_begin/def_instr_end - Walk all defs of the specified register,...
LLVM_ABI Register createVirtualRegister(const TargetRegisterClass *RegClass, StringRef Name="")
createVirtualRegister - Create and return a new virtual register in the function with the specified r...
LLVM_ABI void setType(Register VReg, LLT Ty)
Set the low-level type of VReg to Ty.
Analysis providing profile information.
Holds all the information related to register banks.
Wrapper class representing virtual and physical registers.
constexpr bool isValid() const
constexpr bool isPhysical() const
Return true if the specified register number is in the physical register namespace.
SPIRVType * getSPIRVTypeForVReg(Register VReg, const MachineFunction *MF=nullptr) const
Register getOrCreateConstInt(uint64_t Val, MachineInstr &I, SPIRVType *SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
SPIRVType * getResultType(Register VReg, MachineFunction *MF=nullptr)
SPIRVType * getOrCreateSPIRVBoolType(MachineIRBuilder &MIRBuilder, bool EmitIR)
MachineInstr * getOrAddMemAliasingINTELInst(MachineIRBuilder &MIRBuilder, const MDNode *AliasingListMD)
void assignSPIRVTypeToVReg(SPIRVType *Type, Register VReg, const MachineFunction &MF)
Register getOrCreateUndef(MachineInstr &I, SPIRVType *SpvType, const SPIRVInstrInfo &TII)
Register buildGlobalVariable(Register Reg, SPIRVType *BaseType, StringRef Name, const GlobalValue *GV, SPIRV::StorageClass::StorageClass Storage, const MachineInstr *Init, bool IsConst, const std::optional< SPIRV::LinkageType::LinkageType > &LinkageType, MachineIRBuilder &MIRBuilder, bool IsInstSelector)
SPIRVType * changePointerStorageClass(SPIRVType *PtrType, SPIRV::StorageClass::StorageClass SC, MachineInstr &I)
const Type * getTypeForSPIRVType(const SPIRVType *Ty) const
bool isBitcastCompatible(const SPIRVType *Type1, const SPIRVType *Type2) const
unsigned getScalarOrVectorComponentCount(Register VReg) const
SPIRVType * getOrCreateSPIRVFloatType(unsigned BitWidth, MachineInstr &I, const SPIRVInstrInfo &TII)
bool isScalarOrVectorSigned(const SPIRVType *Type) const
Register getOrCreateGlobalVariableWithBinding(const SPIRVType *VarType, uint32_t Set, uint32_t Binding, StringRef Name, MachineIRBuilder &MIRBuilder)
SPIRVType * getOrCreateSPIRVType(const Type *Type, MachineInstr &I, SPIRV::AccessQualifier::AccessQualifier AQ, bool EmitIR)
unsigned getPointerSize() const
SPIRVType * getOrCreateSPIRVPointerType(const Type *BaseType, MachineIRBuilder &MIRBuilder, SPIRV::StorageClass::StorageClass SC)
Register buildConstantFP(APFloat Val, MachineIRBuilder &MIRBuilder, SPIRVType *SpvType=nullptr)
SPIRVType * getPointeeType(SPIRVType *PtrType)
void invalidateMachineInstr(MachineInstr *MI)
Register getSPIRVTypeID(const SPIRVType *SpirvType) const
bool isScalarOfType(Register VReg, unsigned TypeOpcode) const
bool findValueAttrs(const MachineInstr *Key, Type *&Ty, StringRef &Name)
SPIRVType * getOrCreateOpTypeSampledImage(SPIRVType *ImageType, MachineIRBuilder &MIRBuilder)
void addGlobalObject(const Value *V, const MachineFunction *MF, Register R)
SPIRVType * getScalarOrVectorComponentType(Register VReg) const
void recordFunctionPointer(const MachineOperand *MO, const Function *F)
bool isAggregateType(SPIRVType *Type) const
const TargetRegisterClass * getRegClass(SPIRVType *SpvType) const
SPIRVType * getOrCreateSPIRVVectorType(SPIRVType *BaseType, unsigned NumElements, MachineIRBuilder &MIRBuilder, bool EmitIR)
bool isScalarOrVectorOfType(Register VReg, unsigned TypeOpcode) const
Register getOrCreateConstIntArray(uint64_t Val, size_t Num, MachineInstr &I, SPIRVType *SpvType, const SPIRVInstrInfo &TII)
MachineFunction * setCurrentFunc(MachineFunction &MF)
Register getOrCreateConstVector(uint64_t Val, MachineInstr &I, SPIRVType *SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
SPIRVType * getOrCreateSPIRVIntegerType(unsigned BitWidth, MachineIRBuilder &MIRBuilder)
Type * getDeducedGlobalValueType(const GlobalValue *Global)
LLT getRegType(SPIRVType *SpvType) const
SPIRV::StorageClass::StorageClass getPointerStorageClass(Register VReg) const
Register getOrCreateConstFP(APFloat Val, MachineInstr &I, SPIRVType *SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
Register getOrCreateConstNullPtr(MachineIRBuilder &MIRBuilder, SPIRVType *SpvType)
unsigned getScalarOrVectorBitWidth(const SPIRVType *Type) const
const SPIRVType * retrieveScalarOrVectorIntType(const SPIRVType *Type) const
bool erase(const MachineInstr *MI)
bool add(SPIRV::IRHandle Handle, const MachineInstr *MI)
Register find(SPIRV::IRHandle Handle, const MachineFunction *MF)
bool isPhysicalSPIRV() const
bool isAtLeastSPIRVVer(VersionTuple VerToCompareTo) const
bool canUseExtInstSet(SPIRV::InstructionSet::InstructionSet E) const
bool isLogicalSPIRV() const
bool canUseExtension(SPIRV::Extension::Extension E) const
std::pair< iterator, bool > insert(PtrType Ptr)
Inserts Ptr if and only if there is no element in the container equal to Ptr.
bool contains(ConstPtrType Ptr) const
SmallPtrSet - This class implements a set which is optimized for holding SmallSize or less elements.
reference emplace_back(ArgTypes &&... Args)
void push_back(const T &Elt)
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
StringRef - Represent a constant reference to a string, i.e.
constexpr size_t size() const
size - Get the string size.
static LLVM_ABI StructType * get(LLVMContext &Context, ArrayRef< Type * > Elements, bool isPacked=false)
This static method is the primary way to create a literal StructType.
The instances of the Type class are immutable: once they are created, they are never changed.
@ HalfTyID
16-bit floating point type
@ FloatTyID
32-bit floating point type
@ DoubleTyID
64-bit floating point type
Type * getScalarType() const
If this is a vector type, return the element type, otherwise return 'this'.
bool isStructTy() const
True if this is an instance of StructType.
bool isAggregateType() const
Return true if the type is an aggregate type.
TypeID getTypeID() const
Return the type id for the type.
Value * getOperand(unsigned i) const
LLVM_ABI StringRef getName() const
Return a constant reference to the value's name.
NodeTy * getNextNode()
Get the next node, or nullptr for the list tail.
#define llvm_unreachable(msg)
Marks that the current location is not supposed to be reachable.
constexpr char IsConst[]
Key for Kernel::Arg::Metadata::mIsConst.
unsigned ID
LLVM IR allows to use arbitrary numbers as calling convention identifiers.
Scope
Defines the scope in which this symbol should be visible: Default – Visible in the public interface o...
NodeAddr< DefNode * > Def
NodeAddr< InstrNode * > Instr
NodeAddr< UseNode * > Use
This is an optimization pass for GlobalISel generic memory operations.
void buildOpName(Register Target, const StringRef &Name, MachineIRBuilder &MIRBuilder)
bool all_of(R &&range, UnaryPredicate P)
Provide wrappers to std::all_of which take ranges instead of having to pass begin/end explicitly.
int64_t getIConstValSext(Register ConstReg, const MachineRegisterInfo *MRI)
MachineInstrBuilder BuildMI(MachineFunction &MF, const MIMetadata &MIMD, const MCInstrDesc &MCID)
Builder interface. Specify how to create the initial instruction itself.
bool isTypeFoldingSupported(unsigned Opcode)
decltype(auto) dyn_cast(const From &Val)
dyn_cast<X> - Return the argument parameter cast to the specified type.
FunctionAddr VTableAddr uintptr_t uintptr_t Int32Ty
void addNumImm(const APInt &Imm, MachineInstrBuilder &MIB)
LLVM_ABI void salvageDebugInfo(const MachineRegisterInfo &MRI, MachineInstr &MI)
Assuming the instruction MI is going to be deleted, attempt to salvage debug users of MI by writing t...
LLVM_ABI bool constrainSelectedInstRegOperands(MachineInstr &I, const TargetInstrInfo &TII, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI)
Mutate the newly-selected instruction I to constrain its (possibly generic) virtual register operands...
bool isPreISelGenericOpcode(unsigned Opcode)
Check whether the given Opcode is a generic opcode that is not supposed to appear after ISel.
unsigned getArrayComponentCount(const MachineRegisterInfo *MRI, const MachineInstr *ResType)
uint64_t getIConstVal(Register ConstReg, const MachineRegisterInfo *MRI)
SmallVector< MachineInstr *, 4 > createContinuedInstructions(MachineIRBuilder &MIRBuilder, unsigned Opcode, unsigned MinWC, unsigned ContinuedOpcode, ArrayRef< Register > Args, Register ReturnRegister, Register TypeID)
SPIRV::MemorySemantics::MemorySemantics getMemSemanticsForStorageClass(SPIRV::StorageClass::StorageClass SC)
constexpr unsigned storageClassToAddressSpace(SPIRV::StorageClass::StorageClass SC)
MachineBasicBlock::iterator getFirstValidInstructionInsertPoint(MachineBasicBlock &BB)
void buildOpDecorate(Register Reg, MachineIRBuilder &MIRBuilder, SPIRV::Decoration::Decoration Dec, const std::vector< uint32_t > &DecArgs, StringRef StrImm)
MachineBasicBlock::iterator getOpVariableMBBIt(MachineInstr &I)
Register createVirtualRegister(SPIRVType *SpvType, SPIRVGlobalRegistry *GR, MachineRegisterInfo *MRI, const MachineFunction &MF)
MachineInstr * getImm(const MachineOperand &MO, const MachineRegisterInfo *MRI)
Type * toTypedPointer(Type *Ty)
LLVM_ABI raw_ostream & dbgs()
dbgs() - This returns a reference to a raw_ostream for debugging messages.
LLVM_ABI void report_fatal_error(Error Err, bool gen_crash_diag=true)
const MachineInstr SPIRVType
constexpr bool isGenericCastablePtr(SPIRV::StorageClass::StorageClass SC)
class LLVM_GSL_OWNER SmallVector
Forward declaration of SmallVector so that calculateSmallVectorDefaultInlinedElements can reference s...
MachineInstr * passCopy(MachineInstr *Def, const MachineRegisterInfo *MRI)
bool isa(const From &Val)
isa<X> - Return true if the parameter to the template is an instance of one of the template type argu...
std::optional< SPIRV::LinkageType::LinkageType > getSpirvLinkageTypeFor(const SPIRVSubtarget &ST, const GlobalValue &GV)
LLVM_ABI raw_fd_ostream & errs()
This returns a reference to a raw_ostream for standard error.
SPIRV::StorageClass::StorageClass addressSpaceToStorageClass(unsigned AddrSpace, const SPIRVSubtarget &STI)
AtomicOrdering
Atomic ordering for LLVM's memory model.
SPIRV::Scope::Scope getMemScope(LLVMContext &Ctx, SyncScope::ID Id)
InstructionSelector * createSPIRVInstructionSelector(const SPIRVTargetMachine &TM, const SPIRVSubtarget &Subtarget, const RegisterBankInfo &RBI)
std::string getStringValueFromReg(Register Reg, MachineRegisterInfo &MRI)
int64_t foldImm(const MachineOperand &MO, const MachineRegisterInfo *MRI)
DWARFExpression::Operation Op
MachineInstr * getDefInstrMaybeConstant(Register &ConstReg, const MachineRegisterInfo *MRI)
constexpr unsigned BitWidth
decltype(auto) cast(const From &Val)
cast<X> - Return the argument parameter cast to the specified type.
bool hasInitializer(const GlobalVariable *GV)
void addStringImm(const StringRef &Str, MCInst &Inst)
MachineInstr * getVRegDef(MachineRegisterInfo &MRI, Register Reg)
SPIRV::MemorySemantics::MemorySemantics getMemSemantics(AtomicOrdering Ord)
std::string getLinkStringForBuiltIn(SPIRV::BuiltIn::BuiltIn BuiltInValue)
LLVM_ABI bool isTriviallyDead(const MachineInstr &MI, const MachineRegisterInfo &MRI)
Check whether an instruction MI is dead: it only defines dead virtual registers, and doesn't have oth...