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");
4074 for (
unsigned i =
I.getNumExplicitDefs();
4075 i <
I.getNumExplicitOperands() && IsNullVector; ++i) {
4076 MachineInstr *
Def =
getDef(
I.getOperand(i), MRI);
4081 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4088 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4089 TII.get(IsConst ? SPIRV::OpConstantComposite
4090 : SPIRV::OpCompositeConstruct))
4093 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
4094 MIB.
addUse(
I.getOperand(i).getReg());
4099bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
4100 SPIRVTypeInst ResType,
4101 MachineInstr &
I)
const {
4103 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4105 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4111 if (!
I.getOperand(
OpIdx).isReg())
4118 if (!IsConst &&
N < 2)
4119 return diagnoseUnsupported(
4120 I,
"There must be at least two constituent operands in a vector");
4123 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4124 TII.get(IsConst ? SPIRV::OpConstantComposite
4125 : SPIRV::OpCompositeConstruct))
4128 for (
unsigned i = 0; i <
N; ++i)
4134bool SPIRVInstructionSelector::selectConcatVectors(
Register ResVReg,
4135 SPIRVTypeInst ResType,
4136 MachineInstr &
I)
const {
4140 if (ResType->
getOpcode() != SPIRV::OpTypeVector)
4142 "Cannot select G_CONCAT_VECTORS with a non-vector result");
4144 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4145 TII.get(SPIRV::OpCompositeConstruct))
4155bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
4156 SPIRVTypeInst ResType,
4157 MachineInstr &
I)
const {
4162 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
4164 Opcode = SPIRV::OpDemoteToHelperInvocation;
4166 Opcode = SPIRV::OpKill;
4168 if (MachineInstr *NextI =
I.getNextNode()) {
4170 NextI->eraseFromParent();
4180bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
4181 SPIRVTypeInst ResType,
unsigned CmpOpc,
4182 MachineInstr &
I)
const {
4183 Register Cmp0 =
I.getOperand(2).getReg();
4184 Register Cmp1 =
I.getOperand(3).getReg();
4187 "CMP operands should have the same type");
4188 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
4198bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
4199 SPIRVTypeInst ResType,
4200 MachineInstr &
I)
const {
4201 auto Pred =
I.getOperand(1).getPredicate();
4204 Register CmpOperand =
I.getOperand(2).getReg();
4209 Register Op1 =
I.getOperand(3).getReg();
4213 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
4218 I.getOperand(3).setReg(NewOp1);
4224 return selectCmp(ResVReg, ResType, CmpOpc,
I);
4228SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
4229 SPIRVTypeInst ResType)
const {
4231 SPIRVTypeInst SpvI32Ty =
4234 auto ConstInt = ConstantInt::get(LLVMTy, Val);
4241 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4244 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
4247 .
addImm(APInt(32, Val).getZExtValue());
4249 GR.
add(ConstInt,
MI);
4256Register SPIRVInstructionSelector::buildI32ConstantInEntryBlock(
4257 uint32_t Val, MachineInstr &
I, SPIRVTypeInst ResType)
const {
4259 SPIRVTypeInst SpvI32Ty =
4261 auto *ConstInt = ConstantInt::get(LLVMTy, Val);
4266 MachineBasicBlock &EntryBB = *InsertIt->getParent();
4267 MachineInstr *
MI =
nullptr;
4271 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantNull))
4275 uint64_t ImmVal = APInt(32, Val).getZExtValue();
4276 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantI))
4282 GR.
add(ConstInt,
MI);
4287bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
4288 SPIRVTypeInst ResType,
4289 MachineInstr &
I)
const {
4291 return selectCmp(ResVReg, ResType, CmpOp,
I);
4294bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
4295 SPIRVTypeInst ResType,
4296 MachineInstr &
I)
const {
4298 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
4305 if (ResType->
getOpcode() != SPIRV::OpTypeVector &&
4306 ResType->
getOpcode() != SPIRV::OpTypeFloat)
4309 MachineIRBuilder MIRBuilder(
I);
4316 APFloat ConstVal(3.3219280948873623);
4320 APFloat::rmNearestTiesToEven, &LosesInfo);
4324 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
4325 ? SPIRV::OpVectorTimesScalar
4328 if (!selectOpWithSrcs(ArgReg, ResType,
I,
4329 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
4331 if (!selectExtInst(ResVReg, ResType,
I,
4332 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
4342Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
4343 MachineInstr &
I)
const {
4346 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4351bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
4357 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
4365 if (
Def->getOpcode() == SPIRV::OpConstantNull)
4368 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
4369 Def->getOpcode() == SPIRV::OpConstantI)
4382 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
4383 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
4385 Intrinsic::spv_const_composite)) {
4386 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
4387 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
4388 if (!IsZero(
Def->getOperand(i).getReg()))
4397Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
4398 MachineInstr &
I)
const {
4402 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4407Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
4408 MachineInstr &
I)
const {
4412 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4418 SPIRVTypeInst ResType,
4419 MachineInstr &
I)
const {
4423 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4428bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
4429 SPIRVTypeInst ResType,
4430 MachineInstr &
I)
const {
4431 Register SelectFirstArg =
I.getOperand(2).getReg();
4432 Register SelectSecondArg =
I.getOperand(3).getReg();
4441 SPIRV::OpTypeVector;
4448 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
4449 }
else if (IsPtrTy) {
4450 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
4452 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
4455 assert(IsScalarBool &&
"OpSelect with a scalar result requires a scalar "
4456 "boolean condition");
4458 Opcode = SPIRV::OpSelectSFSCond;
4459 }
else if (IsPtrTy) {
4460 Opcode = SPIRV::OpSelectSPSCond;
4462 Opcode = SPIRV::OpSelectSISCond;
4465 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4468 .
addUse(
I.getOperand(1).getReg())
4477bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
4478 SPIRVTypeInst ResType,
4480 MachineInstr &InsertAt,
4481 bool IsSigned)
const {
4483 Register ZeroReg = buildZerosVal(ResType, InsertAt);
4484 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
4485 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
4487 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4499bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
4500 SPIRVTypeInst ResType,
4501 MachineInstr &
I,
bool IsSigned,
4502 unsigned Opcode)
const {
4503 Register SrcReg =
I.getOperand(1).getReg();
4509 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
4514 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I,
false);
4516 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
4519bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
4520 SPIRVTypeInst ResType, MachineInstr &
I,
4521 bool IsSigned)
const {
4522 Register SrcReg =
I.getOperand(1).getReg();
4524 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
4528 if (ResType == SrcType)
4529 return BuildCOPY(ResVReg, SrcReg,
I);
4531 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4532 return selectUnOp(ResVReg, ResType,
I, Opcode);
4535bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
4536 SPIRVTypeInst ResType,
4538 bool IsSigned)
const {
4539 MachineIRBuilder MIRBuilder(
I);
4540 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
4552 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
4555 .
addUse(
I.getOperand(1).getReg())
4556 .
addUse(
I.getOperand(2).getReg())
4561 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
4564 .
addUse(
I.getOperand(1).getReg())
4565 .
addUse(
I.getOperand(2).getReg())
4573 unsigned SelectOpcode =
4574 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
4579 .
addUse(buildOnesVal(
true, ResType,
I))
4580 .
addUse(buildZerosVal(ResType,
I))
4587 .
addUse(buildOnesVal(
false, ResType,
I))
4592bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
4595 SPIRVTypeInst IntTy,
4596 SPIRVTypeInst BoolTy)
const {
4599 bool IsVectorTy = IntTy->
getOpcode() == SPIRV::OpTypeVector;
4600 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
4602 Register One = buildOnesVal(
false, IntTy,
I);
4610 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
4619bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
4620 SPIRVTypeInst ResType,
4621 MachineInstr &
I)
const {
4622 Register IntReg =
I.getOperand(1).getReg();
4625 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
4626 if (ArgType == ResType)
4627 return BuildCOPY(ResVReg, IntReg,
I);
4629 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4630 return selectUnOp(ResVReg, ResType,
I, Opcode);
4633bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
4634 SPIRVTypeInst ResType,
4635 MachineInstr &
I)
const {
4636 unsigned Opcode =
I.getOpcode();
4637 unsigned TpOpcode = ResType->
getOpcode();
4639 if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {
4640 assert(Opcode == TargetOpcode::G_CONSTANT &&
4641 I.getOperand(1).getCImm()->isZero());
4642 MachineBasicBlock &DepMBB =
I.getMF()->front();
4645 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
4652 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
4655bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
4656 SPIRVTypeInst ResType,
4657 MachineInstr &
I)
const {
4658 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4665bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
4666 SPIRVTypeInst ResType,
4667 MachineInstr &
I)
const {
4669 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
4673 .
addUse(
I.getOperand(3).getReg())
4675 .
addUse(
I.getOperand(2).getReg());
4676 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
4682bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
4683 SPIRVTypeInst ResType,
4684 MachineInstr &
I)
const {
4685 Type *MaybeResTy =
nullptr;
4690 "Expected aggregate type for extractv instruction");
4692 SPIRV::AccessQualifier::ReadWrite,
false);
4696 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
4699 .
addUse(
I.getOperand(2).getReg());
4700 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
4706bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
4707 SPIRVTypeInst ResType,
4708 MachineInstr &
I)
const {
4709 if (
getImm(
I.getOperand(4), MRI))
4710 return selectInsertVal(ResVReg, ResType,
I);
4712 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
4715 .
addUse(
I.getOperand(2).getReg())
4716 .
addUse(
I.getOperand(3).getReg())
4717 .
addUse(
I.getOperand(4).getReg())
4722bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
4723 SPIRVTypeInst ResType,
4724 MachineInstr &
I)
const {
4725 if (
getImm(
I.getOperand(3), MRI))
4726 return selectExtractVal(ResVReg, ResType,
I);
4728 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
4731 .
addUse(
I.getOperand(2).getReg())
4732 .
addUse(
I.getOperand(3).getReg())
4737bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
4738 SPIRVTypeInst ResType,
4739 MachineInstr &
I)
const {
4740 const bool IsGEPInBounds =
I.getOperand(2).getImm();
4746 ? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
4747 : SPIRV::OpAccessChain)
4748 : (IsGEPInBounds ?
SPIRV::OpInBoundsPtrAccessChain
4749 :
SPIRV::OpPtrAccessChain);
4751 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4755 .
addUse(
I.getOperand(3).getReg());
4757 (Opcode == SPIRV::OpPtrAccessChain ||
4758 Opcode == SPIRV::OpInBoundsPtrAccessChain ||
4759 (
getImm(
I.getOperand(4), MRI) &&
foldImm(
I.getOperand(4), MRI) == 0)) &&
4760 "Cannot translate GEP to OpAccessChain. First index must be 0.");
4763 const unsigned StartingIndex =
4764 (Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)
4767 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
4768 Res.addUse(
I.getOperand(i).getReg());
4769 Res.constrainAllUses(
TII,
TRI, RBI);
4774bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
4776 unsigned Lim =
I.getNumExplicitOperands();
4777 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
4778 Register OpReg =
I.getOperand(i).getReg();
4779 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
4781 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
4782 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
4783 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
4790 MachineFunction *MF =
I.getMF();
4802 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4803 TII.get(SPIRV::OpSpecConstantOp))
4806 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
4808 GR.
add(OpDefine, MIB);
4814bool SPIRVInstructionSelector::selectDerivativeInst(
4815 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
4816 const unsigned DPdOpCode)
const {
4819 if (!errorIfInstrOutsideShader(
I))
4825 Register SrcReg =
I.getOperand(2).getReg();
4830 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4833 .
addUse(
I.getOperand(2).getReg());
4835 MachineIRBuilder MIRBuilder(
I);
4838 if (componentCount != 1)
4842 const TargetRegisterClass *RegClass = GR.
getRegClass(SrcType);
4846 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4851 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4856 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4864bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
4865 SPIRVTypeInst ResType,
4866 MachineInstr &
I)
const {
4870 case Intrinsic::spv_load:
4871 return selectLoad(ResVReg, ResType,
I);
4872 case Intrinsic::spv_atomic_load:
4873 return selectAtomicLoad(ResVReg, ResType,
I);
4874 case Intrinsic::spv_store:
4875 return selectStore(
I);
4876 case Intrinsic::spv_atomic_store:
4877 return selectAtomicStore(
I);
4878 case Intrinsic::spv_extractv:
4879 return selectExtractVal(ResVReg, ResType,
I);
4880 case Intrinsic::spv_insertv:
4881 return selectInsertVal(ResVReg, ResType,
I);
4882 case Intrinsic::spv_extractelt:
4883 return selectExtractElt(ResVReg, ResType,
I);
4884 case Intrinsic::spv_insertelt:
4885 return selectInsertElt(ResVReg, ResType,
I);
4886 case Intrinsic::spv_gep:
4887 return selectGEP(ResVReg, ResType,
I);
4888 case Intrinsic::spv_bitcast: {
4889 Register OpReg =
I.getOperand(2).getReg();
4890 SPIRVTypeInst OpType =
4894 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
4896 case Intrinsic::spv_unref_global:
4897 case Intrinsic::spv_init_global: {
4898 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
4903 Register GVarVReg =
MI->getOperand(0).getReg();
4904 if (!selectGlobalValue(GVarVReg, *
MI, Init))
4909 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
4911 MI->eraseFromParent();
4915 case Intrinsic::spv_undef: {
4916 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4922 case Intrinsic::spv_poison:
4923 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpPoisonKHR))
4928 case Intrinsic::spv_freeze:
4929 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
4932 .
addUse(
I.getOperand(2).getReg())
4935 case Intrinsic::spv_named_boolean_spec_constant: {
4936 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
4937 : SPIRV::OpSpecConstantFalse;
4939 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
4940 .
addDef(
I.getOperand(0).getReg())
4943 unsigned SpecId =
I.getOperand(2).getImm();
4945 SPIRV::Decoration::SpecId, {SpecId});
4949 case Intrinsic::spv_const_composite: {
4951 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
4957 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
4959 std::function<bool(
Register)> HasSpecConstOperand =
4969 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
4970 J < Def->getNumExplicitOperands(); ++J) {
4971 if (
Def->getOperand(J).isReg() &&
4972 HasSpecConstOperand(
Def->getOperand(J).getReg()))
4978 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
4979 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
4980 : SPIRV::OpConstantComposite;
4981 unsigned ContinuedOpc = HasSpecConst
4982 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
4983 : SPIRV::OpConstantCompositeContinuedINTEL;
4984 MachineIRBuilder MIR(
I);
4986 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
4988 for (
auto *Instr : Instructions) {
4989 Instr->setDebugLoc(
I.getDebugLoc());
4994 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5001 case Intrinsic::spv_assign_name: {
5002 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
5003 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
5004 for (
unsigned i =
I.getNumExplicitDefs() + 2;
5005 i <
I.getNumExplicitOperands(); ++i) {
5006 MIB.
addImm(
I.getOperand(i).getImm());
5011 case Intrinsic::spv_switch: {
5012 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
5013 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5014 if (
I.getOperand(i).isReg())
5015 MIB.
addReg(
I.getOperand(i).getReg());
5016 else if (
I.getOperand(i).isCImm())
5017 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
5018 else if (
I.getOperand(i).isMBB())
5019 MIB.
addMBB(
I.getOperand(i).getMBB());
5026 case Intrinsic::spv_loop_merge: {
5027 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
5028 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5029 if (
I.getOperand(i).isMBB())
5030 MIB.
addMBB(
I.getOperand(i).getMBB());
5037 case Intrinsic::spv_loop_control_intel: {
5039 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
5040 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
5045 case Intrinsic::spv_selection_merge: {
5047 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
5048 assert(
I.getOperand(1).isMBB() &&
5049 "operand 1 to spv_selection_merge must be a basic block");
5050 MIB.
addMBB(
I.getOperand(1).getMBB());
5051 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
5055 case Intrinsic::spv_cmpxchg:
5056 return selectAtomicCmpXchg(ResVReg, ResType,
I);
5057 case Intrinsic::spv_unreachable:
5058 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
5061 case Intrinsic::spv_abort:
5062 return selectAbort(
I);
5063 case Intrinsic::spv_alloca:
5064 return selectFrameIndex(ResVReg, ResType,
I);
5065 case Intrinsic::spv_alloca_array:
5066 return selectAllocaArray(ResVReg, ResType,
I);
5067 case Intrinsic::spv_assume:
5069 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
5070 .
addUse(
I.getOperand(1).getReg())
5075 case Intrinsic::spv_expect:
5077 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
5080 .
addUse(
I.getOperand(2).getReg())
5081 .
addUse(
I.getOperand(3).getReg())
5086 case Intrinsic::arithmetic_fence:
5087 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
5088 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
5091 .
addUse(
I.getOperand(2).getReg())
5095 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
5097 case Intrinsic::spv_thread_id:
5103 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
5105 case Intrinsic::spv_thread_id_in_group:
5111 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
5113 case Intrinsic::spv_group_id:
5119 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
5121 case Intrinsic::spv_flattened_thread_id_in_group:
5128 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
5130 case Intrinsic::spv_workgroup_size:
5131 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
5133 case Intrinsic::spv_global_size:
5134 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
5136 case Intrinsic::spv_global_offset:
5137 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
5139 case Intrinsic::spv_num_workgroups:
5140 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
5142 case Intrinsic::spv_subgroup_size:
5143 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
5145 case Intrinsic::spv_num_subgroups:
5146 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
5148 case Intrinsic::spv_subgroup_id:
5149 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
5150 case Intrinsic::spv_subgroup_local_invocation_id:
5151 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
5152 ResVReg, ResType,
I);
5153 case Intrinsic::spv_subgroup_max_size:
5154 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
5156 case Intrinsic::spv_fdot:
5157 return selectFloatDot(ResVReg, ResType,
I);
5158 case Intrinsic::spv_udot:
5159 case Intrinsic::spv_sdot:
5160 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5162 return selectIntegerDot(ResVReg, ResType,
I,
5163 IID == Intrinsic::spv_sdot);
5164 return selectIntegerDotExpansion(ResVReg, ResType,
I);
5165 case Intrinsic::spv_dot4add_i8packed:
5166 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5168 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
5169 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
5170 case Intrinsic::spv_dot4add_u8packed:
5171 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5173 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
5174 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
5175 case Intrinsic::spv_all:
5176 return selectAll(ResVReg, ResType,
I);
5177 case Intrinsic::spv_any:
5178 return selectAny(ResVReg, ResType,
I);
5179 case Intrinsic::spv_cross:
5180 return selectExtInst(ResVReg, ResType,
I, CL::cross, GL::Cross);
5181 case Intrinsic::spv_distance:
5182 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
5183 case Intrinsic::spv_lerp:
5184 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
5185 case Intrinsic::spv_length:
5186 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
5187 case Intrinsic::spv_degrees:
5188 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
5189 case Intrinsic::spv_faceforward:
5190 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
5191 case Intrinsic::spv_frac:
5192 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
5193 case Intrinsic::spv_isinf:
5194 return selectOpIsInf(ResVReg, ResType,
I);
5195 case Intrinsic::spv_isnan:
5196 return selectOpIsNan(ResVReg, ResType,
I);
5197 case Intrinsic::spv_isfinite:
5198 return selectOpIsFinite(ResVReg, ResType,
I);
5199 case Intrinsic::spv_isnormal:
5200 return selectOpIsNormal(ResVReg, ResType,
I);
5201 case Intrinsic::spv_normalize:
5202 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
5203 case Intrinsic::spv_refract:
5204 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
5205 case Intrinsic::spv_reflect:
5206 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
5207 case Intrinsic::spv_rsqrt:
5208 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
5209 case Intrinsic::spv_sign:
5210 return selectSign(ResVReg, ResType,
I);
5211 case Intrinsic::spv_smoothstep:
5212 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
5213 case Intrinsic::spv_firstbituhigh:
5214 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
5215 case Intrinsic::spv_firstbitshigh:
5216 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
5217 case Intrinsic::spv_firstbitlow:
5218 return selectFirstBitLow(ResVReg, ResType,
I);
5219 case Intrinsic::spv_all_memory_barrier:
5220 return selectBarrierInst(
I, SPIRV::Scope::Device,
5221 SPIRV::MemorySemantics::UniformMemory |
5222 SPIRV::MemorySemantics::ImageMemory |
5223 SPIRV::MemorySemantics::WorkgroupMemory,
5225 case Intrinsic::spv_all_memory_barrier_with_group_sync:
5226 return selectBarrierInst(
I, SPIRV::Scope::Device,
5227 SPIRV::MemorySemantics::UniformMemory |
5228 SPIRV::MemorySemantics::ImageMemory |
5229 SPIRV::MemorySemantics::WorkgroupMemory,
5231 case Intrinsic::spv_device_memory_barrier:
5232 return selectBarrierInst(
I, SPIRV::Scope::Device,
5233 SPIRV::MemorySemantics::UniformMemory |
5234 SPIRV::MemorySemantics::ImageMemory,
5236 case Intrinsic::spv_device_memory_barrier_with_group_sync:
5237 return selectBarrierInst(
I, SPIRV::Scope::Device,
5238 SPIRV::MemorySemantics::UniformMemory |
5239 SPIRV::MemorySemantics::ImageMemory,
5241 case Intrinsic::spv_group_memory_barrier:
5242 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5243 SPIRV::MemorySemantics::WorkgroupMemory,
5245 case Intrinsic::spv_group_memory_barrier_with_group_sync:
5246 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5247 SPIRV::MemorySemantics::WorkgroupMemory,
5249 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
5250 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
5251 SPIRV::StorageClass::StorageClass ResSC =
5254 return diagnoseUnsupported(
I,
"The target storage class is not castable "
5255 "from the Generic storage class");
5256 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
5264 case Intrinsic::spv_lifetime_start:
5265 case Intrinsic::spv_lifetime_end: {
5266 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
5267 : SPIRV::OpLifetimeStop;
5268 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
5269 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
5278 case Intrinsic::spv_saturate:
5279 return selectSaturate(ResVReg, ResType,
I);
5280 case Intrinsic::spv_nclamp:
5281 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
5282 case Intrinsic::spv_uclamp:
5283 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
5284 case Intrinsic::spv_sclamp:
5285 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
5286 case Intrinsic::spv_subgroup_prefix_bit_count:
5287 return selectWavePrefixBitCount(ResVReg, ResType,
I);
5288 case Intrinsic::spv_wave_active_countbits:
5289 return selectWaveActiveCountBits(ResVReg, ResType,
I);
5290 case Intrinsic::spv_wave_all_equal:
5291 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
5292 case Intrinsic::spv_wave_all:
5293 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
5294 case Intrinsic::spv_wave_any:
5295 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
5296 case Intrinsic::spv_subgroup_ballot:
5297 return selectWaveOpInst(ResVReg, ResType,
I,
5298 SPIRV::OpGroupNonUniformBallot);
5299 case Intrinsic::spv_wave_is_first_lane:
5300 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
5301 case Intrinsic::spv_wave_reduce_or:
5302 return selectWaveReduceOp(ResVReg, ResType,
I,
5303 SPIRV::OpGroupNonUniformBitwiseOr);
5304 case Intrinsic::spv_wave_reduce_xor:
5305 return selectWaveReduceOp(ResVReg, ResType,
I,
5306 SPIRV::OpGroupNonUniformBitwiseXor);
5307 case Intrinsic::spv_wave_reduce_and:
5308 return selectWaveReduceOp(ResVReg, ResType,
I,
5309 SPIRV::OpGroupNonUniformBitwiseAnd);
5310 case Intrinsic::spv_interlocked_add:
5311 return selectInterlockedAdd(ResVReg, ResType,
I);
5312 case Intrinsic::spv_wave_reduce_umax:
5313 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
5314 case Intrinsic::spv_wave_reduce_max:
5315 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
5316 case Intrinsic::spv_wave_reduce_umin:
5317 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
5318 case Intrinsic::spv_wave_reduce_min:
5319 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
5320 case Intrinsic::spv_wave_reduce_sum:
5321 return selectWaveReduceSum(ResVReg, ResType,
I);
5322 case Intrinsic::spv_wave_product:
5323 return selectWaveReduceProduct(ResVReg, ResType,
I);
5324 case Intrinsic::spv_wave_readlane:
5325 return selectWaveOpInst(ResVReg, ResType,
I,
5326 SPIRV::OpGroupNonUniformShuffle);
5327 case Intrinsic::spv_wave_prefix_sum:
5328 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
5329 case Intrinsic::spv_wave_prefix_product:
5330 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
5331 case Intrinsic::spv_quad_read_across_x: {
5332 return selectQuadSwap(ResVReg, ResType,
I, 0);
5334 case Intrinsic::spv_quad_read_across_y: {
5335 return selectQuadSwap(ResVReg, ResType,
I, 1);
5337 case Intrinsic::spv_step:
5338 return selectExtInst(ResVReg, ResType,
I, CL::step, GL::Step);
5339 case Intrinsic::spv_radians:
5340 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
5344 case Intrinsic::instrprof_increment:
5345 case Intrinsic::instrprof_increment_step:
5346 case Intrinsic::instrprof_value_profile:
5349 case Intrinsic::spv_value_md:
5351 case Intrinsic::spv_resource_handlefrombinding: {
5352 return selectHandleFromBinding(ResVReg, ResType,
I);
5354 case Intrinsic::spv_resource_counterhandlefrombinding:
5355 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
5356 case Intrinsic::spv_resource_updatecounter:
5357 return selectUpdateCounter(ResVReg, ResType,
I);
5358 case Intrinsic::spv_resource_store_typedbuffer: {
5359 return selectImageWriteIntrinsic(
I);
5361 case Intrinsic::spv_resource_load_typedbuffer: {
5362 return selectReadImageIntrinsic(ResVReg, ResType,
I);
5364 case Intrinsic::spv_resource_load_level: {
5365 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
5367 case Intrinsic::spv_resource_getdimensions_x:
5368 case Intrinsic::spv_resource_getdimensions_xy:
5369 case Intrinsic::spv_resource_getdimensions_xyz: {
5370 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
5372 case Intrinsic::spv_resource_getdimensions_levels_x:
5373 case Intrinsic::spv_resource_getdimensions_levels_xy:
5374 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
5375 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
5377 case Intrinsic::spv_resource_getdimensions_ms_xy:
5378 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
5379 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
5381 case Intrinsic::spv_resource_calculate_lod:
5382 case Intrinsic::spv_resource_calculate_lod_unclamped:
5383 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
5384 case Intrinsic::spv_resource_sample:
5385 case Intrinsic::spv_resource_sample_clamp:
5386 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
5387 case Intrinsic::spv_resource_samplebias:
5388 case Intrinsic::spv_resource_samplebias_clamp:
5389 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
5390 case Intrinsic::spv_resource_samplegrad:
5391 case Intrinsic::spv_resource_samplegrad_clamp:
5392 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
5393 case Intrinsic::spv_resource_samplelevel:
5394 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
5395 case Intrinsic::spv_resource_samplecmp:
5396 case Intrinsic::spv_resource_samplecmp_clamp:
5397 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
5398 case Intrinsic::spv_resource_samplecmplevelzero:
5399 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
5400 case Intrinsic::spv_resource_gather:
5401 case Intrinsic::spv_resource_gather_cmp:
5402 return selectGatherIntrinsic(ResVReg, ResType,
I);
5403 case Intrinsic::spv_resource_getbasepointer:
5404 case Intrinsic::spv_resource_getpointer: {
5405 return selectResourceGetPointer(ResVReg, ResType,
I);
5407 case Intrinsic::spv_pushconstant_getpointer: {
5408 return selectPushConstantGetPointer(ResVReg, ResType,
I);
5410 case Intrinsic::spv_discard: {
5411 return selectDiscard(ResVReg, ResType,
I);
5413 case Intrinsic::spv_resource_nonuniformindex: {
5414 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
5416 case Intrinsic::spv_unpackhalf2x16: {
5417 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
5419 case Intrinsic::spv_packhalf2x16: {
5420 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
5422 case Intrinsic::spv_ddx:
5423 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
5424 case Intrinsic::spv_ddy:
5425 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
5426 case Intrinsic::spv_ddx_coarse:
5427 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
5428 case Intrinsic::spv_ddy_coarse:
5429 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
5430 case Intrinsic::spv_ddx_fine:
5431 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
5432 case Intrinsic::spv_ddy_fine:
5433 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
5434 case Intrinsic::spv_fwidth:
5435 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
5436 case Intrinsic::spv_masked_gather:
5437 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5438 return selectMaskedGather(ResVReg, ResType,
I);
5439 return diagnoseUnsupported(
5440 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
5441 case Intrinsic::spv_masked_scatter:
5442 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5443 return selectMaskedScatter(
I);
5444 return diagnoseUnsupported(
5445 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
5446 case Intrinsic::returnaddress:
5447 case Intrinsic::frameaddress: {
5449 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5456 return diagnoseUnsupported(
I,
"intrinsic selection not implemented.");
5461bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
5462 SPIRVTypeInst ResType,
5463 MachineInstr &
I)
const {
5466 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
5473bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
5474 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5476 assert(Intr.getIntrinsicID() ==
5477 Intrinsic::spv_resource_counterhandlefrombinding);
5480 Register MainHandleReg = Intr.getOperand(2).getReg();
5482 assert(MainHandleDef->getIntrinsicID() ==
5483 Intrinsic::spv_resource_handlefrombinding);
5487 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
5488 Register IndexReg = MainHandleDef->getOperand(5).getReg();
5489 std::string CounterName =
5494 MachineIRBuilder MIRBuilder(
I);
5496 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
5498 ArraySize, IndexReg, CounterName, MIRBuilder);
5500 return BuildCOPY(ResVReg, CounterVarReg,
I);
5503bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
5504 SPIRVTypeInst ResType,
5505 MachineInstr &
I)
const {
5507 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
5509 Register CounterHandleReg = Intr.getOperand(2).getReg();
5510 Register IncrReg = Intr.getOperand(3).getReg();
5517 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
5518 assert(CounterVarPointeeType &&
5519 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
5520 "Counter variable must be a struct");
5522 SPIRV::StorageClass::StorageBuffer &&
5523 "Counter variable must be in the storage buffer storage class");
5525 "Counter variable must have exactly 1 member in the struct");
5526 const SPIRVTypeInst MemberType =
5529 "Counter variable struct must have a single i32 member");
5533 MachineIRBuilder MIRBuilder(
I);
5535 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
5538 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
5544 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5547 .
addUse(CounterHandleReg)
5554 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
5557 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
5560 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
5569 return BuildCOPY(ResVReg, AtomicRes,
I);
5577 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
5585bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
5586 SPIRVTypeInst ResType,
5587 MachineInstr &
I)
const {
5595 Register ImageReg =
I.getOperand(2).getReg();
5603 Register IdxReg =
I.getOperand(3).getReg();
5605 MachineInstr &Pos =
I;
5607 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
5611bool SPIRVInstructionSelector::generateSampleImage(
5614 DebugLoc Loc, MachineInstr &Pos)
const {
5625 if (!loadHandleBeforePosition(NewSamplerReg,
5631 MachineIRBuilder MIRBuilder(Pos);
5644 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
5645 ImOps.Lod.has_value();
5646 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
5647 : SPIRV::OpImageSampleImplicitLod;
5649 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
5650 : SPIRV::OpImageSampleDrefImplicitLod;
5659 MIB.
addUse(*ImOps.Compare);
5661 uint32_t ImageOperands = 0;
5663 ImageOperands |= SPIRV::ImageOperand::Bias;
5665 ImageOperands |= SPIRV::ImageOperand::Lod;
5666 if (ImOps.GradX && ImOps.GradY)
5667 ImageOperands |= SPIRV::ImageOperand::Grad;
5668 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
5670 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
5673 "Non-constant offsets are not supported in sample instructions.");
5677 ImageOperands |= SPIRV::ImageOperand::MinLod;
5679 if (ImageOperands != 0) {
5680 MIB.
addImm(ImageOperands);
5681 if (ImageOperands & SPIRV::ImageOperand::Bias)
5683 if (ImageOperands & SPIRV::ImageOperand::Lod)
5685 if (ImageOperands & SPIRV::ImageOperand::Grad) {
5686 MIB.
addUse(*ImOps.GradX);
5687 MIB.
addUse(*ImOps.GradY);
5690 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
5691 MIB.
addUse(*ImOps.Offset);
5692 if (ImageOperands & SPIRV::ImageOperand::MinLod)
5693 MIB.
addUse(*ImOps.MinLod);
5700bool SPIRVInstructionSelector::selectImageQuerySize(
5702 std::optional<Register> LodReg)
const {
5704 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
5707 "ImageReg is not an image type.");
5709 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
5711 unsigned NumComponents = 0;
5713 case SPIRV::Dim::DIM_1D:
5714 case SPIRV::Dim::DIM_Buffer:
5715 NumComponents =
IsArray ? 2 : 1;
5717 case SPIRV::Dim::DIM_2D:
5718 case SPIRV::Dim::DIM_Cube:
5719 case SPIRV::Dim::DIM_Rect:
5720 NumComponents =
IsArray ? 3 : 2;
5722 case SPIRV::Dim::DIM_3D:
5726 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
5731 SPIRVTypeInst ResType =
5736 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5746bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
5747 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5748 Register ImageReg =
I.getOperand(2).getReg();
5755 return selectImageQuerySize(NewImageReg, ResVReg,
I);
5758bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
5759 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5760 Register ImageReg =
I.getOperand(2).getReg();
5769 Register LodReg =
I.getOperand(3).getReg();
5772 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
5774 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
5781 TII.get(SPIRV::OpImageQueryLevels))
5788 TII.get(SPIRV::OpCompositeConstruct))
5798bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
5799 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5800 Register ImageReg =
I.getOperand(2).getReg();
5811 "OpImageQuerySamples requires a multisampled image");
5813 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
5821 TII.get(SPIRV::OpImageQuerySamples))
5828 TII.get(SPIRV::OpCompositeConstruct))
5838bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
5839 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5840 Register ImageReg =
I.getOperand(2).getReg();
5841 Register SamplerReg =
I.getOperand(3).getReg();
5842 Register CoordinateReg =
I.getOperand(4).getReg();
5858 if (!loadHandleBeforePosition(
5863 MachineIRBuilder MIRBuilder(
I);
5869 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
5879 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
5886 unsigned ExtractedIndex =
5888 Intrinsic::spv_resource_calculate_lod_unclamped
5892 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5893 TII.get(SPIRV::OpCompositeExtract))
5903bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
5904 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5905 Register ImageReg =
I.getOperand(2).getReg();
5906 Register SamplerReg =
I.getOperand(3).getReg();
5907 Register CoordinateReg =
I.getOperand(4).getReg();
5908 ImageOperands ImOps;
5909 if (
I.getNumOperands() > 5)
5910 ImOps.Offset =
I.getOperand(5).getReg();
5911 if (
I.getNumOperands() > 6)
5912 ImOps.MinLod =
I.getOperand(6).getReg();
5913 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5914 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5917bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
5918 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5919 Register ImageReg =
I.getOperand(2).getReg();
5920 Register SamplerReg =
I.getOperand(3).getReg();
5921 Register CoordinateReg =
I.getOperand(4).getReg();
5922 ImageOperands ImOps;
5923 ImOps.Bias =
I.getOperand(5).getReg();
5924 if (
I.getNumOperands() > 6)
5925 ImOps.Offset =
I.getOperand(6).getReg();
5926 if (
I.getNumOperands() > 7)
5927 ImOps.MinLod =
I.getOperand(7).getReg();
5928 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5929 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5932bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
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.GradX =
I.getOperand(5).getReg();
5939 ImOps.GradY =
I.getOperand(6).getReg();
5940 if (
I.getNumOperands() > 7)
5941 ImOps.Offset =
I.getOperand(7).getReg();
5942 if (
I.getNumOperands() > 8)
5943 ImOps.MinLod =
I.getOperand(8).getReg();
5944 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5945 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5948bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
5949 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5950 Register ImageReg =
I.getOperand(2).getReg();
5951 Register SamplerReg =
I.getOperand(3).getReg();
5952 Register CoordinateReg =
I.getOperand(4).getReg();
5953 ImageOperands ImOps;
5954 ImOps.Lod =
I.getOperand(5).getReg();
5955 if (
I.getNumOperands() > 6)
5956 ImOps.Offset =
I.getOperand(6).getReg();
5957 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5958 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5961bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
5962 SPIRVTypeInst ResType,
5963 MachineInstr &
I)
const {
5964 Register ImageReg =
I.getOperand(2).getReg();
5965 Register SamplerReg =
I.getOperand(3).getReg();
5966 Register CoordinateReg =
I.getOperand(4).getReg();
5967 ImageOperands ImOps;
5968 ImOps.Compare =
I.getOperand(5).getReg();
5969 if (
I.getNumOperands() > 6)
5970 ImOps.Offset =
I.getOperand(6).getReg();
5971 if (
I.getNumOperands() > 7)
5972 ImOps.MinLod =
I.getOperand(7).getReg();
5973 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5974 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5977bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
5978 SPIRVTypeInst ResType,
5979 MachineInstr &
I)
const {
5980 Register ImageReg =
I.getOperand(2).getReg();
5981 Register CoordinateReg =
I.getOperand(3).getReg();
5982 Register LodReg =
I.getOperand(4).getReg();
5984 ImageOperands ImOps;
5986 if (
I.getNumOperands() > 5)
5987 ImOps.Offset =
I.getOperand(5).getReg();
5999 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
6000 I.getDebugLoc(),
I, &ImOps);
6003bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
6004 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6005 Register ImageReg =
I.getOperand(2).getReg();
6006 Register SamplerReg =
I.getOperand(3).getReg();
6007 Register CoordinateReg =
I.getOperand(4).getReg();
6008 ImageOperands ImOps;
6009 ImOps.Compare =
I.getOperand(5).getReg();
6010 if (
I.getNumOperands() > 6)
6011 ImOps.Offset =
I.getOperand(6).getReg();
6014 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6015 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6018bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
6019 SPIRVTypeInst ResType,
6020 MachineInstr &
I)
const {
6021 Register ImageReg =
I.getOperand(2).getReg();
6022 Register SamplerReg =
I.getOperand(3).getReg();
6023 Register CoordinateReg =
I.getOperand(4).getReg();
6026 "ImageReg is not an image type.");
6031 ComponentOrCompareReg =
I.getOperand(5).getReg();
6032 OffsetReg =
I.getOperand(6).getReg();
6035 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
6039 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6040 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
6041 Dim != SPIRV::Dim::DIM_Rect) {
6043 "Gather operations are only supported for 2D, Cube, and Rect images.");
6050 if (!loadHandleBeforePosition(
6055 MachineIRBuilder MIRBuilder(
I);
6056 SPIRVTypeInst SampledImageType =
6061 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6069 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
6071 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
6073 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6078 .
addUse(ComponentOrCompareReg);
6080 uint32_t ImageOperands = 0;
6081 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
6082 if (Dim == SPIRV::Dim::DIM_Cube) {
6084 "Gather operations with offset are not supported for Cube images.");
6088 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6090 ImageOperands |= SPIRV::ImageOperand::Offset;
6094 if (ImageOperands != 0) {
6095 MIB.
addImm(ImageOperands);
6097 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6105bool SPIRVInstructionSelector::generateImageReadOrFetch(
6108 const ImageOperands *ImOps)
const {
6111 "ImageReg is not an image type.");
6113 bool IsSignedInteger =
6118 bool IsFetch = (SampledOp.getImm() == 1);
6120 auto AddOperands = [&](MachineInstrBuilder &MIB) {
6121 uint32_t ImageOperandsMask = 0;
6122 if (IsSignedInteger)
6123 ImageOperandsMask |= 0x1000;
6125 if (IsFetch && ImOps) {
6127 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
6128 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
6130 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
6132 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
6136 if (ImageOperandsMask != 0) {
6137 MIB.
addImm(ImageOperandsMask);
6138 if (IsFetch && ImOps) {
6141 if (ImOps->Offset &&
6142 (ImageOperandsMask &
6143 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
6144 MIB.
addUse(*ImOps->Offset);
6150 if (ResultSize == 4) {
6153 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
6160 BMI.constrainAllUses(
TII,
TRI, RBI);
6164 SPIRVTypeInst ReadType = widenTypeToVec4(ResType, Pos);
6168 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
6174 BMI.constrainAllUses(
TII,
TRI, RBI);
6176 if (ResultSize == 1) {
6185 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
6188bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
6189 SPIRVTypeInst ResType,
6190 MachineInstr &
I)
const {
6191 Register ResourcePtr =
I.getOperand(2).getReg();
6193 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
6202 MachineIRBuilder MIRBuilder(
I);
6207 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6213 if (
I.getNumExplicitOperands() > 3) {
6214 Register IndexReg =
I.getOperand(3).getReg();
6221bool SPIRVInstructionSelector::selectPushConstantGetPointer(
6222 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6227bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
6228 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6229 Register ObjReg =
I.getOperand(2).getReg();
6230 if (!BuildCOPY(ResVReg, ObjReg,
I))
6240 decorateUsesAsNonUniform(ResVReg);
6244void SPIRVInstructionSelector::decorateUsesAsNonUniform(
6247 while (WorkList.
size() > 0) {
6251 bool IsDecorated =
false;
6253 if (
Use.getOpcode() == SPIRV::OpDecorate &&
6254 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
6260 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
6262 if (ResultReg == CurrentReg)
6270 SPIRV::Decoration::NonUniformEXT, {});
6275bool SPIRVInstructionSelector::extractSubvector(
6277 MachineInstr &InsertionPoint)
const {
6279 [[maybe_unused]] uint64_t InputSize =
6282 assert(InputSize > 1 &&
"The input must be a vector.");
6283 assert(ResultSize > 1 &&
"The result must be a vector.");
6284 assert(ResultSize < InputSize &&
6285 "Cannot extract more element than there are in the input.");
6288 const TargetRegisterClass *ScalarRegClass = GR.
getRegClass(ScalarType);
6289 for (uint64_t
I = 0;
I < ResultSize;
I++) {
6292 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6301 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
6303 TII.get(SPIRV::OpCompositeConstruct))
6307 for (
Register ComponentReg : ComponentRegisters)
6308 MIB.
addUse(ComponentReg);
6313bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
6314 MachineInstr &
I)
const {
6321 Register ImageReg =
I.getOperand(1).getReg();
6329 Register CoordinateReg =
I.getOperand(2).getReg();
6330 Register DataReg =
I.getOperand(3).getReg();
6333 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
6341Register SPIRVInstructionSelector::buildPointerToResource(
6342 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
6343 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
6344 StringRef Name, MachineIRBuilder MIRBuilder)
const {
6346 if (ArraySize == 1) {
6347 SPIRVTypeInst PtrType =
6350 "SpirvResType did not have an explicit layout.");
6355 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
6356 SPIRVTypeInst VarPointerType =
6359 VarPointerType, Set,
Binding, Name, MIRBuilder);
6361 SPIRVTypeInst ResPointerType =
6374bool SPIRVInstructionSelector::selectFirstBitSet16(
6375 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
6376 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
6378 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
6382 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
6385bool SPIRVInstructionSelector::selectFirstBitSet32(
6387 unsigned BitSetOpcode)
const {
6388 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6391 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6398bool SPIRVInstructionSelector::selectFirstBitSet64(
6400 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
6413 if (ComponentCount > 2) {
6414 auto Func = [
this, SwapPrimarySide](
Register ResVReg, SPIRVTypeInst ResType,
6416 unsigned Opcode) ->
bool {
6417 return this->selectFirstBitSet64(ResVReg, ResType,
I, SrcReg, Opcode,
6421 return handle64BitOverflow(ResVReg, ResType,
I, SrcReg, BitSetOpcode, Func);
6425 MachineIRBuilder MIRBuilder(
I);
6427 BaseType, 2 * ComponentCount, MIRBuilder,
false);
6431 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
6437 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
6444 bool IsScalarRes = ResType->
getOpcode() != SPIRV::OpTypeVector;
6447 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
6448 SPIRV::OpVectorExtractDynamic))
6450 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
6451 SPIRV::OpVectorExtractDynamic))
6455 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6456 TII.get(SPIRV::OpVectorShuffle))
6464 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
6470 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6471 TII.get(SPIRV::OpVectorShuffle))
6479 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
6499 SelectOp = SPIRV::OpSelectSISCond;
6500 AddOp = SPIRV::OpIAddS;
6508 SelectOp = SPIRV::OpSelectVIVCond;
6509 AddOp = SPIRV::OpIAddV;
6515 Register RegSecondaryOffset = Reg0;
6519 if (SwapPrimarySide) {
6520 PrimaryReg = LowReg;
6521 SecondaryReg = HighReg;
6522 RegPrimaryOffset = Reg0;
6523 RegSecondaryOffset = Reg32;
6528 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
6529 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
6534 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
6535 SPIRV::OpINotEqual))
6542 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
6543 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
6548 if (SwapPrimarySide) {
6550 if (!selectOpWithSrcs(RegAdd, ResType,
I,
6551 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
6562 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
6563 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
6568 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
6569 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
6572 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
6576bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
6577 SPIRVTypeInst ResType,
6579 bool IsSigned)
const {
6581 Register OpReg =
I.getOperand(2).getReg();
6584 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
6585 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
6589 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6591 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6593 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6596 return diagnoseUnsupported(
6598 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
6602bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
6603 SPIRVTypeInst ResType,
6604 MachineInstr &
I)
const {
6606 Register OpReg =
I.getOperand(2).getReg();
6611 unsigned ExtendOpcode = SPIRV::OpUConvert;
6612 unsigned BitSetOpcode = GL::FindILsb;
6616 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6618 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6620 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6623 return diagnoseUnsupported(
I,
6624 "spv_firstbitlow only supports 16,32,64 bits.");
6628bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
6629 SPIRVTypeInst ResType,
6630 MachineInstr &
I)
const {
6634 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
6637 .
addUse(
I.getOperand(2).getReg())
6640 unsigned Alignment =
I.getOperand(3).getImm();
6654 while (!Worklist.
empty()) {
6656 switch (
T->getOpcode()) {
6657 case SPIRV::OpTypeInt:
6658 case SPIRV::OpTypeFloat:
6659 case SPIRV::OpTypePointer:
6661 case SPIRV::OpTypeVector:
6662 case SPIRV::OpTypeMatrix:
6663 case SPIRV::OpTypeArray: {
6664 Register OperandReg =
T->getOperand(1).getReg();
6668 case SPIRV::OpTypeStruct:
6669 for (
unsigned Idx = 1,
E =
T->getNumOperands(); Idx <
E; ++Idx) {
6670 Register OperandReg =
T->getOperand(Idx).getReg();
6682bool SPIRVInstructionSelector::selectAbort(MachineInstr &
I)
const {
6683 assert(
I.getNumExplicitOperands() == 2);
6685 Register MsgReg =
I.getOperand(1).getReg();
6687 assert(MsgType &&
"Message argument of llvm.spv.abort has no SPIR-V type");
6690 return diagnoseUnsupported(
6692 "llvm.spv.abort message type must be a concrete SPIR-V type (numerical "
6693 "scalar, pointer, vector, matrix, or aggregate of such types)");
6696 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
6703bool SPIRVInstructionSelector::selectTrap(MachineInstr &
I)
const {
6712 uint32_t MsgVal = ~0
u;
6713 if (
I.getOpcode() == TargetOpcode::G_UBSANTRAP)
6714 MsgVal =
static_cast<uint32_t
>(
I.getOperand(0).
getImm());
6717 Register MsgReg = buildI32ConstantInEntryBlock(MsgVal,
I, MsgType);
6720 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
6727bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
6728 SPIRVTypeInst ResType,
6729 MachineInstr &
I)
const {
6733 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
6736 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
6739 unsigned Alignment =
I.getOperand(2).getImm();
6746bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
6751 const MachineInstr *PrevI =
I.getPrevNode();
6753 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
6757 .
addMBB(
I.getOperand(0).getMBB())
6762 .
addMBB(
I.getOperand(0).getMBB())
6767bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
6778 const MachineInstr *NextI =
I.getNextNode();
6780 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
6786 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
6788 .
addUse(
I.getOperand(0).getReg())
6789 .
addMBB(
I.getOperand(1).getMBB())
6795bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
6796 MachineInstr &
I)
const {
6798 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
6800 const unsigned NumOps =
I.getNumOperands();
6801 for (
unsigned i = 1; i <
NumOps; i += 2) {
6802 MIB.
addUse(
I.getOperand(i + 0).getReg());
6803 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
6809bool SPIRVInstructionSelector::selectGlobalValue(
6810 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
6812 MachineIRBuilder MIRBuilder(
I);
6813 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
6816 std::string GlobalIdent;
6818 unsigned &
ID = UnnamedGlobalIDs[GV];
6820 ID = UnnamedGlobalIDs.
size();
6821 GlobalIdent =
"__unnamed_" + Twine(
ID).str();
6847 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
6854 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
6859 MachineInstrBuilder MIB1 =
6860 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
6863 MachineInstrBuilder MIB2 =
6865 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
6869 GR.
add(ConstVal, MIB2);
6877 MachineInstrBuilder MIB3 =
6878 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
6881 GR.
add(ConstVal, MIB3);
6887 assert(NewReg != ResVReg);
6888 return BuildCOPY(ResVReg, NewReg,
I);
6898 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
6901 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
6907 SPIRVTypeInst ResType =
6911 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
6916 if (
GlobalVar->isExternallyInitialized() &&
6917 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
6918 constexpr unsigned ReadWriteINTEL = 3u;
6921 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
6927bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
6928 SPIRVTypeInst ResType,
6929 MachineInstr &
I)
const {
6931 return selectExtInst(ResVReg, ResType,
I, CL::log10);
6939 MachineIRBuilder MIRBuilder(
I);
6944 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6947 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6949 .
add(
I.getOperand(1))
6954 ResType->
getOpcode() == SPIRV::OpTypeFloat);
6964 APFloat::rmNearestTiesToEven, &LosesInfo);
6968 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
6969 ? SPIRV::OpVectorTimesScalar
6980bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
6981 SPIRVTypeInst ResType,
6982 MachineInstr &
I)
const {
6985 return selectExtInst(ResVReg, ResType,
I, CL::pown);
6991 Register ExpReg =
I.getOperand(2).getReg();
6993 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
6994 SPIRV::OpConvertSToF))
6996 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
7003bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
7004 SPIRVTypeInst ResType,
7005 MachineInstr &
I)
const {
7021 MachineIRBuilder MIRBuilder(
I);
7024 ResType, MIRBuilder, SPIRV::StorageClass::Function);
7037 MachineBasicBlock &EntryBB =
I.getMF()->
front();
7039 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
7042 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7048 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7051 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
7054 .
add(
I.getOperand(
I.getNumExplicitDefs()))
7058 Register IntegralPartReg =
I.getOperand(1).getReg();
7059 if (IntegralPartReg.
isValid()) {
7061 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7072 assert(
false &&
"GLSL::Modf is deprecated.");
7083bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
7084 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7085 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7086 MachineIRBuilder MIRBuilder(
I);
7087 const SPIRVTypeInst Vec3Ty =
7090 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
7102 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7106 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7112 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7119 assert(
I.getOperand(2).isReg());
7120 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
7124 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
7135bool SPIRVInstructionSelector::loadBuiltinInputID(
7136 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7137 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7138 MachineIRBuilder MIRBuilder(
I);
7140 ResType, MIRBuilder, SPIRV::StorageClass::Input);
7155 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7159 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7168SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
7169 MachineInstr &
I)
const {
7170 MachineIRBuilder MIRBuilder(
I);
7171 if (
Type->getOpcode() != SPIRV::OpTypeVector)
7181bool SPIRVInstructionSelector::loadHandleBeforePosition(
7182 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
7183 MachineInstr &Pos)
const {
7186 Intrinsic::spv_resource_handlefrombinding);
7194 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
7195 MachineIRBuilder MIRBuilder(HandleDef);
7196 SPIRVTypeInst VarType = ResType;
7197 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
7199 if (IsStructuredBuffer) {
7204 if (ResType->
getOpcode() == SPIRV::OpTypeImage && ArraySize == 0)
7206 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
7209 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
7210 ArraySize, IndexReg, Name, MIRBuilder);
7214 uint32_t LoadOpcode =
7215 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
7225bool SPIRVInstructionSelector::errorIfInstrOutsideShader(
7226 MachineInstr &
I)
const {
7228 return diagnoseUnsupported(
7229 I,
"this instruction is only supported in shaders.");
7234InstructionSelector *
7238 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.
MachineInstr * getDef(const MachineOperand &MO, const MachineRegisterInfo *MRI)
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)
LLVM_ABI bool isNullOrNullSplat(const MachineInstr &MI, const MachineRegisterInfo &MRI, bool AllowUndefs=false)
Return true if the value is a constant 0 integer or a splatted vector of a constant 0 integer (with n...
uint64_t getIConstVal(Register ConstReg, const MachineRegisterInfo *MRI)
SmallVector< MachineInstr *, 4 > createContinuedInstructions(MachineIRBuilder &MIRBuilder, unsigned Opcode, unsigned MinWC, unsigned ContinuedOpcode, ArrayRef< Register > Args, Register ReturnRegister, Register TypeID)
SPIRV::MemorySemantics::MemorySemantics getMemSemanticsForStorageClass(SPIRV::StorageClass::StorageClass SC)
constexpr unsigned storageClassToAddressSpace(SPIRV::StorageClass::StorageClass SC)
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...