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 (*II).getOpcode() == SPIRV::OpUntypedVariableKHR) &&
1303 getImm(
I.getOperand(2), MRI));
1305 bool IsGVInit =
false;
1309 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1310 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1311 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1312 (*UseIt).getOpcode() == SPIRV::OpVariable ||
1313 (*UseIt).getOpcode() == SPIRV::OpUntypedVariableKHR) {
1325 const bool UseUntypedPointers =
1326 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1327 if (UseUntypedPointers) {
1328 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1331 .
addImm(
static_cast<uint32_t
>(
1332 SPIRV::Opcode::UntypedInBoundsPtrAccessChainKHR))
1335 .
addUse(
I.getOperand(2).getReg())
1342 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1354 return diagnoseUnsupported(
1355 I,
"incompatible result and operand types in a bitcast");
1357 MachineInstrBuilder MIB =
1358 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1365 : SPIRV::OpInBoundsPtrAccessChain))
1369 .
addUse(
I.getOperand(2).getReg())
1372 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1376 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1378 .
addUse(
I.getOperand(2).getReg())
1387 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1390 .
addImm(
static_cast<uint32_t
>(
1391 SPIRV::Opcode::InBoundsPtrAccessChain))
1394 .
addUse(
I.getOperand(2).getReg());
1399 case TargetOpcode::G_ATOMICRMW_OR:
1400 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1401 case TargetOpcode::G_ATOMICRMW_ADD:
1402 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1403 case TargetOpcode::G_ATOMICRMW_AND:
1404 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1405 case TargetOpcode::G_ATOMICRMW_MAX:
1406 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1407 case TargetOpcode::G_ATOMICRMW_MIN:
1408 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1409 case TargetOpcode::G_ATOMICRMW_SUB:
1410 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1411 case TargetOpcode::G_ATOMICRMW_XOR:
1412 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1413 case TargetOpcode::G_ATOMICRMW_UMAX:
1414 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1415 case TargetOpcode::G_ATOMICRMW_UMIN:
1416 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1417 case TargetOpcode::G_ATOMICRMW_XCHG:
1418 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1420 case TargetOpcode::G_ATOMICRMW_FADD:
1421 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1422 case TargetOpcode::G_ATOMICRMW_FSUB:
1424 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1425 ResType->
getOpcode() == SPIRV::OpTypeVector
1427 : SPIRV::OpFNegate);
1428 case TargetOpcode::G_ATOMICRMW_FMIN:
1429 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1430 case TargetOpcode::G_ATOMICRMW_FMAX:
1431 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1433 case TargetOpcode::G_FENCE:
1434 return selectFence(
I);
1436 case TargetOpcode::G_STACKSAVE:
1437 return selectStackSave(ResVReg, ResType,
I);
1438 case TargetOpcode::G_STACKRESTORE:
1439 return selectStackRestore(
I);
1441 case TargetOpcode::G_UNMERGE_VALUES:
1444 case TargetOpcode::G_TRAP:
1445 case TargetOpcode::G_UBSANTRAP:
1446 return selectTrap(
I);
1451 case TargetOpcode::DBG_LABEL:
1453 case TargetOpcode::G_DEBUGTRAP:
1454 return selectDebugTrap(ResVReg, ResType,
I);
1461bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1462 SPIRVTypeInst ResType,
1463 MachineInstr &
I)
const {
1464 unsigned Opcode = SPIRV::OpNop;
1471bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1472 SPIRVTypeInst ResType,
1474 GL::GLSLExtInst GLInst,
1475 bool setMIFlags,
bool useMISrc,
1478 SPIRV::InstructionSet::InstructionSet::GLSL_std_450))
1479 return diagnoseUnsupported(
1481 "this instruction is only supported with the GLSL extended instruction "
1483 return selectExtInst(ResVReg, ResType,
I,
1484 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1485 setMIFlags, useMISrc, SrcRegs);
1488bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1489 SPIRVTypeInst ResType,
1491 CL::OpenCLExtInst CLInst,
1492 bool setMIFlags,
bool useMISrc,
1494 return selectExtInst(ResVReg, ResType,
I,
1495 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1496 setMIFlags, useMISrc, SrcRegs);
1499bool SPIRVInstructionSelector::selectExtInst(
1500 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1501 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1503 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1504 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1505 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1509bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1510 SPIRVTypeInst ResType,
1513 bool setMIFlags,
bool useMISrc,
1516 for (
const auto &[InstructionSet, Opcode] : Insts) {
1520 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1523 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1528 const unsigned NumOps =
I.getNumOperands();
1531 I.getOperand(Index).getType() ==
1532 MachineOperand::MachineOperandType::MO_IntrinsicID)
1535 MIB.
add(
I.getOperand(Index));
1547bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1548 SPIRVTypeInst ResType,
1549 MachineInstr &
I)
const {
1550 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1551 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1552 for (
const auto &Ex : ExtInsts) {
1553 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1554 uint32_t Opcode = Ex.second;
1558 MachineIRBuilder MIRBuilder(
I);
1561 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1568 const bool IsUntyped =
1569 PointerType->getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1571 BuildMI(*It->getParent(), It, It->getDebugLoc(),
1572 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
1573 : SPIRV::OpVariable))
1576 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
1582 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1585 .
addImm(
static_cast<uint32_t
>(Ex.first))
1587 .
add(
I.getOperand(2))
1591 Register ExpResReg =
I.getOperand(1).getReg();
1593 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1603bool SPIRVInstructionSelector::selectLdexp(
Register ResVReg,
1604 SPIRVTypeInst ResType,
1605 MachineInstr &
I)
const {
1606 Register XReg =
I.getOperand(1).getReg();
1607 Register ExpReg =
I.getOperand(2).getReg();
1613 if (ResType->
getOpcode() == SPIRV::OpTypeVector &&
1614 ExpType->
getOpcode() != SPIRV::OpTypeVector) {
1616 SPIRVTypeInst ExpVecType =
1620 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1621 TII.get(SPIRV::OpCompositeConstruct))
1624 for (
unsigned J = 0; J < NumElts; ++J)
1630 return selectExtInst(ResVReg, ResType,
I, CL::ldexp, GL::Ldexp,
1631 true,
false, {XReg, ExpReg});
1634bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1635 SPIRVTypeInst ResType,
1636 MachineInstr &
I)
const {
1637 Register CosResVReg =
I.getOperand(1).getReg();
1638 unsigned SrcIdx =
I.getNumExplicitDefs();
1643 MachineIRBuilder MIRBuilder(
I);
1645 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1652 const bool IsUntyped =
1653 PointerType->getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1655 BuildMI(*It->getParent(), It, It->getDebugLoc(),
1656 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
1657 : SPIRV::OpVariable))
1660 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
1664 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1667 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1669 .
add(
I.getOperand(SrcIdx))
1672 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1680 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1683 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1685 .
add(
I.getOperand(SrcIdx))
1687 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1690 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1692 .
add(
I.getOperand(SrcIdx))
1699bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1700 SPIRVTypeInst ResType,
1703 unsigned Opcode)
const {
1704 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1714std::optional<SplitParts> SPIRVInstructionSelector::splitEvenOddLanes(
1715 Register PopCountReg,
unsigned ComponentCount, MachineInstr &
I,
1716 SPIRVTypeInst I32Type)
const {
1719 if (ComponentCount == 1) {
1722 Parts.IsScalar =
true;
1723 Parts.Type = I32Type;
1731 if (!selectOpWithSrcs(Parts.High, I32Type,
I, {PopCountReg, IdxOne},
1732 SPIRV::OpVectorExtractDynamic))
1733 return std::nullopt;
1735 if (!selectOpWithSrcs(Parts.Low, I32Type,
I, {PopCountReg, IdxZero},
1736 SPIRV::OpVectorExtractDynamic))
1737 return std::nullopt;
1741 MachineIRBuilder MIRBuilder(
I);
1742 Parts.IsScalar =
false;
1749 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1750 TII.get(SPIRV::OpVectorShuffle))
1755 for (
unsigned J = 1; J < ComponentCount * 2; J += 2)
1760 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1761 TII.get(SPIRV::OpVectorShuffle))
1766 for (
unsigned J = 0; J < ComponentCount * 2; J += 2)
1774bool SPIRVInstructionSelector::selectPopCount16(
Register ResVReg,
1775 SPIRVTypeInst ResType,
1778 unsigned Opcode)
const {
1779 Register OpReg =
I.getOperand(1).getReg();
1782 MachineIRBuilder MIRBuilder(
I);
1784 SPIRVTypeInst I32VectorType =
1787 bool IsVector = NumElems > 1;
1788 SPIRVTypeInst ExtType = IsVector ? I32VectorType : I32Type;
1791 if (!selectOpWithSrcs(ExtReg, ExtType,
I, {OpReg}, SPIRV::OpUConvert))
1795 if (!selectPopCount32(PopCountReg, ExtType,
I, ExtReg, Opcode))
1798 return selectOpWithSrcs(ResVReg, ResType,
I, {PopCountReg}, ExtOpcode);
1801bool SPIRVInstructionSelector::selectPopCount32(
Register ResVReg,
1802 SPIRVTypeInst ResType,
1805 unsigned Opcode)
const {
1806 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
1809bool SPIRVInstructionSelector::selectPopCount64(
Register ResVReg,
1810 SPIRVTypeInst ResType,
1813 unsigned Opcode)
const {
1815 if (ComponentCount > 2)
1816 return handle64BitOverflow(
1817 ResVReg, ResType,
I, SrcReg, Opcode,
1819 unsigned O) {
return this->selectPopCount64(R,
T,
I, S, O); });
1821 MachineIRBuilder MIRBuilder(
I);
1826 I32Type, 2 * ComponentCount, MIRBuilder,
false);
1830 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
1835 if (!selectPopCount32(Pop32, VecI32Type,
I, Vec32, Opcode))
1839 auto MaybeParts = splitEvenOddLanes(Pop32, ComponentCount,
I, I32Type);
1842 SplitParts &Parts = *MaybeParts;
1845 unsigned OpAdd = Parts.IsScalar ? SPIRV::OpIAddS : SPIRV::OpIAddV;
1847 if (!selectOpWithSrcs(Sum, Parts.Type,
I, {Parts.High, Parts.Low}, OpAdd))
1852 unsigned ConvOp = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
1853 return selectOpWithSrcs(ResVReg, ResType,
I, {Sum}, ConvOp);
1856bool SPIRVInstructionSelector::selectPopCount(
Register ResVReg,
1857 SPIRVTypeInst ResType,
1859 unsigned Opcode)
const {
1864 if (!STI.getTargetTriple().isVulkanOS())
1865 return selectUnOp(ResVReg, ResType,
I, Opcode);
1867 Register OpReg =
I.getOperand(1).getReg();
1870 : SPIRV::OpUConvert;
1874 return selectPopCount16(ResVReg, ResType,
I, ExtOpcode, Opcode);
1876 return selectPopCount32(ResVReg, ResType,
I, OpReg, Opcode);
1878 return selectPopCount64(ResVReg, ResType,
I, OpReg, Opcode);
1880 return diagnoseUnsupported(
I,
"unsupported operand bit width for popcount");
1884bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1885 SPIRVTypeInst ResType,
1887 unsigned Opcode)
const {
1889 Register SrcReg =
I.getOperand(1).getReg();
1894 unsigned DefOpCode = DefIt->getOpcode();
1895 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1898 if (
auto *VRD =
getVRegDef(*MRI, DefIt->getOperand(1).getReg()))
1899 DefOpCode = VRD->getOpcode();
1901 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1902 DefOpCode == TargetOpcode::G_CONSTANT ||
1903 DefOpCode == SPIRV::OpVariable ||
1904 DefOpCode == SPIRV::OpUntypedVariableKHR ||
1905 DefOpCode == SPIRV::OpConstantI) {
1911 uint32_t SpecOpcode = 0;
1913 case SPIRV::OpConvertPtrToU:
1914 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1916 case SPIRV::OpConvertUToPtr:
1917 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1922 TII.get(SPIRV::OpSpecConstantOp))
1932 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1936bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1937 SPIRVTypeInst ResType,
1938 MachineInstr &
I)
const {
1939 Register OpReg =
I.getOperand(1).getReg();
1940 SPIRVTypeInst OpType =
1943 return diagnoseUnsupported(
1944 I,
"incompatible result and operand types in a bitcast");
1945 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1956 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1957 if (
MemOp->isNonTemporal())
1958 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1960 if (!ST->isShader() &&
MemOp->getAlign().value())
1961 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1965 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1966 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1970 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1972 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1976 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1980 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1982 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1994 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1996 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1998 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
2002bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
2003 SPIRVTypeInst ResType,
2004 MachineInstr &
I)
const {
2006 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2011 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2012 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2014 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2016 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2020 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2024 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2025 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
2026 I.getDebugLoc(),
I);
2030 MachineIRBuilder MIRBuilder(
I);
2032 if (
I.getNumMemOperands()) {
2033 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2034 if (MemOp->isAtomic())
2035 return selectAtomicLoad(ResVReg, ResType,
I);
2038 auto MIB = MIRBuilder.
buildInstr(SPIRV::OpLoad)
2042 if (!
I.getNumMemOperands()) {
2043 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2045 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2054Register SPIRVInstructionSelector::createPtrSizedIntReg(
2055 MachineIRBuilder &MIRBuilder)
const {
2056 SPIRVTypeInst IntType =
2066SPIRVInstructionSelector::convertPtrToInt(
Register PtrVal,
2067 MachineIRBuilder &MIRBuilder)
const {
2068 SPIRVTypeInst IntType =
2070 Register IntReg = createPtrSizedIntReg(MIRBuilder);
2071 MIRBuilder.
buildInstr(SPIRV::OpConvertPtrToU)
2079Register SPIRVInstructionSelector::castPtrToPtrToInt(
2080 Register Ptr, SPIRV::StorageClass::StorageClass SC,
2081 MachineIRBuilder &MIRBuilder)
const {
2082 SPIRVTypeInst IntType =
2084 SPIRVTypeInst PtrType =
2098bool SPIRVInstructionSelector::selectAtomicPtrValue(
2099 Register ResVReg, SPIRVTypeInst ResType, MachineIRBuilder &MIRBuilder,
2100 function_ref<
Register(SPIRVTypeInst IntType)> EmitAtomic)
const {
2108 Register IntResult = EmitAtomic(IntType);
2110 MIRBuilder.
buildInstr(SPIRV::OpConvertUToPtr)
2118bool SPIRVInstructionSelector::selectAtomicLoad(
Register ResVReg,
2119 SPIRVTypeInst ResType,
2120 MachineInstr &
I)
const {
2121 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2124 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2127 return diagnoseUnsupported(
2128 I,
"Lowering to SPIR-V of atomic load is only "
2129 "allowed for integer, floating point or pointer types");
2131 assert(
I.getNumMemOperands());
2132 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2133 assert(MemOp.isAtomic());
2135 uint32_t
Scope =
static_cast<uint32_t
>(
2137 Register ScopeReg = buildI32Constant(Scope,
I);
2143 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2144 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2147 MachineIRBuilder MIRBuilder(
I);
2151 return diagnoseUnsupported(
2152 I,
"Lowering to SPIR-V of atomic load is only "
2153 "allowed for pointer types for physical addressing model");
2158 SPIRV::StorageClass::StorageClass SC =
2160 return selectAtomicPtrValue(
2161 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2162 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2163 Register IntResult = createPtrSizedIntReg(MIRBuilder);
2174 auto AtomicLoad = MIRBuilder.
buildInstr(SPIRV::OpAtomicLoad)
2185bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
2187 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2188 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2193 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2194 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2196 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2201 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2205 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2206 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2207 SPIRVTypeInst SampledType =
2209 SPIRVTypeInst StoreValCompType =
2211 if (StoreValCompType && StoreValCompType != SampledType) {
2214 SPIRVTypeInst PackedType = widenTypeToVec4(SampledType,
I);
2217 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
2222 StoreVal = PackedReg;
2225 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2226 TII.get(SPIRV::OpImageWrite))
2232 if (sampledTypeIsSignedInteger(LLVMHandleType))
2235 BMI.constrainAllUses(
TII,
TRI, RBI);
2240 if (
I.getNumMemOperands()) {
2241 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2242 if (MemOp->isAtomic())
2243 return selectAtomicStore(
I);
2250 if (PtrSC == SPIRV::StorageClass::UniformConstant ||
2251 PtrSC == SPIRV::StorageClass::Input ||
2252 PtrSC == SPIRV::StorageClass::PushConstant)
2253 return diagnoseUnsupported(
2254 I,
"store into a read-only SPIR-V storage class is not allowed");
2256 MachineIRBuilder MIRBuilder(
I);
2258 if (!
I.getNumMemOperands()) {
2259 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2261 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2270bool SPIRVInstructionSelector::selectAtomicStore(MachineInstr &
I)
const {
2271 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2274 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2275 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2280 if (!PointeeType && PtrType &&
2281 PtrType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR)
2284 return diagnoseUnsupported(
I,
2285 "Lowering to SPIR-V of atomic store is only "
2286 "allowed for integer or floating point types");
2288 assert(
I.getNumMemOperands());
2289 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2290 assert(MemOp.isAtomic());
2292 uint32_t
Scope =
static_cast<uint32_t
>(
2294 Register ScopeReg = buildI32Constant(Scope,
I);
2300 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2301 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2303 MachineIRBuilder MIRBuilder(
I);
2307 return diagnoseUnsupported(
2308 I,
"Lowering to SPIR-V of atomic store is only "
2309 "allowed for pointer types for physical addressing model");
2314 SPIRV::StorageClass::StorageClass SC =
2316 return selectAtomicPtrValue(
2317 Register(), SPIRVTypeInst(), MIRBuilder, [&](SPIRVTypeInst IntType) {
2319 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2332 return diagnoseUnsupported(
I,
2333 "Lowering to SPIR-V of atomic store is only "
2334 "allowed for integer or floating point types");
2336 auto AtomicStore = MIRBuilder.
buildInstr(SPIRV::OpAtomicStore)
2346bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
2347 SPIRVTypeInst ResType,
2348 MachineInstr &
I)
const {
2349 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
2357 const Register PtrsReg =
I.getOperand(2).getReg();
2358 const uint32_t Alignment =
I.getOperand(3).getImm();
2359 const Register MaskReg =
I.getOperand(4).getReg();
2360 const Register PassthruReg =
I.getOperand(5).getReg();
2361 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2365 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
2376bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
2377 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
2384 const Register ValuesReg =
I.getOperand(1).getReg();
2385 const Register PtrsReg =
I.getOperand(2).getReg();
2386 const uint32_t Alignment =
I.getOperand(3).getImm();
2387 const Register MaskReg =
I.getOperand(4).getReg();
2388 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2392 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
2401bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
2402 const Twine &
Msg)
const {
2403 const Function &
F =
I.getMF()->getFunction();
2404 F.getContext().diagnose(
2405 DiagnosticInfoUnsupported(
F,
Msg,
I.getDebugLoc(),
DS_Error));
2409bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
2410 SPIRVTypeInst ResType,
2411 MachineInstr &
I)
const {
2412 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2413 return diagnoseUnsupported(
2414 I,
"llvm.stacksave intrinsic: this instruction requires the following "
2415 "SPIR-V extension: SPV_INTEL_variable_length_array");
2417 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
2424bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
2425 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2426 return diagnoseUnsupported(
2428 "llvm.stackrestore intrinsic: this instruction requires the following "
2429 "SPIR-V extension: SPV_INTEL_variable_length_array");
2430 if (!
I.getOperand(0).isReg())
2433 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
2434 .
addUse(
I.getOperand(0).getReg())
2440SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
2441 MachineIRBuilder MIRBuilder(
I);
2442 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
2449 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
2453 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2454 Type *ArrTy = ArrayType::get(ValTy, Num);
2456 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
2459 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
2470 const bool IsUntyped = VarTy->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
2471 auto MIBVar =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2472 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
2473 : SPIRV::OpVariable))
2476 .
addImm(SPIRV::StorageClass::UniformConstant);
2489bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
2492 Register DstReg =
I.getOperand(0).getReg();
2496 return diagnoseUnsupported(
2497 I,
"OpCopyMemory requires operands to have the same type");
2498 uint64_t CopySize =
getIConstVal(
I.getOperand(2).getReg(), MRI);
2502 return diagnoseUnsupported(
2503 I,
"Unable to determine pointee type size for OpCopyMemory");
2504 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
2505 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
2506 return diagnoseUnsupported(
2507 I,
"OpCopyMemory requires the size to match the pointee type size");
2508 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemory))
2511 if (
I.getNumMemOperands()) {
2512 MachineIRBuilder MIRBuilder(
I);
2519bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &
I,
2522 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemorySized))
2523 .
addUse(
I.getOperand(0).getReg())
2525 .
addUse(
I.getOperand(2).getReg());
2526 if (
I.getNumMemOperands()) {
2527 MachineIRBuilder MIRBuilder(
I);
2534bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
2535 MachineInstr &
I)
const {
2537 Register SizeReg =
I.getOperand(2).getReg();
2539 SizeDef && SizeDef->
getOpcode() == TargetOpcode::G_CONSTANT &&
2543 Register SrcReg =
I.getOperand(1).getReg();
2544 if (
I.getOpcode() == TargetOpcode::G_MEMSET ||
2545 I.getOpcode() == TargetOpcode::G_MEMSET_INLINE) {
2546 Register VarReg = getOrCreateMemSetGlobal(
I);
2549 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2551 ValTy,
I, SPIRV::StorageClass::UniformConstant);
2553 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
2557 if (!selectCopyMemory(
I, SrcReg))
2560 if (!selectCopyMemorySized(
I, SrcReg))
2563 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
2564 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
2569bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
2570 SPIRVTypeInst ResType,
2573 unsigned NegateOpcode)
const {
2575 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2576 uint32_t
Scope =
static_cast<uint32_t
>(
2578 MemOp->getSyncScopeID()));
2579 Register ScopeReg = buildI32Constant(Scope,
I);
2581 Register Ptr =
I.getOperand(1).getReg();
2582 uint32_t ScSem =
static_cast<uint32_t
>(
2586 Register MemSemReg = buildI32Constant(MemSem,
I);
2588 Register ValueReg =
I.getOperand(2).getReg();
2589 if (NegateOpcode != 0) {
2592 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
2598 if (NewOpcode != SPIRV::OpAtomicExchange)
2599 return diagnoseUnsupported(
2600 I,
"Lowering to SPIR-V of this atomic operation is not "
2601 "allowed for pointer types");
2603 return diagnoseUnsupported(
2604 I,
"Lowering to SPIR-V of atomic exchange is only "
2605 "allowed for pointer types for physical addressing model");
2612 MachineIRBuilder MIRBuilder(
I);
2614 return selectAtomicPtrValue(
2615 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2617 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2618 Register ExchangeResReg = createPtrSizedIntReg(MIRBuilder);
2619 MIRBuilder.
buildInstr(SPIRV::OpAtomicExchange)
2627 return ExchangeResReg;
2631 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2642bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2643 unsigned ArgI =
I.getNumOperands() - 1;
2645 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2646 SPIRVTypeInst SrcType =
2648 if (!SrcType || SrcType->
getOpcode() != SPIRV::OpTypeVector)
2650 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2654 unsigned CurrentIndex = 0;
2655 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2656 Register ResVReg =
I.getOperand(i).getReg();
2659 LLT ResLLT = MRI->
getType(ResVReg);
2665 ResType = ScalarType;
2671 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
2674 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2680 for (
unsigned j = 0;
j < NumElements; ++
j) {
2681 MIB.
addImm(CurrentIndex + j);
2683 CurrentIndex += NumElements;
2687 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2699bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2702 Register MemSemReg = buildI32Constant(MemSem,
I);
2706 Register ScopeReg = buildI32Constant(Scope,
I);
2708 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2715bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2716 SPIRVTypeInst ResType,
2718 unsigned Opcode)
const {
2719 Type *ResTy =
nullptr;
2722 return diagnoseUnsupported(
2724 "Not enough info to select the arithmetic with overflow instruction");
2726 return diagnoseUnsupported(
I,
2727 "Expect struct type result for the arithmetic "
2728 "with overflow instruction");
2734 MachineIRBuilder MIRBuilder(
I);
2736 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2737 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2743 Register ZeroReg = buildZerosVal(ResType,
I);
2748 if (ResName.
size() > 0)
2756 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2757 MIB.
addUse(
I.getOperand(i).getReg());
2762 MRI->
setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2763 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2765 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2766 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2773 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2774 .
addDef(
I.getOperand(1).getReg())
2782bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2783 SPIRVTypeInst ResType,
2784 MachineInstr &
I)
const {
2786 "selectAtomicCmpXchg only handles the spv_cmpxchg intrinsic");
2787 Register Ptr =
I.getOperand(2).getReg();
2788 Register ScopeReg =
I.getOperand(5).getReg();
2789 Register MemSemEqReg =
I.getOperand(6).getReg();
2790 Register MemSemNeqReg =
I.getOperand(7).getReg();
2792 Register Val =
I.getOperand(4).getReg();
2796 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2815 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2822 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2834 case SPIRV::StorageClass::DeviceOnlyINTEL:
2835 case SPIRV::StorageClass::HostOnlyINTEL:
2844 bool IsGRef =
false;
2845 bool IsAllowedRefs =
2847 unsigned Opcode = It.getOpcode();
2848 if (Opcode == SPIRV::OpConstantComposite ||
2849 Opcode == SPIRV::OpSpecConstantComposite ||
2850 Opcode == SPIRV::OpVariable ||
2851 Opcode == SPIRV::OpUntypedVariableKHR ||
2852 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2853 return IsGRef = true;
2854 return Opcode == SPIRV::OpName;
2856 return IsAllowedRefs && IsGRef;
2859Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2860 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2862 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2866SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2868 uint32_t Opcode)
const {
2869 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2870 TII.get(SPIRV::OpSpecConstantOp))
2878SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2879 SPIRVTypeInst SrcPtrTy)
const {
2880 SPIRVTypeInst GenericPtrTy =
2884 SPIRV::StorageClass::Generic),
2886 MachineFunction *MF =
I.getParent()->getParent();
2888 MachineInstrBuilder MIB = buildSpecConstantOp(
2890 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2900bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2901 SPIRVTypeInst ResType,
2902 MachineInstr &
I)
const {
2906 Register SrcPtr =
I.getOperand(1).getReg();
2911 return BuildCOPY(ResVReg, SrcPtr,
I);
2921 unsigned SpecOpcode = [&]() ->
unsigned {
2922 if (SrcSC == SPIRV::StorageClass::CodeSectionINTEL)
2923 return static_cast<uint32_t
>(SPIRV::Opcode::Bitcast);
2925 return static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric);
2927 return static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr);
2935 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2937 .constrainAllUses(
TII,
TRI, RBI);
2939 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2941 buildSpecConstantOp(
2943 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2944 .constrainAllUses(
TII,
TRI, RBI);
2951 return BuildCOPY(ResVReg, SrcPtr,
I);
2953 if ((SrcSC == SPIRV::StorageClass::Function &&
2954 DstSC == SPIRV::StorageClass::Private) ||
2955 (DstSC == SPIRV::StorageClass::Function &&
2956 SrcSC == SPIRV::StorageClass::Private))
2957 return BuildCOPY(ResVReg, SrcPtr,
I);
2961 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2964 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2967 SPIRVTypeInst GenericPtrTy =
2986 return selectUnOp(ResVReg, ResType,
I,
2987 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
2989 return selectUnOp(ResVReg, ResType,
I,
2990 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
2992 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2994 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
3004bool SPIRVInstructionSelector::selectPtrMask(
Register ResVReg,
3005 SPIRVTypeInst ResType,
3006 MachineInstr &
I)
const {
3008 return diagnoseUnsupported(
3009 I,
"G_PTRMASK is not supported with logical SPIR-V");
3014 Register PtrReg =
I.getOperand(1).getReg();
3015 Register MaskReg =
I.getOperand(2).getReg();
3034 ? SPIRV::OpBitwiseAndV
3035 : SPIRV::OpBitwiseAndS;
3058 return SPIRV::OpFOrdEqual;
3060 return SPIRV::OpFOrdGreaterThanEqual;
3062 return SPIRV::OpFOrdGreaterThan;
3064 return SPIRV::OpFOrdLessThanEqual;
3066 return SPIRV::OpFOrdLessThan;
3068 return SPIRV::OpFOrdNotEqual;
3070 return SPIRV::OpOrdered;
3072 return SPIRV::OpFUnordEqual;
3074 return SPIRV::OpFUnordGreaterThanEqual;
3076 return SPIRV::OpFUnordGreaterThan;
3078 return SPIRV::OpFUnordLessThanEqual;
3080 return SPIRV::OpFUnordLessThan;
3082 return SPIRV::OpFUnordNotEqual;
3084 return SPIRV::OpUnordered;
3094 return SPIRV::OpIEqual;
3096 return SPIRV::OpINotEqual;
3098 return SPIRV::OpSGreaterThanEqual;
3100 return SPIRV::OpSGreaterThan;
3102 return SPIRV::OpSLessThanEqual;
3104 return SPIRV::OpSLessThan;
3106 return SPIRV::OpUGreaterThanEqual;
3108 return SPIRV::OpUGreaterThan;
3110 return SPIRV::OpULessThanEqual;
3112 return SPIRV::OpULessThan;
3121 return SPIRV::OpPtrEqual;
3123 return SPIRV::OpPtrNotEqual;
3134 return SPIRV::OpLogicalEqual;
3136 return SPIRV::OpLogicalNotEqual;
3174bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
3175 SPIRVTypeInst ResType,
3177 unsigned OpAnyOrAll)
const {
3178 assert(
I.getNumOperands() == 3);
3179 assert(
I.getOperand(2).isReg());
3181 Register InputRegister =
I.getOperand(2).getReg();
3184 assert(InputType &&
"VReg has no type assigned");
3187 bool IsVectorTy = InputType->
getOpcode() == SPIRV::OpTypeVector;
3188 if (IsBoolTy && !IsVectorTy) {
3189 assert(ResVReg ==
I.getOperand(0).getReg());
3190 return BuildCOPY(ResVReg, InputRegister,
I);
3194 unsigned SpirvNotEqualId =
3195 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
3197 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
3202 IsBoolTy ? InputRegister
3210 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
3212 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
3229bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
3230 SPIRVTypeInst ResType,
3231 MachineInstr &
I)
const {
3232 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
3235bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
3236 SPIRVTypeInst ResType,
3237 MachineInstr &
I)
const {
3238 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
3242bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
3243 SPIRVTypeInst ResType,
3244 MachineInstr &
I)
const {
3245 assert(
I.getNumOperands() == 4);
3246 assert(
I.getOperand(2).isReg());
3247 assert(
I.getOperand(3).isReg());
3249 [[maybe_unused]] SPIRVTypeInst VecType =
3254 "dot product requires a vector of at least 2 components");
3256 [[maybe_unused]] SPIRVTypeInst EltType =
3265 .
addUse(
I.getOperand(2).getReg())
3266 .
addUse(
I.getOperand(3).getReg())
3271bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
3272 SPIRVTypeInst ResType,
3275 assert(
I.getNumOperands() == 4);
3276 assert(
I.getOperand(2).isReg());
3277 assert(
I.getOperand(3).isReg());
3280 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3284 .
addUse(
I.getOperand(2).getReg())
3285 .
addUse(
I.getOperand(3).getReg())
3292bool SPIRVInstructionSelector::selectIntegerDotExpansion(
3293 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3294 assert(
I.getNumOperands() == 4);
3295 assert(
I.getOperand(2).isReg());
3296 assert(
I.getOperand(3).isReg());
3300 Register Vec0 =
I.getOperand(2).getReg();
3301 Register Vec1 =
I.getOperand(3).getReg();
3305 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
3314 "dot product requires a vector of at least 2 components");
3317 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3327 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3338 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3350bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
3351 SPIRVTypeInst ResType,
3352 MachineInstr &
I)
const {
3354 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
3357 .
addUse(
I.getOperand(2).getReg())
3362bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
3363 SPIRVTypeInst ResType,
3364 MachineInstr &
I)
const {
3366 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
3369 .
addUse(
I.getOperand(2).getReg())
3374bool SPIRVInstructionSelector::selectOpIsFinite(
Register ResVReg,
3375 SPIRVTypeInst ResType,
3376 MachineInstr &
I)
const {
3378 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsFinite))
3381 .
addUse(
I.getOperand(2).getReg())
3386bool SPIRVInstructionSelector::selectOpIsNormal(
Register ResVReg,
3387 SPIRVTypeInst ResType,
3388 MachineInstr &
I)
const {
3390 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNormal))
3393 .
addUse(
I.getOperand(2).getReg())
3398template <
bool Signed>
3399bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
3400 SPIRVTypeInst ResType,
3401 MachineInstr &
I)
const {
3402 assert(
I.getNumOperands() == 5);
3403 assert(
I.getOperand(2).isReg());
3404 assert(
I.getOperand(3).isReg());
3405 assert(
I.getOperand(4).isReg());
3408 Register Acc =
I.getOperand(2).getReg();
3412 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3414 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
3419 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
3422 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3434template <
bool Signed>
3435bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
3436 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3437 assert(
I.getNumOperands() == 5);
3438 assert(
I.getOperand(2).isReg());
3439 assert(
I.getOperand(3).isReg());
3440 assert(
I.getOperand(4).isReg());
3443 Register Acc =
I.getOperand(2).getReg();
3449 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
3453 for (
unsigned i = 0; i < 4; i++) {
3476 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
3496 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3511bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
3512 SPIRVTypeInst ResType,
3513 MachineInstr &
I)
const {
3514 assert(
I.getNumOperands() == 3);
3515 assert(
I.getOperand(2).isReg());
3517 Register VZero = buildZerosValF(ResType,
I);
3518 Register VOne = buildOnesValF(ResType,
I);
3520 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
3523 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3525 .
addUse(
I.getOperand(2).getReg())
3532bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
3533 SPIRVTypeInst ResType,
3534 MachineInstr &
I)
const {
3535 assert(
I.getNumOperands() == 3);
3536 assert(
I.getOperand(2).isReg());
3538 Register InputRegister =
I.getOperand(2).getReg();
3540 auto &
DL =
I.getDebugLoc();
3543 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3550 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
3552 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
3560 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3565 if (NeedsConversion) {
3566 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
3577bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
3578 SPIRVTypeInst ResType,
3580 unsigned Opcode)
const {
3584 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
3590 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
3591 BMI.addUse(
I.getOperand(J).getReg());
3598bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
3601 bool WithGroupSync)
const {
3603 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
3605 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
3607 assert(((Scope != SPIRV::Scope::Workgroup) ||
3608 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
3609 "Workgroup Scope must set WorkGroupMemory semantic "
3610 "in Barrier instruction");
3612 assert(((Scope != SPIRV::Scope::Device) ||
3613 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
3614 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
3615 "Device Scope must set UniformMemory and ImageMemory semantic "
3616 "in Barrier instruction");
3622 if (WithGroupSync) {
3623 Register ExecReg = buildI32Constant(SPIRV::Scope::Workgroup,
I);
3627 Register ScopeReg = buildI32Constant(Scope,
I);
3628 Register MemSemReg = buildI32Constant(MemSem,
I);
3630 MI.addUse(ScopeReg).addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
3634bool SPIRVInstructionSelector::selectWaveActiveCountBits(
3635 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3640 if (!selectWaveOpInst(BallotReg, BallotType,
I,
3641 SPIRV::OpGroupNonUniformBallot))
3646 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3651 .
addImm(SPIRV::GroupOperation::Reduce)
3658bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
3659 SPIRVTypeInst ResType,
3660 MachineInstr &
I)
const {
3665 Register InputReg =
I.getOperand(2).getReg();
3670 bool IsVector = NumElems > 1;
3683 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
3684 SPIRV::OpGroupNonUniformAllEqual);
3689 ElementResults.
reserve(NumElems);
3691 for (
unsigned Idx = 0; Idx < NumElems; ++Idx) {
3704 ElemInput = Extracted;
3710 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3721 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3732bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3733 SPIRVTypeInst ResType,
3734 MachineInstr &
I)
const {
3736 assert(
I.getNumOperands() == 3);
3738 auto Op =
I.getOperand(2);
3748 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3750 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3751 return diagnoseUnsupported(
I,
"WavePrefixBitCount requires boolean input");
3772 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3776 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3783bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3784 SPIRVTypeInst ResType,
3786 bool IsUnsigned)
const {
3787 return selectWaveReduce(
3788 ResVReg, ResType,
I, IsUnsigned,
3789 [&](
Register InputRegister,
bool IsUnsigned) {
3790 const bool IsFloatTy =
3792 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3793 : SPIRV::OpGroupNonUniformSMax;
3794 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3798bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3799 SPIRVTypeInst ResType,
3801 bool IsUnsigned)
const {
3802 return selectWaveReduce(
3803 ResVReg, ResType,
I, IsUnsigned,
3804 [&](
Register InputRegister,
bool IsUnsigned) {
3805 const bool IsFloatTy =
3807 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3808 : SPIRV::OpGroupNonUniformSMin;
3809 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3813bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3814 SPIRVTypeInst ResType,
3815 MachineInstr &
I)
const {
3816 return selectWaveReduce(ResVReg, ResType,
I,
false,
3817 [&](
Register InputRegister,
bool IsUnsigned) {
3819 InputRegister, SPIRV::OpTypeFloat);
3820 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3821 : SPIRV::OpGroupNonUniformIAdd;
3825bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3826 SPIRVTypeInst ResType,
3827 MachineInstr &
I)
const {
3828 return selectWaveReduce(ResVReg, ResType,
I,
false,
3829 [&](
Register InputRegister,
bool IsUnsigned) {
3831 InputRegister, SPIRV::OpTypeFloat);
3832 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3833 : SPIRV::OpGroupNonUniformIMul;
3837template <
typename PickOpcodeFn>
3838bool SPIRVInstructionSelector::selectWaveReduce(
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::Reduce)
3858 .
addUse(
I.getOperand(2).getReg())
3863bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3864 SPIRVTypeInst ResType,
3866 unsigned Opcode)
const {
3867 return selectWaveReduce(
3868 ResVReg, ResType,
I,
false,
3869 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3872bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3873 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3874 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3875 [&](
Register InputRegister,
bool IsUnsigned) {
3877 InputRegister, SPIRV::OpTypeFloat);
3879 ? SPIRV::OpGroupNonUniformFAdd
3880 : SPIRV::OpGroupNonUniformIAdd;
3884bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3885 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3886 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3887 [&](
Register InputRegister,
bool IsUnsigned) {
3889 InputRegister, SPIRV::OpTypeFloat);
3891 ? SPIRV::OpGroupNonUniformFMul
3892 : SPIRV::OpGroupNonUniformIMul;
3896template <
typename PickOpcodeFn>
3897bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3898 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3899 PickOpcodeFn &&PickOpcode)
const {
3900 assert(
I.getNumOperands() == 3);
3901 assert(
I.getOperand(2).isReg());
3903 Register InputRegister =
I.getOperand(2).getReg();
3907 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3910 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3916 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3917 .
addUse(
I.getOperand(2).getReg())
3922bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
3923 SPIRVTypeInst ResType,
3926 assert(
I.getNumOperands() == 3);
3927 assert(
I.getOperand(2).isReg());
3929 Register InputRegister =
I.getOperand(2).getReg();
3935 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
3946bool SPIRVInstructionSelector::selectBitreverseViaI32(
Register ResVReg,
3947 SPIRVTypeInst ResType,
3954 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
3959 : SPIRV::OpUConvert;
3963 ShiftOp = SPIRV::OpShiftRightLogicalV;
3968 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3969 TII.get(SPIRV::OpConstantComposite))
3972 for (
unsigned It = 0; It <
N; ++It)
3976 ShiftConst = CompositeReg;
3981 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
3986 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
3991 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
3996 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
3999bool SPIRVInstructionSelector::handle64BitOverflow(
4001 unsigned int Opcode,
4008 "handle64BitOverflow should only be used for integer types");
4010 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
4012 MachineIRBuilder MIRBuilder(
I);
4014 SPIRVTypeInst I64x2Type =
4016 SPIRVTypeInst Vec2ResType =
4019 std::vector<Register> PartialRegs;
4021 unsigned CurrentComponent = 0;
4022 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
4026 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4027 TII.get(SPIRV::OpVectorShuffle))
4032 .
addImm(CurrentComponent)
4033 .
addImm(CurrentComponent + 1);
4043 PartialRegs.push_back(SubVecReg);
4046 if (CurrentComponent != ComponentCount) {
4052 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
4053 SPIRV::OpVectorExtractDynamic))
4062 PartialRegs.push_back(FinalElemResReg);
4066 return selectOpWithSrcs(ResVReg, ResType,
I, PartialRegs,
4067 SPIRV::OpCompositeConstruct);
4070bool SPIRVInstructionSelector::selectBitreverse64(
Register ResVReg,
4071 SPIRVTypeInst ResType,
4075 if (ComponentCount > 2)
4076 return handle64BitOverflow(
4077 ResVReg, ResType,
I, SrcReg, SPIRV::OpBitReverse,
4079 unsigned O) {
return this->selectBitreverse64(R,
T,
I, S); });
4081 MachineIRBuilder MIRBuilder(
I);
4085 I32Type, 2 * ComponentCount, MIRBuilder,
false);
4089 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
4094 if (!selectBitreverseNative(Reverse32, VecI32Type,
I, Vec32))
4101 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4102 TII.get(SPIRV::OpVectorShuffle))
4107 for (
unsigned J = 0; J < ComponentCount; ++J) {
4114 return selectOpWithSrcs(ResVReg, ResType,
I, {SwappedVec}, SPIRV::OpBitcast);
4117bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
4118 SPIRVTypeInst ResType,
4122 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
4130bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
4131 SPIRVTypeInst ResType,
4132 MachineInstr &
I)
const {
4133 Register OpReg =
I.getOperand(1).getReg();
4142 return selectBitreverseViaI32(ResVReg, ResType,
I, OpReg);
4144 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4146 return selectBitreverse64(ResVReg, ResType,
I, OpReg);
4148 return SPIRVInstructionSelector::diagnoseUnsupported(
4149 I,
"G_BITREVERSE only support 16,32,64 bits.");
4153 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4164 unsigned AndOp = SPIRV::OpBitwiseAndS;
4165 unsigned OrOp = SPIRV::OpBitwiseOrS;
4166 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
4167 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
4169 AndOp = SPIRV::OpBitwiseAndV;
4170 OrOp = SPIRV::OpBitwiseOrV;
4171 ShlOp = SPIRV::OpShiftLeftLogicalV;
4172 ShrOp = SPIRV::OpShiftRightLogicalV;
4178 const unsigned Shift) ->
Register {
4186 Register MaskReg = CreateConst(Mask);
4187 Register ShiftReg = CreateConst(Shift);
4194 if (!selectOpWithSrcs(
T1, ResType,
I, {Input, ShiftReg}, ShrOp) ||
4195 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
4196 !selectOpWithSrcs(T3, ResType,
I, {Input, MaskReg}, AndOp) ||
4197 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
4198 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
4206 uint64_t
Mask = ~0ull;
4207 while ((Shift >>= 1) > 0) {
4214 return BuildCOPY(ResVReg, Result,
I);
4217bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
4218 SPIRVTypeInst ResType,
4219 MachineInstr &
I)
const {
4220 assert(
I.getOperand(0).isReg() &&
I.getOperand(1).isReg() &&
4221 "G_FREEZE must define and use a register");
4222 Register OpReg =
I.getOperand(1).getReg();
4226 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
4239 if (MachineInstr *Def = MRI->
getVRegDef(OpReg)) {
4240 if (
Def->getOpcode() == TargetOpcode::COPY)
4243 switch (
Def->getOpcode()) {
4244 case SPIRV::ASSIGN_TYPE:
4245 if (MachineInstr *AssignToDef =
4247 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
4248 Reg =
Def->getOperand(2).getReg();
4251 case SPIRV::OpUndef:
4252 Reg =
Def->getOperand(1).getReg();
4255 unsigned DestOpCode;
4257 DestOpCode = SPIRV::OpConstantNull;
4258 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze of a "
4259 "static undef/poison lowered to OpConstantNull\n");
4261 DestOpCode = TargetOpcode::COPY;
4263 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze "
4264 "skipped, lowered as a copy of the operand\n");
4266 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
4267 .
addDef(
I.getOperand(0).getReg())
4275bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
4276 SPIRVTypeInst ResType,
4277 MachineInstr &
I)
const {
4279 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4281 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4285 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
4290 for (
unsigned i =
I.getNumExplicitDefs();
4291 i <
I.getNumExplicitOperands() && IsConst; ++i)
4295 if (!IsConst &&
N < 2)
4296 return diagnoseUnsupported(
4297 I,
"There must be at least two constituent operands in a vector");
4302 for (
unsigned i =
I.getNumExplicitDefs();
4303 i <
I.getNumExplicitOperands() && IsNullVector; ++i) {
4304 MachineInstr *
Def =
getDef(
I.getOperand(i), MRI);
4309 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4316 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4317 TII.get(IsConst ? SPIRV::OpConstantComposite
4318 : SPIRV::OpCompositeConstruct))
4321 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
4322 MIB.
addUse(
I.getOperand(i).getReg());
4327bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
4328 SPIRVTypeInst ResType,
4329 MachineInstr &
I)
const {
4331 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4333 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4339 if (!
I.getOperand(
OpIdx).isReg())
4346 if (!IsConst &&
N < 2)
4347 return diagnoseUnsupported(
4348 I,
"There must be at least two constituent operands in a vector");
4351 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4352 TII.get(IsConst ? SPIRV::OpConstantComposite
4353 : SPIRV::OpCompositeConstruct))
4356 for (
unsigned i = 0; i <
N; ++i)
4362bool SPIRVInstructionSelector::selectConcatVectors(
Register ResVReg,
4363 SPIRVTypeInst ResType,
4364 MachineInstr &
I)
const {
4368 if (ResType->
getOpcode() != SPIRV::OpTypeVector)
4370 "Cannot select G_CONCAT_VECTORS with a non-vector result");
4372 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4373 TII.get(SPIRV::OpCompositeConstruct))
4383bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
4384 SPIRVTypeInst ResType,
4385 MachineInstr &
I)
const {
4390 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
4392 Opcode = SPIRV::OpDemoteToHelperInvocation;
4394 Opcode = SPIRV::OpKill;
4396 if (MachineInstr *NextI =
I.getNextNode()) {
4398 NextI->eraseFromParent();
4408bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
4409 SPIRVTypeInst ResType,
unsigned CmpOpc,
4410 MachineInstr &
I)
const {
4411 Register Cmp0 =
I.getOperand(2).getReg();
4412 Register Cmp1 =
I.getOperand(3).getReg();
4415 "CMP operands should have the same type");
4416 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
4426bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
4427 SPIRVTypeInst ResType,
4428 MachineInstr &
I)
const {
4429 auto Pred =
I.getOperand(1).getPredicate();
4432 Register CmpOperand =
I.getOperand(2).getReg();
4434 bool IsPtrCmp = CmpOperandType && CmpOperandType.
isPointer();
4439 Register Op1 =
I.getOperand(3).getReg();
4443 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
4448 I.getOperand(3).setReg(NewOp1);
4454 return selectCmp(ResVReg, ResType, CmpOpc,
I);
4458SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
4459 SPIRVTypeInst ResType)
const {
4461 SPIRVTypeInst SpvI32Ty =
4464 auto ConstInt = ConstantInt::get(LLVMTy, Val);
4471 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4474 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
4477 .
addImm(APInt(32, Val).getZExtValue());
4479 GR.
add(ConstInt,
MI);
4486Register SPIRVInstructionSelector::buildI32ConstantInEntryBlock(
4487 uint32_t Val, MachineInstr &
I, SPIRVTypeInst ResType)
const {
4489 SPIRVTypeInst SpvI32Ty =
4491 auto *ConstInt = ConstantInt::get(LLVMTy, Val);
4496 MachineBasicBlock &EntryBB = *InsertIt->getParent();
4497 MachineInstr *
MI =
nullptr;
4501 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantNull))
4505 uint64_t ImmVal = APInt(32, Val).getZExtValue();
4506 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantI))
4512 GR.
add(ConstInt,
MI);
4517bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
4518 SPIRVTypeInst ResType,
4519 MachineInstr &
I)
const {
4521 return selectCmp(ResVReg, ResType, CmpOp,
I);
4524bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
4525 SPIRVTypeInst ResType,
4526 MachineInstr &
I)
const {
4528 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
4535 if (ResType->
getOpcode() != SPIRV::OpTypeVector &&
4536 ResType->
getOpcode() != SPIRV::OpTypeFloat)
4539 MachineIRBuilder MIRBuilder(
I);
4546 APFloat ConstVal(3.3219280948873623);
4550 APFloat::rmNearestTiesToEven, &LosesInfo);
4554 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
4555 ? SPIRV::OpVectorTimesScalar
4558 if (!selectOpWithSrcs(ArgReg, ResType,
I,
4559 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
4561 if (!selectExtInst(ResVReg, ResType,
I,
4562 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
4572Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
4573 MachineInstr &
I)
const {
4576 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4581bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
4587 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
4595 if (
Def->getOpcode() == SPIRV::OpConstantNull)
4598 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
4599 Def->getOpcode() == SPIRV::OpConstantI)
4612 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
4613 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
4615 Intrinsic::spv_const_composite)) {
4616 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
4617 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
4618 if (!IsZero(
Def->getOperand(i).getReg()))
4627Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
4628 MachineInstr &
I)
const {
4632 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4637Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
4638 MachineInstr &
I)
const {
4642 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4648 SPIRVTypeInst ResType,
4649 MachineInstr &
I)
const {
4653 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4658bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
4659 SPIRVTypeInst ResType,
4660 MachineInstr &
I)
const {
4661 Register SelectFirstArg =
I.getOperand(2).getReg();
4662 Register SelectSecondArg =
I.getOperand(3).getReg();
4671 SPIRV::OpTypeVector;
4678 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
4679 }
else if (IsPtrTy) {
4680 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
4682 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
4685 assert(IsScalarBool &&
"OpSelect with a scalar result requires a scalar "
4686 "boolean condition");
4688 Opcode = SPIRV::OpSelectSFSCond;
4689 }
else if (IsPtrTy) {
4690 Opcode = SPIRV::OpSelectSPSCond;
4692 Opcode = SPIRV::OpSelectSISCond;
4695 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4698 .
addUse(
I.getOperand(1).getReg())
4707bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
4708 SPIRVTypeInst ResType,
4710 MachineInstr &InsertAt,
4711 bool IsSigned)
const {
4713 Register ZeroReg = buildZerosVal(ResType, InsertAt);
4714 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
4715 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
4717 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4729bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
4730 SPIRVTypeInst ResType,
4731 MachineInstr &
I,
bool IsSigned,
4732 unsigned Opcode)
const {
4733 Register SrcReg =
I.getOperand(1).getReg();
4739 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
4744 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I, IsSigned);
4746 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
4749bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
4750 SPIRVTypeInst ResType, MachineInstr &
I,
4751 bool IsSigned)
const {
4752 Register SrcReg =
I.getOperand(1).getReg();
4754 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
4758 if (ResType == SrcType)
4759 return BuildCOPY(ResVReg, SrcReg,
I);
4761 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4762 return selectUnOp(ResVReg, ResType,
I, Opcode);
4765bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
4766 SPIRVTypeInst ResType,
4768 bool IsSigned)
const {
4769 MachineIRBuilder MIRBuilder(
I);
4770 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
4782 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
4785 .
addUse(
I.getOperand(1).getReg())
4786 .
addUse(
I.getOperand(2).getReg())
4791 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
4794 .
addUse(
I.getOperand(1).getReg())
4795 .
addUse(
I.getOperand(2).getReg())
4803 unsigned SelectOpcode =
4804 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
4809 .
addUse(buildOnesVal(
true, ResType,
I))
4810 .
addUse(buildZerosVal(ResType,
I))
4817 .
addUse(buildOnesVal(
false, ResType,
I))
4822bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
4825 SPIRVTypeInst IntTy,
4826 SPIRVTypeInst BoolTy)
const {
4829 bool IsVectorTy = IntTy->
getOpcode() == SPIRV::OpTypeVector;
4830 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
4832 Register One = buildOnesVal(
false, IntTy,
I);
4840 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
4849bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
4850 SPIRVTypeInst ResType,
4851 MachineInstr &
I)
const {
4852 Register IntReg =
I.getOperand(1).getReg();
4855 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
4856 if (ArgType == ResType)
4857 return BuildCOPY(ResVReg, IntReg,
I);
4859 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4860 return selectUnOp(ResVReg, ResType,
I, Opcode);
4863bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
4864 SPIRVTypeInst ResType,
4865 MachineInstr &
I)
const {
4866 unsigned Opcode =
I.getOpcode();
4867 unsigned TpOpcode = ResType->
getOpcode();
4869 if (ResType.
isPointer() || TpOpcode == SPIRV::OpTypeEvent) {
4870 assert(Opcode == TargetOpcode::G_CONSTANT &&
4871 I.getOperand(1).getCImm()->isZero());
4872 MachineBasicBlock &DepMBB =
I.getMF()->front();
4875 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
4882 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
4885bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
4886 SPIRVTypeInst ResType,
4887 MachineInstr &
I)
const {
4888 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4895bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
4896 SPIRVTypeInst ResType,
4897 MachineInstr &
I)
const {
4899 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
4903 .
addUse(
I.getOperand(3).getReg())
4905 .
addUse(
I.getOperand(2).getReg());
4906 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
4912bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
4913 SPIRVTypeInst ResType,
4914 MachineInstr &
I)
const {
4915 Type *MaybeResTy =
nullptr;
4920 "Expected aggregate type for extractv instruction");
4922 SPIRV::AccessQualifier::ReadWrite,
false);
4926 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
4929 .
addUse(
I.getOperand(2).getReg());
4930 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
4936bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
4937 SPIRVTypeInst ResType,
4938 MachineInstr &
I)
const {
4939 if (
getImm(
I.getOperand(4), MRI))
4940 return selectInsertVal(ResVReg, ResType,
I);
4942 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
4945 .
addUse(
I.getOperand(2).getReg())
4946 .
addUse(
I.getOperand(3).getReg())
4947 .
addUse(
I.getOperand(4).getReg())
4952bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
4953 SPIRVTypeInst ResType,
4954 MachineInstr &
I)
const {
4955 if (
getImm(
I.getOperand(3), MRI))
4956 return selectExtractVal(ResVReg, ResType,
I);
4958 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
4961 .
addUse(
I.getOperand(2).getReg())
4962 .
addUse(
I.getOperand(3).getReg())
4967bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
4968 SPIRVTypeInst ResType,
4969 MachineInstr &
I)
const {
4970 const bool IsGEPInBounds =
I.getOperand(2).getImm();
4973 const bool UseUntypedPointers =
4974 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
4979 if (UseUntypedPointers) {
4981 Opcode = IsGEPInBounds ? SPIRV::OpUntypedInBoundsAccessChainKHR
4982 : SPIRV::OpUntypedAccessChainKHR;
4984 Opcode = IsGEPInBounds ? SPIRV::OpUntypedInBoundsPtrAccessChainKHR
4985 : SPIRV::OpUntypedPtrAccessChainKHR;
4994 IsGEPInBounds ? SPIRV::OpInBoundsAccessChain : SPIRV::OpAccessChain;
4996 Opcode = IsGEPInBounds ? SPIRV::OpInBoundsPtrAccessChain
4997 : SPIRV::OpPtrAccessChain;
5002 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5007 if (UseUntypedPointers) {
5022 while (Def &&
Def->getOpcode() == TargetOpcode::COPY &&
5023 Def->getOperand(1).isReg())
5025 if (Def &&
Def->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)
5026 if (
const auto *GVar =
5029 SPIRV::AccessQualifier::ReadWrite,
5033 return diagnoseUnsupported(
5034 I,
"could not deduce the base type of an untyped access chain");
5039 Res.addUse(BaseReg);
5041 const bool IsAccessChainOpcode =
5042 (Opcode == SPIRV::OpAccessChain ||
5043 Opcode == SPIRV::OpInBoundsAccessChain ||
5044 Opcode == SPIRV::OpUntypedAccessChainKHR ||
5045 Opcode == SPIRV::OpUntypedInBoundsAccessChainKHR);
5047 assert((!IsAccessChainOpcode || (
getImm(
I.getOperand(4), MRI) &&
5048 foldImm(
I.getOperand(4), MRI) == 0)) &&
5049 "Cannot translate GEP to OpAccessChain.");
5052 const unsigned StartingIndex = IsAccessChainOpcode ? 5 : 4;
5053 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
5054 Res.addUse(
I.getOperand(i).getReg());
5055 Res.constrainAllUses(
TII,
TRI, RBI);
5060bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
5062 unsigned Lim =
I.getNumExplicitOperands();
5063 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
5064 Register OpReg =
I.getOperand(i).getReg();
5065 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
5067 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
5068 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
5069 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
5076 MachineFunction *MF =
I.getMF();
5082 SPIRVTypeInst WrapType = OpType;
5083 if (OpType->
getOpcode() == SPIRV::OpTypePointer &&
5085 SPIRV::StorageClass::CodeSectionINTEL) {
5087 SPIRV::StorageClass::Function,
I);
5094 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5095 TII.get(SPIRV::OpSpecConstantOp))
5098 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
5100 GR.
add(OpDefine, MIB);
5106bool SPIRVInstructionSelector::selectDerivativeInst(
5107 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
5108 const unsigned DPdOpCode)
const {
5111 if (!errorIfInstrOutsideShader(
I))
5117 Register SrcReg =
I.getOperand(2).getReg();
5122 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5125 .
addUse(
I.getOperand(2).getReg());
5127 MachineIRBuilder MIRBuilder(
I);
5130 if (componentCount != 1)
5138 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5143 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5148 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5156bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
5157 SPIRVTypeInst ResType,
5158 MachineInstr &
I)
const {
5162 case Intrinsic::spv_load:
5163 return selectLoad(ResVReg, ResType,
I);
5164 case Intrinsic::spv_atomic_load:
5165 return selectAtomicLoad(ResVReg, ResType,
I);
5166 case Intrinsic::spv_store:
5167 return selectStore(
I);
5168 case Intrinsic::spv_atomic_store:
5169 return selectAtomicStore(
I);
5170 case Intrinsic::spv_extractv:
5171 return selectExtractVal(ResVReg, ResType,
I);
5172 case Intrinsic::spv_insertv:
5173 return selectInsertVal(ResVReg, ResType,
I);
5174 case Intrinsic::spv_extractelt:
5175 return selectExtractElt(ResVReg, ResType,
I);
5176 case Intrinsic::spv_insertelt:
5177 return selectInsertElt(ResVReg, ResType,
I);
5178 case Intrinsic::spv_gep:
5179 return selectGEP(ResVReg, ResType,
I);
5180 case Intrinsic::spv_bitcast: {
5181 Register OpReg =
I.getOperand(2).getReg();
5182 SPIRVTypeInst OpType =
5186 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
5188 case Intrinsic::spv_unref_global:
5189 case Intrinsic::spv_init_global: {
5190 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
5195 Register GVarVReg =
MI->getOperand(0).getReg();
5196 if (!selectGlobalValue(GVarVReg, *
MI, Init))
5201 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
5203 MI->eraseFromParent();
5207 case Intrinsic::spv_undef: {
5208 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5214 case Intrinsic::spv_poison:
5215 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpPoisonKHR))
5220 case Intrinsic::spv_freeze:
5221 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
5224 .
addUse(
I.getOperand(2).getReg())
5227 case Intrinsic::spv_named_boolean_spec_constant: {
5228 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
5229 : SPIRV::OpSpecConstantFalse;
5231 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
5232 .
addDef(
I.getOperand(0).getReg())
5235 unsigned SpecId =
I.getOperand(2).getImm();
5237 SPIRV::Decoration::SpecId, {SpecId});
5241 case Intrinsic::spv_const_composite: {
5243 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
5249 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
5251 std::function<bool(
Register)> HasSpecConstOperand =
5261 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
5262 J < Def->getNumExplicitOperands(); ++J) {
5263 if (
Def->getOperand(J).isReg() &&
5264 HasSpecConstOperand(
Def->getOperand(J).getReg()))
5270 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
5271 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
5272 : SPIRV::OpConstantComposite;
5273 unsigned ContinuedOpc = HasSpecConst
5274 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
5275 : SPIRV::OpConstantCompositeContinuedINTEL;
5276 MachineIRBuilder MIR(
I);
5278 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
5280 for (
auto *Instr : Instructions) {
5281 Instr->setDebugLoc(
I.getDebugLoc());
5286 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5293 case Intrinsic::spv_assign_name: {
5294 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
5295 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
5296 for (
unsigned i =
I.getNumExplicitDefs() + 2;
5297 i <
I.getNumExplicitOperands(); ++i) {
5298 MIB.
addImm(
I.getOperand(i).getImm());
5303 case Intrinsic::spv_switch: {
5304 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
5305 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5306 if (
I.getOperand(i).isReg())
5307 MIB.
addReg(
I.getOperand(i).getReg());
5308 else if (
I.getOperand(i).isCImm())
5309 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
5310 else if (
I.getOperand(i).isMBB())
5311 MIB.
addMBB(
I.getOperand(i).getMBB());
5318 case Intrinsic::spv_loop_merge: {
5319 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
5320 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5321 if (
I.getOperand(i).isMBB())
5322 MIB.
addMBB(
I.getOperand(i).getMBB());
5329 case Intrinsic::spv_loop_control_intel: {
5331 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
5332 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
5337 case Intrinsic::spv_selection_merge: {
5339 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
5340 assert(
I.getOperand(1).isMBB() &&
5341 "operand 1 to spv_selection_merge must be a basic block");
5342 MIB.
addMBB(
I.getOperand(1).getMBB());
5343 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
5347 case Intrinsic::spv_cmpxchg:
5348 return selectAtomicCmpXchg(ResVReg, ResType,
I);
5349 case Intrinsic::spv_unreachable:
5350 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
5353 case Intrinsic::spv_abort:
5354 return selectAbort(
I);
5355 case Intrinsic::spv_alloca:
5356 return selectFrameIndex(ResVReg, ResType,
I);
5357 case Intrinsic::spv_alloca_array:
5358 return selectAllocaArray(ResVReg, ResType,
I);
5359 case Intrinsic::spv_assume:
5361 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
5362 .
addUse(
I.getOperand(1).getReg())
5367 case Intrinsic::spv_expect:
5369 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
5372 .
addUse(
I.getOperand(2).getReg())
5373 .
addUse(
I.getOperand(3).getReg())
5378 case Intrinsic::arithmetic_fence:
5379 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
5380 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
5383 .
addUse(
I.getOperand(2).getReg())
5387 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
5389 case Intrinsic::spv_thread_id:
5395 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
5397 case Intrinsic::spv_thread_id_in_group:
5403 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
5405 case Intrinsic::spv_group_id:
5411 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
5413 case Intrinsic::spv_flattened_thread_id_in_group:
5420 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
5422 case Intrinsic::spv_workgroup_size:
5423 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
5425 case Intrinsic::spv_global_size:
5426 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
5428 case Intrinsic::spv_global_offset:
5429 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
5431 case Intrinsic::spv_num_workgroups:
5432 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
5434 case Intrinsic::spv_subgroup_size:
5435 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
5437 case Intrinsic::spv_num_subgroups:
5438 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
5440 case Intrinsic::spv_subgroup_id:
5441 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
5442 case Intrinsic::spv_subgroup_local_invocation_id:
5443 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
5444 ResVReg, ResType,
I);
5445 case Intrinsic::spv_subgroup_max_size:
5446 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
5448 case Intrinsic::spv_fdot:
5449 return selectFloatDot(ResVReg, ResType,
I);
5450 case Intrinsic::spv_udot:
5451 case Intrinsic::spv_sdot:
5452 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5454 return selectIntegerDot(ResVReg, ResType,
I,
5455 IID == Intrinsic::spv_sdot);
5456 return selectIntegerDotExpansion(ResVReg, ResType,
I);
5457 case Intrinsic::spv_dot4add_i8packed:
5458 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5460 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
5461 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
5462 case Intrinsic::spv_dot4add_u8packed:
5463 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5465 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
5466 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
5467 case Intrinsic::spv_all:
5468 return selectAll(ResVReg, ResType,
I);
5469 case Intrinsic::spv_any:
5470 return selectAny(ResVReg, ResType,
I);
5471 case Intrinsic::spv_distance:
5472 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
5473 case Intrinsic::spv_lerp:
5474 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
5475 case Intrinsic::spv_length:
5476 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
5477 case Intrinsic::spv_degrees:
5478 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
5479 case Intrinsic::spv_faceforward:
5480 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
5481 case Intrinsic::spv_frac:
5482 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
5483 case Intrinsic::spv_isinf:
5484 return selectOpIsInf(ResVReg, ResType,
I);
5485 case Intrinsic::spv_isnan:
5486 return selectOpIsNan(ResVReg, ResType,
I);
5487 case Intrinsic::spv_isfinite:
5488 return selectOpIsFinite(ResVReg, ResType,
I);
5489 case Intrinsic::spv_isnormal:
5490 return selectOpIsNormal(ResVReg, ResType,
I);
5491 case Intrinsic::spv_normalize:
5492 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
5493 case Intrinsic::spv_refract:
5494 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
5495 case Intrinsic::spv_reflect:
5496 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
5497 case Intrinsic::spv_rsqrt:
5498 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
5499 case Intrinsic::spv_sign:
5500 return selectSign(ResVReg, ResType,
I);
5501 case Intrinsic::spv_smoothstep:
5502 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
5503 case Intrinsic::spv_firstbituhigh:
5504 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
5505 case Intrinsic::spv_firstbitshigh:
5506 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
5507 case Intrinsic::spv_firstbitlow:
5508 return selectFirstBitLow(ResVReg, ResType,
I);
5509 case Intrinsic::spv_all_memory_barrier:
5510 return selectBarrierInst(
I, SPIRV::Scope::Device,
5511 SPIRV::MemorySemantics::UniformMemory |
5512 SPIRV::MemorySemantics::ImageMemory |
5513 SPIRV::MemorySemantics::WorkgroupMemory,
5515 case Intrinsic::spv_all_memory_barrier_with_group_sync:
5516 return selectBarrierInst(
I, SPIRV::Scope::Device,
5517 SPIRV::MemorySemantics::UniformMemory |
5518 SPIRV::MemorySemantics::ImageMemory |
5519 SPIRV::MemorySemantics::WorkgroupMemory,
5521 case Intrinsic::spv_device_memory_barrier:
5522 return selectBarrierInst(
I, SPIRV::Scope::Device,
5523 SPIRV::MemorySemantics::UniformMemory |
5524 SPIRV::MemorySemantics::ImageMemory,
5526 case Intrinsic::spv_device_memory_barrier_with_group_sync:
5527 return selectBarrierInst(
I, SPIRV::Scope::Device,
5528 SPIRV::MemorySemantics::UniformMemory |
5529 SPIRV::MemorySemantics::ImageMemory,
5531 case Intrinsic::spv_group_memory_barrier:
5532 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5533 SPIRV::MemorySemantics::WorkgroupMemory,
5535 case Intrinsic::spv_group_memory_barrier_with_group_sync:
5536 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5537 SPIRV::MemorySemantics::WorkgroupMemory,
5539 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
5540 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
5541 SPIRV::StorageClass::StorageClass ResSC =
5544 return diagnoseUnsupported(
I,
"The target storage class is not castable "
5545 "from the Generic storage class");
5546 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
5554 case Intrinsic::spv_lifetime_start:
5555 case Intrinsic::spv_lifetime_end: {
5556 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
5557 : SPIRV::OpLifetimeStop;
5558 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
5559 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
5568 case Intrinsic::spv_saturate:
5569 return selectSaturate(ResVReg, ResType,
I);
5570 case Intrinsic::spv_nclamp:
5571 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
5572 case Intrinsic::spv_uclamp:
5573 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
5574 case Intrinsic::spv_sclamp:
5575 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
5576 case Intrinsic::spv_subgroup_prefix_bit_count:
5577 return selectWavePrefixBitCount(ResVReg, ResType,
I);
5578 case Intrinsic::spv_wave_active_countbits:
5579 return selectWaveActiveCountBits(ResVReg, ResType,
I);
5580 case Intrinsic::spv_wave_all_equal:
5581 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
5582 case Intrinsic::spv_wave_all:
5583 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
5584 case Intrinsic::spv_wave_any:
5585 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
5586 case Intrinsic::spv_subgroup_ballot:
5587 return selectWaveOpInst(ResVReg, ResType,
I,
5588 SPIRV::OpGroupNonUniformBallot);
5589 case Intrinsic::spv_wave_is_first_lane:
5590 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
5591 case Intrinsic::spv_wave_reduce_or:
5592 return selectWaveReduceOp(ResVReg, ResType,
I,
5593 SPIRV::OpGroupNonUniformBitwiseOr);
5594 case Intrinsic::spv_wave_reduce_xor:
5595 return selectWaveReduceOp(ResVReg, ResType,
I,
5596 SPIRV::OpGroupNonUniformBitwiseXor);
5597 case Intrinsic::spv_wave_reduce_and:
5598 return selectWaveReduceOp(ResVReg, ResType,
I,
5599 SPIRV::OpGroupNonUniformBitwiseAnd);
5600 case Intrinsic::spv_wave_reduce_umax:
5601 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
5602 case Intrinsic::spv_wave_reduce_max:
5603 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
5604 case Intrinsic::spv_wave_reduce_umin:
5605 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
5606 case Intrinsic::spv_wave_reduce_min:
5607 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
5608 case Intrinsic::spv_wave_reduce_sum:
5609 return selectWaveReduceSum(ResVReg, ResType,
I);
5610 case Intrinsic::spv_wave_product:
5611 return selectWaveReduceProduct(ResVReg, ResType,
I);
5612 case Intrinsic::spv_wave_readlane:
5613 return selectWaveOpInst(ResVReg, ResType,
I,
5614 SPIRV::OpGroupNonUniformShuffle);
5615 case Intrinsic::spv_wave_prefix_sum:
5616 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
5617 case Intrinsic::spv_wave_prefix_product:
5618 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
5619 case Intrinsic::spv_quad_read_across_x: {
5620 return selectQuadSwap(ResVReg, ResType,
I, 0);
5622 case Intrinsic::spv_quad_read_across_y: {
5623 return selectQuadSwap(ResVReg, ResType,
I, 1);
5625 case Intrinsic::spv_quad_read_across_diagonal: {
5626 return selectQuadSwap(ResVReg, ResType,
I, 2);
5628 case Intrinsic::spv_step:
5629 return selectExtInst(ResVReg, ResType,
I, CL::step, GL::Step);
5630 case Intrinsic::spv_radians:
5631 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
5635 case Intrinsic::instrprof_increment:
5636 case Intrinsic::instrprof_increment_step:
5637 case Intrinsic::instrprof_value_profile:
5640 case Intrinsic::spv_value_md:
5642 case Intrinsic::spv_resource_handlefrombinding: {
5643 return selectHandleFromBinding(ResVReg, ResType,
I);
5645 case Intrinsic::spv_resource_counterhandlefrombinding:
5646 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
5647 case Intrinsic::spv_resource_updatecounter:
5648 return selectUpdateCounter(ResVReg, ResType,
I);
5649 case Intrinsic::spv_resource_store_typedbuffer: {
5650 return selectImageWriteIntrinsic(
I);
5652 case Intrinsic::spv_resource_load_typedbuffer: {
5653 return selectReadImageIntrinsic(ResVReg, ResType,
I);
5655 case Intrinsic::spv_resource_load_level: {
5656 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
5658 case Intrinsic::spv_resource_getdimensions_x:
5659 case Intrinsic::spv_resource_getdimensions_xy:
5660 case Intrinsic::spv_resource_getdimensions_xyz: {
5661 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
5663 case Intrinsic::spv_resource_getdimensions_levels_x:
5664 case Intrinsic::spv_resource_getdimensions_levels_xy:
5665 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
5666 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
5668 case Intrinsic::spv_resource_getdimensions_ms_xy:
5669 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
5670 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
5672 case Intrinsic::spv_resource_calculate_lod:
5673 case Intrinsic::spv_resource_calculate_lod_unclamped:
5674 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
5675 case Intrinsic::spv_resource_sample:
5676 case Intrinsic::spv_resource_sample_clamp:
5677 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
5678 case Intrinsic::spv_resource_samplebias:
5679 case Intrinsic::spv_resource_samplebias_clamp:
5680 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
5681 case Intrinsic::spv_resource_samplegrad:
5682 case Intrinsic::spv_resource_samplegrad_clamp:
5683 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
5684 case Intrinsic::spv_resource_samplelevel:
5685 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
5686 case Intrinsic::spv_resource_samplecmp:
5687 case Intrinsic::spv_resource_samplecmp_clamp:
5688 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
5689 case Intrinsic::spv_resource_samplecmplevelzero:
5690 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
5691 case Intrinsic::spv_resource_gather:
5692 case Intrinsic::spv_resource_gather_cmp:
5693 return selectGatherIntrinsic(ResVReg, ResType,
I);
5694 case Intrinsic::spv_resource_getbasepointer:
5695 case Intrinsic::spv_resource_getpointer: {
5696 return selectResourceGetPointer(ResVReg, ResType,
I);
5698 case Intrinsic::spv_pushconstant_getpointer: {
5699 return selectPushConstantGetPointer(ResVReg, ResType,
I);
5701 case Intrinsic::spv_discard: {
5702 return selectDiscard(ResVReg, ResType,
I);
5704 case Intrinsic::spv_resource_nonuniformindex: {
5705 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
5707 case Intrinsic::spv_unpackhalf2x16: {
5708 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
5710 case Intrinsic::spv_packhalf2x16: {
5711 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
5713 case Intrinsic::spv_ddx:
5714 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
5715 case Intrinsic::spv_ddy:
5716 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
5717 case Intrinsic::spv_ddx_coarse:
5718 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
5719 case Intrinsic::spv_ddy_coarse:
5720 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
5721 case Intrinsic::spv_ddx_fine:
5722 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
5723 case Intrinsic::spv_ddy_fine:
5724 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
5725 case Intrinsic::spv_fwidth:
5726 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
5727 case Intrinsic::spv_masked_gather:
5728 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5729 return selectMaskedGather(ResVReg, ResType,
I);
5730 return diagnoseUnsupported(
5731 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
5732 case Intrinsic::spv_masked_scatter:
5733 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5734 return selectMaskedScatter(
I);
5735 return diagnoseUnsupported(
5736 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
5737 case Intrinsic::returnaddress:
5738 case Intrinsic::frameaddress: {
5740 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5747 return diagnoseUnsupported(
I,
"intrinsic selection not implemented.");
5752bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
5753 SPIRVTypeInst ResType,
5754 MachineInstr &
I)
const {
5757 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
5764bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
5765 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5767 assert(Intr.getIntrinsicID() ==
5768 Intrinsic::spv_resource_counterhandlefrombinding);
5771 Register MainHandleReg = Intr.getOperand(2).getReg();
5773 assert(MainHandleDef->getIntrinsicID() ==
5774 Intrinsic::spv_resource_handlefrombinding);
5778 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
5779 Register IndexReg = MainHandleDef->getOperand(5).getReg();
5780 std::string CounterName =
5785 MachineIRBuilder MIRBuilder(
I);
5787 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
5789 ArraySize, IndexReg, CounterName, MIRBuilder);
5791 return BuildCOPY(ResVReg, CounterVarReg,
I);
5794bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
5795 SPIRVTypeInst ResType,
5796 MachineInstr &
I)
const {
5798 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
5800 Register CounterHandleReg = Intr.getOperand(2).getReg();
5801 Register IncrReg = Intr.getOperand(3).getReg();
5808 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
5809 assert(CounterVarPointeeType &&
5810 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
5811 "Counter variable must be a struct");
5813 SPIRV::StorageClass::StorageBuffer &&
5814 "Counter variable must be in the storage buffer storage class");
5816 "Counter variable must have exactly 1 member in the struct");
5817 const SPIRVTypeInst MemberType =
5820 "Counter variable struct must have a single i32 member");
5824 MachineIRBuilder MIRBuilder(
I);
5826 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
5829 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
5835 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5838 .
addUse(CounterHandleReg)
5845 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
5848 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
5851 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
5860 return BuildCOPY(ResVReg, AtomicRes,
I);
5868 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
5876bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
5877 SPIRVTypeInst ResType,
5878 MachineInstr &
I)
const {
5886 Register ImageReg =
I.getOperand(2).getReg();
5894 Register IdxReg =
I.getOperand(3).getReg();
5896 MachineInstr &Pos =
I;
5898 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
5902bool SPIRVInstructionSelector::generateSampleImage(
5905 DebugLoc Loc, MachineInstr &Pos)
const {
5916 if (!loadHandleBeforePosition(NewSamplerReg,
5922 MachineIRBuilder MIRBuilder(Pos);
5935 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
5936 ImOps.Lod.has_value();
5937 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
5938 : SPIRV::OpImageSampleImplicitLod;
5940 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
5941 : SPIRV::OpImageSampleDrefImplicitLod;
5950 MIB.
addUse(*ImOps.Compare);
5952 uint32_t ImageOperands = 0;
5954 ImageOperands |= SPIRV::ImageOperand::Bias;
5956 ImageOperands |= SPIRV::ImageOperand::Lod;
5957 if (ImOps.GradX && ImOps.GradY)
5958 ImageOperands |= SPIRV::ImageOperand::Grad;
5959 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
5961 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
5964 "Non-constant offsets are not supported in sample instructions.");
5969 ImageOperands |= SPIRV::ImageOperand::MinLod;
5971 if (ImageOperands != 0) {
5972 MIB.
addImm(ImageOperands);
5973 if (ImageOperands & SPIRV::ImageOperand::Bias)
5975 if (ImageOperands & SPIRV::ImageOperand::Lod)
5977 if (ImageOperands & SPIRV::ImageOperand::Grad) {
5978 MIB.
addUse(*ImOps.GradX);
5979 MIB.
addUse(*ImOps.GradY);
5982 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
5983 MIB.
addUse(*ImOps.Offset);
5984 if (ImageOperands & SPIRV::ImageOperand::MinLod)
5985 MIB.
addUse(*ImOps.MinLod);
5992bool SPIRVInstructionSelector::selectImageQuerySize(
5994 std::optional<Register> LodReg)
const {
5996 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
5999 "ImageReg is not an image type.");
6001 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6003 unsigned NumComponents = 0;
6005 case SPIRV::Dim::DIM_1D:
6006 case SPIRV::Dim::DIM_Buffer:
6007 NumComponents =
IsArray ? 2 : 1;
6009 case SPIRV::Dim::DIM_2D:
6010 case SPIRV::Dim::DIM_Cube:
6011 case SPIRV::Dim::DIM_Rect:
6012 NumComponents =
IsArray ? 3 : 2;
6014 case SPIRV::Dim::DIM_3D:
6018 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
6023 SPIRVTypeInst ResType =
6028 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6038bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
6039 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6040 Register ImageReg =
I.getOperand(2).getReg();
6047 return selectImageQuerySize(NewImageReg, ResVReg,
I);
6050bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
6051 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6052 Register ImageReg =
I.getOperand(2).getReg();
6061 Register LodReg =
I.getOperand(3).getReg();
6064 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
6066 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
6073 TII.get(SPIRV::OpImageQueryLevels))
6080 TII.get(SPIRV::OpCompositeConstruct))
6090bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
6091 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6092 Register ImageReg =
I.getOperand(2).getReg();
6103 "OpImageQuerySamples requires a multisampled image");
6105 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
6113 TII.get(SPIRV::OpImageQuerySamples))
6120 TII.get(SPIRV::OpCompositeConstruct))
6130bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
6131 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6132 Register ImageReg =
I.getOperand(2).getReg();
6133 Register SamplerReg =
I.getOperand(3).getReg();
6134 Register CoordinateReg =
I.getOperand(4).getReg();
6150 if (!loadHandleBeforePosition(
6155 MachineIRBuilder MIRBuilder(
I);
6161 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6171 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
6178 unsigned ExtractedIndex =
6180 Intrinsic::spv_resource_calculate_lod_unclamped
6184 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6185 TII.get(SPIRV::OpCompositeExtract))
6195bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
6196 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6197 Register ImageReg =
I.getOperand(2).getReg();
6198 Register SamplerReg =
I.getOperand(3).getReg();
6199 Register CoordinateReg =
I.getOperand(4).getReg();
6200 ImageOperands ImOps;
6201 if (
I.getNumOperands() > 5)
6202 ImOps.Offset =
I.getOperand(5).getReg();
6203 if (
I.getNumOperands() > 6)
6204 ImOps.MinLod =
I.getOperand(6).getReg();
6205 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6206 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6209bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
6210 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6211 Register ImageReg =
I.getOperand(2).getReg();
6212 Register SamplerReg =
I.getOperand(3).getReg();
6213 Register CoordinateReg =
I.getOperand(4).getReg();
6214 ImageOperands ImOps;
6215 ImOps.Bias =
I.getOperand(5).getReg();
6216 if (
I.getNumOperands() > 6)
6217 ImOps.Offset =
I.getOperand(6).getReg();
6218 if (
I.getNumOperands() > 7)
6219 ImOps.MinLod =
I.getOperand(7).getReg();
6220 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6221 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6224bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
6225 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6226 Register ImageReg =
I.getOperand(2).getReg();
6227 Register SamplerReg =
I.getOperand(3).getReg();
6228 Register CoordinateReg =
I.getOperand(4).getReg();
6229 ImageOperands ImOps;
6230 ImOps.GradX =
I.getOperand(5).getReg();
6231 ImOps.GradY =
I.getOperand(6).getReg();
6232 if (
I.getNumOperands() > 7)
6233 ImOps.Offset =
I.getOperand(7).getReg();
6234 if (
I.getNumOperands() > 8)
6235 ImOps.MinLod =
I.getOperand(8).getReg();
6236 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6237 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6240bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
6241 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6242 Register ImageReg =
I.getOperand(2).getReg();
6243 Register SamplerReg =
I.getOperand(3).getReg();
6244 Register CoordinateReg =
I.getOperand(4).getReg();
6245 ImageOperands ImOps;
6246 ImOps.Lod =
I.getOperand(5).getReg();
6247 if (
I.getNumOperands() > 6)
6248 ImOps.Offset =
I.getOperand(6).getReg();
6249 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6250 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6253bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
6254 SPIRVTypeInst ResType,
6255 MachineInstr &
I)
const {
6256 Register ImageReg =
I.getOperand(2).getReg();
6257 Register SamplerReg =
I.getOperand(3).getReg();
6258 Register CoordinateReg =
I.getOperand(4).getReg();
6259 ImageOperands ImOps;
6260 ImOps.Compare =
I.getOperand(5).getReg();
6261 if (
I.getNumOperands() > 6)
6262 ImOps.Offset =
I.getOperand(6).getReg();
6263 if (
I.getNumOperands() > 7)
6264 ImOps.MinLod =
I.getOperand(7).getReg();
6265 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6266 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6269bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
6270 SPIRVTypeInst ResType,
6271 MachineInstr &
I)
const {
6272 Register ImageReg =
I.getOperand(2).getReg();
6273 Register CoordinateReg =
I.getOperand(3).getReg();
6274 Register LodReg =
I.getOperand(4).getReg();
6276 ImageOperands ImOps;
6278 if (
I.getNumOperands() > 5)
6279 ImOps.Offset =
I.getOperand(5).getReg();
6291 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
6292 I.getDebugLoc(),
I, &ImOps);
6295bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
6296 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6297 Register ImageReg =
I.getOperand(2).getReg();
6298 Register SamplerReg =
I.getOperand(3).getReg();
6299 Register CoordinateReg =
I.getOperand(4).getReg();
6300 ImageOperands ImOps;
6301 ImOps.Compare =
I.getOperand(5).getReg();
6302 if (
I.getNumOperands() > 6)
6303 ImOps.Offset =
I.getOperand(6).getReg();
6306 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6307 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6310bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
6311 SPIRVTypeInst ResType,
6312 MachineInstr &
I)
const {
6313 Register ImageReg =
I.getOperand(2).getReg();
6314 Register SamplerReg =
I.getOperand(3).getReg();
6315 Register CoordinateReg =
I.getOperand(4).getReg();
6318 "ImageReg is not an image type.");
6323 ComponentOrCompareReg =
I.getOperand(5).getReg();
6324 OffsetReg =
I.getOperand(6).getReg();
6327 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
6331 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6332 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
6333 Dim != SPIRV::Dim::DIM_Rect) {
6335 "Gather operations are only supported for 2D, Cube, and Rect images.");
6342 if (!loadHandleBeforePosition(
6347 MachineIRBuilder MIRBuilder(
I);
6348 SPIRVTypeInst SampledImageType =
6353 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6361 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
6363 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
6365 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6370 .
addUse(ComponentOrCompareReg);
6372 uint32_t ImageOperands = 0;
6373 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
6374 if (Dim == SPIRV::Dim::DIM_Cube) {
6376 "Gather operations with offset are not supported for Cube images.");
6380 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6382 ImageOperands |= SPIRV::ImageOperand::Offset;
6386 if (ImageOperands != 0) {
6387 MIB.
addImm(ImageOperands);
6389 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6397bool SPIRVInstructionSelector::generateImageReadOrFetch(
6400 const ImageOperands *ImOps)
const {
6403 "ImageReg is not an image type.");
6405 bool IsSignedInteger =
6410 bool IsFetch = (SampledOp.getImm() == 1);
6412 auto AddOperands = [&](MachineInstrBuilder &MIB) {
6413 uint32_t ImageOperandsMask = 0;
6414 if (IsSignedInteger)
6415 ImageOperandsMask |= 0x1000;
6417 if (IsFetch && ImOps) {
6419 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
6420 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
6422 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
6424 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
6428 if (ImageOperandsMask != 0) {
6429 MIB.
addImm(ImageOperandsMask);
6430 if (IsFetch && ImOps) {
6433 if (ImOps->Offset &&
6434 (ImageOperandsMask &
6435 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
6436 MIB.
addUse(*ImOps->Offset);
6445 SPIRVTypeInst SampledType =
6448 SPIRVTypeInst ReadType =
6449 widenTypeToVec4(IsPacked ? SampledType : ResType, Pos);
6450 bool ReadTypeMatchesResult = ReadType == ResType;
6452 Register ReadReg = ReadTypeMatchesResult
6458 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
6464 BMI.constrainAllUses(
TII,
TRI, RBI);
6466 if (ReadTypeMatchesResult)
6479 if (ResultSize == 1) {
6488 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
6491bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
6492 SPIRVTypeInst ResType,
6493 MachineInstr &
I)
const {
6494 Register ResourcePtr =
I.getOperand(2).getReg();
6496 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
6505 MachineIRBuilder MIRBuilder(
I);
6510 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6516 if (
I.getNumExplicitOperands() > 3) {
6517 Register IndexReg =
I.getOperand(3).getReg();
6524bool SPIRVInstructionSelector::selectPushConstantGetPointer(
6525 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6530bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
6531 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6532 Register ObjReg =
I.getOperand(2).getReg();
6533 if (!BuildCOPY(ResVReg, ObjReg,
I))
6543 decorateUsesAsNonUniform(ResVReg);
6547void SPIRVInstructionSelector::decorateUsesAsNonUniform(
6550 {NonUniformReg,
nullptr}};
6551 llvm::SmallSet<Register, 8> Visited;
6552 while (WorkList.
size() > 0) {
6555 if (!Visited.
insert(CurrentReg).second)
6558 bool IsDecorated =
false;
6560 if (
Use.getOpcode() == SPIRV::OpDecorate &&
6561 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
6567 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
6569 if (ResultReg == CurrentReg)
6577 MachineInstr &InsertPt =
6580 SPIRV::Decoration::NonUniformEXT, {});
6585bool SPIRVInstructionSelector::extractSubvector(
6587 MachineInstr &InsertionPoint)
const {
6589 [[maybe_unused]] uint64_t InputSize =
6592 assert(InputSize > 1 &&
"The input must be a vector.");
6593 assert(ResultSize > 1 &&
"The result must be a vector.");
6594 assert(ResultSize < InputSize &&
6595 "Cannot extract more element than there are in the input.");
6599 for (uint64_t
I = 0;
I < ResultSize;
I++) {
6602 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6611 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
6613 TII.get(SPIRV::OpCompositeConstruct))
6617 for (
Register ComponentReg : ComponentRegisters)
6618 MIB.
addUse(ComponentReg);
6623bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
6624 MachineInstr &
I)
const {
6631 Register ImageReg =
I.getOperand(1).getReg();
6639 Register CoordinateReg =
I.getOperand(2).getReg();
6640 Register DataReg =
I.getOperand(3).getReg();
6643 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
6651Register SPIRVInstructionSelector::buildPointerToResource(
6652 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
6653 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
6654 StringRef Name, MachineIRBuilder MIRBuilder)
const {
6656 if (ArraySize == 1) {
6657 SPIRVTypeInst PtrType =
6660 "SpirvResType did not have an explicit layout.");
6665 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
6666 SPIRVTypeInst VarPointerType =
6669 VarPointerType, Set,
Binding, Name, MIRBuilder);
6671 SPIRVTypeInst ResPointerType =
6684bool SPIRVInstructionSelector::selectFirstBitSet16(
6685 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
6686 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
6688 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
6692 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
6695bool SPIRVInstructionSelector::selectFirstBitSet32(
6697 unsigned BitSetOpcode)
const {
6698 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6701 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6708bool SPIRVInstructionSelector::selectFirstBitSet64(
6710 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
6723 if (ComponentCount > 2) {
6724 auto Func = [
this, SwapPrimarySide](
Register ResVReg, SPIRVTypeInst ResType,
6726 unsigned Opcode) ->
bool {
6727 return this->selectFirstBitSet64(ResVReg, ResType,
I, SrcReg, Opcode,
6731 return handle64BitOverflow(ResVReg, ResType,
I, SrcReg, BitSetOpcode, Func);
6735 MachineIRBuilder MIRBuilder(
I);
6737 BaseType, 2 * ComponentCount, MIRBuilder,
false);
6741 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
6747 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
6754 bool IsScalarRes = ResType->
getOpcode() != SPIRV::OpTypeVector;
6757 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
6758 SPIRV::OpVectorExtractDynamic))
6760 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
6761 SPIRV::OpVectorExtractDynamic))
6765 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6766 TII.get(SPIRV::OpVectorShuffle))
6774 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
6780 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6781 TII.get(SPIRV::OpVectorShuffle))
6789 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
6809 SelectOp = SPIRV::OpSelectSISCond;
6810 AddOp = SPIRV::OpIAddS;
6818 SelectOp = SPIRV::OpSelectVIVCond;
6819 AddOp = SPIRV::OpIAddV;
6825 Register RegSecondaryOffset = Reg0;
6829 if (SwapPrimarySide) {
6830 PrimaryReg = LowReg;
6831 SecondaryReg = HighReg;
6832 RegPrimaryOffset = Reg0;
6833 RegSecondaryOffset = Reg32;
6838 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
6839 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
6844 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
6845 SPIRV::OpINotEqual))
6852 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
6853 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
6858 if (SwapPrimarySide) {
6860 if (!selectOpWithSrcs(RegAdd, ResType,
I,
6861 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
6872 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
6873 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
6878 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
6879 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
6882 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
6886bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
6887 SPIRVTypeInst ResType,
6889 bool IsSigned)
const {
6891 Register OpReg =
I.getOperand(2).getReg();
6894 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
6895 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
6899 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6901 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6903 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6906 return diagnoseUnsupported(
6908 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
6912bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
6913 SPIRVTypeInst ResType,
6914 MachineInstr &
I)
const {
6916 Register OpReg =
I.getOperand(2).getReg();
6921 unsigned ExtendOpcode = SPIRV::OpUConvert;
6922 unsigned BitSetOpcode = GL::FindILsb;
6926 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6928 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6930 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6933 return diagnoseUnsupported(
I,
6934 "spv_firstbitlow only supports 16,32,64 bits.");
6938bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
6939 SPIRVTypeInst ResType,
6940 MachineInstr &
I)
const {
6944 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
6947 .
addUse(
I.getOperand(2).getReg())
6950 unsigned Alignment =
I.getOperand(3).getImm();
6964 while (!Worklist.
empty()) {
6966 switch (
T->getOpcode()) {
6967 case SPIRV::OpTypeInt:
6968 case SPIRV::OpTypeFloat:
6969 case SPIRV::OpTypePointer:
6971 case SPIRV::OpTypeVector:
6972 case SPIRV::OpTypeMatrix:
6973 case SPIRV::OpTypeArray: {
6974 Register OperandReg =
T->getOperand(1).getReg();
6978 case SPIRV::OpTypeStruct:
6979 for (
unsigned Idx = 1,
E =
T->getNumOperands(); Idx <
E; ++Idx) {
6980 Register OperandReg =
T->getOperand(Idx).getReg();
6992bool SPIRVInstructionSelector::selectAbort(MachineInstr &
I)
const {
6993 assert(
I.getNumExplicitOperands() == 2);
6995 Register MsgReg =
I.getOperand(1).getReg();
6997 assert(MsgType &&
"Message argument of llvm.spv.abort has no SPIR-V type");
7000 return diagnoseUnsupported(
7002 "llvm.spv.abort message type must be a concrete SPIR-V type (numerical "
7003 "scalar, pointer, vector, matrix, or aggregate of such types)");
7006 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7013bool SPIRVInstructionSelector::selectTrap(MachineInstr &
I)
const {
7022 uint32_t MsgVal = ~0
u;
7023 if (
I.getOpcode() == TargetOpcode::G_UBSANTRAP)
7024 MsgVal =
static_cast<uint32_t
>(
I.getOperand(0).
getImm());
7027 Register MsgReg = buildI32ConstantInEntryBlock(MsgVal,
I, MsgType);
7030 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7037bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
7038 SPIRVTypeInst ResType,
7039 MachineInstr &
I)
const {
7046 bool UseUntypedPointers =
7047 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7049 UseUntypedPointers ? SPIRV::OpUntypedVariableKHR : SPIRV::OpVariable;
7051 auto MIB =
BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(Opcode))
7054 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7058 if (UseUntypedPointers) {
7062 return diagnoseUnsupported(
7063 I,
"could not deduce the data type of an untyped variable");
7069 unsigned Alignment =
I.getOperand(2).getImm();
7076bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
7081 const MachineInstr *PrevI =
I.getPrevNode();
7083 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
7087 .
addMBB(
I.getOperand(0).getMBB())
7092 .
addMBB(
I.getOperand(0).getMBB())
7097bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
7108 const MachineInstr *NextI =
I.getNextNode();
7110 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
7116 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
7118 .
addUse(
I.getOperand(0).getReg())
7119 .
addMBB(
I.getOperand(1).getMBB())
7125bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
7126 MachineInstr &
I)
const {
7128 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
7130 const unsigned NumOps =
I.getNumOperands();
7131 for (
unsigned i = 1; i <
NumOps; i += 2) {
7132 MIB.
addUse(
I.getOperand(i + 0).getReg());
7133 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
7139bool SPIRVInstructionSelector::selectGlobalValue(
7140 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
7142 MachineIRBuilder MIRBuilder(
I);
7143 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
7146 std::string GlobalIdent;
7148 unsigned &
ID = UnnamedGlobalIDs[GV];
7150 ID = UnnamedGlobalIDs.
size();
7151 GlobalIdent =
"__unnamed_" + Twine(ID).str();
7177 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
7184 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7189 MachineInstrBuilder MIB1 =
7190 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7193 MachineInstrBuilder MIB2 =
7195 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
7199 GR.
add(ConstVal, MIB2);
7207 MachineInstrBuilder MIB3 =
7208 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7211 GR.
add(ConstVal, MIB3);
7217 assert(NewReg != ResVReg);
7218 return BuildCOPY(ResVReg, NewReg,
I);
7228 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
7231 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
7237 SPIRVTypeInst ResType =
7241 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
7246 if (
GlobalVar->isExternallyInitialized() &&
7247 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
7248 constexpr unsigned ReadWriteINTEL = 3u;
7251 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
7257bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
7258 SPIRVTypeInst ResType,
7259 MachineInstr &
I)
const {
7261 return selectExtInst(ResVReg, ResType,
I, CL::log10);
7269 MachineIRBuilder MIRBuilder(
I);
7274 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7277 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
7279 .
add(
I.getOperand(1))
7284 ResType->
getOpcode() == SPIRV::OpTypeFloat);
7294 APFloat::rmNearestTiesToEven, &LosesInfo);
7298 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
7299 ? SPIRV::OpVectorTimesScalar
7310bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
7311 SPIRVTypeInst ResType,
7312 MachineInstr &
I)
const {
7315 return selectExtInst(ResVReg, ResType,
I, CL::pown);
7321 Register ExpReg =
I.getOperand(2).getReg();
7323 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
7324 SPIRV::OpConvertSToF))
7326 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
7333bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
7334 SPIRVTypeInst ResType,
7335 MachineInstr &
I)
const {
7351 MachineIRBuilder MIRBuilder(
I);
7352 SPIRVTypeInst FloatType =
7356 FloatType, MIRBuilder, SPIRV::StorageClass::Function);
7369 MachineBasicBlock &EntryBB =
I.getMF()->
front();
7370 const bool IsUntyped =
7371 PtrType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7373 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
7374 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
7375 : SPIRV::OpVariable))
7378 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7386 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7389 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
7392 .
add(
I.getOperand(
I.getNumExplicitDefs()))
7396 Register IntegralPartReg =
I.getOperand(1).getReg();
7399 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7409 assert(
false &&
"GLSL::Modf is deprecated.");
7420bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
7421 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7422 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7423 MachineIRBuilder MIRBuilder(
I);
7424 const SPIRVTypeInst Vec3Ty =
7427 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
7439 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7443 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7449 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7456 assert(
I.getOperand(2).isReg());
7457 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
7461 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
7472bool SPIRVInstructionSelector::loadBuiltinInputID(
7473 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7474 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7475 MachineIRBuilder MIRBuilder(
I);
7477 ResType, MIRBuilder, SPIRV::StorageClass::Input);
7492 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7496 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7505SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
7506 MachineInstr &
I)
const {
7507 MachineIRBuilder MIRBuilder(
I);
7508 if (
Type->getOpcode() != SPIRV::OpTypeVector)
7518bool SPIRVInstructionSelector::loadHandleBeforePosition(
7519 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
7520 MachineInstr &Pos)
const {
7523 Intrinsic::spv_resource_handlefrombinding);
7531 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
7532 MachineIRBuilder MIRBuilder(HandleDef);
7533 SPIRVTypeInst VarType = ResType;
7534 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
7536 if (IsStructuredBuffer) {
7541 if (ResType->
getOpcode() == SPIRV::OpTypeImage && ArraySize == 0)
7543 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
7546 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
7547 ArraySize, IndexReg, Name, MIRBuilder);
7551 uint32_t LoadOpcode =
7552 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
7562bool SPIRVInstructionSelector::errorIfInstrOutsideShader(
7563 MachineInstr &
I)
const {
7565 return diagnoseUnsupported(
7566 I,
"this instruction is only supported in shaders.");
7571InstructionSelector *
7575 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)
SPIRVTypeInst getOrCreateSPIRVPointerType(const Type *BaseType, MachineIRBuilder &MIRBuilder, SPIRV::StorageClass::StorageClass SC, bool ForceTyped=false)
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)
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
SPIRVTypeInst getUntypedPtrElementType(Register Reg) 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...
SPIRV::Scope::Scope getMemScope(const Triple &TT, LLVMContext &Ctx, SyncScope::ID Id)
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.
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