35#include "llvm/IR/IntrinsicsSPIRV.h"
41#define DEBUG_TYPE "spirv-isel"
48 std::vector<std::pair<SPIRV::InstructionSet::InstructionSet, uint32_t>>;
53 std::optional<Register> Bias;
54 std::optional<Register>
Offset;
55 std::optional<Register> MinLod;
56 std::optional<Register> GradX;
57 std::optional<Register> GradY;
58 std::optional<Register> Lod;
59 std::optional<Register> Compare;
66 bool IsScalar =
false;
69llvm::SPIRV::SelectionControl::SelectionControl
70getSelectionOperandForImm(
int Imm) {
72 return SPIRV::SelectionControl::Flatten;
74 return SPIRV::SelectionControl::DontFlatten;
76 return SPIRV::SelectionControl::None;
80#define GET_GLOBALISEL_PREDICATE_BITSET
81#include "SPIRVGenGlobalISel.inc"
82#undef GET_GLOBALISEL_PREDICATE_BITSET
109#define GET_GLOBALISEL_PREDICATES_DECL
110#include "SPIRVGenGlobalISel.inc"
111#undef GET_GLOBALISEL_PREDICATES_DECL
113#define GET_GLOBALISEL_TEMPORARIES_DECL
114#include "SPIRVGenGlobalISel.inc"
115#undef GET_GLOBALISEL_TEMPORARIES_DECL
139 unsigned BitSetOpcode)
const;
143 unsigned BitSetOpcode)
const;
147 unsigned BitSetOpcode,
bool SwapPrimarySide)
const;
154 unsigned Opcode)
const;
157 unsigned Opcode)
const;
179 unsigned NewOpcode,
unsigned NegateOpcode = 0)
const;
188 Register castPtrToPtrToInt(
Register Ptr, SPIRV::StorageClass::StorageClass SC,
192 bool selectAtomicPtrValue(
208 unsigned OpType)
const;
275 unsigned Opcode)
const;
279 unsigned Opcode)
const;
283 unsigned Opcode)
const;
287 unsigned Opcode)
const;
289 template <
bool Signed>
292 template <
bool Signed>
299 template <
typename PickOpcodeFn>
302 PickOpcodeFn &&PickOpcode)
const;
319 template <
typename PickOpcodeFn>
322 PickOpcodeFn &&PickOpcode)
const;
340 bool IsSigned)
const;
342 bool IsSigned,
unsigned Opcode)
const;
344 bool IsSigned)
const;
350 bool IsSigned)
const;
391 GL::GLSLExtInst GLInst,
bool setMIFlags =
true,
392 bool useMISrc =
true,
394 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
395 CL::OpenCLExtInst CLInst,
bool setMIFlags =
true,
396 bool useMISrc =
true,
398 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
399 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
400 bool setMIFlags =
true,
bool useMISrc =
true,
402 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
403 const ExtInstList &ExtInsts,
bool setMIFlags =
true,
404 bool useMISrc =
true,
407 bool selectLog10(
Register ResVReg, SPIRVTypeInst ResType,
408 MachineInstr &
I)
const;
410 bool selectFpowi(
Register ResVReg, SPIRVTypeInst ResType,
411 MachineInstr &
I)
const;
413 bool selectSaturate(
Register ResVReg, SPIRVTypeInst ResType,
414 MachineInstr &
I)
const;
416 bool selectWaveOpInst(
Register ResVReg, SPIRVTypeInst ResType,
417 MachineInstr &
I,
unsigned Opcode)
const;
419 bool selectBarrierInst(MachineInstr &
I,
unsigned Scope,
unsigned MemSem,
420 bool WithGroupSync)
const;
422 bool selectWaveActiveCountBits(
Register ResVReg, SPIRVTypeInst ResType,
423 MachineInstr &
I)
const;
425 bool selectWaveActiveAllEqual(
Register ResVReg, SPIRVTypeInst ResType,
426 MachineInstr &
I)
const;
430 bool selectHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
431 MachineInstr &
I)
const;
433 bool selectCounterHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
434 MachineInstr &
I)
const;
436 bool selectReadImageIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
437 MachineInstr &
I)
const;
438 bool selectGetDimensionsIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
439 MachineInstr &
I)
const;
440 bool selectGetDimensionsLevelsIntrinsic(
Register &ResVReg,
441 SPIRVTypeInst ResType,
442 MachineInstr &
I)
const;
443 bool selectGetDimensionsMSIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
444 MachineInstr &
I)
const;
447 std::optional<Register> LodReg = std::nullopt)
const;
448 bool selectSampleBasicIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
449 MachineInstr &
I)
const;
450 bool selectCalculateLodIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
451 MachineInstr &
I)
const;
452 bool selectSampleBiasIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
453 MachineInstr &
I)
const;
454 bool selectSampleGradIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
455 MachineInstr &
I)
const;
456 bool selectSampleLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
457 MachineInstr &
I)
const;
458 bool selectLoadLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
459 MachineInstr &
I)
const;
460 bool selectSampleCmpIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
461 MachineInstr &
I)
const;
462 bool selectSampleCmpLevelZeroIntrinsic(
Register &ResVReg,
463 SPIRVTypeInst ResType,
464 MachineInstr &
I)
const;
465 bool selectGatherIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
466 MachineInstr &
I)
const;
467 bool selectImageWriteIntrinsic(MachineInstr &
I)
const;
468 bool selectResourceGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
469 MachineInstr &
I)
const;
470 bool selectPushConstantGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
471 MachineInstr &
I)
const;
472 bool selectResourceNonUniformIndex(
Register &ResVReg, SPIRVTypeInst ResType,
473 MachineInstr &
I)
const;
474 bool selectModf(
Register ResVReg, SPIRVTypeInst ResType,
475 MachineInstr &
I)
const;
476 bool selectUpdateCounter(
Register &ResVReg, SPIRVTypeInst ResType,
477 MachineInstr &
I)
const;
478 bool selectFrexp(
Register ResVReg, SPIRVTypeInst ResType,
479 MachineInstr &
I)
const;
481 bool selectLdexp(
Register ResVReg, SPIRVTypeInst ResType,
482 MachineInstr &
I)
const;
483 bool selectSincos(
Register ResVReg, SPIRVTypeInst ResType,
484 MachineInstr &
I)
const;
485 bool selectExp10(
Register ResVReg, SPIRVTypeInst ResType,
486 MachineInstr &
I)
const;
487 bool selectDerivativeInst(
Register ResVReg, SPIRVTypeInst ResType,
488 MachineInstr &
I,
const unsigned DPdOpCode)
const;
490 Register buildI32Constant(uint32_t Val, MachineInstr &
I,
491 SPIRVTypeInst ResType =
nullptr)
const;
492 Register buildI32ConstantInEntryBlock(uint32_t Val, MachineInstr &
I,
493 SPIRVTypeInst ResType =
nullptr)
const;
495 Register buildZerosVal(SPIRVTypeInst ResType, MachineInstr &
I)
const;
496 bool isScalarOrVectorIntConstantZero(
Register Reg)
const;
497 Register buildZerosValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
499 MachineInstr &
I)
const;
500 Register buildOnesValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
502 bool wrapIntoSpecConstantOp(MachineInstr &
I,
505 Register getUcharPtrTypeReg(MachineInstr &
I,
506 SPIRV::StorageClass::StorageClass SC)
const;
507 MachineInstrBuilder buildSpecConstantOp(MachineInstr &
I,
Register Dest,
509 uint32_t Opcode)
const;
510 MachineInstrBuilder buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
511 SPIRVTypeInst SrcPtrTy)
const;
512 Register buildPointerToResource(SPIRVTypeInst ResType,
513 SPIRV::StorageClass::StorageClass SC,
514 uint32_t Set, uint32_t
Binding,
515 uint32_t ArraySize,
Register IndexReg,
517 MachineIRBuilder MIRBuilder)
const;
518 SPIRVTypeInst widenTypeToVec4(SPIRVTypeInst
Type, MachineInstr &
I)
const;
519 bool extractSubvector(
Register &ResVReg, SPIRVTypeInst ResType,
520 Register &ReadReg, MachineInstr &InsertionPoint)
const;
521 bool generateImageReadOrFetch(
Register &ResVReg, SPIRVTypeInst ResType,
524 const ImageOperands *ImOps =
nullptr)
const;
525 bool generateSampleImage(
Register ResVReg, SPIRVTypeInst ResType,
527 Register CoordinateReg,
const ImageOperands &ImOps,
530 bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
531 Register ResVReg, SPIRVTypeInst ResType,
532 MachineInstr &
I)
const;
533 bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
534 Register ResVReg, SPIRVTypeInst ResType,
535 MachineInstr &
I)
const;
536 bool loadHandleBeforePosition(
Register &HandleReg, SPIRVTypeInst ResType,
537 GIntrinsic &HandleDef, MachineInstr &Pos)
const;
538 void decorateUsesAsNonUniform(
Register &NonUniformReg)
const;
539 bool errorIfInstrOutsideShader(MachineInstr &
I)
const;
541 std::optional<SplitParts> splitEvenOddLanes(
Register PopCountReg,
542 unsigned ComponentCount,
544 SPIRVTypeInst I32Type)
const;
547 handle64BitOverflow(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
548 Register SrcReg,
unsigned int Opcode,
549 std::function<
bool(
Register, SPIRVTypeInst,
550 MachineInstr &,
Register,
unsigned)>
554bool sampledTypeIsSignedInteger(
const llvm::Type *HandleType) {
556 if (
TET->getTargetExtName() ==
"spirv.Image") {
559 assert(
TET->getTargetExtName() ==
"spirv.SignedImage");
560 return TET->getTypeParameter(0)->isIntegerTy();
564#define GET_GLOBALISEL_IMPL
565#include "SPIRVGenGlobalISel.inc"
566#undef GET_GLOBALISEL_IMPL
572 TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
575#include
"SPIRVGenGlobalISel.inc"
578#include
"SPIRVGenGlobalISel.inc"
590 InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
594void SPIRVInstructionSelector::resetVRegsType(MachineFunction &MF) {
595 if (HasVRegsReset == &MF)
610 for (
const auto &
MBB : MF) {
611 for (
const auto &
MI :
MBB) {
614 if (
MI.getOpcode() != SPIRV::ASSIGN_TYPE)
618 LLT DstType = MRI.
getType(DstReg);
620 LLT SrcType = MRI.
getType(SrcReg);
621 if (DstType != SrcType)
626 if (DstRC != SrcRC && SrcRC)
638 while (!Stack.empty()) {
643 switch (
MI->getOpcode()) {
644 case TargetOpcode::G_INTRINSIC:
645 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
646 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS: {
649 if (IntrID != Intrinsic::spv_const_composite &&
650 IntrID != Intrinsic::spv_undef && IntrID != Intrinsic::spv_poison)
654 case TargetOpcode::G_BUILD_VECTOR:
655 case TargetOpcode::G_SPLAT_VECTOR:
657 i < OpDef->getNumOperands(); i++) {
662 Stack.push_back(OpNestedDef);
665 case TargetOpcode::G_CONSTANT:
666 case TargetOpcode::G_FCONSTANT:
667 case TargetOpcode::G_IMPLICIT_DEF:
668 case SPIRV::OpConstantTrue:
669 case SPIRV::OpConstantFalse:
670 case SPIRV::OpConstantI:
671 case SPIRV::OpConstantF:
672 case SPIRV::OpConstantComposite:
673 case SPIRV::OpConstantCompositeContinuedINTEL:
674 case SPIRV::OpConstantSampler:
675 case SPIRV::OpConstantNull:
677 case SPIRV::OpPoisonKHR:
678 case SPIRV::OpConstantFunctionPointerINTEL:
705 case Intrinsic::spv_all:
706 case Intrinsic::spv_alloca:
707 case Intrinsic::spv_any:
708 case Intrinsic::spv_bitcast:
709 case Intrinsic::spv_const_composite:
710 case Intrinsic::spv_cross:
711 case Intrinsic::spv_degrees:
712 case Intrinsic::spv_distance:
713 case Intrinsic::spv_extractelt:
714 case Intrinsic::spv_extractv:
715 case Intrinsic::spv_faceforward:
716 case Intrinsic::spv_fdot:
717 case Intrinsic::spv_firstbitlow:
718 case Intrinsic::spv_firstbitshigh:
719 case Intrinsic::spv_firstbituhigh:
720 case Intrinsic::spv_frac:
721 case Intrinsic::spv_gep:
722 case Intrinsic::spv_global_offset:
723 case Intrinsic::spv_global_size:
724 case Intrinsic::spv_group_id:
725 case Intrinsic::spv_insertelt:
726 case Intrinsic::spv_insertv:
727 case Intrinsic::spv_isinf:
728 case Intrinsic::spv_isnan:
729 case Intrinsic::spv_isfinite:
730 case Intrinsic::spv_isnormal:
731 case Intrinsic::spv_lerp:
732 case Intrinsic::spv_length:
733 case Intrinsic::spv_normalize:
734 case Intrinsic::spv_num_subgroups:
735 case Intrinsic::spv_num_workgroups:
736 case Intrinsic::spv_ptrcast:
737 case Intrinsic::spv_radians:
738 case Intrinsic::spv_reflect:
739 case Intrinsic::spv_refract:
740 case Intrinsic::spv_resource_getbasepointer:
741 case Intrinsic::spv_resource_getpointer:
742 case Intrinsic::spv_resource_handlefrombinding:
743 case Intrinsic::spv_resource_handlefromimplicitbinding:
744 case Intrinsic::spv_resource_nonuniformindex:
745 case Intrinsic::spv_resource_sample:
746 case Intrinsic::spv_rsqrt:
747 case Intrinsic::spv_saturate:
748 case Intrinsic::spv_sdot:
749 case Intrinsic::spv_sign:
750 case Intrinsic::spv_smoothstep:
751 case Intrinsic::spv_step:
752 case Intrinsic::spv_subgroup_id:
753 case Intrinsic::spv_subgroup_local_invocation_id:
754 case Intrinsic::spv_subgroup_max_size:
755 case Intrinsic::spv_subgroup_size:
756 case Intrinsic::spv_thread_id:
757 case Intrinsic::spv_thread_id_in_group:
758 case Intrinsic::spv_udot:
759 case Intrinsic::spv_undef:
760 case Intrinsic::spv_value_md:
761 case Intrinsic::spv_workgroup_size:
773 case SPIRV::OpTypeVoid:
774 case SPIRV::OpTypeBool:
775 case SPIRV::OpTypeInt:
776 case SPIRV::OpTypeFloat:
777 case SPIRV::OpTypeVector:
778 case SPIRV::OpTypeMatrix:
779 case SPIRV::OpTypeImage:
780 case SPIRV::OpTypeSampler:
781 case SPIRV::OpTypeSampledImage:
782 case SPIRV::OpTypeArray:
783 case SPIRV::OpTypeRuntimeArray:
784 case SPIRV::OpTypeStruct:
785 case SPIRV::OpTypeOpaque:
786 case SPIRV::OpTypePointer:
787 case SPIRV::OpTypeFunction:
788 case SPIRV::OpTypeEvent:
789 case SPIRV::OpTypeDeviceEvent:
790 case SPIRV::OpTypeReserveId:
791 case SPIRV::OpTypeQueue:
792 case SPIRV::OpTypePipe:
793 case SPIRV::OpTypeForwardPointer:
794 case SPIRV::OpTypePipeStorage:
795 case SPIRV::OpTypeNamedBarrier:
796 case SPIRV::OpTypeAccelerationStructureNV:
797 case SPIRV::OpTypeCooperativeMatrixNV:
798 case SPIRV::OpTypeCooperativeMatrixKHR:
808 if (
MI.getNumDefs() == 0)
811 for (
const auto &MO :
MI.all_defs()) {
813 if (
Reg.isPhysical()) {
818 if (
UseMI.getOpcode() != SPIRV::OpName) {
825 if (
MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE ||
MI.isFakeUse() ||
826 MI.isLifetimeMarker()) {
829 <<
"Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
840 if (
MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
841 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
844 LLVM_DEBUG(
dbgs() <<
"Dead: Intrinsic with no real side effects.\n");
849 if (
MI.mayStore() ||
MI.isCall() ||
850 (
MI.mayLoad() &&
MI.hasOrderedMemoryRef()) ||
MI.isPosition() ||
851 MI.isDebugInstr() ||
MI.isTerminator() ||
MI.isJumpTableDebugInfo()) {
852 LLVM_DEBUG(
dbgs() <<
"Not dead: instruction has side effects.\n");
863 LLVM_DEBUG(
dbgs() <<
"Dead: known opcode with no side effects\n");
870void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &
MI)
const {
872 for (
const auto &MO :
MI.all_defs()) {
876 SmallVector<MachineInstr *, 4> UselessOpNames;
879 "There is still a use of the dead function.");
882 for (MachineInstr *OpNameMI : UselessOpNames) {
884 OpNameMI->eraseFromParent();
889void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &
MI)
const {
892 removeOpNamesForDeadMI(
MI);
893 MI.eraseFromParent();
896bool SPIRVInstructionSelector::select(MachineInstr &
I) {
897 resetVRegsType(*
I.getParent()->getParent());
899 assert(
I.getParent() &&
"Instruction should be in a basic block!");
900 assert(
I.getParent()->getParent() &&
"Instruction should be in a function!");
905 removeDeadInstruction(
I);
912 if (Opcode == SPIRV::ASSIGN_TYPE) {
913 Register DstReg =
I.getOperand(0).getReg();
914 Register SrcReg =
I.getOperand(1).getReg();
917 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
918 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
919 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
920 Register SelectDstReg =
Def->getOperand(0).getReg();
921 bool SuccessToSelectSelect [[maybe_unused]] = selectSelect(
923 assert(SuccessToSelectSelect);
925 Def->eraseFromParent();
932 bool Res = selectImpl(
I, *CoverageInfo);
934 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
935 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
939 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
951 }
else if (
I.getNumDefs() == 1) {
963 removeDeadInstruction(
I);
968 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
969 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
975 bool HasDefs =
I.getNumDefs() > 0;
978 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
979 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
980 if (spvSelect(ResVReg, ResType,
I)) {
982 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
993 case TargetOpcode::G_CONSTANT:
994 case TargetOpcode::G_FCONSTANT:
1001 MachineInstr &
I)
const {
1004 if (DstRC != SrcRC && SrcRC)
1006 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
1013bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
1014 SPIRVTypeInst ResType,
1015 MachineInstr &
I)
const {
1016 const unsigned Opcode =
I.getOpcode();
1018 return selectImpl(
I, *CoverageInfo);
1020 case TargetOpcode::G_CONSTANT:
1021 case TargetOpcode::G_FCONSTANT:
1022 return selectConst(ResVReg, ResType,
I);
1023 case TargetOpcode::G_GLOBAL_VALUE:
1024 return selectGlobalValue(ResVReg,
I);
1025 case TargetOpcode::G_IMPLICIT_DEF:
1026 return selectOpUndef(ResVReg, ResType,
I);
1027 case TargetOpcode::G_FREEZE:
1028 return selectFreeze(ResVReg, ResType,
I);
1030 case TargetOpcode::G_INTRINSIC:
1031 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
1032 case TargetOpcode::G_INTRINSIC_CONVERGENT:
1033 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
1034 return selectIntrinsic(ResVReg, ResType,
I);
1035 case TargetOpcode::G_BITREVERSE:
1036 return selectBitreverse(ResVReg, ResType,
I);
1038 case TargetOpcode::G_BUILD_VECTOR:
1039 return selectBuildVector(ResVReg, ResType,
I);
1040 case TargetOpcode::G_SPLAT_VECTOR:
1041 return selectSplatVector(ResVReg, ResType,
I);
1042 case TargetOpcode::G_CONCAT_VECTORS:
1043 return selectConcatVectors(ResVReg, ResType,
I);
1045 case TargetOpcode::G_SHUFFLE_VECTOR: {
1046 MachineBasicBlock &BB = *
I.getParent();
1047 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
1050 .
addUse(
I.getOperand(1).getReg())
1051 .
addUse(
I.getOperand(2).getReg());
1052 for (
auto V :
I.getOperand(3).getShuffleMask())
1057 case TargetOpcode::G_MEMMOVE:
1058 case TargetOpcode::G_MEMCPY:
1059 case TargetOpcode::G_MEMCPY_INLINE:
1060 case TargetOpcode::G_MEMSET:
1061 case TargetOpcode::G_MEMSET_INLINE:
1062 return selectMemOperation(ResVReg,
I);
1064 case TargetOpcode::G_ICMP:
1065 return selectICmp(ResVReg, ResType,
I);
1066 case TargetOpcode::G_FCMP:
1067 return selectFCmp(ResVReg, ResType,
I);
1069 case TargetOpcode::G_FRAME_INDEX:
1070 return selectFrameIndex(ResVReg, ResType,
I);
1072 case TargetOpcode::G_LOAD:
1073 return selectLoad(ResVReg, ResType,
I);
1074 case TargetOpcode::G_STORE:
1075 return selectStore(
I);
1077 case TargetOpcode::G_BR:
1078 return selectBranch(
I);
1079 case TargetOpcode::G_BRCOND:
1080 return selectBranchCond(
I);
1082 case TargetOpcode::G_PHI:
1083 return selectPhi(ResVReg,
I);
1085 case TargetOpcode::G_FPTOSI:
1086 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1087 case TargetOpcode::G_FPTOUI:
1088 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1090 case TargetOpcode::G_FPTOSI_SAT:
1091 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1092 case TargetOpcode::G_FPTOUI_SAT:
1093 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1095 case TargetOpcode::G_SITOFP:
1096 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
1097 case TargetOpcode::G_UITOFP:
1098 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
1100 case TargetOpcode::G_CTPOP:
1101 return selectPopCount(ResVReg, ResType,
I, SPIRV::OpBitCount);
1102 case TargetOpcode::G_SMIN:
1103 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
1104 case TargetOpcode::G_UMIN:
1105 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
1107 case TargetOpcode::G_SMAX:
1108 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
1109 case TargetOpcode::G_UMAX:
1110 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
1112 case TargetOpcode::G_SCMP:
1113 return selectSUCmp(ResVReg, ResType,
I,
true);
1114 case TargetOpcode::G_UCMP:
1115 return selectSUCmp(ResVReg, ResType,
I,
false);
1116 case TargetOpcode::G_LROUND:
1117 case TargetOpcode::G_LLROUND: {
1120 MRI->
setRegClass(regForLround, &SPIRV::iIDRegClass);
1122 regForLround, *(
I.getParent()->getParent()));
1124 CL::round, GL::Round,
false);
1126 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1133 case TargetOpcode::G_STRICT_FMA:
1134 case TargetOpcode::G_FMA: {
1137 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1140 .
addUse(
I.getOperand(1).getReg())
1141 .
addUse(
I.getOperand(2).getReg())
1142 .
addUse(
I.getOperand(3).getReg())
1147 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1150 case TargetOpcode::G_FLDEXP:
1151 case TargetOpcode::G_STRICT_FLDEXP:
1152 return selectLdexp(ResVReg, ResType,
I);
1154 case TargetOpcode::G_FPOW:
1155 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1156 case TargetOpcode::G_FPOWI:
1157 return selectFpowi(ResVReg, ResType,
I);
1159 case TargetOpcode::G_FEXP:
1160 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1161 case TargetOpcode::G_FEXP2:
1162 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1163 case TargetOpcode::G_FEXP10:
1164 return selectExp10(ResVReg, ResType,
I);
1166 case TargetOpcode::G_FMODF:
1167 return selectModf(ResVReg, ResType,
I);
1168 case TargetOpcode::G_FSINCOS:
1169 return selectSincos(ResVReg, ResType,
I);
1171 case TargetOpcode::G_FLOG:
1172 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1173 case TargetOpcode::G_FLOG2:
1174 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1175 case TargetOpcode::G_FLOG10:
1176 return selectLog10(ResVReg, ResType,
I);
1178 case TargetOpcode::G_FABS:
1179 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1180 case TargetOpcode::G_ABS:
1181 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1183 case TargetOpcode::G_FMINNUM:
1184 case TargetOpcode::G_FMINIMUM:
1185 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1186 case TargetOpcode::G_FMAXNUM:
1187 case TargetOpcode::G_FMAXIMUM:
1188 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1190 case TargetOpcode::G_FCOPYSIGN:
1191 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1193 case TargetOpcode::G_FCEIL:
1194 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1195 case TargetOpcode::G_FFLOOR:
1196 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1198 case TargetOpcode::G_FCOS:
1199 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1200 case TargetOpcode::G_FSIN:
1201 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1202 case TargetOpcode::G_FTAN:
1203 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1204 case TargetOpcode::G_FACOS:
1205 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1206 case TargetOpcode::G_FASIN:
1207 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1208 case TargetOpcode::G_FATAN:
1209 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1210 case TargetOpcode::G_FATAN2:
1211 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1212 case TargetOpcode::G_FCOSH:
1213 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1214 case TargetOpcode::G_FSINH:
1215 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1216 case TargetOpcode::G_FTANH:
1217 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1219 case TargetOpcode::G_STRICT_FSQRT:
1220 case TargetOpcode::G_FSQRT:
1221 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1223 case TargetOpcode::G_CTTZ:
1224 case TargetOpcode::G_CTTZ_ZERO_POISON:
1225 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1226 case TargetOpcode::G_CTLZ:
1227 case TargetOpcode::G_CTLZ_ZERO_POISON:
1228 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1230 case TargetOpcode::G_INTRINSIC_ROUND:
1231 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1232 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1233 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1234 case TargetOpcode::G_INTRINSIC_TRUNC:
1235 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1236 case TargetOpcode::G_FRINT:
1237 case TargetOpcode::G_FNEARBYINT:
1238 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1240 case TargetOpcode::G_SMULH:
1241 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1242 case TargetOpcode::G_UMULH:
1243 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1245 case TargetOpcode::G_SADDSAT:
1246 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1247 case TargetOpcode::G_UADDSAT:
1248 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1249 case TargetOpcode::G_SSUBSAT:
1250 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1251 case TargetOpcode::G_USUBSAT:
1252 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1254 case TargetOpcode::G_FFREXP:
1255 return selectFrexp(ResVReg, ResType,
I);
1257 case TargetOpcode::G_UADDO:
1258 return selectOverflowArith(ResVReg, ResType,
I,
1259 ResType->
getOpcode() == SPIRV::OpTypeVector
1260 ? SPIRV::OpIAddCarryV
1261 : SPIRV::OpIAddCarryS);
1262 case TargetOpcode::G_USUBO:
1263 return selectOverflowArith(ResVReg, ResType,
I,
1264 ResType->
getOpcode() == SPIRV::OpTypeVector
1265 ? SPIRV::OpISubBorrowV
1266 : SPIRV::OpISubBorrowS);
1267 case TargetOpcode::G_UMULO:
1268 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1269 case TargetOpcode::G_SMULO:
1270 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1272 case TargetOpcode::G_SEXT:
1273 return selectExt(ResVReg, ResType,
I,
true);
1274 case TargetOpcode::G_ANYEXT:
1275 case TargetOpcode::G_ZEXT:
1276 return selectExt(ResVReg, ResType,
I,
false);
1277 case TargetOpcode::G_TRUNC:
1278 return selectTrunc(ResVReg, ResType,
I);
1279 case TargetOpcode::G_FPTRUNC:
1280 case TargetOpcode::G_FPEXT:
1281 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1283 case TargetOpcode::G_PTRTOINT:
1284 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1285 case TargetOpcode::G_INTTOPTR:
1286 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1287 case TargetOpcode::G_BITCAST:
1288 return selectBitcast(ResVReg, ResType,
I);
1289 case TargetOpcode::G_ADDRSPACE_CAST:
1290 return selectAddrSpaceCast(ResVReg, ResType,
I);
1291 case TargetOpcode::G_PTRMASK:
1292 return selectPtrMask(ResVReg, ResType,
I);
1293 case TargetOpcode::G_PTR_ADD: {
1295 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1299 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1300 (*II).getOpcode() == TargetOpcode::COPY ||
1301 (*II).getOpcode() == SPIRV::OpVariable) &&
1302 getImm(
I.getOperand(2), MRI));
1304 bool IsGVInit =
false;
1308 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1309 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1310 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1311 (*UseIt).getOpcode() == SPIRV::OpVariable) {
1321 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1333 return diagnoseUnsupported(
1334 I,
"incompatible result and operand types in a bitcast");
1336 MachineInstrBuilder MIB =
1337 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1344 : SPIRV::OpInBoundsPtrAccessChain))
1348 .
addUse(
I.getOperand(2).getReg())
1351 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1355 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1357 .
addUse(
I.getOperand(2).getReg())
1366 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1369 .
addImm(
static_cast<uint32_t
>(
1370 SPIRV::Opcode::InBoundsPtrAccessChain))
1373 .
addUse(
I.getOperand(2).getReg());
1378 case TargetOpcode::G_ATOMICRMW_OR:
1379 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1380 case TargetOpcode::G_ATOMICRMW_ADD:
1381 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1382 case TargetOpcode::G_ATOMICRMW_AND:
1383 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1384 case TargetOpcode::G_ATOMICRMW_MAX:
1385 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1386 case TargetOpcode::G_ATOMICRMW_MIN:
1387 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1388 case TargetOpcode::G_ATOMICRMW_SUB:
1389 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1390 case TargetOpcode::G_ATOMICRMW_XOR:
1391 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1392 case TargetOpcode::G_ATOMICRMW_UMAX:
1393 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1394 case TargetOpcode::G_ATOMICRMW_UMIN:
1395 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1396 case TargetOpcode::G_ATOMICRMW_XCHG:
1397 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1399 case TargetOpcode::G_ATOMICRMW_FADD:
1400 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1401 case TargetOpcode::G_ATOMICRMW_FSUB:
1403 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1404 ResType->
getOpcode() == SPIRV::OpTypeVector
1406 : SPIRV::OpFNegate);
1407 case TargetOpcode::G_ATOMICRMW_FMIN:
1408 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1409 case TargetOpcode::G_ATOMICRMW_FMAX:
1410 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1412 case TargetOpcode::G_FENCE:
1413 return selectFence(
I);
1415 case TargetOpcode::G_STACKSAVE:
1416 return selectStackSave(ResVReg, ResType,
I);
1417 case TargetOpcode::G_STACKRESTORE:
1418 return selectStackRestore(
I);
1420 case TargetOpcode::G_UNMERGE_VALUES:
1423 case TargetOpcode::G_TRAP:
1424 case TargetOpcode::G_UBSANTRAP:
1425 return selectTrap(
I);
1430 case TargetOpcode::DBG_LABEL:
1432 case TargetOpcode::G_DEBUGTRAP:
1433 return selectDebugTrap(ResVReg, ResType,
I);
1440bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1441 SPIRVTypeInst ResType,
1442 MachineInstr &
I)
const {
1443 unsigned Opcode = SPIRV::OpNop;
1450bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1451 SPIRVTypeInst ResType,
1453 GL::GLSLExtInst GLInst,
1454 bool setMIFlags,
bool useMISrc,
1457 SPIRV::InstructionSet::InstructionSet::GLSL_std_450))
1458 return diagnoseUnsupported(
1460 "this instruction is only supported with the GLSL extended instruction "
1462 return selectExtInst(ResVReg, ResType,
I,
1463 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1464 setMIFlags, useMISrc, SrcRegs);
1467bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1468 SPIRVTypeInst ResType,
1470 CL::OpenCLExtInst CLInst,
1471 bool setMIFlags,
bool useMISrc,
1473 return selectExtInst(ResVReg, ResType,
I,
1474 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1475 setMIFlags, useMISrc, SrcRegs);
1478bool SPIRVInstructionSelector::selectExtInst(
1479 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1480 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1482 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1483 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1484 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1488bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1489 SPIRVTypeInst ResType,
1492 bool setMIFlags,
bool useMISrc,
1495 for (
const auto &[InstructionSet, Opcode] : Insts) {
1499 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1502 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1507 const unsigned NumOps =
I.getNumOperands();
1510 I.getOperand(Index).getType() ==
1511 MachineOperand::MachineOperandType::MO_IntrinsicID)
1514 MIB.
add(
I.getOperand(Index));
1526bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1527 SPIRVTypeInst ResType,
1528 MachineInstr &
I)
const {
1529 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1530 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1531 for (
const auto &Ex : ExtInsts) {
1532 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1533 uint32_t Opcode = Ex.second;
1537 MachineIRBuilder MIRBuilder(
I);
1540 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1545 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1548 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1552 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1555 .
addImm(
static_cast<uint32_t
>(Ex.first))
1557 .
add(
I.getOperand(2))
1561 Register ExpResReg =
I.getOperand(1).getReg();
1563 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1573bool SPIRVInstructionSelector::selectLdexp(
Register ResVReg,
1574 SPIRVTypeInst ResType,
1575 MachineInstr &
I)
const {
1576 Register XReg =
I.getOperand(1).getReg();
1577 Register ExpReg =
I.getOperand(2).getReg();
1583 if (ResType->
getOpcode() == SPIRV::OpTypeVector &&
1584 ExpType->
getOpcode() != SPIRV::OpTypeVector) {
1586 SPIRVTypeInst ExpVecType =
1590 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1591 TII.get(SPIRV::OpCompositeConstruct))
1594 for (
unsigned J = 0; J < NumElts; ++J)
1600 return selectExtInst(ResVReg, ResType,
I, CL::ldexp, GL::Ldexp,
1601 true,
false, {XReg, ExpReg});
1604bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1605 SPIRVTypeInst ResType,
1606 MachineInstr &
I)
const {
1607 Register CosResVReg =
I.getOperand(1).getReg();
1608 unsigned SrcIdx =
I.getNumExplicitDefs();
1613 MachineIRBuilder MIRBuilder(
I);
1615 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1620 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1623 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1625 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1628 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1630 .
add(
I.getOperand(SrcIdx))
1633 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1641 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1644 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1646 .
add(
I.getOperand(SrcIdx))
1648 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1651 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1653 .
add(
I.getOperand(SrcIdx))
1660bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1661 SPIRVTypeInst ResType,
1664 unsigned Opcode)
const {
1665 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1675std::optional<SplitParts> SPIRVInstructionSelector::splitEvenOddLanes(
1676 Register PopCountReg,
unsigned ComponentCount, MachineInstr &
I,
1677 SPIRVTypeInst I32Type)
const {
1680 if (ComponentCount == 1) {
1683 Parts.IsScalar =
true;
1684 Parts.Type = I32Type;
1692 if (!selectOpWithSrcs(Parts.High, I32Type,
I, {PopCountReg, IdxOne},
1693 SPIRV::OpVectorExtractDynamic))
1694 return std::nullopt;
1696 if (!selectOpWithSrcs(Parts.Low, I32Type,
I, {PopCountReg, IdxZero},
1697 SPIRV::OpVectorExtractDynamic))
1698 return std::nullopt;
1702 MachineIRBuilder MIRBuilder(
I);
1703 Parts.IsScalar =
false;
1710 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1711 TII.get(SPIRV::OpVectorShuffle))
1716 for (
unsigned J = 1; J < ComponentCount * 2; J += 2)
1721 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1722 TII.get(SPIRV::OpVectorShuffle))
1727 for (
unsigned J = 0; J < ComponentCount * 2; J += 2)
1735bool SPIRVInstructionSelector::selectPopCount16(
Register ResVReg,
1736 SPIRVTypeInst ResType,
1739 unsigned Opcode)
const {
1740 Register OpReg =
I.getOperand(1).getReg();
1743 MachineIRBuilder MIRBuilder(
I);
1745 SPIRVTypeInst I32VectorType =
1748 bool IsVector = NumElems > 1;
1749 SPIRVTypeInst ExtType = IsVector ? I32VectorType : I32Type;
1752 if (!selectOpWithSrcs(ExtReg, ExtType,
I, {OpReg}, SPIRV::OpUConvert))
1756 if (!selectPopCount32(PopCountReg, ExtType,
I, ExtReg, Opcode))
1759 return selectOpWithSrcs(ResVReg, ResType,
I, {PopCountReg}, ExtOpcode);
1762bool SPIRVInstructionSelector::selectPopCount32(
Register ResVReg,
1763 SPIRVTypeInst ResType,
1766 unsigned Opcode)
const {
1767 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
1770bool SPIRVInstructionSelector::selectPopCount64(
Register ResVReg,
1771 SPIRVTypeInst ResType,
1774 unsigned Opcode)
const {
1776 if (ComponentCount > 2)
1777 return handle64BitOverflow(
1778 ResVReg, ResType,
I, SrcReg, Opcode,
1780 unsigned O) {
return this->selectPopCount64(R,
T,
I, S, O); });
1782 MachineIRBuilder MIRBuilder(
I);
1787 I32Type, 2 * ComponentCount, MIRBuilder,
false);
1791 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
1796 if (!selectPopCount32(Pop32, VecI32Type,
I, Vec32, Opcode))
1800 auto MaybeParts = splitEvenOddLanes(Pop32, ComponentCount,
I, I32Type);
1803 SplitParts &Parts = *MaybeParts;
1806 unsigned OpAdd = Parts.IsScalar ? SPIRV::OpIAddS : SPIRV::OpIAddV;
1808 if (!selectOpWithSrcs(Sum, Parts.Type,
I, {Parts.High, Parts.Low}, OpAdd))
1813 unsigned ConvOp = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
1814 return selectOpWithSrcs(ResVReg, ResType,
I, {Sum}, ConvOp);
1817bool SPIRVInstructionSelector::selectPopCount(
Register ResVReg,
1818 SPIRVTypeInst ResType,
1820 unsigned Opcode)
const {
1825 if (!STI.getTargetTriple().isVulkanOS())
1826 return selectUnOp(ResVReg, ResType,
I, Opcode);
1828 Register OpReg =
I.getOperand(1).getReg();
1831 : SPIRV::OpUConvert;
1835 return selectPopCount16(ResVReg, ResType,
I, ExtOpcode, Opcode);
1837 return selectPopCount32(ResVReg, ResType,
I, OpReg, Opcode);
1839 return selectPopCount64(ResVReg, ResType,
I, OpReg, Opcode);
1841 return diagnoseUnsupported(
I,
"unsupported operand bit width for popcount");
1845bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1846 SPIRVTypeInst ResType,
1848 unsigned Opcode)
const {
1850 Register SrcReg =
I.getOperand(1).getReg();
1855 unsigned DefOpCode = DefIt->getOpcode();
1856 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1859 if (
auto *VRD =
getVRegDef(*MRI, DefIt->getOperand(1).getReg()))
1860 DefOpCode = VRD->getOpcode();
1862 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1863 DefOpCode == TargetOpcode::G_CONSTANT ||
1864 DefOpCode == SPIRV::OpVariable || DefOpCode == SPIRV::OpConstantI) {
1870 uint32_t SpecOpcode = 0;
1872 case SPIRV::OpConvertPtrToU:
1873 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1875 case SPIRV::OpConvertUToPtr:
1876 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1881 TII.get(SPIRV::OpSpecConstantOp))
1891 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1895bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1896 SPIRVTypeInst ResType,
1897 MachineInstr &
I)
const {
1898 Register OpReg =
I.getOperand(1).getReg();
1899 SPIRVTypeInst OpType =
1902 return diagnoseUnsupported(
1903 I,
"incompatible result and operand types in a bitcast");
1904 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1915 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1916 if (
MemOp->isNonTemporal())
1917 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1919 if (!ST->isShader() &&
MemOp->getAlign().value())
1920 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1924 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1925 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1929 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1931 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1935 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1939 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1941 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1953 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1955 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1957 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
1961bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
1962 SPIRVTypeInst ResType,
1963 MachineInstr &
I)
const {
1965 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1970 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
1971 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
1973 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1975 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1979 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1983 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1984 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
1985 I.getDebugLoc(),
I);
1989 MachineIRBuilder MIRBuilder(
I);
1991 if (
I.getNumMemOperands()) {
1992 const MachineMemOperand *MemOp = *
I.memoperands_begin();
1993 if (MemOp->isAtomic())
1994 return selectAtomicLoad(ResVReg, ResType,
I);
1997 auto MIB = MIRBuilder.
buildInstr(SPIRV::OpLoad)
2001 if (!
I.getNumMemOperands()) {
2002 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2004 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2013Register SPIRVInstructionSelector::createPtrSizedIntReg(
2014 MachineIRBuilder &MIRBuilder)
const {
2015 SPIRVTypeInst IntType =
2025SPIRVInstructionSelector::convertPtrToInt(
Register PtrVal,
2026 MachineIRBuilder &MIRBuilder)
const {
2027 SPIRVTypeInst IntType =
2029 Register IntReg = createPtrSizedIntReg(MIRBuilder);
2030 MIRBuilder.
buildInstr(SPIRV::OpConvertPtrToU)
2038Register SPIRVInstructionSelector::castPtrToPtrToInt(
2039 Register Ptr, SPIRV::StorageClass::StorageClass SC,
2040 MachineIRBuilder &MIRBuilder)
const {
2041 SPIRVTypeInst IntType =
2043 SPIRVTypeInst PtrType =
2057bool SPIRVInstructionSelector::selectAtomicPtrValue(
2058 Register ResVReg, SPIRVTypeInst ResType, MachineIRBuilder &MIRBuilder,
2059 function_ref<
Register(SPIRVTypeInst IntType)> EmitAtomic)
const {
2067 Register IntResult = EmitAtomic(IntType);
2069 MIRBuilder.
buildInstr(SPIRV::OpConvertUToPtr)
2077bool SPIRVInstructionSelector::selectAtomicLoad(
Register ResVReg,
2078 SPIRVTypeInst ResType,
2079 MachineInstr &
I)
const {
2080 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2083 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2086 return diagnoseUnsupported(
2087 I,
"Lowering to SPIR-V of atomic load is only "
2088 "allowed for integer, floating point or pointer types");
2090 assert(
I.getNumMemOperands());
2091 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2092 assert(MemOp.isAtomic());
2096 Register ScopeReg = buildI32Constant(Scope,
I);
2102 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2103 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2106 MachineIRBuilder MIRBuilder(
I);
2110 return diagnoseUnsupported(
2111 I,
"Lowering to SPIR-V of atomic load is only "
2112 "allowed for pointer types for physical addressing model");
2117 SPIRV::StorageClass::StorageClass SC =
2119 return selectAtomicPtrValue(
2120 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2121 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2122 Register IntResult = createPtrSizedIntReg(MIRBuilder);
2133 auto AtomicLoad = MIRBuilder.
buildInstr(SPIRV::OpAtomicLoad)
2144bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
2146 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2147 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2152 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2153 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2155 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2160 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2164 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2165 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2166 SPIRVTypeInst SampledType =
2168 SPIRVTypeInst StoreValCompType =
2170 if (StoreValCompType && StoreValCompType != SampledType) {
2173 SPIRVTypeInst PackedType = widenTypeToVec4(SampledType,
I);
2176 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
2181 StoreVal = PackedReg;
2184 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2185 TII.get(SPIRV::OpImageWrite))
2191 if (sampledTypeIsSignedInteger(LLVMHandleType))
2194 BMI.constrainAllUses(
TII,
TRI, RBI);
2199 if (
I.getNumMemOperands()) {
2200 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2201 if (MemOp->isAtomic())
2202 return selectAtomicStore(
I);
2209 if (PtrSC == SPIRV::StorageClass::UniformConstant ||
2210 PtrSC == SPIRV::StorageClass::Input ||
2211 PtrSC == SPIRV::StorageClass::PushConstant)
2212 return diagnoseUnsupported(
2213 I,
"store into a read-only SPIR-V storage class is not allowed");
2215 MachineIRBuilder MIRBuilder(
I);
2217 if (!
I.getNumMemOperands()) {
2218 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2220 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2229bool SPIRVInstructionSelector::selectAtomicStore(MachineInstr &
I)
const {
2230 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2233 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2234 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2239 assert(
I.getNumMemOperands());
2240 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2241 assert(MemOp.isAtomic());
2245 Register ScopeReg = buildI32Constant(Scope,
I);
2251 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2252 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2254 MachineIRBuilder MIRBuilder(
I);
2258 return diagnoseUnsupported(
2259 I,
"Lowering to SPIR-V of atomic store is only "
2260 "allowed for pointer types for physical addressing model");
2265 SPIRV::StorageClass::StorageClass SC =
2267 return selectAtomicPtrValue(
2268 Register(), SPIRVTypeInst(), MIRBuilder, [&](SPIRVTypeInst IntType) {
2270 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2283 return diagnoseUnsupported(
I,
2284 "Lowering to SPIR-V of atomic store is only "
2285 "allowed for integer or floating point types");
2287 auto AtomicStore = MIRBuilder.
buildInstr(SPIRV::OpAtomicStore)
2297bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
2298 SPIRVTypeInst ResType,
2299 MachineInstr &
I)
const {
2300 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
2308 const Register PtrsReg =
I.getOperand(2).getReg();
2309 const uint32_t Alignment =
I.getOperand(3).getImm();
2310 const Register MaskReg =
I.getOperand(4).getReg();
2311 const Register PassthruReg =
I.getOperand(5).getReg();
2312 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2316 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
2327bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
2328 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
2335 const Register ValuesReg =
I.getOperand(1).getReg();
2336 const Register PtrsReg =
I.getOperand(2).getReg();
2337 const uint32_t Alignment =
I.getOperand(3).getImm();
2338 const Register MaskReg =
I.getOperand(4).getReg();
2339 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2343 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
2352bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
2353 const Twine &
Msg)
const {
2354 const Function &
F =
I.getMF()->getFunction();
2355 F.getContext().diagnose(
2356 DiagnosticInfoUnsupported(
F,
Msg,
I.getDebugLoc(),
DS_Error));
2360bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
2361 SPIRVTypeInst ResType,
2362 MachineInstr &
I)
const {
2363 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2364 return diagnoseUnsupported(
2365 I,
"llvm.stacksave intrinsic: this instruction requires the following "
2366 "SPIR-V extension: SPV_INTEL_variable_length_array");
2368 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
2375bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
2376 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2377 return diagnoseUnsupported(
2379 "llvm.stackrestore intrinsic: this instruction requires the following "
2380 "SPIR-V extension: SPV_INTEL_variable_length_array");
2381 if (!
I.getOperand(0).isReg())
2384 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
2385 .
addUse(
I.getOperand(0).getReg())
2391SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
2392 MachineIRBuilder MIRBuilder(
I);
2393 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
2400 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
2404 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2405 Type *ArrTy = ArrayType::get(ValTy, Num);
2407 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
2410 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
2417 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
2420 .
addImm(SPIRV::StorageClass::UniformConstant)
2431bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
2434 Register DstReg =
I.getOperand(0).getReg();
2438 return diagnoseUnsupported(
2439 I,
"OpCopyMemory requires operands to have the same type");
2440 uint64_t CopySize =
getIConstVal(
I.getOperand(2).getReg(), MRI);
2444 return diagnoseUnsupported(
2445 I,
"Unable to determine pointee type size for OpCopyMemory");
2446 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
2447 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
2448 return diagnoseUnsupported(
2449 I,
"OpCopyMemory requires the size to match the pointee type size");
2450 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemory))
2453 if (
I.getNumMemOperands()) {
2454 MachineIRBuilder MIRBuilder(
I);
2461bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &
I,
2464 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemorySized))
2465 .
addUse(
I.getOperand(0).getReg())
2467 .
addUse(
I.getOperand(2).getReg());
2468 if (
I.getNumMemOperands()) {
2469 MachineIRBuilder MIRBuilder(
I);
2476bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
2477 MachineInstr &
I)
const {
2479 Register SizeReg =
I.getOperand(2).getReg();
2481 SizeDef && SizeDef->
getOpcode() == TargetOpcode::G_CONSTANT &&
2485 Register SrcReg =
I.getOperand(1).getReg();
2486 if (
I.getOpcode() == TargetOpcode::G_MEMSET ||
2487 I.getOpcode() == TargetOpcode::G_MEMSET_INLINE) {
2488 Register VarReg = getOrCreateMemSetGlobal(
I);
2491 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2493 ValTy,
I, SPIRV::StorageClass::UniformConstant);
2495 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
2499 if (!selectCopyMemory(
I, SrcReg))
2502 if (!selectCopyMemorySized(
I, SrcReg))
2505 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
2506 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
2511bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
2512 SPIRVTypeInst ResType,
2515 unsigned NegateOpcode)
const {
2517 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2520 Register ScopeReg = buildI32Constant(Scope,
I);
2522 Register Ptr =
I.getOperand(1).getReg();
2523 uint32_t ScSem =
static_cast<uint32_t
>(
2527 Register MemSemReg = buildI32Constant(MemSem,
I);
2529 Register ValueReg =
I.getOperand(2).getReg();
2530 if (NegateOpcode != 0) {
2533 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
2539 if (NewOpcode != SPIRV::OpAtomicExchange)
2540 return diagnoseUnsupported(
2541 I,
"Lowering to SPIR-V of this atomic operation is not "
2542 "allowed for pointer types");
2544 return diagnoseUnsupported(
2545 I,
"Lowering to SPIR-V of atomic exchange is only "
2546 "allowed for pointer types for physical addressing model");
2553 MachineIRBuilder MIRBuilder(
I);
2555 return selectAtomicPtrValue(
2556 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2558 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2559 Register ExchangeResReg = createPtrSizedIntReg(MIRBuilder);
2560 MIRBuilder.
buildInstr(SPIRV::OpAtomicExchange)
2568 return ExchangeResReg;
2572 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2583bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2584 unsigned ArgI =
I.getNumOperands() - 1;
2586 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2587 SPIRVTypeInst SrcType =
2589 if (!SrcType || SrcType->
getOpcode() != SPIRV::OpTypeVector)
2591 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2595 unsigned CurrentIndex = 0;
2596 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2597 Register ResVReg =
I.getOperand(i).getReg();
2600 LLT ResLLT = MRI->
getType(ResVReg);
2606 ResType = ScalarType;
2612 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
2615 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2621 for (
unsigned j = 0;
j < NumElements; ++
j) {
2622 MIB.
addImm(CurrentIndex + j);
2624 CurrentIndex += NumElements;
2628 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2640bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2643 Register MemSemReg = buildI32Constant(MemSem,
I);
2645 uint32_t
Scope =
static_cast<uint32_t
>(
2647 Register ScopeReg = buildI32Constant(Scope,
I);
2649 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2656bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2657 SPIRVTypeInst ResType,
2659 unsigned Opcode)
const {
2660 Type *ResTy =
nullptr;
2663 return diagnoseUnsupported(
2665 "Not enough info to select the arithmetic with overflow instruction");
2667 return diagnoseUnsupported(
I,
2668 "Expect struct type result for the arithmetic "
2669 "with overflow instruction");
2675 MachineIRBuilder MIRBuilder(
I);
2677 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2678 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2684 Register ZeroReg = buildZerosVal(ResType,
I);
2689 if (ResName.
size() > 0)
2697 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2698 MIB.
addUse(
I.getOperand(i).getReg());
2703 MRI->
setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2704 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2706 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2707 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2714 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2715 .
addDef(
I.getOperand(1).getReg())
2723bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2724 SPIRVTypeInst ResType,
2725 MachineInstr &
I)
const {
2727 "selectAtomicCmpXchg only handles the spv_cmpxchg intrinsic");
2728 Register Ptr =
I.getOperand(2).getReg();
2729 Register ScopeReg =
I.getOperand(5).getReg();
2730 Register MemSemEqReg =
I.getOperand(6).getReg();
2731 Register MemSemNeqReg =
I.getOperand(7).getReg();
2733 Register Val =
I.getOperand(4).getReg();
2737 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2756 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2763 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2775 case SPIRV::StorageClass::DeviceOnlyINTEL:
2776 case SPIRV::StorageClass::HostOnlyINTEL:
2785 bool IsGRef =
false;
2786 bool IsAllowedRefs =
2788 unsigned Opcode = It.getOpcode();
2789 if (Opcode == SPIRV::OpConstantComposite ||
2790 Opcode == SPIRV::OpSpecConstantComposite ||
2791 Opcode == SPIRV::OpVariable ||
2792 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2793 return IsGRef = true;
2794 return Opcode == SPIRV::OpName;
2796 return IsAllowedRefs && IsGRef;
2799Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2800 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2802 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2806SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2808 uint32_t Opcode)
const {
2809 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2810 TII.get(SPIRV::OpSpecConstantOp))
2818SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2819 SPIRVTypeInst SrcPtrTy)
const {
2820 SPIRVTypeInst GenericPtrTy =
2824 SPIRV::StorageClass::Generic),
2826 MachineFunction *MF =
I.getParent()->getParent();
2828 MachineInstrBuilder MIB = buildSpecConstantOp(
2830 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2840bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2841 SPIRVTypeInst ResType,
2842 MachineInstr &
I)
const {
2846 Register SrcPtr =
I.getOperand(1).getReg();
2850 if (SrcPtrTy->
getOpcode() != SPIRV::OpTypePointer ||
2851 ResType->
getOpcode() != SPIRV::OpTypePointer)
2852 return BuildCOPY(ResVReg, SrcPtr,
I);
2862 unsigned SpecOpcode =
2864 ?
static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric)
2867 ? static_cast<uint32_t>(
SPIRV::Opcode::GenericCastToPtr)
2874 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2876 .constrainAllUses(
TII,
TRI, RBI);
2878 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2880 buildSpecConstantOp(
2882 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2883 .constrainAllUses(
TII,
TRI, RBI);
2890 return BuildCOPY(ResVReg, SrcPtr,
I);
2892 if ((SrcSC == SPIRV::StorageClass::Function &&
2893 DstSC == SPIRV::StorageClass::Private) ||
2894 (DstSC == SPIRV::StorageClass::Function &&
2895 SrcSC == SPIRV::StorageClass::Private))
2896 return BuildCOPY(ResVReg, SrcPtr,
I);
2900 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2903 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2906 SPIRVTypeInst GenericPtrTy =
2925 return selectUnOp(ResVReg, ResType,
I,
2926 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
2928 return selectUnOp(ResVReg, ResType,
I,
2929 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
2931 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2933 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2943bool SPIRVInstructionSelector::selectPtrMask(
Register ResVReg,
2944 SPIRVTypeInst ResType,
2945 MachineInstr &
I)
const {
2947 return diagnoseUnsupported(
2948 I,
"G_PTRMASK is not supported with logical SPIR-V");
2953 Register PtrReg =
I.getOperand(1).getReg();
2954 Register MaskReg =
I.getOperand(2).getReg();
2973 ? SPIRV::OpBitwiseAndV
2974 : SPIRV::OpBitwiseAndS;
2997 return SPIRV::OpFOrdEqual;
2999 return SPIRV::OpFOrdGreaterThanEqual;
3001 return SPIRV::OpFOrdGreaterThan;
3003 return SPIRV::OpFOrdLessThanEqual;
3005 return SPIRV::OpFOrdLessThan;
3007 return SPIRV::OpFOrdNotEqual;
3009 return SPIRV::OpOrdered;
3011 return SPIRV::OpFUnordEqual;
3013 return SPIRV::OpFUnordGreaterThanEqual;
3015 return SPIRV::OpFUnordGreaterThan;
3017 return SPIRV::OpFUnordLessThanEqual;
3019 return SPIRV::OpFUnordLessThan;
3021 return SPIRV::OpFUnordNotEqual;
3023 return SPIRV::OpUnordered;
3033 return SPIRV::OpIEqual;
3035 return SPIRV::OpINotEqual;
3037 return SPIRV::OpSGreaterThanEqual;
3039 return SPIRV::OpSGreaterThan;
3041 return SPIRV::OpSLessThanEqual;
3043 return SPIRV::OpSLessThan;
3045 return SPIRV::OpUGreaterThanEqual;
3047 return SPIRV::OpUGreaterThan;
3049 return SPIRV::OpULessThanEqual;
3051 return SPIRV::OpULessThan;
3060 return SPIRV::OpPtrEqual;
3062 return SPIRV::OpPtrNotEqual;
3073 return SPIRV::OpLogicalEqual;
3075 return SPIRV::OpLogicalNotEqual;
3113bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
3114 SPIRVTypeInst ResType,
3116 unsigned OpAnyOrAll)
const {
3117 assert(
I.getNumOperands() == 3);
3118 assert(
I.getOperand(2).isReg());
3120 Register InputRegister =
I.getOperand(2).getReg();
3123 assert(InputType &&
"VReg has no type assigned");
3126 bool IsVectorTy = InputType->
getOpcode() == SPIRV::OpTypeVector;
3127 if (IsBoolTy && !IsVectorTy) {
3128 assert(ResVReg ==
I.getOperand(0).getReg());
3129 return BuildCOPY(ResVReg, InputRegister,
I);
3133 unsigned SpirvNotEqualId =
3134 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
3136 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
3141 IsBoolTy ? InputRegister
3149 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
3151 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
3168bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
3169 SPIRVTypeInst ResType,
3170 MachineInstr &
I)
const {
3171 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
3174bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
3175 SPIRVTypeInst ResType,
3176 MachineInstr &
I)
const {
3177 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
3181bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
3182 SPIRVTypeInst ResType,
3183 MachineInstr &
I)
const {
3184 assert(
I.getNumOperands() == 4);
3185 assert(
I.getOperand(2).isReg());
3186 assert(
I.getOperand(3).isReg());
3188 [[maybe_unused]] SPIRVTypeInst VecType =
3193 "dot product requires a vector of at least 2 components");
3195 [[maybe_unused]] SPIRVTypeInst EltType =
3204 .
addUse(
I.getOperand(2).getReg())
3205 .
addUse(
I.getOperand(3).getReg())
3210bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
3211 SPIRVTypeInst ResType,
3214 assert(
I.getNumOperands() == 4);
3215 assert(
I.getOperand(2).isReg());
3216 assert(
I.getOperand(3).isReg());
3219 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3223 .
addUse(
I.getOperand(2).getReg())
3224 .
addUse(
I.getOperand(3).getReg())
3231bool SPIRVInstructionSelector::selectIntegerDotExpansion(
3232 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3233 assert(
I.getNumOperands() == 4);
3234 assert(
I.getOperand(2).isReg());
3235 assert(
I.getOperand(3).isReg());
3239 Register Vec0 =
I.getOperand(2).getReg();
3240 Register Vec1 =
I.getOperand(3).getReg();
3244 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
3253 "dot product requires a vector of at least 2 components");
3256 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3266 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3277 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3289bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
3290 SPIRVTypeInst ResType,
3291 MachineInstr &
I)
const {
3293 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
3296 .
addUse(
I.getOperand(2).getReg())
3301bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
3302 SPIRVTypeInst ResType,
3303 MachineInstr &
I)
const {
3305 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
3308 .
addUse(
I.getOperand(2).getReg())
3313bool SPIRVInstructionSelector::selectOpIsFinite(
Register ResVReg,
3314 SPIRVTypeInst ResType,
3315 MachineInstr &
I)
const {
3317 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsFinite))
3320 .
addUse(
I.getOperand(2).getReg())
3325bool SPIRVInstructionSelector::selectOpIsNormal(
Register ResVReg,
3326 SPIRVTypeInst ResType,
3327 MachineInstr &
I)
const {
3329 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNormal))
3332 .
addUse(
I.getOperand(2).getReg())
3337template <
bool Signed>
3338bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
3339 SPIRVTypeInst ResType,
3340 MachineInstr &
I)
const {
3341 assert(
I.getNumOperands() == 5);
3342 assert(
I.getOperand(2).isReg());
3343 assert(
I.getOperand(3).isReg());
3344 assert(
I.getOperand(4).isReg());
3347 Register Acc =
I.getOperand(2).getReg();
3351 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3353 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
3358 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
3361 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3373template <
bool Signed>
3374bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
3375 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3376 assert(
I.getNumOperands() == 5);
3377 assert(
I.getOperand(2).isReg());
3378 assert(
I.getOperand(3).isReg());
3379 assert(
I.getOperand(4).isReg());
3382 Register Acc =
I.getOperand(2).getReg();
3388 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
3392 for (
unsigned i = 0; i < 4; i++) {
3415 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
3435 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3450bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
3451 SPIRVTypeInst ResType,
3452 MachineInstr &
I)
const {
3453 assert(
I.getNumOperands() == 3);
3454 assert(
I.getOperand(2).isReg());
3456 Register VZero = buildZerosValF(ResType,
I);
3457 Register VOne = buildOnesValF(ResType,
I);
3459 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
3462 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3464 .
addUse(
I.getOperand(2).getReg())
3471bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
3472 SPIRVTypeInst ResType,
3473 MachineInstr &
I)
const {
3474 assert(
I.getNumOperands() == 3);
3475 assert(
I.getOperand(2).isReg());
3477 Register InputRegister =
I.getOperand(2).getReg();
3479 auto &
DL =
I.getDebugLoc();
3482 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3489 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
3491 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
3499 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3504 if (NeedsConversion) {
3505 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
3516bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
3517 SPIRVTypeInst ResType,
3519 unsigned Opcode)
const {
3523 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
3529 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
3530 BMI.addUse(
I.getOperand(J).getReg());
3537bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
3540 bool WithGroupSync)
const {
3542 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
3544 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
3546 assert(((Scope != SPIRV::Scope::Workgroup) ||
3547 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
3548 "Workgroup Scope must set WorkGroupMemory semantic "
3549 "in Barrier instruction");
3551 assert(((Scope != SPIRV::Scope::Device) ||
3552 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
3553 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
3554 "Device Scope must set UniformMemory and ImageMemory semantic "
3555 "in Barrier instruction");
3561 if (WithGroupSync) {
3562 Register ExecReg = buildI32Constant(SPIRV::Scope::Workgroup,
I);
3566 Register ScopeReg = buildI32Constant(Scope,
I);
3567 Register MemSemReg = buildI32Constant(MemSem,
I);
3569 MI.addUse(ScopeReg).addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
3573bool SPIRVInstructionSelector::selectWaveActiveCountBits(
3574 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3579 if (!selectWaveOpInst(BallotReg, BallotType,
I,
3580 SPIRV::OpGroupNonUniformBallot))
3585 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3590 .
addImm(SPIRV::GroupOperation::Reduce)
3597bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
3598 SPIRVTypeInst ResType,
3599 MachineInstr &
I)
const {
3604 Register InputReg =
I.getOperand(2).getReg();
3609 bool IsVector = NumElems > 1;
3622 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
3623 SPIRV::OpGroupNonUniformAllEqual);
3628 ElementResults.
reserve(NumElems);
3630 for (
unsigned Idx = 0; Idx < NumElems; ++Idx) {
3643 ElemInput = Extracted;
3649 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3660 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3671bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3672 SPIRVTypeInst ResType,
3673 MachineInstr &
I)
const {
3675 assert(
I.getNumOperands() == 3);
3677 auto Op =
I.getOperand(2);
3687 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3689 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3690 return diagnoseUnsupported(
I,
"WavePrefixBitCount requires boolean input");
3711 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3715 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3722bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3723 SPIRVTypeInst ResType,
3725 bool IsUnsigned)
const {
3726 return selectWaveReduce(
3727 ResVReg, ResType,
I, IsUnsigned,
3728 [&](
Register InputRegister,
bool IsUnsigned) {
3729 const bool IsFloatTy =
3731 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3732 : SPIRV::OpGroupNonUniformSMax;
3733 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3737bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3738 SPIRVTypeInst ResType,
3740 bool IsUnsigned)
const {
3741 return selectWaveReduce(
3742 ResVReg, ResType,
I, IsUnsigned,
3743 [&](
Register InputRegister,
bool IsUnsigned) {
3744 const bool IsFloatTy =
3746 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3747 : SPIRV::OpGroupNonUniformSMin;
3748 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3752bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3753 SPIRVTypeInst ResType,
3754 MachineInstr &
I)
const {
3755 return selectWaveReduce(ResVReg, ResType,
I,
false,
3756 [&](
Register InputRegister,
bool IsUnsigned) {
3758 InputRegister, SPIRV::OpTypeFloat);
3759 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3760 : SPIRV::OpGroupNonUniformIAdd;
3764bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3765 SPIRVTypeInst ResType,
3766 MachineInstr &
I)
const {
3767 return selectWaveReduce(ResVReg, ResType,
I,
false,
3768 [&](
Register InputRegister,
bool IsUnsigned) {
3770 InputRegister, SPIRV::OpTypeFloat);
3771 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3772 : SPIRV::OpGroupNonUniformIMul;
3776template <
typename PickOpcodeFn>
3777bool SPIRVInstructionSelector::selectWaveReduce(
3778 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3779 PickOpcodeFn &&PickOpcode)
const {
3780 assert(
I.getNumOperands() == 3);
3781 assert(
I.getOperand(2).isReg());
3783 Register InputRegister =
I.getOperand(2).getReg();
3787 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3790 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3796 .
addImm(SPIRV::GroupOperation::Reduce)
3797 .
addUse(
I.getOperand(2).getReg())
3802bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3803 SPIRVTypeInst ResType,
3805 unsigned Opcode)
const {
3806 return selectWaveReduce(
3807 ResVReg, ResType,
I,
false,
3808 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3811bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3812 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3813 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3814 [&](
Register InputRegister,
bool IsUnsigned) {
3816 InputRegister, SPIRV::OpTypeFloat);
3818 ? SPIRV::OpGroupNonUniformFAdd
3819 : SPIRV::OpGroupNonUniformIAdd;
3823bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3824 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3825 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3826 [&](
Register InputRegister,
bool IsUnsigned) {
3828 InputRegister, SPIRV::OpTypeFloat);
3830 ? SPIRV::OpGroupNonUniformFMul
3831 : SPIRV::OpGroupNonUniformIMul;
3835template <
typename PickOpcodeFn>
3836bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3837 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3838 PickOpcodeFn &&PickOpcode)
const {
3839 assert(
I.getNumOperands() == 3);
3840 assert(
I.getOperand(2).isReg());
3842 Register InputRegister =
I.getOperand(2).getReg();
3846 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3849 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3855 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3856 .
addUse(
I.getOperand(2).getReg())
3861bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
3862 SPIRVTypeInst ResType,
3865 assert(
I.getNumOperands() == 3);
3866 assert(
I.getOperand(2).isReg());
3868 Register InputRegister =
I.getOperand(2).getReg();
3874 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
3885bool SPIRVInstructionSelector::selectBitreverseViaI32(
Register ResVReg,
3886 SPIRVTypeInst ResType,
3893 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
3898 : SPIRV::OpUConvert;
3902 ShiftOp = SPIRV::OpShiftRightLogicalV;
3907 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3908 TII.get(SPIRV::OpConstantComposite))
3911 for (
unsigned It = 0; It <
N; ++It)
3915 ShiftConst = CompositeReg;
3920 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
3925 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
3930 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
3935 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
3938bool SPIRVInstructionSelector::handle64BitOverflow(
3940 unsigned int Opcode,
3947 "handle64BitOverflow should only be used for integer types");
3949 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
3951 MachineIRBuilder MIRBuilder(
I);
3953 SPIRVTypeInst I64x2Type =
3955 SPIRVTypeInst Vec2ResType =
3958 std::vector<Register> PartialRegs;
3960 unsigned CurrentComponent = 0;
3961 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
3965 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3966 TII.get(SPIRV::OpVectorShuffle))
3971 .
addImm(CurrentComponent)
3972 .
addImm(CurrentComponent + 1);
3982 PartialRegs.push_back(SubVecReg);
3985 if (CurrentComponent != ComponentCount) {
3991 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
3992 SPIRV::OpVectorExtractDynamic))
4001 PartialRegs.push_back(FinalElemResReg);
4005 return selectOpWithSrcs(ResVReg, ResType,
I, PartialRegs,
4006 SPIRV::OpCompositeConstruct);
4009bool SPIRVInstructionSelector::selectBitreverse64(
Register ResVReg,
4010 SPIRVTypeInst ResType,
4014 if (ComponentCount > 2)
4015 return handle64BitOverflow(
4016 ResVReg, ResType,
I, SrcReg, SPIRV::OpBitReverse,
4018 unsigned O) {
return this->selectBitreverse64(R,
T,
I, S); });
4020 MachineIRBuilder MIRBuilder(
I);
4024 I32Type, 2 * ComponentCount, MIRBuilder,
false);
4028 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
4033 if (!selectBitreverseNative(Reverse32, VecI32Type,
I, Vec32))
4040 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4041 TII.get(SPIRV::OpVectorShuffle))
4046 for (
unsigned J = 0; J < ComponentCount; ++J) {
4053 return selectOpWithSrcs(ResVReg, ResType,
I, {SwappedVec}, SPIRV::OpBitcast);
4056bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
4057 SPIRVTypeInst ResType,
4061 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
4069bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
4070 SPIRVTypeInst ResType,
4071 MachineInstr &
I)
const {
4072 Register OpReg =
I.getOperand(1).getReg();
4081 return selectBitreverseViaI32(ResVReg, ResType,
I, OpReg);
4083 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4085 return selectBitreverse64(ResVReg, ResType,
I, OpReg);
4087 return SPIRVInstructionSelector::diagnoseUnsupported(
4088 I,
"G_BITREVERSE only support 16,32,64 bits.");
4092 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4103 unsigned AndOp = SPIRV::OpBitwiseAndS;
4104 unsigned OrOp = SPIRV::OpBitwiseOrS;
4105 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
4106 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
4108 AndOp = SPIRV::OpBitwiseAndV;
4109 OrOp = SPIRV::OpBitwiseOrV;
4110 ShlOp = SPIRV::OpShiftLeftLogicalV;
4111 ShrOp = SPIRV::OpShiftRightLogicalV;
4117 const unsigned Shift) ->
Register {
4125 Register MaskReg = CreateConst(Mask);
4126 Register ShiftReg = CreateConst(Shift);
4133 if (!selectOpWithSrcs(
T1, ResType,
I, {Input, ShiftReg}, ShrOp) ||
4134 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
4135 !selectOpWithSrcs(T3, ResType,
I, {Input, MaskReg}, AndOp) ||
4136 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
4137 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
4145 uint64_t
Mask = ~0ull;
4146 while ((Shift >>= 1) > 0) {
4153 return BuildCOPY(ResVReg, Result,
I);
4156bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
4157 SPIRVTypeInst ResType,
4158 MachineInstr &
I)
const {
4159 assert(
I.getOperand(0).isReg() &&
I.getOperand(1).isReg() &&
4160 "G_FREEZE must define and use a register");
4161 Register OpReg =
I.getOperand(1).getReg();
4165 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
4178 if (MachineInstr *Def = MRI->
getVRegDef(OpReg)) {
4179 if (
Def->getOpcode() == TargetOpcode::COPY)
4182 switch (
Def->getOpcode()) {
4183 case SPIRV::ASSIGN_TYPE:
4184 if (MachineInstr *AssignToDef =
4186 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
4187 Reg =
Def->getOperand(2).getReg();
4190 case SPIRV::OpUndef:
4191 Reg =
Def->getOperand(1).getReg();
4194 unsigned DestOpCode;
4196 DestOpCode = SPIRV::OpConstantNull;
4197 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze of a "
4198 "static undef/poison lowered to OpConstantNull\n");
4200 DestOpCode = TargetOpcode::COPY;
4202 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze "
4203 "skipped, lowered as a copy of the operand\n");
4205 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
4206 .
addDef(
I.getOperand(0).getReg())
4214bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
4215 SPIRVTypeInst ResType,
4216 MachineInstr &
I)
const {
4218 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4220 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4224 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
4229 for (
unsigned i =
I.getNumExplicitDefs();
4230 i <
I.getNumExplicitOperands() && IsConst; ++i)
4234 if (!IsConst &&
N < 2)
4235 return diagnoseUnsupported(
4236 I,
"There must be at least two constituent operands in a vector");
4241 for (
unsigned i =
I.getNumExplicitDefs();
4242 i <
I.getNumExplicitOperands() && IsNullVector; ++i) {
4243 MachineInstr *
Def =
getDef(
I.getOperand(i), MRI);
4248 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4255 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4256 TII.get(IsConst ? SPIRV::OpConstantComposite
4257 : SPIRV::OpCompositeConstruct))
4260 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
4261 MIB.
addUse(
I.getOperand(i).getReg());
4266bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
4267 SPIRVTypeInst ResType,
4268 MachineInstr &
I)
const {
4270 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4272 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4278 if (!
I.getOperand(
OpIdx).isReg())
4285 if (!IsConst &&
N < 2)
4286 return diagnoseUnsupported(
4287 I,
"There must be at least two constituent operands in a vector");
4290 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4291 TII.get(IsConst ? SPIRV::OpConstantComposite
4292 : SPIRV::OpCompositeConstruct))
4295 for (
unsigned i = 0; i <
N; ++i)
4301bool SPIRVInstructionSelector::selectConcatVectors(
Register ResVReg,
4302 SPIRVTypeInst ResType,
4303 MachineInstr &
I)
const {
4307 if (ResType->
getOpcode() != SPIRV::OpTypeVector)
4309 "Cannot select G_CONCAT_VECTORS with a non-vector result");
4311 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4312 TII.get(SPIRV::OpCompositeConstruct))
4322bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
4323 SPIRVTypeInst ResType,
4324 MachineInstr &
I)
const {
4329 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
4331 Opcode = SPIRV::OpDemoteToHelperInvocation;
4333 Opcode = SPIRV::OpKill;
4335 if (MachineInstr *NextI =
I.getNextNode()) {
4337 NextI->eraseFromParent();
4347bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
4348 SPIRVTypeInst ResType,
unsigned CmpOpc,
4349 MachineInstr &
I)
const {
4350 Register Cmp0 =
I.getOperand(2).getReg();
4351 Register Cmp1 =
I.getOperand(3).getReg();
4354 "CMP operands should have the same type");
4355 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
4365bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
4366 SPIRVTypeInst ResType,
4367 MachineInstr &
I)
const {
4368 auto Pred =
I.getOperand(1).getPredicate();
4371 Register CmpOperand =
I.getOperand(2).getReg();
4376 Register Op1 =
I.getOperand(3).getReg();
4380 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
4385 I.getOperand(3).setReg(NewOp1);
4391 return selectCmp(ResVReg, ResType, CmpOpc,
I);
4395SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
4396 SPIRVTypeInst ResType)
const {
4398 SPIRVTypeInst SpvI32Ty =
4401 auto ConstInt = ConstantInt::get(LLVMTy, Val);
4408 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4411 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
4414 .
addImm(APInt(32, Val).getZExtValue());
4416 GR.
add(ConstInt,
MI);
4423Register SPIRVInstructionSelector::buildI32ConstantInEntryBlock(
4424 uint32_t Val, MachineInstr &
I, SPIRVTypeInst ResType)
const {
4426 SPIRVTypeInst SpvI32Ty =
4428 auto *ConstInt = ConstantInt::get(LLVMTy, Val);
4433 MachineBasicBlock &EntryBB = *InsertIt->getParent();
4434 MachineInstr *
MI =
nullptr;
4438 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantNull))
4442 uint64_t ImmVal = APInt(32, Val).getZExtValue();
4443 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantI))
4449 GR.
add(ConstInt,
MI);
4454bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
4455 SPIRVTypeInst ResType,
4456 MachineInstr &
I)
const {
4458 return selectCmp(ResVReg, ResType, CmpOp,
I);
4461bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
4462 SPIRVTypeInst ResType,
4463 MachineInstr &
I)
const {
4465 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
4472 if (ResType->
getOpcode() != SPIRV::OpTypeVector &&
4473 ResType->
getOpcode() != SPIRV::OpTypeFloat)
4476 MachineIRBuilder MIRBuilder(
I);
4483 APFloat ConstVal(3.3219280948873623);
4487 APFloat::rmNearestTiesToEven, &LosesInfo);
4491 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
4492 ? SPIRV::OpVectorTimesScalar
4495 if (!selectOpWithSrcs(ArgReg, ResType,
I,
4496 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
4498 if (!selectExtInst(ResVReg, ResType,
I,
4499 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
4509Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
4510 MachineInstr &
I)
const {
4513 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4518bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
4524 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
4532 if (
Def->getOpcode() == SPIRV::OpConstantNull)
4535 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
4536 Def->getOpcode() == SPIRV::OpConstantI)
4549 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
4550 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
4552 Intrinsic::spv_const_composite)) {
4553 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
4554 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
4555 if (!IsZero(
Def->getOperand(i).getReg()))
4564Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
4565 MachineInstr &
I)
const {
4569 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4574Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
4575 MachineInstr &
I)
const {
4579 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4585 SPIRVTypeInst ResType,
4586 MachineInstr &
I)
const {
4590 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4595bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
4596 SPIRVTypeInst ResType,
4597 MachineInstr &
I)
const {
4598 Register SelectFirstArg =
I.getOperand(2).getReg();
4599 Register SelectSecondArg =
I.getOperand(3).getReg();
4608 SPIRV::OpTypeVector;
4615 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
4616 }
else if (IsPtrTy) {
4617 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
4619 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
4622 assert(IsScalarBool &&
"OpSelect with a scalar result requires a scalar "
4623 "boolean condition");
4625 Opcode = SPIRV::OpSelectSFSCond;
4626 }
else if (IsPtrTy) {
4627 Opcode = SPIRV::OpSelectSPSCond;
4629 Opcode = SPIRV::OpSelectSISCond;
4632 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4635 .
addUse(
I.getOperand(1).getReg())
4644bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
4645 SPIRVTypeInst ResType,
4647 MachineInstr &InsertAt,
4648 bool IsSigned)
const {
4650 Register ZeroReg = buildZerosVal(ResType, InsertAt);
4651 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
4652 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
4654 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4666bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
4667 SPIRVTypeInst ResType,
4668 MachineInstr &
I,
bool IsSigned,
4669 unsigned Opcode)
const {
4670 Register SrcReg =
I.getOperand(1).getReg();
4676 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
4681 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I, IsSigned);
4683 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
4686bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
4687 SPIRVTypeInst ResType, MachineInstr &
I,
4688 bool IsSigned)
const {
4689 Register SrcReg =
I.getOperand(1).getReg();
4691 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
4695 if (ResType == SrcType)
4696 return BuildCOPY(ResVReg, SrcReg,
I);
4698 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4699 return selectUnOp(ResVReg, ResType,
I, Opcode);
4702bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
4703 SPIRVTypeInst ResType,
4705 bool IsSigned)
const {
4706 MachineIRBuilder MIRBuilder(
I);
4707 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
4719 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
4722 .
addUse(
I.getOperand(1).getReg())
4723 .
addUse(
I.getOperand(2).getReg())
4728 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
4731 .
addUse(
I.getOperand(1).getReg())
4732 .
addUse(
I.getOperand(2).getReg())
4740 unsigned SelectOpcode =
4741 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
4746 .
addUse(buildOnesVal(
true, ResType,
I))
4747 .
addUse(buildZerosVal(ResType,
I))
4754 .
addUse(buildOnesVal(
false, ResType,
I))
4759bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
4762 SPIRVTypeInst IntTy,
4763 SPIRVTypeInst BoolTy)
const {
4766 bool IsVectorTy = IntTy->
getOpcode() == SPIRV::OpTypeVector;
4767 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
4769 Register One = buildOnesVal(
false, IntTy,
I);
4777 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
4786bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
4787 SPIRVTypeInst ResType,
4788 MachineInstr &
I)
const {
4789 Register IntReg =
I.getOperand(1).getReg();
4792 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
4793 if (ArgType == ResType)
4794 return BuildCOPY(ResVReg, IntReg,
I);
4796 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4797 return selectUnOp(ResVReg, ResType,
I, Opcode);
4800bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
4801 SPIRVTypeInst ResType,
4802 MachineInstr &
I)
const {
4803 unsigned Opcode =
I.getOpcode();
4804 unsigned TpOpcode = ResType->
getOpcode();
4806 if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {
4807 assert(Opcode == TargetOpcode::G_CONSTANT &&
4808 I.getOperand(1).getCImm()->isZero());
4809 MachineBasicBlock &DepMBB =
I.getMF()->front();
4812 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
4819 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
4822bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
4823 SPIRVTypeInst ResType,
4824 MachineInstr &
I)
const {
4825 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4832bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
4833 SPIRVTypeInst ResType,
4834 MachineInstr &
I)
const {
4836 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
4840 .
addUse(
I.getOperand(3).getReg())
4842 .
addUse(
I.getOperand(2).getReg());
4843 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
4849bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
4850 SPIRVTypeInst ResType,
4851 MachineInstr &
I)
const {
4852 Type *MaybeResTy =
nullptr;
4857 "Expected aggregate type for extractv instruction");
4859 SPIRV::AccessQualifier::ReadWrite,
false);
4863 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
4866 .
addUse(
I.getOperand(2).getReg());
4867 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
4873bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
4874 SPIRVTypeInst ResType,
4875 MachineInstr &
I)
const {
4876 if (
getImm(
I.getOperand(4), MRI))
4877 return selectInsertVal(ResVReg, ResType,
I);
4879 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
4882 .
addUse(
I.getOperand(2).getReg())
4883 .
addUse(
I.getOperand(3).getReg())
4884 .
addUse(
I.getOperand(4).getReg())
4889bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
4890 SPIRVTypeInst ResType,
4891 MachineInstr &
I)
const {
4892 if (
getImm(
I.getOperand(3), MRI))
4893 return selectExtractVal(ResVReg, ResType,
I);
4895 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
4898 .
addUse(
I.getOperand(2).getReg())
4899 .
addUse(
I.getOperand(3).getReg())
4904bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
4905 SPIRVTypeInst ResType,
4906 MachineInstr &
I)
const {
4907 const bool IsGEPInBounds =
I.getOperand(2).getImm();
4913 ? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
4914 : SPIRV::OpAccessChain)
4915 : (IsGEPInBounds ?
SPIRV::OpInBoundsPtrAccessChain
4916 :
SPIRV::OpPtrAccessChain);
4918 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4922 .
addUse(
I.getOperand(3).getReg());
4924 (Opcode == SPIRV::OpPtrAccessChain ||
4925 Opcode == SPIRV::OpInBoundsPtrAccessChain ||
4926 (
getImm(
I.getOperand(4), MRI) &&
foldImm(
I.getOperand(4), MRI) == 0)) &&
4927 "Cannot translate GEP to OpAccessChain. First index must be 0.");
4930 const unsigned StartingIndex =
4931 (Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)
4934 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
4935 Res.addUse(
I.getOperand(i).getReg());
4936 Res.constrainAllUses(
TII,
TRI, RBI);
4941bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
4943 unsigned Lim =
I.getNumExplicitOperands();
4944 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
4945 Register OpReg =
I.getOperand(i).getReg();
4946 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
4948 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
4949 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
4950 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
4957 MachineFunction *MF =
I.getMF();
4969 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4970 TII.get(SPIRV::OpSpecConstantOp))
4973 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
4975 GR.
add(OpDefine, MIB);
4981bool SPIRVInstructionSelector::selectDerivativeInst(
4982 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
4983 const unsigned DPdOpCode)
const {
4986 if (!errorIfInstrOutsideShader(
I))
4992 Register SrcReg =
I.getOperand(2).getReg();
4997 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5000 .
addUse(
I.getOperand(2).getReg());
5002 MachineIRBuilder MIRBuilder(
I);
5005 if (componentCount != 1)
5013 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5018 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5023 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5031bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
5032 SPIRVTypeInst ResType,
5033 MachineInstr &
I)
const {
5037 case Intrinsic::spv_load:
5038 return selectLoad(ResVReg, ResType,
I);
5039 case Intrinsic::spv_atomic_load:
5040 return selectAtomicLoad(ResVReg, ResType,
I);
5041 case Intrinsic::spv_store:
5042 return selectStore(
I);
5043 case Intrinsic::spv_atomic_store:
5044 return selectAtomicStore(
I);
5045 case Intrinsic::spv_extractv:
5046 return selectExtractVal(ResVReg, ResType,
I);
5047 case Intrinsic::spv_insertv:
5048 return selectInsertVal(ResVReg, ResType,
I);
5049 case Intrinsic::spv_extractelt:
5050 return selectExtractElt(ResVReg, ResType,
I);
5051 case Intrinsic::spv_insertelt:
5052 return selectInsertElt(ResVReg, ResType,
I);
5053 case Intrinsic::spv_gep:
5054 return selectGEP(ResVReg, ResType,
I);
5055 case Intrinsic::spv_bitcast: {
5056 Register OpReg =
I.getOperand(2).getReg();
5057 SPIRVTypeInst OpType =
5061 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
5063 case Intrinsic::spv_unref_global:
5064 case Intrinsic::spv_init_global: {
5065 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
5070 Register GVarVReg =
MI->getOperand(0).getReg();
5071 if (!selectGlobalValue(GVarVReg, *
MI, Init))
5076 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
5078 MI->eraseFromParent();
5082 case Intrinsic::spv_undef: {
5083 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5089 case Intrinsic::spv_poison:
5090 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpPoisonKHR))
5095 case Intrinsic::spv_freeze:
5096 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
5099 .
addUse(
I.getOperand(2).getReg())
5102 case Intrinsic::spv_named_boolean_spec_constant: {
5103 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
5104 : SPIRV::OpSpecConstantFalse;
5106 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
5107 .
addDef(
I.getOperand(0).getReg())
5110 unsigned SpecId =
I.getOperand(2).getImm();
5112 SPIRV::Decoration::SpecId, {SpecId});
5116 case Intrinsic::spv_const_composite: {
5118 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
5124 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
5126 std::function<bool(
Register)> HasSpecConstOperand =
5136 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
5137 J < Def->getNumExplicitOperands(); ++J) {
5138 if (
Def->getOperand(J).isReg() &&
5139 HasSpecConstOperand(
Def->getOperand(J).getReg()))
5145 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
5146 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
5147 : SPIRV::OpConstantComposite;
5148 unsigned ContinuedOpc = HasSpecConst
5149 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
5150 : SPIRV::OpConstantCompositeContinuedINTEL;
5151 MachineIRBuilder MIR(
I);
5153 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
5155 for (
auto *Instr : Instructions) {
5156 Instr->setDebugLoc(
I.getDebugLoc());
5161 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5168 case Intrinsic::spv_assign_name: {
5169 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
5170 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
5171 for (
unsigned i =
I.getNumExplicitDefs() + 2;
5172 i <
I.getNumExplicitOperands(); ++i) {
5173 MIB.
addImm(
I.getOperand(i).getImm());
5178 case Intrinsic::spv_switch: {
5179 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
5180 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5181 if (
I.getOperand(i).isReg())
5182 MIB.
addReg(
I.getOperand(i).getReg());
5183 else if (
I.getOperand(i).isCImm())
5184 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
5185 else if (
I.getOperand(i).isMBB())
5186 MIB.
addMBB(
I.getOperand(i).getMBB());
5193 case Intrinsic::spv_loop_merge: {
5194 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
5195 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5196 if (
I.getOperand(i).isMBB())
5197 MIB.
addMBB(
I.getOperand(i).getMBB());
5204 case Intrinsic::spv_loop_control_intel: {
5206 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
5207 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
5212 case Intrinsic::spv_selection_merge: {
5214 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
5215 assert(
I.getOperand(1).isMBB() &&
5216 "operand 1 to spv_selection_merge must be a basic block");
5217 MIB.
addMBB(
I.getOperand(1).getMBB());
5218 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
5222 case Intrinsic::spv_cmpxchg:
5223 return selectAtomicCmpXchg(ResVReg, ResType,
I);
5224 case Intrinsic::spv_unreachable:
5225 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
5228 case Intrinsic::spv_abort:
5229 return selectAbort(
I);
5230 case Intrinsic::spv_alloca:
5231 return selectFrameIndex(ResVReg, ResType,
I);
5232 case Intrinsic::spv_alloca_array:
5233 return selectAllocaArray(ResVReg, ResType,
I);
5234 case Intrinsic::spv_assume:
5236 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
5237 .
addUse(
I.getOperand(1).getReg())
5242 case Intrinsic::spv_expect:
5244 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
5247 .
addUse(
I.getOperand(2).getReg())
5248 .
addUse(
I.getOperand(3).getReg())
5253 case Intrinsic::arithmetic_fence:
5254 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
5255 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
5258 .
addUse(
I.getOperand(2).getReg())
5262 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
5264 case Intrinsic::spv_thread_id:
5270 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
5272 case Intrinsic::spv_thread_id_in_group:
5278 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
5280 case Intrinsic::spv_group_id:
5286 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
5288 case Intrinsic::spv_flattened_thread_id_in_group:
5295 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
5297 case Intrinsic::spv_workgroup_size:
5298 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
5300 case Intrinsic::spv_global_size:
5301 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
5303 case Intrinsic::spv_global_offset:
5304 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
5306 case Intrinsic::spv_num_workgroups:
5307 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
5309 case Intrinsic::spv_subgroup_size:
5310 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
5312 case Intrinsic::spv_num_subgroups:
5313 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
5315 case Intrinsic::spv_subgroup_id:
5316 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
5317 case Intrinsic::spv_subgroup_local_invocation_id:
5318 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
5319 ResVReg, ResType,
I);
5320 case Intrinsic::spv_subgroup_max_size:
5321 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
5323 case Intrinsic::spv_fdot:
5324 return selectFloatDot(ResVReg, ResType,
I);
5325 case Intrinsic::spv_udot:
5326 case Intrinsic::spv_sdot:
5327 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5329 return selectIntegerDot(ResVReg, ResType,
I,
5330 IID == Intrinsic::spv_sdot);
5331 return selectIntegerDotExpansion(ResVReg, ResType,
I);
5332 case Intrinsic::spv_dot4add_i8packed:
5333 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5335 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
5336 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
5337 case Intrinsic::spv_dot4add_u8packed:
5338 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5340 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
5341 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
5342 case Intrinsic::spv_all:
5343 return selectAll(ResVReg, ResType,
I);
5344 case Intrinsic::spv_any:
5345 return selectAny(ResVReg, ResType,
I);
5346 case Intrinsic::spv_cross:
5347 return selectExtInst(ResVReg, ResType,
I, CL::cross, GL::Cross);
5348 case Intrinsic::spv_distance:
5349 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
5350 case Intrinsic::spv_lerp:
5351 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
5352 case Intrinsic::spv_length:
5353 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
5354 case Intrinsic::spv_degrees:
5355 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
5356 case Intrinsic::spv_faceforward:
5357 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
5358 case Intrinsic::spv_frac:
5359 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
5360 case Intrinsic::spv_isinf:
5361 return selectOpIsInf(ResVReg, ResType,
I);
5362 case Intrinsic::spv_isnan:
5363 return selectOpIsNan(ResVReg, ResType,
I);
5364 case Intrinsic::spv_isfinite:
5365 return selectOpIsFinite(ResVReg, ResType,
I);
5366 case Intrinsic::spv_isnormal:
5367 return selectOpIsNormal(ResVReg, ResType,
I);
5368 case Intrinsic::spv_normalize:
5369 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
5370 case Intrinsic::spv_refract:
5371 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
5372 case Intrinsic::spv_reflect:
5373 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
5374 case Intrinsic::spv_rsqrt:
5375 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
5376 case Intrinsic::spv_sign:
5377 return selectSign(ResVReg, ResType,
I);
5378 case Intrinsic::spv_smoothstep:
5379 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
5380 case Intrinsic::spv_firstbituhigh:
5381 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
5382 case Intrinsic::spv_firstbitshigh:
5383 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
5384 case Intrinsic::spv_firstbitlow:
5385 return selectFirstBitLow(ResVReg, ResType,
I);
5386 case Intrinsic::spv_all_memory_barrier:
5387 return selectBarrierInst(
I, SPIRV::Scope::Device,
5388 SPIRV::MemorySemantics::UniformMemory |
5389 SPIRV::MemorySemantics::ImageMemory |
5390 SPIRV::MemorySemantics::WorkgroupMemory,
5392 case Intrinsic::spv_all_memory_barrier_with_group_sync:
5393 return selectBarrierInst(
I, SPIRV::Scope::Device,
5394 SPIRV::MemorySemantics::UniformMemory |
5395 SPIRV::MemorySemantics::ImageMemory |
5396 SPIRV::MemorySemantics::WorkgroupMemory,
5398 case Intrinsic::spv_device_memory_barrier:
5399 return selectBarrierInst(
I, SPIRV::Scope::Device,
5400 SPIRV::MemorySemantics::UniformMemory |
5401 SPIRV::MemorySemantics::ImageMemory,
5403 case Intrinsic::spv_device_memory_barrier_with_group_sync:
5404 return selectBarrierInst(
I, SPIRV::Scope::Device,
5405 SPIRV::MemorySemantics::UniformMemory |
5406 SPIRV::MemorySemantics::ImageMemory,
5408 case Intrinsic::spv_group_memory_barrier:
5409 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5410 SPIRV::MemorySemantics::WorkgroupMemory,
5412 case Intrinsic::spv_group_memory_barrier_with_group_sync:
5413 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5414 SPIRV::MemorySemantics::WorkgroupMemory,
5416 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
5417 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
5418 SPIRV::StorageClass::StorageClass ResSC =
5421 return diagnoseUnsupported(
I,
"The target storage class is not castable "
5422 "from the Generic storage class");
5423 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
5431 case Intrinsic::spv_lifetime_start:
5432 case Intrinsic::spv_lifetime_end: {
5433 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
5434 : SPIRV::OpLifetimeStop;
5435 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
5436 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
5445 case Intrinsic::spv_saturate:
5446 return selectSaturate(ResVReg, ResType,
I);
5447 case Intrinsic::spv_nclamp:
5448 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
5449 case Intrinsic::spv_uclamp:
5450 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
5451 case Intrinsic::spv_sclamp:
5452 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
5453 case Intrinsic::spv_subgroup_prefix_bit_count:
5454 return selectWavePrefixBitCount(ResVReg, ResType,
I);
5455 case Intrinsic::spv_wave_active_countbits:
5456 return selectWaveActiveCountBits(ResVReg, ResType,
I);
5457 case Intrinsic::spv_wave_all_equal:
5458 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
5459 case Intrinsic::spv_wave_all:
5460 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
5461 case Intrinsic::spv_wave_any:
5462 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
5463 case Intrinsic::spv_subgroup_ballot:
5464 return selectWaveOpInst(ResVReg, ResType,
I,
5465 SPIRV::OpGroupNonUniformBallot);
5466 case Intrinsic::spv_wave_is_first_lane:
5467 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
5468 case Intrinsic::spv_wave_reduce_or:
5469 return selectWaveReduceOp(ResVReg, ResType,
I,
5470 SPIRV::OpGroupNonUniformBitwiseOr);
5471 case Intrinsic::spv_wave_reduce_xor:
5472 return selectWaveReduceOp(ResVReg, ResType,
I,
5473 SPIRV::OpGroupNonUniformBitwiseXor);
5474 case Intrinsic::spv_wave_reduce_and:
5475 return selectWaveReduceOp(ResVReg, ResType,
I,
5476 SPIRV::OpGroupNonUniformBitwiseAnd);
5477 case Intrinsic::spv_wave_reduce_umax:
5478 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
5479 case Intrinsic::spv_wave_reduce_max:
5480 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
5481 case Intrinsic::spv_wave_reduce_umin:
5482 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
5483 case Intrinsic::spv_wave_reduce_min:
5484 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
5485 case Intrinsic::spv_wave_reduce_sum:
5486 return selectWaveReduceSum(ResVReg, ResType,
I);
5487 case Intrinsic::spv_wave_product:
5488 return selectWaveReduceProduct(ResVReg, ResType,
I);
5489 case Intrinsic::spv_wave_readlane:
5490 return selectWaveOpInst(ResVReg, ResType,
I,
5491 SPIRV::OpGroupNonUniformShuffle);
5492 case Intrinsic::spv_wave_prefix_sum:
5493 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
5494 case Intrinsic::spv_wave_prefix_product:
5495 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
5496 case Intrinsic::spv_quad_read_across_x: {
5497 return selectQuadSwap(ResVReg, ResType,
I, 0);
5499 case Intrinsic::spv_quad_read_across_y: {
5500 return selectQuadSwap(ResVReg, ResType,
I, 1);
5502 case Intrinsic::spv_quad_read_across_diagonal: {
5503 return selectQuadSwap(ResVReg, ResType,
I, 2);
5505 case Intrinsic::spv_step:
5506 return selectExtInst(ResVReg, ResType,
I, CL::step, GL::Step);
5507 case Intrinsic::spv_radians:
5508 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
5512 case Intrinsic::instrprof_increment:
5513 case Intrinsic::instrprof_increment_step:
5514 case Intrinsic::instrprof_value_profile:
5517 case Intrinsic::spv_value_md:
5519 case Intrinsic::spv_resource_handlefrombinding: {
5520 return selectHandleFromBinding(ResVReg, ResType,
I);
5522 case Intrinsic::spv_resource_counterhandlefrombinding:
5523 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
5524 case Intrinsic::spv_resource_updatecounter:
5525 return selectUpdateCounter(ResVReg, ResType,
I);
5526 case Intrinsic::spv_resource_store_typedbuffer: {
5527 return selectImageWriteIntrinsic(
I);
5529 case Intrinsic::spv_resource_load_typedbuffer: {
5530 return selectReadImageIntrinsic(ResVReg, ResType,
I);
5532 case Intrinsic::spv_resource_load_level: {
5533 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
5535 case Intrinsic::spv_resource_getdimensions_x:
5536 case Intrinsic::spv_resource_getdimensions_xy:
5537 case Intrinsic::spv_resource_getdimensions_xyz: {
5538 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
5540 case Intrinsic::spv_resource_getdimensions_levels_x:
5541 case Intrinsic::spv_resource_getdimensions_levels_xy:
5542 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
5543 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
5545 case Intrinsic::spv_resource_getdimensions_ms_xy:
5546 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
5547 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
5549 case Intrinsic::spv_resource_calculate_lod:
5550 case Intrinsic::spv_resource_calculate_lod_unclamped:
5551 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
5552 case Intrinsic::spv_resource_sample:
5553 case Intrinsic::spv_resource_sample_clamp:
5554 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
5555 case Intrinsic::spv_resource_samplebias:
5556 case Intrinsic::spv_resource_samplebias_clamp:
5557 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
5558 case Intrinsic::spv_resource_samplegrad:
5559 case Intrinsic::spv_resource_samplegrad_clamp:
5560 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
5561 case Intrinsic::spv_resource_samplelevel:
5562 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
5563 case Intrinsic::spv_resource_samplecmp:
5564 case Intrinsic::spv_resource_samplecmp_clamp:
5565 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
5566 case Intrinsic::spv_resource_samplecmplevelzero:
5567 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
5568 case Intrinsic::spv_resource_gather:
5569 case Intrinsic::spv_resource_gather_cmp:
5570 return selectGatherIntrinsic(ResVReg, ResType,
I);
5571 case Intrinsic::spv_resource_getbasepointer:
5572 case Intrinsic::spv_resource_getpointer: {
5573 return selectResourceGetPointer(ResVReg, ResType,
I);
5575 case Intrinsic::spv_pushconstant_getpointer: {
5576 return selectPushConstantGetPointer(ResVReg, ResType,
I);
5578 case Intrinsic::spv_discard: {
5579 return selectDiscard(ResVReg, ResType,
I);
5581 case Intrinsic::spv_resource_nonuniformindex: {
5582 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
5584 case Intrinsic::spv_unpackhalf2x16: {
5585 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
5587 case Intrinsic::spv_packhalf2x16: {
5588 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
5590 case Intrinsic::spv_ddx:
5591 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
5592 case Intrinsic::spv_ddy:
5593 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
5594 case Intrinsic::spv_ddx_coarse:
5595 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
5596 case Intrinsic::spv_ddy_coarse:
5597 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
5598 case Intrinsic::spv_ddx_fine:
5599 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
5600 case Intrinsic::spv_ddy_fine:
5601 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
5602 case Intrinsic::spv_fwidth:
5603 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
5604 case Intrinsic::spv_masked_gather:
5605 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5606 return selectMaskedGather(ResVReg, ResType,
I);
5607 return diagnoseUnsupported(
5608 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
5609 case Intrinsic::spv_masked_scatter:
5610 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5611 return selectMaskedScatter(
I);
5612 return diagnoseUnsupported(
5613 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
5614 case Intrinsic::returnaddress:
5615 case Intrinsic::frameaddress: {
5617 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5624 return diagnoseUnsupported(
I,
"intrinsic selection not implemented.");
5629bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
5630 SPIRVTypeInst ResType,
5631 MachineInstr &
I)
const {
5634 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
5641bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
5642 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5644 assert(Intr.getIntrinsicID() ==
5645 Intrinsic::spv_resource_counterhandlefrombinding);
5648 Register MainHandleReg = Intr.getOperand(2).getReg();
5650 assert(MainHandleDef->getIntrinsicID() ==
5651 Intrinsic::spv_resource_handlefrombinding);
5655 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
5656 Register IndexReg = MainHandleDef->getOperand(5).getReg();
5657 std::string CounterName =
5662 MachineIRBuilder MIRBuilder(
I);
5664 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
5666 ArraySize, IndexReg, CounterName, MIRBuilder);
5668 return BuildCOPY(ResVReg, CounterVarReg,
I);
5671bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
5672 SPIRVTypeInst ResType,
5673 MachineInstr &
I)
const {
5675 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
5677 Register CounterHandleReg = Intr.getOperand(2).getReg();
5678 Register IncrReg = Intr.getOperand(3).getReg();
5685 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
5686 assert(CounterVarPointeeType &&
5687 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
5688 "Counter variable must be a struct");
5690 SPIRV::StorageClass::StorageBuffer &&
5691 "Counter variable must be in the storage buffer storage class");
5693 "Counter variable must have exactly 1 member in the struct");
5694 const SPIRVTypeInst MemberType =
5697 "Counter variable struct must have a single i32 member");
5701 MachineIRBuilder MIRBuilder(
I);
5703 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
5706 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
5712 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5715 .
addUse(CounterHandleReg)
5722 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
5725 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
5728 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
5737 return BuildCOPY(ResVReg, AtomicRes,
I);
5745 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
5753bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
5754 SPIRVTypeInst ResType,
5755 MachineInstr &
I)
const {
5763 Register ImageReg =
I.getOperand(2).getReg();
5771 Register IdxReg =
I.getOperand(3).getReg();
5773 MachineInstr &Pos =
I;
5775 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
5779bool SPIRVInstructionSelector::generateSampleImage(
5782 DebugLoc Loc, MachineInstr &Pos)
const {
5793 if (!loadHandleBeforePosition(NewSamplerReg,
5799 MachineIRBuilder MIRBuilder(Pos);
5812 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
5813 ImOps.Lod.has_value();
5814 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
5815 : SPIRV::OpImageSampleImplicitLod;
5817 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
5818 : SPIRV::OpImageSampleDrefImplicitLod;
5827 MIB.
addUse(*ImOps.Compare);
5829 uint32_t ImageOperands = 0;
5831 ImageOperands |= SPIRV::ImageOperand::Bias;
5833 ImageOperands |= SPIRV::ImageOperand::Lod;
5834 if (ImOps.GradX && ImOps.GradY)
5835 ImageOperands |= SPIRV::ImageOperand::Grad;
5836 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
5838 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
5841 "Non-constant offsets are not supported in sample instructions.");
5846 ImageOperands |= SPIRV::ImageOperand::MinLod;
5848 if (ImageOperands != 0) {
5849 MIB.
addImm(ImageOperands);
5850 if (ImageOperands & SPIRV::ImageOperand::Bias)
5852 if (ImageOperands & SPIRV::ImageOperand::Lod)
5854 if (ImageOperands & SPIRV::ImageOperand::Grad) {
5855 MIB.
addUse(*ImOps.GradX);
5856 MIB.
addUse(*ImOps.GradY);
5859 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
5860 MIB.
addUse(*ImOps.Offset);
5861 if (ImageOperands & SPIRV::ImageOperand::MinLod)
5862 MIB.
addUse(*ImOps.MinLod);
5869bool SPIRVInstructionSelector::selectImageQuerySize(
5871 std::optional<Register> LodReg)
const {
5873 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
5876 "ImageReg is not an image type.");
5878 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
5880 unsigned NumComponents = 0;
5882 case SPIRV::Dim::DIM_1D:
5883 case SPIRV::Dim::DIM_Buffer:
5884 NumComponents =
IsArray ? 2 : 1;
5886 case SPIRV::Dim::DIM_2D:
5887 case SPIRV::Dim::DIM_Cube:
5888 case SPIRV::Dim::DIM_Rect:
5889 NumComponents =
IsArray ? 3 : 2;
5891 case SPIRV::Dim::DIM_3D:
5895 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
5900 SPIRVTypeInst ResType =
5905 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5915bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
5916 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5917 Register ImageReg =
I.getOperand(2).getReg();
5924 return selectImageQuerySize(NewImageReg, ResVReg,
I);
5927bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
5928 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5929 Register ImageReg =
I.getOperand(2).getReg();
5938 Register LodReg =
I.getOperand(3).getReg();
5941 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
5943 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
5950 TII.get(SPIRV::OpImageQueryLevels))
5957 TII.get(SPIRV::OpCompositeConstruct))
5967bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
5968 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5969 Register ImageReg =
I.getOperand(2).getReg();
5980 "OpImageQuerySamples requires a multisampled image");
5982 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
5990 TII.get(SPIRV::OpImageQuerySamples))
5997 TII.get(SPIRV::OpCompositeConstruct))
6007bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
6008 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6009 Register ImageReg =
I.getOperand(2).getReg();
6010 Register SamplerReg =
I.getOperand(3).getReg();
6011 Register CoordinateReg =
I.getOperand(4).getReg();
6027 if (!loadHandleBeforePosition(
6032 MachineIRBuilder MIRBuilder(
I);
6038 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6048 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
6055 unsigned ExtractedIndex =
6057 Intrinsic::spv_resource_calculate_lod_unclamped
6061 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6062 TII.get(SPIRV::OpCompositeExtract))
6072bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
6073 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6074 Register ImageReg =
I.getOperand(2).getReg();
6075 Register SamplerReg =
I.getOperand(3).getReg();
6076 Register CoordinateReg =
I.getOperand(4).getReg();
6077 ImageOperands ImOps;
6078 if (
I.getNumOperands() > 5)
6079 ImOps.Offset =
I.getOperand(5).getReg();
6080 if (
I.getNumOperands() > 6)
6081 ImOps.MinLod =
I.getOperand(6).getReg();
6082 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6083 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6086bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
6087 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6088 Register ImageReg =
I.getOperand(2).getReg();
6089 Register SamplerReg =
I.getOperand(3).getReg();
6090 Register CoordinateReg =
I.getOperand(4).getReg();
6091 ImageOperands ImOps;
6092 ImOps.Bias =
I.getOperand(5).getReg();
6093 if (
I.getNumOperands() > 6)
6094 ImOps.Offset =
I.getOperand(6).getReg();
6095 if (
I.getNumOperands() > 7)
6096 ImOps.MinLod =
I.getOperand(7).getReg();
6097 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6098 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6101bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
6102 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6103 Register ImageReg =
I.getOperand(2).getReg();
6104 Register SamplerReg =
I.getOperand(3).getReg();
6105 Register CoordinateReg =
I.getOperand(4).getReg();
6106 ImageOperands ImOps;
6107 ImOps.GradX =
I.getOperand(5).getReg();
6108 ImOps.GradY =
I.getOperand(6).getReg();
6109 if (
I.getNumOperands() > 7)
6110 ImOps.Offset =
I.getOperand(7).getReg();
6111 if (
I.getNumOperands() > 8)
6112 ImOps.MinLod =
I.getOperand(8).getReg();
6113 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6114 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6117bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
6118 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6119 Register ImageReg =
I.getOperand(2).getReg();
6120 Register SamplerReg =
I.getOperand(3).getReg();
6121 Register CoordinateReg =
I.getOperand(4).getReg();
6122 ImageOperands ImOps;
6123 ImOps.Lod =
I.getOperand(5).getReg();
6124 if (
I.getNumOperands() > 6)
6125 ImOps.Offset =
I.getOperand(6).getReg();
6126 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6127 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6130bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
6131 SPIRVTypeInst ResType,
6132 MachineInstr &
I)
const {
6133 Register ImageReg =
I.getOperand(2).getReg();
6134 Register SamplerReg =
I.getOperand(3).getReg();
6135 Register CoordinateReg =
I.getOperand(4).getReg();
6136 ImageOperands ImOps;
6137 ImOps.Compare =
I.getOperand(5).getReg();
6138 if (
I.getNumOperands() > 6)
6139 ImOps.Offset =
I.getOperand(6).getReg();
6140 if (
I.getNumOperands() > 7)
6141 ImOps.MinLod =
I.getOperand(7).getReg();
6142 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6143 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6146bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
6147 SPIRVTypeInst ResType,
6148 MachineInstr &
I)
const {
6149 Register ImageReg =
I.getOperand(2).getReg();
6150 Register CoordinateReg =
I.getOperand(3).getReg();
6151 Register LodReg =
I.getOperand(4).getReg();
6153 ImageOperands ImOps;
6155 if (
I.getNumOperands() > 5)
6156 ImOps.Offset =
I.getOperand(5).getReg();
6168 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
6169 I.getDebugLoc(),
I, &ImOps);
6172bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
6173 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6174 Register ImageReg =
I.getOperand(2).getReg();
6175 Register SamplerReg =
I.getOperand(3).getReg();
6176 Register CoordinateReg =
I.getOperand(4).getReg();
6177 ImageOperands ImOps;
6178 ImOps.Compare =
I.getOperand(5).getReg();
6179 if (
I.getNumOperands() > 6)
6180 ImOps.Offset =
I.getOperand(6).getReg();
6183 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6184 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6187bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
6188 SPIRVTypeInst ResType,
6189 MachineInstr &
I)
const {
6190 Register ImageReg =
I.getOperand(2).getReg();
6191 Register SamplerReg =
I.getOperand(3).getReg();
6192 Register CoordinateReg =
I.getOperand(4).getReg();
6195 "ImageReg is not an image type.");
6200 ComponentOrCompareReg =
I.getOperand(5).getReg();
6201 OffsetReg =
I.getOperand(6).getReg();
6204 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
6208 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6209 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
6210 Dim != SPIRV::Dim::DIM_Rect) {
6212 "Gather operations are only supported for 2D, Cube, and Rect images.");
6219 if (!loadHandleBeforePosition(
6224 MachineIRBuilder MIRBuilder(
I);
6225 SPIRVTypeInst SampledImageType =
6230 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6238 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
6240 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
6242 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6247 .
addUse(ComponentOrCompareReg);
6249 uint32_t ImageOperands = 0;
6250 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
6251 if (Dim == SPIRV::Dim::DIM_Cube) {
6253 "Gather operations with offset are not supported for Cube images.");
6257 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6259 ImageOperands |= SPIRV::ImageOperand::Offset;
6263 if (ImageOperands != 0) {
6264 MIB.
addImm(ImageOperands);
6266 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6274bool SPIRVInstructionSelector::generateImageReadOrFetch(
6277 const ImageOperands *ImOps)
const {
6280 "ImageReg is not an image type.");
6282 bool IsSignedInteger =
6287 bool IsFetch = (SampledOp.getImm() == 1);
6289 auto AddOperands = [&](MachineInstrBuilder &MIB) {
6290 uint32_t ImageOperandsMask = 0;
6291 if (IsSignedInteger)
6292 ImageOperandsMask |= 0x1000;
6294 if (IsFetch && ImOps) {
6296 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
6297 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
6299 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
6301 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
6305 if (ImageOperandsMask != 0) {
6306 MIB.
addImm(ImageOperandsMask);
6307 if (IsFetch && ImOps) {
6310 if (ImOps->Offset &&
6311 (ImageOperandsMask &
6312 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
6313 MIB.
addUse(*ImOps->Offset);
6322 SPIRVTypeInst SampledType =
6325 SPIRVTypeInst ReadType =
6326 widenTypeToVec4(IsPacked ? SampledType : ResType, Pos);
6327 bool ReadTypeMatchesResult = ReadType == ResType;
6329 Register ReadReg = ReadTypeMatchesResult
6335 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
6341 BMI.constrainAllUses(
TII,
TRI, RBI);
6343 if (ReadTypeMatchesResult)
6356 if (ResultSize == 1) {
6365 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
6368bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
6369 SPIRVTypeInst ResType,
6370 MachineInstr &
I)
const {
6371 Register ResourcePtr =
I.getOperand(2).getReg();
6373 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
6382 MachineIRBuilder MIRBuilder(
I);
6387 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6393 if (
I.getNumExplicitOperands() > 3) {
6394 Register IndexReg =
I.getOperand(3).getReg();
6401bool SPIRVInstructionSelector::selectPushConstantGetPointer(
6402 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6407bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
6408 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6409 Register ObjReg =
I.getOperand(2).getReg();
6410 if (!BuildCOPY(ResVReg, ObjReg,
I))
6420 decorateUsesAsNonUniform(ResVReg);
6424void SPIRVInstructionSelector::decorateUsesAsNonUniform(
6427 while (WorkList.
size() > 0) {
6431 bool IsDecorated =
false;
6433 if (
Use.getOpcode() == SPIRV::OpDecorate &&
6434 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
6440 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
6442 if (ResultReg == CurrentReg)
6450 SPIRV::Decoration::NonUniformEXT, {});
6455bool SPIRVInstructionSelector::extractSubvector(
6457 MachineInstr &InsertionPoint)
const {
6459 [[maybe_unused]] uint64_t InputSize =
6462 assert(InputSize > 1 &&
"The input must be a vector.");
6463 assert(ResultSize > 1 &&
"The result must be a vector.");
6464 assert(ResultSize < InputSize &&
6465 "Cannot extract more element than there are in the input.");
6469 for (uint64_t
I = 0;
I < ResultSize;
I++) {
6472 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6481 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
6483 TII.get(SPIRV::OpCompositeConstruct))
6487 for (
Register ComponentReg : ComponentRegisters)
6488 MIB.
addUse(ComponentReg);
6493bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
6494 MachineInstr &
I)
const {
6501 Register ImageReg =
I.getOperand(1).getReg();
6509 Register CoordinateReg =
I.getOperand(2).getReg();
6510 Register DataReg =
I.getOperand(3).getReg();
6513 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
6521Register SPIRVInstructionSelector::buildPointerToResource(
6522 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
6523 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
6524 StringRef Name, MachineIRBuilder MIRBuilder)
const {
6526 if (ArraySize == 1) {
6527 SPIRVTypeInst PtrType =
6530 "SpirvResType did not have an explicit layout.");
6535 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
6536 SPIRVTypeInst VarPointerType =
6539 VarPointerType, Set,
Binding, Name, MIRBuilder);
6541 SPIRVTypeInst ResPointerType =
6554bool SPIRVInstructionSelector::selectFirstBitSet16(
6555 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
6556 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
6558 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
6562 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
6565bool SPIRVInstructionSelector::selectFirstBitSet32(
6567 unsigned BitSetOpcode)
const {
6568 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6571 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6578bool SPIRVInstructionSelector::selectFirstBitSet64(
6580 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
6593 if (ComponentCount > 2) {
6594 auto Func = [
this, SwapPrimarySide](
Register ResVReg, SPIRVTypeInst ResType,
6596 unsigned Opcode) ->
bool {
6597 return this->selectFirstBitSet64(ResVReg, ResType,
I, SrcReg, Opcode,
6601 return handle64BitOverflow(ResVReg, ResType,
I, SrcReg, BitSetOpcode, Func);
6605 MachineIRBuilder MIRBuilder(
I);
6607 BaseType, 2 * ComponentCount, MIRBuilder,
false);
6611 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
6617 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
6624 bool IsScalarRes = ResType->
getOpcode() != SPIRV::OpTypeVector;
6627 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
6628 SPIRV::OpVectorExtractDynamic))
6630 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
6631 SPIRV::OpVectorExtractDynamic))
6635 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6636 TII.get(SPIRV::OpVectorShuffle))
6644 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
6650 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6651 TII.get(SPIRV::OpVectorShuffle))
6659 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
6679 SelectOp = SPIRV::OpSelectSISCond;
6680 AddOp = SPIRV::OpIAddS;
6688 SelectOp = SPIRV::OpSelectVIVCond;
6689 AddOp = SPIRV::OpIAddV;
6695 Register RegSecondaryOffset = Reg0;
6699 if (SwapPrimarySide) {
6700 PrimaryReg = LowReg;
6701 SecondaryReg = HighReg;
6702 RegPrimaryOffset = Reg0;
6703 RegSecondaryOffset = Reg32;
6708 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
6709 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
6714 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
6715 SPIRV::OpINotEqual))
6722 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
6723 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
6728 if (SwapPrimarySide) {
6730 if (!selectOpWithSrcs(RegAdd, ResType,
I,
6731 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
6742 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
6743 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
6748 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
6749 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
6752 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
6756bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
6757 SPIRVTypeInst ResType,
6759 bool IsSigned)
const {
6761 Register OpReg =
I.getOperand(2).getReg();
6764 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
6765 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
6769 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6771 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6773 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6776 return diagnoseUnsupported(
6778 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
6782bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
6783 SPIRVTypeInst ResType,
6784 MachineInstr &
I)
const {
6786 Register OpReg =
I.getOperand(2).getReg();
6791 unsigned ExtendOpcode = SPIRV::OpUConvert;
6792 unsigned BitSetOpcode = GL::FindILsb;
6796 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6798 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6800 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6803 return diagnoseUnsupported(
I,
6804 "spv_firstbitlow only supports 16,32,64 bits.");
6808bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
6809 SPIRVTypeInst ResType,
6810 MachineInstr &
I)
const {
6814 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
6817 .
addUse(
I.getOperand(2).getReg())
6820 unsigned Alignment =
I.getOperand(3).getImm();
6834 while (!Worklist.
empty()) {
6836 switch (
T->getOpcode()) {
6837 case SPIRV::OpTypeInt:
6838 case SPIRV::OpTypeFloat:
6839 case SPIRV::OpTypePointer:
6841 case SPIRV::OpTypeVector:
6842 case SPIRV::OpTypeMatrix:
6843 case SPIRV::OpTypeArray: {
6844 Register OperandReg =
T->getOperand(1).getReg();
6848 case SPIRV::OpTypeStruct:
6849 for (
unsigned Idx = 1,
E =
T->getNumOperands(); Idx <
E; ++Idx) {
6850 Register OperandReg =
T->getOperand(Idx).getReg();
6862bool SPIRVInstructionSelector::selectAbort(MachineInstr &
I)
const {
6863 assert(
I.getNumExplicitOperands() == 2);
6865 Register MsgReg =
I.getOperand(1).getReg();
6867 assert(MsgType &&
"Message argument of llvm.spv.abort has no SPIR-V type");
6870 return diagnoseUnsupported(
6872 "llvm.spv.abort message type must be a concrete SPIR-V type (numerical "
6873 "scalar, pointer, vector, matrix, or aggregate of such types)");
6876 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
6883bool SPIRVInstructionSelector::selectTrap(MachineInstr &
I)
const {
6892 uint32_t MsgVal = ~0
u;
6893 if (
I.getOpcode() == TargetOpcode::G_UBSANTRAP)
6894 MsgVal =
static_cast<uint32_t
>(
I.getOperand(0).
getImm());
6897 Register MsgReg = buildI32ConstantInEntryBlock(MsgVal,
I, MsgType);
6900 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
6907bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
6908 SPIRVTypeInst ResType,
6909 MachineInstr &
I)
const {
6913 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
6916 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
6919 unsigned Alignment =
I.getOperand(2).getImm();
6926bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
6931 const MachineInstr *PrevI =
I.getPrevNode();
6933 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
6937 .
addMBB(
I.getOperand(0).getMBB())
6942 .
addMBB(
I.getOperand(0).getMBB())
6947bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
6958 const MachineInstr *NextI =
I.getNextNode();
6960 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
6966 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
6968 .
addUse(
I.getOperand(0).getReg())
6969 .
addMBB(
I.getOperand(1).getMBB())
6975bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
6976 MachineInstr &
I)
const {
6978 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
6980 const unsigned NumOps =
I.getNumOperands();
6981 for (
unsigned i = 1; i <
NumOps; i += 2) {
6982 MIB.
addUse(
I.getOperand(i + 0).getReg());
6983 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
6989bool SPIRVInstructionSelector::selectGlobalValue(
6990 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
6992 MachineIRBuilder MIRBuilder(
I);
6993 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
6996 std::string GlobalIdent;
6998 unsigned &
ID = UnnamedGlobalIDs[GV];
7000 ID = UnnamedGlobalIDs.
size();
7001 GlobalIdent =
"__unnamed_" + Twine(
ID).str();
7027 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
7034 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7039 MachineInstrBuilder MIB1 =
7040 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7043 MachineInstrBuilder MIB2 =
7045 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
7049 GR.
add(ConstVal, MIB2);
7057 MachineInstrBuilder MIB3 =
7058 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7061 GR.
add(ConstVal, MIB3);
7067 assert(NewReg != ResVReg);
7068 return BuildCOPY(ResVReg, NewReg,
I);
7078 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
7081 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
7087 SPIRVTypeInst ResType =
7091 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
7096 if (
GlobalVar->isExternallyInitialized() &&
7097 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
7098 constexpr unsigned ReadWriteINTEL = 3u;
7101 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
7107bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
7108 SPIRVTypeInst ResType,
7109 MachineInstr &
I)
const {
7111 return selectExtInst(ResVReg, ResType,
I, CL::log10);
7119 MachineIRBuilder MIRBuilder(
I);
7124 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7127 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
7129 .
add(
I.getOperand(1))
7134 ResType->
getOpcode() == SPIRV::OpTypeFloat);
7144 APFloat::rmNearestTiesToEven, &LosesInfo);
7148 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
7149 ? SPIRV::OpVectorTimesScalar
7160bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
7161 SPIRVTypeInst ResType,
7162 MachineInstr &
I)
const {
7165 return selectExtInst(ResVReg, ResType,
I, CL::pown);
7171 Register ExpReg =
I.getOperand(2).getReg();
7173 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
7174 SPIRV::OpConvertSToF))
7176 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
7183bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
7184 SPIRVTypeInst ResType,
7185 MachineInstr &
I)
const {
7201 MachineIRBuilder MIRBuilder(
I);
7202 SPIRVTypeInst FloatType =
7206 FloatType, MIRBuilder, SPIRV::StorageClass::Function);
7219 MachineBasicBlock &EntryBB =
I.getMF()->
front();
7221 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
7224 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7230 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7233 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
7236 .
add(
I.getOperand(
I.getNumExplicitDefs()))
7240 Register IntegralPartReg =
I.getOperand(1).getReg();
7243 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7253 assert(
false &&
"GLSL::Modf is deprecated.");
7264bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
7265 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7266 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7267 MachineIRBuilder MIRBuilder(
I);
7268 const SPIRVTypeInst Vec3Ty =
7271 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
7283 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7287 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7293 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7300 assert(
I.getOperand(2).isReg());
7301 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
7305 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
7316bool SPIRVInstructionSelector::loadBuiltinInputID(
7317 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7318 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7319 MachineIRBuilder MIRBuilder(
I);
7321 ResType, MIRBuilder, SPIRV::StorageClass::Input);
7336 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7340 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7349SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
7350 MachineInstr &
I)
const {
7351 MachineIRBuilder MIRBuilder(
I);
7352 if (
Type->getOpcode() != SPIRV::OpTypeVector)
7362bool SPIRVInstructionSelector::loadHandleBeforePosition(
7363 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
7364 MachineInstr &Pos)
const {
7367 Intrinsic::spv_resource_handlefrombinding);
7375 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
7376 MachineIRBuilder MIRBuilder(HandleDef);
7377 SPIRVTypeInst VarType = ResType;
7378 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
7380 if (IsStructuredBuffer) {
7385 if (ResType->
getOpcode() == SPIRV::OpTypeImage && ArraySize == 0)
7387 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
7390 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
7391 ArraySize, IndexReg, Name, MIRBuilder);
7395 uint32_t LoadOpcode =
7396 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
7406bool SPIRVInstructionSelector::errorIfInstrOutsideShader(
7407 MachineInstr &
I)
const {
7409 return diagnoseUnsupported(
7410 I,
"this instruction is only supported in shaders.");
7415InstructionSelector *
7419 return new SPIRVInstructionSelector(TM, Subtarget, RBI);
MachineInstrBuilder & UseMI
#define GET_GLOBALISEL_PREDICATES_INIT
#define GET_GLOBALISEL_TEMPORARIES_INIT
assert(UImm &&(UImm !=~static_cast< T >(0)) &&"Invalid immediate!")
This file declares a class to represent arbitrary precision floating point values and provide a varie...
static bool selectUnmergeValues(MachineInstrBuilder &MIB, const ARMBaseInstrInfo &TII, MachineRegisterInfo &MRI, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI)
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
static 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.
static TableGen::Emitter::Opt Y("gen-skeleton-entry", EmitSkeleton, "Generate example skeleton entry")
static ManagedStatic< cl::opt< FnT >, OptCreatorT > CallbackFunction
static const fltSemantics & IEEEsingle()
static const fltSemantics & BFloat()
static const fltSemantics & IEEEdouble()
static const fltSemantics & IEEEhalf()
const fltSemantics & getSemantics() const
static APFloat getOne(const fltSemantics &Sem, bool Negative=false)
Factory for Positive and Negative One.
static APFloat getZero(const fltSemantics &Sem, bool Negative=false)
Factory for Positive and Negative Zero.
static APInt getAllOnes(unsigned numBits)
Return an APInt of a specified width with all bits set.
Represent a constant reference to an array (0 or more elements consecutively in memory),...
BlockFrequencyInfo pass uses BlockFrequencyInfoImpl implementation to estimate IR basic block frequen...
Predicate
This enumeration lists the possible predicates for CmpInst subclasses.
@ FCMP_OEQ
0 0 0 1 True if ordered and equal
@ ICMP_SLT
signed less than
@ ICMP_SLE
signed less or equal
@ FCMP_OLT
0 1 0 0 True if ordered and less than
@ FCMP_ULE
1 1 0 1 True if unordered, less than, or equal
@ FCMP_OGT
0 0 1 0 True if ordered and greater than
@ FCMP_OGE
0 0 1 1 True if ordered and greater than or equal
@ ICMP_UGE
unsigned greater or equal
@ ICMP_UGT
unsigned greater than
@ ICMP_SGT
signed greater than
@ FCMP_ULT
1 1 0 0 True if unordered or less than
@ FCMP_ONE
0 1 1 0 True if ordered and operands are unequal
@ FCMP_UEQ
1 0 0 1 True if unordered or equal
@ ICMP_ULT
unsigned less than
@ FCMP_UGT
1 0 1 0 True if unordered or greater than
@ FCMP_OLE
0 1 0 1 True if ordered and less than or equal
@ FCMP_ORD
0 1 1 1 True if ordered (no nans)
@ ICMP_SGE
signed greater or equal
@ FCMP_UNE
1 1 1 0 True if unordered or not equal
@ ICMP_ULE
unsigned less or equal
@ FCMP_UGE
1 0 1 1 True if unordered, greater than, or equal
@ FCMP_UNO
1 0 0 0 True if unordered: isnan(X) | isnan(Y)
static LLVM_ABI Constant * getNullValue(Type *Ty)
Constructor to create a '0' constant of arbitrary type.
LLVMContext & getContext() const
getContext - Return a reference to the LLVMContext associated with this function.
Represents a call to an intrinsic.
Intrinsic::ID getIntrinsicID() const
unsigned getAddressSpace() const
Module * getParent()
Get the module that this global value is contained inside of...
@ InternalLinkage
Rename collisions when linking (static functions).
static LLVM_ABI IntegerType * get(LLVMContext &C, unsigned NumBits)
This static method is the primary way of constructing an IntegerType.
static constexpr LLT scalar(unsigned SizeInBits)
Get a low-level scalar or aggregate "bag of bits".
constexpr bool isValid() const
constexpr uint16_t getNumElements() const
Returns the number of elements in a vector LLT.
constexpr bool isVector() const
static constexpr LLT pointer(unsigned AddressSpace, unsigned SizeInBits)
Get a low-level pointer in the given address space.
static constexpr LLT fixed_vector(unsigned NumElements, unsigned ScalarSizeInBits)
Get a low-level fixed-width vector of some number of elements and element width.
int getNumber() const
MachineBasicBlocks are uniquely numbered at the function level, unless they're not in a MachineFuncti...
LLVM_ABI iterator getFirstNonPHI()
Returns a pointer to the first instruction in this block that is not a PHINode instruction.
const MachineFunction * getParent() const
Return the MachineFunction containing this basic block.
MachineInstrBundleIterator< MachineInstr > iterator
MachineRegisterInfo & getRegInfo()
getRegInfo - Return information about the registers currently in use.
Function & getFunction()
Return the LLVM function that this machine code represents.
Helper class to build MachineInstr.
MachineBasicBlock::iterator getInsertPt()
Current insertion point for new instructions.
MachineInstrBuilder buildInstr(unsigned Opcode)
Build and insert <empty> = Opcode <empty>.
MachineFunction & getMF()
Getter for the function we currently build.
MachineRegisterInfo * getMRI()
Getter for MRI.
void constrainAllUses(const TargetInstrInfo &TII, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI) const
const MachineInstrBuilder & addUse(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a virtual register use operand.
const MachineInstrBuilder & addReg(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a new virtual register operand.
const MachineInstrBuilder & addImm(int64_t Val) const
Add a new immediate operand.
const MachineInstrBuilder & add(const MachineOperand &MO) const
const MachineInstrBuilder & addMBB(MachineBasicBlock *MBB, unsigned TargetFlags=0) const
const MachineInstrBuilder & addDef(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a virtual register definition operand.
const MachineInstrBuilder & setMIFlags(unsigned Flags) const
MachineInstr * getInstr() const
If conversion operators fail, use this method to get the MachineInstr explicitly.
Representation of each machine instruction.
unsigned getOpcode() const
Returns the opcode of this MachineInstr.
const MachineBasicBlock * getParent() const
unsigned getNumOperands() const
Retuns the total number of operands.
LLVM_ABI unsigned getNumExplicitOperands() const
Returns the number of non-implicit operands.
LLVM_ABI unsigned getNumExplicitDefs() const
Returns the number of non-implicit definitions.
LLVM_ABI void emitGenericError(const Twine &ErrMsg) const
LLVM_ABI const MachineFunction * getMF() const
Return the function that contains the basic block that this instruction belongs to.
const DebugLoc & getDebugLoc() const
Returns the debug location id of this MachineInstr.
const MachineOperand & getOperand(unsigned i) const
A description of a memory reference used in the backend.
@ MOVolatile
The memory access is volatile.
@ MONonTemporal
The memory access is non-temporal.
bool isReg() const
isReg - Tests if this is a MO_Register operand.
MachineBasicBlock * getMBB() const
Register getReg() const
getReg - Returns the register number.
MachineRegisterInfo - Keep track of information for virtual and physical registers,...
defusechain_instr_iterator< true, false, false, true > use_instr_iterator
use_instr_iterator/use_instr_begin/use_instr_end - Walk all uses of the specified register,...
const TargetRegisterClass * getRegClass(Register Reg) const
Return the register class of the specified virtual register.
LLVM_ABI MachineInstr * getVRegDef(Register Reg) const
getVRegDef - Return the machine instr that defines the specified virtual register or null if none is ...
use_instr_iterator use_instr_begin(Register RegNo) const
bool use_nodbg_empty(Register RegNo) const
use_nodbg_empty - Return true if there are no non-Debug instructions using the specified register.
static def_instr_iterator def_instr_end()
defusechain_instr_iterator< false, true, false, true > def_instr_iterator
def_instr_iterator/def_instr_begin/def_instr_end - Walk all defs of the specified register,...
LLVM_ABI Register createVirtualRegister(const TargetRegisterClass *RegClass, StringRef Name="")
createVirtualRegister - Create and return a new virtual register in the function with the specified r...
def_instr_iterator def_instr_begin(Register RegNo) const
LLT getType(Register Reg) const
Get the low-level type of Reg or LLT{} if Reg is not a generic (target independent) virtual register.
static use_instr_iterator use_instr_end()
iterator_range< use_instr_nodbg_iterator > use_nodbg_instructions(Register Reg) const
LLVM_ABI void setType(Register VReg, LLT Ty)
Set the low-level type of VReg to Ty.
const MachineFunction & getMF() const
LLVM_ABI void setRegClass(Register Reg, const TargetRegisterClass *RC)
setRegClass - Set the register class of the specified virtual register.
LLVM_ABI Register createGenericVirtualRegister(LLT Ty, StringRef Name="")
Create and return a new generic virtual register with low-level type Ty.
const TargetRegisterClass * getRegClassOrNull(Register Reg) const
Return the register class of Reg, or null if Reg has not been assigned a register class yet.
iterator_range< use_instr_iterator > use_instructions(Register Reg) const
unsigned getNumVirtRegs() const
getNumVirtRegs - Return the number of virtual registers created.
LLVM_ABI void replaceRegWith(Register FromReg, Register ToReg)
replaceRegWith - Replace all instances of FromReg with ToReg in the machine function.
Analysis providing profile information.
Holds all the information related to register banks.
Wrapper class representing virtual and physical registers.
constexpr bool isValid() const
constexpr bool isPhysical() const
Return true if the specified register number is in the physical register namespace.
bool isScalarOrVectorSigned(SPIRVTypeInst Type) const
SPIRVTypeInst getOrCreateOpTypeSampledImage(SPIRVTypeInst ImageType, MachineIRBuilder &MIRBuilder)
void assignSPIRVTypeToVReg(SPIRVTypeInst Type, Register VReg, const MachineFunction &MF)
const TargetRegisterClass * getRegClass(SPIRVTypeInst SpvType) const
MachineInstr * getOrAddMemAliasingINTELInst(MachineIRBuilder &MIRBuilder, const MDNode *AliasingListMD)
bool isAggregateType(SPIRVTypeInst Type) const
unsigned getScalarOrVectorBitWidth(SPIRVTypeInst Type) const
SPIRVTypeInst getOrCreateSPIRVIntegerType(unsigned BitWidth, MachineIRBuilder &MIRBuilder)
SPIRVTypeInst getOrCreateSPIRVVectorType(SPIRVTypeInst BaseType, unsigned NumElements, MachineIRBuilder &MIRBuilder, bool EmitIR)
Register buildGlobalVariable(Register Reg, SPIRVTypeInst BaseType, StringRef Name, const GlobalValue *GV, SPIRV::StorageClass::StorageClass Storage, const MachineInstr *Init, bool IsConst, const std::optional< SPIRV::LinkageType::LinkageType > &LinkageType, MachineIRBuilder &MIRBuilder, bool IsInstSelector)
SPIRVTypeInst getResultType(Register VReg, MachineFunction *MF=nullptr)
unsigned getScalarOrVectorComponentCount(Register VReg) const
const Type * getTypeForSPIRVType(SPIRVTypeInst Ty) const
bool isBitcastCompatible(SPIRVTypeInst Type1, SPIRVTypeInst Type2) const
unsigned getPointerSize() const
Register getOrCreateConstFP(APFloat Val, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
LLT getRegType(SPIRVTypeInst SpvType) const
void invalidateMachineInstr(MachineInstr *MI)
SPIRVTypeInst getOrCreateSPIRVBoolType(MachineIRBuilder &MIRBuilder, bool EmitIR)
SPIRVTypeInst getOrCreateSPIRVPointerType(const Type *BaseType, MachineIRBuilder &MIRBuilder, SPIRV::StorageClass::StorageClass SC)
bool isScalarOfType(Register VReg, unsigned TypeOpcode) const
Register getSPIRVTypeID(SPIRVTypeInst SpirvType) const
Register getOrCreateConstInt(uint64_t Val, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
Register getOrCreateConstIntArray(uint64_t Val, size_t Num, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII)
bool findValueAttrs(const MachineInstr *Key, Type *&Ty, StringRef &Name)
SPIRVTypeInst retrieveScalarOrVectorIntType(SPIRVTypeInst Type) const
Register getOrCreateGlobalVariableWithBinding(SPIRVTypeInst VarType, uint32_t Set, uint32_t Binding, StringRef Name, MachineIRBuilder &MIRBuilder)
SPIRVTypeInst changePointerStorageClass(SPIRVTypeInst PtrType, SPIRV::StorageClass::StorageClass SC, MachineInstr &I)
Register getOrCreateConstVector(uint64_t Val, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
Register buildConstantFP(APFloat Val, MachineIRBuilder &MIRBuilder, SPIRVTypeInst SpvType=nullptr)
void addGlobalObject(const Value *V, const MachineFunction *MF, Register R)
SPIRVTypeInst getScalarOrVectorComponentType(SPIRVTypeInst Type) const
void recordFunctionPointer(const MachineOperand *MO, const Function *F)
SPIRVTypeInst getOrCreateSPIRVFloatType(unsigned BitWidth, MachineInstr &I, const SPIRVInstrInfo &TII)
SPIRVTypeInst getPointeeType(SPIRVTypeInst PtrType)
SPIRVTypeInst getOrCreateSPIRVType(const Type *Type, MachineInstr &I, SPIRV::AccessQualifier::AccessQualifier AQ, bool EmitIR)
bool isScalarOrVectorOfType(Register VReg, unsigned TypeOpcode) const
MachineFunction * setCurrentFunc(MachineFunction &MF)
Register getOrCreateConstNullPtr(MachineIRBuilder &MIRBuilder, SPIRVTypeInst SpvType)
SPIRVTypeInst getSPIRVTypeForVReg(Register VReg, const MachineFunction *MF=nullptr) const
Type * getDeducedGlobalValueType(const GlobalValue *Global)
Register getOrCreateUndef(MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII)
SPIRV::StorageClass::StorageClass getPointerStorageClass(Register VReg) const
bool erase(const MachineInstr *MI)
bool add(SPIRV::IRHandle Handle, const MachineInstr *MI)
Register find(SPIRV::IRHandle Handle, const MachineFunction *MF)
bool isPhysicalSPIRV() const
bool isAtLeastSPIRVVer(VersionTuple VerToCompareTo) const
bool canUseExtInstSet(SPIRV::InstructionSet::InstructionSet E) const
bool isLogicalSPIRV() const
bool canUseExtension(SPIRV::Extension::Extension E) const
bool isTypeIntOrFloat() const
bool erase(PtrType Ptr)
Remove pointer from the set.
std::pair< iterator, bool > insert(PtrType Ptr)
Inserts Ptr if and only if there is no element in the container equal to Ptr.
bool contains(ConstPtrType Ptr) const
SmallPtrSet - This class implements a set which is optimized for holding SmallSize or less elements.
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.
unsigned ID
LLVM IR allows to use arbitrary numbers as calling convention identifiers.
Scope
Defines the scope in which this symbol should be visible: Default – Visible in the public interface o...
NodeAddr< DefNode * > Def
NodeAddr< InstrNode * > Instr
NodeAddr< UseNode * > Use
NodeAddr< FuncNode * > Func
BaseReg
Stack frame base register. Bit 0 of FREInfo.Info.
This is an optimization pass for GlobalISel generic memory operations.
@ Low
Lower the current thread's priority such that it does not affect foreground tasks significantly.
bool all_of(R &&range, UnaryPredicate P)
Provide wrappers to std::all_of which take ranges instead of having to pass begin/end explicitly.
void addStringImm(StringRef Str, MCInst &Inst)
MachineBasicBlock::iterator getOpVariableMBBIt(MachineFunction &MF)
int64_t getIConstValSext(Register ConstReg, const MachineRegisterInfo *MRI)
MachineInstrBuilder BuildMI(MachineFunction &MF, const MIMetadata &MIMD, const MCInstrDesc &MCID)
Builder interface. Specify how to create the initial instruction itself.
bool isTypeFoldingSupported(unsigned Opcode)
decltype(auto) dyn_cast(const From &Val)
dyn_cast<X> - Return the argument parameter cast to the specified type.
MachineInstr * getDef(const MachineOperand &MO, const MachineRegisterInfo *MRI)
void addNumImm(const APInt &Imm, MachineInstrBuilder &MIB)
LLVM_ABI void salvageDebugInfo(const MachineRegisterInfo &MRI, MachineInstr &MI)
Assuming the instruction MI is going to be deleted, attempt to salvage debug users of MI by writing t...
LLVM_ABI void constrainSelectedInstRegOperands(MachineInstr &I, const TargetInstrInfo &TII, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI)
Mutate the newly-selected instruction I to constrain its (possibly generic) virtual register operands...
bool isPreISelGenericOpcode(unsigned Opcode)
Check whether the given Opcode is a generic opcode that is not supposed to appear after ISel.
Register createVirtualRegister(SPIRVTypeInst SpvType, SPIRVGlobalRegistry *GR, MachineRegisterInfo *MRI, const MachineFunction &MF)
unsigned getArrayComponentCount(const MachineRegisterInfo *MRI, const MachineInstr *ResType)
void buildOpDecorate(Register Reg, MachineIRBuilder &MIRBuilder, SPIRV::Decoration::Decoration Dec, ArrayRef< uint32_t > DecArgs, StringRef StrImm)
LLVM_ABI bool isNullOrNullSplat(const MachineInstr &MI, const MachineRegisterInfo &MRI, bool AllowUndefs=false)
Return true if the value is a constant 0 integer or a splatted vector of a constant 0 integer (with n...
uint64_t getIConstVal(Register ConstReg, const MachineRegisterInfo *MRI)
SmallVector< MachineInstr *, 4 > createContinuedInstructions(MachineIRBuilder &MIRBuilder, unsigned Opcode, unsigned MinWC, unsigned ContinuedOpcode, ArrayRef< Register > Args, Register ReturnRegister, Register TypeID)
SPIRV::MemorySemantics::MemorySemantics getMemSemanticsForStorageClass(SPIRV::StorageClass::StorageClass SC)
constexpr unsigned storageClassToAddressSpace(SPIRV::StorageClass::StorageClass SC)
RelativeUniformCounterPtr ValuesPtrExpr VTableAddr Value
void buildOpName(Register Target, StringRef Name, MachineIRBuilder &MIRBuilder)
bool any_of(R &&range, UnaryPredicate P)
Provide wrappers to std::any_of which take ranges instead of having to pass begin/end explicitly.
MachineInstr * getImm(const MachineOperand &MO, const MachineRegisterInfo *MRI)
constexpr bool isPowerOf2_32(uint32_t Value)
Return true if the argument is a power of two > 0.
Type * toTypedPointer(Type *Ty)
LLVM_ABI raw_ostream & dbgs()
dbgs() - This returns a reference to a raw_ostream for debugging messages.
LLVM_ABI void report_fatal_error(Error Err, bool gen_crash_diag=true)
constexpr bool isGenericCastablePtr(SPIRV::StorageClass::StorageClass SC)
class LLVM_GSL_OWNER SmallVector
Forward declaration of SmallVector so that calculateSmallVectorDefaultInlinedElements can reference s...
MachineInstr * passCopy(MachineInstr *Def, const MachineRegisterInfo *MRI)
bool isa(const From &Val)
isa<X> - Return true if the parameter to the template is an instance of one of the template type argu...
std::optional< SPIRV::LinkageType::LinkageType > getSpirvLinkageTypeFor(const SPIRVSubtarget &ST, const GlobalValue &GV)
LLVM_ABI raw_fd_ostream & errs()
This returns a reference to a raw_ostream for standard error.
SPIRV::StorageClass::StorageClass addressSpaceToStorageClass(unsigned AddrSpace, const SPIRVSubtarget &STI)
AtomicOrdering
Atomic ordering for LLVM's memory model.
SPIRV::Scope::Scope getMemScope(LLVMContext &Ctx, SyncScope::ID Id)
InstructionSelector * createSPIRVInstructionSelector(const SPIRVTargetMachine &TM, const SPIRVSubtarget &Subtarget, const RegisterBankInfo &RBI)
std::string getStringValueFromReg(Register Reg, MachineRegisterInfo &MRI)
int64_t foldImm(const MachineOperand &MO, const MachineRegisterInfo *MRI)
DWARFExpression::Operation Op
ArrayRef(const T &OneElt) -> ArrayRef< T >
MachineInstr * getDefInstrMaybeConstant(Register &ConstReg, const MachineRegisterInfo *MRI)
constexpr unsigned BitWidth
decltype(auto) cast(const From &Val)
cast<X> - Return the argument parameter cast to the specified type.
bool hasInitializer(const GlobalVariable *GV)
bool isSpvIntrinsic(const MachineInstr &MI, Intrinsic::ID IntrinsicID)
MachineInstr * getVRegDef(MachineRegisterInfo &MRI, Register Reg)
SPIRV::MemorySemantics::MemorySemantics getMemSemantics(AtomicOrdering Ord)
std::string getLinkStringForBuiltIn(SPIRV::BuiltIn::BuiltIn BuiltInValue)
LLVM_ABI bool isTriviallyDead(const MachineInstr &MI, const MachineRegisterInfo &MRI)
Check whether an instruction MI is dead: it only defines dead virtual registers, and doesn't have oth...
MCRegisterClass TargetRegisterClass