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 return selectMemOperation(ResVReg,
I);
1047 case TargetOpcode::G_ICMP:
1048 return selectICmp(ResVReg, ResType,
I);
1049 case TargetOpcode::G_FCMP:
1050 return selectFCmp(ResVReg, ResType,
I);
1052 case TargetOpcode::G_FRAME_INDEX:
1053 return selectFrameIndex(ResVReg, ResType,
I);
1055 case TargetOpcode::G_LOAD:
1056 return selectLoad(ResVReg, ResType,
I);
1057 case TargetOpcode::G_STORE:
1058 return selectStore(
I);
1060 case TargetOpcode::G_BR:
1061 return selectBranch(
I);
1062 case TargetOpcode::G_BRCOND:
1063 return selectBranchCond(
I);
1065 case TargetOpcode::G_PHI:
1066 return selectPhi(ResVReg,
I);
1068 case TargetOpcode::G_FPTOSI:
1069 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1070 case TargetOpcode::G_FPTOUI:
1071 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1073 case TargetOpcode::G_FPTOSI_SAT:
1074 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1075 case TargetOpcode::G_FPTOUI_SAT:
1076 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1078 case TargetOpcode::G_SITOFP:
1079 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
1080 case TargetOpcode::G_UITOFP:
1081 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
1083 case TargetOpcode::G_CTPOP:
1084 return selectPopCount(ResVReg, ResType,
I, SPIRV::OpBitCount);
1085 case TargetOpcode::G_SMIN:
1086 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
1087 case TargetOpcode::G_UMIN:
1088 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
1090 case TargetOpcode::G_SMAX:
1091 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
1092 case TargetOpcode::G_UMAX:
1093 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
1095 case TargetOpcode::G_SCMP:
1096 return selectSUCmp(ResVReg, ResType,
I,
true);
1097 case TargetOpcode::G_UCMP:
1098 return selectSUCmp(ResVReg, ResType,
I,
false);
1099 case TargetOpcode::G_LROUND:
1100 case TargetOpcode::G_LLROUND: {
1103 MRI->
setRegClass(regForLround, &SPIRV::iIDRegClass);
1105 regForLround, *(
I.getParent()->getParent()));
1107 CL::round, GL::Round,
false);
1109 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1116 case TargetOpcode::G_STRICT_FMA:
1117 case TargetOpcode::G_FMA: {
1120 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1123 .
addUse(
I.getOperand(1).getReg())
1124 .
addUse(
I.getOperand(2).getReg())
1125 .
addUse(
I.getOperand(3).getReg())
1130 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1133 case TargetOpcode::G_STRICT_FLDEXP:
1134 return selectExtInst(ResVReg, ResType,
I, CL::ldexp);
1136 case TargetOpcode::G_FPOW:
1137 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1138 case TargetOpcode::G_FPOWI:
1139 return selectFpowi(ResVReg, ResType,
I);
1141 case TargetOpcode::G_FEXP:
1142 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1143 case TargetOpcode::G_FEXP2:
1144 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1145 case TargetOpcode::G_FEXP10:
1146 return selectExp10(ResVReg, ResType,
I);
1148 case TargetOpcode::G_FMODF:
1149 return selectModf(ResVReg, ResType,
I);
1150 case TargetOpcode::G_FSINCOS:
1151 return selectSincos(ResVReg, ResType,
I);
1153 case TargetOpcode::G_FLOG:
1154 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1155 case TargetOpcode::G_FLOG2:
1156 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1157 case TargetOpcode::G_FLOG10:
1158 return selectLog10(ResVReg, ResType,
I);
1160 case TargetOpcode::G_FABS:
1161 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1162 case TargetOpcode::G_ABS:
1163 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1165 case TargetOpcode::G_FMINNUM:
1166 case TargetOpcode::G_FMINIMUM:
1167 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1168 case TargetOpcode::G_FMAXNUM:
1169 case TargetOpcode::G_FMAXIMUM:
1170 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1172 case TargetOpcode::G_FCOPYSIGN:
1173 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1175 case TargetOpcode::G_FCEIL:
1176 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1177 case TargetOpcode::G_FFLOOR:
1178 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1180 case TargetOpcode::G_FCOS:
1181 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1182 case TargetOpcode::G_FSIN:
1183 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1184 case TargetOpcode::G_FTAN:
1185 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1186 case TargetOpcode::G_FACOS:
1187 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1188 case TargetOpcode::G_FASIN:
1189 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1190 case TargetOpcode::G_FATAN:
1191 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1192 case TargetOpcode::G_FATAN2:
1193 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1194 case TargetOpcode::G_FCOSH:
1195 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1196 case TargetOpcode::G_FSINH:
1197 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1198 case TargetOpcode::G_FTANH:
1199 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1201 case TargetOpcode::G_STRICT_FSQRT:
1202 case TargetOpcode::G_FSQRT:
1203 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1205 case TargetOpcode::G_CTTZ:
1206 case TargetOpcode::G_CTTZ_ZERO_POISON:
1207 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1208 case TargetOpcode::G_CTLZ:
1209 case TargetOpcode::G_CTLZ_ZERO_POISON:
1210 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1212 case TargetOpcode::G_INTRINSIC_ROUND:
1213 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1214 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1215 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1216 case TargetOpcode::G_INTRINSIC_TRUNC:
1217 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1218 case TargetOpcode::G_FRINT:
1219 case TargetOpcode::G_FNEARBYINT:
1220 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1222 case TargetOpcode::G_SMULH:
1223 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1224 case TargetOpcode::G_UMULH:
1225 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1227 case TargetOpcode::G_SADDSAT:
1228 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1229 case TargetOpcode::G_UADDSAT:
1230 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1231 case TargetOpcode::G_SSUBSAT:
1232 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1233 case TargetOpcode::G_USUBSAT:
1234 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1236 case TargetOpcode::G_FFREXP:
1237 return selectFrexp(ResVReg, ResType,
I);
1239 case TargetOpcode::G_UADDO:
1240 return selectOverflowArith(ResVReg, ResType,
I,
1241 ResType->
getOpcode() == SPIRV::OpTypeVector
1242 ? SPIRV::OpIAddCarryV
1243 : SPIRV::OpIAddCarryS);
1244 case TargetOpcode::G_USUBO:
1245 return selectOverflowArith(ResVReg, ResType,
I,
1246 ResType->
getOpcode() == SPIRV::OpTypeVector
1247 ? SPIRV::OpISubBorrowV
1248 : SPIRV::OpISubBorrowS);
1249 case TargetOpcode::G_UMULO:
1250 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1251 case TargetOpcode::G_SMULO:
1252 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1254 case TargetOpcode::G_SEXT:
1255 return selectExt(ResVReg, ResType,
I,
true);
1256 case TargetOpcode::G_ANYEXT:
1257 case TargetOpcode::G_ZEXT:
1258 return selectExt(ResVReg, ResType,
I,
false);
1259 case TargetOpcode::G_TRUNC:
1260 return selectTrunc(ResVReg, ResType,
I);
1261 case TargetOpcode::G_FPTRUNC:
1262 case TargetOpcode::G_FPEXT:
1263 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1265 case TargetOpcode::G_PTRTOINT:
1266 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1267 case TargetOpcode::G_INTTOPTR:
1268 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1269 case TargetOpcode::G_BITCAST:
1270 return selectBitcast(ResVReg, ResType,
I);
1271 case TargetOpcode::G_ADDRSPACE_CAST:
1272 return selectAddrSpaceCast(ResVReg, ResType,
I);
1273 case TargetOpcode::G_PTRMASK:
1274 return selectPtrMask(ResVReg, ResType,
I);
1275 case TargetOpcode::G_PTR_ADD: {
1277 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1281 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1282 (*II).getOpcode() == TargetOpcode::COPY ||
1283 (*II).getOpcode() == SPIRV::OpVariable) &&
1284 getImm(
I.getOperand(2), MRI));
1286 bool IsGVInit =
false;
1290 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1291 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1292 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1293 (*UseIt).getOpcode() == SPIRV::OpVariable) {
1303 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1315 return diagnoseUnsupported(
1316 I,
"incompatible result and operand types in a bitcast");
1318 MachineInstrBuilder MIB =
1319 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1326 : SPIRV::OpInBoundsPtrAccessChain))
1330 .
addUse(
I.getOperand(2).getReg())
1333 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1337 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1339 .
addUse(
I.getOperand(2).getReg())
1348 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1351 .
addImm(
static_cast<uint32_t
>(
1352 SPIRV::Opcode::InBoundsPtrAccessChain))
1355 .
addUse(
I.getOperand(2).getReg());
1360 case TargetOpcode::G_ATOMICRMW_OR:
1361 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1362 case TargetOpcode::G_ATOMICRMW_ADD:
1363 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1364 case TargetOpcode::G_ATOMICRMW_AND:
1365 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1366 case TargetOpcode::G_ATOMICRMW_MAX:
1367 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1368 case TargetOpcode::G_ATOMICRMW_MIN:
1369 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1370 case TargetOpcode::G_ATOMICRMW_SUB:
1371 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1372 case TargetOpcode::G_ATOMICRMW_XOR:
1373 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1374 case TargetOpcode::G_ATOMICRMW_UMAX:
1375 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1376 case TargetOpcode::G_ATOMICRMW_UMIN:
1377 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1378 case TargetOpcode::G_ATOMICRMW_XCHG:
1379 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1381 case TargetOpcode::G_ATOMICRMW_FADD:
1382 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1383 case TargetOpcode::G_ATOMICRMW_FSUB:
1385 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1386 ResType->
getOpcode() == SPIRV::OpTypeVector
1388 : SPIRV::OpFNegate);
1389 case TargetOpcode::G_ATOMICRMW_FMIN:
1390 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1391 case TargetOpcode::G_ATOMICRMW_FMAX:
1392 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1394 case TargetOpcode::G_FENCE:
1395 return selectFence(
I);
1397 case TargetOpcode::G_STACKSAVE:
1398 return selectStackSave(ResVReg, ResType,
I);
1399 case TargetOpcode::G_STACKRESTORE:
1400 return selectStackRestore(
I);
1402 case TargetOpcode::G_UNMERGE_VALUES:
1405 case TargetOpcode::G_TRAP:
1406 case TargetOpcode::G_UBSANTRAP:
1407 return selectTrap(
I);
1412 case TargetOpcode::DBG_LABEL:
1414 case TargetOpcode::G_DEBUGTRAP:
1415 return selectDebugTrap(ResVReg, ResType,
I);
1422bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1423 SPIRVTypeInst ResType,
1424 MachineInstr &
I)
const {
1425 unsigned Opcode = SPIRV::OpNop;
1432bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1433 SPIRVTypeInst ResType,
1435 GL::GLSLExtInst GLInst,
1436 bool setMIFlags,
bool useMISrc,
1439 SPIRV::InstructionSet::InstructionSet::GLSL_std_450))
1440 return diagnoseUnsupported(
1442 "this instruction is only supported with the GLSL extended instruction "
1444 return selectExtInst(ResVReg, ResType,
I,
1445 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1446 setMIFlags, useMISrc, SrcRegs);
1449bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1450 SPIRVTypeInst ResType,
1452 CL::OpenCLExtInst CLInst,
1453 bool setMIFlags,
bool useMISrc,
1455 return selectExtInst(ResVReg, ResType,
I,
1456 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1457 setMIFlags, useMISrc, SrcRegs);
1460bool SPIRVInstructionSelector::selectExtInst(
1461 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1462 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1464 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1465 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1466 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1470bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1471 SPIRVTypeInst ResType,
1474 bool setMIFlags,
bool useMISrc,
1477 for (
const auto &[InstructionSet, Opcode] : Insts) {
1481 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1484 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1489 const unsigned NumOps =
I.getNumOperands();
1492 I.getOperand(Index).getType() ==
1493 MachineOperand::MachineOperandType::MO_IntrinsicID)
1496 MIB.
add(
I.getOperand(Index));
1508bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1509 SPIRVTypeInst ResType,
1510 MachineInstr &
I)
const {
1511 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1512 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1513 for (
const auto &Ex : ExtInsts) {
1514 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1515 uint32_t Opcode = Ex.second;
1519 MachineIRBuilder MIRBuilder(
I);
1522 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1527 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1530 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1533 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1536 .
addImm(
static_cast<uint32_t
>(Ex.first))
1538 .
add(
I.getOperand(2))
1542 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1543 .
addDef(
I.getOperand(1).getReg())
1552bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1553 SPIRVTypeInst ResType,
1554 MachineInstr &
I)
const {
1555 Register CosResVReg =
I.getOperand(1).getReg();
1556 unsigned SrcIdx =
I.getNumExplicitDefs();
1561 MachineIRBuilder MIRBuilder(
I);
1563 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1568 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1571 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1573 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1576 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1578 .
add(
I.getOperand(SrcIdx))
1581 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1589 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1592 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1594 .
add(
I.getOperand(SrcIdx))
1596 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1599 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1601 .
add(
I.getOperand(SrcIdx))
1608bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1609 SPIRVTypeInst ResType,
1611 std::vector<Register> Srcs,
1612 unsigned Opcode)
const {
1613 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1623std::optional<SplitParts> SPIRVInstructionSelector::splitEvenOddLanes(
1624 Register PopCountReg,
unsigned ComponentCount, MachineInstr &
I,
1625 SPIRVTypeInst I32Type)
const {
1628 if (ComponentCount == 1) {
1631 Parts.IsScalar =
true;
1632 Parts.Type = I32Type;
1640 if (!selectOpWithSrcs(Parts.High, I32Type,
I, {PopCountReg, IdxOne},
1641 SPIRV::OpVectorExtractDynamic))
1642 return std::nullopt;
1644 if (!selectOpWithSrcs(Parts.Low, I32Type,
I, {PopCountReg, IdxZero},
1645 SPIRV::OpVectorExtractDynamic))
1646 return std::nullopt;
1650 MachineIRBuilder MIRBuilder(
I);
1651 Parts.IsScalar =
false;
1658 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1659 TII.get(SPIRV::OpVectorShuffle))
1664 for (
unsigned J = 1; J < ComponentCount * 2; J += 2)
1669 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1670 TII.get(SPIRV::OpVectorShuffle))
1675 for (
unsigned J = 0; J < ComponentCount * 2; J += 2)
1683bool SPIRVInstructionSelector::selectPopCount16(
Register ResVReg,
1684 SPIRVTypeInst ResType,
1687 unsigned Opcode)
const {
1688 Register OpReg =
I.getOperand(1).getReg();
1691 MachineIRBuilder MIRBuilder(
I);
1693 SPIRVTypeInst I32VectorType =
1696 bool IsVector = NumElems > 1;
1697 SPIRVTypeInst ExtType = IsVector ? I32VectorType : I32Type;
1700 if (!selectOpWithSrcs(ExtReg, ExtType,
I, {OpReg}, SPIRV::OpUConvert))
1704 if (!selectPopCount32(PopCountReg, ExtType,
I, ExtReg, Opcode))
1707 return selectOpWithSrcs(ResVReg, ResType,
I, {PopCountReg}, ExtOpcode);
1710bool SPIRVInstructionSelector::selectPopCount32(
Register ResVReg,
1711 SPIRVTypeInst ResType,
1714 unsigned Opcode)
const {
1715 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
1718bool SPIRVInstructionSelector::selectPopCount64(
Register ResVReg,
1719 SPIRVTypeInst ResType,
1722 unsigned Opcode)
const {
1724 if (ComponentCount > 2)
1725 return handle64BitOverflow(
1726 ResVReg, ResType,
I, SrcReg, Opcode,
1728 unsigned O) {
return this->selectPopCount64(R,
T,
I, S, O); });
1730 MachineIRBuilder MIRBuilder(
I);
1735 I32Type, 2 * ComponentCount, MIRBuilder,
false);
1739 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
1744 if (!selectPopCount32(Pop32, VecI32Type,
I, Vec32, Opcode))
1748 auto MaybeParts = splitEvenOddLanes(Pop32, ComponentCount,
I, I32Type);
1751 SplitParts &Parts = *MaybeParts;
1754 unsigned OpAdd = Parts.IsScalar ? SPIRV::OpIAddS : SPIRV::OpIAddV;
1756 if (!selectOpWithSrcs(Sum, Parts.Type,
I, {Parts.High, Parts.Low}, OpAdd))
1761 unsigned ConvOp = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
1762 return selectOpWithSrcs(ResVReg, ResType,
I, {Sum}, ConvOp);
1765bool SPIRVInstructionSelector::selectPopCount(
Register ResVReg,
1766 SPIRVTypeInst ResType,
1768 unsigned Opcode)
const {
1773 if (!STI.getTargetTriple().isVulkanOS())
1774 return selectUnOp(ResVReg, ResType,
I, Opcode);
1776 Register OpReg =
I.getOperand(1).getReg();
1779 : SPIRV::OpUConvert;
1783 return selectPopCount16(ResVReg, ResType,
I, ExtOpcode, Opcode);
1785 return selectPopCount32(ResVReg, ResType,
I, OpReg, Opcode);
1787 return selectPopCount64(ResVReg, ResType,
I, OpReg, Opcode);
1789 return diagnoseUnsupported(
I,
"unsupported operand bit width for popcount");
1793bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1794 SPIRVTypeInst ResType,
1796 unsigned Opcode)
const {
1798 Register SrcReg =
I.getOperand(1).getReg();
1803 unsigned DefOpCode = DefIt->getOpcode();
1804 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1807 if (
auto *VRD =
getVRegDef(*MRI, DefIt->getOperand(1).getReg()))
1808 DefOpCode = VRD->getOpcode();
1810 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1811 DefOpCode == TargetOpcode::G_CONSTANT ||
1812 DefOpCode == SPIRV::OpVariable || DefOpCode == SPIRV::OpConstantI) {
1818 uint32_t SpecOpcode = 0;
1820 case SPIRV::OpConvertPtrToU:
1821 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1823 case SPIRV::OpConvertUToPtr:
1824 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1829 TII.get(SPIRV::OpSpecConstantOp))
1839 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1843bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1844 SPIRVTypeInst ResType,
1845 MachineInstr &
I)
const {
1846 Register OpReg =
I.getOperand(1).getReg();
1847 SPIRVTypeInst OpType =
1850 return diagnoseUnsupported(
1851 I,
"incompatible result and operand types in a bitcast");
1852 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1862 if (
MemOp->isVolatile())
1863 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1864 if (
MemOp->isNonTemporal())
1865 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1867 if (!ST->isShader() &&
MemOp->getAlign().value())
1868 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1872 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1873 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1877 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1879 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1883 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1887 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1889 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1901 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1903 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1905 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
1909bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
1910 SPIRVTypeInst ResType,
1911 MachineInstr &
I)
const {
1913 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1918 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
1919 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
1921 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1923 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1927 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1931 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1932 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
1933 I.getDebugLoc(),
I);
1937 MachineIRBuilder MIRBuilder(
I);
1939 if (
I.getNumMemOperands()) {
1940 const MachineMemOperand *MemOp = *
I.memoperands_begin();
1941 if (MemOp->isAtomic())
1942 return selectAtomicLoad(ResVReg, ResType,
I);
1945 auto MIB = MIRBuilder.buildInstr(SPIRV::OpLoad)
1949 if (!
I.getNumMemOperands()) {
1950 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1952 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1961bool SPIRVInstructionSelector::selectAtomicLoad(
Register ResVReg,
1962 SPIRVTypeInst ResType,
1963 MachineInstr &
I)
const {
1964 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
1967 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1970 return diagnoseUnsupported(
I,
1971 "Lowering to SPIR-V of atomic load is only "
1972 "allowed for integer or floating point types");
1974 assert(
I.getNumMemOperands());
1975 const MachineMemOperand &MemOp = **
I.memoperands_begin();
1976 assert(MemOp.isAtomic());
1980 Register ScopeReg = buildI32Constant(Scope,
I);
1986 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
1987 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
1990 MachineIRBuilder MIRBuilder(
I);
1991 auto AtomicLoad = MIRBuilder.buildInstr(SPIRV::OpAtomicLoad)
1997 AtomicLoad.constrainAllUses(
TII,
TRI, RBI);
2001bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
2003 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2004 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2009 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2010 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2012 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2017 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2021 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2022 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2023 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2024 TII.get(SPIRV::OpImageWrite))
2030 if (sampledTypeIsSignedInteger(LLVMHandleType))
2033 BMI.constrainAllUses(
TII,
TRI, RBI);
2038 if (
I.getNumMemOperands()) {
2039 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2040 if (MemOp->isAtomic())
2041 return selectAtomicStore(
I);
2044 MachineIRBuilder MIRBuilder(
I);
2045 auto MIB = MIRBuilder.buildInstr(SPIRV::OpStore).
addUse(Ptr).
addUse(StoreVal);
2046 if (!
I.getNumMemOperands()) {
2047 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2049 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2058bool SPIRVInstructionSelector::selectAtomicStore(MachineInstr &
I)
const {
2059 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2062 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2063 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2068 assert(
I.getNumMemOperands());
2069 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2070 assert(MemOp.isAtomic());
2074 Register ScopeReg = buildI32Constant(Scope,
I);
2080 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2081 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2083 MachineIRBuilder MIRBuilder(
I);
2087 return diagnoseUnsupported(
2088 I,
"Lowering to SPIR-V of atomic store is only "
2089 "allowed for pointer types for physical addressing model");
2095 SPIRVTypeInst PtrAsIntSpirvType =
2102 MIRBuilder.buildInstr(SPIRV::OpConvertPtrToU)
2106 .constrainAllUses(
TII,
TRI, RBI);
2112 PtrAsIntSpirvType, MIRBuilder,
2115 MIRBuilder.getMF());
2117 MIRBuilder.buildInstr(SPIRV::OpBitcast)
2118 .addDef(PtrCastedToMatchValReg)
2121 .constrainAllUses(
TII,
TRI, RBI);
2123 StoreVal = PtrToUVal;
2124 Ptr = PtrCastedToMatchValReg;
2125 PointeeType = PtrAsIntSpirvType;
2129 return diagnoseUnsupported(
I,
2130 "Lowering to SPIR-V of atomic store is only "
2131 "allowed for integer or floating point types");
2133 auto AtomicStore = MIRBuilder.buildInstr(SPIRV::OpAtomicStore)
2138 AtomicStore.constrainAllUses(
TII,
TRI, RBI);
2143bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
2144 SPIRVTypeInst ResType,
2145 MachineInstr &
I)
const {
2146 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
2154 const Register PtrsReg =
I.getOperand(2).getReg();
2155 const uint32_t Alignment =
I.getOperand(3).getImm();
2156 const Register MaskReg =
I.getOperand(4).getReg();
2157 const Register PassthruReg =
I.getOperand(5).getReg();
2158 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2162 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
2173bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
2174 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
2181 const Register ValuesReg =
I.getOperand(1).getReg();
2182 const Register PtrsReg =
I.getOperand(2).getReg();
2183 const uint32_t Alignment =
I.getOperand(3).getImm();
2184 const Register MaskReg =
I.getOperand(4).getReg();
2185 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2189 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
2198bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
2199 const Twine &Msg)
const {
2200 const Function &
F =
I.getMF()->getFunction();
2201 F.getContext().diagnose(
2202 DiagnosticInfoUnsupported(
F, Msg,
I.getDebugLoc(),
DS_Error));
2206bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
2207 SPIRVTypeInst ResType,
2208 MachineInstr &
I)
const {
2209 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2210 return diagnoseUnsupported(
2211 I,
"llvm.stacksave intrinsic: this instruction requires the following "
2212 "SPIR-V extension: SPV_INTEL_variable_length_array");
2214 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
2221bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
2222 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2223 return diagnoseUnsupported(
2225 "llvm.stackrestore intrinsic: this instruction requires the following "
2226 "SPIR-V extension: SPV_INTEL_variable_length_array");
2227 if (!
I.getOperand(0).isReg())
2230 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
2231 .
addUse(
I.getOperand(0).getReg())
2237SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
2238 MachineIRBuilder MIRBuilder(
I);
2239 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
2246 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
2250 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2251 Type *ArrTy = ArrayType::get(ValTy, Num);
2253 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
2256 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
2263 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
2266 .
addImm(SPIRV::StorageClass::UniformConstant)
2277bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
2280 Register DstReg =
I.getOperand(0).getReg();
2284 return diagnoseUnsupported(
2285 I,
"OpCopyMemory requires operands to have the same type");
2286 uint64_t CopySize =
getIConstVal(
I.getOperand(2).getReg(), MRI);
2290 return diagnoseUnsupported(
2291 I,
"Unable to determine pointee type size for OpCopyMemory");
2292 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
2293 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
2294 return diagnoseUnsupported(
2295 I,
"OpCopyMemory requires the size to match the pointee type size");
2296 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemory))
2299 if (
I.getNumMemOperands()) {
2300 MachineIRBuilder MIRBuilder(
I);
2307bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &
I,
2310 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemorySized))
2311 .
addUse(
I.getOperand(0).getReg())
2313 .
addUse(
I.getOperand(2).getReg());
2314 if (
I.getNumMemOperands()) {
2315 MachineIRBuilder MIRBuilder(
I);
2322bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
2323 MachineInstr &
I)
const {
2325 Register SizeReg =
I.getOperand(2).getReg();
2327 SizeDef && SizeDef->
getOpcode() == TargetOpcode::G_CONSTANT &&
2331 Register SrcReg =
I.getOperand(1).getReg();
2332 if (
I.getOpcode() == TargetOpcode::G_MEMSET) {
2333 Register VarReg = getOrCreateMemSetGlobal(
I);
2336 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2338 ValTy,
I, SPIRV::StorageClass::UniformConstant);
2340 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
2344 if (!selectCopyMemory(
I, SrcReg))
2347 if (!selectCopyMemorySized(
I, SrcReg))
2350 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
2351 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
2356bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
2357 SPIRVTypeInst ResType,
2360 unsigned NegateOpcode)
const {
2362 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2365 Register ScopeReg = buildI32Constant(Scope,
I);
2367 Register Ptr =
I.getOperand(1).getReg();
2368 uint32_t ScSem =
static_cast<uint32_t
>(
2372 Register MemSemReg = buildI32Constant(MemSem,
I);
2374 Register ValueReg =
I.getOperand(2).getReg();
2375 if (NegateOpcode != 0) {
2378 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
2383 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2394bool SPIRVInstructionSelector::selectInterlockedAdd(
Register ResVReg,
2395 SPIRVTypeInst ResType,
2396 MachineInstr &
I)
const {
2397 Register Ptr =
I.getOperand(2).getReg();
2401 assert((SC == SPIRV::StorageClass::Workgroup ||
2402 SC == SPIRV::StorageClass::StorageBuffer) &&
2403 "InterlockedAdd requires Workgroup or StorageBuffer storage class");
2404 uint32_t
Scope =
static_cast<uint32_t
>(SC == SPIRV::StorageClass::Workgroup
2405 ? SPIRV::Scope::Workgroup
2406 : SPIRV::Scope::Device);
2407 Register ScopeReg = buildI32Constant(Scope,
I);
2410 Register MemSemReg = buildI32Constant(MemSem,
I);
2412 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
2423bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2424 unsigned ArgI =
I.getNumOperands() - 1;
2426 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2427 SPIRVTypeInst SrcType =
2429 if (!SrcType || SrcType->
getOpcode() != SPIRV::OpTypeVector)
2431 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2435 unsigned CurrentIndex = 0;
2436 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2437 Register ResVReg =
I.getOperand(i).getReg();
2440 LLT ResLLT = MRI->
getType(ResVReg);
2446 ResType = ScalarType;
2452 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
2455 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2461 for (
unsigned j = 0;
j < NumElements; ++
j) {
2462 MIB.
addImm(CurrentIndex + j);
2464 CurrentIndex += NumElements;
2468 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2480bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2483 Register MemSemReg = buildI32Constant(MemSem,
I);
2485 uint32_t
Scope =
static_cast<uint32_t
>(
2487 Register ScopeReg = buildI32Constant(Scope,
I);
2489 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2496bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2497 SPIRVTypeInst ResType,
2499 unsigned Opcode)
const {
2500 Type *ResTy =
nullptr;
2503 return diagnoseUnsupported(
2505 "Not enough info to select the arithmetic with overflow instruction");
2507 return diagnoseUnsupported(
I,
2508 "Expect struct type result for the arithmetic "
2509 "with overflow instruction");
2515 MachineIRBuilder MIRBuilder(
I);
2517 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2518 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2524 Register ZeroReg = buildZerosVal(ResType,
I);
2529 if (ResName.
size() > 0)
2534 BuildMI(BB, MIRBuilder.getInsertPt(),
I.getDebugLoc(),
TII.get(Opcode))
2537 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2538 MIB.
addUse(
I.getOperand(i).getReg());
2543 MRI->
setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2544 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2546 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2547 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2554 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2555 .
addDef(
I.getOperand(1).getReg())
2563bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2564 SPIRVTypeInst ResType,
2565 MachineInstr &
I)
const {
2567 "selectAtomicCmpXchg only handles the spv_cmpxchg intrinsic");
2568 Register Ptr =
I.getOperand(2).getReg();
2569 Register ScopeReg =
I.getOperand(5).getReg();
2570 Register MemSemEqReg =
I.getOperand(6).getReg();
2571 Register MemSemNeqReg =
I.getOperand(7).getReg();
2573 Register Val =
I.getOperand(4).getReg();
2577 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2596 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2603 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2615 case SPIRV::StorageClass::DeviceOnlyINTEL:
2616 case SPIRV::StorageClass::HostOnlyINTEL:
2625 bool IsGRef =
false;
2626 bool IsAllowedRefs =
2628 unsigned Opcode = It.getOpcode();
2629 if (Opcode == SPIRV::OpConstantComposite ||
2630 Opcode == SPIRV::OpSpecConstantComposite ||
2631 Opcode == SPIRV::OpVariable ||
2632 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2633 return IsGRef = true;
2634 return Opcode == SPIRV::OpName;
2636 return IsAllowedRefs && IsGRef;
2639Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2640 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2642 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2646SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2648 uint32_t Opcode)
const {
2649 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2650 TII.get(SPIRV::OpSpecConstantOp))
2658SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2659 SPIRVTypeInst SrcPtrTy)
const {
2660 SPIRVTypeInst GenericPtrTy =
2664 SPIRV::StorageClass::Generic),
2666 MachineFunction *MF =
I.getParent()->getParent();
2668 MachineInstrBuilder MIB = buildSpecConstantOp(
2670 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2680bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2681 SPIRVTypeInst ResType,
2682 MachineInstr &
I)
const {
2686 Register SrcPtr =
I.getOperand(1).getReg();
2690 if (SrcPtrTy->
getOpcode() != SPIRV::OpTypePointer ||
2691 ResType->
getOpcode() != SPIRV::OpTypePointer)
2692 return BuildCOPY(ResVReg, SrcPtr,
I);
2702 unsigned SpecOpcode =
2704 ?
static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric)
2707 ? static_cast<uint32_t>(
SPIRV::Opcode::GenericCastToPtr)
2714 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2716 .constrainAllUses(
TII,
TRI, RBI);
2718 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2720 buildSpecConstantOp(
2722 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2723 .constrainAllUses(
TII,
TRI, RBI);
2730 return BuildCOPY(ResVReg, SrcPtr,
I);
2732 if ((SrcSC == SPIRV::StorageClass::Function &&
2733 DstSC == SPIRV::StorageClass::Private) ||
2734 (DstSC == SPIRV::StorageClass::Function &&
2735 SrcSC == SPIRV::StorageClass::Private))
2736 return BuildCOPY(ResVReg, SrcPtr,
I);
2740 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2743 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2746 SPIRVTypeInst GenericPtrTy =
2765 return selectUnOp(ResVReg, ResType,
I,
2766 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
2768 return selectUnOp(ResVReg, ResType,
I,
2769 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
2771 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2773 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2783bool SPIRVInstructionSelector::selectPtrMask(
Register ResVReg,
2784 SPIRVTypeInst ResType,
2785 MachineInstr &
I)
const {
2787 return diagnoseUnsupported(
2788 I,
"G_PTRMASK is not supported with logical SPIR-V");
2793 Register PtrReg =
I.getOperand(1).getReg();
2794 Register MaskReg =
I.getOperand(2).getReg();
2813 ? SPIRV::OpBitwiseAndV
2814 : SPIRV::OpBitwiseAndS;
2837 return SPIRV::OpFOrdEqual;
2839 return SPIRV::OpFOrdGreaterThanEqual;
2841 return SPIRV::OpFOrdGreaterThan;
2843 return SPIRV::OpFOrdLessThanEqual;
2845 return SPIRV::OpFOrdLessThan;
2847 return SPIRV::OpFOrdNotEqual;
2849 return SPIRV::OpOrdered;
2851 return SPIRV::OpFUnordEqual;
2853 return SPIRV::OpFUnordGreaterThanEqual;
2855 return SPIRV::OpFUnordGreaterThan;
2857 return SPIRV::OpFUnordLessThanEqual;
2859 return SPIRV::OpFUnordLessThan;
2861 return SPIRV::OpFUnordNotEqual;
2863 return SPIRV::OpUnordered;
2873 return SPIRV::OpIEqual;
2875 return SPIRV::OpINotEqual;
2877 return SPIRV::OpSGreaterThanEqual;
2879 return SPIRV::OpSGreaterThan;
2881 return SPIRV::OpSLessThanEqual;
2883 return SPIRV::OpSLessThan;
2885 return SPIRV::OpUGreaterThanEqual;
2887 return SPIRV::OpUGreaterThan;
2889 return SPIRV::OpULessThanEqual;
2891 return SPIRV::OpULessThan;
2900 return SPIRV::OpPtrEqual;
2902 return SPIRV::OpPtrNotEqual;
2913 return SPIRV::OpLogicalEqual;
2915 return SPIRV::OpLogicalNotEqual;
2949bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
2950 SPIRVTypeInst ResType,
2952 unsigned OpAnyOrAll)
const {
2953 assert(
I.getNumOperands() == 3);
2954 assert(
I.getOperand(2).isReg());
2956 Register InputRegister =
I.getOperand(2).getReg();
2959 assert(InputType &&
"VReg has no type assigned");
2962 bool IsVectorTy = InputType->
getOpcode() == SPIRV::OpTypeVector;
2963 if (IsBoolTy && !IsVectorTy) {
2964 assert(ResVReg ==
I.getOperand(0).getReg());
2965 return BuildCOPY(ResVReg, InputRegister,
I);
2969 unsigned SpirvNotEqualId =
2970 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
2972 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
2977 IsBoolTy ? InputRegister
2985 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
2987 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
3004bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
3005 SPIRVTypeInst ResType,
3006 MachineInstr &
I)
const {
3007 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
3010bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
3011 SPIRVTypeInst ResType,
3012 MachineInstr &
I)
const {
3013 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
3017bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
3018 SPIRVTypeInst ResType,
3019 MachineInstr &
I)
const {
3020 assert(
I.getNumOperands() == 4);
3021 assert(
I.getOperand(2).isReg());
3022 assert(
I.getOperand(3).isReg());
3024 [[maybe_unused]] SPIRVTypeInst VecType =
3029 "dot product requires a vector of at least 2 components");
3031 [[maybe_unused]] SPIRVTypeInst EltType =
3040 .
addUse(
I.getOperand(2).getReg())
3041 .
addUse(
I.getOperand(3).getReg())
3046bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
3047 SPIRVTypeInst ResType,
3050 assert(
I.getNumOperands() == 4);
3051 assert(
I.getOperand(2).isReg());
3052 assert(
I.getOperand(3).isReg());
3055 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3059 .
addUse(
I.getOperand(2).getReg())
3060 .
addUse(
I.getOperand(3).getReg())
3067bool SPIRVInstructionSelector::selectIntegerDotExpansion(
3068 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3069 assert(
I.getNumOperands() == 4);
3070 assert(
I.getOperand(2).isReg());
3071 assert(
I.getOperand(3).isReg());
3075 Register Vec0 =
I.getOperand(2).getReg();
3076 Register Vec1 =
I.getOperand(3).getReg();
3080 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
3089 "dot product requires a vector of at least 2 components");
3092 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3102 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3113 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3125bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
3126 SPIRVTypeInst ResType,
3127 MachineInstr &
I)
const {
3129 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
3132 .
addUse(
I.getOperand(2).getReg())
3137bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
3138 SPIRVTypeInst ResType,
3139 MachineInstr &
I)
const {
3141 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
3144 .
addUse(
I.getOperand(2).getReg())
3149bool SPIRVInstructionSelector::selectOpIsFinite(
Register ResVReg,
3150 SPIRVTypeInst ResType,
3151 MachineInstr &
I)
const {
3153 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsFinite))
3156 .
addUse(
I.getOperand(2).getReg())
3161bool SPIRVInstructionSelector::selectOpIsNormal(
Register ResVReg,
3162 SPIRVTypeInst ResType,
3163 MachineInstr &
I)
const {
3165 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNormal))
3168 .
addUse(
I.getOperand(2).getReg())
3173template <
bool Signed>
3174bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
3175 SPIRVTypeInst ResType,
3176 MachineInstr &
I)
const {
3177 assert(
I.getNumOperands() == 5);
3178 assert(
I.getOperand(2).isReg());
3179 assert(
I.getOperand(3).isReg());
3180 assert(
I.getOperand(4).isReg());
3183 Register Acc =
I.getOperand(2).getReg();
3187 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3189 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
3194 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
3197 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3209template <
bool Signed>
3210bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
3211 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3212 assert(
I.getNumOperands() == 5);
3213 assert(
I.getOperand(2).isReg());
3214 assert(
I.getOperand(3).isReg());
3215 assert(
I.getOperand(4).isReg());
3218 Register Acc =
I.getOperand(2).getReg();
3224 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
3228 for (
unsigned i = 0; i < 4; i++) {
3251 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
3271 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3286bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
3287 SPIRVTypeInst ResType,
3288 MachineInstr &
I)
const {
3289 assert(
I.getNumOperands() == 3);
3290 assert(
I.getOperand(2).isReg());
3292 Register VZero = buildZerosValF(ResType,
I);
3293 Register VOne = buildOnesValF(ResType,
I);
3295 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
3298 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3300 .
addUse(
I.getOperand(2).getReg())
3307bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
3308 SPIRVTypeInst ResType,
3309 MachineInstr &
I)
const {
3310 assert(
I.getNumOperands() == 3);
3311 assert(
I.getOperand(2).isReg());
3313 Register InputRegister =
I.getOperand(2).getReg();
3315 auto &
DL =
I.getDebugLoc();
3318 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3325 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
3327 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
3335 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3340 if (NeedsConversion) {
3341 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
3352bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
3353 SPIRVTypeInst ResType,
3355 unsigned Opcode)
const {
3359 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
3365 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
3366 BMI.addUse(
I.getOperand(J).getReg());
3373bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
3376 bool WithGroupSync)
const {
3378 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
3380 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
3382 assert(((Scope != SPIRV::Scope::Workgroup) ||
3383 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
3384 "Workgroup Scope must set WorkGroupMemory semantic "
3385 "in Barrier instruction");
3387 assert(((Scope != SPIRV::Scope::Device) ||
3388 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
3389 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
3390 "Device Scope must set UniformMemory and ImageMemory semantic "
3391 "in Barrier instruction");
3397 if (WithGroupSync) {
3398 Register ExecReg = buildI32Constant(SPIRV::Scope::Workgroup,
I);
3402 Register ScopeReg = buildI32Constant(Scope,
I);
3403 Register MemSemReg = buildI32Constant(MemSem,
I);
3405 MI.addUse(ScopeReg).addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
3409bool SPIRVInstructionSelector::selectWaveActiveCountBits(
3410 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3415 if (!selectWaveOpInst(BallotReg, BallotType,
I,
3416 SPIRV::OpGroupNonUniformBallot))
3421 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3426 .
addImm(SPIRV::GroupOperation::Reduce)
3433bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
3434 SPIRVTypeInst ResType,
3435 MachineInstr &
I)
const {
3440 Register InputReg =
I.getOperand(2).getReg();
3445 bool IsVector = NumElems > 1;
3458 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
3459 SPIRV::OpGroupNonUniformAllEqual);
3464 ElementResults.
reserve(NumElems);
3466 for (
unsigned Idx = 0; Idx < NumElems; ++Idx) {
3479 ElemInput = Extracted;
3485 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3496 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3507bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3508 SPIRVTypeInst ResType,
3509 MachineInstr &
I)
const {
3511 assert(
I.getNumOperands() == 3);
3513 auto Op =
I.getOperand(2);
3523 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3525 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3526 return diagnoseUnsupported(
I,
"WavePrefixBitCount requires boolean input");
3547 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3551 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3558bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3559 SPIRVTypeInst ResType,
3561 bool IsUnsigned)
const {
3562 return selectWaveReduce(
3563 ResVReg, ResType,
I, IsUnsigned,
3564 [&](
Register InputRegister,
bool IsUnsigned) {
3565 const bool IsFloatTy =
3567 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3568 : SPIRV::OpGroupNonUniformSMax;
3569 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3573bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3574 SPIRVTypeInst ResType,
3576 bool IsUnsigned)
const {
3577 return selectWaveReduce(
3578 ResVReg, ResType,
I, IsUnsigned,
3579 [&](
Register InputRegister,
bool IsUnsigned) {
3580 const bool IsFloatTy =
3582 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3583 : SPIRV::OpGroupNonUniformSMin;
3584 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3588bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3589 SPIRVTypeInst ResType,
3590 MachineInstr &
I)
const {
3591 return selectWaveReduce(ResVReg, ResType,
I,
false,
3592 [&](
Register InputRegister,
bool IsUnsigned) {
3594 InputRegister, SPIRV::OpTypeFloat);
3595 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3596 : SPIRV::OpGroupNonUniformIAdd;
3600bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3601 SPIRVTypeInst ResType,
3602 MachineInstr &
I)
const {
3603 return selectWaveReduce(ResVReg, ResType,
I,
false,
3604 [&](
Register InputRegister,
bool IsUnsigned) {
3606 InputRegister, SPIRV::OpTypeFloat);
3607 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3608 : SPIRV::OpGroupNonUniformIMul;
3612template <
typename PickOpcodeFn>
3613bool SPIRVInstructionSelector::selectWaveReduce(
3614 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3615 PickOpcodeFn &&PickOpcode)
const {
3616 assert(
I.getNumOperands() == 3);
3617 assert(
I.getOperand(2).isReg());
3619 Register InputRegister =
I.getOperand(2).getReg();
3623 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3626 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3632 .
addImm(SPIRV::GroupOperation::Reduce)
3633 .
addUse(
I.getOperand(2).getReg())
3638bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3639 SPIRVTypeInst ResType,
3641 unsigned Opcode)
const {
3642 return selectWaveReduce(
3643 ResVReg, ResType,
I,
false,
3644 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3647bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3648 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3649 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3650 [&](
Register InputRegister,
bool IsUnsigned) {
3652 InputRegister, SPIRV::OpTypeFloat);
3654 ? SPIRV::OpGroupNonUniformFAdd
3655 : SPIRV::OpGroupNonUniformIAdd;
3659bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3660 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3661 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3662 [&](
Register InputRegister,
bool IsUnsigned) {
3664 InputRegister, SPIRV::OpTypeFloat);
3666 ? SPIRV::OpGroupNonUniformFMul
3667 : SPIRV::OpGroupNonUniformIMul;
3671template <
typename PickOpcodeFn>
3672bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3673 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3674 PickOpcodeFn &&PickOpcode)
const {
3675 assert(
I.getNumOperands() == 3);
3676 assert(
I.getOperand(2).isReg());
3678 Register InputRegister =
I.getOperand(2).getReg();
3682 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3685 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3691 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3692 .
addUse(
I.getOperand(2).getReg())
3697bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
3698 SPIRVTypeInst ResType,
3701 assert(
I.getNumOperands() == 3);
3702 assert(
I.getOperand(2).isReg());
3704 Register InputRegister =
I.getOperand(2).getReg();
3710 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
3721bool SPIRVInstructionSelector::selectBitreverse16(
Register ResVReg,
3722 SPIRVTypeInst ResType,
3727 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
3732 : SPIRV::OpUConvert;
3736 ShiftOp = SPIRV::OpShiftRightLogicalV;
3741 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3742 TII.get(SPIRV::OpConstantComposite))
3745 for (
unsigned It = 0; It <
N; ++It)
3749 ShiftConst = CompositeReg;
3754 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
3759 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
3764 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
3769 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
3772bool SPIRVInstructionSelector::handle64BitOverflow(
3774 unsigned int Opcode,
3781 "handle64BitOverflow should only be used for integer types");
3783 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
3785 MachineIRBuilder MIRBuilder(
I);
3787 SPIRVTypeInst I64x2Type =
3789 SPIRVTypeInst Vec2ResType =
3792 std::vector<Register> PartialRegs;
3794 unsigned CurrentComponent = 0;
3795 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
3799 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3800 TII.get(SPIRV::OpVectorShuffle))
3805 .
addImm(CurrentComponent)
3806 .
addImm(CurrentComponent + 1);
3816 PartialRegs.push_back(SubVecReg);
3819 if (CurrentComponent != ComponentCount) {
3825 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
3826 SPIRV::OpVectorExtractDynamic))
3835 PartialRegs.push_back(FinalElemResReg);
3839 return selectOpWithSrcs(ResVReg, ResType,
I, std::move(PartialRegs),
3840 SPIRV::OpCompositeConstruct);
3843bool SPIRVInstructionSelector::selectBitreverse64(
Register ResVReg,
3844 SPIRVTypeInst ResType,
3848 if (ComponentCount > 2)
3849 return handle64BitOverflow(
3850 ResVReg, ResType,
I, SrcReg, SPIRV::OpBitReverse,
3852 unsigned O) {
return this->selectBitreverse64(R,
T,
I, S); });
3854 MachineIRBuilder MIRBuilder(
I);
3858 I32Type, 2 * ComponentCount, MIRBuilder,
false);
3862 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
3867 if (!selectBitreverseNative(Reverse32, VecI32Type,
I, Vec32))
3874 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3875 TII.get(SPIRV::OpVectorShuffle))
3880 for (
unsigned J = 0; J < ComponentCount; ++J) {
3887 return selectOpWithSrcs(ResVReg, ResType,
I, {SwappedVec}, SPIRV::OpBitcast);
3890bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
3891 SPIRVTypeInst ResType,
3895 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
3903bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
3904 SPIRVTypeInst ResType,
3905 MachineInstr &
I)
const {
3906 Register OpReg =
I.getOperand(1).getReg();
3914 return selectBitreverse16(ResVReg, ResType,
I, OpReg);
3916 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
3918 return selectBitreverse64(ResVReg, ResType,
I, OpReg);
3920 return SPIRVInstructionSelector::diagnoseUnsupported(
3921 I,
"G_BITREVERSE only support 16,32,64 bits.");
3925 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
3936 unsigned AndOp = SPIRV::OpBitwiseAndS;
3937 unsigned OrOp = SPIRV::OpBitwiseOrS;
3938 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
3939 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
3941 AndOp = SPIRV::OpBitwiseAndV;
3942 OrOp = SPIRV::OpBitwiseOrV;
3943 ShlOp = SPIRV::OpShiftLeftLogicalV;
3944 ShrOp = SPIRV::OpShiftRightLogicalV;
3950 const unsigned Shift) ->
Register {
3958 Register MaskReg = CreateConst(Mask);
3959 Register ShiftReg = CreateConst(Shift);
3966 if (!selectOpWithSrcs(
T1, ResType,
I, {Input, ShiftReg}, ShrOp) ||
3967 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
3968 !selectOpWithSrcs(T3, ResType,
I, {Input, MaskReg}, AndOp) ||
3969 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
3970 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
3978 uint64_t
Mask = ~0ull;
3979 while ((Shift >>= 1) > 0) {
3986 return BuildCOPY(ResVReg, Result,
I);
3989bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
3990 SPIRVTypeInst ResType,
3991 MachineInstr &
I)
const {
3992 assert(
I.getOperand(0).isReg() &&
I.getOperand(1).isReg() &&
3993 "G_FREEZE must define and use a register");
3994 Register OpReg =
I.getOperand(1).getReg();
3998 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
4011 if (MachineInstr *Def = MRI->
getVRegDef(OpReg)) {
4012 if (
Def->getOpcode() == TargetOpcode::COPY)
4015 switch (
Def->getOpcode()) {
4016 case SPIRV::ASSIGN_TYPE:
4017 if (MachineInstr *AssignToDef =
4019 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
4020 Reg =
Def->getOperand(2).getReg();
4023 case SPIRV::OpUndef:
4024 Reg =
Def->getOperand(1).getReg();
4027 unsigned DestOpCode;
4029 DestOpCode = SPIRV::OpConstantNull;
4030 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze of a "
4031 "static undef/poison lowered to OpConstantNull\n");
4033 DestOpCode = TargetOpcode::COPY;
4035 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze "
4036 "skipped, lowered as a copy of the operand\n");
4038 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
4039 .
addDef(
I.getOperand(0).getReg())
4047bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
4048 SPIRVTypeInst ResType,
4049 MachineInstr &
I)
const {
4051 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4053 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4057 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
4062 for (
unsigned i =
I.getNumExplicitDefs();
4063 i <
I.getNumExplicitOperands() && IsConst; ++i)
4067 if (!IsConst &&
N < 2)
4068 return diagnoseUnsupported(
4069 I,
"There must be at least two constituent operands in a vector");
4072 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4073 TII.get(IsConst ? SPIRV::OpConstantComposite
4074 : SPIRV::OpCompositeConstruct))
4077 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
4078 MIB.
addUse(
I.getOperand(i).getReg());
4083bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
4084 SPIRVTypeInst ResType,
4085 MachineInstr &
I)
const {
4087 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4089 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4095 if (!
I.getOperand(
OpIdx).isReg())
4102 if (!IsConst &&
N < 2)
4103 return diagnoseUnsupported(
4104 I,
"There must be at least two constituent operands in a vector");
4107 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4108 TII.get(IsConst ? SPIRV::OpConstantComposite
4109 : SPIRV::OpCompositeConstruct))
4112 for (
unsigned i = 0; i <
N; ++i)
4118bool SPIRVInstructionSelector::selectConcatVectors(
Register ResVReg,
4119 SPIRVTypeInst ResType,
4120 MachineInstr &
I)
const {
4124 if (ResType->
getOpcode() != SPIRV::OpTypeVector)
4126 "Cannot select G_CONCAT_VECTORS with a non-vector result");
4128 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4129 TII.get(SPIRV::OpCompositeConstruct))
4139bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
4140 SPIRVTypeInst ResType,
4141 MachineInstr &
I)
const {
4146 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
4148 Opcode = SPIRV::OpDemoteToHelperInvocation;
4150 Opcode = SPIRV::OpKill;
4152 if (MachineInstr *NextI =
I.getNextNode()) {
4154 NextI->eraseFromParent();
4164bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
4165 SPIRVTypeInst ResType,
unsigned CmpOpc,
4166 MachineInstr &
I)
const {
4167 Register Cmp0 =
I.getOperand(2).getReg();
4168 Register Cmp1 =
I.getOperand(3).getReg();
4171 "CMP operands should have the same type");
4172 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
4182bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
4183 SPIRVTypeInst ResType,
4184 MachineInstr &
I)
const {
4185 auto Pred =
I.getOperand(1).getPredicate();
4188 Register CmpOperand =
I.getOperand(2).getReg();
4193 Register Op1 =
I.getOperand(3).getReg();
4197 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
4202 I.getOperand(3).setReg(NewOp1);
4208 return selectCmp(ResVReg, ResType, CmpOpc,
I);
4212SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
4213 SPIRVTypeInst ResType)
const {
4215 SPIRVTypeInst SpvI32Ty =
4218 auto ConstInt = ConstantInt::get(LLVMTy, Val);
4225 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4228 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
4231 .
addImm(APInt(32, Val).getZExtValue());
4233 GR.
add(ConstInt,
MI);
4240Register SPIRVInstructionSelector::buildI32ConstantInEntryBlock(
4241 uint32_t Val, MachineInstr &
I, SPIRVTypeInst ResType)
const {
4243 SPIRVTypeInst SpvI32Ty =
4245 auto *ConstInt = ConstantInt::get(LLVMTy, Val);
4250 MachineBasicBlock &EntryBB = *InsertIt->getParent();
4251 MachineInstr *
MI =
nullptr;
4255 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantNull))
4259 uint64_t ImmVal = APInt(32, Val).getZExtValue();
4260 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantI))
4266 GR.
add(ConstInt,
MI);
4271bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
4272 SPIRVTypeInst ResType,
4273 MachineInstr &
I)
const {
4275 return selectCmp(ResVReg, ResType, CmpOp,
I);
4278bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
4279 SPIRVTypeInst ResType,
4280 MachineInstr &
I)
const {
4282 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
4289 if (ResType->
getOpcode() != SPIRV::OpTypeVector &&
4290 ResType->
getOpcode() != SPIRV::OpTypeFloat)
4293 MachineIRBuilder MIRBuilder(
I);
4300 APFloat ConstVal(3.3219280948873623);
4304 APFloat::rmNearestTiesToEven, &LosesInfo);
4308 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
4309 ? SPIRV::OpVectorTimesScalar
4312 if (!selectOpWithSrcs(ArgReg, ResType,
I,
4313 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
4315 if (!selectExtInst(ResVReg, ResType,
I,
4316 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
4326Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
4327 MachineInstr &
I)
const {
4330 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4335bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
4341 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
4349 if (
Def->getOpcode() == SPIRV::OpConstantNull)
4352 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
4353 Def->getOpcode() == SPIRV::OpConstantI)
4366 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
4367 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
4369 Intrinsic::spv_const_composite)) {
4370 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
4371 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
4372 if (!IsZero(
Def->getOperand(i).getReg()))
4381Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
4382 MachineInstr &
I)
const {
4386 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4391Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
4392 MachineInstr &
I)
const {
4396 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4402 SPIRVTypeInst ResType,
4403 MachineInstr &
I)
const {
4407 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4412bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
4413 SPIRVTypeInst ResType,
4414 MachineInstr &
I)
const {
4415 Register SelectFirstArg =
I.getOperand(2).getReg();
4416 Register SelectSecondArg =
I.getOperand(3).getReg();
4425 SPIRV::OpTypeVector;
4432 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
4433 }
else if (IsPtrTy) {
4434 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
4436 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
4439 assert(IsScalarBool &&
"OpSelect with a scalar result requires a scalar "
4440 "boolean condition");
4442 Opcode = SPIRV::OpSelectSFSCond;
4443 }
else if (IsPtrTy) {
4444 Opcode = SPIRV::OpSelectSPSCond;
4446 Opcode = SPIRV::OpSelectSISCond;
4449 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4452 .
addUse(
I.getOperand(1).getReg())
4461bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
4462 SPIRVTypeInst ResType,
4464 MachineInstr &InsertAt,
4465 bool IsSigned)
const {
4467 Register ZeroReg = buildZerosVal(ResType, InsertAt);
4468 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
4469 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
4471 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4483bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
4484 SPIRVTypeInst ResType,
4485 MachineInstr &
I,
bool IsSigned,
4486 unsigned Opcode)
const {
4487 Register SrcReg =
I.getOperand(1).getReg();
4493 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
4498 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I,
false);
4500 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
4503bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
4504 SPIRVTypeInst ResType, MachineInstr &
I,
4505 bool IsSigned)
const {
4506 Register SrcReg =
I.getOperand(1).getReg();
4508 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
4512 if (ResType == SrcType)
4513 return BuildCOPY(ResVReg, SrcReg,
I);
4515 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4516 return selectUnOp(ResVReg, ResType,
I, Opcode);
4519bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
4520 SPIRVTypeInst ResType,
4522 bool IsSigned)
const {
4523 MachineIRBuilder MIRBuilder(
I);
4524 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
4536 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
4539 .
addUse(
I.getOperand(1).getReg())
4540 .
addUse(
I.getOperand(2).getReg())
4545 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
4548 .
addUse(
I.getOperand(1).getReg())
4549 .
addUse(
I.getOperand(2).getReg())
4557 unsigned SelectOpcode =
4558 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
4563 .
addUse(buildOnesVal(
true, ResType,
I))
4564 .
addUse(buildZerosVal(ResType,
I))
4571 .
addUse(buildOnesVal(
false, ResType,
I))
4576bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
4579 SPIRVTypeInst IntTy,
4580 SPIRVTypeInst BoolTy)
const {
4583 bool IsVectorTy = IntTy->
getOpcode() == SPIRV::OpTypeVector;
4584 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
4586 Register One = buildOnesVal(
false, IntTy,
I);
4594 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
4603bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
4604 SPIRVTypeInst ResType,
4605 MachineInstr &
I)
const {
4606 Register IntReg =
I.getOperand(1).getReg();
4609 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
4610 if (ArgType == ResType)
4611 return BuildCOPY(ResVReg, IntReg,
I);
4613 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4614 return selectUnOp(ResVReg, ResType,
I, Opcode);
4617bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
4618 SPIRVTypeInst ResType,
4619 MachineInstr &
I)
const {
4620 unsigned Opcode =
I.getOpcode();
4621 unsigned TpOpcode = ResType->
getOpcode();
4623 if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {
4624 assert(Opcode == TargetOpcode::G_CONSTANT &&
4625 I.getOperand(1).getCImm()->isZero());
4626 MachineBasicBlock &DepMBB =
I.getMF()->front();
4629 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
4636 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
4639bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
4640 SPIRVTypeInst ResType,
4641 MachineInstr &
I)
const {
4642 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4649bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
4650 SPIRVTypeInst ResType,
4651 MachineInstr &
I)
const {
4653 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
4657 .
addUse(
I.getOperand(3).getReg())
4659 .
addUse(
I.getOperand(2).getReg());
4660 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
4666bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
4667 SPIRVTypeInst ResType,
4668 MachineInstr &
I)
const {
4669 Type *MaybeResTy =
nullptr;
4674 "Expected aggregate type for extractv instruction");
4676 SPIRV::AccessQualifier::ReadWrite,
false);
4680 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
4683 .
addUse(
I.getOperand(2).getReg());
4684 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
4690bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
4691 SPIRVTypeInst ResType,
4692 MachineInstr &
I)
const {
4693 if (
getImm(
I.getOperand(4), MRI))
4694 return selectInsertVal(ResVReg, ResType,
I);
4696 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
4699 .
addUse(
I.getOperand(2).getReg())
4700 .
addUse(
I.getOperand(3).getReg())
4701 .
addUse(
I.getOperand(4).getReg())
4706bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
4707 SPIRVTypeInst ResType,
4708 MachineInstr &
I)
const {
4709 if (
getImm(
I.getOperand(3), MRI))
4710 return selectExtractVal(ResVReg, ResType,
I);
4712 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
4715 .
addUse(
I.getOperand(2).getReg())
4716 .
addUse(
I.getOperand(3).getReg())
4721bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
4722 SPIRVTypeInst ResType,
4723 MachineInstr &
I)
const {
4724 const bool IsGEPInBounds =
I.getOperand(2).getImm();
4730 ? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
4731 : SPIRV::OpAccessChain)
4732 : (IsGEPInBounds ?
SPIRV::OpInBoundsPtrAccessChain
4733 :
SPIRV::OpPtrAccessChain);
4735 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4739 .
addUse(
I.getOperand(3).getReg());
4741 (Opcode == SPIRV::OpPtrAccessChain ||
4742 Opcode == SPIRV::OpInBoundsPtrAccessChain ||
4743 (
getImm(
I.getOperand(4), MRI) &&
foldImm(
I.getOperand(4), MRI) == 0)) &&
4744 "Cannot translate GEP to OpAccessChain. First index must be 0.");
4747 const unsigned StartingIndex =
4748 (Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)
4751 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
4752 Res.addUse(
I.getOperand(i).getReg());
4753 Res.constrainAllUses(
TII,
TRI, RBI);
4758bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
4760 unsigned Lim =
I.getNumExplicitOperands();
4761 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
4762 Register OpReg =
I.getOperand(i).getReg();
4763 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
4765 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
4766 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
4767 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
4774 MachineFunction *MF =
I.getMF();
4786 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4787 TII.get(SPIRV::OpSpecConstantOp))
4790 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
4792 GR.
add(OpDefine, MIB);
4798bool SPIRVInstructionSelector::selectDerivativeInst(
4799 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
4800 const unsigned DPdOpCode)
const {
4803 if (!errorIfInstrOutsideShader(
I))
4809 Register SrcReg =
I.getOperand(2).getReg();
4814 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4817 .
addUse(
I.getOperand(2).getReg());
4819 MachineIRBuilder MIRBuilder(
I);
4822 if (componentCount != 1)
4826 const TargetRegisterClass *RegClass = GR.
getRegClass(SrcType);
4830 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4835 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4840 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4848bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
4849 SPIRVTypeInst ResType,
4850 MachineInstr &
I)
const {
4854 case Intrinsic::spv_load:
4855 return selectLoad(ResVReg, ResType,
I);
4856 case Intrinsic::spv_atomic_load:
4857 return selectAtomicLoad(ResVReg, ResType,
I);
4858 case Intrinsic::spv_store:
4859 return selectStore(
I);
4860 case Intrinsic::spv_atomic_store:
4861 return selectAtomicStore(
I);
4862 case Intrinsic::spv_extractv:
4863 return selectExtractVal(ResVReg, ResType,
I);
4864 case Intrinsic::spv_insertv:
4865 return selectInsertVal(ResVReg, ResType,
I);
4866 case Intrinsic::spv_extractelt:
4867 return selectExtractElt(ResVReg, ResType,
I);
4868 case Intrinsic::spv_insertelt:
4869 return selectInsertElt(ResVReg, ResType,
I);
4870 case Intrinsic::spv_gep:
4871 return selectGEP(ResVReg, ResType,
I);
4872 case Intrinsic::spv_bitcast: {
4873 Register OpReg =
I.getOperand(2).getReg();
4874 SPIRVTypeInst OpType =
4878 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
4880 case Intrinsic::spv_unref_global:
4881 case Intrinsic::spv_init_global: {
4882 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
4887 Register GVarVReg =
MI->getOperand(0).getReg();
4888 if (!selectGlobalValue(GVarVReg, *
MI, Init))
4893 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
4895 MI->eraseFromParent();
4899 case Intrinsic::spv_undef: {
4900 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4906 case Intrinsic::spv_poison:
4907 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpPoisonKHR))
4912 case Intrinsic::spv_freeze:
4913 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
4916 .
addUse(
I.getOperand(2).getReg())
4919 case Intrinsic::spv_named_boolean_spec_constant: {
4920 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
4921 : SPIRV::OpSpecConstantFalse;
4923 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
4924 .
addDef(
I.getOperand(0).getReg())
4927 unsigned SpecId =
I.getOperand(2).getImm();
4929 SPIRV::Decoration::SpecId, {SpecId});
4933 case Intrinsic::spv_const_composite: {
4935 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
4941 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
4943 std::function<bool(
Register)> HasSpecConstOperand =
4953 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
4954 J < Def->getNumExplicitOperands(); ++J) {
4955 if (
Def->getOperand(J).isReg() &&
4956 HasSpecConstOperand(
Def->getOperand(J).getReg()))
4962 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
4963 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
4964 : SPIRV::OpConstantComposite;
4965 unsigned ContinuedOpc = HasSpecConst
4966 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
4967 : SPIRV::OpConstantCompositeContinuedINTEL;
4968 MachineIRBuilder MIR(
I);
4970 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
4972 for (
auto *Instr : Instructions) {
4973 Instr->setDebugLoc(
I.getDebugLoc());
4978 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4985 case Intrinsic::spv_assign_name: {
4986 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
4987 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
4988 for (
unsigned i =
I.getNumExplicitDefs() + 2;
4989 i <
I.getNumExplicitOperands(); ++i) {
4990 MIB.
addImm(
I.getOperand(i).getImm());
4995 case Intrinsic::spv_switch: {
4996 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
4997 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
4998 if (
I.getOperand(i).isReg())
4999 MIB.
addReg(
I.getOperand(i).getReg());
5000 else if (
I.getOperand(i).isCImm())
5001 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
5002 else if (
I.getOperand(i).isMBB())
5003 MIB.
addMBB(
I.getOperand(i).getMBB());
5010 case Intrinsic::spv_loop_merge: {
5011 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
5012 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5013 if (
I.getOperand(i).isMBB())
5014 MIB.
addMBB(
I.getOperand(i).getMBB());
5021 case Intrinsic::spv_loop_control_intel: {
5023 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
5024 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
5029 case Intrinsic::spv_selection_merge: {
5031 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
5032 assert(
I.getOperand(1).isMBB() &&
5033 "operand 1 to spv_selection_merge must be a basic block");
5034 MIB.
addMBB(
I.getOperand(1).getMBB());
5035 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
5039 case Intrinsic::spv_cmpxchg:
5040 return selectAtomicCmpXchg(ResVReg, ResType,
I);
5041 case Intrinsic::spv_unreachable:
5042 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
5045 case Intrinsic::spv_abort:
5046 return selectAbort(
I);
5047 case Intrinsic::spv_alloca:
5048 return selectFrameIndex(ResVReg, ResType,
I);
5049 case Intrinsic::spv_alloca_array:
5050 return selectAllocaArray(ResVReg, ResType,
I);
5051 case Intrinsic::spv_assume:
5053 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
5054 .
addUse(
I.getOperand(1).getReg())
5059 case Intrinsic::spv_expect:
5061 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
5064 .
addUse(
I.getOperand(2).getReg())
5065 .
addUse(
I.getOperand(3).getReg())
5070 case Intrinsic::arithmetic_fence:
5071 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
5072 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
5075 .
addUse(
I.getOperand(2).getReg())
5079 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
5081 case Intrinsic::spv_thread_id:
5087 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
5089 case Intrinsic::spv_thread_id_in_group:
5095 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
5097 case Intrinsic::spv_group_id:
5103 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
5105 case Intrinsic::spv_flattened_thread_id_in_group:
5112 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
5114 case Intrinsic::spv_workgroup_size:
5115 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
5117 case Intrinsic::spv_global_size:
5118 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
5120 case Intrinsic::spv_global_offset:
5121 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
5123 case Intrinsic::spv_num_workgroups:
5124 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
5126 case Intrinsic::spv_subgroup_size:
5127 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
5129 case Intrinsic::spv_num_subgroups:
5130 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
5132 case Intrinsic::spv_subgroup_id:
5133 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
5134 case Intrinsic::spv_subgroup_local_invocation_id:
5135 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
5136 ResVReg, ResType,
I);
5137 case Intrinsic::spv_subgroup_max_size:
5138 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
5140 case Intrinsic::spv_fdot:
5141 return selectFloatDot(ResVReg, ResType,
I);
5142 case Intrinsic::spv_udot:
5143 case Intrinsic::spv_sdot:
5144 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5146 return selectIntegerDot(ResVReg, ResType,
I,
5147 IID == Intrinsic::spv_sdot);
5148 return selectIntegerDotExpansion(ResVReg, ResType,
I);
5149 case Intrinsic::spv_dot4add_i8packed:
5150 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5152 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
5153 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
5154 case Intrinsic::spv_dot4add_u8packed:
5155 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5157 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
5158 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
5159 case Intrinsic::spv_all:
5160 return selectAll(ResVReg, ResType,
I);
5161 case Intrinsic::spv_any:
5162 return selectAny(ResVReg, ResType,
I);
5163 case Intrinsic::spv_cross:
5164 return selectExtInst(ResVReg, ResType,
I, CL::cross, GL::Cross);
5165 case Intrinsic::spv_distance:
5166 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
5167 case Intrinsic::spv_lerp:
5168 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
5169 case Intrinsic::spv_length:
5170 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
5171 case Intrinsic::spv_degrees:
5172 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
5173 case Intrinsic::spv_faceforward:
5174 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
5175 case Intrinsic::spv_frac:
5176 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
5177 case Intrinsic::spv_isinf:
5178 return selectOpIsInf(ResVReg, ResType,
I);
5179 case Intrinsic::spv_isnan:
5180 return selectOpIsNan(ResVReg, ResType,
I);
5181 case Intrinsic::spv_isfinite:
5182 return selectOpIsFinite(ResVReg, ResType,
I);
5183 case Intrinsic::spv_isnormal:
5184 return selectOpIsNormal(ResVReg, ResType,
I);
5185 case Intrinsic::spv_normalize:
5186 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
5187 case Intrinsic::spv_refract:
5188 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
5189 case Intrinsic::spv_reflect:
5190 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
5191 case Intrinsic::spv_rsqrt:
5192 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
5193 case Intrinsic::spv_sign:
5194 return selectSign(ResVReg, ResType,
I);
5195 case Intrinsic::spv_smoothstep:
5196 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
5197 case Intrinsic::spv_firstbituhigh:
5198 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
5199 case Intrinsic::spv_firstbitshigh:
5200 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
5201 case Intrinsic::spv_firstbitlow:
5202 return selectFirstBitLow(ResVReg, ResType,
I);
5203 case Intrinsic::spv_all_memory_barrier:
5204 return selectBarrierInst(
I, SPIRV::Scope::Device,
5205 SPIRV::MemorySemantics::UniformMemory |
5206 SPIRV::MemorySemantics::ImageMemory |
5207 SPIRV::MemorySemantics::WorkgroupMemory,
5209 case Intrinsic::spv_all_memory_barrier_with_group_sync:
5210 return selectBarrierInst(
I, SPIRV::Scope::Device,
5211 SPIRV::MemorySemantics::UniformMemory |
5212 SPIRV::MemorySemantics::ImageMemory |
5213 SPIRV::MemorySemantics::WorkgroupMemory,
5215 case Intrinsic::spv_device_memory_barrier:
5216 return selectBarrierInst(
I, SPIRV::Scope::Device,
5217 SPIRV::MemorySemantics::UniformMemory |
5218 SPIRV::MemorySemantics::ImageMemory,
5220 case Intrinsic::spv_device_memory_barrier_with_group_sync:
5221 return selectBarrierInst(
I, SPIRV::Scope::Device,
5222 SPIRV::MemorySemantics::UniformMemory |
5223 SPIRV::MemorySemantics::ImageMemory,
5225 case Intrinsic::spv_group_memory_barrier:
5226 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5227 SPIRV::MemorySemantics::WorkgroupMemory,
5229 case Intrinsic::spv_group_memory_barrier_with_group_sync:
5230 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5231 SPIRV::MemorySemantics::WorkgroupMemory,
5233 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
5234 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
5235 SPIRV::StorageClass::StorageClass ResSC =
5238 return diagnoseUnsupported(
I,
"The target storage class is not castable "
5239 "from the Generic storage class");
5240 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
5248 case Intrinsic::spv_lifetime_start:
5249 case Intrinsic::spv_lifetime_end: {
5250 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
5251 : SPIRV::OpLifetimeStop;
5252 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
5253 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
5262 case Intrinsic::spv_saturate:
5263 return selectSaturate(ResVReg, ResType,
I);
5264 case Intrinsic::spv_nclamp:
5265 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
5266 case Intrinsic::spv_uclamp:
5267 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
5268 case Intrinsic::spv_sclamp:
5269 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
5270 case Intrinsic::spv_subgroup_prefix_bit_count:
5271 return selectWavePrefixBitCount(ResVReg, ResType,
I);
5272 case Intrinsic::spv_wave_active_countbits:
5273 return selectWaveActiveCountBits(ResVReg, ResType,
I);
5274 case Intrinsic::spv_wave_all_equal:
5275 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
5276 case Intrinsic::spv_wave_all:
5277 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
5278 case Intrinsic::spv_wave_any:
5279 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
5280 case Intrinsic::spv_subgroup_ballot:
5281 return selectWaveOpInst(ResVReg, ResType,
I,
5282 SPIRV::OpGroupNonUniformBallot);
5283 case Intrinsic::spv_wave_is_first_lane:
5284 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
5285 case Intrinsic::spv_wave_reduce_or:
5286 return selectWaveReduceOp(ResVReg, ResType,
I,
5287 SPIRV::OpGroupNonUniformBitwiseOr);
5288 case Intrinsic::spv_wave_reduce_xor:
5289 return selectWaveReduceOp(ResVReg, ResType,
I,
5290 SPIRV::OpGroupNonUniformBitwiseXor);
5291 case Intrinsic::spv_wave_reduce_and:
5292 return selectWaveReduceOp(ResVReg, ResType,
I,
5293 SPIRV::OpGroupNonUniformBitwiseAnd);
5294 case Intrinsic::spv_interlocked_add:
5295 return selectInterlockedAdd(ResVReg, ResType,
I);
5296 case Intrinsic::spv_wave_reduce_umax:
5297 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
5298 case Intrinsic::spv_wave_reduce_max:
5299 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
5300 case Intrinsic::spv_wave_reduce_umin:
5301 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
5302 case Intrinsic::spv_wave_reduce_min:
5303 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
5304 case Intrinsic::spv_wave_reduce_sum:
5305 return selectWaveReduceSum(ResVReg, ResType,
I);
5306 case Intrinsic::spv_wave_product:
5307 return selectWaveReduceProduct(ResVReg, ResType,
I);
5308 case Intrinsic::spv_wave_readlane:
5309 return selectWaveOpInst(ResVReg, ResType,
I,
5310 SPIRV::OpGroupNonUniformShuffle);
5311 case Intrinsic::spv_wave_prefix_sum:
5312 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
5313 case Intrinsic::spv_wave_prefix_product:
5314 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
5315 case Intrinsic::spv_quad_read_across_x: {
5316 return selectQuadSwap(ResVReg, ResType,
I, 0);
5318 case Intrinsic::spv_quad_read_across_y: {
5319 return selectQuadSwap(ResVReg, ResType,
I, 1);
5321 case Intrinsic::spv_step:
5322 return selectExtInst(ResVReg, ResType,
I, CL::step, GL::Step);
5323 case Intrinsic::spv_radians:
5324 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
5328 case Intrinsic::instrprof_increment:
5329 case Intrinsic::instrprof_increment_step:
5330 case Intrinsic::instrprof_value_profile:
5333 case Intrinsic::spv_value_md:
5335 case Intrinsic::spv_resource_handlefrombinding: {
5336 return selectHandleFromBinding(ResVReg, ResType,
I);
5338 case Intrinsic::spv_resource_counterhandlefrombinding:
5339 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
5340 case Intrinsic::spv_resource_updatecounter:
5341 return selectUpdateCounter(ResVReg, ResType,
I);
5342 case Intrinsic::spv_resource_store_typedbuffer: {
5343 return selectImageWriteIntrinsic(
I);
5345 case Intrinsic::spv_resource_load_typedbuffer: {
5346 return selectReadImageIntrinsic(ResVReg, ResType,
I);
5348 case Intrinsic::spv_resource_load_level: {
5349 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
5351 case Intrinsic::spv_resource_getdimensions_x:
5352 case Intrinsic::spv_resource_getdimensions_xy:
5353 case Intrinsic::spv_resource_getdimensions_xyz: {
5354 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
5356 case Intrinsic::spv_resource_getdimensions_levels_x:
5357 case Intrinsic::spv_resource_getdimensions_levels_xy:
5358 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
5359 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
5361 case Intrinsic::spv_resource_getdimensions_ms_xy:
5362 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
5363 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
5365 case Intrinsic::spv_resource_calculate_lod:
5366 case Intrinsic::spv_resource_calculate_lod_unclamped:
5367 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
5368 case Intrinsic::spv_resource_sample:
5369 case Intrinsic::spv_resource_sample_clamp:
5370 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
5371 case Intrinsic::spv_resource_samplebias:
5372 case Intrinsic::spv_resource_samplebias_clamp:
5373 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
5374 case Intrinsic::spv_resource_samplegrad:
5375 case Intrinsic::spv_resource_samplegrad_clamp:
5376 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
5377 case Intrinsic::spv_resource_samplelevel:
5378 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
5379 case Intrinsic::spv_resource_samplecmp:
5380 case Intrinsic::spv_resource_samplecmp_clamp:
5381 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
5382 case Intrinsic::spv_resource_samplecmplevelzero:
5383 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
5384 case Intrinsic::spv_resource_gather:
5385 case Intrinsic::spv_resource_gather_cmp:
5386 return selectGatherIntrinsic(ResVReg, ResType,
I);
5387 case Intrinsic::spv_resource_getbasepointer:
5388 case Intrinsic::spv_resource_getpointer: {
5389 return selectResourceGetPointer(ResVReg, ResType,
I);
5391 case Intrinsic::spv_pushconstant_getpointer: {
5392 return selectPushConstantGetPointer(ResVReg, ResType,
I);
5394 case Intrinsic::spv_discard: {
5395 return selectDiscard(ResVReg, ResType,
I);
5397 case Intrinsic::spv_resource_nonuniformindex: {
5398 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
5400 case Intrinsic::spv_unpackhalf2x16: {
5401 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
5403 case Intrinsic::spv_packhalf2x16: {
5404 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
5406 case Intrinsic::spv_ddx:
5407 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
5408 case Intrinsic::spv_ddy:
5409 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
5410 case Intrinsic::spv_ddx_coarse:
5411 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
5412 case Intrinsic::spv_ddy_coarse:
5413 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
5414 case Intrinsic::spv_ddx_fine:
5415 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
5416 case Intrinsic::spv_ddy_fine:
5417 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
5418 case Intrinsic::spv_fwidth:
5419 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
5420 case Intrinsic::spv_masked_gather:
5421 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5422 return selectMaskedGather(ResVReg, ResType,
I);
5423 return diagnoseUnsupported(
5424 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
5425 case Intrinsic::spv_masked_scatter:
5426 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5427 return selectMaskedScatter(
I);
5428 return diagnoseUnsupported(
5429 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
5430 case Intrinsic::returnaddress:
5431 case Intrinsic::frameaddress: {
5433 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5440 return diagnoseUnsupported(
I,
"intrinsic selection not implemented.");
5445bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
5446 SPIRVTypeInst ResType,
5447 MachineInstr &
I)
const {
5450 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
5457bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
5458 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5460 assert(Intr.getIntrinsicID() ==
5461 Intrinsic::spv_resource_counterhandlefrombinding);
5464 Register MainHandleReg = Intr.getOperand(2).getReg();
5466 assert(MainHandleDef->getIntrinsicID() ==
5467 Intrinsic::spv_resource_handlefrombinding);
5471 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
5472 Register IndexReg = MainHandleDef->getOperand(5).getReg();
5473 std::string CounterName =
5478 MachineIRBuilder MIRBuilder(
I);
5480 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
5482 ArraySize, IndexReg, CounterName, MIRBuilder);
5484 return BuildCOPY(ResVReg, CounterVarReg,
I);
5487bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
5488 SPIRVTypeInst ResType,
5489 MachineInstr &
I)
const {
5491 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
5493 Register CounterHandleReg = Intr.getOperand(2).getReg();
5494 Register IncrReg = Intr.getOperand(3).getReg();
5501 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
5502 assert(CounterVarPointeeType &&
5503 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
5504 "Counter variable must be a struct");
5506 SPIRV::StorageClass::StorageBuffer &&
5507 "Counter variable must be in the storage buffer storage class");
5509 "Counter variable must have exactly 1 member in the struct");
5510 const SPIRVTypeInst MemberType =
5513 "Counter variable struct must have a single i32 member");
5517 MachineIRBuilder MIRBuilder(
I);
5519 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
5522 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
5528 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5531 .
addUse(CounterHandleReg)
5538 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
5541 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
5544 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
5553 return BuildCOPY(ResVReg, AtomicRes,
I);
5561 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
5569bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
5570 SPIRVTypeInst ResType,
5571 MachineInstr &
I)
const {
5579 Register ImageReg =
I.getOperand(2).getReg();
5587 Register IdxReg =
I.getOperand(3).getReg();
5589 MachineInstr &Pos =
I;
5591 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
5595bool SPIRVInstructionSelector::generateSampleImage(
5598 DebugLoc Loc, MachineInstr &Pos)
const {
5609 if (!loadHandleBeforePosition(NewSamplerReg,
5615 MachineIRBuilder MIRBuilder(Pos);
5628 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
5629 ImOps.Lod.has_value();
5630 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
5631 : SPIRV::OpImageSampleImplicitLod;
5633 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
5634 : SPIRV::OpImageSampleDrefImplicitLod;
5643 MIB.
addUse(*ImOps.Compare);
5645 uint32_t ImageOperands = 0;
5647 ImageOperands |= SPIRV::ImageOperand::Bias;
5649 ImageOperands |= SPIRV::ImageOperand::Lod;
5650 if (ImOps.GradX && ImOps.GradY)
5651 ImageOperands |= SPIRV::ImageOperand::Grad;
5652 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
5654 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
5657 "Non-constant offsets are not supported in sample instructions.");
5661 ImageOperands |= SPIRV::ImageOperand::MinLod;
5663 if (ImageOperands != 0) {
5664 MIB.
addImm(ImageOperands);
5665 if (ImageOperands & SPIRV::ImageOperand::Bias)
5667 if (ImageOperands & SPIRV::ImageOperand::Lod)
5669 if (ImageOperands & SPIRV::ImageOperand::Grad) {
5670 MIB.
addUse(*ImOps.GradX);
5671 MIB.
addUse(*ImOps.GradY);
5674 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
5675 MIB.
addUse(*ImOps.Offset);
5676 if (ImageOperands & SPIRV::ImageOperand::MinLod)
5677 MIB.
addUse(*ImOps.MinLod);
5684bool SPIRVInstructionSelector::selectImageQuerySize(
5686 std::optional<Register> LodReg)
const {
5688 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
5691 "ImageReg is not an image type.");
5693 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
5695 unsigned NumComponents = 0;
5697 case SPIRV::Dim::DIM_1D:
5698 case SPIRV::Dim::DIM_Buffer:
5699 NumComponents =
IsArray ? 2 : 1;
5701 case SPIRV::Dim::DIM_2D:
5702 case SPIRV::Dim::DIM_Cube:
5703 case SPIRV::Dim::DIM_Rect:
5704 NumComponents =
IsArray ? 3 : 2;
5706 case SPIRV::Dim::DIM_3D:
5710 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
5715 SPIRVTypeInst ResType =
5720 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5730bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
5731 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5732 Register ImageReg =
I.getOperand(2).getReg();
5739 return selectImageQuerySize(NewImageReg, ResVReg,
I);
5742bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
5743 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5744 Register ImageReg =
I.getOperand(2).getReg();
5753 Register LodReg =
I.getOperand(3).getReg();
5756 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
5758 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
5765 TII.get(SPIRV::OpImageQueryLevels))
5772 TII.get(SPIRV::OpCompositeConstruct))
5782bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
5783 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5784 Register ImageReg =
I.getOperand(2).getReg();
5795 "OpImageQuerySamples requires a multisampled image");
5797 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
5805 TII.get(SPIRV::OpImageQuerySamples))
5812 TII.get(SPIRV::OpCompositeConstruct))
5822bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
5823 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5824 Register ImageReg =
I.getOperand(2).getReg();
5825 Register SamplerReg =
I.getOperand(3).getReg();
5826 Register CoordinateReg =
I.getOperand(4).getReg();
5842 if (!loadHandleBeforePosition(
5847 MachineIRBuilder MIRBuilder(
I);
5853 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
5863 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
5870 unsigned ExtractedIndex =
5872 Intrinsic::spv_resource_calculate_lod_unclamped
5876 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5877 TII.get(SPIRV::OpCompositeExtract))
5887bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
5888 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5889 Register ImageReg =
I.getOperand(2).getReg();
5890 Register SamplerReg =
I.getOperand(3).getReg();
5891 Register CoordinateReg =
I.getOperand(4).getReg();
5892 ImageOperands ImOps;
5893 if (
I.getNumOperands() > 5)
5894 ImOps.Offset =
I.getOperand(5).getReg();
5895 if (
I.getNumOperands() > 6)
5896 ImOps.MinLod =
I.getOperand(6).getReg();
5897 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5898 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5901bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
5902 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5903 Register ImageReg =
I.getOperand(2).getReg();
5904 Register SamplerReg =
I.getOperand(3).getReg();
5905 Register CoordinateReg =
I.getOperand(4).getReg();
5906 ImageOperands ImOps;
5907 ImOps.Bias =
I.getOperand(5).getReg();
5908 if (
I.getNumOperands() > 6)
5909 ImOps.Offset =
I.getOperand(6).getReg();
5910 if (
I.getNumOperands() > 7)
5911 ImOps.MinLod =
I.getOperand(7).getReg();
5912 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5913 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5916bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
5917 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5918 Register ImageReg =
I.getOperand(2).getReg();
5919 Register SamplerReg =
I.getOperand(3).getReg();
5920 Register CoordinateReg =
I.getOperand(4).getReg();
5921 ImageOperands ImOps;
5922 ImOps.GradX =
I.getOperand(5).getReg();
5923 ImOps.GradY =
I.getOperand(6).getReg();
5924 if (
I.getNumOperands() > 7)
5925 ImOps.Offset =
I.getOperand(7).getReg();
5926 if (
I.getNumOperands() > 8)
5927 ImOps.MinLod =
I.getOperand(8).getReg();
5928 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5929 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5932bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
5933 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5934 Register ImageReg =
I.getOperand(2).getReg();
5935 Register SamplerReg =
I.getOperand(3).getReg();
5936 Register CoordinateReg =
I.getOperand(4).getReg();
5937 ImageOperands ImOps;
5938 ImOps.Lod =
I.getOperand(5).getReg();
5939 if (
I.getNumOperands() > 6)
5940 ImOps.Offset =
I.getOperand(6).getReg();
5941 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5942 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5945bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
5946 SPIRVTypeInst ResType,
5947 MachineInstr &
I)
const {
5948 Register ImageReg =
I.getOperand(2).getReg();
5949 Register SamplerReg =
I.getOperand(3).getReg();
5950 Register CoordinateReg =
I.getOperand(4).getReg();
5951 ImageOperands ImOps;
5952 ImOps.Compare =
I.getOperand(5).getReg();
5953 if (
I.getNumOperands() > 6)
5954 ImOps.Offset =
I.getOperand(6).getReg();
5955 if (
I.getNumOperands() > 7)
5956 ImOps.MinLod =
I.getOperand(7).getReg();
5957 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5958 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5961bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
5962 SPIRVTypeInst ResType,
5963 MachineInstr &
I)
const {
5964 Register ImageReg =
I.getOperand(2).getReg();
5965 Register CoordinateReg =
I.getOperand(3).getReg();
5966 Register LodReg =
I.getOperand(4).getReg();
5968 ImageOperands ImOps;
5970 if (
I.getNumOperands() > 5)
5971 ImOps.Offset =
I.getOperand(5).getReg();
5983 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
5984 I.getDebugLoc(),
I, &ImOps);
5987bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
5988 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5989 Register ImageReg =
I.getOperand(2).getReg();
5990 Register SamplerReg =
I.getOperand(3).getReg();
5991 Register CoordinateReg =
I.getOperand(4).getReg();
5992 ImageOperands ImOps;
5993 ImOps.Compare =
I.getOperand(5).getReg();
5994 if (
I.getNumOperands() > 6)
5995 ImOps.Offset =
I.getOperand(6).getReg();
5998 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5999 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6002bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
6003 SPIRVTypeInst ResType,
6004 MachineInstr &
I)
const {
6005 Register ImageReg =
I.getOperand(2).getReg();
6006 Register SamplerReg =
I.getOperand(3).getReg();
6007 Register CoordinateReg =
I.getOperand(4).getReg();
6010 "ImageReg is not an image type.");
6015 ComponentOrCompareReg =
I.getOperand(5).getReg();
6016 OffsetReg =
I.getOperand(6).getReg();
6019 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
6023 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6024 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
6025 Dim != SPIRV::Dim::DIM_Rect) {
6027 "Gather operations are only supported for 2D, Cube, and Rect images.");
6034 if (!loadHandleBeforePosition(
6039 MachineIRBuilder MIRBuilder(
I);
6040 SPIRVTypeInst SampledImageType =
6045 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6053 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
6055 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
6057 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6062 .
addUse(ComponentOrCompareReg);
6064 uint32_t ImageOperands = 0;
6065 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
6066 if (Dim == SPIRV::Dim::DIM_Cube) {
6068 "Gather operations with offset are not supported for Cube images.");
6072 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6074 ImageOperands |= SPIRV::ImageOperand::Offset;
6078 if (ImageOperands != 0) {
6079 MIB.
addImm(ImageOperands);
6081 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6089bool SPIRVInstructionSelector::generateImageReadOrFetch(
6092 const ImageOperands *ImOps)
const {
6095 "ImageReg is not an image type.");
6097 bool IsSignedInteger =
6102 bool IsFetch = (SampledOp.getImm() == 1);
6104 auto AddOperands = [&](MachineInstrBuilder &MIB) {
6105 uint32_t ImageOperandsMask = 0;
6106 if (IsSignedInteger)
6107 ImageOperandsMask |= 0x1000;
6109 if (IsFetch && ImOps) {
6111 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
6112 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
6114 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
6116 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
6120 if (ImageOperandsMask != 0) {
6121 MIB.
addImm(ImageOperandsMask);
6122 if (IsFetch && ImOps) {
6125 if (ImOps->Offset &&
6126 (ImageOperandsMask &
6127 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
6128 MIB.
addUse(*ImOps->Offset);
6134 if (ResultSize == 4) {
6137 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
6144 BMI.constrainAllUses(
TII,
TRI, RBI);
6148 SPIRVTypeInst ReadType = widenTypeToVec4(ResType, Pos);
6152 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
6158 BMI.constrainAllUses(
TII,
TRI, RBI);
6160 if (ResultSize == 1) {
6169 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
6172bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
6173 SPIRVTypeInst ResType,
6174 MachineInstr &
I)
const {
6175 Register ResourcePtr =
I.getOperand(2).getReg();
6177 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
6186 MachineIRBuilder MIRBuilder(
I);
6191 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6197 if (
I.getNumExplicitOperands() > 3) {
6198 Register IndexReg =
I.getOperand(3).getReg();
6205bool SPIRVInstructionSelector::selectPushConstantGetPointer(
6206 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6211bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
6212 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6213 Register ObjReg =
I.getOperand(2).getReg();
6214 if (!BuildCOPY(ResVReg, ObjReg,
I))
6224 decorateUsesAsNonUniform(ResVReg);
6228void SPIRVInstructionSelector::decorateUsesAsNonUniform(
6231 while (WorkList.
size() > 0) {
6235 bool IsDecorated =
false;
6237 if (
Use.getOpcode() == SPIRV::OpDecorate &&
6238 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
6244 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
6246 if (ResultReg == CurrentReg)
6254 SPIRV::Decoration::NonUniformEXT, {});
6259bool SPIRVInstructionSelector::extractSubvector(
6261 MachineInstr &InsertionPoint)
const {
6263 [[maybe_unused]] uint64_t InputSize =
6266 assert(InputSize > 1 &&
"The input must be a vector.");
6267 assert(ResultSize > 1 &&
"The result must be a vector.");
6268 assert(ResultSize < InputSize &&
6269 "Cannot extract more element than there are in the input.");
6272 const TargetRegisterClass *ScalarRegClass = GR.
getRegClass(ScalarType);
6273 for (uint64_t
I = 0;
I < ResultSize;
I++) {
6276 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6285 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
6287 TII.get(SPIRV::OpCompositeConstruct))
6291 for (
Register ComponentReg : ComponentRegisters)
6292 MIB.
addUse(ComponentReg);
6297bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
6298 MachineInstr &
I)
const {
6305 Register ImageReg =
I.getOperand(1).getReg();
6313 Register CoordinateReg =
I.getOperand(2).getReg();
6314 Register DataReg =
I.getOperand(3).getReg();
6317 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
6325Register SPIRVInstructionSelector::buildPointerToResource(
6326 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
6327 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
6328 StringRef Name, MachineIRBuilder MIRBuilder)
const {
6330 if (ArraySize == 1) {
6331 SPIRVTypeInst PtrType =
6334 "SpirvResType did not have an explicit layout.");
6339 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
6340 SPIRVTypeInst VarPointerType =
6343 VarPointerType, Set,
Binding, Name, MIRBuilder);
6345 SPIRVTypeInst ResPointerType =
6358bool SPIRVInstructionSelector::selectFirstBitSet16(
6359 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
6360 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
6362 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
6366 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
6369bool SPIRVInstructionSelector::selectFirstBitSet32(
6371 unsigned BitSetOpcode)
const {
6372 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6375 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6382bool SPIRVInstructionSelector::selectFirstBitSet64(
6384 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
6397 if (ComponentCount > 2) {
6398 auto Func = [
this, SwapPrimarySide](
Register ResVReg, SPIRVTypeInst ResType,
6400 unsigned Opcode) ->
bool {
6401 return this->selectFirstBitSet64(ResVReg, ResType,
I, SrcReg, Opcode,
6405 return handle64BitOverflow(ResVReg, ResType,
I, SrcReg, BitSetOpcode, Func);
6409 MachineIRBuilder MIRBuilder(
I);
6411 BaseType, 2 * ComponentCount, MIRBuilder,
false);
6415 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
6421 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
6428 bool IsScalarRes = ResType->
getOpcode() != SPIRV::OpTypeVector;
6431 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
6432 SPIRV::OpVectorExtractDynamic))
6434 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
6435 SPIRV::OpVectorExtractDynamic))
6439 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6440 TII.get(SPIRV::OpVectorShuffle))
6448 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
6454 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6455 TII.get(SPIRV::OpVectorShuffle))
6463 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
6483 SelectOp = SPIRV::OpSelectSISCond;
6484 AddOp = SPIRV::OpIAddS;
6492 SelectOp = SPIRV::OpSelectVIVCond;
6493 AddOp = SPIRV::OpIAddV;
6499 Register RegSecondaryOffset = Reg0;
6503 if (SwapPrimarySide) {
6504 PrimaryReg = LowReg;
6505 SecondaryReg = HighReg;
6506 RegPrimaryOffset = Reg0;
6507 RegSecondaryOffset = Reg32;
6512 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
6513 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
6518 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
6519 SPIRV::OpINotEqual))
6526 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
6527 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
6532 if (SwapPrimarySide) {
6534 if (!selectOpWithSrcs(RegAdd, ResType,
I,
6535 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
6546 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
6547 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
6552 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
6553 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
6556 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
6560bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
6561 SPIRVTypeInst ResType,
6563 bool IsSigned)
const {
6565 Register OpReg =
I.getOperand(2).getReg();
6568 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
6569 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
6573 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6575 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6577 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6580 return diagnoseUnsupported(
6582 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
6586bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
6587 SPIRVTypeInst ResType,
6588 MachineInstr &
I)
const {
6590 Register OpReg =
I.getOperand(2).getReg();
6595 unsigned ExtendOpcode = SPIRV::OpUConvert;
6596 unsigned BitSetOpcode = GL::FindILsb;
6600 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6602 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6604 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6607 return diagnoseUnsupported(
I,
6608 "spv_firstbitlow only supports 16,32,64 bits.");
6612bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
6613 SPIRVTypeInst ResType,
6614 MachineInstr &
I)
const {
6618 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
6621 .
addUse(
I.getOperand(2).getReg())
6624 unsigned Alignment =
I.getOperand(3).getImm();
6638 while (!Worklist.
empty()) {
6640 switch (
T->getOpcode()) {
6641 case SPIRV::OpTypeInt:
6642 case SPIRV::OpTypeFloat:
6643 case SPIRV::OpTypePointer:
6645 case SPIRV::OpTypeVector:
6646 case SPIRV::OpTypeMatrix:
6647 case SPIRV::OpTypeArray: {
6648 Register OperandReg =
T->getOperand(1).getReg();
6652 case SPIRV::OpTypeStruct:
6653 for (
unsigned Idx = 1,
E =
T->getNumOperands(); Idx <
E; ++Idx) {
6654 Register OperandReg =
T->getOperand(Idx).getReg();
6666bool SPIRVInstructionSelector::selectAbort(MachineInstr &
I)
const {
6667 assert(
I.getNumExplicitOperands() == 2);
6669 Register MsgReg =
I.getOperand(1).getReg();
6671 assert(MsgType &&
"Message argument of llvm.spv.abort has no SPIR-V type");
6674 return diagnoseUnsupported(
6676 "llvm.spv.abort message type must be a concrete SPIR-V type (numerical "
6677 "scalar, pointer, vector, matrix, or aggregate of such types)");
6680 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
6687bool SPIRVInstructionSelector::selectTrap(MachineInstr &
I)
const {
6696 uint32_t MsgVal = ~0
u;
6697 if (
I.getOpcode() == TargetOpcode::G_UBSANTRAP)
6698 MsgVal =
static_cast<uint32_t
>(
I.getOperand(0).
getImm());
6701 Register MsgReg = buildI32ConstantInEntryBlock(MsgVal,
I, MsgType);
6704 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
6711bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
6712 SPIRVTypeInst ResType,
6713 MachineInstr &
I)
const {
6717 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
6720 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
6723 unsigned Alignment =
I.getOperand(2).getImm();
6730bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
6735 const MachineInstr *PrevI =
I.getPrevNode();
6737 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
6741 .
addMBB(
I.getOperand(0).getMBB())
6746 .
addMBB(
I.getOperand(0).getMBB())
6751bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
6762 const MachineInstr *NextI =
I.getNextNode();
6764 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
6770 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
6772 .
addUse(
I.getOperand(0).getReg())
6773 .
addMBB(
I.getOperand(1).getMBB())
6779bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
6780 MachineInstr &
I)
const {
6782 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
6784 const unsigned NumOps =
I.getNumOperands();
6785 for (
unsigned i = 1; i <
NumOps; i += 2) {
6786 MIB.
addUse(
I.getOperand(i + 0).getReg());
6787 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
6793bool SPIRVInstructionSelector::selectGlobalValue(
6794 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
6796 MachineIRBuilder MIRBuilder(
I);
6797 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
6800 std::string GlobalIdent;
6802 unsigned &
ID = UnnamedGlobalIDs[GV];
6804 ID = UnnamedGlobalIDs.
size();
6805 GlobalIdent =
"__unnamed_" + Twine(
ID).str();
6831 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
6838 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
6843 MachineInstrBuilder MIB1 =
6844 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
6847 MachineInstrBuilder MIB2 =
6849 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
6853 GR.
add(ConstVal, MIB2);
6861 MachineInstrBuilder MIB3 =
6862 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
6865 GR.
add(ConstVal, MIB3);
6871 assert(NewReg != ResVReg);
6872 return BuildCOPY(ResVReg, NewReg,
I);
6882 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
6885 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
6891 SPIRVTypeInst ResType =
6895 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
6900 if (
GlobalVar->isExternallyInitialized() &&
6901 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
6902 constexpr unsigned ReadWriteINTEL = 3u;
6905 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
6911bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
6912 SPIRVTypeInst ResType,
6913 MachineInstr &
I)
const {
6915 return selectExtInst(ResVReg, ResType,
I, CL::log10);
6923 MachineIRBuilder MIRBuilder(
I);
6928 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6931 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6933 .
add(
I.getOperand(1))
6938 ResType->
getOpcode() == SPIRV::OpTypeFloat);
6948 APFloat::rmNearestTiesToEven, &LosesInfo);
6952 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
6953 ? SPIRV::OpVectorTimesScalar
6964bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
6965 SPIRVTypeInst ResType,
6966 MachineInstr &
I)
const {
6969 return selectExtInst(ResVReg, ResType,
I, CL::pown);
6975 Register ExpReg =
I.getOperand(2).getReg();
6977 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
6978 SPIRV::OpConvertSToF))
6980 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
6987bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
6988 SPIRVTypeInst ResType,
6989 MachineInstr &
I)
const {
7005 MachineIRBuilder MIRBuilder(
I);
7008 ResType, MIRBuilder, SPIRV::StorageClass::Function);
7021 MachineBasicBlock &EntryBB =
I.getMF()->
front();
7023 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
7026 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7032 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7035 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
7038 .
add(
I.getOperand(
I.getNumExplicitDefs()))
7042 Register IntegralPartReg =
I.getOperand(1).getReg();
7043 if (IntegralPartReg.
isValid()) {
7045 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7056 assert(
false &&
"GLSL::Modf is deprecated.");
7067bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
7068 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7069 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7070 MachineIRBuilder MIRBuilder(
I);
7071 const SPIRVTypeInst Vec3Ty =
7074 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
7086 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7090 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7096 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7103 assert(
I.getOperand(2).isReg());
7104 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
7108 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
7119bool SPIRVInstructionSelector::loadBuiltinInputID(
7120 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7121 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7122 MachineIRBuilder MIRBuilder(
I);
7124 ResType, MIRBuilder, SPIRV::StorageClass::Input);
7139 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7143 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7152SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
7153 MachineInstr &
I)
const {
7154 MachineIRBuilder MIRBuilder(
I);
7155 if (
Type->getOpcode() != SPIRV::OpTypeVector)
7165bool SPIRVInstructionSelector::loadHandleBeforePosition(
7166 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
7167 MachineInstr &Pos)
const {
7170 Intrinsic::spv_resource_handlefrombinding);
7178 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
7179 MachineIRBuilder MIRBuilder(HandleDef);
7180 SPIRVTypeInst VarType = ResType;
7181 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
7183 if (IsStructuredBuffer) {
7188 if (ResType->
getOpcode() == SPIRV::OpTypeImage && ArraySize == 0)
7190 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
7193 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
7194 ArraySize, IndexReg, Name, MIRBuilder);
7198 uint32_t LoadOpcode =
7199 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
7209bool SPIRVInstructionSelector::errorIfInstrOutsideShader(
7210 MachineInstr &
I)
const {
7212 return diagnoseUnsupported(
7213 I,
"this instruction is only supported in shaders.");
7218InstructionSelector *
7222 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
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.
void buildOpName(Register Target, const StringRef &Name, MachineIRBuilder &MIRBuilder)
@ Low
Lower the current thread's priority such that it does not affect foreground tasks significantly.
FunctionAddr VTableAddr Value
bool all_of(R &&range, UnaryPredicate P)
Provide wrappers to std::all_of which take ranges instead of having to pass begin/end explicitly.
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.
FunctionAddr VTableAddr uintptr_t uintptr_t Int32Ty
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)
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)
bool any_of(R &&range, UnaryPredicate P)
Provide wrappers to std::any_of which take ranges instead of having to pass begin/end explicitly.
void buildOpDecorate(Register Reg, MachineIRBuilder &MIRBuilder, SPIRV::Decoration::Decoration Dec, const std::vector< uint32_t > &DecArgs, StringRef StrImm)
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)
void addStringImm(const StringRef &Str, MCInst &Inst)
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...