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_subgroup_id:
752 case Intrinsic::spv_subgroup_local_invocation_id:
753 case Intrinsic::spv_subgroup_max_size:
754 case Intrinsic::spv_subgroup_size:
755 case Intrinsic::spv_thread_id:
756 case Intrinsic::spv_thread_id_in_group:
757 case Intrinsic::spv_udot:
758 case Intrinsic::spv_undef:
759 case Intrinsic::spv_value_md:
760 case Intrinsic::spv_workgroup_size:
772 case SPIRV::OpTypeVoid:
773 case SPIRV::OpTypeBool:
774 case SPIRV::OpTypeInt:
775 case SPIRV::OpTypeFloat:
776 case SPIRV::OpTypeVector:
777 case SPIRV::OpTypeMatrix:
778 case SPIRV::OpTypeImage:
779 case SPIRV::OpTypeSampler:
780 case SPIRV::OpTypeSampledImage:
781 case SPIRV::OpTypeArray:
782 case SPIRV::OpTypeRuntimeArray:
783 case SPIRV::OpTypeStruct:
784 case SPIRV::OpTypeOpaque:
785 case SPIRV::OpTypePointer:
786 case SPIRV::OpTypeFunction:
787 case SPIRV::OpTypeEvent:
788 case SPIRV::OpTypeDeviceEvent:
789 case SPIRV::OpTypeReserveId:
790 case SPIRV::OpTypeQueue:
791 case SPIRV::OpTypePipe:
792 case SPIRV::OpTypeForwardPointer:
793 case SPIRV::OpTypePipeStorage:
794 case SPIRV::OpTypeNamedBarrier:
795 case SPIRV::OpTypeAccelerationStructureNV:
796 case SPIRV::OpTypeCooperativeMatrixNV:
797 case SPIRV::OpTypeCooperativeMatrixKHR:
807 if (
MI.getNumDefs() == 0)
810 for (
const auto &MO :
MI.all_defs()) {
812 if (
Reg.isPhysical()) {
817 if (
UseMI.getOpcode() != SPIRV::OpName) {
824 if (
MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE ||
MI.isFakeUse() ||
825 MI.isLifetimeMarker()) {
828 <<
"Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
839 if (
MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
840 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
843 LLVM_DEBUG(
dbgs() <<
"Dead: Intrinsic with no real side effects.\n");
848 if (
MI.mayStore() ||
MI.isCall() ||
849 (
MI.mayLoad() &&
MI.hasOrderedMemoryRef()) ||
MI.isPosition() ||
850 MI.isDebugInstr() ||
MI.isTerminator() ||
MI.isJumpTableDebugInfo()) {
851 LLVM_DEBUG(
dbgs() <<
"Not dead: instruction has side effects.\n");
862 LLVM_DEBUG(
dbgs() <<
"Dead: known opcode with no side effects\n");
869void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &
MI)
const {
871 for (
const auto &MO :
MI.all_defs()) {
875 SmallVector<MachineInstr *, 4> UselessOpNames;
878 "There is still a use of the dead function.");
881 for (MachineInstr *OpNameMI : UselessOpNames) {
883 OpNameMI->eraseFromParent();
888void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &
MI)
const {
891 removeOpNamesForDeadMI(
MI);
892 MI.eraseFromParent();
895bool SPIRVInstructionSelector::select(MachineInstr &
I) {
896 resetVRegsType(*
I.getParent()->getParent());
898 assert(
I.getParent() &&
"Instruction should be in a basic block!");
899 assert(
I.getParent()->getParent() &&
"Instruction should be in a function!");
904 removeDeadInstruction(
I);
911 if (Opcode == SPIRV::ASSIGN_TYPE) {
912 Register DstReg =
I.getOperand(0).getReg();
913 Register SrcReg =
I.getOperand(1).getReg();
916 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
917 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
918 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
919 Register SelectDstReg =
Def->getOperand(0).getReg();
920 bool SuccessToSelectSelect [[maybe_unused]] = selectSelect(
922 assert(SuccessToSelectSelect);
924 Def->eraseFromParent();
931 bool Res = selectImpl(
I, *CoverageInfo);
933 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
934 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
938 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
950 }
else if (
I.getNumDefs() == 1) {
962 removeDeadInstruction(
I);
967 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
968 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
974 bool HasDefs =
I.getNumDefs() > 0;
977 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
978 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
979 if (spvSelect(ResVReg, ResType,
I)) {
981 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
992 case TargetOpcode::G_CONSTANT:
993 case TargetOpcode::G_FCONSTANT:
1000 MachineInstr &
I)
const {
1003 if (DstRC != SrcRC && SrcRC)
1005 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
1012bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
1013 SPIRVTypeInst ResType,
1014 MachineInstr &
I)
const {
1015 const unsigned Opcode =
I.getOpcode();
1017 return selectImpl(
I, *CoverageInfo);
1019 case TargetOpcode::G_CONSTANT:
1020 case TargetOpcode::G_FCONSTANT:
1021 return selectConst(ResVReg, ResType,
I);
1022 case TargetOpcode::G_GLOBAL_VALUE:
1023 return selectGlobalValue(ResVReg,
I);
1024 case TargetOpcode::G_IMPLICIT_DEF:
1025 return selectOpUndef(ResVReg, ResType,
I);
1026 case TargetOpcode::G_FREEZE:
1027 return selectFreeze(ResVReg, ResType,
I);
1029 case TargetOpcode::G_INTRINSIC:
1030 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
1031 case TargetOpcode::G_INTRINSIC_CONVERGENT:
1032 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
1033 return selectIntrinsic(ResVReg, ResType,
I);
1034 case TargetOpcode::G_BITREVERSE:
1035 return selectBitreverse(ResVReg, ResType,
I);
1037 case TargetOpcode::G_BUILD_VECTOR:
1038 return selectBuildVector(ResVReg, ResType,
I);
1039 case TargetOpcode::G_SPLAT_VECTOR:
1040 return selectSplatVector(ResVReg, ResType,
I);
1041 case TargetOpcode::G_CONCAT_VECTORS:
1042 return selectConcatVectors(ResVReg, ResType,
I);
1044 case TargetOpcode::G_SHUFFLE_VECTOR: {
1045 MachineBasicBlock &BB = *
I.getParent();
1046 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
1049 .
addUse(
I.getOperand(1).getReg())
1050 .
addUse(
I.getOperand(2).getReg());
1051 for (
auto V :
I.getOperand(3).getShuffleMask())
1056 case TargetOpcode::G_MEMMOVE:
1057 case TargetOpcode::G_MEMCPY:
1058 case TargetOpcode::G_MEMCPY_INLINE:
1059 case TargetOpcode::G_MEMSET:
1060 case TargetOpcode::G_MEMSET_INLINE:
1061 return selectMemOperation(ResVReg,
I);
1063 case TargetOpcode::G_ICMP:
1064 return selectICmp(ResVReg, ResType,
I);
1065 case TargetOpcode::G_FCMP:
1066 return selectFCmp(ResVReg, ResType,
I);
1068 case TargetOpcode::G_FRAME_INDEX:
1069 return selectFrameIndex(ResVReg, ResType,
I);
1071 case TargetOpcode::G_LOAD:
1072 return selectLoad(ResVReg, ResType,
I);
1073 case TargetOpcode::G_STORE:
1074 return selectStore(
I);
1076 case TargetOpcode::G_BR:
1077 return selectBranch(
I);
1078 case TargetOpcode::G_BRCOND:
1079 return selectBranchCond(
I);
1081 case TargetOpcode::G_PHI:
1082 return selectPhi(ResVReg,
I);
1084 case TargetOpcode::G_FPTOSI:
1085 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1086 case TargetOpcode::G_FPTOUI:
1087 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1089 case TargetOpcode::G_FPTOSI_SAT:
1090 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1091 case TargetOpcode::G_FPTOUI_SAT:
1092 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1094 case TargetOpcode::G_SITOFP:
1095 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
1096 case TargetOpcode::G_UITOFP:
1097 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
1099 case TargetOpcode::G_CTPOP:
1100 return selectPopCount(ResVReg, ResType,
I, SPIRV::OpBitCount);
1101 case TargetOpcode::G_SMIN:
1102 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
1103 case TargetOpcode::G_UMIN:
1104 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
1106 case TargetOpcode::G_SMAX:
1107 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
1108 case TargetOpcode::G_UMAX:
1109 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
1111 case TargetOpcode::G_SCMP:
1112 return selectSUCmp(ResVReg, ResType,
I,
true);
1113 case TargetOpcode::G_UCMP:
1114 return selectSUCmp(ResVReg, ResType,
I,
false);
1115 case TargetOpcode::G_LROUND:
1116 case TargetOpcode::G_LLROUND: {
1119 MRI->
setRegClass(regForLround, &SPIRV::iIDRegClass);
1121 regForLround, *(
I.getParent()->getParent()));
1123 CL::round, GL::Round,
false);
1125 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1132 case TargetOpcode::G_STRICT_FMA:
1133 case TargetOpcode::G_FMA: {
1136 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1139 .
addUse(
I.getOperand(1).getReg())
1140 .
addUse(
I.getOperand(2).getReg())
1141 .
addUse(
I.getOperand(3).getReg())
1146 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1149 case TargetOpcode::G_FLDEXP:
1150 case TargetOpcode::G_STRICT_FLDEXP:
1151 return selectLdexp(ResVReg, ResType,
I);
1153 case TargetOpcode::G_FPOW:
1154 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1155 case TargetOpcode::G_FPOWI:
1156 return selectFpowi(ResVReg, ResType,
I);
1158 case TargetOpcode::G_FEXP:
1159 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1160 case TargetOpcode::G_FEXP2:
1161 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1162 case TargetOpcode::G_FEXP10:
1163 return selectExp10(ResVReg, ResType,
I);
1165 case TargetOpcode::G_FMODF:
1166 return selectModf(ResVReg, ResType,
I);
1167 case TargetOpcode::G_FSINCOS:
1168 return selectSincos(ResVReg, ResType,
I);
1170 case TargetOpcode::G_FLOG:
1171 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1172 case TargetOpcode::G_FLOG2:
1173 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1174 case TargetOpcode::G_FLOG10:
1175 return selectLog10(ResVReg, ResType,
I);
1177 case TargetOpcode::G_FABS:
1178 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1179 case TargetOpcode::G_ABS:
1180 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1182 case TargetOpcode::G_FMINNUM:
1183 case TargetOpcode::G_FMINIMUM:
1184 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1185 case TargetOpcode::G_FMAXNUM:
1186 case TargetOpcode::G_FMAXIMUM:
1187 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1189 case TargetOpcode::G_FCOPYSIGN:
1190 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1192 case TargetOpcode::G_FCEIL:
1193 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1194 case TargetOpcode::G_FFLOOR:
1195 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1197 case TargetOpcode::G_FCOS:
1198 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1199 case TargetOpcode::G_FSIN:
1200 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1201 case TargetOpcode::G_FTAN:
1202 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1203 case TargetOpcode::G_FACOS:
1204 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1205 case TargetOpcode::G_FASIN:
1206 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1207 case TargetOpcode::G_FATAN:
1208 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1209 case TargetOpcode::G_FATAN2:
1210 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1211 case TargetOpcode::G_FCOSH:
1212 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1213 case TargetOpcode::G_FSINH:
1214 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1215 case TargetOpcode::G_FTANH:
1216 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1218 case TargetOpcode::G_STRICT_FSQRT:
1219 case TargetOpcode::G_FSQRT:
1220 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1222 case TargetOpcode::G_CTTZ:
1223 case TargetOpcode::G_CTTZ_ZERO_POISON:
1224 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1225 case TargetOpcode::G_CTLZ:
1226 case TargetOpcode::G_CTLZ_ZERO_POISON:
1227 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1229 case TargetOpcode::G_INTRINSIC_ROUND:
1230 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1231 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1232 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1233 case TargetOpcode::G_INTRINSIC_TRUNC:
1234 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1235 case TargetOpcode::G_FRINT:
1236 case TargetOpcode::G_FNEARBYINT:
1237 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1239 case TargetOpcode::G_SMULH:
1240 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1241 case TargetOpcode::G_UMULH:
1242 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1244 case TargetOpcode::G_SADDSAT:
1245 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1246 case TargetOpcode::G_UADDSAT:
1247 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1248 case TargetOpcode::G_SSUBSAT:
1249 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1250 case TargetOpcode::G_USUBSAT:
1251 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1253 case TargetOpcode::G_FFREXP:
1254 return selectFrexp(ResVReg, ResType,
I);
1256 case TargetOpcode::G_UADDO:
1257 return selectOverflowArith(ResVReg, ResType,
I,
1258 ResType->
getOpcode() == SPIRV::OpTypeVector
1259 ? SPIRV::OpIAddCarryV
1260 : SPIRV::OpIAddCarryS);
1261 case TargetOpcode::G_USUBO:
1262 return selectOverflowArith(ResVReg, ResType,
I,
1263 ResType->
getOpcode() == SPIRV::OpTypeVector
1264 ? SPIRV::OpISubBorrowV
1265 : SPIRV::OpISubBorrowS);
1266 case TargetOpcode::G_UMULO:
1267 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1268 case TargetOpcode::G_SMULO:
1269 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1271 case TargetOpcode::G_SEXT:
1272 return selectExt(ResVReg, ResType,
I,
true);
1273 case TargetOpcode::G_ANYEXT:
1274 case TargetOpcode::G_ZEXT:
1275 return selectExt(ResVReg, ResType,
I,
false);
1276 case TargetOpcode::G_TRUNC:
1277 return selectTrunc(ResVReg, ResType,
I);
1278 case TargetOpcode::G_FPTRUNC:
1279 case TargetOpcode::G_FPEXT:
1280 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1282 case TargetOpcode::G_PTRTOINT:
1283 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1284 case TargetOpcode::G_INTTOPTR:
1285 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1286 case TargetOpcode::G_BITCAST:
1287 return selectBitcast(ResVReg, ResType,
I);
1288 case TargetOpcode::G_ADDRSPACE_CAST:
1289 return selectAddrSpaceCast(ResVReg, ResType,
I);
1290 case TargetOpcode::G_PTRMASK:
1291 return selectPtrMask(ResVReg, ResType,
I);
1292 case TargetOpcode::G_PTR_ADD: {
1294 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1298 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1299 (*II).getOpcode() == TargetOpcode::COPY ||
1300 (*II).getOpcode() == SPIRV::OpVariable ||
1301 (*II).getOpcode() == SPIRV::OpUntypedVariableKHR) &&
1302 getImm(
I.getOperand(2), MRI));
1304 bool IsGVInit =
false;
1308 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1309 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1310 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1311 (*UseIt).getOpcode() == SPIRV::OpVariable ||
1312 (*UseIt).getOpcode() == SPIRV::OpUntypedVariableKHR) {
1324 const bool UseUntypedPointers =
1325 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1326 if (UseUntypedPointers) {
1327 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1330 .
addImm(
static_cast<uint32_t
>(
1331 SPIRV::Opcode::UntypedInBoundsPtrAccessChainKHR))
1334 .
addUse(
I.getOperand(2).getReg())
1341 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1353 return diagnoseUnsupported(
1354 I,
"incompatible result and operand types in a bitcast");
1356 MachineInstrBuilder MIB =
1357 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1364 : SPIRV::OpInBoundsPtrAccessChain))
1368 .
addUse(
I.getOperand(2).getReg())
1371 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1375 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1377 .
addUse(
I.getOperand(2).getReg())
1386 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1389 .
addImm(
static_cast<uint32_t
>(
1390 SPIRV::Opcode::InBoundsPtrAccessChain))
1393 .
addUse(
I.getOperand(2).getReg());
1398 case TargetOpcode::G_ATOMICRMW_OR:
1399 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1400 case TargetOpcode::G_ATOMICRMW_ADD:
1401 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1402 case TargetOpcode::G_ATOMICRMW_AND:
1403 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1404 case TargetOpcode::G_ATOMICRMW_MAX:
1405 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1406 case TargetOpcode::G_ATOMICRMW_MIN:
1407 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1408 case TargetOpcode::G_ATOMICRMW_SUB:
1409 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1410 case TargetOpcode::G_ATOMICRMW_XOR:
1411 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1412 case TargetOpcode::G_ATOMICRMW_UMAX:
1413 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1414 case TargetOpcode::G_ATOMICRMW_UMIN:
1415 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1416 case TargetOpcode::G_ATOMICRMW_XCHG:
1417 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1419 case TargetOpcode::G_ATOMICRMW_FADD:
1420 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1421 case TargetOpcode::G_ATOMICRMW_FSUB:
1423 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1424 ResType->
getOpcode() == SPIRV::OpTypeVector
1426 : SPIRV::OpFNegate);
1427 case TargetOpcode::G_ATOMICRMW_FMIN:
1428 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1429 case TargetOpcode::G_ATOMICRMW_FMAX:
1430 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1432 case TargetOpcode::G_FENCE:
1433 return selectFence(
I);
1435 case TargetOpcode::G_STACKSAVE:
1436 return selectStackSave(ResVReg, ResType,
I);
1437 case TargetOpcode::G_STACKRESTORE:
1438 return selectStackRestore(
I);
1440 case TargetOpcode::G_UNMERGE_VALUES:
1443 case TargetOpcode::G_TRAP:
1444 case TargetOpcode::G_UBSANTRAP:
1445 return selectTrap(
I);
1450 case TargetOpcode::DBG_LABEL:
1452 case TargetOpcode::G_DEBUGTRAP:
1453 return selectDebugTrap(ResVReg, ResType,
I);
1460bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1461 SPIRVTypeInst ResType,
1462 MachineInstr &
I)
const {
1463 unsigned Opcode = SPIRV::OpNop;
1470bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1471 SPIRVTypeInst ResType,
1473 GL::GLSLExtInst GLInst,
1474 bool setMIFlags,
bool useMISrc,
1477 SPIRV::InstructionSet::InstructionSet::GLSL_std_450))
1478 return diagnoseUnsupported(
1480 "this instruction is only supported with the GLSL extended instruction "
1482 return selectExtInst(ResVReg, ResType,
I,
1483 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1484 setMIFlags, useMISrc, SrcRegs);
1487bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1488 SPIRVTypeInst ResType,
1490 CL::OpenCLExtInst CLInst,
1491 bool setMIFlags,
bool useMISrc,
1493 return selectExtInst(ResVReg, ResType,
I,
1494 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1495 setMIFlags, useMISrc, SrcRegs);
1498bool SPIRVInstructionSelector::selectExtInst(
1499 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1500 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1502 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1503 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1504 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1508bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1509 SPIRVTypeInst ResType,
1512 bool setMIFlags,
bool useMISrc,
1515 for (
const auto &[InstructionSet, Opcode] : Insts) {
1519 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1522 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1527 const unsigned NumOps =
I.getNumOperands();
1530 I.getOperand(Index).getType() ==
1531 MachineOperand::MachineOperandType::MO_IntrinsicID)
1534 MIB.
add(
I.getOperand(Index));
1546bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1547 SPIRVTypeInst ResType,
1548 MachineInstr &
I)
const {
1549 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1550 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1551 for (
const auto &Ex : ExtInsts) {
1552 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1553 uint32_t Opcode = Ex.second;
1557 MachineIRBuilder MIRBuilder(
I);
1560 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1567 const bool IsUntyped =
1568 PointerType->getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1570 BuildMI(*It->getParent(), It, It->getDebugLoc(),
1571 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
1572 : SPIRV::OpVariable))
1575 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
1581 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1584 .
addImm(
static_cast<uint32_t
>(Ex.first))
1586 .
add(
I.getOperand(2))
1590 Register ExpResReg =
I.getOperand(1).getReg();
1592 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1602bool SPIRVInstructionSelector::selectLdexp(
Register ResVReg,
1603 SPIRVTypeInst ResType,
1604 MachineInstr &
I)
const {
1605 Register XReg =
I.getOperand(1).getReg();
1606 Register ExpReg =
I.getOperand(2).getReg();
1612 if (ResType->
getOpcode() == SPIRV::OpTypeVector &&
1613 ExpType->
getOpcode() != SPIRV::OpTypeVector) {
1615 SPIRVTypeInst ExpVecType =
1619 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1620 TII.get(SPIRV::OpCompositeConstruct))
1623 for (
unsigned J = 0; J < NumElts; ++J)
1629 return selectExtInst(ResVReg, ResType,
I, CL::ldexp, GL::Ldexp,
1630 true,
false, {XReg, ExpReg});
1633bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1634 SPIRVTypeInst ResType,
1635 MachineInstr &
I)
const {
1636 Register CosResVReg =
I.getOperand(1).getReg();
1637 unsigned SrcIdx =
I.getNumExplicitDefs();
1642 MachineIRBuilder MIRBuilder(
I);
1644 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1651 const bool IsUntyped =
1652 PointerType->getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1654 BuildMI(*It->getParent(), It, It->getDebugLoc(),
1655 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
1656 : SPIRV::OpVariable))
1659 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
1663 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1666 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1668 .
add(
I.getOperand(SrcIdx))
1671 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1679 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1682 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1684 .
add(
I.getOperand(SrcIdx))
1686 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1689 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1691 .
add(
I.getOperand(SrcIdx))
1698bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1699 SPIRVTypeInst ResType,
1702 unsigned Opcode)
const {
1703 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1713std::optional<SplitParts> SPIRVInstructionSelector::splitEvenOddLanes(
1714 Register PopCountReg,
unsigned ComponentCount, MachineInstr &
I,
1715 SPIRVTypeInst I32Type)
const {
1718 if (ComponentCount == 1) {
1721 Parts.IsScalar =
true;
1722 Parts.Type = I32Type;
1730 if (!selectOpWithSrcs(Parts.High, I32Type,
I, {PopCountReg, IdxOne},
1731 SPIRV::OpVectorExtractDynamic))
1732 return std::nullopt;
1734 if (!selectOpWithSrcs(Parts.Low, I32Type,
I, {PopCountReg, IdxZero},
1735 SPIRV::OpVectorExtractDynamic))
1736 return std::nullopt;
1740 MachineIRBuilder MIRBuilder(
I);
1741 Parts.IsScalar =
false;
1748 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1749 TII.get(SPIRV::OpVectorShuffle))
1754 for (
unsigned J = 1; J < ComponentCount * 2; J += 2)
1759 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1760 TII.get(SPIRV::OpVectorShuffle))
1765 for (
unsigned J = 0; J < ComponentCount * 2; J += 2)
1773bool SPIRVInstructionSelector::selectPopCount16(
Register ResVReg,
1774 SPIRVTypeInst ResType,
1777 unsigned Opcode)
const {
1778 Register OpReg =
I.getOperand(1).getReg();
1781 MachineIRBuilder MIRBuilder(
I);
1783 SPIRVTypeInst I32VectorType =
1786 bool IsVector = NumElems > 1;
1787 SPIRVTypeInst ExtType = IsVector ? I32VectorType : I32Type;
1790 if (!selectOpWithSrcs(ExtReg, ExtType,
I, {OpReg}, SPIRV::OpUConvert))
1794 if (!selectPopCount32(PopCountReg, ExtType,
I, ExtReg, Opcode))
1797 return selectOpWithSrcs(ResVReg, ResType,
I, {PopCountReg}, ExtOpcode);
1800bool SPIRVInstructionSelector::selectPopCount32(
Register ResVReg,
1801 SPIRVTypeInst ResType,
1804 unsigned Opcode)
const {
1805 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
1808bool SPIRVInstructionSelector::selectPopCount64(
Register ResVReg,
1809 SPIRVTypeInst ResType,
1812 unsigned Opcode)
const {
1814 if (ComponentCount > 2)
1815 return handle64BitOverflow(
1816 ResVReg, ResType,
I, SrcReg, Opcode,
1818 unsigned O) {
return this->selectPopCount64(R,
T,
I, S, O); });
1820 MachineIRBuilder MIRBuilder(
I);
1825 I32Type, 2 * ComponentCount, MIRBuilder,
false);
1829 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
1834 if (!selectPopCount32(Pop32, VecI32Type,
I, Vec32, Opcode))
1838 auto MaybeParts = splitEvenOddLanes(Pop32, ComponentCount,
I, I32Type);
1841 SplitParts &Parts = *MaybeParts;
1844 unsigned OpAdd = Parts.IsScalar ? SPIRV::OpIAddS : SPIRV::OpIAddV;
1846 if (!selectOpWithSrcs(Sum, Parts.Type,
I, {Parts.High, Parts.Low}, OpAdd))
1851 unsigned ConvOp = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
1852 return selectOpWithSrcs(ResVReg, ResType,
I, {Sum}, ConvOp);
1855bool SPIRVInstructionSelector::selectPopCount(
Register ResVReg,
1856 SPIRVTypeInst ResType,
1858 unsigned Opcode)
const {
1863 if (!STI.getTargetTriple().isVulkanOS())
1864 return selectUnOp(ResVReg, ResType,
I, Opcode);
1866 Register OpReg =
I.getOperand(1).getReg();
1869 : SPIRV::OpUConvert;
1873 return selectPopCount16(ResVReg, ResType,
I, ExtOpcode, Opcode);
1875 return selectPopCount32(ResVReg, ResType,
I, OpReg, Opcode);
1877 return selectPopCount64(ResVReg, ResType,
I, OpReg, Opcode);
1879 return diagnoseUnsupported(
I,
"unsupported operand bit width for popcount");
1883bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1884 SPIRVTypeInst ResType,
1886 unsigned Opcode)
const {
1888 Register SrcReg =
I.getOperand(1).getReg();
1893 unsigned DefOpCode = DefIt->getOpcode();
1894 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1897 if (
auto *VRD =
getVRegDef(*MRI, DefIt->getOperand(1).getReg()))
1898 DefOpCode = VRD->getOpcode();
1900 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1901 DefOpCode == TargetOpcode::G_CONSTANT ||
1902 DefOpCode == SPIRV::OpVariable ||
1903 DefOpCode == SPIRV::OpUntypedVariableKHR ||
1904 DefOpCode == SPIRV::OpConstantI) {
1910 uint32_t SpecOpcode = 0;
1912 case SPIRV::OpConvertPtrToU:
1913 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1915 case SPIRV::OpConvertUToPtr:
1916 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1921 TII.get(SPIRV::OpSpecConstantOp))
1931 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1935bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1936 SPIRVTypeInst ResType,
1937 MachineInstr &
I)
const {
1938 Register OpReg =
I.getOperand(1).getReg();
1939 SPIRVTypeInst OpType =
1942 return diagnoseUnsupported(
1943 I,
"incompatible result and operand types in a bitcast");
1944 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1955 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1956 if (
MemOp->isNonTemporal())
1957 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1959 if (!ST->isShader() &&
MemOp->getAlign().value())
1960 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1964 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1965 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1969 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1971 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1975 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1979 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1981 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1993 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1995 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1997 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
2001bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
2002 SPIRVTypeInst ResType,
2003 MachineInstr &
I)
const {
2005 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2010 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2011 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2013 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2015 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2019 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2023 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2024 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
2025 I.getDebugLoc(),
I);
2029 MachineIRBuilder MIRBuilder(
I);
2031 if (
I.getNumMemOperands()) {
2032 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2033 if (MemOp->isAtomic())
2034 return selectAtomicLoad(ResVReg, ResType,
I);
2037 auto MIB = MIRBuilder.
buildInstr(SPIRV::OpLoad)
2041 if (!
I.getNumMemOperands()) {
2042 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2044 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2053Register SPIRVInstructionSelector::createPtrSizedIntReg(
2054 MachineIRBuilder &MIRBuilder)
const {
2055 SPIRVTypeInst IntType =
2065SPIRVInstructionSelector::convertPtrToInt(
Register PtrVal,
2066 MachineIRBuilder &MIRBuilder)
const {
2067 SPIRVTypeInst IntType =
2069 Register IntReg = createPtrSizedIntReg(MIRBuilder);
2070 MIRBuilder.
buildInstr(SPIRV::OpConvertPtrToU)
2078Register SPIRVInstructionSelector::castPtrToPtrToInt(
2079 Register Ptr, SPIRV::StorageClass::StorageClass SC,
2080 MachineIRBuilder &MIRBuilder)
const {
2081 SPIRVTypeInst IntType =
2083 SPIRVTypeInst PtrType =
2097bool SPIRVInstructionSelector::selectAtomicPtrValue(
2098 Register ResVReg, SPIRVTypeInst ResType, MachineIRBuilder &MIRBuilder,
2099 function_ref<
Register(SPIRVTypeInst IntType)> EmitAtomic)
const {
2107 Register IntResult = EmitAtomic(IntType);
2109 MIRBuilder.
buildInstr(SPIRV::OpConvertUToPtr)
2117bool SPIRVInstructionSelector::selectAtomicLoad(
Register ResVReg,
2118 SPIRVTypeInst ResType,
2119 MachineInstr &
I)
const {
2120 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2123 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2126 return diagnoseUnsupported(
2127 I,
"Lowering to SPIR-V of atomic load is only "
2128 "allowed for integer, floating point or pointer types");
2130 assert(
I.getNumMemOperands());
2131 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2132 assert(MemOp.isAtomic());
2134 uint32_t
Scope =
static_cast<uint32_t
>(
2136 Register ScopeReg = buildI32Constant(Scope,
I);
2142 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2143 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2146 MachineIRBuilder MIRBuilder(
I);
2150 return diagnoseUnsupported(
2151 I,
"Lowering to SPIR-V of atomic load is only "
2152 "allowed for pointer types for physical addressing model");
2157 SPIRV::StorageClass::StorageClass SC =
2159 return selectAtomicPtrValue(
2160 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2161 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2162 Register IntResult = createPtrSizedIntReg(MIRBuilder);
2173 auto AtomicLoad = MIRBuilder.
buildInstr(SPIRV::OpAtomicLoad)
2184bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
2186 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2187 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2192 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2193 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2195 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2200 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2204 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2205 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2206 SPIRVTypeInst SampledType =
2208 SPIRVTypeInst StoreValCompType =
2210 if (StoreValCompType && StoreValCompType != SampledType) {
2213 SPIRVTypeInst PackedType = widenTypeToVec4(SampledType,
I);
2216 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
2221 StoreVal = PackedReg;
2224 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2225 TII.get(SPIRV::OpImageWrite))
2231 if (sampledTypeIsSignedInteger(LLVMHandleType))
2234 BMI.constrainAllUses(
TII,
TRI, RBI);
2239 if (
I.getNumMemOperands()) {
2240 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2241 if (MemOp->isAtomic())
2242 return selectAtomicStore(
I);
2249 if (PtrSC == SPIRV::StorageClass::UniformConstant ||
2250 PtrSC == SPIRV::StorageClass::Input ||
2251 PtrSC == SPIRV::StorageClass::PushConstant)
2252 return diagnoseUnsupported(
2253 I,
"store into a read-only SPIR-V storage class is not allowed");
2255 MachineIRBuilder MIRBuilder(
I);
2257 if (!
I.getNumMemOperands()) {
2258 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2260 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2269bool SPIRVInstructionSelector::selectAtomicStore(MachineInstr &
I)
const {
2270 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2273 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2274 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2279 if (!PointeeType && PtrType &&
2280 PtrType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR)
2283 return diagnoseUnsupported(
I,
2284 "Lowering to SPIR-V of atomic store is only "
2285 "allowed for integer or floating point types");
2287 assert(
I.getNumMemOperands());
2288 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2289 assert(MemOp.isAtomic());
2291 uint32_t
Scope =
static_cast<uint32_t
>(
2293 Register ScopeReg = buildI32Constant(Scope,
I);
2299 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2300 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2302 MachineIRBuilder MIRBuilder(
I);
2306 return diagnoseUnsupported(
2307 I,
"Lowering to SPIR-V of atomic store is only "
2308 "allowed for pointer types for physical addressing model");
2313 SPIRV::StorageClass::StorageClass SC =
2315 return selectAtomicPtrValue(
2316 Register(), SPIRVTypeInst(), MIRBuilder, [&](SPIRVTypeInst IntType) {
2318 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2331 return diagnoseUnsupported(
I,
2332 "Lowering to SPIR-V of atomic store is only "
2333 "allowed for integer or floating point types");
2335 auto AtomicStore = MIRBuilder.
buildInstr(SPIRV::OpAtomicStore)
2345bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
2346 SPIRVTypeInst ResType,
2347 MachineInstr &
I)
const {
2348 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
2356 const Register PtrsReg =
I.getOperand(2).getReg();
2357 const uint32_t Alignment =
I.getOperand(3).getImm();
2358 const Register MaskReg =
I.getOperand(4).getReg();
2359 const Register PassthruReg =
I.getOperand(5).getReg();
2360 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2364 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
2375bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
2376 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
2383 const Register ValuesReg =
I.getOperand(1).getReg();
2384 const Register PtrsReg =
I.getOperand(2).getReg();
2385 const uint32_t Alignment =
I.getOperand(3).getImm();
2386 const Register MaskReg =
I.getOperand(4).getReg();
2387 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2391 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
2400bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
2401 const Twine &
Msg)
const {
2402 const Function &
F =
I.getMF()->getFunction();
2403 F.getContext().diagnose(
2404 DiagnosticInfoUnsupported(
F,
Msg,
I.getDebugLoc(),
DS_Error));
2408bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
2409 SPIRVTypeInst ResType,
2410 MachineInstr &
I)
const {
2411 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2412 return diagnoseUnsupported(
2413 I,
"llvm.stacksave intrinsic: this instruction requires the following "
2414 "SPIR-V extension: SPV_INTEL_variable_length_array");
2416 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
2423bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
2424 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2425 return diagnoseUnsupported(
2427 "llvm.stackrestore intrinsic: this instruction requires the following "
2428 "SPIR-V extension: SPV_INTEL_variable_length_array");
2429 if (!
I.getOperand(0).isReg())
2432 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
2433 .
addUse(
I.getOperand(0).getReg())
2439SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
2440 MachineIRBuilder MIRBuilder(
I);
2441 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
2448 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
2452 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2453 Type *ArrTy = ArrayType::get(ValTy, Num);
2455 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
2458 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
2469 const bool IsUntyped = VarTy->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
2470 auto MIBVar =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2471 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
2472 : SPIRV::OpVariable))
2475 .
addImm(SPIRV::StorageClass::UniformConstant);
2488bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
2491 Register DstReg =
I.getOperand(0).getReg();
2495 return diagnoseUnsupported(
2496 I,
"OpCopyMemory requires operands to have the same type");
2497 uint64_t CopySize =
getIConstVal(
I.getOperand(2).getReg(), MRI);
2501 return diagnoseUnsupported(
2502 I,
"Unable to determine pointee type size for OpCopyMemory");
2503 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
2504 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
2505 return diagnoseUnsupported(
2506 I,
"OpCopyMemory requires the size to match the pointee type size");
2507 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemory))
2510 if (
I.getNumMemOperands()) {
2511 MachineIRBuilder MIRBuilder(
I);
2518bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &
I,
2521 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemorySized))
2522 .
addUse(
I.getOperand(0).getReg())
2524 .
addUse(
I.getOperand(2).getReg());
2525 if (
I.getNumMemOperands()) {
2526 MachineIRBuilder MIRBuilder(
I);
2533bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
2534 MachineInstr &
I)
const {
2536 Register SizeReg =
I.getOperand(2).getReg();
2538 SizeDef && SizeDef->
getOpcode() == TargetOpcode::G_CONSTANT &&
2542 Register SrcReg =
I.getOperand(1).getReg();
2543 if (
I.getOpcode() == TargetOpcode::G_MEMSET ||
2544 I.getOpcode() == TargetOpcode::G_MEMSET_INLINE) {
2545 Register VarReg = getOrCreateMemSetGlobal(
I);
2548 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2550 ValTy,
I, SPIRV::StorageClass::UniformConstant);
2552 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
2556 if (!selectCopyMemory(
I, SrcReg))
2559 if (!selectCopyMemorySized(
I, SrcReg))
2562 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
2563 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
2568bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
2569 SPIRVTypeInst ResType,
2572 unsigned NegateOpcode)
const {
2574 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2575 uint32_t
Scope =
static_cast<uint32_t
>(
2577 MemOp->getSyncScopeID()));
2578 Register ScopeReg = buildI32Constant(Scope,
I);
2580 Register Ptr =
I.getOperand(1).getReg();
2581 uint32_t ScSem =
static_cast<uint32_t
>(
2585 Register MemSemReg = buildI32Constant(MemSem,
I);
2587 Register ValueReg =
I.getOperand(2).getReg();
2588 if (NegateOpcode != 0) {
2591 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
2597 if (NewOpcode != SPIRV::OpAtomicExchange)
2598 return diagnoseUnsupported(
2599 I,
"Lowering to SPIR-V of this atomic operation is not "
2600 "allowed for pointer types");
2602 return diagnoseUnsupported(
2603 I,
"Lowering to SPIR-V of atomic exchange is only "
2604 "allowed for pointer types for physical addressing model");
2611 MachineIRBuilder MIRBuilder(
I);
2613 return selectAtomicPtrValue(
2614 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2616 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2617 Register ExchangeResReg = createPtrSizedIntReg(MIRBuilder);
2618 MIRBuilder.
buildInstr(SPIRV::OpAtomicExchange)
2626 return ExchangeResReg;
2630 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2641bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2642 unsigned ArgI =
I.getNumOperands() - 1;
2644 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2645 SPIRVTypeInst SrcType =
2647 if (!SrcType || SrcType->
getOpcode() != SPIRV::OpTypeVector)
2649 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2653 unsigned CurrentIndex = 0;
2654 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2655 Register ResVReg =
I.getOperand(i).getReg();
2658 LLT ResLLT = MRI->
getType(ResVReg);
2664 ResType = ScalarType;
2670 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
2673 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2679 for (
unsigned j = 0;
j < NumElements; ++
j) {
2680 MIB.
addImm(CurrentIndex + j);
2682 CurrentIndex += NumElements;
2686 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2698bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2701 Register MemSemReg = buildI32Constant(MemSem,
I);
2705 Register ScopeReg = buildI32Constant(Scope,
I);
2707 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2714bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2715 SPIRVTypeInst ResType,
2717 unsigned Opcode)
const {
2718 Type *ResTy =
nullptr;
2721 return diagnoseUnsupported(
2723 "Not enough info to select the arithmetic with overflow instruction");
2725 return diagnoseUnsupported(
I,
2726 "Expect struct type result for the arithmetic "
2727 "with overflow instruction");
2733 MachineIRBuilder MIRBuilder(
I);
2735 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2736 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2742 Register ZeroReg = buildZerosVal(ResType,
I);
2747 if (ResName.
size() > 0)
2755 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2756 MIB.
addUse(
I.getOperand(i).getReg());
2761 MRI->
setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2762 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2764 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2765 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2772 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2773 .
addDef(
I.getOperand(1).getReg())
2781bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2782 SPIRVTypeInst ResType,
2783 MachineInstr &
I)
const {
2785 "selectAtomicCmpXchg only handles the spv_cmpxchg intrinsic");
2786 Register Ptr =
I.getOperand(2).getReg();
2787 Register ScopeReg =
I.getOperand(5).getReg();
2788 Register MemSemEqReg =
I.getOperand(6).getReg();
2789 Register MemSemNeqReg =
I.getOperand(7).getReg();
2791 Register Val =
I.getOperand(4).getReg();
2795 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2814 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2821 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2833 case SPIRV::StorageClass::DeviceOnlyINTEL:
2834 case SPIRV::StorageClass::HostOnlyINTEL:
2843 bool IsGRef =
false;
2844 bool IsAllowedRefs =
2846 unsigned Opcode = It.getOpcode();
2847 if (Opcode == SPIRV::OpConstantComposite ||
2848 Opcode == SPIRV::OpSpecConstantComposite ||
2849 Opcode == SPIRV::OpVariable ||
2850 Opcode == SPIRV::OpUntypedVariableKHR ||
2851 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2852 return IsGRef = true;
2853 return Opcode == SPIRV::OpName;
2855 return IsAllowedRefs && IsGRef;
2858Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2859 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2861 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2865SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2867 uint32_t Opcode)
const {
2868 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2869 TII.get(SPIRV::OpSpecConstantOp))
2877SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2878 SPIRVTypeInst SrcPtrTy)
const {
2879 SPIRVTypeInst GenericPtrTy =
2883 SPIRV::StorageClass::Generic),
2885 MachineFunction *MF =
I.getParent()->getParent();
2887 MachineInstrBuilder MIB = buildSpecConstantOp(
2889 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2899bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2900 SPIRVTypeInst ResType,
2901 MachineInstr &
I)
const {
2905 Register SrcPtr =
I.getOperand(1).getReg();
2910 return BuildCOPY(ResVReg, SrcPtr,
I);
2920 unsigned SpecOpcode = [&]() ->
unsigned {
2921 if (SrcSC == SPIRV::StorageClass::CodeSectionINTEL)
2922 return static_cast<uint32_t
>(SPIRV::Opcode::Bitcast);
2924 return static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric);
2926 return static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr);
2934 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2936 .constrainAllUses(
TII,
TRI, RBI);
2938 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2940 buildSpecConstantOp(
2942 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2943 .constrainAllUses(
TII,
TRI, RBI);
2950 return BuildCOPY(ResVReg, SrcPtr,
I);
2952 if ((SrcSC == SPIRV::StorageClass::Function &&
2953 DstSC == SPIRV::StorageClass::Private) ||
2954 (DstSC == SPIRV::StorageClass::Function &&
2955 SrcSC == SPIRV::StorageClass::Private))
2956 return BuildCOPY(ResVReg, SrcPtr,
I);
2960 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2963 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2966 SPIRVTypeInst GenericPtrTy =
2985 return selectUnOp(ResVReg, ResType,
I,
2986 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
2988 return selectUnOp(ResVReg, ResType,
I,
2989 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
2991 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2993 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
3003bool SPIRVInstructionSelector::selectPtrMask(
Register ResVReg,
3004 SPIRVTypeInst ResType,
3005 MachineInstr &
I)
const {
3007 return diagnoseUnsupported(
3008 I,
"G_PTRMASK is not supported with logical SPIR-V");
3013 Register PtrReg =
I.getOperand(1).getReg();
3014 Register MaskReg =
I.getOperand(2).getReg();
3033 ? SPIRV::OpBitwiseAndV
3034 : SPIRV::OpBitwiseAndS;
3057 return SPIRV::OpFOrdEqual;
3059 return SPIRV::OpFOrdGreaterThanEqual;
3061 return SPIRV::OpFOrdGreaterThan;
3063 return SPIRV::OpFOrdLessThanEqual;
3065 return SPIRV::OpFOrdLessThan;
3067 return SPIRV::OpFOrdNotEqual;
3069 return SPIRV::OpOrdered;
3071 return SPIRV::OpFUnordEqual;
3073 return SPIRV::OpFUnordGreaterThanEqual;
3075 return SPIRV::OpFUnordGreaterThan;
3077 return SPIRV::OpFUnordLessThanEqual;
3079 return SPIRV::OpFUnordLessThan;
3081 return SPIRV::OpFUnordNotEqual;
3083 return SPIRV::OpUnordered;
3093 return SPIRV::OpIEqual;
3095 return SPIRV::OpINotEqual;
3097 return SPIRV::OpSGreaterThanEqual;
3099 return SPIRV::OpSGreaterThan;
3101 return SPIRV::OpSLessThanEqual;
3103 return SPIRV::OpSLessThan;
3105 return SPIRV::OpUGreaterThanEqual;
3107 return SPIRV::OpUGreaterThan;
3109 return SPIRV::OpULessThanEqual;
3111 return SPIRV::OpULessThan;
3120 return SPIRV::OpPtrEqual;
3122 return SPIRV::OpPtrNotEqual;
3133 return SPIRV::OpLogicalEqual;
3135 return SPIRV::OpLogicalNotEqual;
3173bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
3174 SPIRVTypeInst ResType,
3176 unsigned OpAnyOrAll)
const {
3177 assert(
I.getNumOperands() == 3);
3178 assert(
I.getOperand(2).isReg());
3180 Register InputRegister =
I.getOperand(2).getReg();
3183 assert(InputType &&
"VReg has no type assigned");
3186 bool IsVectorTy = InputType->
getOpcode() == SPIRV::OpTypeVector;
3187 if (IsBoolTy && !IsVectorTy) {
3188 assert(ResVReg ==
I.getOperand(0).getReg());
3189 return BuildCOPY(ResVReg, InputRegister,
I);
3193 unsigned SpirvNotEqualId =
3194 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
3196 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
3201 IsBoolTy ? InputRegister
3209 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
3211 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
3228bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
3229 SPIRVTypeInst ResType,
3230 MachineInstr &
I)
const {
3231 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
3234bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
3235 SPIRVTypeInst ResType,
3236 MachineInstr &
I)
const {
3237 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
3241bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
3242 SPIRVTypeInst ResType,
3243 MachineInstr &
I)
const {
3244 assert(
I.getNumOperands() == 4);
3245 assert(
I.getOperand(2).isReg());
3246 assert(
I.getOperand(3).isReg());
3248 [[maybe_unused]] SPIRVTypeInst VecType =
3253 "dot product requires a vector of at least 2 components");
3255 [[maybe_unused]] SPIRVTypeInst EltType =
3264 .
addUse(
I.getOperand(2).getReg())
3265 .
addUse(
I.getOperand(3).getReg())
3270bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
3271 SPIRVTypeInst ResType,
3274 assert(
I.getNumOperands() == 4);
3275 assert(
I.getOperand(2).isReg());
3276 assert(
I.getOperand(3).isReg());
3279 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3283 .
addUse(
I.getOperand(2).getReg())
3284 .
addUse(
I.getOperand(3).getReg())
3291bool SPIRVInstructionSelector::selectIntegerDotExpansion(
3292 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3293 assert(
I.getNumOperands() == 4);
3294 assert(
I.getOperand(2).isReg());
3295 assert(
I.getOperand(3).isReg());
3299 Register Vec0 =
I.getOperand(2).getReg();
3300 Register Vec1 =
I.getOperand(3).getReg();
3304 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
3313 "dot product requires a vector of at least 2 components");
3316 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3326 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3337 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3349bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
3350 SPIRVTypeInst ResType,
3351 MachineInstr &
I)
const {
3353 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
3356 .
addUse(
I.getOperand(2).getReg())
3361bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
3362 SPIRVTypeInst ResType,
3363 MachineInstr &
I)
const {
3365 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
3368 .
addUse(
I.getOperand(2).getReg())
3373bool SPIRVInstructionSelector::selectOpIsFinite(
Register ResVReg,
3374 SPIRVTypeInst ResType,
3375 MachineInstr &
I)
const {
3377 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsFinite))
3380 .
addUse(
I.getOperand(2).getReg())
3385bool SPIRVInstructionSelector::selectOpIsNormal(
Register ResVReg,
3386 SPIRVTypeInst ResType,
3387 MachineInstr &
I)
const {
3389 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNormal))
3392 .
addUse(
I.getOperand(2).getReg())
3397template <
bool Signed>
3398bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
3399 SPIRVTypeInst ResType,
3400 MachineInstr &
I)
const {
3401 assert(
I.getNumOperands() == 5);
3402 assert(
I.getOperand(2).isReg());
3403 assert(
I.getOperand(3).isReg());
3404 assert(
I.getOperand(4).isReg());
3407 Register Acc =
I.getOperand(2).getReg();
3411 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3413 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
3418 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
3421 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3433template <
bool Signed>
3434bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
3435 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3436 assert(
I.getNumOperands() == 5);
3437 assert(
I.getOperand(2).isReg());
3438 assert(
I.getOperand(3).isReg());
3439 assert(
I.getOperand(4).isReg());
3442 Register Acc =
I.getOperand(2).getReg();
3448 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
3452 for (
unsigned i = 0; i < 4; i++) {
3475 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
3495 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3510bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
3511 SPIRVTypeInst ResType,
3512 MachineInstr &
I)
const {
3513 assert(
I.getNumOperands() == 3);
3514 assert(
I.getOperand(2).isReg());
3516 Register VZero = buildZerosValF(ResType,
I);
3517 Register VOne = buildOnesValF(ResType,
I);
3519 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
3522 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3524 .
addUse(
I.getOperand(2).getReg())
3531bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
3532 SPIRVTypeInst ResType,
3533 MachineInstr &
I)
const {
3534 assert(
I.getNumOperands() == 3);
3535 assert(
I.getOperand(2).isReg());
3537 Register InputRegister =
I.getOperand(2).getReg();
3539 auto &
DL =
I.getDebugLoc();
3542 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3549 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
3551 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
3559 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3564 if (NeedsConversion) {
3565 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
3576bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
3577 SPIRVTypeInst ResType,
3579 unsigned Opcode)
const {
3583 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
3589 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
3590 BMI.addUse(
I.getOperand(J).getReg());
3597bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
3600 bool WithGroupSync)
const {
3602 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
3604 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
3606 assert(((Scope != SPIRV::Scope::Workgroup) ||
3607 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
3608 "Workgroup Scope must set WorkGroupMemory semantic "
3609 "in Barrier instruction");
3611 assert(((Scope != SPIRV::Scope::Device) ||
3612 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
3613 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
3614 "Device Scope must set UniformMemory and ImageMemory semantic "
3615 "in Barrier instruction");
3621 if (WithGroupSync) {
3622 Register ExecReg = buildI32Constant(SPIRV::Scope::Workgroup,
I);
3626 Register ScopeReg = buildI32Constant(Scope,
I);
3627 Register MemSemReg = buildI32Constant(MemSem,
I);
3629 MI.addUse(ScopeReg).addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
3633bool SPIRVInstructionSelector::selectWaveActiveCountBits(
3634 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3639 if (!selectWaveOpInst(BallotReg, BallotType,
I,
3640 SPIRV::OpGroupNonUniformBallot))
3645 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3650 .
addImm(SPIRV::GroupOperation::Reduce)
3657bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
3658 SPIRVTypeInst ResType,
3659 MachineInstr &
I)
const {
3664 Register InputReg =
I.getOperand(2).getReg();
3669 bool IsVector = NumElems > 1;
3682 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
3683 SPIRV::OpGroupNonUniformAllEqual);
3688 ElementResults.
reserve(NumElems);
3690 for (
unsigned Idx = 0; Idx < NumElems; ++Idx) {
3703 ElemInput = Extracted;
3709 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3720 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3731bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3732 SPIRVTypeInst ResType,
3733 MachineInstr &
I)
const {
3735 assert(
I.getNumOperands() == 3);
3737 auto Op =
I.getOperand(2);
3747 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3749 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3750 return diagnoseUnsupported(
I,
"WavePrefixBitCount requires boolean input");
3771 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3775 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3782bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3783 SPIRVTypeInst ResType,
3785 bool IsUnsigned)
const {
3786 return selectWaveReduce(
3787 ResVReg, ResType,
I, IsUnsigned,
3788 [&](
Register InputRegister,
bool IsUnsigned) {
3789 const bool IsFloatTy =
3791 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3792 : SPIRV::OpGroupNonUniformSMax;
3793 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3797bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3798 SPIRVTypeInst ResType,
3800 bool IsUnsigned)
const {
3801 return selectWaveReduce(
3802 ResVReg, ResType,
I, IsUnsigned,
3803 [&](
Register InputRegister,
bool IsUnsigned) {
3804 const bool IsFloatTy =
3806 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3807 : SPIRV::OpGroupNonUniformSMin;
3808 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3812bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3813 SPIRVTypeInst ResType,
3814 MachineInstr &
I)
const {
3815 return selectWaveReduce(ResVReg, ResType,
I,
false,
3816 [&](
Register InputRegister,
bool IsUnsigned) {
3818 InputRegister, SPIRV::OpTypeFloat);
3819 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3820 : SPIRV::OpGroupNonUniformIAdd;
3824bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3825 SPIRVTypeInst ResType,
3826 MachineInstr &
I)
const {
3827 return selectWaveReduce(ResVReg, ResType,
I,
false,
3828 [&](
Register InputRegister,
bool IsUnsigned) {
3830 InputRegister, SPIRV::OpTypeFloat);
3831 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3832 : SPIRV::OpGroupNonUniformIMul;
3836template <
typename PickOpcodeFn>
3837bool SPIRVInstructionSelector::selectWaveReduce(
3838 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3839 PickOpcodeFn &&PickOpcode)
const {
3840 assert(
I.getNumOperands() == 3);
3841 assert(
I.getOperand(2).isReg());
3843 Register InputRegister =
I.getOperand(2).getReg();
3847 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3850 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3856 .
addImm(SPIRV::GroupOperation::Reduce)
3857 .
addUse(
I.getOperand(2).getReg())
3862bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3863 SPIRVTypeInst ResType,
3865 unsigned Opcode)
const {
3866 return selectWaveReduce(
3867 ResVReg, ResType,
I,
false,
3868 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3871bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3872 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3873 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3874 [&](
Register InputRegister,
bool IsUnsigned) {
3876 InputRegister, SPIRV::OpTypeFloat);
3878 ? SPIRV::OpGroupNonUniformFAdd
3879 : SPIRV::OpGroupNonUniformIAdd;
3883bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3884 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3885 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3886 [&](
Register InputRegister,
bool IsUnsigned) {
3888 InputRegister, SPIRV::OpTypeFloat);
3890 ? SPIRV::OpGroupNonUniformFMul
3891 : SPIRV::OpGroupNonUniformIMul;
3895template <
typename PickOpcodeFn>
3896bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3897 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3898 PickOpcodeFn &&PickOpcode)
const {
3899 assert(
I.getNumOperands() == 3);
3900 assert(
I.getOperand(2).isReg());
3902 Register InputRegister =
I.getOperand(2).getReg();
3906 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3909 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3915 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3916 .
addUse(
I.getOperand(2).getReg())
3921bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
3922 SPIRVTypeInst ResType,
3925 assert(
I.getNumOperands() == 3);
3926 assert(
I.getOperand(2).isReg());
3928 Register InputRegister =
I.getOperand(2).getReg();
3934 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
3945bool SPIRVInstructionSelector::selectBitreverseViaI32(
Register ResVReg,
3946 SPIRVTypeInst ResType,
3953 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
3958 : SPIRV::OpUConvert;
3962 ShiftOp = SPIRV::OpShiftRightLogicalV;
3967 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3968 TII.get(SPIRV::OpConstantComposite))
3971 for (
unsigned It = 0; It <
N; ++It)
3975 ShiftConst = CompositeReg;
3980 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
3985 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
3990 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
3995 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
3998bool SPIRVInstructionSelector::handle64BitOverflow(
4000 unsigned int Opcode,
4007 "handle64BitOverflow should only be used for integer types");
4009 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
4011 MachineIRBuilder MIRBuilder(
I);
4013 SPIRVTypeInst I64x2Type =
4015 SPIRVTypeInst Vec2ResType =
4018 std::vector<Register> PartialRegs;
4020 unsigned CurrentComponent = 0;
4021 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
4025 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4026 TII.get(SPIRV::OpVectorShuffle))
4031 .
addImm(CurrentComponent)
4032 .
addImm(CurrentComponent + 1);
4042 PartialRegs.push_back(SubVecReg);
4045 if (CurrentComponent != ComponentCount) {
4051 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
4052 SPIRV::OpVectorExtractDynamic))
4061 PartialRegs.push_back(FinalElemResReg);
4065 return selectOpWithSrcs(ResVReg, ResType,
I, PartialRegs,
4066 SPIRV::OpCompositeConstruct);
4069bool SPIRVInstructionSelector::selectBitreverse64(
Register ResVReg,
4070 SPIRVTypeInst ResType,
4074 if (ComponentCount > 2)
4075 return handle64BitOverflow(
4076 ResVReg, ResType,
I, SrcReg, SPIRV::OpBitReverse,
4078 unsigned O) {
return this->selectBitreverse64(R,
T,
I, S); });
4080 MachineIRBuilder MIRBuilder(
I);
4084 I32Type, 2 * ComponentCount, MIRBuilder,
false);
4088 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
4093 if (!selectBitreverseNative(Reverse32, VecI32Type,
I, Vec32))
4100 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4101 TII.get(SPIRV::OpVectorShuffle))
4106 for (
unsigned J = 0; J < ComponentCount; ++J) {
4113 return selectOpWithSrcs(ResVReg, ResType,
I, {SwappedVec}, SPIRV::OpBitcast);
4116bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
4117 SPIRVTypeInst ResType,
4121 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
4129bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
4130 SPIRVTypeInst ResType,
4131 MachineInstr &
I)
const {
4132 Register OpReg =
I.getOperand(1).getReg();
4141 return selectBitreverseViaI32(ResVReg, ResType,
I, OpReg);
4143 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4145 return selectBitreverse64(ResVReg, ResType,
I, OpReg);
4147 return SPIRVInstructionSelector::diagnoseUnsupported(
4148 I,
"G_BITREVERSE only support 16,32,64 bits.");
4152 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4163 unsigned AndOp = SPIRV::OpBitwiseAndS;
4164 unsigned OrOp = SPIRV::OpBitwiseOrS;
4165 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
4166 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
4168 AndOp = SPIRV::OpBitwiseAndV;
4169 OrOp = SPIRV::OpBitwiseOrV;
4170 ShlOp = SPIRV::OpShiftLeftLogicalV;
4171 ShrOp = SPIRV::OpShiftRightLogicalV;
4177 const unsigned Shift) ->
Register {
4185 Register MaskReg = CreateConst(Mask);
4186 Register ShiftReg = CreateConst(Shift);
4193 if (!selectOpWithSrcs(
T1, ResType,
I, {Input, ShiftReg}, ShrOp) ||
4194 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
4195 !selectOpWithSrcs(T3, ResType,
I, {Input, MaskReg}, AndOp) ||
4196 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
4197 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
4205 uint64_t
Mask = ~0ull;
4206 while ((Shift >>= 1) > 0) {
4213 return BuildCOPY(ResVReg, Result,
I);
4216bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
4217 SPIRVTypeInst ResType,
4218 MachineInstr &
I)
const {
4219 assert(
I.getOperand(0).isReg() &&
I.getOperand(1).isReg() &&
4220 "G_FREEZE must define and use a register");
4221 Register OpReg =
I.getOperand(1).getReg();
4225 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
4238 if (MachineInstr *Def = MRI->
getVRegDef(OpReg)) {
4239 if (
Def->getOpcode() == TargetOpcode::COPY)
4242 switch (
Def->getOpcode()) {
4243 case SPIRV::ASSIGN_TYPE:
4244 if (MachineInstr *AssignToDef =
4246 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
4247 Reg =
Def->getOperand(2).getReg();
4250 case SPIRV::OpUndef:
4251 Reg =
Def->getOperand(1).getReg();
4254 unsigned DestOpCode;
4256 DestOpCode = SPIRV::OpConstantNull;
4257 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze of a "
4258 "static undef/poison lowered to OpConstantNull\n");
4260 DestOpCode = TargetOpcode::COPY;
4262 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze "
4263 "skipped, lowered as a copy of the operand\n");
4265 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
4266 .
addDef(
I.getOperand(0).getReg())
4274bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
4275 SPIRVTypeInst ResType,
4276 MachineInstr &
I)
const {
4278 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4280 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4284 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
4289 for (
unsigned i =
I.getNumExplicitDefs();
4290 i <
I.getNumExplicitOperands() && IsConst; ++i)
4294 if (!IsConst &&
N < 2)
4295 return diagnoseUnsupported(
4296 I,
"There must be at least two constituent operands in a vector");
4301 for (
unsigned i =
I.getNumExplicitDefs();
4302 i <
I.getNumExplicitOperands() && IsNullVector; ++i) {
4303 MachineInstr *
Def =
getDef(
I.getOperand(i), MRI);
4308 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4315 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4316 TII.get(IsConst ? SPIRV::OpConstantComposite
4317 : SPIRV::OpCompositeConstruct))
4320 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
4321 MIB.
addUse(
I.getOperand(i).getReg());
4326bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
4327 SPIRVTypeInst ResType,
4328 MachineInstr &
I)
const {
4330 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4332 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4338 if (!
I.getOperand(
OpIdx).isReg())
4345 if (!IsConst &&
N < 2)
4346 return diagnoseUnsupported(
4347 I,
"There must be at least two constituent operands in a vector");
4350 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4351 TII.get(IsConst ? SPIRV::OpConstantComposite
4352 : SPIRV::OpCompositeConstruct))
4355 for (
unsigned i = 0; i <
N; ++i)
4361bool SPIRVInstructionSelector::selectConcatVectors(
Register ResVReg,
4362 SPIRVTypeInst ResType,
4363 MachineInstr &
I)
const {
4367 if (ResType->
getOpcode() != SPIRV::OpTypeVector)
4369 "Cannot select G_CONCAT_VECTORS with a non-vector result");
4371 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4372 TII.get(SPIRV::OpCompositeConstruct))
4382bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
4383 SPIRVTypeInst ResType,
4384 MachineInstr &
I)
const {
4389 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
4391 Opcode = SPIRV::OpDemoteToHelperInvocation;
4393 Opcode = SPIRV::OpKill;
4395 if (MachineInstr *NextI =
I.getNextNode()) {
4397 NextI->eraseFromParent();
4407bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
4408 SPIRVTypeInst ResType,
unsigned CmpOpc,
4409 MachineInstr &
I)
const {
4410 Register Cmp0 =
I.getOperand(2).getReg();
4411 Register Cmp1 =
I.getOperand(3).getReg();
4414 "CMP operands should have the same type");
4415 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
4425bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
4426 SPIRVTypeInst ResType,
4427 MachineInstr &
I)
const {
4428 auto Pred =
I.getOperand(1).getPredicate();
4431 Register CmpOperand =
I.getOperand(2).getReg();
4433 bool IsPtrCmp = CmpOperandType && CmpOperandType.
isPointer();
4438 Register Op1 =
I.getOperand(3).getReg();
4442 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
4447 I.getOperand(3).setReg(NewOp1);
4453 return selectCmp(ResVReg, ResType, CmpOpc,
I);
4457SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
4458 SPIRVTypeInst ResType)
const {
4460 SPIRVTypeInst SpvI32Ty =
4463 auto ConstInt = ConstantInt::get(LLVMTy, Val);
4470 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4473 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
4476 .
addImm(APInt(32, Val).getZExtValue());
4478 GR.
add(ConstInt,
MI);
4485Register SPIRVInstructionSelector::buildI32ConstantInEntryBlock(
4486 uint32_t Val, MachineInstr &
I, SPIRVTypeInst ResType)
const {
4488 SPIRVTypeInst SpvI32Ty =
4490 auto *ConstInt = ConstantInt::get(LLVMTy, Val);
4495 MachineBasicBlock &EntryBB = *InsertIt->getParent();
4496 MachineInstr *
MI =
nullptr;
4500 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantNull))
4504 uint64_t ImmVal = APInt(32, Val).getZExtValue();
4505 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantI))
4511 GR.
add(ConstInt,
MI);
4516bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
4517 SPIRVTypeInst ResType,
4518 MachineInstr &
I)
const {
4520 return selectCmp(ResVReg, ResType, CmpOp,
I);
4523bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
4524 SPIRVTypeInst ResType,
4525 MachineInstr &
I)
const {
4527 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
4534 if (ResType->
getOpcode() != SPIRV::OpTypeVector &&
4535 ResType->
getOpcode() != SPIRV::OpTypeFloat)
4538 MachineIRBuilder MIRBuilder(
I);
4545 APFloat ConstVal(3.3219280948873623);
4549 APFloat::rmNearestTiesToEven, &LosesInfo);
4553 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
4554 ? SPIRV::OpVectorTimesScalar
4557 if (!selectOpWithSrcs(ArgReg, ResType,
I,
4558 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
4560 if (!selectExtInst(ResVReg, ResType,
I,
4561 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
4571Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
4572 MachineInstr &
I)
const {
4575 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4580bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
4586 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
4594 if (
Def->getOpcode() == SPIRV::OpConstantNull)
4597 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
4598 Def->getOpcode() == SPIRV::OpConstantI)
4611 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
4612 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
4614 Intrinsic::spv_const_composite)) {
4615 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
4616 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
4617 if (!IsZero(
Def->getOperand(i).getReg()))
4626Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
4627 MachineInstr &
I)
const {
4631 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4636Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
4637 MachineInstr &
I)
const {
4641 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4647 SPIRVTypeInst ResType,
4648 MachineInstr &
I)
const {
4652 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4657bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
4658 SPIRVTypeInst ResType,
4659 MachineInstr &
I)
const {
4660 Register SelectFirstArg =
I.getOperand(2).getReg();
4661 Register SelectSecondArg =
I.getOperand(3).getReg();
4670 SPIRV::OpTypeVector;
4677 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
4678 }
else if (IsPtrTy) {
4679 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
4681 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
4684 assert(IsScalarBool &&
"OpSelect with a scalar result requires a scalar "
4685 "boolean condition");
4687 Opcode = SPIRV::OpSelectSFSCond;
4688 }
else if (IsPtrTy) {
4689 Opcode = SPIRV::OpSelectSPSCond;
4691 Opcode = SPIRV::OpSelectSISCond;
4694 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4697 .
addUse(
I.getOperand(1).getReg())
4706bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
4707 SPIRVTypeInst ResType,
4709 MachineInstr &InsertAt,
4710 bool IsSigned)
const {
4712 Register ZeroReg = buildZerosVal(ResType, InsertAt);
4713 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
4714 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
4716 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4728bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
4729 SPIRVTypeInst ResType,
4730 MachineInstr &
I,
bool IsSigned,
4731 unsigned Opcode)
const {
4732 Register SrcReg =
I.getOperand(1).getReg();
4738 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
4743 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I, IsSigned);
4745 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
4748bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
4749 SPIRVTypeInst ResType, MachineInstr &
I,
4750 bool IsSigned)
const {
4751 Register SrcReg =
I.getOperand(1).getReg();
4753 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
4757 if (ResType == SrcType)
4758 return BuildCOPY(ResVReg, SrcReg,
I);
4760 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4761 return selectUnOp(ResVReg, ResType,
I, Opcode);
4764bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
4765 SPIRVTypeInst ResType,
4767 bool IsSigned)
const {
4768 MachineIRBuilder MIRBuilder(
I);
4769 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
4781 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
4784 .
addUse(
I.getOperand(1).getReg())
4785 .
addUse(
I.getOperand(2).getReg())
4790 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
4793 .
addUse(
I.getOperand(1).getReg())
4794 .
addUse(
I.getOperand(2).getReg())
4802 unsigned SelectOpcode =
4803 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
4808 .
addUse(buildOnesVal(
true, ResType,
I))
4809 .
addUse(buildZerosVal(ResType,
I))
4816 .
addUse(buildOnesVal(
false, ResType,
I))
4821bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
4824 SPIRVTypeInst IntTy,
4825 SPIRVTypeInst BoolTy)
const {
4828 bool IsVectorTy = IntTy->
getOpcode() == SPIRV::OpTypeVector;
4829 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
4831 Register One = buildOnesVal(
false, IntTy,
I);
4839 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
4848bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
4849 SPIRVTypeInst ResType,
4850 MachineInstr &
I)
const {
4851 Register IntReg =
I.getOperand(1).getReg();
4854 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
4855 if (ArgType == ResType)
4856 return BuildCOPY(ResVReg, IntReg,
I);
4858 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4859 return selectUnOp(ResVReg, ResType,
I, Opcode);
4862bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
4863 SPIRVTypeInst ResType,
4864 MachineInstr &
I)
const {
4865 unsigned Opcode =
I.getOpcode();
4866 unsigned TpOpcode = ResType->
getOpcode();
4868 if (ResType.
isPointer() || TpOpcode == SPIRV::OpTypeEvent) {
4869 assert(Opcode == TargetOpcode::G_CONSTANT &&
4870 I.getOperand(1).getCImm()->isZero());
4871 MachineBasicBlock &DepMBB =
I.getMF()->front();
4874 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
4881 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
4884bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
4885 SPIRVTypeInst ResType,
4886 MachineInstr &
I)
const {
4887 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4894bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
4895 SPIRVTypeInst ResType,
4896 MachineInstr &
I)
const {
4898 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
4902 .
addUse(
I.getOperand(3).getReg())
4904 .
addUse(
I.getOperand(2).getReg());
4905 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
4911bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
4912 SPIRVTypeInst ResType,
4913 MachineInstr &
I)
const {
4914 Type *MaybeResTy =
nullptr;
4919 "Expected aggregate type for extractv instruction");
4921 SPIRV::AccessQualifier::ReadWrite,
false);
4925 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
4928 .
addUse(
I.getOperand(2).getReg());
4929 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
4935bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
4936 SPIRVTypeInst ResType,
4937 MachineInstr &
I)
const {
4938 if (
getImm(
I.getOperand(4), MRI))
4939 return selectInsertVal(ResVReg, ResType,
I);
4941 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
4944 .
addUse(
I.getOperand(2).getReg())
4945 .
addUse(
I.getOperand(3).getReg())
4946 .
addUse(
I.getOperand(4).getReg())
4951bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
4952 SPIRVTypeInst ResType,
4953 MachineInstr &
I)
const {
4954 if (
getImm(
I.getOperand(3), MRI))
4955 return selectExtractVal(ResVReg, ResType,
I);
4957 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
4960 .
addUse(
I.getOperand(2).getReg())
4961 .
addUse(
I.getOperand(3).getReg())
4966bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
4967 SPIRVTypeInst ResType,
4968 MachineInstr &
I)
const {
4969 const bool IsGEPInBounds =
I.getOperand(2).getImm();
4972 const bool UseUntypedPointers =
4973 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
4978 if (UseUntypedPointers) {
4980 Opcode = IsGEPInBounds ? SPIRV::OpUntypedInBoundsAccessChainKHR
4981 : SPIRV::OpUntypedAccessChainKHR;
4983 Opcode = IsGEPInBounds ? SPIRV::OpUntypedInBoundsPtrAccessChainKHR
4984 : SPIRV::OpUntypedPtrAccessChainKHR;
4993 IsGEPInBounds ? SPIRV::OpInBoundsAccessChain : SPIRV::OpAccessChain;
4995 Opcode = IsGEPInBounds ? SPIRV::OpInBoundsPtrAccessChain
4996 : SPIRV::OpPtrAccessChain;
5001 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5006 if (UseUntypedPointers) {
5021 while (Def &&
Def->getOpcode() == TargetOpcode::COPY &&
5022 Def->getOperand(1).isReg())
5024 if (Def &&
Def->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)
5025 if (
const auto *GVar =
5028 SPIRV::AccessQualifier::ReadWrite,
5032 return diagnoseUnsupported(
5033 I,
"could not deduce the base type of an untyped access chain");
5038 Res.addUse(BaseReg);
5040 const bool IsAccessChainOpcode =
5041 (Opcode == SPIRV::OpAccessChain ||
5042 Opcode == SPIRV::OpInBoundsAccessChain ||
5043 Opcode == SPIRV::OpUntypedAccessChainKHR ||
5044 Opcode == SPIRV::OpUntypedInBoundsAccessChainKHR);
5046 assert((!IsAccessChainOpcode || (
getImm(
I.getOperand(4), MRI) &&
5047 foldImm(
I.getOperand(4), MRI) == 0)) &&
5048 "Cannot translate GEP to OpAccessChain.");
5051 const unsigned StartingIndex = IsAccessChainOpcode ? 5 : 4;
5052 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
5053 Res.addUse(
I.getOperand(i).getReg());
5054 Res.constrainAllUses(
TII,
TRI, RBI);
5059bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
5061 unsigned Lim =
I.getNumExplicitOperands();
5062 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
5063 Register OpReg =
I.getOperand(i).getReg();
5064 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
5066 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
5067 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
5068 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
5075 MachineFunction *MF =
I.getMF();
5081 SPIRVTypeInst WrapType = OpType;
5082 if (OpType->
getOpcode() == SPIRV::OpTypePointer &&
5084 SPIRV::StorageClass::CodeSectionINTEL) {
5086 SPIRV::StorageClass::Function,
I);
5093 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5094 TII.get(SPIRV::OpSpecConstantOp))
5097 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
5099 GR.
add(OpDefine, MIB);
5105bool SPIRVInstructionSelector::selectDerivativeInst(
5106 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
5107 const unsigned DPdOpCode)
const {
5110 if (!errorIfInstrOutsideShader(
I))
5116 Register SrcReg =
I.getOperand(2).getReg();
5121 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5124 .
addUse(
I.getOperand(2).getReg());
5126 MachineIRBuilder MIRBuilder(
I);
5129 if (componentCount != 1)
5137 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5142 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5147 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5155bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
5156 SPIRVTypeInst ResType,
5157 MachineInstr &
I)
const {
5161 case Intrinsic::spv_load:
5162 return selectLoad(ResVReg, ResType,
I);
5163 case Intrinsic::spv_atomic_load:
5164 return selectAtomicLoad(ResVReg, ResType,
I);
5165 case Intrinsic::spv_store:
5166 return selectStore(
I);
5167 case Intrinsic::spv_atomic_store:
5168 return selectAtomicStore(
I);
5169 case Intrinsic::spv_extractv:
5170 return selectExtractVal(ResVReg, ResType,
I);
5171 case Intrinsic::spv_insertv:
5172 return selectInsertVal(ResVReg, ResType,
I);
5173 case Intrinsic::spv_extractelt:
5174 return selectExtractElt(ResVReg, ResType,
I);
5175 case Intrinsic::spv_insertelt:
5176 return selectInsertElt(ResVReg, ResType,
I);
5177 case Intrinsic::spv_gep:
5178 return selectGEP(ResVReg, ResType,
I);
5179 case Intrinsic::spv_bitcast: {
5180 Register OpReg =
I.getOperand(2).getReg();
5181 SPIRVTypeInst OpType =
5185 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
5187 case Intrinsic::spv_unref_global:
5188 case Intrinsic::spv_init_global: {
5189 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
5194 Register GVarVReg =
MI->getOperand(0).getReg();
5195 if (!selectGlobalValue(GVarVReg, *
MI, Init))
5200 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
5202 MI->eraseFromParent();
5206 case Intrinsic::spv_undef: {
5207 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5213 case Intrinsic::spv_poison:
5214 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpPoisonKHR))
5219 case Intrinsic::spv_freeze:
5220 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
5223 .
addUse(
I.getOperand(2).getReg())
5226 case Intrinsic::spv_named_boolean_spec_constant: {
5227 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
5228 : SPIRV::OpSpecConstantFalse;
5230 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
5231 .
addDef(
I.getOperand(0).getReg())
5234 unsigned SpecId =
I.getOperand(2).getImm();
5236 SPIRV::Decoration::SpecId, {SpecId});
5240 case Intrinsic::spv_const_composite: {
5242 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
5248 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
5250 std::function<bool(
Register)> HasSpecConstOperand =
5260 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
5261 J < Def->getNumExplicitOperands(); ++J) {
5262 if (
Def->getOperand(J).isReg() &&
5263 HasSpecConstOperand(
Def->getOperand(J).getReg()))
5269 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
5270 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
5271 : SPIRV::OpConstantComposite;
5272 unsigned ContinuedOpc = HasSpecConst
5273 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
5274 : SPIRV::OpConstantCompositeContinuedINTEL;
5275 MachineIRBuilder MIR(
I);
5277 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
5279 for (
auto *Instr : Instructions) {
5280 Instr->setDebugLoc(
I.getDebugLoc());
5285 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5292 case Intrinsic::spv_assign_name: {
5293 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
5294 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
5295 for (
unsigned i =
I.getNumExplicitDefs() + 2;
5296 i <
I.getNumExplicitOperands(); ++i) {
5297 MIB.
addImm(
I.getOperand(i).getImm());
5302 case Intrinsic::spv_switch: {
5303 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
5304 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5305 if (
I.getOperand(i).isReg())
5306 MIB.
addReg(
I.getOperand(i).getReg());
5307 else if (
I.getOperand(i).isCImm())
5308 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
5309 else if (
I.getOperand(i).isMBB())
5310 MIB.
addMBB(
I.getOperand(i).getMBB());
5317 case Intrinsic::spv_loop_merge: {
5318 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
5319 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5320 if (
I.getOperand(i).isMBB())
5321 MIB.
addMBB(
I.getOperand(i).getMBB());
5328 case Intrinsic::spv_loop_control_intel: {
5330 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
5331 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
5336 case Intrinsic::spv_selection_merge: {
5338 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
5339 assert(
I.getOperand(1).isMBB() &&
5340 "operand 1 to spv_selection_merge must be a basic block");
5341 MIB.
addMBB(
I.getOperand(1).getMBB());
5342 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
5346 case Intrinsic::spv_cmpxchg:
5347 return selectAtomicCmpXchg(ResVReg, ResType,
I);
5348 case Intrinsic::spv_unreachable:
5349 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
5352 case Intrinsic::spv_abort:
5353 return selectAbort(
I);
5354 case Intrinsic::spv_alloca:
5355 return selectFrameIndex(ResVReg, ResType,
I);
5356 case Intrinsic::spv_alloca_array:
5357 return selectAllocaArray(ResVReg, ResType,
I);
5358 case Intrinsic::spv_assume:
5360 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
5361 .
addUse(
I.getOperand(1).getReg())
5366 case Intrinsic::spv_expect:
5368 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
5371 .
addUse(
I.getOperand(2).getReg())
5372 .
addUse(
I.getOperand(3).getReg())
5377 case Intrinsic::arithmetic_fence:
5378 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
5379 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
5382 .
addUse(
I.getOperand(2).getReg())
5386 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
5388 case Intrinsic::spv_thread_id:
5394 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
5396 case Intrinsic::spv_thread_id_in_group:
5402 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
5404 case Intrinsic::spv_group_id:
5410 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
5412 case Intrinsic::spv_flattened_thread_id_in_group:
5419 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
5421 case Intrinsic::spv_workgroup_size:
5422 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
5424 case Intrinsic::spv_global_size:
5425 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
5427 case Intrinsic::spv_global_offset:
5428 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
5430 case Intrinsic::spv_num_workgroups:
5431 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
5433 case Intrinsic::spv_subgroup_size:
5434 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
5436 case Intrinsic::spv_num_subgroups:
5437 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
5439 case Intrinsic::spv_subgroup_id:
5440 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
5441 case Intrinsic::spv_subgroup_local_invocation_id:
5442 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
5443 ResVReg, ResType,
I);
5444 case Intrinsic::spv_subgroup_max_size:
5445 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
5447 case Intrinsic::spv_fdot:
5448 return selectFloatDot(ResVReg, ResType,
I);
5449 case Intrinsic::spv_udot:
5450 case Intrinsic::spv_sdot:
5451 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5453 return selectIntegerDot(ResVReg, ResType,
I,
5454 IID == Intrinsic::spv_sdot);
5455 return selectIntegerDotExpansion(ResVReg, ResType,
I);
5456 case Intrinsic::spv_dot4add_i8packed:
5457 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5459 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
5460 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
5461 case Intrinsic::spv_dot4add_u8packed:
5462 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5464 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
5465 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
5466 case Intrinsic::spv_all:
5467 return selectAll(ResVReg, ResType,
I);
5468 case Intrinsic::spv_any:
5469 return selectAny(ResVReg, ResType,
I);
5470 case Intrinsic::spv_distance:
5471 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
5472 case Intrinsic::spv_lerp:
5473 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
5474 case Intrinsic::spv_length:
5475 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
5476 case Intrinsic::spv_degrees:
5477 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
5478 case Intrinsic::spv_faceforward:
5479 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
5480 case Intrinsic::spv_frac:
5481 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
5482 case Intrinsic::spv_isinf:
5483 return selectOpIsInf(ResVReg, ResType,
I);
5484 case Intrinsic::spv_isnan:
5485 return selectOpIsNan(ResVReg, ResType,
I);
5486 case Intrinsic::spv_isfinite:
5487 return selectOpIsFinite(ResVReg, ResType,
I);
5488 case Intrinsic::spv_isnormal:
5489 return selectOpIsNormal(ResVReg, ResType,
I);
5490 case Intrinsic::spv_normalize:
5491 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
5492 case Intrinsic::spv_refract:
5493 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
5494 case Intrinsic::spv_reflect:
5495 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
5496 case Intrinsic::spv_rsqrt:
5497 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
5498 case Intrinsic::spv_sign:
5499 return selectSign(ResVReg, ResType,
I);
5500 case Intrinsic::spv_smoothstep:
5501 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
5502 case Intrinsic::spv_firstbituhigh:
5503 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
5504 case Intrinsic::spv_firstbitshigh:
5505 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
5506 case Intrinsic::spv_firstbitlow:
5507 return selectFirstBitLow(ResVReg, ResType,
I);
5508 case Intrinsic::spv_all_memory_barrier:
5509 return selectBarrierInst(
I, SPIRV::Scope::Device,
5510 SPIRV::MemorySemantics::UniformMemory |
5511 SPIRV::MemorySemantics::ImageMemory |
5512 SPIRV::MemorySemantics::WorkgroupMemory,
5514 case Intrinsic::spv_all_memory_barrier_with_group_sync:
5515 return selectBarrierInst(
I, SPIRV::Scope::Device,
5516 SPIRV::MemorySemantics::UniformMemory |
5517 SPIRV::MemorySemantics::ImageMemory |
5518 SPIRV::MemorySemantics::WorkgroupMemory,
5520 case Intrinsic::spv_device_memory_barrier:
5521 return selectBarrierInst(
I, SPIRV::Scope::Device,
5522 SPIRV::MemorySemantics::UniformMemory |
5523 SPIRV::MemorySemantics::ImageMemory,
5525 case Intrinsic::spv_device_memory_barrier_with_group_sync:
5526 return selectBarrierInst(
I, SPIRV::Scope::Device,
5527 SPIRV::MemorySemantics::UniformMemory |
5528 SPIRV::MemorySemantics::ImageMemory,
5530 case Intrinsic::spv_group_memory_barrier:
5531 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5532 SPIRV::MemorySemantics::WorkgroupMemory,
5534 case Intrinsic::spv_group_memory_barrier_with_group_sync:
5535 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5536 SPIRV::MemorySemantics::WorkgroupMemory,
5538 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
5539 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
5540 SPIRV::StorageClass::StorageClass ResSC =
5543 return diagnoseUnsupported(
I,
"The target storage class is not castable "
5544 "from the Generic storage class");
5545 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
5553 case Intrinsic::spv_lifetime_start:
5554 case Intrinsic::spv_lifetime_end: {
5555 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
5556 : SPIRV::OpLifetimeStop;
5557 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
5558 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
5567 case Intrinsic::spv_saturate:
5568 return selectSaturate(ResVReg, ResType,
I);
5569 case Intrinsic::spv_nclamp:
5570 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
5571 case Intrinsic::spv_uclamp:
5572 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
5573 case Intrinsic::spv_sclamp:
5574 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
5575 case Intrinsic::spv_subgroup_prefix_bit_count:
5576 return selectWavePrefixBitCount(ResVReg, ResType,
I);
5577 case Intrinsic::spv_wave_active_countbits:
5578 return selectWaveActiveCountBits(ResVReg, ResType,
I);
5579 case Intrinsic::spv_wave_all_equal:
5580 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
5581 case Intrinsic::spv_wave_all:
5582 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
5583 case Intrinsic::spv_wave_any:
5584 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
5585 case Intrinsic::spv_subgroup_ballot:
5586 return selectWaveOpInst(ResVReg, ResType,
I,
5587 SPIRV::OpGroupNonUniformBallot);
5588 case Intrinsic::spv_wave_is_first_lane:
5589 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
5590 case Intrinsic::spv_wave_reduce_or:
5591 return selectWaveReduceOp(ResVReg, ResType,
I,
5592 SPIRV::OpGroupNonUniformBitwiseOr);
5593 case Intrinsic::spv_wave_reduce_xor:
5594 return selectWaveReduceOp(ResVReg, ResType,
I,
5595 SPIRV::OpGroupNonUniformBitwiseXor);
5596 case Intrinsic::spv_wave_reduce_and:
5597 return selectWaveReduceOp(ResVReg, ResType,
I,
5598 SPIRV::OpGroupNonUniformBitwiseAnd);
5599 case Intrinsic::spv_wave_reduce_umax:
5600 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
5601 case Intrinsic::spv_wave_reduce_max:
5602 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
5603 case Intrinsic::spv_wave_reduce_umin:
5604 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
5605 case Intrinsic::spv_wave_reduce_min:
5606 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
5607 case Intrinsic::spv_wave_reduce_sum:
5608 return selectWaveReduceSum(ResVReg, ResType,
I);
5609 case Intrinsic::spv_wave_product:
5610 return selectWaveReduceProduct(ResVReg, ResType,
I);
5611 case Intrinsic::spv_wave_readlane:
5612 return selectWaveOpInst(ResVReg, ResType,
I,
5613 SPIRV::OpGroupNonUniformShuffle);
5614 case Intrinsic::spv_wave_prefix_sum:
5615 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
5616 case Intrinsic::spv_wave_prefix_product:
5617 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
5618 case Intrinsic::spv_quad_read_across_x: {
5619 return selectQuadSwap(ResVReg, ResType,
I, 0);
5621 case Intrinsic::spv_quad_read_across_y: {
5622 return selectQuadSwap(ResVReg, ResType,
I, 1);
5624 case Intrinsic::spv_quad_read_across_diagonal: {
5625 return selectQuadSwap(ResVReg, ResType,
I, 2);
5627 case Intrinsic::spv_radians:
5628 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
5632 case Intrinsic::instrprof_increment:
5633 case Intrinsic::instrprof_increment_step:
5634 case Intrinsic::instrprof_value_profile:
5637 case Intrinsic::spv_value_md:
5639 case Intrinsic::spv_resource_handlefrombinding: {
5640 return selectHandleFromBinding(ResVReg, ResType,
I);
5642 case Intrinsic::spv_resource_counterhandlefrombinding:
5643 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
5644 case Intrinsic::spv_resource_updatecounter:
5645 return selectUpdateCounter(ResVReg, ResType,
I);
5646 case Intrinsic::spv_resource_store_typedbuffer: {
5647 return selectImageWriteIntrinsic(
I);
5649 case Intrinsic::spv_resource_load_typedbuffer: {
5650 return selectReadImageIntrinsic(ResVReg, ResType,
I);
5652 case Intrinsic::spv_resource_load_level: {
5653 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
5655 case Intrinsic::spv_resource_getdimensions_x:
5656 case Intrinsic::spv_resource_getdimensions_xy:
5657 case Intrinsic::spv_resource_getdimensions_xyz: {
5658 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
5660 case Intrinsic::spv_resource_getdimensions_levels_x:
5661 case Intrinsic::spv_resource_getdimensions_levels_xy:
5662 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
5663 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
5665 case Intrinsic::spv_resource_getdimensions_ms_xy:
5666 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
5667 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
5669 case Intrinsic::spv_resource_calculate_lod:
5670 case Intrinsic::spv_resource_calculate_lod_unclamped:
5671 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
5672 case Intrinsic::spv_resource_sample:
5673 case Intrinsic::spv_resource_sample_clamp:
5674 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
5675 case Intrinsic::spv_resource_samplebias:
5676 case Intrinsic::spv_resource_samplebias_clamp:
5677 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
5678 case Intrinsic::spv_resource_samplegrad:
5679 case Intrinsic::spv_resource_samplegrad_clamp:
5680 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
5681 case Intrinsic::spv_resource_samplelevel:
5682 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
5683 case Intrinsic::spv_resource_samplecmp:
5684 case Intrinsic::spv_resource_samplecmp_clamp:
5685 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
5686 case Intrinsic::spv_resource_samplecmplevelzero:
5687 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
5688 case Intrinsic::spv_resource_gather:
5689 case Intrinsic::spv_resource_gather_cmp:
5690 return selectGatherIntrinsic(ResVReg, ResType,
I);
5691 case Intrinsic::spv_resource_getbasepointer:
5692 case Intrinsic::spv_resource_getpointer: {
5693 return selectResourceGetPointer(ResVReg, ResType,
I);
5695 case Intrinsic::spv_pushconstant_getpointer: {
5696 return selectPushConstantGetPointer(ResVReg, ResType,
I);
5698 case Intrinsic::spv_discard: {
5699 return selectDiscard(ResVReg, ResType,
I);
5701 case Intrinsic::spv_resource_nonuniformindex: {
5702 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
5704 case Intrinsic::spv_unpackhalf2x16: {
5705 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
5707 case Intrinsic::spv_packhalf2x16: {
5708 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
5710 case Intrinsic::spv_ddx:
5711 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
5712 case Intrinsic::spv_ddy:
5713 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
5714 case Intrinsic::spv_ddx_coarse:
5715 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
5716 case Intrinsic::spv_ddy_coarse:
5717 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
5718 case Intrinsic::spv_ddx_fine:
5719 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
5720 case Intrinsic::spv_ddy_fine:
5721 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
5722 case Intrinsic::spv_fwidth:
5723 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
5724 case Intrinsic::spv_masked_gather:
5725 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5726 return selectMaskedGather(ResVReg, ResType,
I);
5727 return diagnoseUnsupported(
5728 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
5729 case Intrinsic::spv_masked_scatter:
5730 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5731 return selectMaskedScatter(
I);
5732 return diagnoseUnsupported(
5733 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
5734 case Intrinsic::returnaddress:
5735 case Intrinsic::frameaddress: {
5737 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5744 return diagnoseUnsupported(
I,
"intrinsic selection not implemented.");
5749bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
5750 SPIRVTypeInst ResType,
5751 MachineInstr &
I)
const {
5754 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
5761bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
5762 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5764 assert(Intr.getIntrinsicID() ==
5765 Intrinsic::spv_resource_counterhandlefrombinding);
5768 Register MainHandleReg = Intr.getOperand(2).getReg();
5770 assert(MainHandleDef->getIntrinsicID() ==
5771 Intrinsic::spv_resource_handlefrombinding);
5775 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
5776 Register IndexReg = MainHandleDef->getOperand(5).getReg();
5777 std::string CounterName =
5782 MachineIRBuilder MIRBuilder(
I);
5784 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
5786 ArraySize, IndexReg, CounterName, MIRBuilder);
5788 return BuildCOPY(ResVReg, CounterVarReg,
I);
5791bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
5792 SPIRVTypeInst ResType,
5793 MachineInstr &
I)
const {
5795 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
5797 Register CounterHandleReg = Intr.getOperand(2).getReg();
5798 Register IncrReg = Intr.getOperand(3).getReg();
5805 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
5806 assert(CounterVarPointeeType &&
5807 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
5808 "Counter variable must be a struct");
5810 SPIRV::StorageClass::StorageBuffer &&
5811 "Counter variable must be in the storage buffer storage class");
5813 "Counter variable must have exactly 1 member in the struct");
5814 const SPIRVTypeInst MemberType =
5817 "Counter variable struct must have a single i32 member");
5821 MachineIRBuilder MIRBuilder(
I);
5823 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
5826 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
5832 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5835 .
addUse(CounterHandleReg)
5842 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
5845 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
5848 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
5857 return BuildCOPY(ResVReg, AtomicRes,
I);
5865 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
5873bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
5874 SPIRVTypeInst ResType,
5875 MachineInstr &
I)
const {
5883 Register ImageReg =
I.getOperand(2).getReg();
5891 Register IdxReg =
I.getOperand(3).getReg();
5893 MachineInstr &Pos =
I;
5895 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
5899bool SPIRVInstructionSelector::generateSampleImage(
5902 DebugLoc Loc, MachineInstr &Pos)
const {
5913 if (!loadHandleBeforePosition(NewSamplerReg,
5919 MachineIRBuilder MIRBuilder(Pos);
5932 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
5933 ImOps.Lod.has_value();
5934 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
5935 : SPIRV::OpImageSampleImplicitLod;
5937 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
5938 : SPIRV::OpImageSampleDrefImplicitLod;
5947 MIB.
addUse(*ImOps.Compare);
5949 uint32_t ImageOperands = 0;
5951 ImageOperands |= SPIRV::ImageOperand::Bias;
5953 ImageOperands |= SPIRV::ImageOperand::Lod;
5954 if (ImOps.GradX && ImOps.GradY)
5955 ImageOperands |= SPIRV::ImageOperand::Grad;
5956 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
5958 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
5961 "Non-constant offsets are not supported in sample instructions.");
5966 ImageOperands |= SPIRV::ImageOperand::MinLod;
5968 if (ImageOperands != 0) {
5969 MIB.
addImm(ImageOperands);
5970 if (ImageOperands & SPIRV::ImageOperand::Bias)
5972 if (ImageOperands & SPIRV::ImageOperand::Lod)
5974 if (ImageOperands & SPIRV::ImageOperand::Grad) {
5975 MIB.
addUse(*ImOps.GradX);
5976 MIB.
addUse(*ImOps.GradY);
5979 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
5980 MIB.
addUse(*ImOps.Offset);
5981 if (ImageOperands & SPIRV::ImageOperand::MinLod)
5982 MIB.
addUse(*ImOps.MinLod);
5989bool SPIRVInstructionSelector::selectImageQuerySize(
5991 std::optional<Register> LodReg)
const {
5993 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
5996 "ImageReg is not an image type.");
5998 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6000 unsigned NumComponents = 0;
6002 case SPIRV::Dim::DIM_1D:
6003 case SPIRV::Dim::DIM_Buffer:
6004 NumComponents =
IsArray ? 2 : 1;
6006 case SPIRV::Dim::DIM_2D:
6007 case SPIRV::Dim::DIM_Cube:
6008 case SPIRV::Dim::DIM_Rect:
6009 NumComponents =
IsArray ? 3 : 2;
6011 case SPIRV::Dim::DIM_3D:
6015 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
6020 SPIRVTypeInst ResType =
6025 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6035bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
6036 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6037 Register ImageReg =
I.getOperand(2).getReg();
6044 return selectImageQuerySize(NewImageReg, ResVReg,
I);
6047bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
6048 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6049 Register ImageReg =
I.getOperand(2).getReg();
6058 Register LodReg =
I.getOperand(3).getReg();
6061 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
6063 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
6070 TII.get(SPIRV::OpImageQueryLevels))
6077 TII.get(SPIRV::OpCompositeConstruct))
6087bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
6088 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6089 Register ImageReg =
I.getOperand(2).getReg();
6100 "OpImageQuerySamples requires a multisampled image");
6102 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
6110 TII.get(SPIRV::OpImageQuerySamples))
6117 TII.get(SPIRV::OpCompositeConstruct))
6127bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
6128 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6129 Register ImageReg =
I.getOperand(2).getReg();
6130 Register SamplerReg =
I.getOperand(3).getReg();
6131 Register CoordinateReg =
I.getOperand(4).getReg();
6147 if (!loadHandleBeforePosition(
6152 MachineIRBuilder MIRBuilder(
I);
6158 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6168 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
6175 unsigned ExtractedIndex =
6177 Intrinsic::spv_resource_calculate_lod_unclamped
6181 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6182 TII.get(SPIRV::OpCompositeExtract))
6192bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
6193 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6194 Register ImageReg =
I.getOperand(2).getReg();
6195 Register SamplerReg =
I.getOperand(3).getReg();
6196 Register CoordinateReg =
I.getOperand(4).getReg();
6197 ImageOperands ImOps;
6198 if (
I.getNumOperands() > 5)
6199 ImOps.Offset =
I.getOperand(5).getReg();
6200 if (
I.getNumOperands() > 6)
6201 ImOps.MinLod =
I.getOperand(6).getReg();
6202 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6203 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6206bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
6207 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6208 Register ImageReg =
I.getOperand(2).getReg();
6209 Register SamplerReg =
I.getOperand(3).getReg();
6210 Register CoordinateReg =
I.getOperand(4).getReg();
6211 ImageOperands ImOps;
6212 ImOps.Bias =
I.getOperand(5).getReg();
6213 if (
I.getNumOperands() > 6)
6214 ImOps.Offset =
I.getOperand(6).getReg();
6215 if (
I.getNumOperands() > 7)
6216 ImOps.MinLod =
I.getOperand(7).getReg();
6217 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6218 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6221bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
6222 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6223 Register ImageReg =
I.getOperand(2).getReg();
6224 Register SamplerReg =
I.getOperand(3).getReg();
6225 Register CoordinateReg =
I.getOperand(4).getReg();
6226 ImageOperands ImOps;
6227 ImOps.GradX =
I.getOperand(5).getReg();
6228 ImOps.GradY =
I.getOperand(6).getReg();
6229 if (
I.getNumOperands() > 7)
6230 ImOps.Offset =
I.getOperand(7).getReg();
6231 if (
I.getNumOperands() > 8)
6232 ImOps.MinLod =
I.getOperand(8).getReg();
6233 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6234 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6237bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
6238 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6239 Register ImageReg =
I.getOperand(2).getReg();
6240 Register SamplerReg =
I.getOperand(3).getReg();
6241 Register CoordinateReg =
I.getOperand(4).getReg();
6242 ImageOperands ImOps;
6243 ImOps.Lod =
I.getOperand(5).getReg();
6244 if (
I.getNumOperands() > 6)
6245 ImOps.Offset =
I.getOperand(6).getReg();
6246 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6247 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6250bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
6251 SPIRVTypeInst ResType,
6252 MachineInstr &
I)
const {
6253 Register ImageReg =
I.getOperand(2).getReg();
6254 Register SamplerReg =
I.getOperand(3).getReg();
6255 Register CoordinateReg =
I.getOperand(4).getReg();
6256 ImageOperands ImOps;
6257 ImOps.Compare =
I.getOperand(5).getReg();
6258 if (
I.getNumOperands() > 6)
6259 ImOps.Offset =
I.getOperand(6).getReg();
6260 if (
I.getNumOperands() > 7)
6261 ImOps.MinLod =
I.getOperand(7).getReg();
6262 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6263 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6266bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
6267 SPIRVTypeInst ResType,
6268 MachineInstr &
I)
const {
6269 Register ImageReg =
I.getOperand(2).getReg();
6270 Register CoordinateReg =
I.getOperand(3).getReg();
6271 Register LodReg =
I.getOperand(4).getReg();
6273 ImageOperands ImOps;
6275 if (
I.getNumOperands() > 5)
6276 ImOps.Offset =
I.getOperand(5).getReg();
6288 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
6289 I.getDebugLoc(),
I, &ImOps);
6292bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
6293 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6294 Register ImageReg =
I.getOperand(2).getReg();
6295 Register SamplerReg =
I.getOperand(3).getReg();
6296 Register CoordinateReg =
I.getOperand(4).getReg();
6297 ImageOperands ImOps;
6298 ImOps.Compare =
I.getOperand(5).getReg();
6299 if (
I.getNumOperands() > 6)
6300 ImOps.Offset =
I.getOperand(6).getReg();
6303 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6304 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6307bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
6308 SPIRVTypeInst ResType,
6309 MachineInstr &
I)
const {
6310 Register ImageReg =
I.getOperand(2).getReg();
6311 Register SamplerReg =
I.getOperand(3).getReg();
6312 Register CoordinateReg =
I.getOperand(4).getReg();
6315 "ImageReg is not an image type.");
6320 ComponentOrCompareReg =
I.getOperand(5).getReg();
6321 OffsetReg =
I.getOperand(6).getReg();
6324 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
6328 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6329 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
6330 Dim != SPIRV::Dim::DIM_Rect) {
6332 "Gather operations are only supported for 2D, Cube, and Rect images.");
6339 if (!loadHandleBeforePosition(
6344 MachineIRBuilder MIRBuilder(
I);
6345 SPIRVTypeInst SampledImageType =
6350 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6358 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
6360 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
6362 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6367 .
addUse(ComponentOrCompareReg);
6369 uint32_t ImageOperands = 0;
6370 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
6371 if (Dim == SPIRV::Dim::DIM_Cube) {
6373 "Gather operations with offset are not supported for Cube images.");
6377 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6379 ImageOperands |= SPIRV::ImageOperand::Offset;
6383 if (ImageOperands != 0) {
6384 MIB.
addImm(ImageOperands);
6386 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6394bool SPIRVInstructionSelector::generateImageReadOrFetch(
6397 const ImageOperands *ImOps)
const {
6400 "ImageReg is not an image type.");
6402 bool IsSignedInteger =
6407 bool IsFetch = (SampledOp.getImm() == 1);
6409 auto AddOperands = [&](MachineInstrBuilder &MIB) {
6410 uint32_t ImageOperandsMask = 0;
6411 if (IsSignedInteger)
6412 ImageOperandsMask |= 0x1000;
6414 if (IsFetch && ImOps) {
6416 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
6417 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
6419 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
6421 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
6425 if (ImageOperandsMask != 0) {
6426 MIB.
addImm(ImageOperandsMask);
6427 if (IsFetch && ImOps) {
6430 if (ImOps->Offset &&
6431 (ImageOperandsMask &
6432 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
6433 MIB.
addUse(*ImOps->Offset);
6442 SPIRVTypeInst SampledType =
6445 SPIRVTypeInst ReadType =
6446 widenTypeToVec4(IsPacked ? SampledType : ResType, Pos);
6447 bool ReadTypeMatchesResult = ReadType == ResType;
6449 Register ReadReg = ReadTypeMatchesResult
6455 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
6461 BMI.constrainAllUses(
TII,
TRI, RBI);
6463 if (ReadTypeMatchesResult)
6476 if (ResultSize == 1) {
6485 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
6488bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
6489 SPIRVTypeInst ResType,
6490 MachineInstr &
I)
const {
6491 Register ResourcePtr =
I.getOperand(2).getReg();
6493 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
6502 MachineIRBuilder MIRBuilder(
I);
6507 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6513 if (
I.getNumExplicitOperands() > 3) {
6514 Register IndexReg =
I.getOperand(3).getReg();
6521bool SPIRVInstructionSelector::selectPushConstantGetPointer(
6522 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6527bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
6528 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6529 Register ObjReg =
I.getOperand(2).getReg();
6530 if (!BuildCOPY(ResVReg, ObjReg,
I))
6540 decorateUsesAsNonUniform(ResVReg);
6544void SPIRVInstructionSelector::decorateUsesAsNonUniform(
6547 {NonUniformReg,
nullptr}};
6548 llvm::SmallSet<Register, 8> Visited;
6549 while (WorkList.
size() > 0) {
6552 if (!Visited.
insert(CurrentReg).second)
6555 bool IsDecorated =
false;
6557 if (
Use.getOpcode() == SPIRV::OpDecorate &&
6558 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
6564 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
6566 if (ResultReg == CurrentReg)
6574 MachineInstr &InsertPt =
6577 SPIRV::Decoration::NonUniformEXT, {});
6582bool SPIRVInstructionSelector::extractSubvector(
6584 MachineInstr &InsertionPoint)
const {
6586 [[maybe_unused]] uint64_t InputSize =
6589 assert(InputSize > 1 &&
"The input must be a vector.");
6590 assert(ResultSize > 1 &&
"The result must be a vector.");
6591 assert(ResultSize < InputSize &&
6592 "Cannot extract more element than there are in the input.");
6596 for (uint64_t
I = 0;
I < ResultSize;
I++) {
6599 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6608 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
6610 TII.get(SPIRV::OpCompositeConstruct))
6614 for (
Register ComponentReg : ComponentRegisters)
6615 MIB.
addUse(ComponentReg);
6620bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
6621 MachineInstr &
I)
const {
6628 Register ImageReg =
I.getOperand(1).getReg();
6636 Register CoordinateReg =
I.getOperand(2).getReg();
6637 Register DataReg =
I.getOperand(3).getReg();
6640 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
6648Register SPIRVInstructionSelector::buildPointerToResource(
6649 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
6650 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
6651 StringRef Name, MachineIRBuilder MIRBuilder)
const {
6653 if (ArraySize == 1) {
6654 SPIRVTypeInst PtrType =
6657 "SpirvResType did not have an explicit layout.");
6662 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
6663 SPIRVTypeInst VarPointerType =
6666 VarPointerType, Set,
Binding, Name, MIRBuilder);
6668 SPIRVTypeInst ResPointerType =
6681bool SPIRVInstructionSelector::selectFirstBitSet16(
6682 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
6683 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
6685 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
6689 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
6692bool SPIRVInstructionSelector::selectFirstBitSet32(
6694 unsigned BitSetOpcode)
const {
6695 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6698 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6705bool SPIRVInstructionSelector::selectFirstBitSet64(
6707 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
6720 if (ComponentCount > 2) {
6721 auto Func = [
this, SwapPrimarySide](
Register ResVReg, SPIRVTypeInst ResType,
6723 unsigned Opcode) ->
bool {
6724 return this->selectFirstBitSet64(ResVReg, ResType,
I, SrcReg, Opcode,
6728 return handle64BitOverflow(ResVReg, ResType,
I, SrcReg, BitSetOpcode, Func);
6732 MachineIRBuilder MIRBuilder(
I);
6734 BaseType, 2 * ComponentCount, MIRBuilder,
false);
6738 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
6744 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
6751 bool IsScalarRes = ResType->
getOpcode() != SPIRV::OpTypeVector;
6754 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
6755 SPIRV::OpVectorExtractDynamic))
6757 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
6758 SPIRV::OpVectorExtractDynamic))
6762 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6763 TII.get(SPIRV::OpVectorShuffle))
6771 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
6777 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6778 TII.get(SPIRV::OpVectorShuffle))
6786 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
6806 SelectOp = SPIRV::OpSelectSISCond;
6807 AddOp = SPIRV::OpIAddS;
6815 SelectOp = SPIRV::OpSelectVIVCond;
6816 AddOp = SPIRV::OpIAddV;
6822 Register RegSecondaryOffset = Reg0;
6826 if (SwapPrimarySide) {
6827 PrimaryReg = LowReg;
6828 SecondaryReg = HighReg;
6829 RegPrimaryOffset = Reg0;
6830 RegSecondaryOffset = Reg32;
6835 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
6836 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
6841 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
6842 SPIRV::OpINotEqual))
6849 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
6850 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
6855 if (SwapPrimarySide) {
6857 if (!selectOpWithSrcs(RegAdd, ResType,
I,
6858 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
6869 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
6870 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
6875 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
6876 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
6879 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
6883bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
6884 SPIRVTypeInst ResType,
6886 bool IsSigned)
const {
6888 Register OpReg =
I.getOperand(2).getReg();
6891 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
6892 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
6896 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6898 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6900 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6903 return diagnoseUnsupported(
6905 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
6909bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
6910 SPIRVTypeInst ResType,
6911 MachineInstr &
I)
const {
6913 Register OpReg =
I.getOperand(2).getReg();
6918 unsigned ExtendOpcode = SPIRV::OpUConvert;
6919 unsigned BitSetOpcode = GL::FindILsb;
6923 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6925 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6927 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6930 return diagnoseUnsupported(
I,
6931 "spv_firstbitlow only supports 16,32,64 bits.");
6935bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
6936 SPIRVTypeInst ResType,
6937 MachineInstr &
I)
const {
6941 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
6944 .
addUse(
I.getOperand(2).getReg())
6947 unsigned Alignment =
I.getOperand(3).getImm();
6961 while (!Worklist.
empty()) {
6963 switch (
T->getOpcode()) {
6964 case SPIRV::OpTypeInt:
6965 case SPIRV::OpTypeFloat:
6966 case SPIRV::OpTypePointer:
6968 case SPIRV::OpTypeVector:
6969 case SPIRV::OpTypeMatrix:
6970 case SPIRV::OpTypeArray: {
6971 Register OperandReg =
T->getOperand(1).getReg();
6975 case SPIRV::OpTypeStruct:
6976 for (
unsigned Idx = 1,
E =
T->getNumOperands(); Idx <
E; ++Idx) {
6977 Register OperandReg =
T->getOperand(Idx).getReg();
6989bool SPIRVInstructionSelector::selectAbort(MachineInstr &
I)
const {
6990 assert(
I.getNumExplicitOperands() == 2);
6992 Register MsgReg =
I.getOperand(1).getReg();
6994 assert(MsgType &&
"Message argument of llvm.spv.abort has no SPIR-V type");
6997 return diagnoseUnsupported(
6999 "llvm.spv.abort message type must be a concrete SPIR-V type (numerical "
7000 "scalar, pointer, vector, matrix, or aggregate of such types)");
7003 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7010bool SPIRVInstructionSelector::selectTrap(MachineInstr &
I)
const {
7019 uint32_t MsgVal = ~0
u;
7020 if (
I.getOpcode() == TargetOpcode::G_UBSANTRAP)
7021 MsgVal =
static_cast<uint32_t
>(
I.getOperand(0).
getImm());
7024 Register MsgReg = buildI32ConstantInEntryBlock(MsgVal,
I, MsgType);
7027 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7034bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
7035 SPIRVTypeInst ResType,
7036 MachineInstr &
I)
const {
7043 bool UseUntypedPointers =
7044 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7046 UseUntypedPointers ? SPIRV::OpUntypedVariableKHR : SPIRV::OpVariable;
7048 auto MIB =
BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(Opcode))
7051 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7055 if (UseUntypedPointers) {
7059 return diagnoseUnsupported(
7060 I,
"could not deduce the data type of an untyped variable");
7066 unsigned Alignment =
I.getOperand(2).getImm();
7073bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
7078 const MachineInstr *PrevI =
I.getPrevNode();
7080 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
7084 .
addMBB(
I.getOperand(0).getMBB())
7089 .
addMBB(
I.getOperand(0).getMBB())
7094bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
7105 const MachineInstr *NextI =
I.getNextNode();
7107 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
7113 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
7115 .
addUse(
I.getOperand(0).getReg())
7116 .
addMBB(
I.getOperand(1).getMBB())
7122bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
7123 MachineInstr &
I)
const {
7125 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
7127 const unsigned NumOps =
I.getNumOperands();
7128 for (
unsigned i = 1; i <
NumOps; i += 2) {
7129 MIB.
addUse(
I.getOperand(i + 0).getReg());
7130 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
7136bool SPIRVInstructionSelector::selectGlobalValue(
7137 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
7139 MachineIRBuilder MIRBuilder(
I);
7140 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
7143 std::string GlobalIdent;
7145 unsigned &
ID = UnnamedGlobalIDs[GV];
7147 ID = UnnamedGlobalIDs.
size();
7148 GlobalIdent =
"__unnamed_" + Twine(ID).str();
7174 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
7181 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7186 MachineInstrBuilder MIB1 =
7187 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7190 MachineInstrBuilder MIB2 =
7192 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
7196 GR.
add(ConstVal, MIB2);
7204 MachineInstrBuilder MIB3 =
7205 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7208 GR.
add(ConstVal, MIB3);
7214 assert(NewReg != ResVReg);
7215 return BuildCOPY(ResVReg, NewReg,
I);
7225 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
7228 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
7234 SPIRVTypeInst ResType =
7238 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
7243 if (
GlobalVar->isExternallyInitialized() &&
7244 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
7245 constexpr unsigned ReadWriteINTEL = 3u;
7248 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
7254bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
7255 SPIRVTypeInst ResType,
7256 MachineInstr &
I)
const {
7258 return selectExtInst(ResVReg, ResType,
I, CL::log10);
7266 MachineIRBuilder MIRBuilder(
I);
7271 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7274 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
7276 .
add(
I.getOperand(1))
7281 ResType->
getOpcode() == SPIRV::OpTypeFloat);
7291 APFloat::rmNearestTiesToEven, &LosesInfo);
7295 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
7296 ? SPIRV::OpVectorTimesScalar
7307bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
7308 SPIRVTypeInst ResType,
7309 MachineInstr &
I)
const {
7312 return selectExtInst(ResVReg, ResType,
I, CL::pown);
7318 Register ExpReg =
I.getOperand(2).getReg();
7320 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
7321 SPIRV::OpConvertSToF))
7323 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
7330bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
7331 SPIRVTypeInst ResType,
7332 MachineInstr &
I)
const {
7348 MachineIRBuilder MIRBuilder(
I);
7349 SPIRVTypeInst FloatType =
7353 FloatType, MIRBuilder, SPIRV::StorageClass::Function);
7366 MachineBasicBlock &EntryBB =
I.getMF()->
front();
7367 const bool IsUntyped =
7368 PtrType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7370 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
7371 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
7372 : SPIRV::OpVariable))
7375 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7383 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7386 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
7389 .
add(
I.getOperand(
I.getNumExplicitDefs()))
7393 Register IntegralPartReg =
I.getOperand(1).getReg();
7396 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7406 assert(
false &&
"GLSL::Modf is deprecated.");
7417bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
7418 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7419 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7420 MachineIRBuilder MIRBuilder(
I);
7421 const SPIRVTypeInst Vec3Ty =
7424 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
7436 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7440 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7446 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7453 assert(
I.getOperand(2).isReg());
7454 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
7458 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
7469bool SPIRVInstructionSelector::loadBuiltinInputID(
7470 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7471 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7472 MachineIRBuilder MIRBuilder(
I);
7474 ResType, MIRBuilder, SPIRV::StorageClass::Input);
7489 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7493 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7502SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
7503 MachineInstr &
I)
const {
7504 MachineIRBuilder MIRBuilder(
I);
7505 if (
Type->getOpcode() != SPIRV::OpTypeVector)
7515bool SPIRVInstructionSelector::loadHandleBeforePosition(
7516 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
7517 MachineInstr &Pos)
const {
7520 Intrinsic::spv_resource_handlefrombinding);
7528 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
7529 MachineIRBuilder MIRBuilder(HandleDef);
7530 SPIRVTypeInst VarType = ResType;
7531 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
7533 if (IsStructuredBuffer) {
7538 if (ResType->
getOpcode() == SPIRV::OpTypeImage && ArraySize == 0)
7540 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
7543 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
7544 ArraySize, IndexReg, Name, MIRBuilder);
7548 uint32_t LoadOpcode =
7549 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
7559bool SPIRVInstructionSelector::errorIfInstrOutsideShader(
7560 MachineInstr &
I)
const {
7562 return diagnoseUnsupported(
7563 I,
"this instruction is only supported in shaders.");
7568InstructionSelector *
7572 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