34#include "llvm/IR/IntrinsicsSPIRV.h"
40#define DEBUG_TYPE "spirv-isel"
47 std::vector<std::pair<SPIRV::InstructionSet::InstructionSet, uint32_t>>;
52 std::optional<Register> Bias;
53 std::optional<Register>
Offset;
54 std::optional<Register> MinLod;
55 std::optional<Register> GradX;
56 std::optional<Register> GradY;
57 std::optional<Register> Lod;
58 std::optional<Register> Compare;
65 bool IsScalar =
false;
68llvm::SPIRV::SelectionControl::SelectionControl
69getSelectionOperandForImm(
int Imm) {
71 return SPIRV::SelectionControl::Flatten;
73 return SPIRV::SelectionControl::DontFlatten;
75 return SPIRV::SelectionControl::None;
79#define GET_GLOBALISEL_PREDICATE_BITSET
80#include "SPIRVGenGlobalISel.inc"
81#undef GET_GLOBALISEL_PREDICATE_BITSET
108#define GET_GLOBALISEL_PREDICATES_DECL
109#include "SPIRVGenGlobalISel.inc"
110#undef GET_GLOBALISEL_PREDICATES_DECL
112#define GET_GLOBALISEL_TEMPORARIES_DECL
113#include "SPIRVGenGlobalISel.inc"
114#undef GET_GLOBALISEL_TEMPORARIES_DECL
138 unsigned BitSetOpcode)
const;
142 unsigned BitSetOpcode)
const;
146 unsigned BitSetOpcode,
bool SwapPrimarySide)
const;
153 unsigned Opcode)
const;
156 unsigned Opcode)
const;
178 unsigned NewOpcode,
unsigned NegateOpcode = 0)
const;
195 unsigned OpType)
const;
262 unsigned Opcode)
const;
266 unsigned Opcode)
const;
270 unsigned Opcode)
const;
274 unsigned Opcode)
const;
276 template <
bool Signed>
279 template <
bool Signed>
286 template <
typename PickOpcodeFn>
289 PickOpcodeFn &&PickOpcode)
const;
306 template <
typename PickOpcodeFn>
309 PickOpcodeFn &&PickOpcode)
const;
327 bool IsSigned)
const;
329 bool IsSigned,
unsigned Opcode)
const;
331 bool IsSigned)
const;
337 bool IsSigned)
const;
378 GL::GLSLExtInst GLInst,
bool setMIFlags =
true,
379 bool useMISrc =
true,
381 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
382 CL::OpenCLExtInst CLInst,
bool setMIFlags =
true,
383 bool useMISrc =
true,
385 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
386 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
387 bool setMIFlags =
true,
bool useMISrc =
true,
389 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
390 const ExtInstList &ExtInsts,
bool setMIFlags =
true,
391 bool useMISrc =
true,
394 bool selectLog10(
Register ResVReg, SPIRVTypeInst ResType,
395 MachineInstr &
I)
const;
397 bool selectFpowi(
Register ResVReg, SPIRVTypeInst ResType,
398 MachineInstr &
I)
const;
400 bool selectSaturate(
Register ResVReg, SPIRVTypeInst ResType,
401 MachineInstr &
I)
const;
403 bool selectWaveOpInst(
Register ResVReg, SPIRVTypeInst ResType,
404 MachineInstr &
I,
unsigned Opcode)
const;
406 bool selectBarrierInst(MachineInstr &
I,
unsigned Scope,
unsigned MemSem,
407 bool WithGroupSync)
const;
409 bool selectWaveActiveCountBits(
Register ResVReg, SPIRVTypeInst ResType,
410 MachineInstr &
I)
const;
412 bool selectWaveActiveAllEqual(
Register ResVReg, SPIRVTypeInst ResType,
413 MachineInstr &
I)
const;
417 bool selectHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
418 MachineInstr &
I)
const;
420 bool selectCounterHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
421 MachineInstr &
I)
const;
423 bool selectReadImageIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
424 MachineInstr &
I)
const;
425 bool selectGetDimensionsIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
426 MachineInstr &
I)
const;
427 bool selectGetDimensionsLevelsIntrinsic(
Register &ResVReg,
428 SPIRVTypeInst ResType,
429 MachineInstr &
I)
const;
430 bool selectGetDimensionsMSIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
431 MachineInstr &
I)
const;
434 std::optional<Register> LodReg = std::nullopt)
const;
435 bool selectSampleBasicIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
436 MachineInstr &
I)
const;
437 bool selectCalculateLodIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
438 MachineInstr &
I)
const;
439 bool selectSampleBiasIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
440 MachineInstr &
I)
const;
441 bool selectSampleGradIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
442 MachineInstr &
I)
const;
443 bool selectSampleLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
444 MachineInstr &
I)
const;
445 bool selectLoadLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
446 MachineInstr &
I)
const;
447 bool selectSampleCmpIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
448 MachineInstr &
I)
const;
449 bool selectSampleCmpLevelZeroIntrinsic(
Register &ResVReg,
450 SPIRVTypeInst ResType,
451 MachineInstr &
I)
const;
452 bool selectGatherIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
453 MachineInstr &
I)
const;
454 bool selectImageWriteIntrinsic(MachineInstr &
I)
const;
455 bool selectResourceGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
456 MachineInstr &
I)
const;
457 bool selectPushConstantGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
458 MachineInstr &
I)
const;
459 bool selectResourceNonUniformIndex(
Register &ResVReg, SPIRVTypeInst ResType,
460 MachineInstr &
I)
const;
461 bool selectModf(
Register ResVReg, SPIRVTypeInst ResType,
462 MachineInstr &
I)
const;
463 bool selectUpdateCounter(
Register &ResVReg, SPIRVTypeInst ResType,
464 MachineInstr &
I)
const;
465 bool selectFrexp(
Register ResVReg, SPIRVTypeInst ResType,
466 MachineInstr &
I)
const;
467 bool selectSincos(
Register ResVReg, SPIRVTypeInst ResType,
468 MachineInstr &
I)
const;
469 bool selectExp10(
Register ResVReg, SPIRVTypeInst ResType,
470 MachineInstr &
I)
const;
471 bool selectDerivativeInst(
Register ResVReg, SPIRVTypeInst ResType,
472 MachineInstr &
I,
const unsigned DPdOpCode)
const;
474 Register buildI32Constant(uint32_t Val, MachineInstr &
I,
475 SPIRVTypeInst ResType =
nullptr)
const;
476 Register buildI32ConstantInEntryBlock(uint32_t Val, MachineInstr &
I,
477 SPIRVTypeInst ResType =
nullptr)
const;
479 Register buildZerosVal(SPIRVTypeInst ResType, MachineInstr &
I)
const;
480 bool isScalarOrVectorIntConstantZero(
Register Reg)
const;
481 Register buildZerosValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
483 MachineInstr &
I)
const;
484 Register buildOnesValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
486 bool wrapIntoSpecConstantOp(MachineInstr &
I,
489 Register getUcharPtrTypeReg(MachineInstr &
I,
490 SPIRV::StorageClass::StorageClass SC)
const;
491 MachineInstrBuilder buildSpecConstantOp(MachineInstr &
I,
Register Dest,
493 uint32_t Opcode)
const;
494 MachineInstrBuilder buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
495 SPIRVTypeInst SrcPtrTy)
const;
496 Register buildPointerToResource(SPIRVTypeInst ResType,
497 SPIRV::StorageClass::StorageClass SC,
498 uint32_t Set, uint32_t
Binding,
499 uint32_t ArraySize,
Register IndexReg,
501 MachineIRBuilder MIRBuilder)
const;
502 SPIRVTypeInst widenTypeToVec4(SPIRVTypeInst
Type, MachineInstr &
I)
const;
503 bool extractSubvector(
Register &ResVReg, SPIRVTypeInst ResType,
504 Register &ReadReg, MachineInstr &InsertionPoint)
const;
505 bool generateImageReadOrFetch(
Register &ResVReg, SPIRVTypeInst ResType,
508 const ImageOperands *ImOps =
nullptr)
const;
509 bool generateSampleImage(
Register ResVReg, SPIRVTypeInst ResType,
511 Register CoordinateReg,
const ImageOperands &ImOps,
514 bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
515 Register ResVReg, SPIRVTypeInst ResType,
516 MachineInstr &
I)
const;
517 bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
518 Register ResVReg, SPIRVTypeInst ResType,
519 MachineInstr &
I)
const;
520 bool loadHandleBeforePosition(
Register &HandleReg, SPIRVTypeInst ResType,
521 GIntrinsic &HandleDef, MachineInstr &Pos)
const;
522 void decorateUsesAsNonUniform(
Register &NonUniformReg)
const;
523 bool errorIfInstrOutsideShader(MachineInstr &
I)
const;
525 std::optional<SplitParts> splitEvenOddLanes(
Register PopCountReg,
526 unsigned ComponentCount,
528 SPIRVTypeInst I32Type)
const;
531 handle64BitOverflow(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
532 Register SrcReg,
unsigned int Opcode,
533 std::function<
bool(
Register, SPIRVTypeInst,
534 MachineInstr &,
Register,
unsigned)>
538bool sampledTypeIsSignedInteger(
const llvm::Type *HandleType) {
540 if (
TET->getTargetExtName() ==
"spirv.Image") {
543 assert(
TET->getTargetExtName() ==
"spirv.SignedImage");
544 return TET->getTypeParameter(0)->isIntegerTy();
548#define GET_GLOBALISEL_IMPL
549#include "SPIRVGenGlobalISel.inc"
550#undef GET_GLOBALISEL_IMPL
556 TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
559#include
"SPIRVGenGlobalISel.inc"
562#include
"SPIRVGenGlobalISel.inc"
574 InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
578void SPIRVInstructionSelector::resetVRegsType(MachineFunction &MF) {
579 if (HasVRegsReset == &MF)
594 for (
const auto &
MBB : MF) {
595 for (
const auto &
MI :
MBB) {
598 if (
MI.getOpcode() != SPIRV::ASSIGN_TYPE)
602 LLT DstType = MRI.
getType(DstReg);
604 LLT SrcType = MRI.
getType(SrcReg);
605 if (DstType != SrcType)
610 if (DstRC != SrcRC && SrcRC)
622 while (!Stack.empty()) {
627 switch (
MI->getOpcode()) {
628 case TargetOpcode::G_INTRINSIC:
629 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
630 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS: {
633 if (IntrID != Intrinsic::spv_const_composite &&
634 IntrID != Intrinsic::spv_undef && IntrID != Intrinsic::spv_poison)
638 case TargetOpcode::G_BUILD_VECTOR:
639 case TargetOpcode::G_SPLAT_VECTOR:
641 i < OpDef->getNumOperands(); i++) {
646 Stack.push_back(OpNestedDef);
649 case TargetOpcode::G_CONSTANT:
650 case TargetOpcode::G_FCONSTANT:
651 case TargetOpcode::G_IMPLICIT_DEF:
652 case SPIRV::OpConstantTrue:
653 case SPIRV::OpConstantFalse:
654 case SPIRV::OpConstantI:
655 case SPIRV::OpConstantF:
656 case SPIRV::OpConstantComposite:
657 case SPIRV::OpConstantCompositeContinuedINTEL:
658 case SPIRV::OpConstantSampler:
659 case SPIRV::OpConstantNull:
661 case SPIRV::OpPoisonKHR:
662 case SPIRV::OpConstantFunctionPointerINTEL:
689 case Intrinsic::spv_all:
690 case Intrinsic::spv_alloca:
691 case Intrinsic::spv_any:
692 case Intrinsic::spv_bitcast:
693 case Intrinsic::spv_const_composite:
694 case Intrinsic::spv_cross:
695 case Intrinsic::spv_degrees:
696 case Intrinsic::spv_distance:
697 case Intrinsic::spv_extractelt:
698 case Intrinsic::spv_extractv:
699 case Intrinsic::spv_faceforward:
700 case Intrinsic::spv_fdot:
701 case Intrinsic::spv_firstbitlow:
702 case Intrinsic::spv_firstbitshigh:
703 case Intrinsic::spv_firstbituhigh:
704 case Intrinsic::spv_frac:
705 case Intrinsic::spv_gep:
706 case Intrinsic::spv_global_offset:
707 case Intrinsic::spv_global_size:
708 case Intrinsic::spv_group_id:
709 case Intrinsic::spv_insertelt:
710 case Intrinsic::spv_insertv:
711 case Intrinsic::spv_isinf:
712 case Intrinsic::spv_isnan:
713 case Intrinsic::spv_isfinite:
714 case Intrinsic::spv_isnormal:
715 case Intrinsic::spv_lerp:
716 case Intrinsic::spv_length:
717 case Intrinsic::spv_normalize:
718 case Intrinsic::spv_num_subgroups:
719 case Intrinsic::spv_num_workgroups:
720 case Intrinsic::spv_ptrcast:
721 case Intrinsic::spv_radians:
722 case Intrinsic::spv_reflect:
723 case Intrinsic::spv_refract:
724 case Intrinsic::spv_resource_getbasepointer:
725 case Intrinsic::spv_resource_getpointer:
726 case Intrinsic::spv_resource_handlefrombinding:
727 case Intrinsic::spv_resource_handlefromimplicitbinding:
728 case Intrinsic::spv_resource_nonuniformindex:
729 case Intrinsic::spv_resource_sample:
730 case Intrinsic::spv_rsqrt:
731 case Intrinsic::spv_saturate:
732 case Intrinsic::spv_sdot:
733 case Intrinsic::spv_sign:
734 case Intrinsic::spv_smoothstep:
735 case Intrinsic::spv_step:
736 case Intrinsic::spv_subgroup_id:
737 case Intrinsic::spv_subgroup_local_invocation_id:
738 case Intrinsic::spv_subgroup_max_size:
739 case Intrinsic::spv_subgroup_size:
740 case Intrinsic::spv_thread_id:
741 case Intrinsic::spv_thread_id_in_group:
742 case Intrinsic::spv_udot:
743 case Intrinsic::spv_undef:
744 case Intrinsic::spv_value_md:
745 case Intrinsic::spv_workgroup_size:
757 case SPIRV::OpTypeVoid:
758 case SPIRV::OpTypeBool:
759 case SPIRV::OpTypeInt:
760 case SPIRV::OpTypeFloat:
761 case SPIRV::OpTypeVector:
762 case SPIRV::OpTypeMatrix:
763 case SPIRV::OpTypeImage:
764 case SPIRV::OpTypeSampler:
765 case SPIRV::OpTypeSampledImage:
766 case SPIRV::OpTypeArray:
767 case SPIRV::OpTypeRuntimeArray:
768 case SPIRV::OpTypeStruct:
769 case SPIRV::OpTypeOpaque:
770 case SPIRV::OpTypePointer:
771 case SPIRV::OpTypeFunction:
772 case SPIRV::OpTypeEvent:
773 case SPIRV::OpTypeDeviceEvent:
774 case SPIRV::OpTypeReserveId:
775 case SPIRV::OpTypeQueue:
776 case SPIRV::OpTypePipe:
777 case SPIRV::OpTypeForwardPointer:
778 case SPIRV::OpTypePipeStorage:
779 case SPIRV::OpTypeNamedBarrier:
780 case SPIRV::OpTypeAccelerationStructureNV:
781 case SPIRV::OpTypeCooperativeMatrixNV:
782 case SPIRV::OpTypeCooperativeMatrixKHR:
792 if (
MI.getNumDefs() == 0)
795 for (
const auto &MO :
MI.all_defs()) {
797 if (
Reg.isPhysical()) {
802 if (
UseMI.getOpcode() != SPIRV::OpName) {
809 if (
MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE ||
MI.isFakeUse() ||
810 MI.isLifetimeMarker()) {
813 <<
"Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
824 if (
MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
825 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
828 LLVM_DEBUG(
dbgs() <<
"Dead: Intrinsic with no real side effects.\n");
833 if (
MI.mayStore() ||
MI.isCall() ||
834 (
MI.mayLoad() &&
MI.hasOrderedMemoryRef()) ||
MI.isPosition() ||
835 MI.isDebugInstr() ||
MI.isTerminator() ||
MI.isJumpTableDebugInfo()) {
836 LLVM_DEBUG(
dbgs() <<
"Not dead: instruction has side effects.\n");
847 LLVM_DEBUG(
dbgs() <<
"Dead: known opcode with no side effects\n");
854void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &
MI)
const {
856 for (
const auto &MO :
MI.all_defs()) {
860 SmallVector<MachineInstr *, 4> UselessOpNames;
863 "There is still a use of the dead function.");
866 for (MachineInstr *OpNameMI : UselessOpNames) {
868 OpNameMI->eraseFromParent();
873void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &
MI)
const {
876 removeOpNamesForDeadMI(
MI);
877 MI.eraseFromParent();
880bool SPIRVInstructionSelector::select(MachineInstr &
I) {
881 resetVRegsType(*
I.getParent()->getParent());
883 assert(
I.getParent() &&
"Instruction should be in a basic block!");
884 assert(
I.getParent()->getParent() &&
"Instruction should be in a function!");
889 removeDeadInstruction(
I);
896 if (Opcode == SPIRV::ASSIGN_TYPE) {
897 Register DstReg =
I.getOperand(0).getReg();
898 Register SrcReg =
I.getOperand(1).getReg();
901 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
902 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
903 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
904 Register SelectDstReg =
Def->getOperand(0).getReg();
905 bool SuccessToSelectSelect [[maybe_unused]] = selectSelect(
907 assert(SuccessToSelectSelect);
909 Def->eraseFromParent();
916 bool Res = selectImpl(
I, *CoverageInfo);
918 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
919 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
923 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
935 }
else if (
I.getNumDefs() == 1) {
947 removeDeadInstruction(
I);
952 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
953 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
959 bool HasDefs =
I.getNumDefs() > 0;
962 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
963 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
964 if (spvSelect(ResVReg, ResType,
I)) {
966 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
977 case TargetOpcode::G_CONSTANT:
978 case TargetOpcode::G_FCONSTANT:
985 MachineInstr &
I)
const {
988 if (DstRC != SrcRC && SrcRC)
990 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
997bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
998 SPIRVTypeInst ResType,
999 MachineInstr &
I)
const {
1000 const unsigned Opcode =
I.getOpcode();
1002 return selectImpl(
I, *CoverageInfo);
1004 case TargetOpcode::G_CONSTANT:
1005 case TargetOpcode::G_FCONSTANT:
1006 return selectConst(ResVReg, ResType,
I);
1007 case TargetOpcode::G_GLOBAL_VALUE:
1008 return selectGlobalValue(ResVReg,
I);
1009 case TargetOpcode::G_IMPLICIT_DEF:
1010 return selectOpUndef(ResVReg, ResType,
I);
1011 case TargetOpcode::G_FREEZE:
1012 return selectFreeze(ResVReg, ResType,
I);
1014 case TargetOpcode::G_INTRINSIC:
1015 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
1016 case TargetOpcode::G_INTRINSIC_CONVERGENT:
1017 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
1018 return selectIntrinsic(ResVReg, ResType,
I);
1019 case TargetOpcode::G_BITREVERSE:
1020 return selectBitreverse(ResVReg, ResType,
I);
1022 case TargetOpcode::G_BUILD_VECTOR:
1023 return selectBuildVector(ResVReg, ResType,
I);
1024 case TargetOpcode::G_SPLAT_VECTOR:
1025 return selectSplatVector(ResVReg, ResType,
I);
1026 case TargetOpcode::G_CONCAT_VECTORS:
1027 return selectConcatVectors(ResVReg, ResType,
I);
1029 case TargetOpcode::G_SHUFFLE_VECTOR: {
1030 MachineBasicBlock &BB = *
I.getParent();
1031 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
1034 .
addUse(
I.getOperand(1).getReg())
1035 .
addUse(
I.getOperand(2).getReg());
1036 for (
auto V :
I.getOperand(3).getShuffleMask())
1041 case TargetOpcode::G_MEMMOVE:
1042 case TargetOpcode::G_MEMCPY:
1043 case TargetOpcode::G_MEMCPY_INLINE:
1044 case TargetOpcode::G_MEMSET:
1045 case TargetOpcode::G_MEMSET_INLINE:
1046 return selectMemOperation(ResVReg,
I);
1048 case TargetOpcode::G_ICMP:
1049 return selectICmp(ResVReg, ResType,
I);
1050 case TargetOpcode::G_FCMP:
1051 return selectFCmp(ResVReg, ResType,
I);
1053 case TargetOpcode::G_FRAME_INDEX:
1054 return selectFrameIndex(ResVReg, ResType,
I);
1056 case TargetOpcode::G_LOAD:
1057 return selectLoad(ResVReg, ResType,
I);
1058 case TargetOpcode::G_STORE:
1059 return selectStore(
I);
1061 case TargetOpcode::G_BR:
1062 return selectBranch(
I);
1063 case TargetOpcode::G_BRCOND:
1064 return selectBranchCond(
I);
1066 case TargetOpcode::G_PHI:
1067 return selectPhi(ResVReg,
I);
1069 case TargetOpcode::G_FPTOSI:
1070 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1071 case TargetOpcode::G_FPTOUI:
1072 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1074 case TargetOpcode::G_FPTOSI_SAT:
1075 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1076 case TargetOpcode::G_FPTOUI_SAT:
1077 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1079 case TargetOpcode::G_SITOFP:
1080 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
1081 case TargetOpcode::G_UITOFP:
1082 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
1084 case TargetOpcode::G_CTPOP:
1085 return selectPopCount(ResVReg, ResType,
I, SPIRV::OpBitCount);
1086 case TargetOpcode::G_SMIN:
1087 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
1088 case TargetOpcode::G_UMIN:
1089 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
1091 case TargetOpcode::G_SMAX:
1092 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
1093 case TargetOpcode::G_UMAX:
1094 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
1096 case TargetOpcode::G_SCMP:
1097 return selectSUCmp(ResVReg, ResType,
I,
true);
1098 case TargetOpcode::G_UCMP:
1099 return selectSUCmp(ResVReg, ResType,
I,
false);
1100 case TargetOpcode::G_LROUND:
1101 case TargetOpcode::G_LLROUND: {
1104 MRI->
setRegClass(regForLround, &SPIRV::iIDRegClass);
1106 regForLround, *(
I.getParent()->getParent()));
1108 CL::round, GL::Round,
false);
1110 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1117 case TargetOpcode::G_STRICT_FMA:
1118 case TargetOpcode::G_FMA: {
1121 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1124 .
addUse(
I.getOperand(1).getReg())
1125 .
addUse(
I.getOperand(2).getReg())
1126 .
addUse(
I.getOperand(3).getReg())
1131 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1134 case TargetOpcode::G_STRICT_FLDEXP:
1135 return selectExtInst(ResVReg, ResType,
I, CL::ldexp);
1137 case TargetOpcode::G_FPOW:
1138 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1139 case TargetOpcode::G_FPOWI:
1140 return selectFpowi(ResVReg, ResType,
I);
1142 case TargetOpcode::G_FEXP:
1143 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1144 case TargetOpcode::G_FEXP2:
1145 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1146 case TargetOpcode::G_FEXP10:
1147 return selectExp10(ResVReg, ResType,
I);
1149 case TargetOpcode::G_FMODF:
1150 return selectModf(ResVReg, ResType,
I);
1151 case TargetOpcode::G_FSINCOS:
1152 return selectSincos(ResVReg, ResType,
I);
1154 case TargetOpcode::G_FLOG:
1155 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1156 case TargetOpcode::G_FLOG2:
1157 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1158 case TargetOpcode::G_FLOG10:
1159 return selectLog10(ResVReg, ResType,
I);
1161 case TargetOpcode::G_FABS:
1162 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1163 case TargetOpcode::G_ABS:
1164 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1166 case TargetOpcode::G_FMINNUM:
1167 case TargetOpcode::G_FMINIMUM:
1168 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1169 case TargetOpcode::G_FMAXNUM:
1170 case TargetOpcode::G_FMAXIMUM:
1171 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1173 case TargetOpcode::G_FCOPYSIGN:
1174 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1176 case TargetOpcode::G_FCEIL:
1177 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1178 case TargetOpcode::G_FFLOOR:
1179 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1181 case TargetOpcode::G_FCOS:
1182 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1183 case TargetOpcode::G_FSIN:
1184 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1185 case TargetOpcode::G_FTAN:
1186 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1187 case TargetOpcode::G_FACOS:
1188 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1189 case TargetOpcode::G_FASIN:
1190 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1191 case TargetOpcode::G_FATAN:
1192 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1193 case TargetOpcode::G_FATAN2:
1194 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1195 case TargetOpcode::G_FCOSH:
1196 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1197 case TargetOpcode::G_FSINH:
1198 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1199 case TargetOpcode::G_FTANH:
1200 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1202 case TargetOpcode::G_STRICT_FSQRT:
1203 case TargetOpcode::G_FSQRT:
1204 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1206 case TargetOpcode::G_CTTZ:
1207 case TargetOpcode::G_CTTZ_ZERO_POISON:
1208 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1209 case TargetOpcode::G_CTLZ:
1210 case TargetOpcode::G_CTLZ_ZERO_POISON:
1211 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1213 case TargetOpcode::G_INTRINSIC_ROUND:
1214 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1215 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1216 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1217 case TargetOpcode::G_INTRINSIC_TRUNC:
1218 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1219 case TargetOpcode::G_FRINT:
1220 case TargetOpcode::G_FNEARBYINT:
1221 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1223 case TargetOpcode::G_SMULH:
1224 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1225 case TargetOpcode::G_UMULH:
1226 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1228 case TargetOpcode::G_SADDSAT:
1229 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1230 case TargetOpcode::G_UADDSAT:
1231 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1232 case TargetOpcode::G_SSUBSAT:
1233 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1234 case TargetOpcode::G_USUBSAT:
1235 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1237 case TargetOpcode::G_FFREXP:
1238 return selectFrexp(ResVReg, ResType,
I);
1240 case TargetOpcode::G_UADDO:
1241 return selectOverflowArith(ResVReg, ResType,
I,
1242 ResType->
getOpcode() == SPIRV::OpTypeVector
1243 ? SPIRV::OpIAddCarryV
1244 : SPIRV::OpIAddCarryS);
1245 case TargetOpcode::G_USUBO:
1246 return selectOverflowArith(ResVReg, ResType,
I,
1247 ResType->
getOpcode() == SPIRV::OpTypeVector
1248 ? SPIRV::OpISubBorrowV
1249 : SPIRV::OpISubBorrowS);
1250 case TargetOpcode::G_UMULO:
1251 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1252 case TargetOpcode::G_SMULO:
1253 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1255 case TargetOpcode::G_SEXT:
1256 return selectExt(ResVReg, ResType,
I,
true);
1257 case TargetOpcode::G_ANYEXT:
1258 case TargetOpcode::G_ZEXT:
1259 return selectExt(ResVReg, ResType,
I,
false);
1260 case TargetOpcode::G_TRUNC:
1261 return selectTrunc(ResVReg, ResType,
I);
1262 case TargetOpcode::G_FPTRUNC:
1263 case TargetOpcode::G_FPEXT:
1264 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1266 case TargetOpcode::G_PTRTOINT:
1267 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1268 case TargetOpcode::G_INTTOPTR:
1269 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1270 case TargetOpcode::G_BITCAST:
1271 return selectBitcast(ResVReg, ResType,
I);
1272 case TargetOpcode::G_ADDRSPACE_CAST:
1273 return selectAddrSpaceCast(ResVReg, ResType,
I);
1274 case TargetOpcode::G_PTRMASK:
1275 return selectPtrMask(ResVReg, ResType,
I);
1276 case TargetOpcode::G_PTR_ADD: {
1278 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1282 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1283 (*II).getOpcode() == TargetOpcode::COPY ||
1284 (*II).getOpcode() == SPIRV::OpVariable) &&
1285 getImm(
I.getOperand(2), MRI));
1287 bool IsGVInit =
false;
1291 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1292 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1293 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1294 (*UseIt).getOpcode() == SPIRV::OpVariable) {
1304 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1316 return diagnoseUnsupported(
1317 I,
"incompatible result and operand types in a bitcast");
1319 MachineInstrBuilder MIB =
1320 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1327 : SPIRV::OpInBoundsPtrAccessChain))
1331 .
addUse(
I.getOperand(2).getReg())
1334 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1338 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1340 .
addUse(
I.getOperand(2).getReg())
1349 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1352 .
addImm(
static_cast<uint32_t
>(
1353 SPIRV::Opcode::InBoundsPtrAccessChain))
1356 .
addUse(
I.getOperand(2).getReg());
1361 case TargetOpcode::G_ATOMICRMW_OR:
1362 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1363 case TargetOpcode::G_ATOMICRMW_ADD:
1364 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1365 case TargetOpcode::G_ATOMICRMW_AND:
1366 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1367 case TargetOpcode::G_ATOMICRMW_MAX:
1368 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1369 case TargetOpcode::G_ATOMICRMW_MIN:
1370 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1371 case TargetOpcode::G_ATOMICRMW_SUB:
1372 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1373 case TargetOpcode::G_ATOMICRMW_XOR:
1374 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1375 case TargetOpcode::G_ATOMICRMW_UMAX:
1376 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1377 case TargetOpcode::G_ATOMICRMW_UMIN:
1378 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1379 case TargetOpcode::G_ATOMICRMW_XCHG:
1380 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1382 case TargetOpcode::G_ATOMICRMW_FADD:
1383 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1384 case TargetOpcode::G_ATOMICRMW_FSUB:
1386 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1387 ResType->
getOpcode() == SPIRV::OpTypeVector
1389 : SPIRV::OpFNegate);
1390 case TargetOpcode::G_ATOMICRMW_FMIN:
1391 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1392 case TargetOpcode::G_ATOMICRMW_FMAX:
1393 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1395 case TargetOpcode::G_FENCE:
1396 return selectFence(
I);
1398 case TargetOpcode::G_STACKSAVE:
1399 return selectStackSave(ResVReg, ResType,
I);
1400 case TargetOpcode::G_STACKRESTORE:
1401 return selectStackRestore(
I);
1403 case TargetOpcode::G_UNMERGE_VALUES:
1406 case TargetOpcode::G_TRAP:
1407 case TargetOpcode::G_UBSANTRAP:
1408 return selectTrap(
I);
1413 case TargetOpcode::DBG_LABEL:
1415 case TargetOpcode::G_DEBUGTRAP:
1416 return selectDebugTrap(ResVReg, ResType,
I);
1423bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1424 SPIRVTypeInst ResType,
1425 MachineInstr &
I)
const {
1426 unsigned Opcode = SPIRV::OpNop;
1433bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1434 SPIRVTypeInst ResType,
1436 GL::GLSLExtInst GLInst,
1437 bool setMIFlags,
bool useMISrc,
1440 SPIRV::InstructionSet::InstructionSet::GLSL_std_450))
1441 return diagnoseUnsupported(
1443 "this instruction is only supported with the GLSL extended instruction "
1445 return selectExtInst(ResVReg, ResType,
I,
1446 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1447 setMIFlags, useMISrc, SrcRegs);
1450bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1451 SPIRVTypeInst ResType,
1453 CL::OpenCLExtInst CLInst,
1454 bool setMIFlags,
bool useMISrc,
1456 return selectExtInst(ResVReg, ResType,
I,
1457 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1458 setMIFlags, useMISrc, SrcRegs);
1461bool SPIRVInstructionSelector::selectExtInst(
1462 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1463 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1465 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1466 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1467 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1471bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1472 SPIRVTypeInst ResType,
1475 bool setMIFlags,
bool useMISrc,
1478 for (
const auto &[InstructionSet, Opcode] : Insts) {
1482 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1485 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1490 const unsigned NumOps =
I.getNumOperands();
1493 I.getOperand(Index).getType() ==
1494 MachineOperand::MachineOperandType::MO_IntrinsicID)
1497 MIB.
add(
I.getOperand(Index));
1509bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1510 SPIRVTypeInst ResType,
1511 MachineInstr &
I)
const {
1512 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1513 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1514 for (
const auto &Ex : ExtInsts) {
1515 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1516 uint32_t Opcode = Ex.second;
1520 MachineIRBuilder MIRBuilder(
I);
1523 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1528 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1531 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1535 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1538 .
addImm(
static_cast<uint32_t
>(Ex.first))
1540 .
add(
I.getOperand(2))
1544 Register ExpResReg =
I.getOperand(1).getReg();
1546 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1556bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1557 SPIRVTypeInst ResType,
1558 MachineInstr &
I)
const {
1559 Register CosResVReg =
I.getOperand(1).getReg();
1560 unsigned SrcIdx =
I.getNumExplicitDefs();
1565 MachineIRBuilder MIRBuilder(
I);
1567 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1572 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1575 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1577 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1580 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1582 .
add(
I.getOperand(SrcIdx))
1585 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1593 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1596 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1598 .
add(
I.getOperand(SrcIdx))
1600 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1603 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1605 .
add(
I.getOperand(SrcIdx))
1612bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1613 SPIRVTypeInst ResType,
1616 unsigned Opcode)
const {
1617 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1627std::optional<SplitParts> SPIRVInstructionSelector::splitEvenOddLanes(
1628 Register PopCountReg,
unsigned ComponentCount, MachineInstr &
I,
1629 SPIRVTypeInst I32Type)
const {
1632 if (ComponentCount == 1) {
1635 Parts.IsScalar =
true;
1636 Parts.Type = I32Type;
1644 if (!selectOpWithSrcs(Parts.High, I32Type,
I, {PopCountReg, IdxOne},
1645 SPIRV::OpVectorExtractDynamic))
1646 return std::nullopt;
1648 if (!selectOpWithSrcs(Parts.Low, I32Type,
I, {PopCountReg, IdxZero},
1649 SPIRV::OpVectorExtractDynamic))
1650 return std::nullopt;
1654 MachineIRBuilder MIRBuilder(
I);
1655 Parts.IsScalar =
false;
1662 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1663 TII.get(SPIRV::OpVectorShuffle))
1668 for (
unsigned J = 1; J < ComponentCount * 2; J += 2)
1673 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1674 TII.get(SPIRV::OpVectorShuffle))
1679 for (
unsigned J = 0; J < ComponentCount * 2; J += 2)
1687bool SPIRVInstructionSelector::selectPopCount16(
Register ResVReg,
1688 SPIRVTypeInst ResType,
1691 unsigned Opcode)
const {
1692 Register OpReg =
I.getOperand(1).getReg();
1695 MachineIRBuilder MIRBuilder(
I);
1697 SPIRVTypeInst I32VectorType =
1700 bool IsVector = NumElems > 1;
1701 SPIRVTypeInst ExtType = IsVector ? I32VectorType : I32Type;
1704 if (!selectOpWithSrcs(ExtReg, ExtType,
I, {OpReg}, SPIRV::OpUConvert))
1708 if (!selectPopCount32(PopCountReg, ExtType,
I, ExtReg, Opcode))
1711 return selectOpWithSrcs(ResVReg, ResType,
I, {PopCountReg}, ExtOpcode);
1714bool SPIRVInstructionSelector::selectPopCount32(
Register ResVReg,
1715 SPIRVTypeInst ResType,
1718 unsigned Opcode)
const {
1719 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
1722bool SPIRVInstructionSelector::selectPopCount64(
Register ResVReg,
1723 SPIRVTypeInst ResType,
1726 unsigned Opcode)
const {
1728 if (ComponentCount > 2)
1729 return handle64BitOverflow(
1730 ResVReg, ResType,
I, SrcReg, Opcode,
1732 unsigned O) {
return this->selectPopCount64(R,
T,
I, S, O); });
1734 MachineIRBuilder MIRBuilder(
I);
1739 I32Type, 2 * ComponentCount, MIRBuilder,
false);
1743 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
1748 if (!selectPopCount32(Pop32, VecI32Type,
I, Vec32, Opcode))
1752 auto MaybeParts = splitEvenOddLanes(Pop32, ComponentCount,
I, I32Type);
1755 SplitParts &Parts = *MaybeParts;
1758 unsigned OpAdd = Parts.IsScalar ? SPIRV::OpIAddS : SPIRV::OpIAddV;
1760 if (!selectOpWithSrcs(Sum, Parts.Type,
I, {Parts.High, Parts.Low}, OpAdd))
1765 unsigned ConvOp = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
1766 return selectOpWithSrcs(ResVReg, ResType,
I, {Sum}, ConvOp);
1769bool SPIRVInstructionSelector::selectPopCount(
Register ResVReg,
1770 SPIRVTypeInst ResType,
1772 unsigned Opcode)
const {
1777 if (!STI.getTargetTriple().isVulkanOS())
1778 return selectUnOp(ResVReg, ResType,
I, Opcode);
1780 Register OpReg =
I.getOperand(1).getReg();
1783 : SPIRV::OpUConvert;
1787 return selectPopCount16(ResVReg, ResType,
I, ExtOpcode, Opcode);
1789 return selectPopCount32(ResVReg, ResType,
I, OpReg, Opcode);
1791 return selectPopCount64(ResVReg, ResType,
I, OpReg, Opcode);
1793 return diagnoseUnsupported(
I,
"unsupported operand bit width for popcount");
1797bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1798 SPIRVTypeInst ResType,
1800 unsigned Opcode)
const {
1802 Register SrcReg =
I.getOperand(1).getReg();
1807 unsigned DefOpCode = DefIt->getOpcode();
1808 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1811 if (
auto *VRD =
getVRegDef(*MRI, DefIt->getOperand(1).getReg()))
1812 DefOpCode = VRD->getOpcode();
1814 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1815 DefOpCode == TargetOpcode::G_CONSTANT ||
1816 DefOpCode == SPIRV::OpVariable || DefOpCode == SPIRV::OpConstantI) {
1822 uint32_t SpecOpcode = 0;
1824 case SPIRV::OpConvertPtrToU:
1825 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1827 case SPIRV::OpConvertUToPtr:
1828 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1833 TII.get(SPIRV::OpSpecConstantOp))
1843 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1847bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1848 SPIRVTypeInst ResType,
1849 MachineInstr &
I)
const {
1850 Register OpReg =
I.getOperand(1).getReg();
1851 SPIRVTypeInst OpType =
1854 return diagnoseUnsupported(
1855 I,
"incompatible result and operand types in a bitcast");
1856 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1866 if (
MemOp->isVolatile())
1867 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1868 if (
MemOp->isNonTemporal())
1869 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1871 if (!ST->isShader() &&
MemOp->getAlign().value())
1872 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1876 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1877 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1881 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1883 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1887 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1891 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1893 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1905 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1907 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1909 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
1913bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
1914 SPIRVTypeInst ResType,
1915 MachineInstr &
I)
const {
1917 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1922 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
1923 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
1925 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1927 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1931 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1935 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1936 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
1937 I.getDebugLoc(),
I);
1941 MachineIRBuilder MIRBuilder(
I);
1943 if (
I.getNumMemOperands()) {
1944 const MachineMemOperand *MemOp = *
I.memoperands_begin();
1945 if (MemOp->isAtomic())
1946 return selectAtomicLoad(ResVReg, ResType,
I);
1949 auto MIB = MIRBuilder.buildInstr(SPIRV::OpLoad)
1953 if (!
I.getNumMemOperands()) {
1954 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1956 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1965bool SPIRVInstructionSelector::selectAtomicLoad(
Register ResVReg,
1966 SPIRVTypeInst ResType,
1967 MachineInstr &
I)
const {
1968 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
1971 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1974 return diagnoseUnsupported(
1975 I,
"Lowering to SPIR-V of atomic load is only "
1976 "allowed for integer, floating point or pointer types");
1978 assert(
I.getNumMemOperands());
1979 const MachineMemOperand &MemOp = **
I.memoperands_begin();
1980 assert(MemOp.isAtomic());
1984 Register ScopeReg = buildI32Constant(Scope,
I);
1990 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
1991 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
1994 MachineIRBuilder MIRBuilder(
I);
1998 return diagnoseUnsupported(
1999 I,
"Lowering to SPIR-V of atomic load is only "
2000 "allowed for pointer types for physical addressing model");
2007 SPIRVTypeInst PtrAsIntSpirvType =
2018 PtrAsIntSpirvType, MIRBuilder,
2021 MIRBuilder.getMF());
2023 MIRBuilder.buildInstr(SPIRV::OpBitcast)
2024 .addDef(PtrCastedToMatchValReg)
2027 .constrainAllUses(
TII,
TRI, RBI);
2029 MIRBuilder.buildInstr(SPIRV::OpAtomicLoad)
2032 .addUse(PtrCastedToMatchValReg)
2035 .constrainAllUses(
TII,
TRI, RBI);
2036 MIRBuilder.buildInstr(SPIRV::OpConvertUToPtr)
2040 .constrainAllUses(
TII,
TRI, RBI);
2043 auto AtomicLoad = MIRBuilder.buildInstr(SPIRV::OpAtomicLoad)
2049 AtomicLoad.constrainAllUses(
TII,
TRI, RBI);
2054bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
2056 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2057 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2062 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2063 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2065 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2070 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2074 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2075 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2076 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2077 TII.get(SPIRV::OpImageWrite))
2083 if (sampledTypeIsSignedInteger(LLVMHandleType))
2086 BMI.constrainAllUses(
TII,
TRI, RBI);
2091 if (
I.getNumMemOperands()) {
2092 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2093 if (MemOp->isAtomic())
2094 return selectAtomicStore(
I);
2101 if (PtrSC == SPIRV::StorageClass::UniformConstant ||
2102 PtrSC == SPIRV::StorageClass::Input ||
2103 PtrSC == SPIRV::StorageClass::PushConstant)
2104 return diagnoseUnsupported(
2105 I,
"store into a read-only SPIR-V storage class is not allowed");
2107 MachineIRBuilder MIRBuilder(
I);
2108 auto MIB = MIRBuilder.buildInstr(SPIRV::OpStore).
addUse(Ptr).
addUse(StoreVal);
2109 if (!
I.getNumMemOperands()) {
2110 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2112 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2121bool SPIRVInstructionSelector::selectAtomicStore(MachineInstr &
I)
const {
2122 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2125 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2126 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2131 assert(
I.getNumMemOperands());
2132 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2133 assert(MemOp.isAtomic());
2137 Register ScopeReg = buildI32Constant(Scope,
I);
2143 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2144 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2146 MachineIRBuilder MIRBuilder(
I);
2150 return diagnoseUnsupported(
2151 I,
"Lowering to SPIR-V of atomic store is only "
2152 "allowed for pointer types for physical addressing model");
2158 SPIRVTypeInst PtrAsIntSpirvType =
2165 MIRBuilder.buildInstr(SPIRV::OpConvertPtrToU)
2169 .constrainAllUses(
TII,
TRI, RBI);
2175 PtrAsIntSpirvType, MIRBuilder,
2178 MIRBuilder.getMF());
2180 MIRBuilder.buildInstr(SPIRV::OpBitcast)
2181 .addDef(PtrCastedToMatchValReg)
2184 .constrainAllUses(
TII,
TRI, RBI);
2186 StoreVal = PtrToUVal;
2187 Ptr = PtrCastedToMatchValReg;
2188 PointeeType = PtrAsIntSpirvType;
2192 return diagnoseUnsupported(
I,
2193 "Lowering to SPIR-V of atomic store is only "
2194 "allowed for integer or floating point types");
2196 auto AtomicStore = MIRBuilder.buildInstr(SPIRV::OpAtomicStore)
2201 AtomicStore.constrainAllUses(
TII,
TRI, RBI);
2206bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
2207 SPIRVTypeInst ResType,
2208 MachineInstr &
I)
const {
2209 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
2217 const Register PtrsReg =
I.getOperand(2).getReg();
2218 const uint32_t Alignment =
I.getOperand(3).getImm();
2219 const Register MaskReg =
I.getOperand(4).getReg();
2220 const Register PassthruReg =
I.getOperand(5).getReg();
2221 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2225 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
2236bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
2237 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
2244 const Register ValuesReg =
I.getOperand(1).getReg();
2245 const Register PtrsReg =
I.getOperand(2).getReg();
2246 const uint32_t Alignment =
I.getOperand(3).getImm();
2247 const Register MaskReg =
I.getOperand(4).getReg();
2248 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2252 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
2261bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
2262 const Twine &
Msg)
const {
2263 const Function &
F =
I.getMF()->getFunction();
2264 F.getContext().diagnose(
2265 DiagnosticInfoUnsupported(
F,
Msg,
I.getDebugLoc(),
DS_Error));
2269bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
2270 SPIRVTypeInst ResType,
2271 MachineInstr &
I)
const {
2272 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2273 return diagnoseUnsupported(
2274 I,
"llvm.stacksave intrinsic: this instruction requires the following "
2275 "SPIR-V extension: SPV_INTEL_variable_length_array");
2277 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
2284bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
2285 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2286 return diagnoseUnsupported(
2288 "llvm.stackrestore intrinsic: this instruction requires the following "
2289 "SPIR-V extension: SPV_INTEL_variable_length_array");
2290 if (!
I.getOperand(0).isReg())
2293 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
2294 .
addUse(
I.getOperand(0).getReg())
2300SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
2301 MachineIRBuilder MIRBuilder(
I);
2302 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
2309 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
2313 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2314 Type *ArrTy = ArrayType::get(ValTy, Num);
2316 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
2319 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
2326 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
2329 .
addImm(SPIRV::StorageClass::UniformConstant)
2340bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
2343 Register DstReg =
I.getOperand(0).getReg();
2347 return diagnoseUnsupported(
2348 I,
"OpCopyMemory requires operands to have the same type");
2349 uint64_t CopySize =
getIConstVal(
I.getOperand(2).getReg(), MRI);
2353 return diagnoseUnsupported(
2354 I,
"Unable to determine pointee type size for OpCopyMemory");
2355 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
2356 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
2357 return diagnoseUnsupported(
2358 I,
"OpCopyMemory requires the size to match the pointee type size");
2359 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemory))
2362 if (
I.getNumMemOperands()) {
2363 MachineIRBuilder MIRBuilder(
I);
2370bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &
I,
2373 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemorySized))
2374 .
addUse(
I.getOperand(0).getReg())
2376 .
addUse(
I.getOperand(2).getReg());
2377 if (
I.getNumMemOperands()) {
2378 MachineIRBuilder MIRBuilder(
I);
2385bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
2386 MachineInstr &
I)
const {
2388 Register SizeReg =
I.getOperand(2).getReg();
2390 SizeDef && SizeDef->
getOpcode() == TargetOpcode::G_CONSTANT &&
2394 Register SrcReg =
I.getOperand(1).getReg();
2395 if (
I.getOpcode() == TargetOpcode::G_MEMSET ||
2396 I.getOpcode() == TargetOpcode::G_MEMSET_INLINE) {
2397 Register VarReg = getOrCreateMemSetGlobal(
I);
2400 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2402 ValTy,
I, SPIRV::StorageClass::UniformConstant);
2404 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
2408 if (!selectCopyMemory(
I, SrcReg))
2411 if (!selectCopyMemorySized(
I, SrcReg))
2414 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
2415 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
2420bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
2421 SPIRVTypeInst ResType,
2424 unsigned NegateOpcode)
const {
2426 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2429 Register ScopeReg = buildI32Constant(Scope,
I);
2431 Register Ptr =
I.getOperand(1).getReg();
2432 uint32_t ScSem =
static_cast<uint32_t
>(
2436 Register MemSemReg = buildI32Constant(MemSem,
I);
2438 Register ValueReg =
I.getOperand(2).getReg();
2439 if (NegateOpcode != 0) {
2442 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
2448 if (NewOpcode != SPIRV::OpAtomicExchange)
2449 return diagnoseUnsupported(
2450 I,
"Lowering to SPIR-V of this atomic operation is not "
2451 "allowed for pointer types");
2453 return diagnoseUnsupported(
2454 I,
"Lowering to SPIR-V of atomic exchange is only "
2455 "allowed for pointer types for physical addressing model");
2462 MachineIRBuilder MIRBuilder(
I);
2464 SPIRVTypeInst PtrAsIntSpirvType =
2471 MIRBuilder.getMF());
2472 MIRBuilder.buildInstr(SPIRV::OpConvertPtrToU)
2473 .addDef(ValueAsIntReg)
2476 .constrainAllUses(
TII,
TRI, RBI);
2484 MIRBuilder.getMF());
2485 MIRBuilder.buildInstr(SPIRV::OpBitcast)
2486 .addDef(PtrCastedToMatchValReg)
2489 .constrainAllUses(
TII,
TRI, RBI);
2495 MIRBuilder.getMF());
2496 MIRBuilder.buildInstr(SPIRV::OpAtomicExchange)
2497 .addDef(ExchangeResReg)
2499 .addUse(PtrCastedToMatchValReg)
2502 .addUse(ValueAsIntReg)
2503 .constrainAllUses(
TII,
TRI, RBI);
2504 MIRBuilder.buildInstr(SPIRV::OpConvertUToPtr)
2507 .addUse(ExchangeResReg)
2508 .constrainAllUses(
TII,
TRI, RBI);
2512 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2523bool SPIRVInstructionSelector::selectInterlockedOp(
Register ResVReg,
2524 SPIRVTypeInst ResType,
2526 unsigned Opcode)
const {
2527 Register Ptr =
I.getOperand(2).getReg();
2531 assert((SC == SPIRV::StorageClass::Workgroup ||
2532 SC == SPIRV::StorageClass::StorageBuffer) &&
2533 "InterlockedAdd requires Workgroup or StorageBuffer storage class");
2534 uint32_t
Scope =
static_cast<uint32_t
>(SC == SPIRV::StorageClass::Workgroup
2535 ? SPIRV::Scope::Workgroup
2536 : SPIRV::Scope::Device);
2537 Register ScopeReg = buildI32Constant(Scope,
I);
2540 Register MemSemReg = buildI32Constant(MemSem,
I);
2542 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
2553bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2554 unsigned ArgI =
I.getNumOperands() - 1;
2556 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2557 SPIRVTypeInst SrcType =
2559 if (!SrcType || SrcType->
getOpcode() != SPIRV::OpTypeVector)
2561 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2565 unsigned CurrentIndex = 0;
2566 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2567 Register ResVReg =
I.getOperand(i).getReg();
2570 LLT ResLLT = MRI->
getType(ResVReg);
2576 ResType = ScalarType;
2582 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
2585 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2591 for (
unsigned j = 0;
j < NumElements; ++
j) {
2592 MIB.
addImm(CurrentIndex + j);
2594 CurrentIndex += NumElements;
2598 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2610bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2613 Register MemSemReg = buildI32Constant(MemSem,
I);
2615 uint32_t
Scope =
static_cast<uint32_t
>(
2617 Register ScopeReg = buildI32Constant(Scope,
I);
2619 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2626bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2627 SPIRVTypeInst ResType,
2629 unsigned Opcode)
const {
2630 Type *ResTy =
nullptr;
2633 return diagnoseUnsupported(
2635 "Not enough info to select the arithmetic with overflow instruction");
2637 return diagnoseUnsupported(
I,
2638 "Expect struct type result for the arithmetic "
2639 "with overflow instruction");
2645 MachineIRBuilder MIRBuilder(
I);
2647 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2648 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2654 Register ZeroReg = buildZerosVal(ResType,
I);
2659 if (ResName.
size() > 0)
2664 BuildMI(BB, MIRBuilder.getInsertPt(),
I.getDebugLoc(),
TII.get(Opcode))
2667 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2668 MIB.
addUse(
I.getOperand(i).getReg());
2673 MRI->
setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2674 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2676 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2677 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2684 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2685 .
addDef(
I.getOperand(1).getReg())
2693bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2694 SPIRVTypeInst ResType,
2695 MachineInstr &
I)
const {
2697 "selectAtomicCmpXchg only handles the spv_cmpxchg intrinsic");
2698 Register Ptr =
I.getOperand(2).getReg();
2699 Register ScopeReg =
I.getOperand(5).getReg();
2700 Register MemSemEqReg =
I.getOperand(6).getReg();
2701 Register MemSemNeqReg =
I.getOperand(7).getReg();
2703 Register Val =
I.getOperand(4).getReg();
2707 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2726 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2733 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2745 case SPIRV::StorageClass::DeviceOnlyINTEL:
2746 case SPIRV::StorageClass::HostOnlyINTEL:
2755 bool IsGRef =
false;
2756 bool IsAllowedRefs =
2758 unsigned Opcode = It.getOpcode();
2759 if (Opcode == SPIRV::OpConstantComposite ||
2760 Opcode == SPIRV::OpSpecConstantComposite ||
2761 Opcode == SPIRV::OpVariable ||
2762 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2763 return IsGRef = true;
2764 return Opcode == SPIRV::OpName;
2766 return IsAllowedRefs && IsGRef;
2769Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2770 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2772 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2776SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2778 uint32_t Opcode)
const {
2779 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2780 TII.get(SPIRV::OpSpecConstantOp))
2788SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2789 SPIRVTypeInst SrcPtrTy)
const {
2790 SPIRVTypeInst GenericPtrTy =
2794 SPIRV::StorageClass::Generic),
2796 MachineFunction *MF =
I.getParent()->getParent();
2798 MachineInstrBuilder MIB = buildSpecConstantOp(
2800 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2810bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2811 SPIRVTypeInst ResType,
2812 MachineInstr &
I)
const {
2816 Register SrcPtr =
I.getOperand(1).getReg();
2820 if (SrcPtrTy->
getOpcode() != SPIRV::OpTypePointer ||
2821 ResType->
getOpcode() != SPIRV::OpTypePointer)
2822 return BuildCOPY(ResVReg, SrcPtr,
I);
2832 unsigned SpecOpcode =
2834 ?
static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric)
2837 ? static_cast<uint32_t>(
SPIRV::Opcode::GenericCastToPtr)
2844 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2846 .constrainAllUses(
TII,
TRI, RBI);
2848 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2850 buildSpecConstantOp(
2852 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2853 .constrainAllUses(
TII,
TRI, RBI);
2860 return BuildCOPY(ResVReg, SrcPtr,
I);
2862 if ((SrcSC == SPIRV::StorageClass::Function &&
2863 DstSC == SPIRV::StorageClass::Private) ||
2864 (DstSC == SPIRV::StorageClass::Function &&
2865 SrcSC == SPIRV::StorageClass::Private))
2866 return BuildCOPY(ResVReg, SrcPtr,
I);
2870 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2873 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2876 SPIRVTypeInst GenericPtrTy =
2895 return selectUnOp(ResVReg, ResType,
I,
2896 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
2898 return selectUnOp(ResVReg, ResType,
I,
2899 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
2901 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2903 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2913bool SPIRVInstructionSelector::selectPtrMask(
Register ResVReg,
2914 SPIRVTypeInst ResType,
2915 MachineInstr &
I)
const {
2917 return diagnoseUnsupported(
2918 I,
"G_PTRMASK is not supported with logical SPIR-V");
2923 Register PtrReg =
I.getOperand(1).getReg();
2924 Register MaskReg =
I.getOperand(2).getReg();
2943 ? SPIRV::OpBitwiseAndV
2944 : SPIRV::OpBitwiseAndS;
2967 return SPIRV::OpFOrdEqual;
2969 return SPIRV::OpFOrdGreaterThanEqual;
2971 return SPIRV::OpFOrdGreaterThan;
2973 return SPIRV::OpFOrdLessThanEqual;
2975 return SPIRV::OpFOrdLessThan;
2977 return SPIRV::OpFOrdNotEqual;
2979 return SPIRV::OpOrdered;
2981 return SPIRV::OpFUnordEqual;
2983 return SPIRV::OpFUnordGreaterThanEqual;
2985 return SPIRV::OpFUnordGreaterThan;
2987 return SPIRV::OpFUnordLessThanEqual;
2989 return SPIRV::OpFUnordLessThan;
2991 return SPIRV::OpFUnordNotEqual;
2993 return SPIRV::OpUnordered;
3003 return SPIRV::OpIEqual;
3005 return SPIRV::OpINotEqual;
3007 return SPIRV::OpSGreaterThanEqual;
3009 return SPIRV::OpSGreaterThan;
3011 return SPIRV::OpSLessThanEqual;
3013 return SPIRV::OpSLessThan;
3015 return SPIRV::OpUGreaterThanEqual;
3017 return SPIRV::OpUGreaterThan;
3019 return SPIRV::OpULessThanEqual;
3021 return SPIRV::OpULessThan;
3030 return SPIRV::OpPtrEqual;
3032 return SPIRV::OpPtrNotEqual;
3043 return SPIRV::OpLogicalEqual;
3045 return SPIRV::OpLogicalNotEqual;
3079bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
3080 SPIRVTypeInst ResType,
3082 unsigned OpAnyOrAll)
const {
3083 assert(
I.getNumOperands() == 3);
3084 assert(
I.getOperand(2).isReg());
3086 Register InputRegister =
I.getOperand(2).getReg();
3089 assert(InputType &&
"VReg has no type assigned");
3092 bool IsVectorTy = InputType->
getOpcode() == SPIRV::OpTypeVector;
3093 if (IsBoolTy && !IsVectorTy) {
3094 assert(ResVReg ==
I.getOperand(0).getReg());
3095 return BuildCOPY(ResVReg, InputRegister,
I);
3099 unsigned SpirvNotEqualId =
3100 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
3102 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
3107 IsBoolTy ? InputRegister
3115 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
3117 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
3134bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
3135 SPIRVTypeInst ResType,
3136 MachineInstr &
I)
const {
3137 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
3140bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
3141 SPIRVTypeInst ResType,
3142 MachineInstr &
I)
const {
3143 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
3147bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
3148 SPIRVTypeInst ResType,
3149 MachineInstr &
I)
const {
3150 assert(
I.getNumOperands() == 4);
3151 assert(
I.getOperand(2).isReg());
3152 assert(
I.getOperand(3).isReg());
3154 [[maybe_unused]] SPIRVTypeInst VecType =
3159 "dot product requires a vector of at least 2 components");
3161 [[maybe_unused]] SPIRVTypeInst EltType =
3170 .
addUse(
I.getOperand(2).getReg())
3171 .
addUse(
I.getOperand(3).getReg())
3176bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
3177 SPIRVTypeInst ResType,
3180 assert(
I.getNumOperands() == 4);
3181 assert(
I.getOperand(2).isReg());
3182 assert(
I.getOperand(3).isReg());
3185 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3189 .
addUse(
I.getOperand(2).getReg())
3190 .
addUse(
I.getOperand(3).getReg())
3197bool SPIRVInstructionSelector::selectIntegerDotExpansion(
3198 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3199 assert(
I.getNumOperands() == 4);
3200 assert(
I.getOperand(2).isReg());
3201 assert(
I.getOperand(3).isReg());
3205 Register Vec0 =
I.getOperand(2).getReg();
3206 Register Vec1 =
I.getOperand(3).getReg();
3210 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
3219 "dot product requires a vector of at least 2 components");
3222 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3232 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3243 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3255bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
3256 SPIRVTypeInst ResType,
3257 MachineInstr &
I)
const {
3259 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
3262 .
addUse(
I.getOperand(2).getReg())
3267bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
3268 SPIRVTypeInst ResType,
3269 MachineInstr &
I)
const {
3271 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
3274 .
addUse(
I.getOperand(2).getReg())
3279bool SPIRVInstructionSelector::selectOpIsFinite(
Register ResVReg,
3280 SPIRVTypeInst ResType,
3281 MachineInstr &
I)
const {
3283 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsFinite))
3286 .
addUse(
I.getOperand(2).getReg())
3291bool SPIRVInstructionSelector::selectOpIsNormal(
Register ResVReg,
3292 SPIRVTypeInst ResType,
3293 MachineInstr &
I)
const {
3295 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNormal))
3298 .
addUse(
I.getOperand(2).getReg())
3303template <
bool Signed>
3304bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
3305 SPIRVTypeInst ResType,
3306 MachineInstr &
I)
const {
3307 assert(
I.getNumOperands() == 5);
3308 assert(
I.getOperand(2).isReg());
3309 assert(
I.getOperand(3).isReg());
3310 assert(
I.getOperand(4).isReg());
3313 Register Acc =
I.getOperand(2).getReg();
3317 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3319 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
3324 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
3327 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3339template <
bool Signed>
3340bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
3341 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3342 assert(
I.getNumOperands() == 5);
3343 assert(
I.getOperand(2).isReg());
3344 assert(
I.getOperand(3).isReg());
3345 assert(
I.getOperand(4).isReg());
3348 Register Acc =
I.getOperand(2).getReg();
3354 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
3358 for (
unsigned i = 0; i < 4; i++) {
3381 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
3401 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3416bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
3417 SPIRVTypeInst ResType,
3418 MachineInstr &
I)
const {
3419 assert(
I.getNumOperands() == 3);
3420 assert(
I.getOperand(2).isReg());
3422 Register VZero = buildZerosValF(ResType,
I);
3423 Register VOne = buildOnesValF(ResType,
I);
3425 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
3428 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3430 .
addUse(
I.getOperand(2).getReg())
3437bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
3438 SPIRVTypeInst ResType,
3439 MachineInstr &
I)
const {
3440 assert(
I.getNumOperands() == 3);
3441 assert(
I.getOperand(2).isReg());
3443 Register InputRegister =
I.getOperand(2).getReg();
3445 auto &
DL =
I.getDebugLoc();
3448 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3455 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
3457 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
3465 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3470 if (NeedsConversion) {
3471 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
3482bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
3483 SPIRVTypeInst ResType,
3485 unsigned Opcode)
const {
3489 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
3495 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
3496 BMI.addUse(
I.getOperand(J).getReg());
3503bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
3506 bool WithGroupSync)
const {
3508 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
3510 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
3512 assert(((Scope != SPIRV::Scope::Workgroup) ||
3513 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
3514 "Workgroup Scope must set WorkGroupMemory semantic "
3515 "in Barrier instruction");
3517 assert(((Scope != SPIRV::Scope::Device) ||
3518 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
3519 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
3520 "Device Scope must set UniformMemory and ImageMemory semantic "
3521 "in Barrier instruction");
3527 if (WithGroupSync) {
3528 Register ExecReg = buildI32Constant(SPIRV::Scope::Workgroup,
I);
3532 Register ScopeReg = buildI32Constant(Scope,
I);
3533 Register MemSemReg = buildI32Constant(MemSem,
I);
3535 MI.addUse(ScopeReg).addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
3539bool SPIRVInstructionSelector::selectWaveActiveCountBits(
3540 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3545 if (!selectWaveOpInst(BallotReg, BallotType,
I,
3546 SPIRV::OpGroupNonUniformBallot))
3551 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3556 .
addImm(SPIRV::GroupOperation::Reduce)
3563bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
3564 SPIRVTypeInst ResType,
3565 MachineInstr &
I)
const {
3570 Register InputReg =
I.getOperand(2).getReg();
3575 bool IsVector = NumElems > 1;
3588 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
3589 SPIRV::OpGroupNonUniformAllEqual);
3594 ElementResults.
reserve(NumElems);
3596 for (
unsigned Idx = 0; Idx < NumElems; ++Idx) {
3609 ElemInput = Extracted;
3615 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3626 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3637bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3638 SPIRVTypeInst ResType,
3639 MachineInstr &
I)
const {
3641 assert(
I.getNumOperands() == 3);
3643 auto Op =
I.getOperand(2);
3653 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3655 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3656 return diagnoseUnsupported(
I,
"WavePrefixBitCount requires boolean input");
3677 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3681 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3688bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3689 SPIRVTypeInst ResType,
3691 bool IsUnsigned)
const {
3692 return selectWaveReduce(
3693 ResVReg, ResType,
I, IsUnsigned,
3694 [&](
Register InputRegister,
bool IsUnsigned) {
3695 const bool IsFloatTy =
3697 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3698 : SPIRV::OpGroupNonUniformSMax;
3699 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3703bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3704 SPIRVTypeInst ResType,
3706 bool IsUnsigned)
const {
3707 return selectWaveReduce(
3708 ResVReg, ResType,
I, IsUnsigned,
3709 [&](
Register InputRegister,
bool IsUnsigned) {
3710 const bool IsFloatTy =
3712 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3713 : SPIRV::OpGroupNonUniformSMin;
3714 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3718bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3719 SPIRVTypeInst ResType,
3720 MachineInstr &
I)
const {
3721 return selectWaveReduce(ResVReg, ResType,
I,
false,
3722 [&](
Register InputRegister,
bool IsUnsigned) {
3724 InputRegister, SPIRV::OpTypeFloat);
3725 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3726 : SPIRV::OpGroupNonUniformIAdd;
3730bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3731 SPIRVTypeInst ResType,
3732 MachineInstr &
I)
const {
3733 return selectWaveReduce(ResVReg, ResType,
I,
false,
3734 [&](
Register InputRegister,
bool IsUnsigned) {
3736 InputRegister, SPIRV::OpTypeFloat);
3737 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3738 : SPIRV::OpGroupNonUniformIMul;
3742template <
typename PickOpcodeFn>
3743bool SPIRVInstructionSelector::selectWaveReduce(
3744 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3745 PickOpcodeFn &&PickOpcode)
const {
3746 assert(
I.getNumOperands() == 3);
3747 assert(
I.getOperand(2).isReg());
3749 Register InputRegister =
I.getOperand(2).getReg();
3753 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3756 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3762 .
addImm(SPIRV::GroupOperation::Reduce)
3763 .
addUse(
I.getOperand(2).getReg())
3768bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3769 SPIRVTypeInst ResType,
3771 unsigned Opcode)
const {
3772 return selectWaveReduce(
3773 ResVReg, ResType,
I,
false,
3774 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3777bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3778 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3779 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3780 [&](
Register InputRegister,
bool IsUnsigned) {
3782 InputRegister, SPIRV::OpTypeFloat);
3784 ? SPIRV::OpGroupNonUniformFAdd
3785 : SPIRV::OpGroupNonUniformIAdd;
3789bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3790 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3791 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3792 [&](
Register InputRegister,
bool IsUnsigned) {
3794 InputRegister, SPIRV::OpTypeFloat);
3796 ? SPIRV::OpGroupNonUniformFMul
3797 : SPIRV::OpGroupNonUniformIMul;
3801template <
typename PickOpcodeFn>
3802bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3803 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3804 PickOpcodeFn &&PickOpcode)
const {
3805 assert(
I.getNumOperands() == 3);
3806 assert(
I.getOperand(2).isReg());
3808 Register InputRegister =
I.getOperand(2).getReg();
3812 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3815 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3821 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3822 .
addUse(
I.getOperand(2).getReg())
3827bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
3828 SPIRVTypeInst ResType,
3831 assert(
I.getNumOperands() == 3);
3832 assert(
I.getOperand(2).isReg());
3834 Register InputRegister =
I.getOperand(2).getReg();
3840 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
3851bool SPIRVInstructionSelector::selectBitreverseViaI32(
Register ResVReg,
3852 SPIRVTypeInst ResType,
3859 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
3864 : SPIRV::OpUConvert;
3868 ShiftOp = SPIRV::OpShiftRightLogicalV;
3873 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3874 TII.get(SPIRV::OpConstantComposite))
3877 for (
unsigned It = 0; It <
N; ++It)
3881 ShiftConst = CompositeReg;
3886 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
3891 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
3896 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
3901 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
3904bool SPIRVInstructionSelector::handle64BitOverflow(
3906 unsigned int Opcode,
3913 "handle64BitOverflow should only be used for integer types");
3915 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
3917 MachineIRBuilder MIRBuilder(
I);
3919 SPIRVTypeInst I64x2Type =
3921 SPIRVTypeInst Vec2ResType =
3924 std::vector<Register> PartialRegs;
3926 unsigned CurrentComponent = 0;
3927 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
3931 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3932 TII.get(SPIRV::OpVectorShuffle))
3937 .
addImm(CurrentComponent)
3938 .
addImm(CurrentComponent + 1);
3948 PartialRegs.push_back(SubVecReg);
3951 if (CurrentComponent != ComponentCount) {
3957 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
3958 SPIRV::OpVectorExtractDynamic))
3967 PartialRegs.push_back(FinalElemResReg);
3971 return selectOpWithSrcs(ResVReg, ResType,
I, PartialRegs,
3972 SPIRV::OpCompositeConstruct);
3975bool SPIRVInstructionSelector::selectBitreverse64(
Register ResVReg,
3976 SPIRVTypeInst ResType,
3980 if (ComponentCount > 2)
3981 return handle64BitOverflow(
3982 ResVReg, ResType,
I, SrcReg, SPIRV::OpBitReverse,
3984 unsigned O) {
return this->selectBitreverse64(R,
T,
I, S); });
3986 MachineIRBuilder MIRBuilder(
I);
3990 I32Type, 2 * ComponentCount, MIRBuilder,
false);
3994 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
3999 if (!selectBitreverseNative(Reverse32, VecI32Type,
I, Vec32))
4006 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4007 TII.get(SPIRV::OpVectorShuffle))
4012 for (
unsigned J = 0; J < ComponentCount; ++J) {
4019 return selectOpWithSrcs(ResVReg, ResType,
I, {SwappedVec}, SPIRV::OpBitcast);
4022bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
4023 SPIRVTypeInst ResType,
4027 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
4035bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
4036 SPIRVTypeInst ResType,
4037 MachineInstr &
I)
const {
4038 Register OpReg =
I.getOperand(1).getReg();
4047 return selectBitreverseViaI32(ResVReg, ResType,
I, OpReg);
4049 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4051 return selectBitreverse64(ResVReg, ResType,
I, OpReg);
4053 return SPIRVInstructionSelector::diagnoseUnsupported(
4054 I,
"G_BITREVERSE only support 16,32,64 bits.");
4058 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4069 unsigned AndOp = SPIRV::OpBitwiseAndS;
4070 unsigned OrOp = SPIRV::OpBitwiseOrS;
4071 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
4072 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
4074 AndOp = SPIRV::OpBitwiseAndV;
4075 OrOp = SPIRV::OpBitwiseOrV;
4076 ShlOp = SPIRV::OpShiftLeftLogicalV;
4077 ShrOp = SPIRV::OpShiftRightLogicalV;
4083 const unsigned Shift) ->
Register {
4091 Register MaskReg = CreateConst(Mask);
4092 Register ShiftReg = CreateConst(Shift);
4099 if (!selectOpWithSrcs(
T1, ResType,
I, {Input, ShiftReg}, ShrOp) ||
4100 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
4101 !selectOpWithSrcs(T3, ResType,
I, {Input, MaskReg}, AndOp) ||
4102 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
4103 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
4111 uint64_t
Mask = ~0ull;
4112 while ((Shift >>= 1) > 0) {
4119 return BuildCOPY(ResVReg, Result,
I);
4122bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
4123 SPIRVTypeInst ResType,
4124 MachineInstr &
I)
const {
4125 assert(
I.getOperand(0).isReg() &&
I.getOperand(1).isReg() &&
4126 "G_FREEZE must define and use a register");
4127 Register OpReg =
I.getOperand(1).getReg();
4131 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
4144 if (MachineInstr *Def = MRI->
getVRegDef(OpReg)) {
4145 if (
Def->getOpcode() == TargetOpcode::COPY)
4148 switch (
Def->getOpcode()) {
4149 case SPIRV::ASSIGN_TYPE:
4150 if (MachineInstr *AssignToDef =
4152 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
4153 Reg =
Def->getOperand(2).getReg();
4156 case SPIRV::OpUndef:
4157 Reg =
Def->getOperand(1).getReg();
4160 unsigned DestOpCode;
4162 DestOpCode = SPIRV::OpConstantNull;
4163 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze of a "
4164 "static undef/poison lowered to OpConstantNull\n");
4166 DestOpCode = TargetOpcode::COPY;
4168 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze "
4169 "skipped, lowered as a copy of the operand\n");
4171 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
4172 .
addDef(
I.getOperand(0).getReg())
4180bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
4181 SPIRVTypeInst ResType,
4182 MachineInstr &
I)
const {
4184 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4186 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4190 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
4195 for (
unsigned i =
I.getNumExplicitDefs();
4196 i <
I.getNumExplicitOperands() && IsConst; ++i)
4200 if (!IsConst &&
N < 2)
4201 return diagnoseUnsupported(
4202 I,
"There must be at least two constituent operands in a vector");
4207 for (
unsigned i =
I.getNumExplicitDefs();
4208 i <
I.getNumExplicitOperands() && IsNullVector; ++i) {
4209 MachineInstr *
Def =
getDef(
I.getOperand(i), MRI);
4214 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4221 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4222 TII.get(IsConst ? SPIRV::OpConstantComposite
4223 : SPIRV::OpCompositeConstruct))
4226 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
4227 MIB.
addUse(
I.getOperand(i).getReg());
4232bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
4233 SPIRVTypeInst ResType,
4234 MachineInstr &
I)
const {
4236 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4238 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4244 if (!
I.getOperand(
OpIdx).isReg())
4251 if (!IsConst &&
N < 2)
4252 return diagnoseUnsupported(
4253 I,
"There must be at least two constituent operands in a vector");
4256 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4257 TII.get(IsConst ? SPIRV::OpConstantComposite
4258 : SPIRV::OpCompositeConstruct))
4261 for (
unsigned i = 0; i <
N; ++i)
4267bool SPIRVInstructionSelector::selectConcatVectors(
Register ResVReg,
4268 SPIRVTypeInst ResType,
4269 MachineInstr &
I)
const {
4273 if (ResType->
getOpcode() != SPIRV::OpTypeVector)
4275 "Cannot select G_CONCAT_VECTORS with a non-vector result");
4277 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4278 TII.get(SPIRV::OpCompositeConstruct))
4288bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
4289 SPIRVTypeInst ResType,
4290 MachineInstr &
I)
const {
4295 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
4297 Opcode = SPIRV::OpDemoteToHelperInvocation;
4299 Opcode = SPIRV::OpKill;
4301 if (MachineInstr *NextI =
I.getNextNode()) {
4303 NextI->eraseFromParent();
4313bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
4314 SPIRVTypeInst ResType,
unsigned CmpOpc,
4315 MachineInstr &
I)
const {
4316 Register Cmp0 =
I.getOperand(2).getReg();
4317 Register Cmp1 =
I.getOperand(3).getReg();
4320 "CMP operands should have the same type");
4321 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
4331bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
4332 SPIRVTypeInst ResType,
4333 MachineInstr &
I)
const {
4334 auto Pred =
I.getOperand(1).getPredicate();
4337 Register CmpOperand =
I.getOperand(2).getReg();
4342 Register Op1 =
I.getOperand(3).getReg();
4346 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
4351 I.getOperand(3).setReg(NewOp1);
4357 return selectCmp(ResVReg, ResType, CmpOpc,
I);
4361SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
4362 SPIRVTypeInst ResType)
const {
4364 SPIRVTypeInst SpvI32Ty =
4367 auto ConstInt = ConstantInt::get(LLVMTy, Val);
4374 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4377 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
4380 .
addImm(APInt(32, Val).getZExtValue());
4382 GR.
add(ConstInt,
MI);
4389Register SPIRVInstructionSelector::buildI32ConstantInEntryBlock(
4390 uint32_t Val, MachineInstr &
I, SPIRVTypeInst ResType)
const {
4392 SPIRVTypeInst SpvI32Ty =
4394 auto *ConstInt = ConstantInt::get(LLVMTy, Val);
4399 MachineBasicBlock &EntryBB = *InsertIt->getParent();
4400 MachineInstr *
MI =
nullptr;
4404 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantNull))
4408 uint64_t ImmVal = APInt(32, Val).getZExtValue();
4409 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantI))
4415 GR.
add(ConstInt,
MI);
4420bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
4421 SPIRVTypeInst ResType,
4422 MachineInstr &
I)
const {
4424 return selectCmp(ResVReg, ResType, CmpOp,
I);
4427bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
4428 SPIRVTypeInst ResType,
4429 MachineInstr &
I)
const {
4431 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
4438 if (ResType->
getOpcode() != SPIRV::OpTypeVector &&
4439 ResType->
getOpcode() != SPIRV::OpTypeFloat)
4442 MachineIRBuilder MIRBuilder(
I);
4449 APFloat ConstVal(3.3219280948873623);
4453 APFloat::rmNearestTiesToEven, &LosesInfo);
4457 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
4458 ? SPIRV::OpVectorTimesScalar
4461 if (!selectOpWithSrcs(ArgReg, ResType,
I,
4462 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
4464 if (!selectExtInst(ResVReg, ResType,
I,
4465 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
4475Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
4476 MachineInstr &
I)
const {
4479 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4484bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
4490 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
4498 if (
Def->getOpcode() == SPIRV::OpConstantNull)
4501 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
4502 Def->getOpcode() == SPIRV::OpConstantI)
4515 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
4516 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
4518 Intrinsic::spv_const_composite)) {
4519 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
4520 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
4521 if (!IsZero(
Def->getOperand(i).getReg()))
4530Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
4531 MachineInstr &
I)
const {
4535 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4540Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
4541 MachineInstr &
I)
const {
4545 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4551 SPIRVTypeInst ResType,
4552 MachineInstr &
I)
const {
4556 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4561bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
4562 SPIRVTypeInst ResType,
4563 MachineInstr &
I)
const {
4564 Register SelectFirstArg =
I.getOperand(2).getReg();
4565 Register SelectSecondArg =
I.getOperand(3).getReg();
4574 SPIRV::OpTypeVector;
4581 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
4582 }
else if (IsPtrTy) {
4583 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
4585 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
4588 assert(IsScalarBool &&
"OpSelect with a scalar result requires a scalar "
4589 "boolean condition");
4591 Opcode = SPIRV::OpSelectSFSCond;
4592 }
else if (IsPtrTy) {
4593 Opcode = SPIRV::OpSelectSPSCond;
4595 Opcode = SPIRV::OpSelectSISCond;
4598 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4601 .
addUse(
I.getOperand(1).getReg())
4610bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
4611 SPIRVTypeInst ResType,
4613 MachineInstr &InsertAt,
4614 bool IsSigned)
const {
4616 Register ZeroReg = buildZerosVal(ResType, InsertAt);
4617 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
4618 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
4620 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4632bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
4633 SPIRVTypeInst ResType,
4634 MachineInstr &
I,
bool IsSigned,
4635 unsigned Opcode)
const {
4636 Register SrcReg =
I.getOperand(1).getReg();
4642 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
4647 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I,
false);
4649 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
4652bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
4653 SPIRVTypeInst ResType, MachineInstr &
I,
4654 bool IsSigned)
const {
4655 Register SrcReg =
I.getOperand(1).getReg();
4657 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
4661 if (ResType == SrcType)
4662 return BuildCOPY(ResVReg, SrcReg,
I);
4664 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4665 return selectUnOp(ResVReg, ResType,
I, Opcode);
4668bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
4669 SPIRVTypeInst ResType,
4671 bool IsSigned)
const {
4672 MachineIRBuilder MIRBuilder(
I);
4673 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
4685 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
4688 .
addUse(
I.getOperand(1).getReg())
4689 .
addUse(
I.getOperand(2).getReg())
4694 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
4697 .
addUse(
I.getOperand(1).getReg())
4698 .
addUse(
I.getOperand(2).getReg())
4706 unsigned SelectOpcode =
4707 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
4712 .
addUse(buildOnesVal(
true, ResType,
I))
4713 .
addUse(buildZerosVal(ResType,
I))
4720 .
addUse(buildOnesVal(
false, ResType,
I))
4725bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
4728 SPIRVTypeInst IntTy,
4729 SPIRVTypeInst BoolTy)
const {
4732 bool IsVectorTy = IntTy->
getOpcode() == SPIRV::OpTypeVector;
4733 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
4735 Register One = buildOnesVal(
false, IntTy,
I);
4743 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
4752bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
4753 SPIRVTypeInst ResType,
4754 MachineInstr &
I)
const {
4755 Register IntReg =
I.getOperand(1).getReg();
4758 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
4759 if (ArgType == ResType)
4760 return BuildCOPY(ResVReg, IntReg,
I);
4762 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4763 return selectUnOp(ResVReg, ResType,
I, Opcode);
4766bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
4767 SPIRVTypeInst ResType,
4768 MachineInstr &
I)
const {
4769 unsigned Opcode =
I.getOpcode();
4770 unsigned TpOpcode = ResType->
getOpcode();
4772 if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {
4773 assert(Opcode == TargetOpcode::G_CONSTANT &&
4774 I.getOperand(1).getCImm()->isZero());
4775 MachineBasicBlock &DepMBB =
I.getMF()->front();
4778 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
4785 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
4788bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
4789 SPIRVTypeInst ResType,
4790 MachineInstr &
I)
const {
4791 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4798bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
4799 SPIRVTypeInst ResType,
4800 MachineInstr &
I)
const {
4802 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
4806 .
addUse(
I.getOperand(3).getReg())
4808 .
addUse(
I.getOperand(2).getReg());
4809 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
4815bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
4816 SPIRVTypeInst ResType,
4817 MachineInstr &
I)
const {
4818 Type *MaybeResTy =
nullptr;
4823 "Expected aggregate type for extractv instruction");
4825 SPIRV::AccessQualifier::ReadWrite,
false);
4829 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
4832 .
addUse(
I.getOperand(2).getReg());
4833 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
4839bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
4840 SPIRVTypeInst ResType,
4841 MachineInstr &
I)
const {
4842 if (
getImm(
I.getOperand(4), MRI))
4843 return selectInsertVal(ResVReg, ResType,
I);
4845 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
4848 .
addUse(
I.getOperand(2).getReg())
4849 .
addUse(
I.getOperand(3).getReg())
4850 .
addUse(
I.getOperand(4).getReg())
4855bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
4856 SPIRVTypeInst ResType,
4857 MachineInstr &
I)
const {
4858 if (
getImm(
I.getOperand(3), MRI))
4859 return selectExtractVal(ResVReg, ResType,
I);
4861 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
4864 .
addUse(
I.getOperand(2).getReg())
4865 .
addUse(
I.getOperand(3).getReg())
4870bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
4871 SPIRVTypeInst ResType,
4872 MachineInstr &
I)
const {
4873 const bool IsGEPInBounds =
I.getOperand(2).getImm();
4879 ? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
4880 : SPIRV::OpAccessChain)
4881 : (IsGEPInBounds ?
SPIRV::OpInBoundsPtrAccessChain
4882 :
SPIRV::OpPtrAccessChain);
4884 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4888 .
addUse(
I.getOperand(3).getReg());
4890 (Opcode == SPIRV::OpPtrAccessChain ||
4891 Opcode == SPIRV::OpInBoundsPtrAccessChain ||
4892 (
getImm(
I.getOperand(4), MRI) &&
foldImm(
I.getOperand(4), MRI) == 0)) &&
4893 "Cannot translate GEP to OpAccessChain. First index must be 0.");
4896 const unsigned StartingIndex =
4897 (Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)
4900 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
4901 Res.addUse(
I.getOperand(i).getReg());
4902 Res.constrainAllUses(
TII,
TRI, RBI);
4907bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
4909 unsigned Lim =
I.getNumExplicitOperands();
4910 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
4911 Register OpReg =
I.getOperand(i).getReg();
4912 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
4914 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
4915 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
4916 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
4923 MachineFunction *MF =
I.getMF();
4935 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4936 TII.get(SPIRV::OpSpecConstantOp))
4939 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
4941 GR.
add(OpDefine, MIB);
4947bool SPIRVInstructionSelector::selectDerivativeInst(
4948 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
4949 const unsigned DPdOpCode)
const {
4952 if (!errorIfInstrOutsideShader(
I))
4958 Register SrcReg =
I.getOperand(2).getReg();
4963 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4966 .
addUse(
I.getOperand(2).getReg());
4968 MachineIRBuilder MIRBuilder(
I);
4971 if (componentCount != 1)
4979 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4984 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4989 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4997bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
4998 SPIRVTypeInst ResType,
4999 MachineInstr &
I)
const {
5003 case Intrinsic::spv_load:
5004 return selectLoad(ResVReg, ResType,
I);
5005 case Intrinsic::spv_atomic_load:
5006 return selectAtomicLoad(ResVReg, ResType,
I);
5007 case Intrinsic::spv_store:
5008 return selectStore(
I);
5009 case Intrinsic::spv_atomic_store:
5010 return selectAtomicStore(
I);
5011 case Intrinsic::spv_extractv:
5012 return selectExtractVal(ResVReg, ResType,
I);
5013 case Intrinsic::spv_insertv:
5014 return selectInsertVal(ResVReg, ResType,
I);
5015 case Intrinsic::spv_extractelt:
5016 return selectExtractElt(ResVReg, ResType,
I);
5017 case Intrinsic::spv_insertelt:
5018 return selectInsertElt(ResVReg, ResType,
I);
5019 case Intrinsic::spv_gep:
5020 return selectGEP(ResVReg, ResType,
I);
5021 case Intrinsic::spv_bitcast: {
5022 Register OpReg =
I.getOperand(2).getReg();
5023 SPIRVTypeInst OpType =
5027 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
5029 case Intrinsic::spv_unref_global:
5030 case Intrinsic::spv_init_global: {
5031 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
5036 Register GVarVReg =
MI->getOperand(0).getReg();
5037 if (!selectGlobalValue(GVarVReg, *
MI, Init))
5042 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
5044 MI->eraseFromParent();
5048 case Intrinsic::spv_undef: {
5049 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5055 case Intrinsic::spv_poison:
5056 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpPoisonKHR))
5061 case Intrinsic::spv_freeze:
5062 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
5065 .
addUse(
I.getOperand(2).getReg())
5068 case Intrinsic::spv_named_boolean_spec_constant: {
5069 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
5070 : SPIRV::OpSpecConstantFalse;
5072 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
5073 .
addDef(
I.getOperand(0).getReg())
5076 unsigned SpecId =
I.getOperand(2).getImm();
5078 SPIRV::Decoration::SpecId, {SpecId});
5082 case Intrinsic::spv_const_composite: {
5084 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
5090 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
5092 std::function<bool(
Register)> HasSpecConstOperand =
5102 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
5103 J < Def->getNumExplicitOperands(); ++J) {
5104 if (
Def->getOperand(J).isReg() &&
5105 HasSpecConstOperand(
Def->getOperand(J).getReg()))
5111 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
5112 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
5113 : SPIRV::OpConstantComposite;
5114 unsigned ContinuedOpc = HasSpecConst
5115 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
5116 : SPIRV::OpConstantCompositeContinuedINTEL;
5117 MachineIRBuilder MIR(
I);
5119 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
5121 for (
auto *Instr : Instructions) {
5122 Instr->setDebugLoc(
I.getDebugLoc());
5127 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5134 case Intrinsic::spv_assign_name: {
5135 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
5136 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
5137 for (
unsigned i =
I.getNumExplicitDefs() + 2;
5138 i <
I.getNumExplicitOperands(); ++i) {
5139 MIB.
addImm(
I.getOperand(i).getImm());
5144 case Intrinsic::spv_switch: {
5145 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
5146 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5147 if (
I.getOperand(i).isReg())
5148 MIB.
addReg(
I.getOperand(i).getReg());
5149 else if (
I.getOperand(i).isCImm())
5150 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
5151 else if (
I.getOperand(i).isMBB())
5152 MIB.
addMBB(
I.getOperand(i).getMBB());
5159 case Intrinsic::spv_loop_merge: {
5160 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
5161 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5162 if (
I.getOperand(i).isMBB())
5163 MIB.
addMBB(
I.getOperand(i).getMBB());
5170 case Intrinsic::spv_loop_control_intel: {
5172 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
5173 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
5178 case Intrinsic::spv_selection_merge: {
5180 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
5181 assert(
I.getOperand(1).isMBB() &&
5182 "operand 1 to spv_selection_merge must be a basic block");
5183 MIB.
addMBB(
I.getOperand(1).getMBB());
5184 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
5188 case Intrinsic::spv_cmpxchg:
5189 return selectAtomicCmpXchg(ResVReg, ResType,
I);
5190 case Intrinsic::spv_unreachable:
5191 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
5194 case Intrinsic::spv_abort:
5195 return selectAbort(
I);
5196 case Intrinsic::spv_alloca:
5197 return selectFrameIndex(ResVReg, ResType,
I);
5198 case Intrinsic::spv_alloca_array:
5199 return selectAllocaArray(ResVReg, ResType,
I);
5200 case Intrinsic::spv_assume:
5202 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
5203 .
addUse(
I.getOperand(1).getReg())
5208 case Intrinsic::spv_expect:
5210 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
5213 .
addUse(
I.getOperand(2).getReg())
5214 .
addUse(
I.getOperand(3).getReg())
5219 case Intrinsic::arithmetic_fence:
5220 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
5221 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
5224 .
addUse(
I.getOperand(2).getReg())
5228 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
5230 case Intrinsic::spv_thread_id:
5236 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
5238 case Intrinsic::spv_thread_id_in_group:
5244 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
5246 case Intrinsic::spv_group_id:
5252 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
5254 case Intrinsic::spv_flattened_thread_id_in_group:
5261 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
5263 case Intrinsic::spv_workgroup_size:
5264 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
5266 case Intrinsic::spv_global_size:
5267 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
5269 case Intrinsic::spv_global_offset:
5270 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
5272 case Intrinsic::spv_num_workgroups:
5273 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
5275 case Intrinsic::spv_subgroup_size:
5276 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
5278 case Intrinsic::spv_num_subgroups:
5279 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
5281 case Intrinsic::spv_subgroup_id:
5282 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
5283 case Intrinsic::spv_subgroup_local_invocation_id:
5284 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
5285 ResVReg, ResType,
I);
5286 case Intrinsic::spv_subgroup_max_size:
5287 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
5289 case Intrinsic::spv_fdot:
5290 return selectFloatDot(ResVReg, ResType,
I);
5291 case Intrinsic::spv_udot:
5292 case Intrinsic::spv_sdot:
5293 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5295 return selectIntegerDot(ResVReg, ResType,
I,
5296 IID == Intrinsic::spv_sdot);
5297 return selectIntegerDotExpansion(ResVReg, ResType,
I);
5298 case Intrinsic::spv_dot4add_i8packed:
5299 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5301 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
5302 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
5303 case Intrinsic::spv_dot4add_u8packed:
5304 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5306 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
5307 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
5308 case Intrinsic::spv_all:
5309 return selectAll(ResVReg, ResType,
I);
5310 case Intrinsic::spv_any:
5311 return selectAny(ResVReg, ResType,
I);
5312 case Intrinsic::spv_cross:
5313 return selectExtInst(ResVReg, ResType,
I, CL::cross, GL::Cross);
5314 case Intrinsic::spv_distance:
5315 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
5316 case Intrinsic::spv_lerp:
5317 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
5318 case Intrinsic::spv_length:
5319 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
5320 case Intrinsic::spv_degrees:
5321 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
5322 case Intrinsic::spv_faceforward:
5323 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
5324 case Intrinsic::spv_frac:
5325 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
5326 case Intrinsic::spv_isinf:
5327 return selectOpIsInf(ResVReg, ResType,
I);
5328 case Intrinsic::spv_isnan:
5329 return selectOpIsNan(ResVReg, ResType,
I);
5330 case Intrinsic::spv_isfinite:
5331 return selectOpIsFinite(ResVReg, ResType,
I);
5332 case Intrinsic::spv_isnormal:
5333 return selectOpIsNormal(ResVReg, ResType,
I);
5334 case Intrinsic::spv_normalize:
5335 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
5336 case Intrinsic::spv_refract:
5337 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
5338 case Intrinsic::spv_reflect:
5339 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
5340 case Intrinsic::spv_rsqrt:
5341 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
5342 case Intrinsic::spv_sign:
5343 return selectSign(ResVReg, ResType,
I);
5344 case Intrinsic::spv_smoothstep:
5345 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
5346 case Intrinsic::spv_firstbituhigh:
5347 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
5348 case Intrinsic::spv_firstbitshigh:
5349 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
5350 case Intrinsic::spv_firstbitlow:
5351 return selectFirstBitLow(ResVReg, ResType,
I);
5352 case Intrinsic::spv_all_memory_barrier:
5353 return selectBarrierInst(
I, SPIRV::Scope::Device,
5354 SPIRV::MemorySemantics::UniformMemory |
5355 SPIRV::MemorySemantics::ImageMemory |
5356 SPIRV::MemorySemantics::WorkgroupMemory,
5358 case Intrinsic::spv_all_memory_barrier_with_group_sync:
5359 return selectBarrierInst(
I, SPIRV::Scope::Device,
5360 SPIRV::MemorySemantics::UniformMemory |
5361 SPIRV::MemorySemantics::ImageMemory |
5362 SPIRV::MemorySemantics::WorkgroupMemory,
5364 case Intrinsic::spv_device_memory_barrier:
5365 return selectBarrierInst(
I, SPIRV::Scope::Device,
5366 SPIRV::MemorySemantics::UniformMemory |
5367 SPIRV::MemorySemantics::ImageMemory,
5369 case Intrinsic::spv_device_memory_barrier_with_group_sync:
5370 return selectBarrierInst(
I, SPIRV::Scope::Device,
5371 SPIRV::MemorySemantics::UniformMemory |
5372 SPIRV::MemorySemantics::ImageMemory,
5374 case Intrinsic::spv_group_memory_barrier:
5375 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5376 SPIRV::MemorySemantics::WorkgroupMemory,
5378 case Intrinsic::spv_group_memory_barrier_with_group_sync:
5379 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5380 SPIRV::MemorySemantics::WorkgroupMemory,
5382 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
5383 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
5384 SPIRV::StorageClass::StorageClass ResSC =
5387 return diagnoseUnsupported(
I,
"The target storage class is not castable "
5388 "from the Generic storage class");
5389 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
5397 case Intrinsic::spv_lifetime_start:
5398 case Intrinsic::spv_lifetime_end: {
5399 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
5400 : SPIRV::OpLifetimeStop;
5401 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
5402 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
5411 case Intrinsic::spv_saturate:
5412 return selectSaturate(ResVReg, ResType,
I);
5413 case Intrinsic::spv_nclamp:
5414 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
5415 case Intrinsic::spv_uclamp:
5416 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
5417 case Intrinsic::spv_sclamp:
5418 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
5419 case Intrinsic::spv_subgroup_prefix_bit_count:
5420 return selectWavePrefixBitCount(ResVReg, ResType,
I);
5421 case Intrinsic::spv_wave_active_countbits:
5422 return selectWaveActiveCountBits(ResVReg, ResType,
I);
5423 case Intrinsic::spv_wave_all_equal:
5424 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
5425 case Intrinsic::spv_wave_all:
5426 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
5427 case Intrinsic::spv_wave_any:
5428 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
5429 case Intrinsic::spv_subgroup_ballot:
5430 return selectWaveOpInst(ResVReg, ResType,
I,
5431 SPIRV::OpGroupNonUniformBallot);
5432 case Intrinsic::spv_wave_is_first_lane:
5433 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
5434 case Intrinsic::spv_wave_reduce_or:
5435 return selectWaveReduceOp(ResVReg, ResType,
I,
5436 SPIRV::OpGroupNonUniformBitwiseOr);
5437 case Intrinsic::spv_wave_reduce_xor:
5438 return selectWaveReduceOp(ResVReg, ResType,
I,
5439 SPIRV::OpGroupNonUniformBitwiseXor);
5440 case Intrinsic::spv_wave_reduce_and:
5441 return selectWaveReduceOp(ResVReg, ResType,
I,
5442 SPIRV::OpGroupNonUniformBitwiseAnd);
5443 case Intrinsic::spv_interlocked_add:
5444 return selectInterlockedOp(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
5445 case Intrinsic::spv_interlocked_or:
5446 return selectInterlockedOp(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
5447 case Intrinsic::spv_wave_reduce_umax:
5448 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
5449 case Intrinsic::spv_wave_reduce_max:
5450 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
5451 case Intrinsic::spv_wave_reduce_umin:
5452 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
5453 case Intrinsic::spv_wave_reduce_min:
5454 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
5455 case Intrinsic::spv_wave_reduce_sum:
5456 return selectWaveReduceSum(ResVReg, ResType,
I);
5457 case Intrinsic::spv_wave_product:
5458 return selectWaveReduceProduct(ResVReg, ResType,
I);
5459 case Intrinsic::spv_wave_readlane:
5460 return selectWaveOpInst(ResVReg, ResType,
I,
5461 SPIRV::OpGroupNonUniformShuffle);
5462 case Intrinsic::spv_wave_prefix_sum:
5463 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
5464 case Intrinsic::spv_wave_prefix_product:
5465 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
5466 case Intrinsic::spv_quad_read_across_x: {
5467 return selectQuadSwap(ResVReg, ResType,
I, 0);
5469 case Intrinsic::spv_quad_read_across_y: {
5470 return selectQuadSwap(ResVReg, ResType,
I, 1);
5472 case Intrinsic::spv_quad_read_across_diagonal: {
5473 return selectQuadSwap(ResVReg, ResType,
I, 2);
5475 case Intrinsic::spv_step:
5476 return selectExtInst(ResVReg, ResType,
I, CL::step, GL::Step);
5477 case Intrinsic::spv_radians:
5478 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
5482 case Intrinsic::instrprof_increment:
5483 case Intrinsic::instrprof_increment_step:
5484 case Intrinsic::instrprof_value_profile:
5487 case Intrinsic::spv_value_md:
5489 case Intrinsic::spv_resource_handlefrombinding: {
5490 return selectHandleFromBinding(ResVReg, ResType,
I);
5492 case Intrinsic::spv_resource_counterhandlefrombinding:
5493 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
5494 case Intrinsic::spv_resource_updatecounter:
5495 return selectUpdateCounter(ResVReg, ResType,
I);
5496 case Intrinsic::spv_resource_store_typedbuffer: {
5497 return selectImageWriteIntrinsic(
I);
5499 case Intrinsic::spv_resource_load_typedbuffer: {
5500 return selectReadImageIntrinsic(ResVReg, ResType,
I);
5502 case Intrinsic::spv_resource_load_level: {
5503 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
5505 case Intrinsic::spv_resource_getdimensions_x:
5506 case Intrinsic::spv_resource_getdimensions_xy:
5507 case Intrinsic::spv_resource_getdimensions_xyz: {
5508 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
5510 case Intrinsic::spv_resource_getdimensions_levels_x:
5511 case Intrinsic::spv_resource_getdimensions_levels_xy:
5512 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
5513 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
5515 case Intrinsic::spv_resource_getdimensions_ms_xy:
5516 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
5517 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
5519 case Intrinsic::spv_resource_calculate_lod:
5520 case Intrinsic::spv_resource_calculate_lod_unclamped:
5521 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
5522 case Intrinsic::spv_resource_sample:
5523 case Intrinsic::spv_resource_sample_clamp:
5524 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
5525 case Intrinsic::spv_resource_samplebias:
5526 case Intrinsic::spv_resource_samplebias_clamp:
5527 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
5528 case Intrinsic::spv_resource_samplegrad:
5529 case Intrinsic::spv_resource_samplegrad_clamp:
5530 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
5531 case Intrinsic::spv_resource_samplelevel:
5532 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
5533 case Intrinsic::spv_resource_samplecmp:
5534 case Intrinsic::spv_resource_samplecmp_clamp:
5535 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
5536 case Intrinsic::spv_resource_samplecmplevelzero:
5537 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
5538 case Intrinsic::spv_resource_gather:
5539 case Intrinsic::spv_resource_gather_cmp:
5540 return selectGatherIntrinsic(ResVReg, ResType,
I);
5541 case Intrinsic::spv_resource_getbasepointer:
5542 case Intrinsic::spv_resource_getpointer: {
5543 return selectResourceGetPointer(ResVReg, ResType,
I);
5545 case Intrinsic::spv_pushconstant_getpointer: {
5546 return selectPushConstantGetPointer(ResVReg, ResType,
I);
5548 case Intrinsic::spv_discard: {
5549 return selectDiscard(ResVReg, ResType,
I);
5551 case Intrinsic::spv_resource_nonuniformindex: {
5552 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
5554 case Intrinsic::spv_unpackhalf2x16: {
5555 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
5557 case Intrinsic::spv_packhalf2x16: {
5558 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
5560 case Intrinsic::spv_ddx:
5561 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
5562 case Intrinsic::spv_ddy:
5563 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
5564 case Intrinsic::spv_ddx_coarse:
5565 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
5566 case Intrinsic::spv_ddy_coarse:
5567 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
5568 case Intrinsic::spv_ddx_fine:
5569 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
5570 case Intrinsic::spv_ddy_fine:
5571 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
5572 case Intrinsic::spv_fwidth:
5573 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
5574 case Intrinsic::spv_masked_gather:
5575 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5576 return selectMaskedGather(ResVReg, ResType,
I);
5577 return diagnoseUnsupported(
5578 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
5579 case Intrinsic::spv_masked_scatter:
5580 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5581 return selectMaskedScatter(
I);
5582 return diagnoseUnsupported(
5583 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
5584 case Intrinsic::returnaddress:
5585 case Intrinsic::frameaddress: {
5587 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5594 return diagnoseUnsupported(
I,
"intrinsic selection not implemented.");
5599bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
5600 SPIRVTypeInst ResType,
5601 MachineInstr &
I)
const {
5604 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
5611bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
5612 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5614 assert(Intr.getIntrinsicID() ==
5615 Intrinsic::spv_resource_counterhandlefrombinding);
5618 Register MainHandleReg = Intr.getOperand(2).getReg();
5620 assert(MainHandleDef->getIntrinsicID() ==
5621 Intrinsic::spv_resource_handlefrombinding);
5625 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
5626 Register IndexReg = MainHandleDef->getOperand(5).getReg();
5627 std::string CounterName =
5632 MachineIRBuilder MIRBuilder(
I);
5634 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
5636 ArraySize, IndexReg, CounterName, MIRBuilder);
5638 return BuildCOPY(ResVReg, CounterVarReg,
I);
5641bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
5642 SPIRVTypeInst ResType,
5643 MachineInstr &
I)
const {
5645 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
5647 Register CounterHandleReg = Intr.getOperand(2).getReg();
5648 Register IncrReg = Intr.getOperand(3).getReg();
5655 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
5656 assert(CounterVarPointeeType &&
5657 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
5658 "Counter variable must be a struct");
5660 SPIRV::StorageClass::StorageBuffer &&
5661 "Counter variable must be in the storage buffer storage class");
5663 "Counter variable must have exactly 1 member in the struct");
5664 const SPIRVTypeInst MemberType =
5667 "Counter variable struct must have a single i32 member");
5671 MachineIRBuilder MIRBuilder(
I);
5673 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
5676 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
5682 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5685 .
addUse(CounterHandleReg)
5692 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
5695 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
5698 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
5707 return BuildCOPY(ResVReg, AtomicRes,
I);
5715 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
5723bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
5724 SPIRVTypeInst ResType,
5725 MachineInstr &
I)
const {
5733 Register ImageReg =
I.getOperand(2).getReg();
5741 Register IdxReg =
I.getOperand(3).getReg();
5743 MachineInstr &Pos =
I;
5745 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
5749bool SPIRVInstructionSelector::generateSampleImage(
5752 DebugLoc Loc, MachineInstr &Pos)
const {
5763 if (!loadHandleBeforePosition(NewSamplerReg,
5769 MachineIRBuilder MIRBuilder(Pos);
5782 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
5783 ImOps.Lod.has_value();
5784 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
5785 : SPIRV::OpImageSampleImplicitLod;
5787 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
5788 : SPIRV::OpImageSampleDrefImplicitLod;
5797 MIB.
addUse(*ImOps.Compare);
5799 uint32_t ImageOperands = 0;
5801 ImageOperands |= SPIRV::ImageOperand::Bias;
5803 ImageOperands |= SPIRV::ImageOperand::Lod;
5804 if (ImOps.GradX && ImOps.GradY)
5805 ImageOperands |= SPIRV::ImageOperand::Grad;
5806 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
5808 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
5811 "Non-constant offsets are not supported in sample instructions.");
5816 ImageOperands |= SPIRV::ImageOperand::MinLod;
5818 if (ImageOperands != 0) {
5819 MIB.
addImm(ImageOperands);
5820 if (ImageOperands & SPIRV::ImageOperand::Bias)
5822 if (ImageOperands & SPIRV::ImageOperand::Lod)
5824 if (ImageOperands & SPIRV::ImageOperand::Grad) {
5825 MIB.
addUse(*ImOps.GradX);
5826 MIB.
addUse(*ImOps.GradY);
5829 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
5830 MIB.
addUse(*ImOps.Offset);
5831 if (ImageOperands & SPIRV::ImageOperand::MinLod)
5832 MIB.
addUse(*ImOps.MinLod);
5839bool SPIRVInstructionSelector::selectImageQuerySize(
5841 std::optional<Register> LodReg)
const {
5843 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
5846 "ImageReg is not an image type.");
5848 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
5850 unsigned NumComponents = 0;
5852 case SPIRV::Dim::DIM_1D:
5853 case SPIRV::Dim::DIM_Buffer:
5854 NumComponents =
IsArray ? 2 : 1;
5856 case SPIRV::Dim::DIM_2D:
5857 case SPIRV::Dim::DIM_Cube:
5858 case SPIRV::Dim::DIM_Rect:
5859 NumComponents =
IsArray ? 3 : 2;
5861 case SPIRV::Dim::DIM_3D:
5865 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
5870 SPIRVTypeInst ResType =
5875 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5885bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
5886 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5887 Register ImageReg =
I.getOperand(2).getReg();
5894 return selectImageQuerySize(NewImageReg, ResVReg,
I);
5897bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
5898 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5899 Register ImageReg =
I.getOperand(2).getReg();
5908 Register LodReg =
I.getOperand(3).getReg();
5911 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
5913 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
5920 TII.get(SPIRV::OpImageQueryLevels))
5927 TII.get(SPIRV::OpCompositeConstruct))
5937bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
5938 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5939 Register ImageReg =
I.getOperand(2).getReg();
5950 "OpImageQuerySamples requires a multisampled image");
5952 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
5960 TII.get(SPIRV::OpImageQuerySamples))
5967 TII.get(SPIRV::OpCompositeConstruct))
5977bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
5978 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5979 Register ImageReg =
I.getOperand(2).getReg();
5980 Register SamplerReg =
I.getOperand(3).getReg();
5981 Register CoordinateReg =
I.getOperand(4).getReg();
5997 if (!loadHandleBeforePosition(
6002 MachineIRBuilder MIRBuilder(
I);
6008 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6018 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
6025 unsigned ExtractedIndex =
6027 Intrinsic::spv_resource_calculate_lod_unclamped
6031 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6032 TII.get(SPIRV::OpCompositeExtract))
6042bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
6043 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6044 Register ImageReg =
I.getOperand(2).getReg();
6045 Register SamplerReg =
I.getOperand(3).getReg();
6046 Register CoordinateReg =
I.getOperand(4).getReg();
6047 ImageOperands ImOps;
6048 if (
I.getNumOperands() > 5)
6049 ImOps.Offset =
I.getOperand(5).getReg();
6050 if (
I.getNumOperands() > 6)
6051 ImOps.MinLod =
I.getOperand(6).getReg();
6052 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6053 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6056bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
6057 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6058 Register ImageReg =
I.getOperand(2).getReg();
6059 Register SamplerReg =
I.getOperand(3).getReg();
6060 Register CoordinateReg =
I.getOperand(4).getReg();
6061 ImageOperands ImOps;
6062 ImOps.Bias =
I.getOperand(5).getReg();
6063 if (
I.getNumOperands() > 6)
6064 ImOps.Offset =
I.getOperand(6).getReg();
6065 if (
I.getNumOperands() > 7)
6066 ImOps.MinLod =
I.getOperand(7).getReg();
6067 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6068 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6071bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
6072 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6073 Register ImageReg =
I.getOperand(2).getReg();
6074 Register SamplerReg =
I.getOperand(3).getReg();
6075 Register CoordinateReg =
I.getOperand(4).getReg();
6076 ImageOperands ImOps;
6077 ImOps.GradX =
I.getOperand(5).getReg();
6078 ImOps.GradY =
I.getOperand(6).getReg();
6079 if (
I.getNumOperands() > 7)
6080 ImOps.Offset =
I.getOperand(7).getReg();
6081 if (
I.getNumOperands() > 8)
6082 ImOps.MinLod =
I.getOperand(8).getReg();
6083 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6084 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6087bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
6088 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6089 Register ImageReg =
I.getOperand(2).getReg();
6090 Register SamplerReg =
I.getOperand(3).getReg();
6091 Register CoordinateReg =
I.getOperand(4).getReg();
6092 ImageOperands ImOps;
6093 ImOps.Lod =
I.getOperand(5).getReg();
6094 if (
I.getNumOperands() > 6)
6095 ImOps.Offset =
I.getOperand(6).getReg();
6096 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6097 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6100bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
6101 SPIRVTypeInst ResType,
6102 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.Compare =
I.getOperand(5).getReg();
6108 if (
I.getNumOperands() > 6)
6109 ImOps.Offset =
I.getOperand(6).getReg();
6110 if (
I.getNumOperands() > 7)
6111 ImOps.MinLod =
I.getOperand(7).getReg();
6112 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6113 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6116bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
6117 SPIRVTypeInst ResType,
6118 MachineInstr &
I)
const {
6119 Register ImageReg =
I.getOperand(2).getReg();
6120 Register CoordinateReg =
I.getOperand(3).getReg();
6121 Register LodReg =
I.getOperand(4).getReg();
6123 ImageOperands ImOps;
6125 if (
I.getNumOperands() > 5)
6126 ImOps.Offset =
I.getOperand(5).getReg();
6138 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
6139 I.getDebugLoc(),
I, &ImOps);
6142bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
6143 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6144 Register ImageReg =
I.getOperand(2).getReg();
6145 Register SamplerReg =
I.getOperand(3).getReg();
6146 Register CoordinateReg =
I.getOperand(4).getReg();
6147 ImageOperands ImOps;
6148 ImOps.Compare =
I.getOperand(5).getReg();
6149 if (
I.getNumOperands() > 6)
6150 ImOps.Offset =
I.getOperand(6).getReg();
6153 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6154 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6157bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
6158 SPIRVTypeInst ResType,
6159 MachineInstr &
I)
const {
6160 Register ImageReg =
I.getOperand(2).getReg();
6161 Register SamplerReg =
I.getOperand(3).getReg();
6162 Register CoordinateReg =
I.getOperand(4).getReg();
6165 "ImageReg is not an image type.");
6170 ComponentOrCompareReg =
I.getOperand(5).getReg();
6171 OffsetReg =
I.getOperand(6).getReg();
6174 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
6178 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6179 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
6180 Dim != SPIRV::Dim::DIM_Rect) {
6182 "Gather operations are only supported for 2D, Cube, and Rect images.");
6189 if (!loadHandleBeforePosition(
6194 MachineIRBuilder MIRBuilder(
I);
6195 SPIRVTypeInst SampledImageType =
6200 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6208 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
6210 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
6212 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6217 .
addUse(ComponentOrCompareReg);
6219 uint32_t ImageOperands = 0;
6220 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
6221 if (Dim == SPIRV::Dim::DIM_Cube) {
6223 "Gather operations with offset are not supported for Cube images.");
6227 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6229 ImageOperands |= SPIRV::ImageOperand::Offset;
6233 if (ImageOperands != 0) {
6234 MIB.
addImm(ImageOperands);
6236 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6244bool SPIRVInstructionSelector::generateImageReadOrFetch(
6247 const ImageOperands *ImOps)
const {
6250 "ImageReg is not an image type.");
6252 bool IsSignedInteger =
6257 bool IsFetch = (SampledOp.getImm() == 1);
6259 auto AddOperands = [&](MachineInstrBuilder &MIB) {
6260 uint32_t ImageOperandsMask = 0;
6261 if (IsSignedInteger)
6262 ImageOperandsMask |= 0x1000;
6264 if (IsFetch && ImOps) {
6266 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
6267 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
6269 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
6271 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
6275 if (ImageOperandsMask != 0) {
6276 MIB.
addImm(ImageOperandsMask);
6277 if (IsFetch && ImOps) {
6280 if (ImOps->Offset &&
6281 (ImageOperandsMask &
6282 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
6283 MIB.
addUse(*ImOps->Offset);
6289 if (ResultSize == 4) {
6292 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
6299 BMI.constrainAllUses(
TII,
TRI, RBI);
6303 SPIRVTypeInst ReadType = widenTypeToVec4(ResType, Pos);
6307 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
6313 BMI.constrainAllUses(
TII,
TRI, RBI);
6315 if (ResultSize == 1) {
6324 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
6327bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
6328 SPIRVTypeInst ResType,
6329 MachineInstr &
I)
const {
6330 Register ResourcePtr =
I.getOperand(2).getReg();
6332 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
6341 MachineIRBuilder MIRBuilder(
I);
6346 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6352 if (
I.getNumExplicitOperands() > 3) {
6353 Register IndexReg =
I.getOperand(3).getReg();
6360bool SPIRVInstructionSelector::selectPushConstantGetPointer(
6361 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6366bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
6367 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6368 Register ObjReg =
I.getOperand(2).getReg();
6369 if (!BuildCOPY(ResVReg, ObjReg,
I))
6379 decorateUsesAsNonUniform(ResVReg);
6383void SPIRVInstructionSelector::decorateUsesAsNonUniform(
6386 while (WorkList.
size() > 0) {
6390 bool IsDecorated =
false;
6392 if (
Use.getOpcode() == SPIRV::OpDecorate &&
6393 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
6399 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
6401 if (ResultReg == CurrentReg)
6409 SPIRV::Decoration::NonUniformEXT, {});
6414bool SPIRVInstructionSelector::extractSubvector(
6416 MachineInstr &InsertionPoint)
const {
6418 [[maybe_unused]] uint64_t InputSize =
6421 assert(InputSize > 1 &&
"The input must be a vector.");
6422 assert(ResultSize > 1 &&
"The result must be a vector.");
6423 assert(ResultSize < InputSize &&
6424 "Cannot extract more element than there are in the input.");
6428 for (uint64_t
I = 0;
I < ResultSize;
I++) {
6431 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6440 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
6442 TII.get(SPIRV::OpCompositeConstruct))
6446 for (
Register ComponentReg : ComponentRegisters)
6447 MIB.
addUse(ComponentReg);
6452bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
6453 MachineInstr &
I)
const {
6460 Register ImageReg =
I.getOperand(1).getReg();
6468 Register CoordinateReg =
I.getOperand(2).getReg();
6469 Register DataReg =
I.getOperand(3).getReg();
6472 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
6480Register SPIRVInstructionSelector::buildPointerToResource(
6481 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
6482 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
6483 StringRef Name, MachineIRBuilder MIRBuilder)
const {
6485 if (ArraySize == 1) {
6486 SPIRVTypeInst PtrType =
6489 "SpirvResType did not have an explicit layout.");
6494 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
6495 SPIRVTypeInst VarPointerType =
6498 VarPointerType, Set,
Binding, Name, MIRBuilder);
6500 SPIRVTypeInst ResPointerType =
6513bool SPIRVInstructionSelector::selectFirstBitSet16(
6514 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
6515 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
6517 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
6521 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
6524bool SPIRVInstructionSelector::selectFirstBitSet32(
6526 unsigned BitSetOpcode)
const {
6527 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6530 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6537bool SPIRVInstructionSelector::selectFirstBitSet64(
6539 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
6552 if (ComponentCount > 2) {
6553 auto Func = [
this, SwapPrimarySide](
Register ResVReg, SPIRVTypeInst ResType,
6555 unsigned Opcode) ->
bool {
6556 return this->selectFirstBitSet64(ResVReg, ResType,
I, SrcReg, Opcode,
6560 return handle64BitOverflow(ResVReg, ResType,
I, SrcReg, BitSetOpcode, Func);
6564 MachineIRBuilder MIRBuilder(
I);
6566 BaseType, 2 * ComponentCount, MIRBuilder,
false);
6570 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
6576 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
6583 bool IsScalarRes = ResType->
getOpcode() != SPIRV::OpTypeVector;
6586 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
6587 SPIRV::OpVectorExtractDynamic))
6589 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
6590 SPIRV::OpVectorExtractDynamic))
6594 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6595 TII.get(SPIRV::OpVectorShuffle))
6603 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
6609 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6610 TII.get(SPIRV::OpVectorShuffle))
6618 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
6638 SelectOp = SPIRV::OpSelectSISCond;
6639 AddOp = SPIRV::OpIAddS;
6647 SelectOp = SPIRV::OpSelectVIVCond;
6648 AddOp = SPIRV::OpIAddV;
6654 Register RegSecondaryOffset = Reg0;
6658 if (SwapPrimarySide) {
6659 PrimaryReg = LowReg;
6660 SecondaryReg = HighReg;
6661 RegPrimaryOffset = Reg0;
6662 RegSecondaryOffset = Reg32;
6667 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
6668 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
6673 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
6674 SPIRV::OpINotEqual))
6681 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
6682 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
6687 if (SwapPrimarySide) {
6689 if (!selectOpWithSrcs(RegAdd, ResType,
I,
6690 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
6701 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
6702 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
6707 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
6708 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
6711 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
6715bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
6716 SPIRVTypeInst ResType,
6718 bool IsSigned)
const {
6720 Register OpReg =
I.getOperand(2).getReg();
6723 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
6724 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
6728 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6730 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6732 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6735 return diagnoseUnsupported(
6737 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
6741bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
6742 SPIRVTypeInst ResType,
6743 MachineInstr &
I)
const {
6745 Register OpReg =
I.getOperand(2).getReg();
6750 unsigned ExtendOpcode = SPIRV::OpUConvert;
6751 unsigned BitSetOpcode = GL::FindILsb;
6755 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6757 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6759 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6762 return diagnoseUnsupported(
I,
6763 "spv_firstbitlow only supports 16,32,64 bits.");
6767bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
6768 SPIRVTypeInst ResType,
6769 MachineInstr &
I)
const {
6773 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
6776 .
addUse(
I.getOperand(2).getReg())
6779 unsigned Alignment =
I.getOperand(3).getImm();
6793 while (!Worklist.
empty()) {
6795 switch (
T->getOpcode()) {
6796 case SPIRV::OpTypeInt:
6797 case SPIRV::OpTypeFloat:
6798 case SPIRV::OpTypePointer:
6800 case SPIRV::OpTypeVector:
6801 case SPIRV::OpTypeMatrix:
6802 case SPIRV::OpTypeArray: {
6803 Register OperandReg =
T->getOperand(1).getReg();
6807 case SPIRV::OpTypeStruct:
6808 for (
unsigned Idx = 1,
E =
T->getNumOperands(); Idx <
E; ++Idx) {
6809 Register OperandReg =
T->getOperand(Idx).getReg();
6821bool SPIRVInstructionSelector::selectAbort(MachineInstr &
I)
const {
6822 assert(
I.getNumExplicitOperands() == 2);
6824 Register MsgReg =
I.getOperand(1).getReg();
6826 assert(MsgType &&
"Message argument of llvm.spv.abort has no SPIR-V type");
6829 return diagnoseUnsupported(
6831 "llvm.spv.abort message type must be a concrete SPIR-V type (numerical "
6832 "scalar, pointer, vector, matrix, or aggregate of such types)");
6835 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
6842bool SPIRVInstructionSelector::selectTrap(MachineInstr &
I)
const {
6851 uint32_t MsgVal = ~0
u;
6852 if (
I.getOpcode() == TargetOpcode::G_UBSANTRAP)
6853 MsgVal =
static_cast<uint32_t
>(
I.getOperand(0).
getImm());
6856 Register MsgReg = buildI32ConstantInEntryBlock(MsgVal,
I, MsgType);
6859 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
6866bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
6867 SPIRVTypeInst ResType,
6868 MachineInstr &
I)
const {
6872 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
6875 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
6878 unsigned Alignment =
I.getOperand(2).getImm();
6885bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
6890 const MachineInstr *PrevI =
I.getPrevNode();
6892 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
6896 .
addMBB(
I.getOperand(0).getMBB())
6901 .
addMBB(
I.getOperand(0).getMBB())
6906bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
6917 const MachineInstr *NextI =
I.getNextNode();
6919 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
6925 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
6927 .
addUse(
I.getOperand(0).getReg())
6928 .
addMBB(
I.getOperand(1).getMBB())
6934bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
6935 MachineInstr &
I)
const {
6937 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
6939 const unsigned NumOps =
I.getNumOperands();
6940 for (
unsigned i = 1; i <
NumOps; i += 2) {
6941 MIB.
addUse(
I.getOperand(i + 0).getReg());
6942 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
6948bool SPIRVInstructionSelector::selectGlobalValue(
6949 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
6951 MachineIRBuilder MIRBuilder(
I);
6952 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
6955 std::string GlobalIdent;
6957 unsigned &
ID = UnnamedGlobalIDs[GV];
6959 ID = UnnamedGlobalIDs.
size();
6960 GlobalIdent =
"__unnamed_" + Twine(
ID).str();
6986 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
6993 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
6998 MachineInstrBuilder MIB1 =
6999 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7002 MachineInstrBuilder MIB2 =
7004 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
7008 GR.
add(ConstVal, MIB2);
7016 MachineInstrBuilder MIB3 =
7017 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7020 GR.
add(ConstVal, MIB3);
7026 assert(NewReg != ResVReg);
7027 return BuildCOPY(ResVReg, NewReg,
I);
7037 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
7040 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
7046 SPIRVTypeInst ResType =
7050 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
7055 if (
GlobalVar->isExternallyInitialized() &&
7056 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
7057 constexpr unsigned ReadWriteINTEL = 3u;
7060 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
7066bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
7067 SPIRVTypeInst ResType,
7068 MachineInstr &
I)
const {
7070 return selectExtInst(ResVReg, ResType,
I, CL::log10);
7078 MachineIRBuilder MIRBuilder(
I);
7083 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7086 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
7088 .
add(
I.getOperand(1))
7093 ResType->
getOpcode() == SPIRV::OpTypeFloat);
7103 APFloat::rmNearestTiesToEven, &LosesInfo);
7107 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
7108 ? SPIRV::OpVectorTimesScalar
7119bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
7120 SPIRVTypeInst ResType,
7121 MachineInstr &
I)
const {
7124 return selectExtInst(ResVReg, ResType,
I, CL::pown);
7130 Register ExpReg =
I.getOperand(2).getReg();
7132 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
7133 SPIRV::OpConvertSToF))
7135 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
7142bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
7143 SPIRVTypeInst ResType,
7144 MachineInstr &
I)
const {
7160 MachineIRBuilder MIRBuilder(
I);
7161 SPIRVTypeInst FloatType =
7165 FloatType, MIRBuilder, SPIRV::StorageClass::Function);
7178 MachineBasicBlock &EntryBB =
I.getMF()->
front();
7180 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
7183 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7189 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7192 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
7195 .
add(
I.getOperand(
I.getNumExplicitDefs()))
7199 Register IntegralPartReg =
I.getOperand(1).getReg();
7202 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7212 assert(
false &&
"GLSL::Modf is deprecated.");
7223bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
7224 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7225 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7226 MachineIRBuilder MIRBuilder(
I);
7227 const SPIRVTypeInst Vec3Ty =
7230 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
7242 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7246 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7252 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7259 assert(
I.getOperand(2).isReg());
7260 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
7264 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
7275bool SPIRVInstructionSelector::loadBuiltinInputID(
7276 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7277 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7278 MachineIRBuilder MIRBuilder(
I);
7280 ResType, MIRBuilder, SPIRV::StorageClass::Input);
7295 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7299 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7308SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
7309 MachineInstr &
I)
const {
7310 MachineIRBuilder MIRBuilder(
I);
7311 if (
Type->getOpcode() != SPIRV::OpTypeVector)
7321bool SPIRVInstructionSelector::loadHandleBeforePosition(
7322 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
7323 MachineInstr &Pos)
const {
7326 Intrinsic::spv_resource_handlefrombinding);
7334 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
7335 MachineIRBuilder MIRBuilder(HandleDef);
7336 SPIRVTypeInst VarType = ResType;
7337 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
7339 if (IsStructuredBuffer) {
7344 if (ResType->
getOpcode() == SPIRV::OpTypeImage && ArraySize == 0)
7346 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
7349 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
7350 ArraySize, IndexReg, Name, MIRBuilder);
7354 uint32_t LoadOpcode =
7355 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
7365bool SPIRVInstructionSelector::errorIfInstrOutsideShader(
7366 MachineInstr &
I)
const {
7368 return diagnoseUnsupported(
7369 I,
"this instruction is only supported in shaders.");
7374InstructionSelector *
7378 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)
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 & 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
@ 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.
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