36#include "llvm/IR/IntrinsicsSPIRV.h"
42#define DEBUG_TYPE "spirv-isel"
49 std::vector<std::pair<SPIRV::InstructionSet::InstructionSet, uint32_t>>;
54 std::optional<Register> Bias;
55 std::optional<Register>
Offset;
56 std::optional<Register> MinLod;
57 std::optional<Register> GradX;
58 std::optional<Register> GradY;
59 std::optional<Register> Lod;
60 std::optional<Register> Compare;
67 bool IsScalar =
false;
70llvm::SPIRV::SelectionControl::SelectionControl
71getSelectionOperandForImm(
int Imm) {
73 return SPIRV::SelectionControl::Flatten;
75 return SPIRV::SelectionControl::DontFlatten;
77 return SPIRV::SelectionControl::None;
81#define GET_GLOBALISEL_PREDICATE_BITSET
82#include "SPIRVGenGlobalISel.inc"
83#undef GET_GLOBALISEL_PREDICATE_BITSET
110#define GET_GLOBALISEL_PREDICATES_DECL
111#include "SPIRVGenGlobalISel.inc"
112#undef GET_GLOBALISEL_PREDICATES_DECL
114#define GET_GLOBALISEL_TEMPORARIES_DECL
115#include "SPIRVGenGlobalISel.inc"
116#undef GET_GLOBALISEL_TEMPORARIES_DECL
140 unsigned BitSetOpcode)
const;
144 unsigned BitSetOpcode)
const;
148 unsigned BitSetOpcode,
bool SwapPrimarySide)
const;
155 unsigned Opcode)
const;
158 unsigned Opcode)
const;
180 unsigned NewOpcode,
unsigned NegateOpcode = 0)
const;
189 Register castPtrToPtrToInt(
Register Ptr, SPIRV::StorageClass::StorageClass SC,
193 bool selectAtomicPtrValue(
209 unsigned OpType)
const;
276 unsigned Opcode)
const;
280 unsigned Opcode)
const;
284 unsigned Opcode)
const;
288 unsigned Opcode)
const;
290 template <
bool Signed>
293 template <
bool Signed>
300 template <
typename PickOpcodeFn>
303 PickOpcodeFn &&PickOpcode)
const;
320 template <
typename PickOpcodeFn>
323 PickOpcodeFn &&PickOpcode)
const;
341 bool IsSigned)
const;
343 bool IsSigned,
unsigned Opcode)
const;
345 bool IsSigned)
const;
351 bool IsSigned)
const;
392 GL::GLSLExtInst GLInst,
bool setMIFlags =
true,
393 bool useMISrc =
true,
395 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
396 CL::OpenCLExtInst CLInst,
bool setMIFlags =
true,
397 bool useMISrc =
true,
399 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
400 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
401 bool setMIFlags =
true,
bool useMISrc =
true,
403 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
404 const ExtInstList &ExtInsts,
bool setMIFlags =
true,
405 bool useMISrc =
true,
408 bool selectLog10(
Register ResVReg, SPIRVTypeInst ResType,
409 MachineInstr &
I)
const;
411 bool selectFpowi(
Register ResVReg, SPIRVTypeInst ResType,
412 MachineInstr &
I)
const;
414 bool selectSaturate(
Register ResVReg, SPIRVTypeInst ResType,
415 MachineInstr &
I)
const;
417 bool selectWaveOpInst(
Register ResVReg, SPIRVTypeInst ResType,
418 MachineInstr &
I,
unsigned Opcode)
const;
420 bool selectBarrierInst(MachineInstr &
I,
unsigned Scope,
unsigned MemSem,
421 bool WithGroupSync)
const;
423 bool selectWaveActiveCountBits(
Register ResVReg, SPIRVTypeInst ResType,
424 MachineInstr &
I)
const;
426 bool selectWaveActiveAllEqual(
Register ResVReg, SPIRVTypeInst ResType,
427 MachineInstr &
I)
const;
431 bool selectHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
432 MachineInstr &
I)
const;
434 bool selectCounterHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
435 MachineInstr &
I)
const;
437 bool selectReadImageIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
438 MachineInstr &
I)
const;
439 bool selectGetDimensionsIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
440 MachineInstr &
I)
const;
441 bool selectGetDimensionsLevelsIntrinsic(
Register &ResVReg,
442 SPIRVTypeInst ResType,
443 MachineInstr &
I)
const;
444 bool selectGetDimensionsMSIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
445 MachineInstr &
I)
const;
448 std::optional<Register> LodReg = std::nullopt)
const;
449 bool selectSampleBasicIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
450 MachineInstr &
I)
const;
451 bool selectCalculateLodIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
452 MachineInstr &
I)
const;
453 bool selectSampleBiasIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
454 MachineInstr &
I)
const;
455 bool selectSampleGradIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
456 MachineInstr &
I)
const;
457 bool selectSampleLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
458 MachineInstr &
I)
const;
459 bool selectLoadLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
460 MachineInstr &
I)
const;
461 bool selectSampleCmpIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
462 MachineInstr &
I)
const;
463 bool selectSampleCmpLevelZeroIntrinsic(
Register &ResVReg,
464 SPIRVTypeInst ResType,
465 MachineInstr &
I)
const;
466 bool selectGatherIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
467 MachineInstr &
I)
const;
468 bool selectImageWriteIntrinsic(MachineInstr &
I)
const;
469 bool selectResourceGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
470 MachineInstr &
I)
const;
471 bool selectPushConstantGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
472 MachineInstr &
I)
const;
473 bool selectResourceNonUniformIndex(
Register &ResVReg, SPIRVTypeInst ResType,
474 MachineInstr &
I)
const;
475 bool selectModf(
Register ResVReg, SPIRVTypeInst ResType,
476 MachineInstr &
I)
const;
477 bool selectUpdateCounter(
Register &ResVReg, SPIRVTypeInst ResType,
478 MachineInstr &
I)
const;
479 bool selectFrexp(
Register ResVReg, SPIRVTypeInst ResType,
480 MachineInstr &
I)
const;
482 bool selectLdexp(
Register ResVReg, SPIRVTypeInst ResType,
483 MachineInstr &
I)
const;
484 bool selectSincos(
Register ResVReg, SPIRVTypeInst ResType,
485 MachineInstr &
I)
const;
486 bool selectExp10(
Register ResVReg, SPIRVTypeInst ResType,
487 MachineInstr &
I)
const;
488 bool selectDerivativeInst(
Register ResVReg, SPIRVTypeInst ResType,
489 MachineInstr &
I,
const unsigned DPdOpCode)
const;
491 Register buildI32Constant(uint32_t Val, MachineInstr &
I,
492 SPIRVTypeInst ResType =
nullptr)
const;
493 Register buildI32ConstantInEntryBlock(uint32_t Val, MachineInstr &
I,
494 SPIRVTypeInst ResType =
nullptr)
const;
496 Register buildZerosVal(SPIRVTypeInst ResType, MachineInstr &
I)
const;
497 bool isScalarOrVectorIntConstantZero(
Register Reg)
const;
498 Register buildZerosValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
500 MachineInstr &
I)
const;
501 Register buildOnesValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
503 bool wrapIntoSpecConstantOp(MachineInstr &
I,
506 Register getUcharPtrTypeReg(MachineInstr &
I,
507 SPIRV::StorageClass::StorageClass SC)
const;
508 MachineInstrBuilder buildSpecConstantOp(MachineInstr &
I,
Register Dest,
510 uint32_t Opcode)
const;
511 MachineInstrBuilder buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
512 SPIRVTypeInst SrcPtrTy)
const;
513 Register buildPointerToResource(SPIRVTypeInst ResType,
514 SPIRV::StorageClass::StorageClass SC,
515 uint32_t Set, uint32_t
Binding,
516 uint32_t ArraySize,
Register IndexReg,
518 MachineIRBuilder MIRBuilder)
const;
519 SPIRVTypeInst widenTypeToVec4(SPIRVTypeInst
Type, MachineInstr &
I)
const;
520 bool extractSubvector(
Register &ResVReg, SPIRVTypeInst ResType,
521 Register &ReadReg, MachineInstr &InsertionPoint)
const;
522 bool generateImageReadOrFetch(
Register &ResVReg, SPIRVTypeInst ResType,
525 const ImageOperands *ImOps =
nullptr)
const;
526 bool generateSampleImage(
Register ResVReg, SPIRVTypeInst ResType,
528 Register CoordinateReg,
const ImageOperands &ImOps,
531 bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
532 Register ResVReg, SPIRVTypeInst ResType,
533 MachineInstr &
I)
const;
534 bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
535 Register ResVReg, SPIRVTypeInst ResType,
536 MachineInstr &
I)
const;
537 bool loadHandleBeforePosition(
Register &HandleReg, SPIRVTypeInst ResType,
538 GIntrinsic &HandleDef, MachineInstr &Pos)
const;
539 void decorateUsesAsNonUniform(
Register &NonUniformReg)
const;
540 bool errorIfInstrOutsideShader(MachineInstr &
I)
const;
542 std::optional<SplitParts> splitEvenOddLanes(
Register PopCountReg,
543 unsigned ComponentCount,
545 SPIRVTypeInst I32Type)
const;
548 handle64BitOverflow(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
549 Register SrcReg,
unsigned int Opcode,
550 std::function<
bool(
Register, SPIRVTypeInst,
551 MachineInstr &,
Register,
unsigned)>
555bool sampledTypeIsSignedInteger(
const llvm::Type *HandleType) {
557 if (
TET->getTargetExtName() ==
"spirv.Image") {
560 assert(
TET->getTargetExtName() ==
"spirv.SignedImage");
561 return TET->getTypeParameter(0)->isIntegerTy();
565#define GET_GLOBALISEL_IMPL
566#include "SPIRVGenGlobalISel.inc"
567#undef GET_GLOBALISEL_IMPL
573 TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
576#include
"SPIRVGenGlobalISel.inc"
579#include
"SPIRVGenGlobalISel.inc"
591 InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
595void SPIRVInstructionSelector::resetVRegsType(MachineFunction &MF) {
596 if (HasVRegsReset == &MF)
611 for (
const auto &
MBB : MF) {
612 for (
const auto &
MI :
MBB) {
615 if (
MI.getOpcode() != SPIRV::ASSIGN_TYPE)
619 LLT DstType = MRI.
getType(DstReg);
621 LLT SrcType = MRI.
getType(SrcReg);
622 if (DstType != SrcType)
627 if (DstRC != SrcRC && SrcRC)
639 while (!Stack.empty()) {
644 switch (
MI->getOpcode()) {
645 case TargetOpcode::G_INTRINSIC:
646 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
647 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS: {
650 if (IntrID != Intrinsic::spv_const_composite &&
651 IntrID != Intrinsic::spv_undef && IntrID != Intrinsic::spv_poison)
655 case TargetOpcode::G_BUILD_VECTOR:
656 case TargetOpcode::G_SPLAT_VECTOR:
658 i < OpDef->getNumOperands(); i++) {
663 Stack.push_back(OpNestedDef);
666 case TargetOpcode::G_CONSTANT:
667 case TargetOpcode::G_FCONSTANT:
668 case TargetOpcode::G_IMPLICIT_DEF:
669 case SPIRV::OpConstantTrue:
670 case SPIRV::OpConstantFalse:
671 case SPIRV::OpConstantI:
672 case SPIRV::OpConstantF:
673 case SPIRV::OpConstantComposite:
674 case SPIRV::OpConstantCompositeContinuedINTEL:
675 case SPIRV::OpConstantSampler:
676 case SPIRV::OpConstantNull:
678 case SPIRV::OpPoisonKHR:
679 case SPIRV::OpConstantFunctionPointerINTEL:
706 case Intrinsic::spv_all:
707 case Intrinsic::spv_alloca:
708 case Intrinsic::spv_any:
709 case Intrinsic::spv_bitcast:
710 case Intrinsic::spv_const_composite:
711 case Intrinsic::spv_degrees:
712 case Intrinsic::spv_distance:
713 case Intrinsic::spv_extractelt:
714 case Intrinsic::spv_extractv:
715 case Intrinsic::spv_faceforward:
716 case Intrinsic::spv_fdot:
717 case Intrinsic::spv_firstbitlow:
718 case Intrinsic::spv_firstbitshigh:
719 case Intrinsic::spv_firstbituhigh:
720 case Intrinsic::spv_frac:
721 case Intrinsic::spv_gep:
722 case Intrinsic::spv_global_offset:
723 case Intrinsic::spv_global_size:
724 case Intrinsic::spv_group_id:
725 case Intrinsic::spv_insertelt:
726 case Intrinsic::spv_insertv:
727 case Intrinsic::spv_isinf:
728 case Intrinsic::spv_isnan:
729 case Intrinsic::spv_isfinite:
730 case Intrinsic::spv_isnormal:
731 case Intrinsic::spv_lerp:
732 case Intrinsic::spv_length:
733 case Intrinsic::spv_normalize:
734 case Intrinsic::spv_num_subgroups:
735 case Intrinsic::spv_num_workgroups:
736 case Intrinsic::spv_ptrcast:
737 case Intrinsic::spv_radians:
738 case Intrinsic::spv_reflect:
739 case Intrinsic::spv_refract:
740 case Intrinsic::spv_resource_getbasepointer:
741 case Intrinsic::spv_resource_getpointer:
742 case Intrinsic::spv_resource_handlefrombinding:
743 case Intrinsic::spv_resource_handlefromimplicitbinding:
744 case Intrinsic::spv_resource_nonuniformindex:
745 case Intrinsic::spv_resource_sample:
746 case Intrinsic::spv_rsqrt:
747 case Intrinsic::spv_saturate:
748 case Intrinsic::spv_sdot:
749 case Intrinsic::spv_sign:
750 case Intrinsic::spv_smoothstep:
751 case Intrinsic::spv_step:
752 case Intrinsic::spv_subgroup_id:
753 case Intrinsic::spv_subgroup_local_invocation_id:
754 case Intrinsic::spv_subgroup_max_size:
755 case Intrinsic::spv_subgroup_size:
756 case Intrinsic::spv_thread_id:
757 case Intrinsic::spv_thread_id_in_group:
758 case Intrinsic::spv_udot:
759 case Intrinsic::spv_undef:
760 case Intrinsic::spv_value_md:
761 case Intrinsic::spv_workgroup_size:
773 case SPIRV::OpTypeVoid:
774 case SPIRV::OpTypeBool:
775 case SPIRV::OpTypeInt:
776 case SPIRV::OpTypeFloat:
777 case SPIRV::OpTypeVector:
778 case SPIRV::OpTypeMatrix:
779 case SPIRV::OpTypeImage:
780 case SPIRV::OpTypeSampler:
781 case SPIRV::OpTypeSampledImage:
782 case SPIRV::OpTypeArray:
783 case SPIRV::OpTypeRuntimeArray:
784 case SPIRV::OpTypeStruct:
785 case SPIRV::OpTypeOpaque:
786 case SPIRV::OpTypePointer:
787 case SPIRV::OpTypeFunction:
788 case SPIRV::OpTypeEvent:
789 case SPIRV::OpTypeDeviceEvent:
790 case SPIRV::OpTypeReserveId:
791 case SPIRV::OpTypeQueue:
792 case SPIRV::OpTypePipe:
793 case SPIRV::OpTypeForwardPointer:
794 case SPIRV::OpTypePipeStorage:
795 case SPIRV::OpTypeNamedBarrier:
796 case SPIRV::OpTypeAccelerationStructureNV:
797 case SPIRV::OpTypeCooperativeMatrixNV:
798 case SPIRV::OpTypeCooperativeMatrixKHR:
808 if (
MI.getNumDefs() == 0)
811 for (
const auto &MO :
MI.all_defs()) {
813 if (
Reg.isPhysical()) {
818 if (
UseMI.getOpcode() != SPIRV::OpName) {
825 if (
MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE ||
MI.isFakeUse() ||
826 MI.isLifetimeMarker()) {
829 <<
"Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
840 if (
MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
841 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
844 LLVM_DEBUG(
dbgs() <<
"Dead: Intrinsic with no real side effects.\n");
849 if (
MI.mayStore() ||
MI.isCall() ||
850 (
MI.mayLoad() &&
MI.hasOrderedMemoryRef()) ||
MI.isPosition() ||
851 MI.isDebugInstr() ||
MI.isTerminator() ||
MI.isJumpTableDebugInfo()) {
852 LLVM_DEBUG(
dbgs() <<
"Not dead: instruction has side effects.\n");
863 LLVM_DEBUG(
dbgs() <<
"Dead: known opcode with no side effects\n");
870void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &
MI)
const {
872 for (
const auto &MO :
MI.all_defs()) {
876 SmallVector<MachineInstr *, 4> UselessOpNames;
879 "There is still a use of the dead function.");
882 for (MachineInstr *OpNameMI : UselessOpNames) {
884 OpNameMI->eraseFromParent();
889void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &
MI)
const {
892 removeOpNamesForDeadMI(
MI);
893 MI.eraseFromParent();
896bool SPIRVInstructionSelector::select(MachineInstr &
I) {
897 resetVRegsType(*
I.getParent()->getParent());
899 assert(
I.getParent() &&
"Instruction should be in a basic block!");
900 assert(
I.getParent()->getParent() &&
"Instruction should be in a function!");
905 removeDeadInstruction(
I);
912 if (Opcode == SPIRV::ASSIGN_TYPE) {
913 Register DstReg =
I.getOperand(0).getReg();
914 Register SrcReg =
I.getOperand(1).getReg();
917 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
918 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
919 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
920 Register SelectDstReg =
Def->getOperand(0).getReg();
921 bool SuccessToSelectSelect [[maybe_unused]] = selectSelect(
923 assert(SuccessToSelectSelect);
925 Def->eraseFromParent();
932 bool Res = selectImpl(
I, *CoverageInfo);
934 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
935 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
939 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
951 }
else if (
I.getNumDefs() == 1) {
963 removeDeadInstruction(
I);
968 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
969 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
975 bool HasDefs =
I.getNumDefs() > 0;
978 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
979 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
980 if (spvSelect(ResVReg, ResType,
I)) {
982 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
993 case TargetOpcode::G_CONSTANT:
994 case TargetOpcode::G_FCONSTANT:
1001 MachineInstr &
I)
const {
1004 if (DstRC != SrcRC && SrcRC)
1006 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
1013bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
1014 SPIRVTypeInst ResType,
1015 MachineInstr &
I)
const {
1016 const unsigned Opcode =
I.getOpcode();
1018 return selectImpl(
I, *CoverageInfo);
1020 case TargetOpcode::G_CONSTANT:
1021 case TargetOpcode::G_FCONSTANT:
1022 return selectConst(ResVReg, ResType,
I);
1023 case TargetOpcode::G_GLOBAL_VALUE:
1024 return selectGlobalValue(ResVReg,
I);
1025 case TargetOpcode::G_IMPLICIT_DEF:
1026 return selectOpUndef(ResVReg, ResType,
I);
1027 case TargetOpcode::G_FREEZE:
1028 return selectFreeze(ResVReg, ResType,
I);
1030 case TargetOpcode::G_INTRINSIC:
1031 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
1032 case TargetOpcode::G_INTRINSIC_CONVERGENT:
1033 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
1034 return selectIntrinsic(ResVReg, ResType,
I);
1035 case TargetOpcode::G_BITREVERSE:
1036 return selectBitreverse(ResVReg, ResType,
I);
1038 case TargetOpcode::G_BUILD_VECTOR:
1039 return selectBuildVector(ResVReg, ResType,
I);
1040 case TargetOpcode::G_SPLAT_VECTOR:
1041 return selectSplatVector(ResVReg, ResType,
I);
1042 case TargetOpcode::G_CONCAT_VECTORS:
1043 return selectConcatVectors(ResVReg, ResType,
I);
1045 case TargetOpcode::G_SHUFFLE_VECTOR: {
1046 MachineBasicBlock &BB = *
I.getParent();
1047 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
1050 .
addUse(
I.getOperand(1).getReg())
1051 .
addUse(
I.getOperand(2).getReg());
1052 for (
auto V :
I.getOperand(3).getShuffleMask())
1057 case TargetOpcode::G_MEMMOVE:
1058 case TargetOpcode::G_MEMCPY:
1059 case TargetOpcode::G_MEMCPY_INLINE:
1060 case TargetOpcode::G_MEMSET:
1061 case TargetOpcode::G_MEMSET_INLINE:
1062 return selectMemOperation(ResVReg,
I);
1064 case TargetOpcode::G_ICMP:
1065 return selectICmp(ResVReg, ResType,
I);
1066 case TargetOpcode::G_FCMP:
1067 return selectFCmp(ResVReg, ResType,
I);
1069 case TargetOpcode::G_FRAME_INDEX:
1070 return selectFrameIndex(ResVReg, ResType,
I);
1072 case TargetOpcode::G_LOAD:
1073 return selectLoad(ResVReg, ResType,
I);
1074 case TargetOpcode::G_STORE:
1075 return selectStore(
I);
1077 case TargetOpcode::G_BR:
1078 return selectBranch(
I);
1079 case TargetOpcode::G_BRCOND:
1080 return selectBranchCond(
I);
1082 case TargetOpcode::G_PHI:
1083 return selectPhi(ResVReg,
I);
1085 case TargetOpcode::G_FPTOSI:
1086 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1087 case TargetOpcode::G_FPTOUI:
1088 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1090 case TargetOpcode::G_FPTOSI_SAT:
1091 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1092 case TargetOpcode::G_FPTOUI_SAT:
1093 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1095 case TargetOpcode::G_SITOFP:
1096 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
1097 case TargetOpcode::G_UITOFP:
1098 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
1100 case TargetOpcode::G_CTPOP:
1101 return selectPopCount(ResVReg, ResType,
I, SPIRV::OpBitCount);
1102 case TargetOpcode::G_SMIN:
1103 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
1104 case TargetOpcode::G_UMIN:
1105 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
1107 case TargetOpcode::G_SMAX:
1108 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
1109 case TargetOpcode::G_UMAX:
1110 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
1112 case TargetOpcode::G_SCMP:
1113 return selectSUCmp(ResVReg, ResType,
I,
true);
1114 case TargetOpcode::G_UCMP:
1115 return selectSUCmp(ResVReg, ResType,
I,
false);
1116 case TargetOpcode::G_LROUND:
1117 case TargetOpcode::G_LLROUND: {
1120 MRI->
setRegClass(regForLround, &SPIRV::iIDRegClass);
1122 regForLround, *(
I.getParent()->getParent()));
1124 CL::round, GL::Round,
false);
1126 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1133 case TargetOpcode::G_STRICT_FMA:
1134 case TargetOpcode::G_FMA: {
1137 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1140 .
addUse(
I.getOperand(1).getReg())
1141 .
addUse(
I.getOperand(2).getReg())
1142 .
addUse(
I.getOperand(3).getReg())
1147 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1150 case TargetOpcode::G_FLDEXP:
1151 case TargetOpcode::G_STRICT_FLDEXP:
1152 return selectLdexp(ResVReg, ResType,
I);
1154 case TargetOpcode::G_FPOW:
1155 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1156 case TargetOpcode::G_FPOWI:
1157 return selectFpowi(ResVReg, ResType,
I);
1159 case TargetOpcode::G_FEXP:
1160 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1161 case TargetOpcode::G_FEXP2:
1162 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1163 case TargetOpcode::G_FEXP10:
1164 return selectExp10(ResVReg, ResType,
I);
1166 case TargetOpcode::G_FMODF:
1167 return selectModf(ResVReg, ResType,
I);
1168 case TargetOpcode::G_FSINCOS:
1169 return selectSincos(ResVReg, ResType,
I);
1171 case TargetOpcode::G_FLOG:
1172 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1173 case TargetOpcode::G_FLOG2:
1174 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1175 case TargetOpcode::G_FLOG10:
1176 return selectLog10(ResVReg, ResType,
I);
1178 case TargetOpcode::G_FABS:
1179 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1180 case TargetOpcode::G_ABS:
1181 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1183 case TargetOpcode::G_FMINNUM:
1184 case TargetOpcode::G_FMINIMUM:
1185 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1186 case TargetOpcode::G_FMAXNUM:
1187 case TargetOpcode::G_FMAXIMUM:
1188 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1190 case TargetOpcode::G_FCOPYSIGN:
1191 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1193 case TargetOpcode::G_FCEIL:
1194 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1195 case TargetOpcode::G_FFLOOR:
1196 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1198 case TargetOpcode::G_FCOS:
1199 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1200 case TargetOpcode::G_FSIN:
1201 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1202 case TargetOpcode::G_FTAN:
1203 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1204 case TargetOpcode::G_FACOS:
1205 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1206 case TargetOpcode::G_FASIN:
1207 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1208 case TargetOpcode::G_FATAN:
1209 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1210 case TargetOpcode::G_FATAN2:
1211 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1212 case TargetOpcode::G_FCOSH:
1213 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1214 case TargetOpcode::G_FSINH:
1215 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1216 case TargetOpcode::G_FTANH:
1217 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1219 case TargetOpcode::G_STRICT_FSQRT:
1220 case TargetOpcode::G_FSQRT:
1221 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1223 case TargetOpcode::G_CTTZ:
1224 case TargetOpcode::G_CTTZ_ZERO_POISON:
1225 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1226 case TargetOpcode::G_CTLZ:
1227 case TargetOpcode::G_CTLZ_ZERO_POISON:
1228 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1230 case TargetOpcode::G_INTRINSIC_ROUND:
1231 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1232 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1233 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1234 case TargetOpcode::G_INTRINSIC_TRUNC:
1235 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1236 case TargetOpcode::G_FRINT:
1237 case TargetOpcode::G_FNEARBYINT:
1238 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1240 case TargetOpcode::G_SMULH:
1241 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1242 case TargetOpcode::G_UMULH:
1243 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1245 case TargetOpcode::G_SADDSAT:
1246 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1247 case TargetOpcode::G_UADDSAT:
1248 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1249 case TargetOpcode::G_SSUBSAT:
1250 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1251 case TargetOpcode::G_USUBSAT:
1252 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1254 case TargetOpcode::G_FFREXP:
1255 return selectFrexp(ResVReg, ResType,
I);
1257 case TargetOpcode::G_UADDO:
1258 return selectOverflowArith(ResVReg, ResType,
I,
1259 ResType->
getOpcode() == SPIRV::OpTypeVector
1260 ? SPIRV::OpIAddCarryV
1261 : SPIRV::OpIAddCarryS);
1262 case TargetOpcode::G_USUBO:
1263 return selectOverflowArith(ResVReg, ResType,
I,
1264 ResType->
getOpcode() == SPIRV::OpTypeVector
1265 ? SPIRV::OpISubBorrowV
1266 : SPIRV::OpISubBorrowS);
1267 case TargetOpcode::G_UMULO:
1268 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1269 case TargetOpcode::G_SMULO:
1270 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1272 case TargetOpcode::G_SEXT:
1273 return selectExt(ResVReg, ResType,
I,
true);
1274 case TargetOpcode::G_ANYEXT:
1275 case TargetOpcode::G_ZEXT:
1276 return selectExt(ResVReg, ResType,
I,
false);
1277 case TargetOpcode::G_TRUNC:
1278 return selectTrunc(ResVReg, ResType,
I);
1279 case TargetOpcode::G_FPTRUNC:
1280 case TargetOpcode::G_FPEXT:
1281 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1283 case TargetOpcode::G_PTRTOINT:
1284 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1285 case TargetOpcode::G_INTTOPTR:
1286 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1287 case TargetOpcode::G_BITCAST:
1288 return selectBitcast(ResVReg, ResType,
I);
1289 case TargetOpcode::G_ADDRSPACE_CAST:
1290 return selectAddrSpaceCast(ResVReg, ResType,
I);
1291 case TargetOpcode::G_PTRMASK:
1292 return selectPtrMask(ResVReg, ResType,
I);
1293 case TargetOpcode::G_PTR_ADD: {
1295 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1299 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1300 (*II).getOpcode() == TargetOpcode::COPY ||
1301 (*II).getOpcode() == SPIRV::OpVariable) &&
1302 getImm(
I.getOperand(2), MRI));
1304 bool IsGVInit =
false;
1308 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1309 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1310 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1311 (*UseIt).getOpcode() == SPIRV::OpVariable) {
1321 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1333 return diagnoseUnsupported(
1334 I,
"incompatible result and operand types in a bitcast");
1336 MachineInstrBuilder MIB =
1337 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1344 : SPIRV::OpInBoundsPtrAccessChain))
1348 .
addUse(
I.getOperand(2).getReg())
1351 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1355 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1357 .
addUse(
I.getOperand(2).getReg())
1366 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1369 .
addImm(
static_cast<uint32_t
>(
1370 SPIRV::Opcode::InBoundsPtrAccessChain))
1373 .
addUse(
I.getOperand(2).getReg());
1378 case TargetOpcode::G_ATOMICRMW_OR:
1379 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1380 case TargetOpcode::G_ATOMICRMW_ADD:
1381 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1382 case TargetOpcode::G_ATOMICRMW_AND:
1383 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1384 case TargetOpcode::G_ATOMICRMW_MAX:
1385 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1386 case TargetOpcode::G_ATOMICRMW_MIN:
1387 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1388 case TargetOpcode::G_ATOMICRMW_SUB:
1389 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1390 case TargetOpcode::G_ATOMICRMW_XOR:
1391 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1392 case TargetOpcode::G_ATOMICRMW_UMAX:
1393 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1394 case TargetOpcode::G_ATOMICRMW_UMIN:
1395 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1396 case TargetOpcode::G_ATOMICRMW_XCHG:
1397 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1399 case TargetOpcode::G_ATOMICRMW_FADD:
1400 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1401 case TargetOpcode::G_ATOMICRMW_FSUB:
1403 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1404 ResType->
getOpcode() == SPIRV::OpTypeVector
1406 : SPIRV::OpFNegate);
1407 case TargetOpcode::G_ATOMICRMW_FMIN:
1408 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1409 case TargetOpcode::G_ATOMICRMW_FMAX:
1410 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1412 case TargetOpcode::G_FENCE:
1413 return selectFence(
I);
1415 case TargetOpcode::G_STACKSAVE:
1416 return selectStackSave(ResVReg, ResType,
I);
1417 case TargetOpcode::G_STACKRESTORE:
1418 return selectStackRestore(
I);
1420 case TargetOpcode::G_UNMERGE_VALUES:
1423 case TargetOpcode::G_TRAP:
1424 case TargetOpcode::G_UBSANTRAP:
1425 return selectTrap(
I);
1430 case TargetOpcode::DBG_LABEL:
1432 case TargetOpcode::G_DEBUGTRAP:
1433 return selectDebugTrap(ResVReg, ResType,
I);
1440bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1441 SPIRVTypeInst ResType,
1442 MachineInstr &
I)
const {
1443 unsigned Opcode = SPIRV::OpNop;
1450bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1451 SPIRVTypeInst ResType,
1453 GL::GLSLExtInst GLInst,
1454 bool setMIFlags,
bool useMISrc,
1457 SPIRV::InstructionSet::InstructionSet::GLSL_std_450))
1458 return diagnoseUnsupported(
1460 "this instruction is only supported with the GLSL extended instruction "
1462 return selectExtInst(ResVReg, ResType,
I,
1463 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1464 setMIFlags, useMISrc, SrcRegs);
1467bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1468 SPIRVTypeInst ResType,
1470 CL::OpenCLExtInst CLInst,
1471 bool setMIFlags,
bool useMISrc,
1473 return selectExtInst(ResVReg, ResType,
I,
1474 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1475 setMIFlags, useMISrc, SrcRegs);
1478bool SPIRVInstructionSelector::selectExtInst(
1479 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1480 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1482 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1483 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1484 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1488bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1489 SPIRVTypeInst ResType,
1492 bool setMIFlags,
bool useMISrc,
1495 for (
const auto &[InstructionSet, Opcode] : Insts) {
1499 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1502 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1507 const unsigned NumOps =
I.getNumOperands();
1510 I.getOperand(Index).getType() ==
1511 MachineOperand::MachineOperandType::MO_IntrinsicID)
1514 MIB.
add(
I.getOperand(Index));
1526bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1527 SPIRVTypeInst ResType,
1528 MachineInstr &
I)
const {
1529 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1530 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1531 for (
const auto &Ex : ExtInsts) {
1532 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1533 uint32_t Opcode = Ex.second;
1537 MachineIRBuilder MIRBuilder(
I);
1540 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1545 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1548 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1552 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1555 .
addImm(
static_cast<uint32_t
>(Ex.first))
1557 .
add(
I.getOperand(2))
1561 Register ExpResReg =
I.getOperand(1).getReg();
1563 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1573bool SPIRVInstructionSelector::selectLdexp(
Register ResVReg,
1574 SPIRVTypeInst ResType,
1575 MachineInstr &
I)
const {
1576 Register XReg =
I.getOperand(1).getReg();
1577 Register ExpReg =
I.getOperand(2).getReg();
1583 if (ResType->
getOpcode() == SPIRV::OpTypeVector &&
1584 ExpType->
getOpcode() != SPIRV::OpTypeVector) {
1586 SPIRVTypeInst ExpVecType =
1590 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1591 TII.get(SPIRV::OpCompositeConstruct))
1594 for (
unsigned J = 0; J < NumElts; ++J)
1600 return selectExtInst(ResVReg, ResType,
I, CL::ldexp, GL::Ldexp,
1601 true,
false, {XReg, ExpReg});
1604bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1605 SPIRVTypeInst ResType,
1606 MachineInstr &
I)
const {
1607 Register CosResVReg =
I.getOperand(1).getReg();
1608 unsigned SrcIdx =
I.getNumExplicitDefs();
1613 MachineIRBuilder MIRBuilder(
I);
1615 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1620 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1623 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1625 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1628 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1630 .
add(
I.getOperand(SrcIdx))
1633 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1641 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1644 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1646 .
add(
I.getOperand(SrcIdx))
1648 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1651 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1653 .
add(
I.getOperand(SrcIdx))
1660bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1661 SPIRVTypeInst ResType,
1664 unsigned Opcode)
const {
1665 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1675std::optional<SplitParts> SPIRVInstructionSelector::splitEvenOddLanes(
1676 Register PopCountReg,
unsigned ComponentCount, MachineInstr &
I,
1677 SPIRVTypeInst I32Type)
const {
1680 if (ComponentCount == 1) {
1683 Parts.IsScalar =
true;
1684 Parts.Type = I32Type;
1692 if (!selectOpWithSrcs(Parts.High, I32Type,
I, {PopCountReg, IdxOne},
1693 SPIRV::OpVectorExtractDynamic))
1694 return std::nullopt;
1696 if (!selectOpWithSrcs(Parts.Low, I32Type,
I, {PopCountReg, IdxZero},
1697 SPIRV::OpVectorExtractDynamic))
1698 return std::nullopt;
1702 MachineIRBuilder MIRBuilder(
I);
1703 Parts.IsScalar =
false;
1710 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1711 TII.get(SPIRV::OpVectorShuffle))
1716 for (
unsigned J = 1; J < ComponentCount * 2; J += 2)
1721 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1722 TII.get(SPIRV::OpVectorShuffle))
1727 for (
unsigned J = 0; J < ComponentCount * 2; J += 2)
1735bool SPIRVInstructionSelector::selectPopCount16(
Register ResVReg,
1736 SPIRVTypeInst ResType,
1739 unsigned Opcode)
const {
1740 Register OpReg =
I.getOperand(1).getReg();
1743 MachineIRBuilder MIRBuilder(
I);
1745 SPIRVTypeInst I32VectorType =
1748 bool IsVector = NumElems > 1;
1749 SPIRVTypeInst ExtType = IsVector ? I32VectorType : I32Type;
1752 if (!selectOpWithSrcs(ExtReg, ExtType,
I, {OpReg}, SPIRV::OpUConvert))
1756 if (!selectPopCount32(PopCountReg, ExtType,
I, ExtReg, Opcode))
1759 return selectOpWithSrcs(ResVReg, ResType,
I, {PopCountReg}, ExtOpcode);
1762bool SPIRVInstructionSelector::selectPopCount32(
Register ResVReg,
1763 SPIRVTypeInst ResType,
1766 unsigned Opcode)
const {
1767 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
1770bool SPIRVInstructionSelector::selectPopCount64(
Register ResVReg,
1771 SPIRVTypeInst ResType,
1774 unsigned Opcode)
const {
1776 if (ComponentCount > 2)
1777 return handle64BitOverflow(
1778 ResVReg, ResType,
I, SrcReg, Opcode,
1780 unsigned O) {
return this->selectPopCount64(R,
T,
I, S, O); });
1782 MachineIRBuilder MIRBuilder(
I);
1787 I32Type, 2 * ComponentCount, MIRBuilder,
false);
1791 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
1796 if (!selectPopCount32(Pop32, VecI32Type,
I, Vec32, Opcode))
1800 auto MaybeParts = splitEvenOddLanes(Pop32, ComponentCount,
I, I32Type);
1803 SplitParts &Parts = *MaybeParts;
1806 unsigned OpAdd = Parts.IsScalar ? SPIRV::OpIAddS : SPIRV::OpIAddV;
1808 if (!selectOpWithSrcs(Sum, Parts.Type,
I, {Parts.High, Parts.Low}, OpAdd))
1813 unsigned ConvOp = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
1814 return selectOpWithSrcs(ResVReg, ResType,
I, {Sum}, ConvOp);
1817bool SPIRVInstructionSelector::selectPopCount(
Register ResVReg,
1818 SPIRVTypeInst ResType,
1820 unsigned Opcode)
const {
1825 if (!STI.getTargetTriple().isVulkanOS())
1826 return selectUnOp(ResVReg, ResType,
I, Opcode);
1828 Register OpReg =
I.getOperand(1).getReg();
1831 : SPIRV::OpUConvert;
1835 return selectPopCount16(ResVReg, ResType,
I, ExtOpcode, Opcode);
1837 return selectPopCount32(ResVReg, ResType,
I, OpReg, Opcode);
1839 return selectPopCount64(ResVReg, ResType,
I, OpReg, Opcode);
1841 return diagnoseUnsupported(
I,
"unsupported operand bit width for popcount");
1845bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1846 SPIRVTypeInst ResType,
1848 unsigned Opcode)
const {
1850 Register SrcReg =
I.getOperand(1).getReg();
1855 unsigned DefOpCode = DefIt->getOpcode();
1856 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1859 if (
auto *VRD =
getVRegDef(*MRI, DefIt->getOperand(1).getReg()))
1860 DefOpCode = VRD->getOpcode();
1862 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1863 DefOpCode == TargetOpcode::G_CONSTANT ||
1864 DefOpCode == SPIRV::OpVariable || DefOpCode == SPIRV::OpConstantI) {
1870 uint32_t SpecOpcode = 0;
1872 case SPIRV::OpConvertPtrToU:
1873 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1875 case SPIRV::OpConvertUToPtr:
1876 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1881 TII.get(SPIRV::OpSpecConstantOp))
1891 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1895bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1896 SPIRVTypeInst ResType,
1897 MachineInstr &
I)
const {
1898 Register OpReg =
I.getOperand(1).getReg();
1899 SPIRVTypeInst OpType =
1902 return diagnoseUnsupported(
1903 I,
"incompatible result and operand types in a bitcast");
1904 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1915 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1916 if (
MemOp->isNonTemporal())
1917 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1919 if (!ST->isShader() &&
MemOp->getAlign().value())
1920 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1924 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1925 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1929 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1931 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1935 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1939 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1941 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1953 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1955 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1957 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
1961bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
1962 SPIRVTypeInst ResType,
1963 MachineInstr &
I)
const {
1965 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1970 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
1971 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
1973 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1975 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1979 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1983 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1984 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
1985 I.getDebugLoc(),
I);
1989 MachineIRBuilder MIRBuilder(
I);
1991 if (
I.getNumMemOperands()) {
1992 const MachineMemOperand *MemOp = *
I.memoperands_begin();
1993 if (MemOp->isAtomic())
1994 return selectAtomicLoad(ResVReg, ResType,
I);
1997 auto MIB = MIRBuilder.
buildInstr(SPIRV::OpLoad)
2001 if (!
I.getNumMemOperands()) {
2002 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2004 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2013Register SPIRVInstructionSelector::createPtrSizedIntReg(
2014 MachineIRBuilder &MIRBuilder)
const {
2015 SPIRVTypeInst IntType =
2025SPIRVInstructionSelector::convertPtrToInt(
Register PtrVal,
2026 MachineIRBuilder &MIRBuilder)
const {
2027 SPIRVTypeInst IntType =
2029 Register IntReg = createPtrSizedIntReg(MIRBuilder);
2030 MIRBuilder.
buildInstr(SPIRV::OpConvertPtrToU)
2038Register SPIRVInstructionSelector::castPtrToPtrToInt(
2039 Register Ptr, SPIRV::StorageClass::StorageClass SC,
2040 MachineIRBuilder &MIRBuilder)
const {
2041 SPIRVTypeInst IntType =
2043 SPIRVTypeInst PtrType =
2057bool SPIRVInstructionSelector::selectAtomicPtrValue(
2058 Register ResVReg, SPIRVTypeInst ResType, MachineIRBuilder &MIRBuilder,
2059 function_ref<
Register(SPIRVTypeInst IntType)> EmitAtomic)
const {
2067 Register IntResult = EmitAtomic(IntType);
2069 MIRBuilder.
buildInstr(SPIRV::OpConvertUToPtr)
2077bool SPIRVInstructionSelector::selectAtomicLoad(
Register ResVReg,
2078 SPIRVTypeInst ResType,
2079 MachineInstr &
I)
const {
2080 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2083 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2086 return diagnoseUnsupported(
2087 I,
"Lowering to SPIR-V of atomic load is only "
2088 "allowed for integer, floating point or pointer types");
2090 assert(
I.getNumMemOperands());
2091 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2092 assert(MemOp.isAtomic());
2096 Register ScopeReg = buildI32Constant(Scope,
I);
2102 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2103 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2106 MachineIRBuilder MIRBuilder(
I);
2110 return diagnoseUnsupported(
2111 I,
"Lowering to SPIR-V of atomic load is only "
2112 "allowed for pointer types for physical addressing model");
2117 SPIRV::StorageClass::StorageClass SC =
2119 return selectAtomicPtrValue(
2120 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2121 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2122 Register IntResult = createPtrSizedIntReg(MIRBuilder);
2133 auto AtomicLoad = MIRBuilder.
buildInstr(SPIRV::OpAtomicLoad)
2144bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
2146 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2147 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2152 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2153 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2155 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2160 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2164 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2165 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2166 SPIRVTypeInst SampledType =
2168 SPIRVTypeInst StoreValCompType =
2170 if (StoreValCompType && StoreValCompType != SampledType) {
2173 SPIRVTypeInst PackedType = widenTypeToVec4(SampledType,
I);
2176 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
2181 StoreVal = PackedReg;
2184 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2185 TII.get(SPIRV::OpImageWrite))
2191 if (sampledTypeIsSignedInteger(LLVMHandleType))
2194 BMI.constrainAllUses(
TII,
TRI, RBI);
2199 if (
I.getNumMemOperands()) {
2200 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2201 if (MemOp->isAtomic())
2202 return selectAtomicStore(
I);
2209 if (PtrSC == SPIRV::StorageClass::UniformConstant ||
2210 PtrSC == SPIRV::StorageClass::Input ||
2211 PtrSC == SPIRV::StorageClass::PushConstant)
2212 return diagnoseUnsupported(
2213 I,
"store into a read-only SPIR-V storage class is not allowed");
2215 MachineIRBuilder MIRBuilder(
I);
2217 if (!
I.getNumMemOperands()) {
2218 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2220 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2229bool SPIRVInstructionSelector::selectAtomicStore(MachineInstr &
I)
const {
2230 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2233 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2234 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2239 assert(
I.getNumMemOperands());
2240 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2241 assert(MemOp.isAtomic());
2245 Register ScopeReg = buildI32Constant(Scope,
I);
2251 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2252 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2254 MachineIRBuilder MIRBuilder(
I);
2258 return diagnoseUnsupported(
2259 I,
"Lowering to SPIR-V of atomic store is only "
2260 "allowed for pointer types for physical addressing model");
2265 SPIRV::StorageClass::StorageClass SC =
2267 return selectAtomicPtrValue(
2268 Register(), SPIRVTypeInst(), MIRBuilder, [&](SPIRVTypeInst IntType) {
2270 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2283 return diagnoseUnsupported(
I,
2284 "Lowering to SPIR-V of atomic store is only "
2285 "allowed for integer or floating point types");
2287 auto AtomicStore = MIRBuilder.
buildInstr(SPIRV::OpAtomicStore)
2297bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
2298 SPIRVTypeInst ResType,
2299 MachineInstr &
I)
const {
2300 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
2308 const Register PtrsReg =
I.getOperand(2).getReg();
2309 const uint32_t Alignment =
I.getOperand(3).getImm();
2310 const Register MaskReg =
I.getOperand(4).getReg();
2311 const Register PassthruReg =
I.getOperand(5).getReg();
2312 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2316 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
2327bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
2328 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
2335 const Register ValuesReg =
I.getOperand(1).getReg();
2336 const Register PtrsReg =
I.getOperand(2).getReg();
2337 const uint32_t Alignment =
I.getOperand(3).getImm();
2338 const Register MaskReg =
I.getOperand(4).getReg();
2339 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2343 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
2352bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
2353 const Twine &
Msg)
const {
2354 const Function &
F =
I.getMF()->getFunction();
2355 F.getContext().diagnose(
2356 DiagnosticInfoUnsupported(
F,
Msg,
I.getDebugLoc(),
DS_Error));
2360bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
2361 SPIRVTypeInst ResType,
2362 MachineInstr &
I)
const {
2363 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2364 return diagnoseUnsupported(
2365 I,
"llvm.stacksave intrinsic: this instruction requires the following "
2366 "SPIR-V extension: SPV_INTEL_variable_length_array");
2368 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
2375bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
2376 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2377 return diagnoseUnsupported(
2379 "llvm.stackrestore intrinsic: this instruction requires the following "
2380 "SPIR-V extension: SPV_INTEL_variable_length_array");
2381 if (!
I.getOperand(0).isReg())
2384 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
2385 .
addUse(
I.getOperand(0).getReg())
2391SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
2392 MachineIRBuilder MIRBuilder(
I);
2393 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
2400 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
2404 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2405 Type *ArrTy = ArrayType::get(ValTy, Num);
2407 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
2410 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
2417 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
2420 .
addImm(SPIRV::StorageClass::UniformConstant)
2431bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
2434 Register DstReg =
I.getOperand(0).getReg();
2438 return diagnoseUnsupported(
2439 I,
"OpCopyMemory requires operands to have the same type");
2440 uint64_t CopySize =
getIConstVal(
I.getOperand(2).getReg(), MRI);
2444 return diagnoseUnsupported(
2445 I,
"Unable to determine pointee type size for OpCopyMemory");
2446 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
2447 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
2448 return diagnoseUnsupported(
2449 I,
"OpCopyMemory requires the size to match the pointee type size");
2450 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemory))
2453 if (
I.getNumMemOperands()) {
2454 MachineIRBuilder MIRBuilder(
I);
2461bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &
I,
2464 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemorySized))
2465 .
addUse(
I.getOperand(0).getReg())
2467 .
addUse(
I.getOperand(2).getReg());
2468 if (
I.getNumMemOperands()) {
2469 MachineIRBuilder MIRBuilder(
I);
2476bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
2477 MachineInstr &
I)
const {
2479 Register SizeReg =
I.getOperand(2).getReg();
2481 SizeDef && SizeDef->
getOpcode() == TargetOpcode::G_CONSTANT &&
2485 Register SrcReg =
I.getOperand(1).getReg();
2486 if (
I.getOpcode() == TargetOpcode::G_MEMSET ||
2487 I.getOpcode() == TargetOpcode::G_MEMSET_INLINE) {
2488 Register VarReg = getOrCreateMemSetGlobal(
I);
2491 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2493 ValTy,
I, SPIRV::StorageClass::UniformConstant);
2495 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
2499 if (!selectCopyMemory(
I, SrcReg))
2502 if (!selectCopyMemorySized(
I, SrcReg))
2505 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
2506 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
2511bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
2512 SPIRVTypeInst ResType,
2515 unsigned NegateOpcode)
const {
2517 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2520 Register ScopeReg = buildI32Constant(Scope,
I);
2522 Register Ptr =
I.getOperand(1).getReg();
2523 uint32_t ScSem =
static_cast<uint32_t
>(
2527 Register MemSemReg = buildI32Constant(MemSem,
I);
2529 Register ValueReg =
I.getOperand(2).getReg();
2530 if (NegateOpcode != 0) {
2533 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
2539 if (NewOpcode != SPIRV::OpAtomicExchange)
2540 return diagnoseUnsupported(
2541 I,
"Lowering to SPIR-V of this atomic operation is not "
2542 "allowed for pointer types");
2544 return diagnoseUnsupported(
2545 I,
"Lowering to SPIR-V of atomic exchange is only "
2546 "allowed for pointer types for physical addressing model");
2553 MachineIRBuilder MIRBuilder(
I);
2555 return selectAtomicPtrValue(
2556 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2558 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2559 Register ExchangeResReg = createPtrSizedIntReg(MIRBuilder);
2560 MIRBuilder.
buildInstr(SPIRV::OpAtomicExchange)
2568 return ExchangeResReg;
2572 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2583bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2584 unsigned ArgI =
I.getNumOperands() - 1;
2586 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2587 SPIRVTypeInst SrcType =
2589 if (!SrcType || SrcType->
getOpcode() != SPIRV::OpTypeVector)
2591 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2595 unsigned CurrentIndex = 0;
2596 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2597 Register ResVReg =
I.getOperand(i).getReg();
2600 LLT ResLLT = MRI->
getType(ResVReg);
2606 ResType = ScalarType;
2612 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
2615 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2621 for (
unsigned j = 0;
j < NumElements; ++
j) {
2622 MIB.
addImm(CurrentIndex + j);
2624 CurrentIndex += NumElements;
2628 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2640bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2643 Register MemSemReg = buildI32Constant(MemSem,
I);
2645 uint32_t
Scope =
static_cast<uint32_t
>(
2647 Register ScopeReg = buildI32Constant(Scope,
I);
2649 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2656bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2657 SPIRVTypeInst ResType,
2659 unsigned Opcode)
const {
2660 Type *ResTy =
nullptr;
2663 return diagnoseUnsupported(
2665 "Not enough info to select the arithmetic with overflow instruction");
2667 return diagnoseUnsupported(
I,
2668 "Expect struct type result for the arithmetic "
2669 "with overflow instruction");
2675 MachineIRBuilder MIRBuilder(
I);
2677 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2678 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2684 Register ZeroReg = buildZerosVal(ResType,
I);
2689 if (ResName.
size() > 0)
2697 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2698 MIB.
addUse(
I.getOperand(i).getReg());
2703 MRI->
setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2704 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2706 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2707 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2714 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2715 .
addDef(
I.getOperand(1).getReg())
2723bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2724 SPIRVTypeInst ResType,
2725 MachineInstr &
I)
const {
2727 "selectAtomicCmpXchg only handles the spv_cmpxchg intrinsic");
2728 Register Ptr =
I.getOperand(2).getReg();
2729 Register ScopeReg =
I.getOperand(5).getReg();
2730 Register MemSemEqReg =
I.getOperand(6).getReg();
2731 Register MemSemNeqReg =
I.getOperand(7).getReg();
2733 Register Val =
I.getOperand(4).getReg();
2737 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2756 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2763 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2775 case SPIRV::StorageClass::DeviceOnlyINTEL:
2776 case SPIRV::StorageClass::HostOnlyINTEL:
2785 bool IsGRef =
false;
2786 bool IsAllowedRefs =
2788 unsigned Opcode = It.getOpcode();
2789 if (Opcode == SPIRV::OpConstantComposite ||
2790 Opcode == SPIRV::OpSpecConstantComposite ||
2791 Opcode == SPIRV::OpVariable ||
2792 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2793 return IsGRef = true;
2794 return Opcode == SPIRV::OpName;
2796 return IsAllowedRefs && IsGRef;
2799Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2800 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2802 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2806SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2808 uint32_t Opcode)
const {
2809 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2810 TII.get(SPIRV::OpSpecConstantOp))
2818SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2819 SPIRVTypeInst SrcPtrTy)
const {
2820 SPIRVTypeInst GenericPtrTy =
2824 SPIRV::StorageClass::Generic),
2826 MachineFunction *MF =
I.getParent()->getParent();
2828 MachineInstrBuilder MIB = buildSpecConstantOp(
2830 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2840bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2841 SPIRVTypeInst ResType,
2842 MachineInstr &
I)
const {
2846 Register SrcPtr =
I.getOperand(1).getReg();
2850 if (SrcPtrTy->
getOpcode() != SPIRV::OpTypePointer ||
2851 ResType->
getOpcode() != SPIRV::OpTypePointer)
2852 return BuildCOPY(ResVReg, SrcPtr,
I);
2862 unsigned SpecOpcode = [&]() ->
unsigned {
2863 if (SrcSC == SPIRV::StorageClass::CodeSectionINTEL)
2864 return static_cast<uint32_t
>(SPIRV::Opcode::Bitcast);
2866 return static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric);
2868 return static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr);
2876 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2878 .constrainAllUses(
TII,
TRI, RBI);
2880 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2882 buildSpecConstantOp(
2884 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2885 .constrainAllUses(
TII,
TRI, RBI);
2892 return BuildCOPY(ResVReg, SrcPtr,
I);
2894 if ((SrcSC == SPIRV::StorageClass::Function &&
2895 DstSC == SPIRV::StorageClass::Private) ||
2896 (DstSC == SPIRV::StorageClass::Function &&
2897 SrcSC == SPIRV::StorageClass::Private))
2898 return BuildCOPY(ResVReg, SrcPtr,
I);
2902 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2905 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2908 SPIRVTypeInst GenericPtrTy =
2927 return selectUnOp(ResVReg, ResType,
I,
2928 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
2930 return selectUnOp(ResVReg, ResType,
I,
2931 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
2933 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2935 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2945bool SPIRVInstructionSelector::selectPtrMask(
Register ResVReg,
2946 SPIRVTypeInst ResType,
2947 MachineInstr &
I)
const {
2949 return diagnoseUnsupported(
2950 I,
"G_PTRMASK is not supported with logical SPIR-V");
2955 Register PtrReg =
I.getOperand(1).getReg();
2956 Register MaskReg =
I.getOperand(2).getReg();
2975 ? SPIRV::OpBitwiseAndV
2976 : SPIRV::OpBitwiseAndS;
2999 return SPIRV::OpFOrdEqual;
3001 return SPIRV::OpFOrdGreaterThanEqual;
3003 return SPIRV::OpFOrdGreaterThan;
3005 return SPIRV::OpFOrdLessThanEqual;
3007 return SPIRV::OpFOrdLessThan;
3009 return SPIRV::OpFOrdNotEqual;
3011 return SPIRV::OpOrdered;
3013 return SPIRV::OpFUnordEqual;
3015 return SPIRV::OpFUnordGreaterThanEqual;
3017 return SPIRV::OpFUnordGreaterThan;
3019 return SPIRV::OpFUnordLessThanEqual;
3021 return SPIRV::OpFUnordLessThan;
3023 return SPIRV::OpFUnordNotEqual;
3025 return SPIRV::OpUnordered;
3035 return SPIRV::OpIEqual;
3037 return SPIRV::OpINotEqual;
3039 return SPIRV::OpSGreaterThanEqual;
3041 return SPIRV::OpSGreaterThan;
3043 return SPIRV::OpSLessThanEqual;
3045 return SPIRV::OpSLessThan;
3047 return SPIRV::OpUGreaterThanEqual;
3049 return SPIRV::OpUGreaterThan;
3051 return SPIRV::OpULessThanEqual;
3053 return SPIRV::OpULessThan;
3062 return SPIRV::OpPtrEqual;
3064 return SPIRV::OpPtrNotEqual;
3075 return SPIRV::OpLogicalEqual;
3077 return SPIRV::OpLogicalNotEqual;
3115bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
3116 SPIRVTypeInst ResType,
3118 unsigned OpAnyOrAll)
const {
3119 assert(
I.getNumOperands() == 3);
3120 assert(
I.getOperand(2).isReg());
3122 Register InputRegister =
I.getOperand(2).getReg();
3125 assert(InputType &&
"VReg has no type assigned");
3128 bool IsVectorTy = InputType->
getOpcode() == SPIRV::OpTypeVector;
3129 if (IsBoolTy && !IsVectorTy) {
3130 assert(ResVReg ==
I.getOperand(0).getReg());
3131 return BuildCOPY(ResVReg, InputRegister,
I);
3135 unsigned SpirvNotEqualId =
3136 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
3138 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
3143 IsBoolTy ? InputRegister
3151 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
3153 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
3170bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
3171 SPIRVTypeInst ResType,
3172 MachineInstr &
I)
const {
3173 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
3176bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
3177 SPIRVTypeInst ResType,
3178 MachineInstr &
I)
const {
3179 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
3183bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
3184 SPIRVTypeInst ResType,
3185 MachineInstr &
I)
const {
3186 assert(
I.getNumOperands() == 4);
3187 assert(
I.getOperand(2).isReg());
3188 assert(
I.getOperand(3).isReg());
3190 [[maybe_unused]] SPIRVTypeInst VecType =
3195 "dot product requires a vector of at least 2 components");
3197 [[maybe_unused]] SPIRVTypeInst EltType =
3206 .
addUse(
I.getOperand(2).getReg())
3207 .
addUse(
I.getOperand(3).getReg())
3212bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
3213 SPIRVTypeInst ResType,
3216 assert(
I.getNumOperands() == 4);
3217 assert(
I.getOperand(2).isReg());
3218 assert(
I.getOperand(3).isReg());
3221 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3225 .
addUse(
I.getOperand(2).getReg())
3226 .
addUse(
I.getOperand(3).getReg())
3233bool SPIRVInstructionSelector::selectIntegerDotExpansion(
3234 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3235 assert(
I.getNumOperands() == 4);
3236 assert(
I.getOperand(2).isReg());
3237 assert(
I.getOperand(3).isReg());
3241 Register Vec0 =
I.getOperand(2).getReg();
3242 Register Vec1 =
I.getOperand(3).getReg();
3246 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
3255 "dot product requires a vector of at least 2 components");
3258 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3268 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3279 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3291bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
3292 SPIRVTypeInst ResType,
3293 MachineInstr &
I)
const {
3295 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
3298 .
addUse(
I.getOperand(2).getReg())
3303bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
3304 SPIRVTypeInst ResType,
3305 MachineInstr &
I)
const {
3307 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
3310 .
addUse(
I.getOperand(2).getReg())
3315bool SPIRVInstructionSelector::selectOpIsFinite(
Register ResVReg,
3316 SPIRVTypeInst ResType,
3317 MachineInstr &
I)
const {
3319 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsFinite))
3322 .
addUse(
I.getOperand(2).getReg())
3327bool SPIRVInstructionSelector::selectOpIsNormal(
Register ResVReg,
3328 SPIRVTypeInst ResType,
3329 MachineInstr &
I)
const {
3331 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNormal))
3334 .
addUse(
I.getOperand(2).getReg())
3339template <
bool Signed>
3340bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
3341 SPIRVTypeInst ResType,
3342 MachineInstr &
I)
const {
3343 assert(
I.getNumOperands() == 5);
3344 assert(
I.getOperand(2).isReg());
3345 assert(
I.getOperand(3).isReg());
3346 assert(
I.getOperand(4).isReg());
3349 Register Acc =
I.getOperand(2).getReg();
3353 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3355 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
3360 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
3363 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3375template <
bool Signed>
3376bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
3377 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3378 assert(
I.getNumOperands() == 5);
3379 assert(
I.getOperand(2).isReg());
3380 assert(
I.getOperand(3).isReg());
3381 assert(
I.getOperand(4).isReg());
3384 Register Acc =
I.getOperand(2).getReg();
3390 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
3394 for (
unsigned i = 0; i < 4; i++) {
3417 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
3437 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3452bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
3453 SPIRVTypeInst ResType,
3454 MachineInstr &
I)
const {
3455 assert(
I.getNumOperands() == 3);
3456 assert(
I.getOperand(2).isReg());
3458 Register VZero = buildZerosValF(ResType,
I);
3459 Register VOne = buildOnesValF(ResType,
I);
3461 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
3464 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3466 .
addUse(
I.getOperand(2).getReg())
3473bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
3474 SPIRVTypeInst ResType,
3475 MachineInstr &
I)
const {
3476 assert(
I.getNumOperands() == 3);
3477 assert(
I.getOperand(2).isReg());
3479 Register InputRegister =
I.getOperand(2).getReg();
3481 auto &
DL =
I.getDebugLoc();
3484 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3491 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
3493 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
3501 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3506 if (NeedsConversion) {
3507 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
3518bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
3519 SPIRVTypeInst ResType,
3521 unsigned Opcode)
const {
3525 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
3531 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
3532 BMI.addUse(
I.getOperand(J).getReg());
3539bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
3542 bool WithGroupSync)
const {
3544 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
3546 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
3548 assert(((Scope != SPIRV::Scope::Workgroup) ||
3549 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
3550 "Workgroup Scope must set WorkGroupMemory semantic "
3551 "in Barrier instruction");
3553 assert(((Scope != SPIRV::Scope::Device) ||
3554 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
3555 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
3556 "Device Scope must set UniformMemory and ImageMemory semantic "
3557 "in Barrier instruction");
3563 if (WithGroupSync) {
3564 Register ExecReg = buildI32Constant(SPIRV::Scope::Workgroup,
I);
3568 Register ScopeReg = buildI32Constant(Scope,
I);
3569 Register MemSemReg = buildI32Constant(MemSem,
I);
3571 MI.addUse(ScopeReg).addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
3575bool SPIRVInstructionSelector::selectWaveActiveCountBits(
3576 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3581 if (!selectWaveOpInst(BallotReg, BallotType,
I,
3582 SPIRV::OpGroupNonUniformBallot))
3587 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3592 .
addImm(SPIRV::GroupOperation::Reduce)
3599bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
3600 SPIRVTypeInst ResType,
3601 MachineInstr &
I)
const {
3606 Register InputReg =
I.getOperand(2).getReg();
3611 bool IsVector = NumElems > 1;
3624 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
3625 SPIRV::OpGroupNonUniformAllEqual);
3630 ElementResults.
reserve(NumElems);
3632 for (
unsigned Idx = 0; Idx < NumElems; ++Idx) {
3645 ElemInput = Extracted;
3651 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3662 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3673bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3674 SPIRVTypeInst ResType,
3675 MachineInstr &
I)
const {
3677 assert(
I.getNumOperands() == 3);
3679 auto Op =
I.getOperand(2);
3689 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3691 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3692 return diagnoseUnsupported(
I,
"WavePrefixBitCount requires boolean input");
3713 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3717 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3724bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3725 SPIRVTypeInst ResType,
3727 bool IsUnsigned)
const {
3728 return selectWaveReduce(
3729 ResVReg, ResType,
I, IsUnsigned,
3730 [&](
Register InputRegister,
bool IsUnsigned) {
3731 const bool IsFloatTy =
3733 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3734 : SPIRV::OpGroupNonUniformSMax;
3735 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3739bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3740 SPIRVTypeInst ResType,
3742 bool IsUnsigned)
const {
3743 return selectWaveReduce(
3744 ResVReg, ResType,
I, IsUnsigned,
3745 [&](
Register InputRegister,
bool IsUnsigned) {
3746 const bool IsFloatTy =
3748 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3749 : SPIRV::OpGroupNonUniformSMin;
3750 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3754bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3755 SPIRVTypeInst ResType,
3756 MachineInstr &
I)
const {
3757 return selectWaveReduce(ResVReg, ResType,
I,
false,
3758 [&](
Register InputRegister,
bool IsUnsigned) {
3760 InputRegister, SPIRV::OpTypeFloat);
3761 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3762 : SPIRV::OpGroupNonUniformIAdd;
3766bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3767 SPIRVTypeInst ResType,
3768 MachineInstr &
I)
const {
3769 return selectWaveReduce(ResVReg, ResType,
I,
false,
3770 [&](
Register InputRegister,
bool IsUnsigned) {
3772 InputRegister, SPIRV::OpTypeFloat);
3773 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3774 : SPIRV::OpGroupNonUniformIMul;
3778template <
typename PickOpcodeFn>
3779bool SPIRVInstructionSelector::selectWaveReduce(
3780 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3781 PickOpcodeFn &&PickOpcode)
const {
3782 assert(
I.getNumOperands() == 3);
3783 assert(
I.getOperand(2).isReg());
3785 Register InputRegister =
I.getOperand(2).getReg();
3789 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3792 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3798 .
addImm(SPIRV::GroupOperation::Reduce)
3799 .
addUse(
I.getOperand(2).getReg())
3804bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3805 SPIRVTypeInst ResType,
3807 unsigned Opcode)
const {
3808 return selectWaveReduce(
3809 ResVReg, ResType,
I,
false,
3810 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3813bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3814 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3815 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3816 [&](
Register InputRegister,
bool IsUnsigned) {
3818 InputRegister, SPIRV::OpTypeFloat);
3820 ? SPIRV::OpGroupNonUniformFAdd
3821 : SPIRV::OpGroupNonUniformIAdd;
3825bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3826 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3827 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3828 [&](
Register InputRegister,
bool IsUnsigned) {
3830 InputRegister, SPIRV::OpTypeFloat);
3832 ? SPIRV::OpGroupNonUniformFMul
3833 : SPIRV::OpGroupNonUniformIMul;
3837template <
typename PickOpcodeFn>
3838bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3839 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3840 PickOpcodeFn &&PickOpcode)
const {
3841 assert(
I.getNumOperands() == 3);
3842 assert(
I.getOperand(2).isReg());
3844 Register InputRegister =
I.getOperand(2).getReg();
3848 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3851 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3857 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3858 .
addUse(
I.getOperand(2).getReg())
3863bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
3864 SPIRVTypeInst ResType,
3867 assert(
I.getNumOperands() == 3);
3868 assert(
I.getOperand(2).isReg());
3870 Register InputRegister =
I.getOperand(2).getReg();
3876 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
3887bool SPIRVInstructionSelector::selectBitreverseViaI32(
Register ResVReg,
3888 SPIRVTypeInst ResType,
3895 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
3900 : SPIRV::OpUConvert;
3904 ShiftOp = SPIRV::OpShiftRightLogicalV;
3909 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3910 TII.get(SPIRV::OpConstantComposite))
3913 for (
unsigned It = 0; It <
N; ++It)
3917 ShiftConst = CompositeReg;
3922 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
3927 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
3932 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
3937 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
3940bool SPIRVInstructionSelector::handle64BitOverflow(
3942 unsigned int Opcode,
3949 "handle64BitOverflow should only be used for integer types");
3951 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
3953 MachineIRBuilder MIRBuilder(
I);
3955 SPIRVTypeInst I64x2Type =
3957 SPIRVTypeInst Vec2ResType =
3960 std::vector<Register> PartialRegs;
3962 unsigned CurrentComponent = 0;
3963 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
3967 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3968 TII.get(SPIRV::OpVectorShuffle))
3973 .
addImm(CurrentComponent)
3974 .
addImm(CurrentComponent + 1);
3984 PartialRegs.push_back(SubVecReg);
3987 if (CurrentComponent != ComponentCount) {
3993 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
3994 SPIRV::OpVectorExtractDynamic))
4003 PartialRegs.push_back(FinalElemResReg);
4007 return selectOpWithSrcs(ResVReg, ResType,
I, PartialRegs,
4008 SPIRV::OpCompositeConstruct);
4011bool SPIRVInstructionSelector::selectBitreverse64(
Register ResVReg,
4012 SPIRVTypeInst ResType,
4016 if (ComponentCount > 2)
4017 return handle64BitOverflow(
4018 ResVReg, ResType,
I, SrcReg, SPIRV::OpBitReverse,
4020 unsigned O) {
return this->selectBitreverse64(R,
T,
I, S); });
4022 MachineIRBuilder MIRBuilder(
I);
4026 I32Type, 2 * ComponentCount, MIRBuilder,
false);
4030 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
4035 if (!selectBitreverseNative(Reverse32, VecI32Type,
I, Vec32))
4042 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4043 TII.get(SPIRV::OpVectorShuffle))
4048 for (
unsigned J = 0; J < ComponentCount; ++J) {
4055 return selectOpWithSrcs(ResVReg, ResType,
I, {SwappedVec}, SPIRV::OpBitcast);
4058bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
4059 SPIRVTypeInst ResType,
4063 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
4071bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
4072 SPIRVTypeInst ResType,
4073 MachineInstr &
I)
const {
4074 Register OpReg =
I.getOperand(1).getReg();
4083 return selectBitreverseViaI32(ResVReg, ResType,
I, OpReg);
4085 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4087 return selectBitreverse64(ResVReg, ResType,
I, OpReg);
4089 return SPIRVInstructionSelector::diagnoseUnsupported(
4090 I,
"G_BITREVERSE only support 16,32,64 bits.");
4094 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4105 unsigned AndOp = SPIRV::OpBitwiseAndS;
4106 unsigned OrOp = SPIRV::OpBitwiseOrS;
4107 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
4108 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
4110 AndOp = SPIRV::OpBitwiseAndV;
4111 OrOp = SPIRV::OpBitwiseOrV;
4112 ShlOp = SPIRV::OpShiftLeftLogicalV;
4113 ShrOp = SPIRV::OpShiftRightLogicalV;
4119 const unsigned Shift) ->
Register {
4127 Register MaskReg = CreateConst(Mask);
4128 Register ShiftReg = CreateConst(Shift);
4135 if (!selectOpWithSrcs(
T1, ResType,
I, {Input, ShiftReg}, ShrOp) ||
4136 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
4137 !selectOpWithSrcs(T3, ResType,
I, {Input, MaskReg}, AndOp) ||
4138 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
4139 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
4147 uint64_t
Mask = ~0ull;
4148 while ((Shift >>= 1) > 0) {
4155 return BuildCOPY(ResVReg, Result,
I);
4158bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
4159 SPIRVTypeInst ResType,
4160 MachineInstr &
I)
const {
4161 assert(
I.getOperand(0).isReg() &&
I.getOperand(1).isReg() &&
4162 "G_FREEZE must define and use a register");
4163 Register OpReg =
I.getOperand(1).getReg();
4167 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
4180 if (MachineInstr *Def = MRI->
getVRegDef(OpReg)) {
4181 if (
Def->getOpcode() == TargetOpcode::COPY)
4184 switch (
Def->getOpcode()) {
4185 case SPIRV::ASSIGN_TYPE:
4186 if (MachineInstr *AssignToDef =
4188 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
4189 Reg =
Def->getOperand(2).getReg();
4192 case SPIRV::OpUndef:
4193 Reg =
Def->getOperand(1).getReg();
4196 unsigned DestOpCode;
4198 DestOpCode = SPIRV::OpConstantNull;
4199 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze of a "
4200 "static undef/poison lowered to OpConstantNull\n");
4202 DestOpCode = TargetOpcode::COPY;
4204 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze "
4205 "skipped, lowered as a copy of the operand\n");
4207 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
4208 .
addDef(
I.getOperand(0).getReg())
4216bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
4217 SPIRVTypeInst ResType,
4218 MachineInstr &
I)
const {
4220 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4222 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4226 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
4231 for (
unsigned i =
I.getNumExplicitDefs();
4232 i <
I.getNumExplicitOperands() && IsConst; ++i)
4236 if (!IsConst &&
N < 2)
4237 return diagnoseUnsupported(
4238 I,
"There must be at least two constituent operands in a vector");
4243 for (
unsigned i =
I.getNumExplicitDefs();
4244 i <
I.getNumExplicitOperands() && IsNullVector; ++i) {
4245 MachineInstr *
Def =
getDef(
I.getOperand(i), MRI);
4250 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4257 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4258 TII.get(IsConst ? SPIRV::OpConstantComposite
4259 : SPIRV::OpCompositeConstruct))
4262 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
4263 MIB.
addUse(
I.getOperand(i).getReg());
4268bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
4269 SPIRVTypeInst ResType,
4270 MachineInstr &
I)
const {
4272 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4274 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4280 if (!
I.getOperand(
OpIdx).isReg())
4287 if (!IsConst &&
N < 2)
4288 return diagnoseUnsupported(
4289 I,
"There must be at least two constituent operands in a vector");
4292 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4293 TII.get(IsConst ? SPIRV::OpConstantComposite
4294 : SPIRV::OpCompositeConstruct))
4297 for (
unsigned i = 0; i <
N; ++i)
4303bool SPIRVInstructionSelector::selectConcatVectors(
Register ResVReg,
4304 SPIRVTypeInst ResType,
4305 MachineInstr &
I)
const {
4309 if (ResType->
getOpcode() != SPIRV::OpTypeVector)
4311 "Cannot select G_CONCAT_VECTORS with a non-vector result");
4313 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4314 TII.get(SPIRV::OpCompositeConstruct))
4324bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
4325 SPIRVTypeInst ResType,
4326 MachineInstr &
I)
const {
4331 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
4333 Opcode = SPIRV::OpDemoteToHelperInvocation;
4335 Opcode = SPIRV::OpKill;
4337 if (MachineInstr *NextI =
I.getNextNode()) {
4339 NextI->eraseFromParent();
4349bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
4350 SPIRVTypeInst ResType,
unsigned CmpOpc,
4351 MachineInstr &
I)
const {
4352 Register Cmp0 =
I.getOperand(2).getReg();
4353 Register Cmp1 =
I.getOperand(3).getReg();
4356 "CMP operands should have the same type");
4357 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
4367bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
4368 SPIRVTypeInst ResType,
4369 MachineInstr &
I)
const {
4370 auto Pred =
I.getOperand(1).getPredicate();
4373 Register CmpOperand =
I.getOperand(2).getReg();
4378 Register Op1 =
I.getOperand(3).getReg();
4382 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
4387 I.getOperand(3).setReg(NewOp1);
4393 return selectCmp(ResVReg, ResType, CmpOpc,
I);
4397SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
4398 SPIRVTypeInst ResType)
const {
4400 SPIRVTypeInst SpvI32Ty =
4403 auto ConstInt = ConstantInt::get(LLVMTy, Val);
4410 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4413 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
4416 .
addImm(APInt(32, Val).getZExtValue());
4418 GR.
add(ConstInt,
MI);
4425Register SPIRVInstructionSelector::buildI32ConstantInEntryBlock(
4426 uint32_t Val, MachineInstr &
I, SPIRVTypeInst ResType)
const {
4428 SPIRVTypeInst SpvI32Ty =
4430 auto *ConstInt = ConstantInt::get(LLVMTy, Val);
4435 MachineBasicBlock &EntryBB = *InsertIt->getParent();
4436 MachineInstr *
MI =
nullptr;
4440 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantNull))
4444 uint64_t ImmVal = APInt(32, Val).getZExtValue();
4445 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantI))
4451 GR.
add(ConstInt,
MI);
4456bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
4457 SPIRVTypeInst ResType,
4458 MachineInstr &
I)
const {
4460 return selectCmp(ResVReg, ResType, CmpOp,
I);
4463bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
4464 SPIRVTypeInst ResType,
4465 MachineInstr &
I)
const {
4467 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
4474 if (ResType->
getOpcode() != SPIRV::OpTypeVector &&
4475 ResType->
getOpcode() != SPIRV::OpTypeFloat)
4478 MachineIRBuilder MIRBuilder(
I);
4485 APFloat ConstVal(3.3219280948873623);
4489 APFloat::rmNearestTiesToEven, &LosesInfo);
4493 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
4494 ? SPIRV::OpVectorTimesScalar
4497 if (!selectOpWithSrcs(ArgReg, ResType,
I,
4498 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
4500 if (!selectExtInst(ResVReg, ResType,
I,
4501 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
4511Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
4512 MachineInstr &
I)
const {
4515 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4520bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
4526 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
4534 if (
Def->getOpcode() == SPIRV::OpConstantNull)
4537 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
4538 Def->getOpcode() == SPIRV::OpConstantI)
4551 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
4552 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
4554 Intrinsic::spv_const_composite)) {
4555 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
4556 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
4557 if (!IsZero(
Def->getOperand(i).getReg()))
4566Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
4567 MachineInstr &
I)
const {
4571 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4576Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
4577 MachineInstr &
I)
const {
4581 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4587 SPIRVTypeInst ResType,
4588 MachineInstr &
I)
const {
4592 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4597bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
4598 SPIRVTypeInst ResType,
4599 MachineInstr &
I)
const {
4600 Register SelectFirstArg =
I.getOperand(2).getReg();
4601 Register SelectSecondArg =
I.getOperand(3).getReg();
4610 SPIRV::OpTypeVector;
4617 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
4618 }
else if (IsPtrTy) {
4619 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
4621 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
4624 assert(IsScalarBool &&
"OpSelect with a scalar result requires a scalar "
4625 "boolean condition");
4627 Opcode = SPIRV::OpSelectSFSCond;
4628 }
else if (IsPtrTy) {
4629 Opcode = SPIRV::OpSelectSPSCond;
4631 Opcode = SPIRV::OpSelectSISCond;
4634 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4637 .
addUse(
I.getOperand(1).getReg())
4646bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
4647 SPIRVTypeInst ResType,
4649 MachineInstr &InsertAt,
4650 bool IsSigned)
const {
4652 Register ZeroReg = buildZerosVal(ResType, InsertAt);
4653 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
4654 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
4656 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4668bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
4669 SPIRVTypeInst ResType,
4670 MachineInstr &
I,
bool IsSigned,
4671 unsigned Opcode)
const {
4672 Register SrcReg =
I.getOperand(1).getReg();
4678 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
4683 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I, IsSigned);
4685 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
4688bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
4689 SPIRVTypeInst ResType, MachineInstr &
I,
4690 bool IsSigned)
const {
4691 Register SrcReg =
I.getOperand(1).getReg();
4693 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
4697 if (ResType == SrcType)
4698 return BuildCOPY(ResVReg, SrcReg,
I);
4700 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4701 return selectUnOp(ResVReg, ResType,
I, Opcode);
4704bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
4705 SPIRVTypeInst ResType,
4707 bool IsSigned)
const {
4708 MachineIRBuilder MIRBuilder(
I);
4709 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
4721 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
4724 .
addUse(
I.getOperand(1).getReg())
4725 .
addUse(
I.getOperand(2).getReg())
4730 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
4733 .
addUse(
I.getOperand(1).getReg())
4734 .
addUse(
I.getOperand(2).getReg())
4742 unsigned SelectOpcode =
4743 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
4748 .
addUse(buildOnesVal(
true, ResType,
I))
4749 .
addUse(buildZerosVal(ResType,
I))
4756 .
addUse(buildOnesVal(
false, ResType,
I))
4761bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
4764 SPIRVTypeInst IntTy,
4765 SPIRVTypeInst BoolTy)
const {
4768 bool IsVectorTy = IntTy->
getOpcode() == SPIRV::OpTypeVector;
4769 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
4771 Register One = buildOnesVal(
false, IntTy,
I);
4779 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
4788bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
4789 SPIRVTypeInst ResType,
4790 MachineInstr &
I)
const {
4791 Register IntReg =
I.getOperand(1).getReg();
4794 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
4795 if (ArgType == ResType)
4796 return BuildCOPY(ResVReg, IntReg,
I);
4798 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4799 return selectUnOp(ResVReg, ResType,
I, Opcode);
4802bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
4803 SPIRVTypeInst ResType,
4804 MachineInstr &
I)
const {
4805 unsigned Opcode =
I.getOpcode();
4806 unsigned TpOpcode = ResType->
getOpcode();
4808 if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {
4809 assert(Opcode == TargetOpcode::G_CONSTANT &&
4810 I.getOperand(1).getCImm()->isZero());
4811 MachineBasicBlock &DepMBB =
I.getMF()->front();
4814 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
4821 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
4824bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
4825 SPIRVTypeInst ResType,
4826 MachineInstr &
I)
const {
4827 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4834bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
4835 SPIRVTypeInst ResType,
4836 MachineInstr &
I)
const {
4838 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
4842 .
addUse(
I.getOperand(3).getReg())
4844 .
addUse(
I.getOperand(2).getReg());
4845 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
4851bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
4852 SPIRVTypeInst ResType,
4853 MachineInstr &
I)
const {
4854 Type *MaybeResTy =
nullptr;
4859 "Expected aggregate type for extractv instruction");
4861 SPIRV::AccessQualifier::ReadWrite,
false);
4865 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
4868 .
addUse(
I.getOperand(2).getReg());
4869 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
4875bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
4876 SPIRVTypeInst ResType,
4877 MachineInstr &
I)
const {
4878 if (
getImm(
I.getOperand(4), MRI))
4879 return selectInsertVal(ResVReg, ResType,
I);
4881 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
4884 .
addUse(
I.getOperand(2).getReg())
4885 .
addUse(
I.getOperand(3).getReg())
4886 .
addUse(
I.getOperand(4).getReg())
4891bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
4892 SPIRVTypeInst ResType,
4893 MachineInstr &
I)
const {
4894 if (
getImm(
I.getOperand(3), MRI))
4895 return selectExtractVal(ResVReg, ResType,
I);
4897 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
4900 .
addUse(
I.getOperand(2).getReg())
4901 .
addUse(
I.getOperand(3).getReg())
4906bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
4907 SPIRVTypeInst ResType,
4908 MachineInstr &
I)
const {
4909 const bool IsGEPInBounds =
I.getOperand(2).getImm();
4915 ? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
4916 : SPIRV::OpAccessChain)
4917 : (IsGEPInBounds ?
SPIRV::OpInBoundsPtrAccessChain
4918 :
SPIRV::OpPtrAccessChain);
4920 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4924 .
addUse(
I.getOperand(3).getReg());
4926 (Opcode == SPIRV::OpPtrAccessChain ||
4927 Opcode == SPIRV::OpInBoundsPtrAccessChain ||
4928 (
getImm(
I.getOperand(4), MRI) &&
foldImm(
I.getOperand(4), MRI) == 0)) &&
4929 "Cannot translate GEP to OpAccessChain. First index must be 0.");
4932 const unsigned StartingIndex =
4933 (Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)
4936 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
4937 Res.addUse(
I.getOperand(i).getReg());
4938 Res.constrainAllUses(
TII,
TRI, RBI);
4943bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
4945 unsigned Lim =
I.getNumExplicitOperands();
4946 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
4947 Register OpReg =
I.getOperand(i).getReg();
4948 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
4950 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
4951 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
4952 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
4959 MachineFunction *MF =
I.getMF();
4965 SPIRVTypeInst WrapType = OpType;
4966 if (OpType->
getOpcode() == SPIRV::OpTypePointer &&
4968 SPIRV::StorageClass::CodeSectionINTEL) {
4970 SPIRV::StorageClass::Function,
I);
4977 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4978 TII.get(SPIRV::OpSpecConstantOp))
4981 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
4983 GR.
add(OpDefine, MIB);
4989bool SPIRVInstructionSelector::selectDerivativeInst(
4990 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
4991 const unsigned DPdOpCode)
const {
4994 if (!errorIfInstrOutsideShader(
I))
5000 Register SrcReg =
I.getOperand(2).getReg();
5005 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5008 .
addUse(
I.getOperand(2).getReg());
5010 MachineIRBuilder MIRBuilder(
I);
5013 if (componentCount != 1)
5021 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5026 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5031 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5039bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
5040 SPIRVTypeInst ResType,
5041 MachineInstr &
I)
const {
5045 case Intrinsic::spv_load:
5046 return selectLoad(ResVReg, ResType,
I);
5047 case Intrinsic::spv_atomic_load:
5048 return selectAtomicLoad(ResVReg, ResType,
I);
5049 case Intrinsic::spv_store:
5050 return selectStore(
I);
5051 case Intrinsic::spv_atomic_store:
5052 return selectAtomicStore(
I);
5053 case Intrinsic::spv_extractv:
5054 return selectExtractVal(ResVReg, ResType,
I);
5055 case Intrinsic::spv_insertv:
5056 return selectInsertVal(ResVReg, ResType,
I);
5057 case Intrinsic::spv_extractelt:
5058 return selectExtractElt(ResVReg, ResType,
I);
5059 case Intrinsic::spv_insertelt:
5060 return selectInsertElt(ResVReg, ResType,
I);
5061 case Intrinsic::spv_gep:
5062 return selectGEP(ResVReg, ResType,
I);
5063 case Intrinsic::spv_bitcast: {
5064 Register OpReg =
I.getOperand(2).getReg();
5065 SPIRVTypeInst OpType =
5069 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
5071 case Intrinsic::spv_unref_global:
5072 case Intrinsic::spv_init_global: {
5073 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
5078 Register GVarVReg =
MI->getOperand(0).getReg();
5079 if (!selectGlobalValue(GVarVReg, *
MI, Init))
5084 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
5086 MI->eraseFromParent();
5090 case Intrinsic::spv_undef: {
5091 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5097 case Intrinsic::spv_poison:
5098 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpPoisonKHR))
5103 case Intrinsic::spv_freeze:
5104 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
5107 .
addUse(
I.getOperand(2).getReg())
5110 case Intrinsic::spv_named_boolean_spec_constant: {
5111 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
5112 : SPIRV::OpSpecConstantFalse;
5114 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
5115 .
addDef(
I.getOperand(0).getReg())
5118 unsigned SpecId =
I.getOperand(2).getImm();
5120 SPIRV::Decoration::SpecId, {SpecId});
5124 case Intrinsic::spv_const_composite: {
5126 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
5132 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
5134 std::function<bool(
Register)> HasSpecConstOperand =
5144 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
5145 J < Def->getNumExplicitOperands(); ++J) {
5146 if (
Def->getOperand(J).isReg() &&
5147 HasSpecConstOperand(
Def->getOperand(J).getReg()))
5153 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
5154 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
5155 : SPIRV::OpConstantComposite;
5156 unsigned ContinuedOpc = HasSpecConst
5157 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
5158 : SPIRV::OpConstantCompositeContinuedINTEL;
5159 MachineIRBuilder MIR(
I);
5161 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
5163 for (
auto *Instr : Instructions) {
5164 Instr->setDebugLoc(
I.getDebugLoc());
5169 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5176 case Intrinsic::spv_assign_name: {
5177 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
5178 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
5179 for (
unsigned i =
I.getNumExplicitDefs() + 2;
5180 i <
I.getNumExplicitOperands(); ++i) {
5181 MIB.
addImm(
I.getOperand(i).getImm());
5186 case Intrinsic::spv_switch: {
5187 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
5188 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5189 if (
I.getOperand(i).isReg())
5190 MIB.
addReg(
I.getOperand(i).getReg());
5191 else if (
I.getOperand(i).isCImm())
5192 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
5193 else if (
I.getOperand(i).isMBB())
5194 MIB.
addMBB(
I.getOperand(i).getMBB());
5201 case Intrinsic::spv_loop_merge: {
5202 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
5203 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5204 if (
I.getOperand(i).isMBB())
5205 MIB.
addMBB(
I.getOperand(i).getMBB());
5212 case Intrinsic::spv_loop_control_intel: {
5214 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
5215 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
5220 case Intrinsic::spv_selection_merge: {
5222 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
5223 assert(
I.getOperand(1).isMBB() &&
5224 "operand 1 to spv_selection_merge must be a basic block");
5225 MIB.
addMBB(
I.getOperand(1).getMBB());
5226 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
5230 case Intrinsic::spv_cmpxchg:
5231 return selectAtomicCmpXchg(ResVReg, ResType,
I);
5232 case Intrinsic::spv_unreachable:
5233 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
5236 case Intrinsic::spv_abort:
5237 return selectAbort(
I);
5238 case Intrinsic::spv_alloca:
5239 return selectFrameIndex(ResVReg, ResType,
I);
5240 case Intrinsic::spv_alloca_array:
5241 return selectAllocaArray(ResVReg, ResType,
I);
5242 case Intrinsic::spv_assume:
5244 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
5245 .
addUse(
I.getOperand(1).getReg())
5250 case Intrinsic::spv_expect:
5252 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
5255 .
addUse(
I.getOperand(2).getReg())
5256 .
addUse(
I.getOperand(3).getReg())
5261 case Intrinsic::arithmetic_fence:
5262 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
5263 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
5266 .
addUse(
I.getOperand(2).getReg())
5270 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
5272 case Intrinsic::spv_thread_id:
5278 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
5280 case Intrinsic::spv_thread_id_in_group:
5286 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
5288 case Intrinsic::spv_group_id:
5294 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
5296 case Intrinsic::spv_flattened_thread_id_in_group:
5303 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
5305 case Intrinsic::spv_workgroup_size:
5306 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
5308 case Intrinsic::spv_global_size:
5309 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
5311 case Intrinsic::spv_global_offset:
5312 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
5314 case Intrinsic::spv_num_workgroups:
5315 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
5317 case Intrinsic::spv_subgroup_size:
5318 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
5320 case Intrinsic::spv_num_subgroups:
5321 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
5323 case Intrinsic::spv_subgroup_id:
5324 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
5325 case Intrinsic::spv_subgroup_local_invocation_id:
5326 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
5327 ResVReg, ResType,
I);
5328 case Intrinsic::spv_subgroup_max_size:
5329 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
5331 case Intrinsic::spv_fdot:
5332 return selectFloatDot(ResVReg, ResType,
I);
5333 case Intrinsic::spv_udot:
5334 case Intrinsic::spv_sdot:
5335 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5337 return selectIntegerDot(ResVReg, ResType,
I,
5338 IID == Intrinsic::spv_sdot);
5339 return selectIntegerDotExpansion(ResVReg, ResType,
I);
5340 case Intrinsic::spv_dot4add_i8packed:
5341 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5343 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
5344 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
5345 case Intrinsic::spv_dot4add_u8packed:
5346 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5348 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
5349 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
5350 case Intrinsic::spv_all:
5351 return selectAll(ResVReg, ResType,
I);
5352 case Intrinsic::spv_any:
5353 return selectAny(ResVReg, ResType,
I);
5354 case Intrinsic::spv_distance:
5355 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
5356 case Intrinsic::spv_lerp:
5357 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
5358 case Intrinsic::spv_length:
5359 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
5360 case Intrinsic::spv_degrees:
5361 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
5362 case Intrinsic::spv_faceforward:
5363 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
5364 case Intrinsic::spv_frac:
5365 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
5366 case Intrinsic::spv_isinf:
5367 return selectOpIsInf(ResVReg, ResType,
I);
5368 case Intrinsic::spv_isnan:
5369 return selectOpIsNan(ResVReg, ResType,
I);
5370 case Intrinsic::spv_isfinite:
5371 return selectOpIsFinite(ResVReg, ResType,
I);
5372 case Intrinsic::spv_isnormal:
5373 return selectOpIsNormal(ResVReg, ResType,
I);
5374 case Intrinsic::spv_normalize:
5375 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
5376 case Intrinsic::spv_refract:
5377 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
5378 case Intrinsic::spv_reflect:
5379 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
5380 case Intrinsic::spv_rsqrt:
5381 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
5382 case Intrinsic::spv_sign:
5383 return selectSign(ResVReg, ResType,
I);
5384 case Intrinsic::spv_smoothstep:
5385 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
5386 case Intrinsic::spv_firstbituhigh:
5387 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
5388 case Intrinsic::spv_firstbitshigh:
5389 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
5390 case Intrinsic::spv_firstbitlow:
5391 return selectFirstBitLow(ResVReg, ResType,
I);
5392 case Intrinsic::spv_all_memory_barrier:
5393 return selectBarrierInst(
I, SPIRV::Scope::Device,
5394 SPIRV::MemorySemantics::UniformMemory |
5395 SPIRV::MemorySemantics::ImageMemory |
5396 SPIRV::MemorySemantics::WorkgroupMemory,
5398 case Intrinsic::spv_all_memory_barrier_with_group_sync:
5399 return selectBarrierInst(
I, SPIRV::Scope::Device,
5400 SPIRV::MemorySemantics::UniformMemory |
5401 SPIRV::MemorySemantics::ImageMemory |
5402 SPIRV::MemorySemantics::WorkgroupMemory,
5404 case Intrinsic::spv_device_memory_barrier:
5405 return selectBarrierInst(
I, SPIRV::Scope::Device,
5406 SPIRV::MemorySemantics::UniformMemory |
5407 SPIRV::MemorySemantics::ImageMemory,
5409 case Intrinsic::spv_device_memory_barrier_with_group_sync:
5410 return selectBarrierInst(
I, SPIRV::Scope::Device,
5411 SPIRV::MemorySemantics::UniformMemory |
5412 SPIRV::MemorySemantics::ImageMemory,
5414 case Intrinsic::spv_group_memory_barrier:
5415 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5416 SPIRV::MemorySemantics::WorkgroupMemory,
5418 case Intrinsic::spv_group_memory_barrier_with_group_sync:
5419 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5420 SPIRV::MemorySemantics::WorkgroupMemory,
5422 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
5423 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
5424 SPIRV::StorageClass::StorageClass ResSC =
5427 return diagnoseUnsupported(
I,
"The target storage class is not castable "
5428 "from the Generic storage class");
5429 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
5437 case Intrinsic::spv_lifetime_start:
5438 case Intrinsic::spv_lifetime_end: {
5439 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
5440 : SPIRV::OpLifetimeStop;
5441 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
5442 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
5451 case Intrinsic::spv_saturate:
5452 return selectSaturate(ResVReg, ResType,
I);
5453 case Intrinsic::spv_nclamp:
5454 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
5455 case Intrinsic::spv_uclamp:
5456 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
5457 case Intrinsic::spv_sclamp:
5458 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
5459 case Intrinsic::spv_subgroup_prefix_bit_count:
5460 return selectWavePrefixBitCount(ResVReg, ResType,
I);
5461 case Intrinsic::spv_wave_active_countbits:
5462 return selectWaveActiveCountBits(ResVReg, ResType,
I);
5463 case Intrinsic::spv_wave_all_equal:
5464 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
5465 case Intrinsic::spv_wave_all:
5466 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
5467 case Intrinsic::spv_wave_any:
5468 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
5469 case Intrinsic::spv_subgroup_ballot:
5470 return selectWaveOpInst(ResVReg, ResType,
I,
5471 SPIRV::OpGroupNonUniformBallot);
5472 case Intrinsic::spv_wave_is_first_lane:
5473 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
5474 case Intrinsic::spv_wave_reduce_or:
5475 return selectWaveReduceOp(ResVReg, ResType,
I,
5476 SPIRV::OpGroupNonUniformBitwiseOr);
5477 case Intrinsic::spv_wave_reduce_xor:
5478 return selectWaveReduceOp(ResVReg, ResType,
I,
5479 SPIRV::OpGroupNonUniformBitwiseXor);
5480 case Intrinsic::spv_wave_reduce_and:
5481 return selectWaveReduceOp(ResVReg, ResType,
I,
5482 SPIRV::OpGroupNonUniformBitwiseAnd);
5483 case Intrinsic::spv_wave_reduce_umax:
5484 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
5485 case Intrinsic::spv_wave_reduce_max:
5486 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
5487 case Intrinsic::spv_wave_reduce_umin:
5488 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
5489 case Intrinsic::spv_wave_reduce_min:
5490 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
5491 case Intrinsic::spv_wave_reduce_sum:
5492 return selectWaveReduceSum(ResVReg, ResType,
I);
5493 case Intrinsic::spv_wave_product:
5494 return selectWaveReduceProduct(ResVReg, ResType,
I);
5495 case Intrinsic::spv_wave_readlane:
5496 return selectWaveOpInst(ResVReg, ResType,
I,
5497 SPIRV::OpGroupNonUniformShuffle);
5498 case Intrinsic::spv_wave_prefix_sum:
5499 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
5500 case Intrinsic::spv_wave_prefix_product:
5501 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
5502 case Intrinsic::spv_quad_read_across_x: {
5503 return selectQuadSwap(ResVReg, ResType,
I, 0);
5505 case Intrinsic::spv_quad_read_across_y: {
5506 return selectQuadSwap(ResVReg, ResType,
I, 1);
5508 case Intrinsic::spv_quad_read_across_diagonal: {
5509 return selectQuadSwap(ResVReg, ResType,
I, 2);
5511 case Intrinsic::spv_step:
5512 return selectExtInst(ResVReg, ResType,
I, CL::step, GL::Step);
5513 case Intrinsic::spv_radians:
5514 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
5518 case Intrinsic::instrprof_increment:
5519 case Intrinsic::instrprof_increment_step:
5520 case Intrinsic::instrprof_value_profile:
5523 case Intrinsic::spv_value_md:
5525 case Intrinsic::spv_resource_handlefrombinding: {
5526 return selectHandleFromBinding(ResVReg, ResType,
I);
5528 case Intrinsic::spv_resource_counterhandlefrombinding:
5529 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
5530 case Intrinsic::spv_resource_updatecounter:
5531 return selectUpdateCounter(ResVReg, ResType,
I);
5532 case Intrinsic::spv_resource_store_typedbuffer: {
5533 return selectImageWriteIntrinsic(
I);
5535 case Intrinsic::spv_resource_load_typedbuffer: {
5536 return selectReadImageIntrinsic(ResVReg, ResType,
I);
5538 case Intrinsic::spv_resource_load_level: {
5539 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
5541 case Intrinsic::spv_resource_getdimensions_x:
5542 case Intrinsic::spv_resource_getdimensions_xy:
5543 case Intrinsic::spv_resource_getdimensions_xyz: {
5544 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
5546 case Intrinsic::spv_resource_getdimensions_levels_x:
5547 case Intrinsic::spv_resource_getdimensions_levels_xy:
5548 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
5549 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
5551 case Intrinsic::spv_resource_getdimensions_ms_xy:
5552 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
5553 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
5555 case Intrinsic::spv_resource_calculate_lod:
5556 case Intrinsic::spv_resource_calculate_lod_unclamped:
5557 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
5558 case Intrinsic::spv_resource_sample:
5559 case Intrinsic::spv_resource_sample_clamp:
5560 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
5561 case Intrinsic::spv_resource_samplebias:
5562 case Intrinsic::spv_resource_samplebias_clamp:
5563 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
5564 case Intrinsic::spv_resource_samplegrad:
5565 case Intrinsic::spv_resource_samplegrad_clamp:
5566 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
5567 case Intrinsic::spv_resource_samplelevel:
5568 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
5569 case Intrinsic::spv_resource_samplecmp:
5570 case Intrinsic::spv_resource_samplecmp_clamp:
5571 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
5572 case Intrinsic::spv_resource_samplecmplevelzero:
5573 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
5574 case Intrinsic::spv_resource_gather:
5575 case Intrinsic::spv_resource_gather_cmp:
5576 return selectGatherIntrinsic(ResVReg, ResType,
I);
5577 case Intrinsic::spv_resource_getbasepointer:
5578 case Intrinsic::spv_resource_getpointer: {
5579 return selectResourceGetPointer(ResVReg, ResType,
I);
5581 case Intrinsic::spv_pushconstant_getpointer: {
5582 return selectPushConstantGetPointer(ResVReg, ResType,
I);
5584 case Intrinsic::spv_discard: {
5585 return selectDiscard(ResVReg, ResType,
I);
5587 case Intrinsic::spv_resource_nonuniformindex: {
5588 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
5590 case Intrinsic::spv_unpackhalf2x16: {
5591 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
5593 case Intrinsic::spv_packhalf2x16: {
5594 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
5596 case Intrinsic::spv_ddx:
5597 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
5598 case Intrinsic::spv_ddy:
5599 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
5600 case Intrinsic::spv_ddx_coarse:
5601 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
5602 case Intrinsic::spv_ddy_coarse:
5603 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
5604 case Intrinsic::spv_ddx_fine:
5605 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
5606 case Intrinsic::spv_ddy_fine:
5607 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
5608 case Intrinsic::spv_fwidth:
5609 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
5610 case Intrinsic::spv_masked_gather:
5611 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5612 return selectMaskedGather(ResVReg, ResType,
I);
5613 return diagnoseUnsupported(
5614 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
5615 case Intrinsic::spv_masked_scatter:
5616 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5617 return selectMaskedScatter(
I);
5618 return diagnoseUnsupported(
5619 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
5620 case Intrinsic::returnaddress:
5621 case Intrinsic::frameaddress: {
5623 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5630 return diagnoseUnsupported(
I,
"intrinsic selection not implemented.");
5635bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
5636 SPIRVTypeInst ResType,
5637 MachineInstr &
I)
const {
5640 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
5647bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
5648 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5650 assert(Intr.getIntrinsicID() ==
5651 Intrinsic::spv_resource_counterhandlefrombinding);
5654 Register MainHandleReg = Intr.getOperand(2).getReg();
5656 assert(MainHandleDef->getIntrinsicID() ==
5657 Intrinsic::spv_resource_handlefrombinding);
5661 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
5662 Register IndexReg = MainHandleDef->getOperand(5).getReg();
5663 std::string CounterName =
5668 MachineIRBuilder MIRBuilder(
I);
5670 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
5672 ArraySize, IndexReg, CounterName, MIRBuilder);
5674 return BuildCOPY(ResVReg, CounterVarReg,
I);
5677bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
5678 SPIRVTypeInst ResType,
5679 MachineInstr &
I)
const {
5681 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
5683 Register CounterHandleReg = Intr.getOperand(2).getReg();
5684 Register IncrReg = Intr.getOperand(3).getReg();
5691 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
5692 assert(CounterVarPointeeType &&
5693 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
5694 "Counter variable must be a struct");
5696 SPIRV::StorageClass::StorageBuffer &&
5697 "Counter variable must be in the storage buffer storage class");
5699 "Counter variable must have exactly 1 member in the struct");
5700 const SPIRVTypeInst MemberType =
5703 "Counter variable struct must have a single i32 member");
5707 MachineIRBuilder MIRBuilder(
I);
5709 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
5712 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
5718 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5721 .
addUse(CounterHandleReg)
5728 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
5731 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
5734 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
5743 return BuildCOPY(ResVReg, AtomicRes,
I);
5751 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
5759bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
5760 SPIRVTypeInst ResType,
5761 MachineInstr &
I)
const {
5769 Register ImageReg =
I.getOperand(2).getReg();
5777 Register IdxReg =
I.getOperand(3).getReg();
5779 MachineInstr &Pos =
I;
5781 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
5785bool SPIRVInstructionSelector::generateSampleImage(
5788 DebugLoc Loc, MachineInstr &Pos)
const {
5799 if (!loadHandleBeforePosition(NewSamplerReg,
5805 MachineIRBuilder MIRBuilder(Pos);
5818 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
5819 ImOps.Lod.has_value();
5820 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
5821 : SPIRV::OpImageSampleImplicitLod;
5823 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
5824 : SPIRV::OpImageSampleDrefImplicitLod;
5833 MIB.
addUse(*ImOps.Compare);
5835 uint32_t ImageOperands = 0;
5837 ImageOperands |= SPIRV::ImageOperand::Bias;
5839 ImageOperands |= SPIRV::ImageOperand::Lod;
5840 if (ImOps.GradX && ImOps.GradY)
5841 ImageOperands |= SPIRV::ImageOperand::Grad;
5842 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
5844 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
5847 "Non-constant offsets are not supported in sample instructions.");
5852 ImageOperands |= SPIRV::ImageOperand::MinLod;
5854 if (ImageOperands != 0) {
5855 MIB.
addImm(ImageOperands);
5856 if (ImageOperands & SPIRV::ImageOperand::Bias)
5858 if (ImageOperands & SPIRV::ImageOperand::Lod)
5860 if (ImageOperands & SPIRV::ImageOperand::Grad) {
5861 MIB.
addUse(*ImOps.GradX);
5862 MIB.
addUse(*ImOps.GradY);
5865 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
5866 MIB.
addUse(*ImOps.Offset);
5867 if (ImageOperands & SPIRV::ImageOperand::MinLod)
5868 MIB.
addUse(*ImOps.MinLod);
5875bool SPIRVInstructionSelector::selectImageQuerySize(
5877 std::optional<Register> LodReg)
const {
5879 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
5882 "ImageReg is not an image type.");
5884 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
5886 unsigned NumComponents = 0;
5888 case SPIRV::Dim::DIM_1D:
5889 case SPIRV::Dim::DIM_Buffer:
5890 NumComponents =
IsArray ? 2 : 1;
5892 case SPIRV::Dim::DIM_2D:
5893 case SPIRV::Dim::DIM_Cube:
5894 case SPIRV::Dim::DIM_Rect:
5895 NumComponents =
IsArray ? 3 : 2;
5897 case SPIRV::Dim::DIM_3D:
5901 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
5906 SPIRVTypeInst ResType =
5911 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5921bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
5922 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5923 Register ImageReg =
I.getOperand(2).getReg();
5930 return selectImageQuerySize(NewImageReg, ResVReg,
I);
5933bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
5934 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5935 Register ImageReg =
I.getOperand(2).getReg();
5944 Register LodReg =
I.getOperand(3).getReg();
5947 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
5949 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
5956 TII.get(SPIRV::OpImageQueryLevels))
5963 TII.get(SPIRV::OpCompositeConstruct))
5973bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
5974 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5975 Register ImageReg =
I.getOperand(2).getReg();
5986 "OpImageQuerySamples requires a multisampled image");
5988 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
5996 TII.get(SPIRV::OpImageQuerySamples))
6003 TII.get(SPIRV::OpCompositeConstruct))
6013bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
6014 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6015 Register ImageReg =
I.getOperand(2).getReg();
6016 Register SamplerReg =
I.getOperand(3).getReg();
6017 Register CoordinateReg =
I.getOperand(4).getReg();
6033 if (!loadHandleBeforePosition(
6038 MachineIRBuilder MIRBuilder(
I);
6044 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6054 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
6061 unsigned ExtractedIndex =
6063 Intrinsic::spv_resource_calculate_lod_unclamped
6067 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6068 TII.get(SPIRV::OpCompositeExtract))
6078bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
6079 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6080 Register ImageReg =
I.getOperand(2).getReg();
6081 Register SamplerReg =
I.getOperand(3).getReg();
6082 Register CoordinateReg =
I.getOperand(4).getReg();
6083 ImageOperands ImOps;
6084 if (
I.getNumOperands() > 5)
6085 ImOps.Offset =
I.getOperand(5).getReg();
6086 if (
I.getNumOperands() > 6)
6087 ImOps.MinLod =
I.getOperand(6).getReg();
6088 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6089 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6092bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
6093 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6094 Register ImageReg =
I.getOperand(2).getReg();
6095 Register SamplerReg =
I.getOperand(3).getReg();
6096 Register CoordinateReg =
I.getOperand(4).getReg();
6097 ImageOperands ImOps;
6098 ImOps.Bias =
I.getOperand(5).getReg();
6099 if (
I.getNumOperands() > 6)
6100 ImOps.Offset =
I.getOperand(6).getReg();
6101 if (
I.getNumOperands() > 7)
6102 ImOps.MinLod =
I.getOperand(7).getReg();
6103 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6104 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6107bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
6108 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6109 Register ImageReg =
I.getOperand(2).getReg();
6110 Register SamplerReg =
I.getOperand(3).getReg();
6111 Register CoordinateReg =
I.getOperand(4).getReg();
6112 ImageOperands ImOps;
6113 ImOps.GradX =
I.getOperand(5).getReg();
6114 ImOps.GradY =
I.getOperand(6).getReg();
6115 if (
I.getNumOperands() > 7)
6116 ImOps.Offset =
I.getOperand(7).getReg();
6117 if (
I.getNumOperands() > 8)
6118 ImOps.MinLod =
I.getOperand(8).getReg();
6119 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6120 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6123bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
6124 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6125 Register ImageReg =
I.getOperand(2).getReg();
6126 Register SamplerReg =
I.getOperand(3).getReg();
6127 Register CoordinateReg =
I.getOperand(4).getReg();
6128 ImageOperands ImOps;
6129 ImOps.Lod =
I.getOperand(5).getReg();
6130 if (
I.getNumOperands() > 6)
6131 ImOps.Offset =
I.getOperand(6).getReg();
6132 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6133 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6136bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
6137 SPIRVTypeInst ResType,
6138 MachineInstr &
I)
const {
6139 Register ImageReg =
I.getOperand(2).getReg();
6140 Register SamplerReg =
I.getOperand(3).getReg();
6141 Register CoordinateReg =
I.getOperand(4).getReg();
6142 ImageOperands ImOps;
6143 ImOps.Compare =
I.getOperand(5).getReg();
6144 if (
I.getNumOperands() > 6)
6145 ImOps.Offset =
I.getOperand(6).getReg();
6146 if (
I.getNumOperands() > 7)
6147 ImOps.MinLod =
I.getOperand(7).getReg();
6148 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6149 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6152bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
6153 SPIRVTypeInst ResType,
6154 MachineInstr &
I)
const {
6155 Register ImageReg =
I.getOperand(2).getReg();
6156 Register CoordinateReg =
I.getOperand(3).getReg();
6157 Register LodReg =
I.getOperand(4).getReg();
6159 ImageOperands ImOps;
6161 if (
I.getNumOperands() > 5)
6162 ImOps.Offset =
I.getOperand(5).getReg();
6174 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
6175 I.getDebugLoc(),
I, &ImOps);
6178bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
6179 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6180 Register ImageReg =
I.getOperand(2).getReg();
6181 Register SamplerReg =
I.getOperand(3).getReg();
6182 Register CoordinateReg =
I.getOperand(4).getReg();
6183 ImageOperands ImOps;
6184 ImOps.Compare =
I.getOperand(5).getReg();
6185 if (
I.getNumOperands() > 6)
6186 ImOps.Offset =
I.getOperand(6).getReg();
6189 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6190 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6193bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
6194 SPIRVTypeInst ResType,
6195 MachineInstr &
I)
const {
6196 Register ImageReg =
I.getOperand(2).getReg();
6197 Register SamplerReg =
I.getOperand(3).getReg();
6198 Register CoordinateReg =
I.getOperand(4).getReg();
6201 "ImageReg is not an image type.");
6206 ComponentOrCompareReg =
I.getOperand(5).getReg();
6207 OffsetReg =
I.getOperand(6).getReg();
6210 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
6214 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6215 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
6216 Dim != SPIRV::Dim::DIM_Rect) {
6218 "Gather operations are only supported for 2D, Cube, and Rect images.");
6225 if (!loadHandleBeforePosition(
6230 MachineIRBuilder MIRBuilder(
I);
6231 SPIRVTypeInst SampledImageType =
6236 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6244 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
6246 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
6248 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6253 .
addUse(ComponentOrCompareReg);
6255 uint32_t ImageOperands = 0;
6256 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
6257 if (Dim == SPIRV::Dim::DIM_Cube) {
6259 "Gather operations with offset are not supported for Cube images.");
6263 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6265 ImageOperands |= SPIRV::ImageOperand::Offset;
6269 if (ImageOperands != 0) {
6270 MIB.
addImm(ImageOperands);
6272 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6280bool SPIRVInstructionSelector::generateImageReadOrFetch(
6283 const ImageOperands *ImOps)
const {
6286 "ImageReg is not an image type.");
6288 bool IsSignedInteger =
6293 bool IsFetch = (SampledOp.getImm() == 1);
6295 auto AddOperands = [&](MachineInstrBuilder &MIB) {
6296 uint32_t ImageOperandsMask = 0;
6297 if (IsSignedInteger)
6298 ImageOperandsMask |= 0x1000;
6300 if (IsFetch && ImOps) {
6302 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
6303 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
6305 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
6307 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
6311 if (ImageOperandsMask != 0) {
6312 MIB.
addImm(ImageOperandsMask);
6313 if (IsFetch && ImOps) {
6316 if (ImOps->Offset &&
6317 (ImageOperandsMask &
6318 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
6319 MIB.
addUse(*ImOps->Offset);
6328 SPIRVTypeInst SampledType =
6331 SPIRVTypeInst ReadType =
6332 widenTypeToVec4(IsPacked ? SampledType : ResType, Pos);
6333 bool ReadTypeMatchesResult = ReadType == ResType;
6335 Register ReadReg = ReadTypeMatchesResult
6341 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
6347 BMI.constrainAllUses(
TII,
TRI, RBI);
6349 if (ReadTypeMatchesResult)
6362 if (ResultSize == 1) {
6371 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
6374bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
6375 SPIRVTypeInst ResType,
6376 MachineInstr &
I)
const {
6377 Register ResourcePtr =
I.getOperand(2).getReg();
6379 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
6388 MachineIRBuilder MIRBuilder(
I);
6393 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6399 if (
I.getNumExplicitOperands() > 3) {
6400 Register IndexReg =
I.getOperand(3).getReg();
6407bool SPIRVInstructionSelector::selectPushConstantGetPointer(
6408 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6413bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
6414 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6415 Register ObjReg =
I.getOperand(2).getReg();
6416 if (!BuildCOPY(ResVReg, ObjReg,
I))
6426 decorateUsesAsNonUniform(ResVReg);
6430void SPIRVInstructionSelector::decorateUsesAsNonUniform(
6433 {NonUniformReg,
nullptr}};
6434 llvm::SmallSet<Register, 8> Visited;
6435 while (WorkList.
size() > 0) {
6438 if (!Visited.
insert(CurrentReg).second)
6441 bool IsDecorated =
false;
6443 if (
Use.getOpcode() == SPIRV::OpDecorate &&
6444 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
6450 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
6452 if (ResultReg == CurrentReg)
6460 MachineInstr &InsertPt =
6463 SPIRV::Decoration::NonUniformEXT, {});
6468bool SPIRVInstructionSelector::extractSubvector(
6470 MachineInstr &InsertionPoint)
const {
6472 [[maybe_unused]] uint64_t InputSize =
6475 assert(InputSize > 1 &&
"The input must be a vector.");
6476 assert(ResultSize > 1 &&
"The result must be a vector.");
6477 assert(ResultSize < InputSize &&
6478 "Cannot extract more element than there are in the input.");
6482 for (uint64_t
I = 0;
I < ResultSize;
I++) {
6485 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6494 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
6496 TII.get(SPIRV::OpCompositeConstruct))
6500 for (
Register ComponentReg : ComponentRegisters)
6501 MIB.
addUse(ComponentReg);
6506bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
6507 MachineInstr &
I)
const {
6514 Register ImageReg =
I.getOperand(1).getReg();
6522 Register CoordinateReg =
I.getOperand(2).getReg();
6523 Register DataReg =
I.getOperand(3).getReg();
6526 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
6534Register SPIRVInstructionSelector::buildPointerToResource(
6535 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
6536 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
6537 StringRef Name, MachineIRBuilder MIRBuilder)
const {
6539 if (ArraySize == 1) {
6540 SPIRVTypeInst PtrType =
6543 "SpirvResType did not have an explicit layout.");
6548 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
6549 SPIRVTypeInst VarPointerType =
6552 VarPointerType, Set,
Binding, Name, MIRBuilder);
6554 SPIRVTypeInst ResPointerType =
6567bool SPIRVInstructionSelector::selectFirstBitSet16(
6568 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
6569 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
6571 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
6575 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
6578bool SPIRVInstructionSelector::selectFirstBitSet32(
6580 unsigned BitSetOpcode)
const {
6581 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6584 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6591bool SPIRVInstructionSelector::selectFirstBitSet64(
6593 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
6606 if (ComponentCount > 2) {
6607 auto Func = [
this, SwapPrimarySide](
Register ResVReg, SPIRVTypeInst ResType,
6609 unsigned Opcode) ->
bool {
6610 return this->selectFirstBitSet64(ResVReg, ResType,
I, SrcReg, Opcode,
6614 return handle64BitOverflow(ResVReg, ResType,
I, SrcReg, BitSetOpcode, Func);
6618 MachineIRBuilder MIRBuilder(
I);
6620 BaseType, 2 * ComponentCount, MIRBuilder,
false);
6624 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
6630 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
6637 bool IsScalarRes = ResType->
getOpcode() != SPIRV::OpTypeVector;
6640 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
6641 SPIRV::OpVectorExtractDynamic))
6643 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
6644 SPIRV::OpVectorExtractDynamic))
6648 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6649 TII.get(SPIRV::OpVectorShuffle))
6657 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
6663 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6664 TII.get(SPIRV::OpVectorShuffle))
6672 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
6692 SelectOp = SPIRV::OpSelectSISCond;
6693 AddOp = SPIRV::OpIAddS;
6701 SelectOp = SPIRV::OpSelectVIVCond;
6702 AddOp = SPIRV::OpIAddV;
6708 Register RegSecondaryOffset = Reg0;
6712 if (SwapPrimarySide) {
6713 PrimaryReg = LowReg;
6714 SecondaryReg = HighReg;
6715 RegPrimaryOffset = Reg0;
6716 RegSecondaryOffset = Reg32;
6721 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
6722 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
6727 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
6728 SPIRV::OpINotEqual))
6735 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
6736 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
6741 if (SwapPrimarySide) {
6743 if (!selectOpWithSrcs(RegAdd, ResType,
I,
6744 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
6755 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
6756 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
6761 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
6762 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
6765 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
6769bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
6770 SPIRVTypeInst ResType,
6772 bool IsSigned)
const {
6774 Register OpReg =
I.getOperand(2).getReg();
6777 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
6778 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
6782 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6784 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6786 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6789 return diagnoseUnsupported(
6791 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
6795bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
6796 SPIRVTypeInst ResType,
6797 MachineInstr &
I)
const {
6799 Register OpReg =
I.getOperand(2).getReg();
6804 unsigned ExtendOpcode = SPIRV::OpUConvert;
6805 unsigned BitSetOpcode = GL::FindILsb;
6809 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6811 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6813 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6816 return diagnoseUnsupported(
I,
6817 "spv_firstbitlow only supports 16,32,64 bits.");
6821bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
6822 SPIRVTypeInst ResType,
6823 MachineInstr &
I)
const {
6827 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
6830 .
addUse(
I.getOperand(2).getReg())
6833 unsigned Alignment =
I.getOperand(3).getImm();
6847 while (!Worklist.
empty()) {
6849 switch (
T->getOpcode()) {
6850 case SPIRV::OpTypeInt:
6851 case SPIRV::OpTypeFloat:
6852 case SPIRV::OpTypePointer:
6854 case SPIRV::OpTypeVector:
6855 case SPIRV::OpTypeMatrix:
6856 case SPIRV::OpTypeArray: {
6857 Register OperandReg =
T->getOperand(1).getReg();
6861 case SPIRV::OpTypeStruct:
6862 for (
unsigned Idx = 1,
E =
T->getNumOperands(); Idx <
E; ++Idx) {
6863 Register OperandReg =
T->getOperand(Idx).getReg();
6875bool SPIRVInstructionSelector::selectAbort(MachineInstr &
I)
const {
6876 assert(
I.getNumExplicitOperands() == 2);
6878 Register MsgReg =
I.getOperand(1).getReg();
6880 assert(MsgType &&
"Message argument of llvm.spv.abort has no SPIR-V type");
6883 return diagnoseUnsupported(
6885 "llvm.spv.abort message type must be a concrete SPIR-V type (numerical "
6886 "scalar, pointer, vector, matrix, or aggregate of such types)");
6889 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
6896bool SPIRVInstructionSelector::selectTrap(MachineInstr &
I)
const {
6905 uint32_t MsgVal = ~0
u;
6906 if (
I.getOpcode() == TargetOpcode::G_UBSANTRAP)
6907 MsgVal =
static_cast<uint32_t
>(
I.getOperand(0).
getImm());
6910 Register MsgReg = buildI32ConstantInEntryBlock(MsgVal,
I, MsgType);
6913 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
6920bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
6921 SPIRVTypeInst ResType,
6922 MachineInstr &
I)
const {
6926 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
6929 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
6932 unsigned Alignment =
I.getOperand(2).getImm();
6939bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
6944 const MachineInstr *PrevI =
I.getPrevNode();
6946 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
6950 .
addMBB(
I.getOperand(0).getMBB())
6955 .
addMBB(
I.getOperand(0).getMBB())
6960bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
6971 const MachineInstr *NextI =
I.getNextNode();
6973 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
6979 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
6981 .
addUse(
I.getOperand(0).getReg())
6982 .
addMBB(
I.getOperand(1).getMBB())
6988bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
6989 MachineInstr &
I)
const {
6991 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
6993 const unsigned NumOps =
I.getNumOperands();
6994 for (
unsigned i = 1; i <
NumOps; i += 2) {
6995 MIB.
addUse(
I.getOperand(i + 0).getReg());
6996 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
7002bool SPIRVInstructionSelector::selectGlobalValue(
7003 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
7005 MachineIRBuilder MIRBuilder(
I);
7006 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
7009 std::string GlobalIdent;
7011 unsigned &
ID = UnnamedGlobalIDs[GV];
7013 ID = UnnamedGlobalIDs.
size();
7014 GlobalIdent =
"__unnamed_" + Twine(ID).str();
7040 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
7047 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7052 MachineInstrBuilder MIB1 =
7053 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7056 MachineInstrBuilder MIB2 =
7058 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
7062 GR.
add(ConstVal, MIB2);
7070 MachineInstrBuilder MIB3 =
7071 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7074 GR.
add(ConstVal, MIB3);
7080 assert(NewReg != ResVReg);
7081 return BuildCOPY(ResVReg, NewReg,
I);
7091 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
7094 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
7100 SPIRVTypeInst ResType =
7104 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
7109 if (
GlobalVar->isExternallyInitialized() &&
7110 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
7111 constexpr unsigned ReadWriteINTEL = 3u;
7114 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
7120bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
7121 SPIRVTypeInst ResType,
7122 MachineInstr &
I)
const {
7124 return selectExtInst(ResVReg, ResType,
I, CL::log10);
7132 MachineIRBuilder MIRBuilder(
I);
7137 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7140 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
7142 .
add(
I.getOperand(1))
7147 ResType->
getOpcode() == SPIRV::OpTypeFloat);
7157 APFloat::rmNearestTiesToEven, &LosesInfo);
7161 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
7162 ? SPIRV::OpVectorTimesScalar
7173bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
7174 SPIRVTypeInst ResType,
7175 MachineInstr &
I)
const {
7178 return selectExtInst(ResVReg, ResType,
I, CL::pown);
7184 Register ExpReg =
I.getOperand(2).getReg();
7186 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
7187 SPIRV::OpConvertSToF))
7189 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
7196bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
7197 SPIRVTypeInst ResType,
7198 MachineInstr &
I)
const {
7214 MachineIRBuilder MIRBuilder(
I);
7215 SPIRVTypeInst FloatType =
7219 FloatType, MIRBuilder, SPIRV::StorageClass::Function);
7232 MachineBasicBlock &EntryBB =
I.getMF()->
front();
7234 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
7237 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7243 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7246 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
7249 .
add(
I.getOperand(
I.getNumExplicitDefs()))
7253 Register IntegralPartReg =
I.getOperand(1).getReg();
7256 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7266 assert(
false &&
"GLSL::Modf is deprecated.");
7277bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
7278 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7279 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7280 MachineIRBuilder MIRBuilder(
I);
7281 const SPIRVTypeInst Vec3Ty =
7284 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
7296 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7300 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7306 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7313 assert(
I.getOperand(2).isReg());
7314 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
7318 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
7329bool SPIRVInstructionSelector::loadBuiltinInputID(
7330 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7331 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7332 MachineIRBuilder MIRBuilder(
I);
7334 ResType, MIRBuilder, SPIRV::StorageClass::Input);
7349 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7353 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7362SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
7363 MachineInstr &
I)
const {
7364 MachineIRBuilder MIRBuilder(
I);
7365 if (
Type->getOpcode() != SPIRV::OpTypeVector)
7375bool SPIRVInstructionSelector::loadHandleBeforePosition(
7376 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
7377 MachineInstr &Pos)
const {
7380 Intrinsic::spv_resource_handlefrombinding);
7388 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
7389 MachineIRBuilder MIRBuilder(HandleDef);
7390 SPIRVTypeInst VarType = ResType;
7391 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
7393 if (IsStructuredBuffer) {
7398 if (ResType->
getOpcode() == SPIRV::OpTypeImage && ArraySize == 0)
7400 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
7403 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
7404 ArraySize, IndexReg, Name, MIRBuilder);
7408 uint32_t LoadOpcode =
7409 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
7419bool SPIRVInstructionSelector::errorIfInstrOutsideShader(
7420 MachineInstr &
I)
const {
7422 return diagnoseUnsupported(
7423 I,
"this instruction is only supported in shaders.");
7428InstructionSelector *
7432 return new SPIRVInstructionSelector(TM, Subtarget, RBI);
MachineInstrBuilder & UseMI
MachineInstrBuilder MachineInstrBuilder & DefMI
#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 uint8_t SwapBits(uint8_t Val)
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
Loop::LoopBounds::Direction Direction
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 bool isConcreteSPIRVType(SPIRVTypeInst Ty, const SPIRVGlobalRegistry &GR)
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)
static unsigned getPtrCmpOpcode(unsigned Pred)
bool isDead(const MachineInstr &MI, const MachineRegisterInfo &MRI)
static Register convertPtrToInt(Register Reg, LLT ConvTy, SPIRVTypeInst SpvType, LegalizerHelper &Helper, MachineRegisterInfo &MRI, SPIRVGlobalRegistry *GR)
BaseType
A given derived pointer can have multiple base pointers through phi/selects.
This file defines the SmallSet class.
static TableGen::Emitter::Opt Y("gen-skeleton-entry", EmitSkeleton, "Generate example skeleton entry")
static ManagedStatic< cl::opt< FnT >, OptCreatorT > CallbackFunction
static const fltSemantics & IEEEsingle()
static const fltSemantics & BFloat()
static const fltSemantics & IEEEdouble()
static const fltSemantics & IEEEhalf()
const fltSemantics & getSemantics() const
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.
Represent a constant reference to an array (0 or more elements consecutively in memory),...
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.
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.
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.
void constrainAllUses(const TargetInstrInfo &TII, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI) const
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
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 unsigned getNumExplicitOperands() const
Returns the number of non-implicit operands.
LLVM_ABI unsigned getNumExplicitDefs() const
Returns the number of non-implicit definitions.
LLVM_ABI void emitGenericError(const Twine &ErrMsg) const
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.
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,...
const TargetRegisterClass * getRegClass(Register Reg) const
Return the register class of the specified virtual register.
LLVM_ABI MachineInstr * getVRegDef(Register Reg) const
getVRegDef - Return the machine instr that defines the specified virtual register or null if none is ...
use_instr_iterator use_instr_begin(Register RegNo) const
bool use_nodbg_empty(Register RegNo) const
use_nodbg_empty - Return true if there are no non-Debug instructions using the specified register.
static def_instr_iterator def_instr_end()
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...
def_instr_iterator def_instr_begin(Register RegNo) const
LLT getType(Register Reg) const
Get the low-level type of Reg or LLT{} if Reg is not a generic (target independent) virtual register.
static use_instr_iterator use_instr_end()
iterator_range< use_instr_nodbg_iterator > use_nodbg_instructions(Register Reg) const
LLVM_ABI void setType(Register VReg, LLT Ty)
Set the low-level type of VReg to Ty.
const MachineFunction & getMF() const
LLVM_ABI void setRegClass(Register Reg, const TargetRegisterClass *RC)
setRegClass - Set the register class of the specified virtual register.
LLVM_ABI Register createGenericVirtualRegister(LLT Ty, StringRef Name="")
Create and return a new generic virtual register with low-level type Ty.
const TargetRegisterClass * getRegClassOrNull(Register Reg) const
Return the register class of Reg, or null if Reg has not been assigned a register class yet.
iterator_range< use_instr_iterator > use_instructions(Register Reg) const
unsigned getNumVirtRegs() const
getNumVirtRegs - Return the number of virtual registers created.
LLVM_ABI void replaceRegWith(Register FromReg, Register ToReg)
replaceRegWith - Replace all instances of FromReg with ToReg in the machine function.
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.
bool isScalarOrVectorSigned(SPIRVTypeInst Type) const
SPIRVTypeInst getOrCreateOpTypeSampledImage(SPIRVTypeInst ImageType, MachineIRBuilder &MIRBuilder)
void assignSPIRVTypeToVReg(SPIRVTypeInst Type, Register VReg, const MachineFunction &MF)
const TargetRegisterClass * getRegClass(SPIRVTypeInst SpvType) const
MachineInstr * getOrAddMemAliasingINTELInst(MachineIRBuilder &MIRBuilder, const MDNode *AliasingListMD)
bool isAggregateType(SPIRVTypeInst Type) const
unsigned getScalarOrVectorBitWidth(SPIRVTypeInst Type) const
SPIRVTypeInst getOrCreateSPIRVIntegerType(unsigned BitWidth, MachineIRBuilder &MIRBuilder)
SPIRVTypeInst getOrCreateSPIRVVectorType(SPIRVTypeInst BaseType, unsigned NumElements, MachineIRBuilder &MIRBuilder, bool EmitIR)
Register buildGlobalVariable(Register Reg, SPIRVTypeInst 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)
SPIRVTypeInst getResultType(Register VReg, MachineFunction *MF=nullptr)
unsigned getScalarOrVectorComponentCount(Register VReg) const
const Type * getTypeForSPIRVType(SPIRVTypeInst Ty) const
bool isBitcastCompatible(SPIRVTypeInst Type1, SPIRVTypeInst Type2) const
unsigned getPointerSize() const
Register getOrCreateConstFP(APFloat Val, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
LLT getRegType(SPIRVTypeInst SpvType) const
void invalidateMachineInstr(MachineInstr *MI)
SPIRVTypeInst getOrCreateSPIRVBoolType(MachineIRBuilder &MIRBuilder, bool EmitIR)
SPIRVTypeInst getOrCreateSPIRVPointerType(const Type *BaseType, MachineIRBuilder &MIRBuilder, SPIRV::StorageClass::StorageClass SC)
bool isScalarOfType(Register VReg, unsigned TypeOpcode) const
Register getSPIRVTypeID(SPIRVTypeInst SpirvType) const
Register getOrCreateConstInt(uint64_t Val, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
Register getOrCreateConstIntArray(uint64_t Val, size_t Num, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII)
bool findValueAttrs(const MachineInstr *Key, Type *&Ty, StringRef &Name)
SPIRVTypeInst retrieveScalarOrVectorIntType(SPIRVTypeInst Type) const
Register getOrCreateGlobalVariableWithBinding(SPIRVTypeInst VarType, uint32_t Set, uint32_t Binding, StringRef Name, MachineIRBuilder &MIRBuilder)
SPIRVTypeInst changePointerStorageClass(SPIRVTypeInst PtrType, SPIRV::StorageClass::StorageClass SC, MachineInstr &I)
Register getOrCreateConstVector(uint64_t Val, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
Register buildConstantFP(APFloat Val, MachineIRBuilder &MIRBuilder, SPIRVTypeInst SpvType=nullptr)
void addGlobalObject(const Value *V, const MachineFunction *MF, Register R)
SPIRVTypeInst getScalarOrVectorComponentType(SPIRVTypeInst Type) const
void recordFunctionPointer(const MachineOperand *MO, const Function *F)
SPIRVTypeInst getOrCreateSPIRVFloatType(unsigned BitWidth, MachineInstr &I, const SPIRVInstrInfo &TII)
SPIRVTypeInst getPointeeType(SPIRVTypeInst PtrType)
SPIRVTypeInst getOrCreateSPIRVType(const Type *Type, MachineInstr &I, SPIRV::AccessQualifier::AccessQualifier AQ, bool EmitIR)
bool isScalarOrVectorOfType(Register VReg, unsigned TypeOpcode) const
MachineFunction * setCurrentFunc(MachineFunction &MF)
Register getOrCreateConstNullPtr(MachineIRBuilder &MIRBuilder, SPIRVTypeInst SpvType)
SPIRVTypeInst getSPIRVTypeForVReg(Register VReg, const MachineFunction *MF=nullptr) const
Type * getDeducedGlobalValueType(const GlobalValue *Global)
Register getOrCreateUndef(MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII)
SPIRV::StorageClass::StorageClass getPointerStorageClass(Register VReg) 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
bool isTypeIntOrFloat() const
bool erase(PtrType Ptr)
Remove pointer from the set.
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.
std::pair< const_iterator, bool > insert(const T &V)
insert - Insert an element into the set if it isn't already there.
reference emplace_back(ArgTypes &&... Args)
void reserve(size_type N)
void push_back(const T &Elt)
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
constexpr size_t size() const
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.
Twine - A lightweight data structure for efficiently representing the concatenation of temporary valu...
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
@ BFloatTyID
16-bit floating point type (7-bit significand)
@ 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.
An efficient, type-erasing, non-owning reference to a callable.
self_iterator getIterator()
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.
constexpr std::underlying_type_t< E > Mask()
Get a bitmask with 1s in all places up to the high-order bit of E's largest value.
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
NodeAddr< FuncNode * > Func
BaseReg
Stack frame base register. Bit 0 of FREInfo.Info.
This is an optimization pass for GlobalISel generic memory operations.
@ Low
Lower the current thread's priority such that it does not affect foreground tasks significantly.
bool all_of(R &&range, UnaryPredicate P)
Provide wrappers to std::all_of which take ranges instead of having to pass begin/end explicitly.
void addStringImm(StringRef Str, MCInst &Inst)
MachineBasicBlock::iterator getOpVariableMBBIt(MachineFunction &MF)
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.
MachineInstr * getDef(const MachineOperand &MO, const MachineRegisterInfo *MRI)
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 void 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.
Register createVirtualRegister(SPIRVTypeInst SpvType, SPIRVGlobalRegistry *GR, MachineRegisterInfo *MRI, const MachineFunction &MF)
unsigned getArrayComponentCount(const MachineRegisterInfo *MRI, const MachineInstr *ResType)
void buildOpDecorate(Register Reg, MachineIRBuilder &MIRBuilder, SPIRV::Decoration::Decoration Dec, ArrayRef< uint32_t > DecArgs, StringRef StrImm)
LLVM_ABI bool isNullOrNullSplat(const MachineInstr &MI, const MachineRegisterInfo &MRI, bool AllowUndefs=false)
Return true if the value is a constant 0 integer or a splatted vector of a constant 0 integer (with n...
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)
RelativeUniformCounterPtr ValuesPtrExpr VTableAddr Value
void buildOpName(Register Target, StringRef Name, MachineIRBuilder &MIRBuilder)
bool any_of(R &&range, UnaryPredicate P)
Provide wrappers to std::any_of which take ranges instead of having to pass begin/end explicitly.
MachineInstr * getImm(const MachineOperand &MO, const MachineRegisterInfo *MRI)
constexpr bool isPowerOf2_32(uint32_t Value)
Return true if the argument is a power of two > 0.
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)
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
ArrayRef(const T &OneElt) -> ArrayRef< T >
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)
bool isSpvIntrinsic(const MachineInstr &MI, Intrinsic::ID IntrinsicID)
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...
MCRegisterClass TargetRegisterClass