37#include "llvm/IR/IntrinsicsSPIRV.h"
43#define DEBUG_TYPE "spirv-isel"
50 std::vector<std::pair<SPIRV::InstructionSet::InstructionSet, uint32_t>>;
55 std::optional<Register> Bias;
56 std::optional<Register>
Offset;
57 std::optional<Register> MinLod;
58 std::optional<Register> GradX;
59 std::optional<Register> GradY;
60 std::optional<Register> Lod;
61 std::optional<Register> Compare;
68 bool IsScalar =
false;
71llvm::SPIRV::SelectionControl::SelectionControl
72getSelectionOperandForImm(
int Imm) {
74 return SPIRV::SelectionControl::Flatten;
76 return SPIRV::SelectionControl::DontFlatten;
78 return SPIRV::SelectionControl::None;
82#define GET_GLOBALISEL_PREDICATE_BITSET
83#include "SPIRVGenGlobalISel.inc"
84#undef GET_GLOBALISEL_PREDICATE_BITSET
111#define GET_GLOBALISEL_PREDICATES_DECL
112#include "SPIRVGenGlobalISel.inc"
113#undef GET_GLOBALISEL_PREDICATES_DECL
115#define GET_GLOBALISEL_TEMPORARIES_DECL
116#include "SPIRVGenGlobalISel.inc"
117#undef GET_GLOBALISEL_TEMPORARIES_DECL
141 unsigned BitSetOpcode)
const;
145 unsigned BitSetOpcode)
const;
149 unsigned BitSetOpcode,
bool SwapPrimarySide)
const;
156 unsigned Opcode)
const;
159 unsigned Opcode)
const;
181 unsigned NewOpcode,
unsigned NegateOpcode = 0)
const;
190 Register castPtrToPtrToInt(
Register Ptr, SPIRV::StorageClass::StorageClass SC,
194 bool selectAtomicPtrValue(
210 unsigned OpType)
const;
278 unsigned Opcode)
const;
282 unsigned Opcode)
const;
286 unsigned Opcode)
const;
290 unsigned Opcode)
const;
292 template <
bool Signed>
295 template <
bool Signed>
302 template <
typename PickOpcodeFn>
305 PickOpcodeFn &&PickOpcode)
const;
322 template <
typename PickOpcodeFn>
325 PickOpcodeFn &&PickOpcode)
const;
343 bool IsSigned)
const;
345 bool IsSigned,
unsigned Opcode)
const;
347 bool IsSigned)
const;
353 bool IsSigned)
const;
394 GL::GLSLExtInst GLInst,
bool setMIFlags =
true,
395 bool useMISrc =
true,
397 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
398 CL::OpenCLExtInst CLInst,
bool setMIFlags =
true,
399 bool useMISrc =
true,
401 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
402 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
403 bool setMIFlags =
true,
bool useMISrc =
true,
405 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
406 const ExtInstList &ExtInsts,
bool setMIFlags =
true,
407 bool useMISrc =
true,
410 bool selectLog10(
Register ResVReg, SPIRVTypeInst ResType,
411 MachineInstr &
I)
const;
413 bool selectFpowi(
Register ResVReg, SPIRVTypeInst ResType,
414 MachineInstr &
I)
const;
416 bool selectSaturate(
Register ResVReg, SPIRVTypeInst ResType,
417 MachineInstr &
I)
const;
419 bool selectWaveOpInst(
Register ResVReg, SPIRVTypeInst ResType,
420 MachineInstr &
I,
unsigned Opcode)
const;
422 bool selectBarrierInst(MachineInstr &
I,
unsigned Scope,
unsigned MemSem,
423 bool WithGroupSync)
const;
425 bool selectWaveActiveCountBits(
Register ResVReg, SPIRVTypeInst ResType,
426 MachineInstr &
I)
const;
428 bool selectWaveActiveAllEqual(
Register ResVReg, SPIRVTypeInst ResType,
429 MachineInstr &
I)
const;
433 bool selectHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
434 MachineInstr &
I)
const;
436 bool selectCounterHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
437 MachineInstr &
I)
const;
439 bool selectReadImageIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
440 MachineInstr &
I)
const;
441 bool selectGetDimensionsIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
442 MachineInstr &
I)
const;
443 bool selectGetDimensionsLevelsIntrinsic(
Register &ResVReg,
444 SPIRVTypeInst ResType,
445 MachineInstr &
I)
const;
446 bool selectGetDimensionsMSIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
447 MachineInstr &
I)
const;
450 std::optional<Register> LodReg = std::nullopt)
const;
451 bool selectSampleBasicIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
452 MachineInstr &
I)
const;
453 bool selectCalculateLodIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
454 MachineInstr &
I)
const;
455 bool selectSampleBiasIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
456 MachineInstr &
I)
const;
457 bool selectSampleGradIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
458 MachineInstr &
I)
const;
459 bool selectSampleLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
460 MachineInstr &
I)
const;
461 bool selectLoadLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
462 MachineInstr &
I)
const;
463 bool selectSampleCmpIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
464 MachineInstr &
I)
const;
465 bool selectSampleCmpLevelZeroIntrinsic(
Register &ResVReg,
466 SPIRVTypeInst ResType,
467 MachineInstr &
I)
const;
468 bool selectGatherIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
469 MachineInstr &
I)
const;
470 bool selectImageWriteIntrinsic(MachineInstr &
I)
const;
471 bool selectResourceGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
472 MachineInstr &
I)
const;
473 bool selectPushConstantGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
474 MachineInstr &
I)
const;
475 bool selectResourceNonUniformIndex(
Register &ResVReg, SPIRVTypeInst ResType,
476 MachineInstr &
I)
const;
477 bool selectModf(
Register ResVReg, SPIRVTypeInst ResType,
478 MachineInstr &
I)
const;
479 bool selectUpdateCounter(
Register &ResVReg, SPIRVTypeInst ResType,
480 MachineInstr &
I)
const;
481 bool selectFrexp(
Register ResVReg, SPIRVTypeInst ResType,
482 MachineInstr &
I)
const;
484 bool selectCopySign(
Register ResVReg, SPIRVTypeInst ResType,
485 MachineInstr &
I)
const;
487 bool selectLdexp(
Register ResVReg, SPIRVTypeInst ResType,
488 MachineInstr &
I)
const;
489 bool selectSincos(
Register ResVReg, SPIRVTypeInst ResType,
490 MachineInstr &
I)
const;
491 bool selectExp10(
Register ResVReg, SPIRVTypeInst ResType,
492 MachineInstr &
I)
const;
493 bool selectDerivativeInst(
Register ResVReg, SPIRVTypeInst ResType,
494 MachineInstr &
I,
const unsigned DPdOpCode)
const;
496 Register buildI32Constant(uint32_t Val, MachineInstr &
I,
497 SPIRVTypeInst ResType =
nullptr)
const;
498 Register buildI32ConstantInEntryBlock(uint32_t Val, MachineInstr &
I,
499 SPIRVTypeInst ResType =
nullptr)
const;
501 Register buildZerosVal(SPIRVTypeInst ResType, MachineInstr &
I)
const;
502 bool isScalarOrVectorIntConstantZero(
Register Reg)
const;
503 Register buildZerosValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
505 MachineInstr &
I)
const;
506 Register buildOnesValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
508 bool wrapIntoSpecConstantOp(MachineInstr &
I,
511 Register getUcharPtrTypeReg(MachineInstr &
I,
512 SPIRV::StorageClass::StorageClass SC)
const;
513 MachineInstrBuilder buildSpecConstantOp(MachineInstr &
I,
Register Dest,
515 uint32_t Opcode)
const;
516 MachineInstrBuilder buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
517 SPIRVTypeInst SrcPtrTy)
const;
518 Register buildPointerToResource(SPIRVTypeInst ResType,
519 SPIRV::StorageClass::StorageClass SC,
520 uint32_t Set, uint32_t
Binding,
521 uint32_t ArraySize,
Register IndexReg,
523 MachineIRBuilder MIRBuilder)
const;
524 SPIRVTypeInst widenTypeToVec4(SPIRVTypeInst
Type, MachineInstr &
I)
const;
525 bool extractSubvector(
Register &ResVReg, SPIRVTypeInst ResType,
526 Register &ReadReg, MachineInstr &InsertionPoint)
const;
527 bool generateImageReadOrFetch(
Register &ResVReg, SPIRVTypeInst ResType,
530 const ImageOperands *ImOps =
nullptr)
const;
531 bool generateSampleImage(
Register ResVReg, SPIRVTypeInst ResType,
533 Register CoordinateReg,
const ImageOperands &ImOps,
536 bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
537 Register ResVReg, SPIRVTypeInst ResType,
538 MachineInstr &
I)
const;
539 bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
540 Register ResVReg, SPIRVTypeInst ResType,
541 MachineInstr &
I)
const;
542 bool loadHandleBeforePosition(
Register &HandleReg, SPIRVTypeInst ResType,
543 GIntrinsic &HandleDef, MachineInstr &Pos)
const;
544 void decorateUsesAsNonUniform(
Register &NonUniformReg)
const;
545 bool errorIfInstrOutsideShader(MachineInstr &
I)
const;
547 std::optional<SplitParts> splitEvenOddLanes(
Register PopCountReg,
548 unsigned ComponentCount,
550 SPIRVTypeInst I32Type)
const;
553 handle64BitOverflow(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
554 Register SrcReg,
unsigned int Opcode,
555 std::function<
bool(
Register, SPIRVTypeInst,
556 MachineInstr &,
Register,
unsigned)>
560bool sampledTypeIsSignedInteger(
const llvm::Type *HandleType) {
562 if (
TET->getTargetExtName() ==
"spirv.Image") {
565 assert(
TET->getTargetExtName() ==
"spirv.SignedImage");
566 return TET->getTypeParameter(0)->isIntegerTy();
570#define GET_GLOBALISEL_IMPL
571#include "SPIRVGenGlobalISel.inc"
572#undef GET_GLOBALISEL_IMPL
578 TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
581#include
"SPIRVGenGlobalISel.inc"
584#include
"SPIRVGenGlobalISel.inc"
596 InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
601 if (HasVRegsReset == &MF)
616 for (
const auto &
MBB : MF) {
617 for (
const auto &
MI :
MBB) {
620 if (
MI.getOpcode() != SPIRV::ASSIGN_TYPE)
624 LLT DstType = MRI.
getType(DstReg);
626 LLT SrcType = MRI.
getType(SrcReg);
627 if (DstType != SrcType)
632 if (DstRC != SrcRC && SrcRC)
644 while (!Stack.empty()) {
649 switch (
MI->getOpcode()) {
650 case TargetOpcode::G_INTRINSIC:
651 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
652 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS: {
655 if (IntrID != Intrinsic::spv_const_composite &&
656 IntrID != Intrinsic::spv_undef && IntrID != Intrinsic::spv_poison)
660 case TargetOpcode::G_BUILD_VECTOR:
661 case TargetOpcode::G_SPLAT_VECTOR:
663 i < OpDef->getNumOperands(); i++) {
668 Stack.push_back(OpNestedDef);
671 case TargetOpcode::G_CONSTANT:
672 case TargetOpcode::G_FCONSTANT:
673 case TargetOpcode::G_IMPLICIT_DEF:
674 case SPIRV::OpConstantTrue:
675 case SPIRV::OpConstantFalse:
676 case SPIRV::OpConstantI:
677 case SPIRV::OpConstantF:
678 case SPIRV::OpConstantComposite:
679 case SPIRV::OpConstantCompositeContinuedINTEL:
680 case SPIRV::OpConstantSampler:
681 case SPIRV::OpConstantNull:
683 case SPIRV::OpPoisonKHR:
684 case SPIRV::OpConstantFunctionPointerINTEL:
711 case Intrinsic::spv_all:
712 case Intrinsic::spv_alloca:
713 case Intrinsic::spv_any:
714 case Intrinsic::spv_bitcast:
715 case Intrinsic::spv_const_composite:
716 case Intrinsic::spv_degrees:
717 case Intrinsic::spv_distance:
718 case Intrinsic::spv_extractelt:
719 case Intrinsic::spv_extractv:
720 case Intrinsic::spv_faceforward:
721 case Intrinsic::spv_fdot:
722 case Intrinsic::spv_firstbitlow:
723 case Intrinsic::spv_firstbitshigh:
724 case Intrinsic::spv_firstbituhigh:
725 case Intrinsic::spv_frac:
726 case Intrinsic::spv_gep:
727 case Intrinsic::spv_global_offset:
728 case Intrinsic::spv_global_size:
729 case Intrinsic::spv_group_id:
730 case Intrinsic::spv_insertelt:
731 case Intrinsic::spv_insertv:
732 case Intrinsic::spv_isinf:
733 case Intrinsic::spv_isnan:
734 case Intrinsic::spv_isfinite:
735 case Intrinsic::spv_isnormal:
736 case Intrinsic::spv_lerp:
737 case Intrinsic::spv_length:
738 case Intrinsic::spv_normalize:
739 case Intrinsic::spv_num_subgroups:
740 case Intrinsic::spv_num_workgroups:
741 case Intrinsic::spv_ptrcast:
742 case Intrinsic::spv_radians:
743 case Intrinsic::spv_reflect:
744 case Intrinsic::spv_refract:
745 case Intrinsic::spv_resource_getbasepointer:
746 case Intrinsic::spv_resource_getpointer:
747 case Intrinsic::spv_resource_handlefrombinding:
748 case Intrinsic::spv_resource_handlefromimplicitbinding:
749 case Intrinsic::spv_resource_nonuniformindex:
750 case Intrinsic::spv_resource_sample:
751 case Intrinsic::spv_rsqrt:
752 case Intrinsic::spv_saturate:
753 case Intrinsic::spv_sdot:
754 case Intrinsic::spv_sign:
755 case Intrinsic::spv_smoothstep:
756 case Intrinsic::spv_subgroup_id:
757 case Intrinsic::spv_subgroup_local_invocation_id:
758 case Intrinsic::spv_subgroup_max_size:
759 case Intrinsic::spv_subgroup_size:
760 case Intrinsic::spv_thread_id:
761 case Intrinsic::spv_thread_id_in_group:
762 case Intrinsic::spv_udot:
763 case Intrinsic::spv_undef:
764 case Intrinsic::spv_value_md:
765 case Intrinsic::spv_workgroup_size:
777 case SPIRV::OpTypeVoid:
778 case SPIRV::OpTypeBool:
779 case SPIRV::OpTypeInt:
780 case SPIRV::OpTypeFloat:
781 case SPIRV::OpTypeVector:
782 case SPIRV::OpTypeVectorIdEXT:
783 case SPIRV::OpTypeMatrix:
784 case SPIRV::OpTypeImage:
785 case SPIRV::OpTypeSampler:
786 case SPIRV::OpTypeSampledImage:
787 case SPIRV::OpTypeArray:
788 case SPIRV::OpTypeRuntimeArray:
789 case SPIRV::OpTypeStruct:
790 case SPIRV::OpTypeOpaque:
791 case SPIRV::OpTypePointer:
792 case SPIRV::OpTypeFunction:
793 case SPIRV::OpTypeEvent:
794 case SPIRV::OpTypeDeviceEvent:
795 case SPIRV::OpTypeReserveId:
796 case SPIRV::OpTypeQueue:
797 case SPIRV::OpTypePipe:
798 case SPIRV::OpTypeForwardPointer:
799 case SPIRV::OpTypePipeStorage:
800 case SPIRV::OpTypeNamedBarrier:
801 case SPIRV::OpTypeAccelerationStructureNV:
802 case SPIRV::OpTypeCooperativeMatrixNV:
803 case SPIRV::OpTypeCooperativeMatrixKHR:
813 if (
MI.getNumDefs() == 0)
816 for (
const auto &MO :
MI.all_defs()) {
818 if (
Reg.isPhysical()) {
823 if (
UseMI.getOpcode() != SPIRV::OpName) {
830 if (
MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE ||
MI.isFakeUse() ||
831 MI.isLifetimeMarker()) {
834 <<
"Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
845 if (
MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
846 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
849 LLVM_DEBUG(
dbgs() <<
"Dead: Intrinsic with no real side effects.\n");
854 if (
MI.mayStore() ||
MI.isCall() ||
855 (
MI.mayLoad() &&
MI.hasOrderedMemoryRef()) ||
MI.isPosition() ||
856 MI.isDebugInstr() ||
MI.isTerminator() ||
MI.isJumpTableDebugInfo()) {
857 LLVM_DEBUG(
dbgs() <<
"Not dead: instruction has side effects.\n");
868 LLVM_DEBUG(
dbgs() <<
"Dead: known opcode with no side effects\n");
875void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &
MI)
const {
877 for (
const auto &MO :
MI.all_defs()) {
881 SmallVector<MachineInstr *, 4> UselessOpNames;
884 "There is still a use of the dead function.");
887 for (MachineInstr *OpNameMI : UselessOpNames) {
889 OpNameMI->eraseFromParent();
894void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &
MI)
const {
897 removeOpNamesForDeadMI(
MI);
898 MI.eraseFromParent();
901bool SPIRVInstructionSelector::select(MachineInstr &
I) {
902 resetVRegsType(*
I.getParent()->getParent());
904 assert(
I.getParent() &&
"Instruction should be in a basic block!");
905 assert(
I.getParent()->getParent() &&
"Instruction should be in a function!");
910 removeDeadInstruction(
I);
917 if (Opcode == SPIRV::ASSIGN_TYPE) {
918 Register DstReg =
I.getOperand(0).getReg();
919 Register SrcReg =
I.getOperand(1).getReg();
922 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
923 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
924 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
925 Register SelectDstReg =
Def->getOperand(0).getReg();
926 bool SuccessToSelectSelect [[maybe_unused]] = selectSelect(
928 assert(SuccessToSelectSelect);
930 Def->eraseFromParent();
937 bool Res = selectImpl(
I, *CoverageInfo);
939 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
940 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
944 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
956 }
else if (
I.getNumDefs() == 1) {
968 removeDeadInstruction(
I);
973 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
974 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
980 bool HasDefs =
I.getNumDefs() > 0;
983 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
984 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
985 if (spvSelect(ResVReg, ResType,
I)) {
987 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
998 case TargetOpcode::G_CONSTANT:
999 case TargetOpcode::G_FCONSTANT:
1006 MachineInstr &
I)
const {
1009 if (DstRC != SrcRC && SrcRC)
1011 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
1018bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
1019 SPIRVTypeInst ResType,
1020 MachineInstr &
I)
const {
1021 const unsigned Opcode =
I.getOpcode();
1023 return selectImpl(
I, *CoverageInfo);
1025 case TargetOpcode::G_CONSTANT:
1026 case TargetOpcode::G_FCONSTANT:
1027 return selectConst(ResVReg, ResType,
I);
1028 case TargetOpcode::G_GLOBAL_VALUE:
1029 return selectGlobalValue(ResVReg,
I);
1030 case TargetOpcode::G_IMPLICIT_DEF:
1031 return selectOpUndef(ResVReg, ResType,
I);
1032 case TargetOpcode::G_FREEZE:
1033 return selectFreeze(ResVReg, ResType,
I);
1035 case TargetOpcode::G_INTRINSIC:
1036 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
1037 case TargetOpcode::G_INTRINSIC_CONVERGENT:
1038 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
1039 return selectIntrinsic(ResVReg, ResType,
I);
1040 case TargetOpcode::G_BITREVERSE:
1041 return selectBitreverse(ResVReg, ResType,
I);
1043 case TargetOpcode::G_BUILD_VECTOR:
1044 return selectBuildVector(ResVReg, ResType,
I);
1045 case TargetOpcode::G_SPLAT_VECTOR:
1046 return selectSplatVector(ResVReg, ResType,
I);
1047 case TargetOpcode::G_CONCAT_VECTORS:
1048 return selectConcatVectors(ResVReg, ResType,
I);
1050 case TargetOpcode::G_SHUFFLE_VECTOR: {
1051 MachineBasicBlock &BB = *
I.getParent();
1052 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
1055 .
addUse(
I.getOperand(1).getReg())
1056 .
addUse(
I.getOperand(2).getReg());
1057 for (
auto V :
I.getOperand(3).getShuffleMask())
1062 case TargetOpcode::G_MEMMOVE:
1063 case TargetOpcode::G_MEMCPY:
1064 case TargetOpcode::G_MEMCPY_INLINE:
1065 case TargetOpcode::G_MEMSET:
1066 case TargetOpcode::G_MEMSET_INLINE:
1067 return selectMemOperation(ResVReg,
I);
1069 case TargetOpcode::G_ICMP:
1070 return selectICmp(ResVReg, ResType,
I);
1071 case TargetOpcode::G_FCMP:
1072 return selectFCmp(ResVReg, ResType,
I);
1074 case TargetOpcode::G_FRAME_INDEX:
1075 return selectFrameIndex(ResVReg, ResType,
I);
1077 case TargetOpcode::G_LOAD:
1078 return selectLoad(ResVReg, ResType,
I);
1079 case TargetOpcode::G_STORE:
1080 return selectStore(
I);
1082 case TargetOpcode::G_BR:
1083 return selectBranch(
I);
1084 case TargetOpcode::G_BRCOND:
1085 return selectBranchCond(
I);
1087 case TargetOpcode::G_PHI:
1088 return selectPhi(ResVReg,
I);
1090 case TargetOpcode::G_FPTOSI:
1091 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1092 case TargetOpcode::G_FPTOUI:
1093 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1095 case TargetOpcode::G_FPTOSI_SAT:
1096 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1097 case TargetOpcode::G_FPTOUI_SAT:
1098 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1100 case TargetOpcode::G_SITOFP:
1101 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
1102 case TargetOpcode::G_UITOFP:
1103 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
1105 case TargetOpcode::G_CTPOP:
1106 return selectPopCount(ResVReg, ResType,
I, SPIRV::OpBitCount);
1107 case TargetOpcode::G_SMIN:
1108 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
1109 case TargetOpcode::G_UMIN:
1110 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
1112 case TargetOpcode::G_SMAX:
1113 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
1114 case TargetOpcode::G_UMAX:
1115 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
1117 case TargetOpcode::G_SCMP:
1118 return selectSUCmp(ResVReg, ResType,
I,
true);
1119 case TargetOpcode::G_UCMP:
1120 return selectSUCmp(ResVReg, ResType,
I,
false);
1121 case TargetOpcode::G_LROUND:
1122 case TargetOpcode::G_LLROUND: {
1125 MRI->
setRegClass(regForLround, &SPIRV::iIDRegClass);
1127 regForLround, *(
I.getParent()->getParent()));
1129 CL::round, GL::Round,
false);
1131 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1138 case TargetOpcode::G_STRICT_FMA:
1139 case TargetOpcode::G_FMA: {
1142 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1145 .
addUse(
I.getOperand(1).getReg())
1146 .
addUse(
I.getOperand(2).getReg())
1147 .
addUse(
I.getOperand(3).getReg())
1152 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1155 case TargetOpcode::G_FLDEXP:
1156 case TargetOpcode::G_STRICT_FLDEXP:
1157 return selectLdexp(ResVReg, ResType,
I);
1159 case TargetOpcode::G_FPOW:
1160 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1161 case TargetOpcode::G_FPOWI:
1162 return selectFpowi(ResVReg, ResType,
I);
1164 case TargetOpcode::G_FEXP:
1165 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1166 case TargetOpcode::G_FEXP2:
1167 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1168 case TargetOpcode::G_FEXP10:
1169 return selectExp10(ResVReg, ResType,
I);
1171 case TargetOpcode::G_FMODF:
1172 return selectModf(ResVReg, ResType,
I);
1173 case TargetOpcode::G_FSINCOS:
1174 return selectSincos(ResVReg, ResType,
I);
1176 case TargetOpcode::G_FLOG:
1177 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1178 case TargetOpcode::G_FLOG2:
1179 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1180 case TargetOpcode::G_FLOG10:
1181 return selectLog10(ResVReg, ResType,
I);
1183 case TargetOpcode::G_FABS:
1184 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1185 case TargetOpcode::G_ABS:
1186 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1188 case TargetOpcode::G_FMINNUM:
1189 case TargetOpcode::G_FMINIMUM:
1190 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1191 case TargetOpcode::G_FMAXNUM:
1192 case TargetOpcode::G_FMAXIMUM:
1193 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1195 case TargetOpcode::G_FCOPYSIGN:
1196 return selectCopySign(ResVReg, ResType,
I);
1198 case TargetOpcode::G_FCEIL:
1199 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1200 case TargetOpcode::G_FFLOOR:
1201 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1203 case TargetOpcode::G_FCOS:
1204 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1205 case TargetOpcode::G_FSIN:
1206 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1207 case TargetOpcode::G_FTAN:
1208 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1209 case TargetOpcode::G_FACOS:
1210 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1211 case TargetOpcode::G_FASIN:
1212 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1213 case TargetOpcode::G_FATAN:
1214 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1215 case TargetOpcode::G_FATAN2:
1216 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1217 case TargetOpcode::G_FCOSH:
1218 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1219 case TargetOpcode::G_FSINH:
1220 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1221 case TargetOpcode::G_FTANH:
1222 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1224 case TargetOpcode::G_STRICT_FSQRT:
1225 case TargetOpcode::G_FSQRT:
1226 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1228 case TargetOpcode::G_CTTZ:
1229 case TargetOpcode::G_CTTZ_ZERO_POISON:
1230 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1231 case TargetOpcode::G_CTLZ:
1232 case TargetOpcode::G_CTLZ_ZERO_POISON:
1233 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1235 case TargetOpcode::G_INTRINSIC_ROUND:
1236 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1237 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1238 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1239 case TargetOpcode::G_INTRINSIC_TRUNC:
1240 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1241 case TargetOpcode::G_FRINT:
1242 case TargetOpcode::G_FNEARBYINT:
1243 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1245 case TargetOpcode::G_SMULH:
1246 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1247 case TargetOpcode::G_UMULH:
1248 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1250 case TargetOpcode::G_SADDSAT:
1251 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1252 case TargetOpcode::G_UADDSAT:
1253 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1254 case TargetOpcode::G_SSUBSAT:
1255 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1256 case TargetOpcode::G_USUBSAT:
1257 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1259 case TargetOpcode::G_FFREXP:
1260 return selectFrexp(ResVReg, ResType,
I);
1262 case TargetOpcode::G_UADDO:
1263 return selectOverflowArith(ResVReg, ResType,
I,
1265 : SPIRV::OpIAddCarryS);
1266 case TargetOpcode::G_USUBO:
1267 return selectOverflowArith(ResVReg, ResType,
I,
1269 : SPIRV::OpISubBorrowS);
1270 case TargetOpcode::G_UMULO:
1271 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1272 case TargetOpcode::G_SMULO:
1273 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1275 case TargetOpcode::G_SEXT:
1276 return selectExt(ResVReg, ResType,
I,
true);
1277 case TargetOpcode::G_ANYEXT:
1278 case TargetOpcode::G_ZEXT:
1279 return selectExt(ResVReg, ResType,
I,
false);
1280 case TargetOpcode::G_TRUNC:
1281 return selectTrunc(ResVReg, ResType,
I);
1282 case TargetOpcode::G_FPTRUNC:
1283 case TargetOpcode::G_FPEXT:
1284 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1286 case TargetOpcode::G_PTRTOINT:
1287 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1288 case TargetOpcode::G_INTTOPTR:
1289 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1290 case TargetOpcode::G_BITCAST:
1291 return selectBitcast(ResVReg, ResType,
I);
1292 case TargetOpcode::G_ADDRSPACE_CAST:
1293 return selectAddrSpaceCast(ResVReg, ResType,
I);
1294 case TargetOpcode::G_PTRMASK:
1295 return selectPtrMask(ResVReg, ResType,
I);
1296 case TargetOpcode::G_PTR_ADD: {
1298 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1302 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1303 (*II).getOpcode() == TargetOpcode::COPY ||
1304 (*II).getOpcode() == SPIRV::OpVariable ||
1305 (*II).getOpcode() == SPIRV::OpUntypedVariableKHR) &&
1306 getImm(
I.getOperand(2), MRI));
1308 bool IsGVInit =
false;
1312 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1313 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1314 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1315 (*UseIt).getOpcode() == SPIRV::OpVariable ||
1316 (*UseIt).getOpcode() == SPIRV::OpUntypedVariableKHR) {
1328 const bool UseUntypedPointers =
1329 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1330 if (UseUntypedPointers) {
1331 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1334 .
addImm(
static_cast<uint32_t
>(
1335 SPIRV::Opcode::UntypedInBoundsPtrAccessChainKHR))
1338 .
addUse(
I.getOperand(2).getReg())
1345 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1357 return diagnoseUnsupported(
1358 I,
"incompatible result and operand types in a bitcast");
1360 MachineInstrBuilder MIB =
1361 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1368 : SPIRV::OpInBoundsPtrAccessChain))
1372 .
addUse(
I.getOperand(2).getReg())
1375 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1379 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1381 .
addUse(
I.getOperand(2).getReg())
1390 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1393 .
addImm(
static_cast<uint32_t
>(
1394 SPIRV::Opcode::InBoundsPtrAccessChain))
1397 .
addUse(
I.getOperand(2).getReg());
1402 case TargetOpcode::G_ATOMICRMW_OR:
1403 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1404 case TargetOpcode::G_ATOMICRMW_ADD:
1405 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1406 case TargetOpcode::G_ATOMICRMW_AND:
1407 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1408 case TargetOpcode::G_ATOMICRMW_MAX:
1409 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1410 case TargetOpcode::G_ATOMICRMW_MIN:
1411 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1412 case TargetOpcode::G_ATOMICRMW_SUB:
1413 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1414 case TargetOpcode::G_ATOMICRMW_XOR:
1415 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1416 case TargetOpcode::G_ATOMICRMW_UMAX:
1417 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1418 case TargetOpcode::G_ATOMICRMW_UMIN:
1419 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1420 case TargetOpcode::G_ATOMICRMW_XCHG:
1421 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1423 case TargetOpcode::G_ATOMICRMW_FADD:
1424 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1425 case TargetOpcode::G_ATOMICRMW_FSUB:
1427 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1429 : SPIRV::OpFNegate);
1430 case TargetOpcode::G_ATOMICRMW_FMIN:
1431 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1432 case TargetOpcode::G_ATOMICRMW_FMAX:
1433 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1435 case TargetOpcode::G_FENCE:
1436 return selectFence(
I);
1438 case TargetOpcode::G_STACKSAVE:
1439 return selectStackSave(ResVReg, ResType,
I);
1440 case TargetOpcode::G_STACKRESTORE:
1441 return selectStackRestore(
I);
1443 case TargetOpcode::G_UNMERGE_VALUES:
1446 case TargetOpcode::G_TRAP:
1447 case TargetOpcode::G_UBSANTRAP:
1448 return selectTrap(
I);
1453 case TargetOpcode::DBG_LABEL:
1455 case TargetOpcode::G_DEBUGTRAP:
1456 return selectDebugTrap(ResVReg, ResType,
I);
1457 case TargetOpcode::G_PREFETCH:
1458 return selectPrefetch(
I);
1465bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1466 SPIRVTypeInst ResType,
1467 MachineInstr &
I)
const {
1468 unsigned Opcode = SPIRV::OpNop;
1475bool SPIRVInstructionSelector::selectPrefetch(MachineInstr &
I)
const {
1484 MachineIRBuilder MIRBuilder(
I);
1486 const SPIRVTypeInst PointerSizeType =
1494 Register AddrVal =
I.getOperand(0).getReg();
1497 return selectExtInst(ExtReg, GR.
getOpTypeVoid(MIRBuilder),
I, CL::prefetch,
1499 {AddrVal, ConstIntOne});
1504bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1505 SPIRVTypeInst ResType,
1507 GL::GLSLExtInst GLInst,
1508 bool setMIFlags,
bool useMISrc,
1511 SPIRV::InstructionSet::InstructionSet::GLSL_std_450))
1512 return diagnoseUnsupported(
1514 "this instruction is only supported with the GLSL extended instruction "
1516 return selectExtInst(ResVReg, ResType,
I,
1517 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1518 setMIFlags, useMISrc, SrcRegs);
1521bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1522 SPIRVTypeInst ResType,
1524 CL::OpenCLExtInst CLInst,
1525 bool setMIFlags,
bool useMISrc,
1527 return selectExtInst(ResVReg, ResType,
I,
1528 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1529 setMIFlags, useMISrc, SrcRegs);
1532bool SPIRVInstructionSelector::selectExtInst(
1533 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1534 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1536 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1537 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1538 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1542bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1543 SPIRVTypeInst ResType,
1546 bool setMIFlags,
bool useMISrc,
1549 for (
const auto &[InstructionSet, Opcode] : Insts) {
1553 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1556 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1561 const unsigned NumOps =
I.getNumOperands();
1564 I.getOperand(Index).getType() ==
1565 MachineOperand::MachineOperandType::MO_IntrinsicID)
1568 MIB.
add(
I.getOperand(Index));
1580bool SPIRVInstructionSelector::selectCopySign(
Register ResVReg,
1581 SPIRVTypeInst ResType,
1582 MachineInstr &
I)
const {
1584 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1589 Register MagnitudeReg =
I.getOperand(1).getReg();
1590 Register SignReg =
I.getOperand(2).getReg();
1598 unsigned AndOpcode, OrOpcode;
1599 if (ComponentCount > 1) {
1603 AndOpcode = SPIRV::OpBitwiseAndV;
1604 OrOpcode = SPIRV::OpBitwiseOrV;
1608 AndOpcode = SPIRV::OpBitwiseAndS;
1609 OrOpcode = SPIRV::OpBitwiseOrS;
1615 return selectOpWithSrcs(ResReg, IntType,
I, SrcRegs, Opcode);
1618 Register MagnitudeInt, SignInt, MagnitudeBits, SignBits, CombinedInt;
1619 if (!EmitBitOp(MagnitudeInt, {MagnitudeReg}, SPIRV::OpBitcast) ||
1620 !EmitBitOp(SignInt, {SignReg}, SPIRV::OpBitcast) ||
1621 !EmitBitOp(MagnitudeBits, {MagnitudeInt, NotSignMask}, AndOpcode) ||
1622 !EmitBitOp(SignBits, {SignInt, SignMask}, AndOpcode) ||
1623 !EmitBitOp(CombinedInt, {MagnitudeBits, SignBits}, OrOpcode))
1626 return selectOpWithSrcs(ResVReg, ResType,
I, {CombinedInt}, SPIRV::OpBitcast);
1629bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1630 SPIRVTypeInst ResType,
1631 MachineInstr &
I)
const {
1632 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1633 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1634 for (
const auto &Ex : ExtInsts) {
1635 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1636 uint32_t Opcode = Ex.second;
1640 MachineIRBuilder MIRBuilder(
I);
1643 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1650 const bool IsUntyped =
1651 PointerType->getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1653 BuildMI(*It->getParent(), It, It->getDebugLoc(),
1654 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
1655 : SPIRV::OpVariable))
1658 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
1664 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1667 .
addImm(
static_cast<uint32_t
>(Ex.first))
1669 .
add(
I.getOperand(2))
1673 Register ExpResReg =
I.getOperand(1).getReg();
1675 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1685bool SPIRVInstructionSelector::selectLdexp(
Register ResVReg,
1686 SPIRVTypeInst ResType,
1687 MachineInstr &
I)
const {
1688 Register XReg =
I.getOperand(1).getReg();
1689 Register ExpReg =
I.getOperand(2).getReg();
1695 if (ResType->
getOpcode() == SPIRV::OpTypeVector &&
1696 ExpType->
getOpcode() != SPIRV::OpTypeVector) {
1698 SPIRVTypeInst ExpVecType =
1702 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1703 TII.get(SPIRV::OpCompositeConstruct))
1706 for (
unsigned J = 0; J < NumElts; ++J)
1712 return selectExtInst(ResVReg, ResType,
I, CL::ldexp, GL::Ldexp,
1713 true,
false, {XReg, ExpReg});
1716bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1717 SPIRVTypeInst ResType,
1718 MachineInstr &
I)
const {
1719 Register CosResVReg =
I.getOperand(1).getReg();
1720 unsigned SrcIdx =
I.getNumExplicitDefs();
1725 MachineIRBuilder MIRBuilder(
I);
1727 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1734 const bool IsUntyped =
1735 PointerType->getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1737 BuildMI(*It->getParent(), It, It->getDebugLoc(),
1738 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
1739 : SPIRV::OpVariable))
1742 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
1746 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1749 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1751 .
add(
I.getOperand(SrcIdx))
1755 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1763 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1766 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1768 .
add(
I.getOperand(SrcIdx))
1770 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1773 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1775 .
add(
I.getOperand(SrcIdx))
1782bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1783 SPIRVTypeInst ResType,
1786 unsigned Opcode)
const {
1787 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1797std::optional<SplitParts> SPIRVInstructionSelector::splitEvenOddLanes(
1798 Register PopCountReg,
unsigned ComponentCount, MachineInstr &
I,
1799 SPIRVTypeInst I32Type)
const {
1802 if (ComponentCount == 1) {
1805 Parts.IsScalar =
true;
1806 Parts.Type = I32Type;
1814 if (!selectOpWithSrcs(Parts.High, I32Type,
I, {PopCountReg, IdxOne},
1815 SPIRV::OpVectorExtractDynamic))
1816 return std::nullopt;
1818 if (!selectOpWithSrcs(Parts.Low, I32Type,
I, {PopCountReg, IdxZero},
1819 SPIRV::OpVectorExtractDynamic))
1820 return std::nullopt;
1824 MachineIRBuilder MIRBuilder(
I);
1825 Parts.IsScalar =
false;
1832 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1833 TII.get(SPIRV::OpVectorShuffle))
1838 for (
unsigned J = 1; J < ComponentCount * 2; J += 2)
1843 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1844 TII.get(SPIRV::OpVectorShuffle))
1849 for (
unsigned J = 0; J < ComponentCount * 2; J += 2)
1857bool SPIRVInstructionSelector::selectPopCount16(
Register ResVReg,
1858 SPIRVTypeInst ResType,
1861 unsigned Opcode)
const {
1862 Register OpReg =
I.getOperand(1).getReg();
1865 MachineIRBuilder MIRBuilder(
I);
1867 SPIRVTypeInst I32VectorType =
1870 bool IsVector = NumElems > 1;
1871 SPIRVTypeInst ExtType = IsVector ? I32VectorType : I32Type;
1874 if (!selectOpWithSrcs(ExtReg, ExtType,
I, {OpReg}, SPIRV::OpUConvert))
1878 if (!selectPopCount32(PopCountReg, ExtType,
I, ExtReg, Opcode))
1881 return selectOpWithSrcs(ResVReg, ResType,
I, {PopCountReg}, ExtOpcode);
1884bool SPIRVInstructionSelector::selectPopCount32(
Register ResVReg,
1885 SPIRVTypeInst ResType,
1888 unsigned Opcode)
const {
1889 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
1892bool SPIRVInstructionSelector::selectPopCount64(
Register ResVReg,
1893 SPIRVTypeInst ResType,
1896 unsigned Opcode)
const {
1898 if (ComponentCount > 2)
1899 return handle64BitOverflow(
1900 ResVReg, ResType,
I, SrcReg, Opcode,
1902 unsigned O) {
return this->selectPopCount64(R,
T,
I, S, O); });
1904 MachineIRBuilder MIRBuilder(
I);
1909 I32Type, 2 * ComponentCount, MIRBuilder,
false);
1913 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
1918 if (!selectPopCount32(Pop32, VecI32Type,
I, Vec32, Opcode))
1922 auto MaybeParts = splitEvenOddLanes(Pop32, ComponentCount,
I, I32Type);
1925 SplitParts &Parts = *MaybeParts;
1928 unsigned OpAdd = Parts.IsScalar ? SPIRV::OpIAddS : SPIRV::OpIAddV;
1930 if (!selectOpWithSrcs(Sum, Parts.Type,
I, {Parts.High, Parts.Low}, OpAdd))
1935 unsigned ConvOp = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
1936 return selectOpWithSrcs(ResVReg, ResType,
I, {Sum}, ConvOp);
1939bool SPIRVInstructionSelector::selectPopCount(
Register ResVReg,
1940 SPIRVTypeInst ResType,
1942 unsigned Opcode)
const {
1947 if (!STI.getTargetTriple().isVulkanOS())
1948 return selectUnOp(ResVReg, ResType,
I, Opcode);
1950 Register OpReg =
I.getOperand(1).getReg();
1953 : SPIRV::OpUConvert;
1957 return selectPopCount16(ResVReg, ResType,
I, ExtOpcode, Opcode);
1959 return selectPopCount32(ResVReg, ResType,
I, OpReg, Opcode);
1961 return selectPopCount64(ResVReg, ResType,
I, OpReg, Opcode);
1963 return diagnoseUnsupported(
I,
"unsupported operand bit width for popcount");
1967bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1968 SPIRVTypeInst ResType,
1970 unsigned Opcode)
const {
1972 Register SrcReg =
I.getOperand(1).getReg();
1977 unsigned DefOpCode = DefIt->getOpcode();
1978 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1981 if (
auto *VRD =
getVRegDef(*MRI, DefIt->getOperand(1).getReg()))
1982 DefOpCode = VRD->getOpcode();
1984 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1985 DefOpCode == TargetOpcode::G_CONSTANT ||
1986 DefOpCode == SPIRV::OpVariable ||
1987 DefOpCode == SPIRV::OpUntypedVariableKHR ||
1988 DefOpCode == SPIRV::OpConstantI) {
1994 uint32_t SpecOpcode = 0;
1996 case SPIRV::OpConvertPtrToU:
1997 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1999 case SPIRV::OpConvertUToPtr:
2000 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
2005 TII.get(SPIRV::OpSpecConstantOp))
2015 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
2019bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
2020 SPIRVTypeInst ResType,
2021 MachineInstr &
I)
const {
2022 Register OpReg =
I.getOperand(1).getReg();
2023 SPIRVTypeInst OpType =
2026 return diagnoseUnsupported(
2027 I,
"incompatible result and operand types in a bitcast");
2028 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
2039 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
2040 if (
MemOp->isNonTemporal())
2041 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
2043 if (!ST->isShader() &&
MemOp->getAlign().value())
2044 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
2048 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
2049 if (
auto *MD =
MemOp->getAAInfo().Scope) {
2053 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
2055 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
2059 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
2063 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
2065 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
2077 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
2079 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
2081 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
2085bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
2086 SPIRVTypeInst ResType,
2087 MachineInstr &
I)
const {
2089 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2094 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2095 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2097 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2099 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2103 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2107 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2108 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
2109 I.getDebugLoc(),
I);
2113 MachineIRBuilder MIRBuilder(
I);
2115 if (
I.getNumMemOperands()) {
2116 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2117 if (MemOp->isAtomic())
2118 return selectAtomicLoad(ResVReg, ResType,
I);
2121 auto MIB = MIRBuilder.
buildInstr(SPIRV::OpLoad)
2125 if (!
I.getNumMemOperands()) {
2126 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2128 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2137Register SPIRVInstructionSelector::createPtrSizedIntReg(
2138 MachineIRBuilder &MIRBuilder)
const {
2139 SPIRVTypeInst IntType =
2149SPIRVInstructionSelector::convertPtrToInt(
Register PtrVal,
2150 MachineIRBuilder &MIRBuilder)
const {
2151 SPIRVTypeInst IntType =
2153 Register IntReg = createPtrSizedIntReg(MIRBuilder);
2154 MIRBuilder.
buildInstr(SPIRV::OpConvertPtrToU)
2162Register SPIRVInstructionSelector::castPtrToPtrToInt(
2163 Register Ptr, SPIRV::StorageClass::StorageClass SC,
2164 MachineIRBuilder &MIRBuilder)
const {
2165 SPIRVTypeInst IntType =
2167 SPIRVTypeInst PtrType =
2181bool SPIRVInstructionSelector::selectAtomicPtrValue(
2182 Register ResVReg, SPIRVTypeInst ResType, MachineIRBuilder &MIRBuilder,
2183 function_ref<
Register(SPIRVTypeInst IntType)> EmitAtomic)
const {
2191 Register IntResult = EmitAtomic(IntType);
2193 MIRBuilder.
buildInstr(SPIRV::OpConvertUToPtr)
2201bool SPIRVInstructionSelector::selectAtomicLoad(
Register ResVReg,
2202 SPIRVTypeInst ResType,
2203 MachineInstr &
I)
const {
2204 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2207 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2210 return diagnoseUnsupported(
2211 I,
"Lowering to SPIR-V of atomic load is only "
2212 "allowed for integer, floating point or pointer types");
2214 assert(
I.getNumMemOperands());
2215 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2216 assert(MemOp.isAtomic());
2218 uint32_t
Scope =
static_cast<uint32_t
>(
2220 Register ScopeReg = buildI32Constant(Scope,
I);
2226 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2227 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2230 Register MemSemReg = buildI32Constant(Sem,
I);
2232 MachineIRBuilder MIRBuilder(
I);
2236 return diagnoseUnsupported(
2237 I,
"Lowering to SPIR-V of atomic load is only "
2238 "allowed for pointer types for physical addressing model");
2243 SPIRV::StorageClass::StorageClass SC =
2245 return selectAtomicPtrValue(
2246 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2247 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2248 Register IntResult = createPtrSizedIntReg(MIRBuilder);
2259 auto AtomicLoad = MIRBuilder.
buildInstr(SPIRV::OpAtomicLoad)
2270bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
2272 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2273 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2278 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2279 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2281 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2286 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2290 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2291 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2292 SPIRVTypeInst SampledType =
2294 SPIRVTypeInst StoreValCompType =
2296 if (StoreValCompType && StoreValCompType != SampledType) {
2299 SPIRVTypeInst PackedType = widenTypeToVec4(SampledType,
I);
2302 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
2307 StoreVal = PackedReg;
2310 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2311 TII.get(SPIRV::OpImageWrite))
2317 if (sampledTypeIsSignedInteger(LLVMHandleType))
2320 BMI.constrainAllUses(
TII,
TRI, RBI);
2327 if (PointeeTy && PointeeTy->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
2328 StoreTy->
getOpcode() != SPIRV::OpTypeVectorIdEXT &&
2330 MachineInstr *StoreValDef =
getVRegDef(*MRI, StoreVal);
2342 if (
I.getNumMemOperands()) {
2343 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2344 if (MemOp->isAtomic())
2345 return selectAtomicStore(
I);
2352 if (PtrSC == SPIRV::StorageClass::UniformConstant ||
2353 PtrSC == SPIRV::StorageClass::Input ||
2354 PtrSC == SPIRV::StorageClass::PushConstant)
2355 return diagnoseUnsupported(
2356 I,
"store into a read-only SPIR-V storage class is not allowed");
2358 MachineIRBuilder MIRBuilder(
I);
2360 if (!
I.getNumMemOperands()) {
2361 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2363 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2372bool SPIRVInstructionSelector::selectAtomicStore(MachineInstr &
I)
const {
2373 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2376 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2377 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2382 if (!PointeeType && PtrType &&
2383 PtrType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR)
2386 return diagnoseUnsupported(
I,
2387 "Lowering to SPIR-V of atomic store is only "
2388 "allowed for integer or floating point types");
2390 assert(
I.getNumMemOperands());
2391 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2392 assert(MemOp.isAtomic());
2394 uint32_t
Scope =
static_cast<uint32_t
>(
2396 Register ScopeReg = buildI32Constant(Scope,
I);
2402 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2403 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2406 Register MemSemReg = buildI32Constant(Sem,
I);
2407 MachineIRBuilder MIRBuilder(
I);
2411 return diagnoseUnsupported(
2412 I,
"Lowering to SPIR-V of atomic store is only "
2413 "allowed for pointer types for physical addressing model");
2418 SPIRV::StorageClass::StorageClass SC =
2420 return selectAtomicPtrValue(
2421 Register(), SPIRVTypeInst(), MIRBuilder, [&](SPIRVTypeInst IntType) {
2423 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2436 return diagnoseUnsupported(
I,
2437 "Lowering to SPIR-V of atomic store is only "
2438 "allowed for integer or floating point types");
2440 auto AtomicStore = MIRBuilder.
buildInstr(SPIRV::OpAtomicStore)
2450bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
2451 SPIRVTypeInst ResType,
2452 MachineInstr &
I)
const {
2453 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
2461 const Register PtrsReg =
I.getOperand(2).getReg();
2462 const uint32_t
Alignment =
I.getOperand(3).getImm();
2463 const Register MaskReg =
I.getOperand(4).getReg();
2464 const Register PassthruReg =
I.getOperand(5).getReg();
2465 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2469 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
2480bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
2481 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
2488 const Register ValuesReg =
I.getOperand(1).getReg();
2489 const Register PtrsReg =
I.getOperand(2).getReg();
2490 const uint32_t
Alignment =
I.getOperand(3).getImm();
2491 const Register MaskReg =
I.getOperand(4).getReg();
2492 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2496 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
2505bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
2506 const Twine &
Msg)
const {
2507 const Function &
F =
I.getMF()->getFunction();
2508 F.getContext().diagnose(
2509 DiagnosticInfoUnsupported(
F,
Msg,
I.getDebugLoc(),
DS_Error));
2513bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
2514 SPIRVTypeInst ResType,
2515 MachineInstr &
I)
const {
2516 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2517 return diagnoseUnsupported(
2518 I,
"llvm.stacksave intrinsic: this instruction requires the following "
2519 "SPIR-V extension: SPV_INTEL_variable_length_array");
2521 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
2528bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
2529 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2530 return diagnoseUnsupported(
2532 "llvm.stackrestore intrinsic: this instruction requires the following "
2533 "SPIR-V extension: SPV_INTEL_variable_length_array");
2534 if (!
I.getOperand(0).isReg())
2537 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
2538 .
addUse(
I.getOperand(0).getReg())
2544SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
2545 MachineIRBuilder MIRBuilder(
I);
2546 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
2553 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
2557 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2558 Type *ArrTy = ArrayType::get(ValTy, Num);
2560 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
2563 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
2574 const bool IsUntyped = VarTy->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
2575 auto MIBVar =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2576 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
2577 : SPIRV::OpVariable))
2580 .
addImm(SPIRV::StorageClass::UniformConstant);
2593bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
2596 Register DstReg =
I.getOperand(0).getReg();
2600 return diagnoseUnsupported(
2601 I,
"OpCopyMemory requires operands to have the same type");
2606 return diagnoseUnsupported(
2607 I,
"Unable to determine pointee type size for OpCopyMemory");
2608 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
2609 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
2610 return diagnoseUnsupported(
2611 I,
"OpCopyMemory requires the size to match the pointee type size");
2612 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemory))
2615 if (
I.getNumMemOperands()) {
2616 MachineIRBuilder MIRBuilder(
I);
2623bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &
I,
2626 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemorySized))
2627 .
addUse(
I.getOperand(0).getReg())
2629 .
addUse(
I.getOperand(2).getReg());
2630 if (
I.getNumMemOperands()) {
2631 MachineIRBuilder MIRBuilder(
I);
2638bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
2639 MachineInstr &
I)
const {
2641 Register SizeReg =
I.getOperand(2).getReg();
2643 SizeDef && SizeDef->
getOpcode() == TargetOpcode::G_CONSTANT &&
2647 Register SrcReg =
I.getOperand(1).getReg();
2648 if (
I.getOpcode() == TargetOpcode::G_MEMSET ||
2649 I.getOpcode() == TargetOpcode::G_MEMSET_INLINE) {
2650 Register VarReg = getOrCreateMemSetGlobal(
I);
2653 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2655 ValTy,
I, SPIRV::StorageClass::UniformConstant);
2657 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
2661 if (!selectCopyMemory(
I, SrcReg))
2664 if (!selectCopyMemorySized(
I, SrcReg))
2667 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
2668 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
2673bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
2674 SPIRVTypeInst ResType,
2677 unsigned NegateOpcode)
const {
2679 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2680 uint32_t
Scope =
static_cast<uint32_t
>(
2682 MemOp->getSyncScopeID()));
2683 Register ScopeReg = buildI32Constant(Scope,
I);
2685 Register Ptr =
I.getOperand(1).getReg();
2686 uint32_t ScSem =
static_cast<uint32_t
>(
2690 Register MemSemReg = buildI32Constant(
2694 Register ValueReg =
I.getOperand(2).getReg();
2695 if (NegateOpcode != 0) {
2698 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
2704 if (NewOpcode != SPIRV::OpAtomicExchange)
2705 return diagnoseUnsupported(
2706 I,
"Lowering to SPIR-V of this atomic operation is not "
2707 "allowed for pointer types");
2709 return diagnoseUnsupported(
2710 I,
"Lowering to SPIR-V of atomic exchange is only "
2711 "allowed for pointer types for physical addressing model");
2718 MachineIRBuilder MIRBuilder(
I);
2720 return selectAtomicPtrValue(
2721 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2723 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2724 Register ExchangeResReg = createPtrSizedIntReg(MIRBuilder);
2725 MIRBuilder.
buildInstr(SPIRV::OpAtomicExchange)
2733 return ExchangeResReg;
2737 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2748bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2749 unsigned ArgI =
I.getNumOperands() - 1;
2751 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2752 SPIRVTypeInst SrcType =
2756 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2760 unsigned CurrentIndex = 0;
2761 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2762 Register ResVReg =
I.getOperand(i).getReg();
2765 LLT ResLLT = MRI->
getType(ResVReg);
2771 ResType = ScalarType;
2780 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2786 for (
unsigned j = 0;
j < NumElements; ++
j) {
2787 MIB.
addImm(CurrentIndex + j);
2789 CurrentIndex += NumElements;
2793 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2805bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2808 Register MemSemReg = buildI32ConstantInEntryBlock(MemSem,
I);
2812 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
2814 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2821bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2822 SPIRVTypeInst ResType,
2824 unsigned Opcode)
const {
2825 Type *ResTy =
nullptr;
2828 return diagnoseUnsupported(
2830 "Not enough info to select the arithmetic with overflow instruction");
2832 return diagnoseUnsupported(
I,
2833 "Expect struct type result for the arithmetic "
2834 "with overflow instruction");
2840 MachineIRBuilder MIRBuilder(
I);
2842 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2843 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2850 Register ZeroReg = buildZerosVal(ResType,
I);
2855 if (ResName.
size() > 0)
2863 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2864 MIB.
addUse(
I.getOperand(i).getReg());
2869 MRI->
setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2870 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2872 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2873 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2880 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2881 .
addDef(
I.getOperand(1).getReg())
2889bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2890 SPIRVTypeInst ResType,
2891 MachineInstr &
I)
const {
2893 "selectAtomicCmpXchg only handles the spv_cmpxchg intrinsic");
2894 Register Ptr =
I.getOperand(2).getReg();
2895 Register ScopeReg =
I.getOperand(5).getReg();
2896 Register MemSemEqReg =
I.getOperand(6).getReg();
2897 Register MemSemNeqReg =
I.getOperand(7).getReg();
2899 Register Val =
I.getOperand(4).getReg();
2903 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2922 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2929 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2941 case SPIRV::StorageClass::DeviceOnlyINTEL:
2942 case SPIRV::StorageClass::HostOnlyINTEL:
2951 bool IsGRef =
false;
2952 bool IsAllowedRefs =
2954 unsigned Opcode = It.getOpcode();
2955 if (Opcode == SPIRV::OpConstantComposite ||
2956 Opcode == SPIRV::OpSpecConstantComposite ||
2957 Opcode == SPIRV::OpVariable ||
2958 Opcode == SPIRV::OpUntypedVariableKHR ||
2959 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2960 return IsGRef = true;
2961 return Opcode == SPIRV::OpName;
2963 return IsAllowedRefs && IsGRef;
2966Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2967 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2969 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2973SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2975 uint32_t Opcode)
const {
2976 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2977 TII.get(SPIRV::OpSpecConstantOp))
2985SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2986 SPIRVTypeInst SrcPtrTy)
const {
2987 SPIRVTypeInst GenericPtrTy =
2991 SPIRV::StorageClass::Generic),
2995 MachineInstrBuilder MIB = buildSpecConstantOp(
2997 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
3007bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
3008 SPIRVTypeInst ResType,
3009 MachineInstr &
I)
const {
3013 Register SrcPtr =
I.getOperand(1).getReg();
3018 return BuildCOPY(ResVReg, SrcPtr,
I);
3028 unsigned SpecOpcode = [&]() ->
unsigned {
3029 if (SrcSC == SPIRV::StorageClass::CodeSectionINTEL)
3030 return static_cast<uint32_t
>(SPIRV::Opcode::Bitcast);
3032 return static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric);
3034 return static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr);
3042 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
3044 .constrainAllUses(
TII,
TRI, RBI);
3046 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
3048 buildSpecConstantOp(
3050 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
3051 .constrainAllUses(
TII,
TRI, RBI);
3058 return BuildCOPY(ResVReg, SrcPtr,
I);
3060 if ((SrcSC == SPIRV::StorageClass::Function &&
3061 DstSC == SPIRV::StorageClass::Private) ||
3062 (DstSC == SPIRV::StorageClass::Function &&
3063 SrcSC == SPIRV::StorageClass::Private))
3064 return BuildCOPY(ResVReg, SrcPtr,
I);
3068 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
3071 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
3074 SPIRVTypeInst GenericPtrTy =
3093 return selectUnOp(ResVReg, ResType,
I,
3094 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
3096 return selectUnOp(ResVReg, ResType,
I,
3097 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
3099 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
3101 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
3111bool SPIRVInstructionSelector::selectPtrMask(
Register ResVReg,
3112 SPIRVTypeInst ResType,
3113 MachineInstr &
I)
const {
3115 return diagnoseUnsupported(
3116 I,
"G_PTRMASK is not supported with logical SPIR-V");
3121 Register PtrReg =
I.getOperand(1).getReg();
3122 Register MaskReg =
I.getOperand(2).getReg();
3141 ? SPIRV::OpBitwiseAndV
3142 : SPIRV::OpBitwiseAndS;
3165 return SPIRV::OpFOrdEqual;
3167 return SPIRV::OpFOrdGreaterThanEqual;
3169 return SPIRV::OpFOrdGreaterThan;
3171 return SPIRV::OpFOrdLessThanEqual;
3173 return SPIRV::OpFOrdLessThan;
3175 return SPIRV::OpFOrdNotEqual;
3177 return SPIRV::OpOrdered;
3179 return SPIRV::OpFUnordEqual;
3181 return SPIRV::OpFUnordGreaterThanEqual;
3183 return SPIRV::OpFUnordGreaterThan;
3185 return SPIRV::OpFUnordLessThanEqual;
3187 return SPIRV::OpFUnordLessThan;
3189 return SPIRV::OpFUnordNotEqual;
3191 return SPIRV::OpUnordered;
3201 return SPIRV::OpIEqual;
3203 return SPIRV::OpINotEqual;
3205 return SPIRV::OpSGreaterThanEqual;
3207 return SPIRV::OpSGreaterThan;
3209 return SPIRV::OpSLessThanEqual;
3211 return SPIRV::OpSLessThan;
3213 return SPIRV::OpUGreaterThanEqual;
3215 return SPIRV::OpUGreaterThan;
3217 return SPIRV::OpULessThanEqual;
3219 return SPIRV::OpULessThan;
3228 return SPIRV::OpPtrEqual;
3230 return SPIRV::OpPtrNotEqual;
3241 return SPIRV::OpLogicalEqual;
3243 return SPIRV::OpLogicalNotEqual;
3281bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
3282 SPIRVTypeInst ResType,
3284 unsigned OpAnyOrAll)
const {
3285 assert(
I.getNumOperands() == 3);
3286 assert(
I.getOperand(2).isReg());
3288 Register InputRegister =
I.getOperand(2).getReg();
3291 assert(InputType &&
"VReg has no type assigned");
3295 assert(ResVReg ==
I.getOperand(0).getReg());
3296 return BuildCOPY(ResVReg, InputRegister,
I);
3300 unsigned SpirvNotEqualId =
3301 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
3303 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
3308 IsBoolTy ? InputRegister
3316 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
3318 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
3335bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
3336 SPIRVTypeInst ResType,
3337 MachineInstr &
I)
const {
3338 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
3341bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
3342 SPIRVTypeInst ResType,
3343 MachineInstr &
I)
const {
3344 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
3348bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
3349 SPIRVTypeInst ResType,
3350 MachineInstr &
I)
const {
3351 assert(
I.getNumOperands() == 4);
3352 assert(
I.getOperand(2).isReg());
3353 assert(
I.getOperand(3).isReg());
3355 [[maybe_unused]] SPIRVTypeInst VecType =
3360 VecType->
getOpcode() == SPIRV::OpTypeVectorIdEXT) &&
3361 "dot product requires either a vector of at least 2 components or"
3362 " the SPV_EXT_long vector extension.");
3364 [[maybe_unused]] SPIRVTypeInst EltType =
3373 .
addUse(
I.getOperand(2).getReg())
3374 .
addUse(
I.getOperand(3).getReg())
3379bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
3380 SPIRVTypeInst ResType,
3383 assert(
I.getNumOperands() == 4);
3384 assert(
I.getOperand(2).isReg());
3385 assert(
I.getOperand(3).isReg());
3388 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3392 .
addUse(
I.getOperand(2).getReg())
3393 .
addUse(
I.getOperand(3).getReg())
3400bool SPIRVInstructionSelector::selectIntegerDotExpansion(
3401 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3402 assert(
I.getNumOperands() == 4);
3403 assert(
I.getOperand(2).isReg());
3404 assert(
I.getOperand(3).isReg());
3408 Register Vec0 =
I.getOperand(2).getReg();
3409 Register Vec1 =
I.getOperand(3).getReg();
3413 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
3422 VecType->
getOpcode() == SPIRV::OpTypeVectorIdEXT) &&
3423 "dot product requires either a vector of at least 2 components "
3424 "or the SPV_EXT_long_vector extension.");
3427 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3437 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3448 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3460bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
3461 SPIRVTypeInst ResType,
3462 MachineInstr &
I)
const {
3464 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
3467 .
addUse(
I.getOperand(2).getReg())
3472bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
3473 SPIRVTypeInst ResType,
3474 MachineInstr &
I)
const {
3476 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
3479 .
addUse(
I.getOperand(2).getReg())
3484bool SPIRVInstructionSelector::selectOpIsFinite(
Register ResVReg,
3485 SPIRVTypeInst ResType,
3486 MachineInstr &
I)
const {
3488 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsFinite))
3491 .
addUse(
I.getOperand(2).getReg())
3496bool SPIRVInstructionSelector::selectOpIsNormal(
Register ResVReg,
3497 SPIRVTypeInst ResType,
3498 MachineInstr &
I)
const {
3500 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNormal))
3503 .
addUse(
I.getOperand(2).getReg())
3508template <
bool Signed>
3509bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
3510 SPIRVTypeInst ResType,
3511 MachineInstr &
I)
const {
3512 assert(
I.getNumOperands() == 5);
3513 assert(
I.getOperand(2).isReg());
3514 assert(
I.getOperand(3).isReg());
3515 assert(
I.getOperand(4).isReg());
3518 Register Acc =
I.getOperand(2).getReg();
3522 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3524 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
3529 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
3532 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3544template <
bool Signed>
3545bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
3546 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3547 assert(
I.getNumOperands() == 5);
3548 assert(
I.getOperand(2).isReg());
3549 assert(
I.getOperand(3).isReg());
3550 assert(
I.getOperand(4).isReg());
3553 Register Acc =
I.getOperand(2).getReg();
3559 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
3563 for (
unsigned i = 0; i < 4; i++) {
3586 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
3606 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3621bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
3622 SPIRVTypeInst ResType,
3623 MachineInstr &
I)
const {
3624 assert(
I.getNumOperands() == 3);
3625 assert(
I.getOperand(2).isReg());
3627 Register VZero = buildZerosValF(ResType,
I);
3628 Register VOne = buildOnesValF(ResType,
I);
3630 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
3633 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3635 .
addUse(
I.getOperand(2).getReg())
3642bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
3643 SPIRVTypeInst ResType,
3644 MachineInstr &
I)
const {
3645 assert(
I.getNumOperands() == 3);
3646 assert(
I.getOperand(2).isReg());
3648 Register InputRegister =
I.getOperand(2).getReg();
3650 auto &
DL =
I.getDebugLoc();
3653 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3660 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
3662 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
3670 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3675 if (NeedsConversion) {
3676 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
3687bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
3688 SPIRVTypeInst ResType,
3690 unsigned Opcode)
const {
3694 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
3700 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
3701 BMI.addUse(
I.getOperand(J).getReg());
3708bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
3711 bool WithGroupSync)
const {
3713 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
3715 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
3717 assert(((Scope != SPIRV::Scope::Workgroup) ||
3718 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
3719 "Workgroup Scope must set WorkGroupMemory semantic "
3720 "in Barrier instruction");
3722 assert(((Scope != SPIRV::Scope::Device) ||
3723 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
3724 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
3725 "Device Scope must set UniformMemory and ImageMemory semantic "
3726 "in Barrier instruction");
3732 if (WithGroupSync) {
3733 Register ExecReg = buildI32Constant(SPIRV::Scope::Workgroup,
I);
3737 Register ScopeReg = buildI32Constant(Scope,
I);
3738 Register MemSemReg = buildI32Constant(MemSem,
I);
3740 MI.addUse(ScopeReg).addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
3744bool SPIRVInstructionSelector::selectWaveActiveCountBits(
3745 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3750 if (!selectWaveOpInst(BallotReg, BallotType,
I,
3751 SPIRV::OpGroupNonUniformBallot))
3756 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3761 .
addImm(SPIRV::GroupOperation::Reduce)
3768bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
3769 SPIRVTypeInst ResType,
3770 MachineInstr &
I)
const {
3775 Register InputReg =
I.getOperand(2).getReg();
3780 bool IsVector = NumElems > 1 ||
3781 (InputType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
3795 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
3796 SPIRV::OpGroupNonUniformAllEqual);
3801 ElementResults.
reserve(NumElems);
3803 for (
unsigned Idx = 0; Idx < NumElems; ++Idx) {
3816 ElemInput = Extracted;
3822 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3833 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3844bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3845 SPIRVTypeInst ResType,
3846 MachineInstr &
I)
const {
3848 assert(
I.getNumOperands() == 3);
3850 auto Op =
I.getOperand(2);
3860 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3862 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3863 return diagnoseUnsupported(
I,
"WavePrefixBitCount requires boolean input");
3884 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3888 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3895bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3896 SPIRVTypeInst ResType,
3898 bool IsUnsigned)
const {
3899 return selectWaveReduce(
3900 ResVReg, ResType,
I, IsUnsigned,
3901 [&](
Register InputRegister,
bool IsUnsigned) {
3902 const bool IsFloatTy =
3904 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3905 : SPIRV::OpGroupNonUniformSMax;
3906 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3910bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3911 SPIRVTypeInst ResType,
3913 bool IsUnsigned)
const {
3914 return selectWaveReduce(
3915 ResVReg, ResType,
I, IsUnsigned,
3916 [&](
Register InputRegister,
bool IsUnsigned) {
3917 const bool IsFloatTy =
3919 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3920 : SPIRV::OpGroupNonUniformSMin;
3921 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3925bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3926 SPIRVTypeInst ResType,
3927 MachineInstr &
I)
const {
3928 return selectWaveReduce(ResVReg, ResType,
I,
false,
3929 [&](
Register InputRegister,
bool IsUnsigned) {
3931 InputRegister, SPIRV::OpTypeFloat);
3932 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3933 : SPIRV::OpGroupNonUniformIAdd;
3937bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3938 SPIRVTypeInst ResType,
3939 MachineInstr &
I)
const {
3940 return selectWaveReduce(ResVReg, ResType,
I,
false,
3941 [&](
Register InputRegister,
bool IsUnsigned) {
3943 InputRegister, SPIRV::OpTypeFloat);
3944 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3945 : SPIRV::OpGroupNonUniformIMul;
3949template <
typename PickOpcodeFn>
3950bool SPIRVInstructionSelector::selectWaveReduce(
3951 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3952 PickOpcodeFn &&PickOpcode)
const {
3953 assert(
I.getNumOperands() == 3);
3954 assert(
I.getOperand(2).isReg());
3956 Register InputRegister =
I.getOperand(2).getReg();
3960 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3963 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3969 .
addImm(SPIRV::GroupOperation::Reduce)
3970 .
addUse(
I.getOperand(2).getReg())
3975bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3976 SPIRVTypeInst ResType,
3978 unsigned Opcode)
const {
3979 return selectWaveReduce(
3980 ResVReg, ResType,
I,
false,
3981 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3984bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3985 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3986 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3987 [&](
Register InputRegister,
bool IsUnsigned) {
3989 InputRegister, SPIRV::OpTypeFloat);
3991 ? SPIRV::OpGroupNonUniformFAdd
3992 : SPIRV::OpGroupNonUniformIAdd;
3996bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3997 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3998 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3999 [&](
Register InputRegister,
bool IsUnsigned) {
4001 InputRegister, SPIRV::OpTypeFloat);
4003 ? SPIRV::OpGroupNonUniformFMul
4004 : SPIRV::OpGroupNonUniformIMul;
4008template <
typename PickOpcodeFn>
4009bool SPIRVInstructionSelector::selectWaveExclusiveScan(
4010 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
4011 PickOpcodeFn &&PickOpcode)
const {
4012 assert(
I.getNumOperands() == 3);
4013 assert(
I.getOperand(2).isReg());
4015 Register InputRegister =
I.getOperand(2).getReg();
4019 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
4022 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
4028 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
4029 .
addUse(
I.getOperand(2).getReg())
4034bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
4035 SPIRVTypeInst ResType,
4038 assert(
I.getNumOperands() == 3);
4039 assert(
I.getOperand(2).isReg());
4041 Register InputRegister =
I.getOperand(2).getReg();
4047 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
4058bool SPIRVInstructionSelector::selectBitreverseViaI32(
Register ResVReg,
4059 SPIRVTypeInst ResType,
4066 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
4071 : SPIRV::OpUConvert;
4073 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4076 ShiftOp = SPIRV::OpShiftRightLogicalV;
4081 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4082 TII.get(SPIRV::OpConstantComposite))
4085 for (
unsigned It = 0; It <
N; ++It)
4089 ShiftConst = CompositeReg;
4094 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
4099 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
4104 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
4109 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
4112bool SPIRVInstructionSelector::handle64BitOverflow(
4114 unsigned int Opcode,
4121 "handle64BitOverflow should only be used for integer types");
4123 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
4125 MachineIRBuilder MIRBuilder(
I);
4127 SPIRVTypeInst I64x2Type =
4129 SPIRVTypeInst Vec2ResType =
4132 std::vector<Register> PartialRegs;
4134 unsigned CurrentComponent = 0;
4135 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
4139 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4140 TII.get(SPIRV::OpVectorShuffle))
4145 .
addImm(CurrentComponent)
4146 .
addImm(CurrentComponent + 1);
4156 PartialRegs.push_back(SubVecReg);
4159 if (CurrentComponent != ComponentCount) {
4165 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
4166 SPIRV::OpVectorExtractDynamic))
4175 PartialRegs.push_back(FinalElemResReg);
4179 return selectOpWithSrcs(ResVReg, ResType,
I, PartialRegs,
4180 SPIRV::OpCompositeConstruct);
4183bool SPIRVInstructionSelector::selectBitreverse64(
Register ResVReg,
4184 SPIRVTypeInst ResType,
4188 if (ComponentCount > 2)
4189 return handle64BitOverflow(
4190 ResVReg, ResType,
I, SrcReg, SPIRV::OpBitReverse,
4192 unsigned O) {
return this->selectBitreverse64(R,
T,
I, S); });
4194 MachineIRBuilder MIRBuilder(
I);
4198 I32Type, 2 * ComponentCount, MIRBuilder,
false);
4202 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
4207 if (!selectBitreverseNative(Reverse32, VecI32Type,
I, Vec32))
4214 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4215 TII.get(SPIRV::OpVectorShuffle))
4220 for (
unsigned J = 0; J < ComponentCount; ++J) {
4227 return selectOpWithSrcs(ResVReg, ResType,
I, {SwappedVec}, SPIRV::OpBitcast);
4230bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
4231 SPIRVTypeInst ResType,
4235 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
4243bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
4244 SPIRVTypeInst ResType,
4245 MachineInstr &
I)
const {
4246 Register OpReg =
I.getOperand(1).getReg();
4255 return selectBitreverseViaI32(ResVReg, ResType,
I, OpReg);
4257 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4259 return selectBitreverse64(ResVReg, ResType,
I, OpReg);
4261 return SPIRVInstructionSelector::diagnoseUnsupported(
4262 I,
"G_BITREVERSE only support 16,32,64 bits.");
4266 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4277 unsigned AndOp = SPIRV::OpBitwiseAndS;
4278 unsigned OrOp = SPIRV::OpBitwiseOrS;
4279 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
4280 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
4281 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4283 AndOp = SPIRV::OpBitwiseAndV;
4284 OrOp = SPIRV::OpBitwiseOrV;
4285 ShlOp = SPIRV::OpShiftLeftLogicalV;
4286 ShrOp = SPIRV::OpShiftRightLogicalV;
4292 const unsigned Shift) ->
Register {
4295 (ResType->
getOpcode() != SPIRV::OpTypeVectorIdEXT ||
4302 Register MaskReg = CreateConst(Mask);
4303 Register ShiftReg = CreateConst(Shift);
4310 if (!selectOpWithSrcs(
T1, ResType,
I, {Input, ShiftReg}, ShrOp) ||
4311 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
4312 !selectOpWithSrcs(T3, ResType,
I, {Input, MaskReg}, AndOp) ||
4313 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
4314 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
4323 while ((Shift >>= 1) > 0) {
4330 return BuildCOPY(ResVReg, Result,
I);
4333bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
4334 SPIRVTypeInst ResType,
4335 MachineInstr &
I)
const {
4336 assert(
I.getOperand(0).isReg() &&
I.getOperand(1).isReg() &&
4337 "G_FREEZE must define and use a register");
4338 Register OpReg =
I.getOperand(1).getReg();
4342 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
4355 if (MachineInstr *Def = MRI->
getVRegDef(OpReg)) {
4356 if (
Def->getOpcode() == TargetOpcode::COPY)
4359 switch (
Def->getOpcode()) {
4360 case SPIRV::ASSIGN_TYPE:
4361 if (MachineInstr *AssignToDef =
4363 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
4364 Reg =
Def->getOperand(2).getReg();
4367 case SPIRV::OpUndef:
4368 Reg =
Def->getOperand(1).getReg();
4371 unsigned DestOpCode;
4373 DestOpCode = SPIRV::OpConstantNull;
4374 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze of a "
4375 "static undef/poison lowered to OpConstantNull\n");
4377 DestOpCode = TargetOpcode::COPY;
4379 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze "
4380 "skipped, lowered as a copy of the operand\n");
4382 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
4383 .
addDef(
I.getOperand(0).getReg())
4391bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
4392 SPIRVTypeInst ResType,
4393 MachineInstr &
I)
const {
4397 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4401 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
4406 for (
unsigned i =
I.getNumExplicitDefs();
4407 i <
I.getNumExplicitOperands() && IsConst; ++i)
4412 return diagnoseUnsupported(
4413 I,
"There must be at least two constituent operands in a vector");
4418 for (
unsigned i =
I.getNumExplicitDefs();
4419 i <
I.getNumExplicitOperands() && IsNullVector; ++i) {
4420 MachineInstr *
Def =
getDef(
I.getOperand(i), MRI);
4425 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4432 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4433 TII.get(IsConst ? SPIRV::OpConstantComposite
4434 : SPIRV::OpCompositeConstruct))
4437 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
4438 MIB.
addUse(
I.getOperand(i).getReg());
4443bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
4444 SPIRVTypeInst ResType,
4445 MachineInstr &
I)
const {
4449 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4454 unsigned OpIdx =
I.getNumExplicitDefs();
4455 if (!
I.getOperand(OpIdx).isReg())
4459 Register OpReg =
I.getOperand(OpIdx).getReg();
4463 return diagnoseUnsupported(
4464 I,
"There must be at least two constituent operands in a vector");
4467 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4468 TII.get(IsConst ? SPIRV::OpConstantComposite
4469 : SPIRV::OpCompositeConstruct))
4472 for (
unsigned i = 0; i <
N; ++i)
4478bool SPIRVInstructionSelector::selectConcatVectors(
Register ResVReg,
4479 SPIRVTypeInst ResType,
4480 MachineInstr &
I)
const {
4486 "Cannot select G_CONCAT_VECTORS with a non-vector result");
4488 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4489 TII.get(SPIRV::OpCompositeConstruct))
4492 for (
unsigned OpIdx =
I.getNumExplicitDefs();
4493 OpIdx <
I.getNumExplicitOperands(); ++OpIdx)
4494 MIB.
addUse(
I.getOperand(OpIdx).getReg());
4499bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
4500 SPIRVTypeInst ResType,
4501 MachineInstr &
I)
const {
4507 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
4509 Opcode = SPIRV::OpDemoteToHelperInvocation;
4511 Opcode = SPIRV::OpKill;
4516 ToErase.eraseFromParent();
4525bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
4526 SPIRVTypeInst ResType,
unsigned CmpOpc,
4527 MachineInstr &
I)
const {
4528 Register Cmp0 =
I.getOperand(2).getReg();
4529 Register Cmp1 =
I.getOperand(3).getReg();
4532 "CMP operands should have the same type");
4533 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
4543bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
4544 SPIRVTypeInst ResType,
4545 MachineInstr &
I)
const {
4546 auto Pred =
I.getOperand(1).getPredicate();
4549 Register CmpOperand =
I.getOperand(2).getReg();
4551 bool IsPtrCmp = CmpOperandType && CmpOperandType.
isPointer();
4556 Register Op1 =
I.getOperand(3).getReg();
4560 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
4565 I.getOperand(3).setReg(NewOp1);
4571 return selectCmp(ResVReg, ResType, CmpOpc,
I);
4575SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
4576 SPIRVTypeInst ResType)
const {
4578 SPIRVTypeInst SpvI32Ty =
4581 auto ConstInt = ConstantInt::get(LLVMTy, Val);
4588 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4591 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
4594 .
addImm(APInt(32, Val).getZExtValue());
4596 GR.
add(ConstInt,
MI);
4603Register SPIRVInstructionSelector::buildI32ConstantInEntryBlock(
4604 uint32_t Val, MachineInstr &
I, SPIRVTypeInst ResType)
const {
4606 SPIRVTypeInst SpvI32Ty =
4608 auto *ConstInt = ConstantInt::get(LLVMTy, Val);
4613 MachineBasicBlock &EntryBB = *InsertIt->getParent();
4614 MachineInstr *
MI =
nullptr;
4618 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantNull))
4622 uint64_t ImmVal = APInt(32, Val).getZExtValue();
4623 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantI))
4629 GR.
add(ConstInt,
MI);
4634bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
4635 SPIRVTypeInst ResType,
4636 MachineInstr &
I)
const {
4638 return selectCmp(ResVReg, ResType, CmpOp,
I);
4641bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
4642 SPIRVTypeInst ResType,
4643 MachineInstr &
I)
const {
4645 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
4655 MachineIRBuilder MIRBuilder(
I);
4662 APFloat ConstVal(3.3219280948873623);
4666 APFloat::rmNearestTiesToEven, &LosesInfo);
4671 isVectorType(ResType) ? SPIRV::OpVectorTimesScalar : SPIRV::OpFMulS;
4673 if (!selectOpWithSrcs(ArgReg, ResType,
I,
4674 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
4676 if (!selectExtInst(ResVReg, ResType,
I,
4677 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
4687Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
4688 MachineInstr &
I)
const {
4696bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
4702 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
4710 if (
Def->getOpcode() == SPIRV::OpConstantNull)
4713 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
4714 Def->getOpcode() == SPIRV::OpConstantI)
4727 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
4728 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
4730 Intrinsic::spv_const_composite)) {
4731 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
4732 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
4733 if (!IsZero(
Def->getOperand(i).getReg()))
4742Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
4743 MachineInstr &
I)
const {
4752Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
4753 MachineInstr &
I)
const {
4763 SPIRVTypeInst ResType,
4764 MachineInstr &
I)
const {
4773bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
4774 SPIRVTypeInst ResType,
4775 MachineInstr &
I)
const {
4776 Register SelectFirstArg =
I.getOperand(2).getReg();
4777 Register SelectSecondArg =
I.getOperand(3).getReg();
4791 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
4792 }
else if (IsPtrTy) {
4793 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
4795 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
4798 assert(IsScalarBool &&
"OpSelect with a scalar result requires a scalar "
4799 "boolean condition");
4801 Opcode = SPIRV::OpSelectSFSCond;
4802 }
else if (IsPtrTy) {
4803 Opcode = SPIRV::OpSelectSPSCond;
4805 Opcode = SPIRV::OpSelectSISCond;
4808 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4811 .
addUse(
I.getOperand(1).getReg())
4820bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
4821 SPIRVTypeInst ResType,
4823 MachineInstr &InsertAt,
4824 bool IsSigned)
const {
4826 Register ZeroReg = buildZerosVal(ResType, InsertAt);
4827 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
4828 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
4830 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4842bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
4843 SPIRVTypeInst ResType,
4844 MachineInstr &
I,
bool IsSigned,
4845 unsigned Opcode)
const {
4846 Register SrcReg =
I.getOperand(1).getReg();
4857 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I, IsSigned);
4859 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
4862bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
4863 SPIRVTypeInst ResType, MachineInstr &
I,
4864 bool IsSigned)
const {
4865 Register SrcReg =
I.getOperand(1).getReg();
4867 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
4871 if (ResType == SrcType)
4872 return BuildCOPY(ResVReg, SrcReg,
I);
4874 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4875 return selectUnOp(ResVReg, ResType,
I, Opcode);
4878bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
4879 SPIRVTypeInst ResType,
4881 bool IsSigned)
const {
4882 MachineIRBuilder MIRBuilder(
I);
4883 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
4888 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4896 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
4899 .
addUse(
I.getOperand(1).getReg())
4900 .
addUse(
I.getOperand(2).getReg())
4905 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
4908 .
addUse(
I.getOperand(1).getReg())
4909 .
addUse(
I.getOperand(2).getReg())
4917 unsigned SelectOpcode =
4918 (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4920 ? SPIRV::OpSelectVIVCond
4921 : SPIRV::OpSelectSISCond;
4926 .
addUse(buildOnesVal(
true, ResType,
I))
4927 .
addUse(buildZerosVal(ResType,
I))
4934 .
addUse(buildOnesVal(
false, ResType,
I))
4939bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
4942 SPIRVTypeInst IntTy,
4943 SPIRVTypeInst BoolTy)
const {
4947 isVectorType(IntTy) ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
4949 Register One = buildOnesVal(
false, IntTy,
I);
4957 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
4966bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
4967 SPIRVTypeInst ResType,
4968 MachineInstr &
I)
const {
4969 Register IntReg =
I.getOperand(1).getReg();
4972 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
4973 if (ArgType == ResType)
4974 return BuildCOPY(ResVReg, IntReg,
I);
4976 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4977 return selectUnOp(ResVReg, ResType,
I, Opcode);
4980bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
4981 SPIRVTypeInst ResType,
4982 MachineInstr &
I)
const {
4983 unsigned Opcode =
I.getOpcode();
4984 unsigned TpOpcode = ResType->
getOpcode();
4986 if (ResType.
isPointer() || TpOpcode == SPIRV::OpTypeEvent) {
4987 assert(Opcode == TargetOpcode::G_CONSTANT &&
4988 I.getOperand(1).getCImm()->isZero());
4989 MachineBasicBlock &DepMBB =
I.getMF()->front();
4992 }
else if (TpOpcode == SPIRV::OpTypeVectorIdEXT) {
4997 "Expected <1 x T> Vector!");
4998 if (Opcode == TargetOpcode::G_FCONSTANT)
5004 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
5012 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
5015bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
5016 SPIRVTypeInst ResType,
5017 MachineInstr &
I)
const {
5018 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5025bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
5026 SPIRVTypeInst ResType,
5027 MachineInstr &
I)
const {
5029 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
5033 .
addUse(
I.getOperand(3).getReg())
5035 .
addUse(
I.getOperand(2).getReg());
5036 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
5042bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
5043 SPIRVTypeInst ResType,
5044 MachineInstr &
I)
const {
5045 Type *MaybeResTy =
nullptr;
5050 "Expected aggregate type for extractv instruction");
5052 SPIRV::AccessQualifier::ReadWrite,
false);
5056 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
5059 .
addUse(
I.getOperand(2).getReg());
5060 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
5066bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
5067 SPIRVTypeInst ResType,
5068 MachineInstr &
I)
const {
5069 if (
getImm(
I.getOperand(4), MRI))
5070 return selectInsertVal(ResVReg, ResType,
I);
5072 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
5075 .
addUse(
I.getOperand(2).getReg())
5076 .
addUse(
I.getOperand(3).getReg())
5077 .
addUse(
I.getOperand(4).getReg())
5082bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
5083 SPIRVTypeInst ResType,
5084 MachineInstr &
I)
const {
5085 if (
getImm(
I.getOperand(3), MRI))
5086 return selectExtractVal(ResVReg, ResType,
I);
5088 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
5091 .
addUse(
I.getOperand(2).getReg())
5092 .
addUse(
I.getOperand(3).getReg())
5097bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
5098 SPIRVTypeInst ResType,
5099 MachineInstr &
I)
const {
5100 const bool IsGEPInBounds =
I.getOperand(2).getImm();
5103 const bool UseUntypedPointers =
5104 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
5109 if (UseUntypedPointers) {
5111 Opcode = IsGEPInBounds ? SPIRV::OpUntypedInBoundsAccessChainKHR
5112 : SPIRV::OpUntypedAccessChainKHR;
5114 Opcode = IsGEPInBounds ? SPIRV::OpUntypedInBoundsPtrAccessChainKHR
5115 : SPIRV::OpUntypedPtrAccessChainKHR;
5124 IsGEPInBounds ? SPIRV::OpInBoundsAccessChain : SPIRV::OpAccessChain;
5126 Opcode = IsGEPInBounds ? SPIRV::OpInBoundsPtrAccessChain
5127 : SPIRV::OpPtrAccessChain;
5132 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5137 if (UseUntypedPointers) {
5152 while (Def &&
Def->getOpcode() == TargetOpcode::COPY &&
5153 Def->getOperand(1).isReg())
5155 if (Def &&
Def->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)
5156 if (
const auto *GVar =
5159 SPIRV::AccessQualifier::ReadWrite,
5163 return diagnoseUnsupported(
5164 I,
"could not deduce the base type of an untyped access chain");
5169 Res.addUse(BaseReg);
5171 const bool IsAccessChainOpcode =
5172 (Opcode == SPIRV::OpAccessChain ||
5173 Opcode == SPIRV::OpInBoundsAccessChain ||
5174 Opcode == SPIRV::OpUntypedAccessChainKHR ||
5175 Opcode == SPIRV::OpUntypedInBoundsAccessChainKHR);
5177 assert((!IsAccessChainOpcode || (
getImm(
I.getOperand(4), MRI) &&
5178 foldImm(
I.getOperand(4), MRI) == 0)) &&
5179 "Cannot translate GEP to OpAccessChain.");
5182 const unsigned StartingIndex = IsAccessChainOpcode ? 5 : 4;
5183 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
5184 Res.addUse(
I.getOperand(i).getReg());
5185 Res.constrainAllUses(
TII,
TRI, RBI);
5194 if (Extract.
getOpcode() == SPIRV::OpCompositeExtract) {
5200 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5206 TII.get(SPIRV::OpCompositeInsert))
5220bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
5222 unsigned Lim =
I.getNumExplicitOperands();
5223 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
5224 Register OpReg =
I.getOperand(i).getReg();
5225 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
5227 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
5228 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
5229 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
5242 SPIRVTypeInst WrapType = OpType;
5243 if (OpType->
getOpcode() == SPIRV::OpTypePointer &&
5245 SPIRV::StorageClass::CodeSectionINTEL) {
5247 SPIRV::StorageClass::Function,
I);
5254 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5255 TII.get(SPIRV::OpSpecConstantOp))
5258 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
5260 GR.
add(OpDefine, MIB);
5266bool SPIRVInstructionSelector::selectDerivativeInst(
5267 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
5268 const unsigned DPdOpCode)
const {
5271 if (!errorIfInstrOutsideShader(
I))
5277 Register SrcReg =
I.getOperand(2).getReg();
5282 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5285 .
addUse(
I.getOperand(2).getReg());
5287 MachineIRBuilder MIRBuilder(
I);
5290 if (componentCount != 1)
5298 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5303 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5308 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5316bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
5317 SPIRVTypeInst ResType,
5318 MachineInstr &
I)
const {
5322 case Intrinsic::spv_load:
5323 return selectLoad(ResVReg, ResType,
I);
5324 case Intrinsic::spv_atomic_load:
5325 return selectAtomicLoad(ResVReg, ResType,
I);
5326 case Intrinsic::spv_store:
5327 return selectStore(
I);
5328 case Intrinsic::spv_atomic_store:
5329 return selectAtomicStore(
I);
5330 case Intrinsic::spv_extractv:
5331 return selectExtractVal(ResVReg, ResType,
I);
5332 case Intrinsic::spv_insertv:
5333 return selectInsertVal(ResVReg, ResType,
I);
5334 case Intrinsic::spv_extractelt:
5335 return selectExtractElt(ResVReg, ResType,
I);
5336 case Intrinsic::spv_insertelt:
5337 return selectInsertElt(ResVReg, ResType,
I);
5338 case Intrinsic::spv_gep:
5339 return selectGEP(ResVReg, ResType,
I);
5340 case Intrinsic::spv_bitcast: {
5341 Register OpReg =
I.getOperand(2).getReg();
5342 SPIRVTypeInst OpType =
5346 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
5348 case Intrinsic::spv_unref_global:
5349 case Intrinsic::spv_init_global: {
5350 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
5355 Register GVarVReg =
MI->getOperand(0).getReg();
5356 if (!selectGlobalValue(GVarVReg, *
MI, Init))
5361 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
5363 MI->eraseFromParent();
5367 case Intrinsic::spv_undef: {
5368 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5374 case Intrinsic::spv_poison:
5375 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpPoisonKHR))
5380 case Intrinsic::spv_freeze:
5381 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
5384 .
addUse(
I.getOperand(2).getReg())
5387 case Intrinsic::spv_named_boolean_spec_constant: {
5388 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
5389 : SPIRV::OpSpecConstantFalse;
5391 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
5392 .
addDef(
I.getOperand(0).getReg())
5395 unsigned SpecId =
I.getOperand(2).getImm();
5397 SPIRV::Decoration::SpecId, {SpecId});
5401 case Intrinsic::spv_const_composite: {
5403 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
5409 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
5411 std::function<bool(
Register)> HasSpecConstOperand =
5421 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
5422 J < Def->getNumExplicitOperands(); ++J) {
5423 if (
Def->getOperand(J).isReg() &&
5424 HasSpecConstOperand(
Def->getOperand(J).getReg()))
5430 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
5431 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
5432 : SPIRV::OpConstantComposite;
5433 unsigned ContinuedOpc = HasSpecConst
5434 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
5435 : SPIRV::OpConstantCompositeContinuedINTEL;
5436 MachineIRBuilder MIR(
I);
5438 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
5440 for (
auto *Instr : Instructions) {
5441 Instr->setDebugLoc(
I.getDebugLoc());
5446 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5453 case Intrinsic::spv_assign_name: {
5454 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
5455 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
5456 for (
unsigned i =
I.getNumExplicitDefs() + 2;
5457 i <
I.getNumExplicitOperands(); ++i) {
5458 MIB.
addImm(
I.getOperand(i).getImm());
5463 case Intrinsic::spv_switch: {
5464 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
5465 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5466 if (
I.getOperand(i).isReg())
5467 MIB.
addReg(
I.getOperand(i).getReg());
5468 else if (
I.getOperand(i).isCImm())
5469 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
5470 else if (
I.getOperand(i).isMBB())
5471 MIB.
addMBB(
I.getOperand(i).getMBB());
5478 case Intrinsic::spv_loop_merge: {
5479 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
5480 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5481 if (
I.getOperand(i).isMBB())
5482 MIB.
addMBB(
I.getOperand(i).getMBB());
5489 case Intrinsic::spv_loop_control_intel: {
5491 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
5492 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
5497 case Intrinsic::spv_selection_merge: {
5499 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
5500 assert(
I.getOperand(1).isMBB() &&
5501 "operand 1 to spv_selection_merge must be a basic block");
5502 MIB.
addMBB(
I.getOperand(1).getMBB());
5503 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
5507 case Intrinsic::spv_cmpxchg:
5508 return selectAtomicCmpXchg(ResVReg, ResType,
I);
5509 case Intrinsic::spv_unreachable:
5510 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
5513 case Intrinsic::spv_abort:
5514 return selectAbort(
I);
5515 case Intrinsic::spv_alloca:
5516 return selectFrameIndex(ResVReg, ResType,
I);
5517 case Intrinsic::spv_alloca_array:
5518 return selectAllocaArray(ResVReg, ResType,
I);
5519 case Intrinsic::spv_assume:
5521 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
5522 .
addUse(
I.getOperand(1).getReg())
5527 case Intrinsic::spv_expect:
5529 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
5532 .
addUse(
I.getOperand(2).getReg())
5533 .
addUse(
I.getOperand(3).getReg())
5538 case Intrinsic::arithmetic_fence:
5539 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
5540 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
5543 .
addUse(
I.getOperand(2).getReg())
5547 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
5549 case Intrinsic::spv_thread_id:
5555 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
5557 case Intrinsic::spv_thread_id_in_group:
5563 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
5565 case Intrinsic::spv_group_id:
5571 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
5573 case Intrinsic::spv_flattened_thread_id_in_group:
5580 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
5582 case Intrinsic::spv_workgroup_size:
5583 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
5585 case Intrinsic::spv_global_size:
5586 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
5588 case Intrinsic::spv_global_offset:
5589 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
5591 case Intrinsic::spv_num_workgroups:
5592 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
5594 case Intrinsic::spv_subgroup_size:
5595 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
5597 case Intrinsic::spv_num_subgroups:
5598 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
5600 case Intrinsic::spv_subgroup_id:
5601 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
5602 case Intrinsic::spv_subgroup_local_invocation_id:
5603 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
5604 ResVReg, ResType,
I);
5605 case Intrinsic::spv_subgroup_max_size:
5606 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
5608 case Intrinsic::spv_fdot:
5609 return selectFloatDot(ResVReg, ResType,
I);
5610 case Intrinsic::spv_udot:
5611 case Intrinsic::spv_sdot:
5612 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5614 return selectIntegerDot(ResVReg, ResType,
I,
5615 IID == Intrinsic::spv_sdot);
5616 return selectIntegerDotExpansion(ResVReg, ResType,
I);
5617 case Intrinsic::spv_dot4add_i8packed:
5618 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5620 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
5621 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
5622 case Intrinsic::spv_dot4add_u8packed:
5623 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5625 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
5626 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
5627 case Intrinsic::spv_all:
5628 return selectAll(ResVReg, ResType,
I);
5629 case Intrinsic::spv_any:
5630 return selectAny(ResVReg, ResType,
I);
5631 case Intrinsic::spv_distance:
5632 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
5633 case Intrinsic::spv_lerp:
5634 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
5635 case Intrinsic::spv_length:
5636 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
5637 case Intrinsic::spv_degrees:
5638 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
5639 case Intrinsic::spv_faceforward:
5640 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
5641 case Intrinsic::spv_frac:
5642 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
5643 case Intrinsic::spv_isinf:
5644 return selectOpIsInf(ResVReg, ResType,
I);
5645 case Intrinsic::spv_isnan:
5646 return selectOpIsNan(ResVReg, ResType,
I);
5647 case Intrinsic::spv_isfinite:
5648 return selectOpIsFinite(ResVReg, ResType,
I);
5649 case Intrinsic::spv_isnormal:
5650 return selectOpIsNormal(ResVReg, ResType,
I);
5651 case Intrinsic::spv_normalize:
5652 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
5653 case Intrinsic::spv_refract:
5654 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
5655 case Intrinsic::spv_reflect:
5656 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
5657 case Intrinsic::spv_rsqrt:
5658 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
5659 case Intrinsic::spv_sign:
5660 return selectSign(ResVReg, ResType,
I);
5661 case Intrinsic::spv_smoothstep:
5662 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
5663 case Intrinsic::spv_firstbituhigh:
5664 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
5665 case Intrinsic::spv_firstbitshigh:
5666 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
5667 case Intrinsic::spv_firstbitlow:
5668 return selectFirstBitLow(ResVReg, ResType,
I);
5669 case Intrinsic::spv_all_memory_barrier:
5670 return selectBarrierInst(
I, SPIRV::Scope::Device,
5671 SPIRV::MemorySemantics::UniformMemory |
5672 SPIRV::MemorySemantics::ImageMemory |
5673 SPIRV::MemorySemantics::WorkgroupMemory,
5675 case Intrinsic::spv_all_memory_barrier_with_group_sync:
5676 return selectBarrierInst(
I, SPIRV::Scope::Device,
5677 SPIRV::MemorySemantics::UniformMemory |
5678 SPIRV::MemorySemantics::ImageMemory |
5679 SPIRV::MemorySemantics::WorkgroupMemory,
5681 case Intrinsic::spv_device_memory_barrier:
5682 return selectBarrierInst(
I, SPIRV::Scope::Device,
5683 SPIRV::MemorySemantics::UniformMemory |
5684 SPIRV::MemorySemantics::ImageMemory,
5686 case Intrinsic::spv_device_memory_barrier_with_group_sync:
5687 return selectBarrierInst(
I, SPIRV::Scope::Device,
5688 SPIRV::MemorySemantics::UniformMemory |
5689 SPIRV::MemorySemantics::ImageMemory,
5691 case Intrinsic::spv_group_memory_barrier:
5692 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5693 SPIRV::MemorySemantics::WorkgroupMemory,
5695 case Intrinsic::spv_group_memory_barrier_with_group_sync:
5696 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5697 SPIRV::MemorySemantics::WorkgroupMemory,
5699 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
5700 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
5701 SPIRV::StorageClass::StorageClass ResSC =
5704 return diagnoseUnsupported(
I,
"The target storage class is not castable "
5705 "from the Generic storage class");
5706 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
5714 case Intrinsic::spv_lifetime_start:
5715 case Intrinsic::spv_lifetime_end: {
5716 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
5717 : SPIRV::OpLifetimeStop;
5718 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
5719 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
5728 case Intrinsic::spv_saturate:
5729 return selectSaturate(ResVReg, ResType,
I);
5730 case Intrinsic::spv_nclamp:
5731 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
5732 case Intrinsic::spv_uclamp:
5733 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
5734 case Intrinsic::spv_sclamp:
5735 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
5736 case Intrinsic::spv_subgroup_prefix_bit_count:
5737 return selectWavePrefixBitCount(ResVReg, ResType,
I);
5738 case Intrinsic::spv_wave_active_countbits:
5739 return selectWaveActiveCountBits(ResVReg, ResType,
I);
5740 case Intrinsic::spv_wave_all_equal:
5741 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
5742 case Intrinsic::spv_wave_all:
5743 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
5744 case Intrinsic::spv_wave_any:
5745 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
5746 case Intrinsic::spv_subgroup_ballot:
5747 return selectWaveOpInst(ResVReg, ResType,
I,
5748 SPIRV::OpGroupNonUniformBallot);
5749 case Intrinsic::spv_wave_is_first_lane:
5750 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
5751 case Intrinsic::spv_wave_reduce_or:
5752 return selectWaveReduceOp(ResVReg, ResType,
I,
5753 SPIRV::OpGroupNonUniformBitwiseOr);
5754 case Intrinsic::spv_wave_reduce_xor:
5755 return selectWaveReduceOp(ResVReg, ResType,
I,
5756 SPIRV::OpGroupNonUniformBitwiseXor);
5757 case Intrinsic::spv_wave_reduce_and:
5758 return selectWaveReduceOp(ResVReg, ResType,
I,
5759 SPIRV::OpGroupNonUniformBitwiseAnd);
5760 case Intrinsic::spv_wave_reduce_umax:
5761 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
5762 case Intrinsic::spv_wave_reduce_max:
5763 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
5764 case Intrinsic::spv_wave_reduce_umin:
5765 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
5766 case Intrinsic::spv_wave_reduce_min:
5767 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
5768 case Intrinsic::spv_wave_reduce_sum:
5769 return selectWaveReduceSum(ResVReg, ResType,
I);
5770 case Intrinsic::spv_wave_product:
5771 return selectWaveReduceProduct(ResVReg, ResType,
I);
5772 case Intrinsic::spv_wave_readlane:
5773 return selectWaveOpInst(ResVReg, ResType,
I,
5774 SPIRV::OpGroupNonUniformShuffle);
5775 case Intrinsic::spv_wave_prefix_sum:
5776 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
5777 case Intrinsic::spv_wave_prefix_product:
5778 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
5779 case Intrinsic::spv_quad_read_across_x: {
5780 return selectQuadSwap(ResVReg, ResType,
I, 0);
5782 case Intrinsic::spv_quad_read_across_y: {
5783 return selectQuadSwap(ResVReg, ResType,
I, 1);
5785 case Intrinsic::spv_quad_read_across_diagonal: {
5786 return selectQuadSwap(ResVReg, ResType,
I, 2);
5788 case Intrinsic::spv_radians:
5789 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
5793 case Intrinsic::instrprof_increment:
5794 case Intrinsic::instrprof_increment_step:
5795 case Intrinsic::instrprof_value_profile:
5798 case Intrinsic::spv_value_md:
5800 case Intrinsic::spv_resource_handlefrombinding: {
5801 return selectHandleFromBinding(ResVReg, ResType,
I);
5803 case Intrinsic::spv_resource_counterhandlefrombinding:
5804 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
5805 case Intrinsic::spv_resource_updatecounter:
5806 return selectUpdateCounter(ResVReg, ResType,
I);
5807 case Intrinsic::spv_resource_store_typedbuffer: {
5808 return selectImageWriteIntrinsic(
I);
5810 case Intrinsic::spv_resource_load_typedbuffer: {
5811 return selectReadImageIntrinsic(ResVReg, ResType,
I);
5813 case Intrinsic::spv_resource_load_level: {
5814 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
5816 case Intrinsic::spv_resource_getdimensions_x:
5817 case Intrinsic::spv_resource_getdimensions_xy:
5818 case Intrinsic::spv_resource_getdimensions_xyz: {
5819 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
5821 case Intrinsic::spv_resource_getdimensions_levels_x:
5822 case Intrinsic::spv_resource_getdimensions_levels_xy:
5823 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
5824 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
5826 case Intrinsic::spv_resource_getdimensions_ms_xy:
5827 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
5828 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
5830 case Intrinsic::spv_resource_calculate_lod:
5831 case Intrinsic::spv_resource_calculate_lod_unclamped:
5832 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
5833 case Intrinsic::spv_resource_sample:
5834 case Intrinsic::spv_resource_sample_clamp:
5835 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
5836 case Intrinsic::spv_resource_samplebias:
5837 case Intrinsic::spv_resource_samplebias_clamp:
5838 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
5839 case Intrinsic::spv_resource_samplegrad:
5840 case Intrinsic::spv_resource_samplegrad_clamp:
5841 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
5842 case Intrinsic::spv_resource_samplelevel:
5843 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
5844 case Intrinsic::spv_resource_samplecmp:
5845 case Intrinsic::spv_resource_samplecmp_clamp:
5846 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
5847 case Intrinsic::spv_resource_samplecmplevelzero:
5848 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
5849 case Intrinsic::spv_resource_gather:
5850 case Intrinsic::spv_resource_gather_cmp:
5851 return selectGatherIntrinsic(ResVReg, ResType,
I);
5852 case Intrinsic::spv_resource_getbasepointer:
5853 case Intrinsic::spv_resource_getpointer: {
5854 return selectResourceGetPointer(ResVReg, ResType,
I);
5856 case Intrinsic::spv_pushconstant_getpointer: {
5857 return selectPushConstantGetPointer(ResVReg, ResType,
I);
5859 case Intrinsic::spv_discard: {
5860 return selectDiscard(ResVReg, ResType,
I);
5862 case Intrinsic::spv_resource_nonuniformindex: {
5863 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
5865 case Intrinsic::spv_unpackhalf2x16: {
5866 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
5868 case Intrinsic::spv_packhalf2x16: {
5869 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
5871 case Intrinsic::spv_ddx:
5872 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
5873 case Intrinsic::spv_ddy:
5874 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
5875 case Intrinsic::spv_ddx_coarse:
5876 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
5877 case Intrinsic::spv_ddy_coarse:
5878 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
5879 case Intrinsic::spv_ddx_fine:
5880 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
5881 case Intrinsic::spv_ddy_fine:
5882 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
5883 case Intrinsic::spv_fwidth:
5884 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
5885 case Intrinsic::spv_masked_gather:
5886 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5887 return selectMaskedGather(ResVReg, ResType,
I);
5888 return diagnoseUnsupported(
5889 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
5890 case Intrinsic::spv_masked_scatter:
5891 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5892 return selectMaskedScatter(
I);
5893 return diagnoseUnsupported(
5894 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
5895 case Intrinsic::returnaddress:
5896 case Intrinsic::frameaddress: {
5898 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5905 return diagnoseUnsupported(
I,
"intrinsic selection not implemented.");
5910bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
5911 SPIRVTypeInst ResType,
5912 MachineInstr &
I)
const {
5915 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
5922bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
5923 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5925 assert(Intr.getIntrinsicID() ==
5926 Intrinsic::spv_resource_counterhandlefrombinding);
5929 Register MainHandleReg = Intr.getOperand(2).getReg();
5931 assert(MainHandleDef->getIntrinsicID() ==
5932 Intrinsic::spv_resource_handlefrombinding);
5936 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
5937 Register IndexReg = MainHandleDef->getOperand(5).getReg();
5938 std::string CounterName =
5943 MachineIRBuilder MIRBuilder(
I);
5945 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
5947 ArraySize, IndexReg, CounterName, MIRBuilder);
5949 return BuildCOPY(ResVReg, CounterVarReg,
I);
5952bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
5953 SPIRVTypeInst ResType,
5954 MachineInstr &
I)
const {
5956 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
5958 Register CounterHandleReg = Intr.getOperand(2).getReg();
5959 Register IncrReg = Intr.getOperand(3).getReg();
5966 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
5967 assert(CounterVarPointeeType &&
5968 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
5969 "Counter variable must be a struct");
5971 SPIRV::StorageClass::StorageBuffer &&
5972 "Counter variable must be in the storage buffer storage class");
5974 "Counter variable must have exactly 1 member in the struct");
5975 const SPIRVTypeInst MemberType =
5978 "Counter variable struct must have a single i32 member");
5982 MachineIRBuilder MIRBuilder(
I);
5984 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
5987 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
5993 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5996 .
addUse(CounterHandleReg)
6003 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
6006 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
6009 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
6018 return BuildCOPY(ResVReg, AtomicRes,
I);
6026 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
6034bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
6035 SPIRVTypeInst ResType,
6036 MachineInstr &
I)
const {
6044 Register ImageReg =
I.getOperand(2).getReg();
6052 Register IdxReg =
I.getOperand(3).getReg();
6054 MachineInstr &Pos =
I;
6056 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
6060bool SPIRVInstructionSelector::generateSampleImage(
6063 DebugLoc Loc, MachineInstr &Pos)
const {
6074 if (!loadHandleBeforePosition(NewSamplerReg,
6080 MachineIRBuilder MIRBuilder(Pos);
6093 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
6094 ImOps.Lod.has_value();
6095 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
6096 : SPIRV::OpImageSampleImplicitLod;
6098 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
6099 : SPIRV::OpImageSampleDrefImplicitLod;
6108 MIB.
addUse(*ImOps.Compare);
6110 uint32_t ImageOperands = 0;
6112 ImageOperands |= SPIRV::ImageOperand::Bias;
6114 ImageOperands |= SPIRV::ImageOperand::Lod;
6115 if (ImOps.GradX && ImOps.GradY)
6116 ImageOperands |= SPIRV::ImageOperand::Grad;
6117 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
6119 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6122 "Non-constant offsets are not supported in sample instructions.");
6127 ImageOperands |= SPIRV::ImageOperand::MinLod;
6129 if (ImageOperands != 0) {
6130 MIB.
addImm(ImageOperands);
6131 if (ImageOperands & SPIRV::ImageOperand::Bias)
6133 if (ImageOperands & SPIRV::ImageOperand::Lod)
6135 if (ImageOperands & SPIRV::ImageOperand::Grad) {
6136 MIB.
addUse(*ImOps.GradX);
6137 MIB.
addUse(*ImOps.GradY);
6140 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6141 MIB.
addUse(*ImOps.Offset);
6142 if (ImageOperands & SPIRV::ImageOperand::MinLod)
6143 MIB.
addUse(*ImOps.MinLod);
6150bool SPIRVInstructionSelector::selectImageQuerySize(
6152 std::optional<Register> LodReg)
const {
6154 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
6157 "ImageReg is not an image type.");
6159 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6161 unsigned NumComponents = 0;
6163 case SPIRV::Dim::DIM_1D:
6164 case SPIRV::Dim::DIM_Buffer:
6165 NumComponents =
IsArray ? 2 : 1;
6167 case SPIRV::Dim::DIM_2D:
6168 case SPIRV::Dim::DIM_Cube:
6169 case SPIRV::Dim::DIM_Rect:
6170 NumComponents =
IsArray ? 3 : 2;
6172 case SPIRV::Dim::DIM_3D:
6176 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
6181 SPIRVTypeInst ResType =
6186 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6196bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
6197 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6198 Register ImageReg =
I.getOperand(2).getReg();
6205 return selectImageQuerySize(NewImageReg, ResVReg,
I);
6208bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
6209 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6210 Register ImageReg =
I.getOperand(2).getReg();
6219 Register LodReg =
I.getOperand(3).getReg();
6222 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
6224 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
6231 TII.get(SPIRV::OpImageQueryLevels))
6238 TII.get(SPIRV::OpCompositeConstruct))
6248bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
6249 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6250 Register ImageReg =
I.getOperand(2).getReg();
6261 "OpImageQuerySamples requires a multisampled image");
6263 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
6271 TII.get(SPIRV::OpImageQuerySamples))
6278 TII.get(SPIRV::OpCompositeConstruct))
6288bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
6289 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6290 Register ImageReg =
I.getOperand(2).getReg();
6291 Register SamplerReg =
I.getOperand(3).getReg();
6292 Register CoordinateReg =
I.getOperand(4).getReg();
6308 if (!loadHandleBeforePosition(
6313 MachineIRBuilder MIRBuilder(
I);
6319 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6329 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
6336 unsigned ExtractedIndex =
6338 Intrinsic::spv_resource_calculate_lod_unclamped
6342 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6343 TII.get(SPIRV::OpCompositeExtract))
6353bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
6354 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6355 Register ImageReg =
I.getOperand(2).getReg();
6356 Register SamplerReg =
I.getOperand(3).getReg();
6357 Register CoordinateReg =
I.getOperand(4).getReg();
6358 ImageOperands ImOps;
6359 if (
I.getNumOperands() > 5)
6360 ImOps.Offset =
I.getOperand(5).getReg();
6361 if (
I.getNumOperands() > 6)
6362 ImOps.MinLod =
I.getOperand(6).getReg();
6363 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6364 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6367bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
6368 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6369 Register ImageReg =
I.getOperand(2).getReg();
6370 Register SamplerReg =
I.getOperand(3).getReg();
6371 Register CoordinateReg =
I.getOperand(4).getReg();
6372 ImageOperands ImOps;
6373 ImOps.Bias =
I.getOperand(5).getReg();
6374 if (
I.getNumOperands() > 6)
6375 ImOps.Offset =
I.getOperand(6).getReg();
6376 if (
I.getNumOperands() > 7)
6377 ImOps.MinLod =
I.getOperand(7).getReg();
6378 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6379 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6382bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
6383 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6384 Register ImageReg =
I.getOperand(2).getReg();
6385 Register SamplerReg =
I.getOperand(3).getReg();
6386 Register CoordinateReg =
I.getOperand(4).getReg();
6387 ImageOperands ImOps;
6388 ImOps.GradX =
I.getOperand(5).getReg();
6389 ImOps.GradY =
I.getOperand(6).getReg();
6390 if (
I.getNumOperands() > 7)
6391 ImOps.Offset =
I.getOperand(7).getReg();
6392 if (
I.getNumOperands() > 8)
6393 ImOps.MinLod =
I.getOperand(8).getReg();
6394 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6395 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6398bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
6399 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6400 Register ImageReg =
I.getOperand(2).getReg();
6401 Register SamplerReg =
I.getOperand(3).getReg();
6402 Register CoordinateReg =
I.getOperand(4).getReg();
6403 ImageOperands ImOps;
6404 ImOps.Lod =
I.getOperand(5).getReg();
6405 if (
I.getNumOperands() > 6)
6406 ImOps.Offset =
I.getOperand(6).getReg();
6407 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6408 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6411bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
6412 SPIRVTypeInst ResType,
6413 MachineInstr &
I)
const {
6414 Register ImageReg =
I.getOperand(2).getReg();
6415 Register SamplerReg =
I.getOperand(3).getReg();
6416 Register CoordinateReg =
I.getOperand(4).getReg();
6417 ImageOperands ImOps;
6418 ImOps.Compare =
I.getOperand(5).getReg();
6419 if (
I.getNumOperands() > 6)
6420 ImOps.Offset =
I.getOperand(6).getReg();
6421 if (
I.getNumOperands() > 7)
6422 ImOps.MinLod =
I.getOperand(7).getReg();
6423 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6424 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6427bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
6428 SPIRVTypeInst ResType,
6429 MachineInstr &
I)
const {
6430 Register ImageReg =
I.getOperand(2).getReg();
6431 Register CoordinateReg =
I.getOperand(3).getReg();
6432 Register LodReg =
I.getOperand(4).getReg();
6434 ImageOperands ImOps;
6436 if (
I.getNumOperands() > 5)
6437 ImOps.Offset =
I.getOperand(5).getReg();
6449 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
6450 I.getDebugLoc(),
I, &ImOps);
6453bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
6454 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6455 Register ImageReg =
I.getOperand(2).getReg();
6456 Register SamplerReg =
I.getOperand(3).getReg();
6457 Register CoordinateReg =
I.getOperand(4).getReg();
6458 ImageOperands ImOps;
6459 ImOps.Compare =
I.getOperand(5).getReg();
6460 if (
I.getNumOperands() > 6)
6461 ImOps.Offset =
I.getOperand(6).getReg();
6464 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6465 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6468bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
6469 SPIRVTypeInst ResType,
6470 MachineInstr &
I)
const {
6471 Register ImageReg =
I.getOperand(2).getReg();
6472 Register SamplerReg =
I.getOperand(3).getReg();
6473 Register CoordinateReg =
I.getOperand(4).getReg();
6476 "ImageReg is not an image type.");
6481 ComponentOrCompareReg =
I.getOperand(5).getReg();
6482 OffsetReg =
I.getOperand(6).getReg();
6485 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
6489 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6490 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
6491 Dim != SPIRV::Dim::DIM_Rect) {
6493 "Gather operations are only supported for 2D, Cube, and Rect images.");
6500 if (!loadHandleBeforePosition(
6505 MachineIRBuilder MIRBuilder(
I);
6506 SPIRVTypeInst SampledImageType =
6511 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6519 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
6521 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
6523 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6528 .
addUse(ComponentOrCompareReg);
6530 uint32_t ImageOperands = 0;
6531 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
6532 if (Dim == SPIRV::Dim::DIM_Cube) {
6534 "Gather operations with offset are not supported for Cube images.");
6538 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6540 ImageOperands |= SPIRV::ImageOperand::Offset;
6544 if (ImageOperands != 0) {
6545 MIB.
addImm(ImageOperands);
6547 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6555bool SPIRVInstructionSelector::generateImageReadOrFetch(
6558 const ImageOperands *ImOps)
const {
6561 "ImageReg is not an image type.");
6563 bool IsSignedInteger =
6568 bool IsFetch = (SampledOp.getImm() == 1);
6570 auto AddOperands = [&](MachineInstrBuilder &MIB) {
6571 uint32_t ImageOperandsMask = 0;
6572 if (IsSignedInteger)
6573 ImageOperandsMask |= 0x1000;
6575 if (IsFetch && ImOps) {
6577 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
6578 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
6580 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
6582 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
6586 if (ImageOperandsMask != 0) {
6587 MIB.
addImm(ImageOperandsMask);
6588 if (IsFetch && ImOps) {
6591 if (ImOps->Offset &&
6592 (ImageOperandsMask &
6593 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
6594 MIB.
addUse(*ImOps->Offset);
6603 SPIRVTypeInst SampledType =
6606 SPIRVTypeInst ReadType =
6607 widenTypeToVec4(IsPacked ? SampledType : ResType, Pos);
6608 bool ReadTypeMatchesResult = ReadType == ResType;
6610 Register ReadReg = ReadTypeMatchesResult
6616 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
6622 BMI.constrainAllUses(
TII,
TRI, RBI);
6624 if (ReadTypeMatchesResult)
6637 if (ResultSize == 1) {
6646 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
6649bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
6650 SPIRVTypeInst ResType,
6651 MachineInstr &
I)
const {
6652 Register ResourcePtr =
I.getOperand(2).getReg();
6654 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
6663 MachineIRBuilder MIRBuilder(
I);
6668 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6674 if (
I.getNumExplicitOperands() > 3) {
6675 Register IndexReg =
I.getOperand(3).getReg();
6682bool SPIRVInstructionSelector::selectPushConstantGetPointer(
6683 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6688bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
6689 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6690 Register ObjReg =
I.getOperand(2).getReg();
6691 if (!BuildCOPY(ResVReg, ObjReg,
I))
6701 decorateUsesAsNonUniform(ResVReg);
6705void SPIRVInstructionSelector::decorateUsesAsNonUniform(
6708 {NonUniformReg,
nullptr}};
6709 llvm::SmallSet<Register, 8> Visited;
6710 while (WorkList.
size() > 0) {
6713 if (!Visited.
insert(CurrentReg).second)
6716 bool IsDecorated =
false;
6718 if (
Use.getOpcode() == SPIRV::OpDecorate &&
6719 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
6725 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
6727 if (ResultReg == CurrentReg)
6735 MachineInstr &InsertPt =
6738 SPIRV::Decoration::NonUniformEXT, {});
6743bool SPIRVInstructionSelector::extractSubvector(
6745 MachineInstr &InsertionPoint)
const {
6747 [[maybe_unused]]
uint64_t InputSize =
6750 [[maybe_unused]]
bool IsLongVectorEXT =
6752 assert((InputSize > 1 || IsLongVectorEXT) &&
"The input must be a vector.");
6753 assert((ResultSize > 1 || IsLongVectorEXT) &&
"The result must be a vector.");
6754 assert(ResultSize < InputSize &&
6755 "Cannot extract more element than there are in the input.");
6762 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6771 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
6773 TII.get(SPIRV::OpCompositeConstruct))
6777 for (
Register ComponentReg : ComponentRegisters)
6778 MIB.
addUse(ComponentReg);
6783bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
6784 MachineInstr &
I)
const {
6791 Register ImageReg =
I.getOperand(1).getReg();
6799 Register CoordinateReg =
I.getOperand(2).getReg();
6800 Register DataReg =
I.getOperand(3).getReg();
6803 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
6811Register SPIRVInstructionSelector::buildPointerToResource(
6812 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
6813 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
6814 StringRef Name, MachineIRBuilder MIRBuilder)
const {
6816 if (ArraySize == 1) {
6817 SPIRVTypeInst PtrType =
6820 "SpirvResType did not have an explicit layout.");
6825 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
6826 SPIRVTypeInst VarPointerType =
6829 VarPointerType, Set,
Binding, Name, MIRBuilder);
6831 SPIRVTypeInst ResPointerType =
6844bool SPIRVInstructionSelector::selectFirstBitSet16(
6845 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
6846 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
6848 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
6852 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
6855bool SPIRVInstructionSelector::selectFirstBitSet32(
6857 unsigned BitSetOpcode)
const {
6858 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6861 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6868bool SPIRVInstructionSelector::selectFirstBitSet64(
6870 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
6884 if (ComponentCount > 2) {
6885 auto Func = [
this, SwapPrimarySide](
Register ResVReg, SPIRVTypeInst ResType,
6887 unsigned Opcode) ->
bool {
6888 return this->selectFirstBitSet64(ResVReg, ResType,
I, SrcReg, Opcode,
6892 return handle64BitOverflow(ResVReg, ResType,
I, SrcReg, BitSetOpcode, Func);
6896 MachineIRBuilder MIRBuilder(
I);
6898 BaseType, 2 * ComponentCount, MIRBuilder,
false);
6902 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
6908 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
6918 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
6919 SPIRV::OpVectorExtractDynamic))
6921 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
6922 SPIRV::OpVectorExtractDynamic))
6926 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6927 TII.get(SPIRV::OpVectorShuffle))
6935 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
6941 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6942 TII.get(SPIRV::OpVectorShuffle))
6950 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
6970 SelectOp = SPIRV::OpSelectSISCond;
6971 AddOp = SPIRV::OpIAddS;
6979 SelectOp = SPIRV::OpSelectVIVCond;
6980 AddOp = SPIRV::OpIAddV;
6986 Register RegSecondaryOffset = Reg0;
6990 if (SwapPrimarySide) {
6991 PrimaryReg = LowReg;
6992 SecondaryReg = HighReg;
6993 RegPrimaryOffset = Reg0;
6994 RegSecondaryOffset = Reg32;
6999 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
7000 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
7005 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
7006 SPIRV::OpINotEqual))
7013 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
7014 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
7019 if (SwapPrimarySide) {
7021 if (!selectOpWithSrcs(RegAdd, ResType,
I,
7022 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
7033 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
7034 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
7039 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
7040 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
7043 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
7047bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
7048 SPIRVTypeInst ResType,
7050 bool IsSigned)
const {
7052 Register OpReg =
I.getOperand(2).getReg();
7055 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
7056 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
7060 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
7062 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
7064 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
7067 return diagnoseUnsupported(
7069 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
7073bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
7074 SPIRVTypeInst ResType,
7075 MachineInstr &
I)
const {
7077 Register OpReg =
I.getOperand(2).getReg();
7082 unsigned ExtendOpcode = SPIRV::OpUConvert;
7083 unsigned BitSetOpcode = GL::FindILsb;
7087 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
7089 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
7091 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
7094 return diagnoseUnsupported(
I,
7095 "spv_firstbitlow only supports 16,32,64 bits.");
7099bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
7100 SPIRVTypeInst ResType,
7101 MachineInstr &
I)
const {
7105 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
7108 .
addUse(
I.getOperand(2).getReg())
7111 unsigned Alignment =
I.getOperand(3).getImm();
7125 while (!Worklist.
empty()) {
7127 switch (
T->getOpcode()) {
7128 case SPIRV::OpTypeInt:
7129 case SPIRV::OpTypeFloat:
7130 case SPIRV::OpTypePointer:
7132 case SPIRV::OpTypeVector:
7133 case SPIRV::OpTypeVectorIdEXT:
7134 case SPIRV::OpTypeMatrix:
7135 case SPIRV::OpTypeArray: {
7136 Register OperandReg =
T->getOperand(1).getReg();
7140 case SPIRV::OpTypeStruct:
7141 for (
unsigned Idx = 1,
E =
T->getNumOperands(); Idx <
E; ++Idx) {
7142 Register OperandReg =
T->getOperand(Idx).getReg();
7154bool SPIRVInstructionSelector::selectAbort(MachineInstr &
I)
const {
7155 assert(
I.getNumExplicitOperands() == 2);
7157 Register MsgReg =
I.getOperand(1).getReg();
7159 assert(MsgType &&
"Message argument of llvm.spv.abort has no SPIR-V type");
7162 return diagnoseUnsupported(
7164 "llvm.spv.abort message type must be a concrete SPIR-V type (numerical "
7165 "scalar, pointer, vector, matrix, or aggregate of such types)");
7168 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7175bool SPIRVInstructionSelector::selectTrap(MachineInstr &
I)
const {
7184 uint32_t MsgVal = ~0
u;
7185 if (
I.getOpcode() == TargetOpcode::G_UBSANTRAP)
7186 MsgVal =
static_cast<uint32_t
>(
I.getOperand(0).
getImm());
7189 Register MsgReg = buildI32ConstantInEntryBlock(MsgVal,
I, MsgType);
7192 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7199bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
7200 SPIRVTypeInst ResType,
7201 MachineInstr &
I)
const {
7208 bool UseUntypedPointers =
7209 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7211 UseUntypedPointers ? SPIRV::OpUntypedVariableKHR : SPIRV::OpVariable;
7213 auto MIB =
BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(Opcode))
7216 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7220 if (UseUntypedPointers) {
7224 return diagnoseUnsupported(
7225 I,
"could not deduce the data type of an untyped variable");
7231 unsigned Alignment =
I.getOperand(2).getImm();
7238bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
7243 const MachineInstr *PrevI =
I.getPrevNode();
7245 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
7249 .
addMBB(
I.getOperand(0).getMBB())
7254 .
addMBB(
I.getOperand(0).getMBB())
7259bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
7270 const MachineInstr *NextI =
I.getNextNode();
7272 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
7278 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
7280 .
addUse(
I.getOperand(0).getReg())
7281 .
addMBB(
I.getOperand(1).getMBB())
7287bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
7288 MachineInstr &
I)
const {
7290 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
7292 const unsigned NumOps =
I.getNumOperands();
7293 for (
unsigned i = 1; i <
NumOps; i += 2) {
7294 MIB.
addUse(
I.getOperand(i + 0).getReg());
7295 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
7301bool SPIRVInstructionSelector::selectGlobalValue(
7302 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
7304 MachineIRBuilder MIRBuilder(
I);
7305 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
7308 std::string GlobalIdent;
7310 unsigned &
ID = UnnamedGlobalIDs[GV];
7312 ID = UnnamedGlobalIDs.
size();
7313 GlobalIdent =
"__unnamed_" + Twine(ID).str();
7339 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
7346 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7351 MachineInstrBuilder MIB1 =
7352 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7355 MachineInstrBuilder MIB2 =
7357 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
7361 GR.
add(ConstVal, MIB2);
7369 MachineInstrBuilder MIB3 =
7370 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7373 GR.
add(ConstVal, MIB3);
7379 assert(NewReg != ResVReg);
7380 return BuildCOPY(ResVReg, NewReg,
I);
7390 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
7393 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
7399 SPIRVTypeInst ResType =
7403 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
7408 if (
GlobalVar->isExternallyInitialized() &&
7409 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
7410 constexpr unsigned ReadWriteINTEL = 3u;
7413 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
7419bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
7420 SPIRVTypeInst ResType,
7421 MachineInstr &
I)
const {
7423 return selectExtInst(ResVReg, ResType,
I, CL::log10);
7431 MachineIRBuilder MIRBuilder(
I);
7436 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7439 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
7441 .
add(
I.getOperand(1))
7455 APFloat::rmNearestTiesToEven, &LosesInfo);
7460 isVectorType(ResType) ? SPIRV::OpVectorTimesScalar : SPIRV::OpFMulS;
7470bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
7471 SPIRVTypeInst ResType,
7472 MachineInstr &
I)
const {
7475 return selectExtInst(ResVReg, ResType,
I, CL::pown);
7481 Register ExpReg =
I.getOperand(2).getReg();
7483 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
7484 SPIRV::OpConvertSToF))
7486 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
7493bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
7494 SPIRVTypeInst ResType,
7495 MachineInstr &
I)
const {
7511 MachineIRBuilder MIRBuilder(
I);
7512 SPIRVTypeInst FloatType =
7516 FloatType, MIRBuilder, SPIRV::StorageClass::Function);
7529 MachineBasicBlock &EntryBB =
I.getMF()->
front();
7530 const bool IsUntyped =
7531 PtrType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7533 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
7534 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
7535 : SPIRV::OpVariable))
7538 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7546 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7549 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
7552 .
add(
I.getOperand(
I.getNumExplicitDefs()))
7556 Register IntegralPartReg =
I.getOperand(1).getReg();
7559 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7569 assert(
false &&
"GLSL::Modf is deprecated.");
7580bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
7581 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7582 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7583 MachineIRBuilder MIRBuilder(
I);
7584 const SPIRVTypeInst Vec3Ty =
7587 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
7599 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7603 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7609 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7616 assert(
I.getOperand(2).isReg());
7617 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
7621 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
7632bool SPIRVInstructionSelector::loadBuiltinInputID(
7633 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7634 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7635 MachineIRBuilder MIRBuilder(
I);
7637 ResType, MIRBuilder, SPIRV::StorageClass::Input);
7652 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7656 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7665SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
7666 MachineInstr &
I)
const {
7667 MachineIRBuilder MIRBuilder(
I);
7678bool SPIRVInstructionSelector::loadHandleBeforePosition(
7679 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
7680 MachineInstr &Pos)
const {
7683 Intrinsic::spv_resource_handlefrombinding);
7691 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
7692 MachineIRBuilder MIRBuilder(HandleDef);
7693 SPIRVTypeInst VarType = ResType;
7694 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
7696 if (IsStructuredBuffer) {
7701 if (ResType->
getOpcode() == SPIRV::OpTypeImage && ArraySize == 0)
7703 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
7706 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
7707 ArraySize, IndexReg, Name, MIRBuilder);
7711 uint32_t LoadOpcode =
7712 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
7722bool SPIRVInstructionSelector::errorIfInstrOutsideShader(
7723 MachineInstr &
I)
const {
7725 return diagnoseUnsupported(
7726 I,
"this instruction is only supported in shaders.");
7731InstructionSelector *
7735 return new SPIRVInstructionSelector(TM, Subtarget, RBI);
MachineInstrBuilder & UseMI
MachineInstrBuilder MachineInstrBuilder & DefMI
#define GET_GLOBALISEL_PREDICATES_INIT
#define GET_GLOBALISEL_TEMPORARIES_INIT
assert(UImm &&(UImm !=~static_cast< T >(0)) &&"Invalid immediate!")
This file declares a class to represent arbitrary precision floating point values and provide a varie...
static bool selectUnmergeValues(MachineInstrBuilder &MIB, const ARMBaseInstrInfo &TII, MachineRegisterInfo &MRI, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI)
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
static uint8_t SwapBits(uint8_t Val)
static GCRegistry::Add< CoreCLRGC > E("coreclr", "CoreCLR-compatible GC")
DXIL Resource Implicit Binding
Declares convenience wrapper classes for interpreting MachineInstr instances as specific generic oper...
const HexagonInstrInfo * TII
LLVMTypeRef LLVMIntType(unsigned NumBits)
const size_t AbstractManglingParser< Derived, Alloc >::NumOps
Loop::LoopBounds::Direction Direction
Register const TargetRegisterInfo * TRI
Promote Memory to Register
uint64_t IntrinsicInst * II
static StringRef getName(Value *V)
static unsigned getFCmpOpcode(CmpInst::Predicate Pred, unsigned Size)
static bool isConcreteSPIRVType(SPIRVTypeInst Ty, const SPIRVGlobalRegistry &GR)
static APFloat getOneFP(const Type *LLVMFloatTy)
static bool isUSMStorageClass(SPIRV::StorageClass::StorageClass SC)
static bool isASCastInGVar(MachineRegisterInfo *MRI, Register ResVReg)
static bool mayApplyGenericSelection(unsigned Opcode)
static APFloat getZeroFP(const Type *LLVMFloatTy)
std::vector< std::pair< SPIRV::InstructionSet::InstructionSet, uint32_t > > ExtInstList
static bool intrinsicHasSideEffects(Intrinsic::ID ID)
static unsigned getBoolCmpOpcode(unsigned PredNum)
static unsigned getICmpOpcode(unsigned PredNum)
static bool isOpcodeWithNoSideEffects(unsigned Opcode)
static void addMemoryOperands(MachineMemOperand *MemOp, MachineInstrBuilder &MIB, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry &GR)
static bool isConstReg(MachineRegisterInfo *MRI, MachineInstr *OpDef)
static unsigned getPtrCmpOpcode(unsigned Pred)
bool isDead(const MachineInstr &MI, const MachineRegisterInfo &MRI)
static Register convertPtrToInt(Register Reg, LLT ConvTy, SPIRVTypeInst SpvType, LegalizerHelper &Helper, MachineRegisterInfo &MRI, SPIRVGlobalRegistry *GR)
BaseType
A given derived pointer can have multiple base pointers through phi/selects.
This file defines the SmallSet class.
static TableGen::Emitter::Opt Y("gen-skeleton-entry", EmitSkeleton, "Generate example skeleton entry")
static ManagedStatic< cl::opt< FnT >, OptCreatorT > CallbackFunction
static const fltSemantics & IEEEsingle()
static const fltSemantics & BFloat()
static const fltSemantics & IEEEdouble()
static const fltSemantics & IEEEhalf()
const fltSemantics & getSemantics() const
static APFloat getOne(const fltSemantics &Sem, bool Negative=false)
Factory for Positive and Negative One.
static APFloat getZero(const fltSemantics &Sem, bool Negative=false)
Factory for Positive and Negative Zero.
static APInt getAllOnes(unsigned numBits)
Return an APInt of a specified width with all bits set.
static APInt getSignMask(unsigned BitWidth)
Get the SignMask for a specific bit width.
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.
static LLVM_ABI std::optional< GFConstant > getConstant(Register Const, const MachineRegisterInfo &MRI)
static LLVM_ABI std::optional< GIConstant > getConstant(Register Const, const MachineRegisterInfo &MRI)
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.
instr_iterator instr_end()
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 substituteRegister(Register FromReg, Register ToReg, unsigned SubIdx, const TargetRegisterInfo &RegInfo)
Replace all occurrences of FromReg with ToReg:SubIdx, properly composing subreg indices where necessa...
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 LLVM_READONLY MachineInstr * getVRegDef(Register Reg) const
getVRegDef - Return the machine instr that defines the specified virtual register or null if none is ...
use_instr_iterator use_instr_begin(Register RegNo) const
bool use_nodbg_empty(Register RegNo) const
use_nodbg_empty - Return true if there are no non-Debug instructions using the specified register.
static def_instr_iterator def_instr_end()
defusechain_instr_iterator< false, true, false, true > def_instr_iterator
def_instr_iterator/def_instr_begin/def_instr_end - Walk all defs of the specified register,...
LLVM_ABI Register createVirtualRegister(const TargetRegisterClass *RegClass, StringRef Name="")
createVirtualRegister - Create and return a new virtual register in the function with the specified r...
def_instr_iterator def_instr_begin(Register RegNo) const
LLT getType(Register Reg) const
Get the low-level type of Reg or LLT{} if Reg is not a generic (target independent) virtual register.
bool hasOneUse(Register RegNo) const
hasOneUse - Return true if there is exactly one instruction using the specified register.
static use_instr_iterator use_instr_end()
iterator_range< use_instr_nodbg_iterator > use_nodbg_instructions(Register Reg) const
LLVM_ABI void setType(Register VReg, LLT Ty)
Set the low-level type of VReg to Ty.
const MachineFunction & getMF() const
LLVM_ABI void setRegClass(Register Reg, const TargetRegisterClass *RC)
setRegClass - Set the register class of the specified virtual register.
LLVM_ABI Register createGenericVirtualRegister(LLT Ty, StringRef Name="")
Create and return a new generic virtual register with low-level type Ty.
const TargetRegisterClass * getRegClassOrNull(Register Reg) const
Return the register class of Reg, or null if Reg has not been assigned a register class yet.
iterator_range< use_instr_iterator > use_instructions(Register Reg) const
unsigned getNumVirtRegs() const
getNumVirtRegs - Return the number of virtual registers created.
LLVM_ABI void replaceRegWith(Register FromReg, Register ToReg)
replaceRegWith - Replace all instances of FromReg with ToReg in the machine function.
Analysis providing profile information.
Holds all the information related to register banks.
Wrapper class representing virtual and physical registers.
constexpr bool isValid() const
constexpr bool isPhysical() const
Return true if the specified register number is in the physical register namespace.
bool isScalarOrVectorSigned(SPIRVTypeInst Type) const
SPIRVTypeInst getOrCreateOpTypeSampledImage(SPIRVTypeInst ImageType, MachineIRBuilder &MIRBuilder)
void assignSPIRVTypeToVReg(SPIRVTypeInst Type, Register VReg, const MachineFunction &MF)
SPIRVTypeInst getOrCreateSPIRVPointerType(const Type *BaseType, MachineIRBuilder &MIRBuilder, SPIRV::StorageClass::StorageClass SC, bool ForceTyped=false)
const TargetRegisterClass * getRegClass(SPIRVTypeInst SpvType) const
MachineInstr * getOrAddMemAliasingINTELInst(MachineIRBuilder &MIRBuilder, const MDNode *AliasingListMD)
bool isAggregateType(SPIRVTypeInst Type) const
unsigned getScalarOrVectorBitWidth(SPIRVTypeInst Type) const
SPIRVTypeInst getOrCreateSPIRVIntegerType(unsigned BitWidth, MachineIRBuilder &MIRBuilder)
SPIRVTypeInst getOrCreateSPIRVVectorType(SPIRVTypeInst BaseType, unsigned NumElements, MachineIRBuilder &MIRBuilder, bool EmitIR)
Register buildGlobalVariable(Register Reg, SPIRVTypeInst BaseType, StringRef Name, const GlobalValue *GV, SPIRV::StorageClass::StorageClass Storage, const MachineInstr *Init, bool IsConst, const std::optional< SPIRV::LinkageType::LinkageType > &LinkageType, MachineIRBuilder &MIRBuilder, bool IsInstSelector)
SPIRVTypeInst getResultType(Register VReg, MachineFunction *MF=nullptr)
unsigned getScalarOrVectorComponentCount(Register VReg) const
const Type * getTypeForSPIRVType(SPIRVTypeInst Ty) const
bool isBitcastCompatible(SPIRVTypeInst Type1, SPIRVTypeInst Type2) const
unsigned getPointerSize() const
Register getOrCreateConstFP(APFloat Val, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
LLT getRegType(SPIRVTypeInst SpvType) const
SPIRVTypeInst getOpTypeVoid(MachineIRBuilder &MIRBuilder)
void invalidateMachineInstr(MachineInstr *MI)
SPIRVTypeInst getOrCreateSPIRVBoolType(MachineIRBuilder &MIRBuilder, bool EmitIR)
bool isScalarOfType(Register VReg, unsigned TypeOpcode) const
Register getSPIRVTypeID(SPIRVTypeInst SpirvType) const
Register getOrCreateConstInt(uint64_t Val, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
Register getOrCreateConstIntArray(uint64_t Val, size_t Num, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII)
bool findValueAttrs(const MachineInstr *Key, Type *&Ty, StringRef &Name)
SPIRVTypeInst retrieveScalarOrVectorIntType(SPIRVTypeInst Type) const
Register getOrCreateGlobalVariableWithBinding(SPIRVTypeInst VarType, uint32_t Set, uint32_t Binding, StringRef Name, MachineIRBuilder &MIRBuilder)
SPIRVTypeInst changePointerStorageClass(SPIRVTypeInst PtrType, SPIRV::StorageClass::StorageClass SC, MachineInstr &I)
Register getOrCreateConstVector(uint64_t Val, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
Register buildConstantFP(APFloat Val, MachineIRBuilder &MIRBuilder, SPIRVTypeInst SpvType=nullptr)
void addGlobalObject(const Value *V, const MachineFunction *MF, Register R)
SPIRVTypeInst getScalarOrVectorComponentType(SPIRVTypeInst Type) const
void recordFunctionPointer(const MachineOperand *MO, const Function *F)
SPIRVTypeInst getOrCreateSPIRVFloatType(unsigned BitWidth, MachineInstr &I, const SPIRVInstrInfo &TII)
SPIRVTypeInst getPointeeType(SPIRVTypeInst PtrType)
SPIRVTypeInst getOrCreateSPIRVType(const Type *Type, MachineInstr &I, SPIRV::AccessQualifier::AccessQualifier AQ, bool EmitIR)
bool isScalarOrVectorOfType(Register VReg, unsigned TypeOpcode) const
MachineFunction * setCurrentFunc(MachineFunction &MF)
Register getOrCreateConstNullPtr(MachineIRBuilder &MIRBuilder, SPIRVTypeInst SpvType)
SPIRVTypeInst getSPIRVTypeForVReg(Register VReg, const MachineFunction *MF=nullptr) const
Type * getDeducedGlobalValueType(const GlobalValue *Global)
Register getOrCreateUndef(MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII)
SPIRV::StorageClass::StorageClass getPointerStorageClass(Register VReg) const
SPIRVTypeInst getUntypedPtrElementType(Register Reg) const
bool erase(const MachineInstr *MI)
bool add(SPIRV::IRHandle Handle, const MachineInstr *MI)
Register find(SPIRV::IRHandle Handle, const MachineFunction *MF)
bool isPhysicalSPIRV() const
bool isAtLeastSPIRVVer(VersionTuple VerToCompareTo) const
bool canUseExtInstSet(SPIRV::InstructionSet::InstructionSet E) const
bool isLogicalSPIRV() const
bool canUseExtension(SPIRV::Extension::Extension E) const
bool isTypeIntOrFloat() const
bool isAnyTypeFloat() const
bool erase(PtrType Ptr)
Remove pointer from the set.
std::pair< iterator, bool > insert(PtrType Ptr)
Inserts Ptr if and only if there is no element in the container equal to Ptr.
bool contains(ConstPtrType Ptr) const
SmallPtrSet - This class implements a set which is optimized for holding SmallSize or less elements.
std::pair< const_iterator, bool > insert(const T &V)
insert - Insert an element into the set if it isn't already there.
reference emplace_back(ArgTypes &&... Args)
void reserve(size_type N)
void push_back(const T &Elt)
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
constexpr size_t size() const
Get the string size.
static LLVM_ABI StructType * get(LLVMContext &Context, ArrayRef< Type * > Elements, bool isPacked=false)
This static method is the primary way to create a literal StructType.
Twine - A lightweight data structure for efficiently representing the concatenation of temporary valu...
The instances of the Type class are immutable: once they are created, they are never changed.
@ HalfTyID
16-bit floating point type
@ FloatTyID
32-bit floating point type
@ BFloatTyID
16-bit floating point type (7-bit significand)
@ DoubleTyID
64-bit floating point type
Type * getScalarType() const
If this is a vector type, return the element type, otherwise return 'this'.
bool isStructTy() const
True if this is an instance of StructType.
bool isAggregateType() const
Return true if the type is an aggregate type.
TypeID getTypeID() const
Return the type id for the type.
Value * getOperand(unsigned i) const
LLVM_ABI StringRef getName() const
Return a constant reference to the value's name.
An efficient, type-erasing, non-owning reference to a callable.
self_iterator getIterator()
NodeTy * getNextNode()
Get the next node, or nullptr for the list tail.
#define llvm_unreachable(msg)
Marks that the current location is not supposed to be reachable.
constexpr char IsConst[]
Key for Kernel::Arg::Metadata::mIsConst.
constexpr std::underlying_type_t< E > Mask()
Get a bitmask with 1s in all places up to the high-order bit of E's largest value.
constexpr uint64_t PointerSize
aarch64 pointer size.
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.
unsigned getOpcode(const VPValue *V)
Return the instruction opcode for the recipe defining V or 0 for unsupported recipes and VPValues not...
This is an optimization pass for GlobalISel generic memory operations.
@ Low
Lower the current thread's priority such that it does not affect foreground tasks significantly.
bool all_of(R &&range, UnaryPredicate P)
Provide wrappers to std::all_of which take ranges instead of having to pass begin/end explicitly.
void addStringImm(StringRef Str, MCInst &Inst)
MachineBasicBlock::iterator getOpVariableMBBIt(MachineFunction &MF)
int64_t getIConstValSext(Register ConstReg, const MachineRegisterInfo *MRI)
MachineInstrBuilder BuildMI(MachineFunction &MF, const MIMetadata &MIMD, const MCInstrDesc &MCID)
Builder interface. Specify how to create the initial instruction itself.
bool isTypeFoldingSupported(unsigned Opcode)
uint32_t getMemSemanticsWithStorageClass(const Triple &TT, uint32_t OrderSem, uint32_t StorageClassSem)
decltype(auto) dyn_cast(const From &Val)
dyn_cast<X> - Return the argument parameter cast to the specified type.
MachineInstr * getDef(const MachineOperand &MO, const MachineRegisterInfo *MRI)
void addNumImm(const APInt &Imm, MachineInstrBuilder &MIB)
LLVM_ABI void salvageDebugInfo(const MachineRegisterInfo &MRI, MachineInstr &MI)
Assuming the instruction MI is going to be deleted, attempt to salvage debug users of MI by writing t...
LLVM_ABI void constrainSelectedInstRegOperands(MachineInstr &I, const TargetInstrInfo &TII, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI)
Mutate the newly-selected instruction I to constrain its (possibly generic) virtual register operands...
bool isPreISelGenericOpcode(unsigned Opcode)
Check whether the given Opcode is a generic opcode that is not supposed to appear after ISel.
iterator_range< T > make_range(T x, T y)
Convenience function for iterating over sub-ranges.
Register createVirtualRegister(SPIRVTypeInst SpvType, SPIRVGlobalRegistry *GR, MachineRegisterInfo *MRI, const MachineFunction &MF)
unsigned getArrayComponentCount(const MachineRegisterInfo *MRI, const MachineInstr *ResType)
iterator_range< early_inc_iterator_impl< detail::IterOfRange< RangeT > > > make_early_inc_range(RangeT &&Range)
Make a range that does early increment to allow mutation of the underlying range without disrupting i...
void buildOpDecorate(Register Reg, MachineIRBuilder &MIRBuilder, SPIRV::Decoration::Decoration Dec, ArrayRef< uint32_t > DecArgs, StringRef StrImm)
LLVM_ABI bool isNullOrNullSplat(const MachineInstr &MI, const MachineRegisterInfo &MRI, bool AllowUndefs=false)
Return true if the value is a constant 0 integer or a splatted vector of a constant 0 integer (with n...
SPIRV::Scope::Scope getMemScope(const Triple &TT, LLVMContext &Ctx, SyncScope::ID Id)
uint64_t getIConstVal(Register ConstReg, const MachineRegisterInfo *MRI)
SmallVector< MachineInstr *, 4 > createContinuedInstructions(MachineIRBuilder &MIRBuilder, unsigned Opcode, unsigned MinWC, unsigned ContinuedOpcode, ArrayRef< Register > Args, Register ReturnRegister, Register TypeID)
SPIRV::MemorySemantics::MemorySemantics getMemSemanticsForStorageClass(SPIRV::StorageClass::StorageClass SC)
constexpr unsigned storageClassToAddressSpace(SPIRV::StorageClass::StorageClass SC)
bool isVectorType(SPIRVTypeInst SPVTy)
RelativeUniformCounterPtr ValuesPtrExpr VTableAddr Value
void buildOpName(Register Target, StringRef Name, MachineIRBuilder &MIRBuilder)
bool any_of(R &&range, UnaryPredicate P)
Provide wrappers to std::any_of which take ranges instead of having to pass begin/end explicitly.
MachineInstr * getImm(const MachineOperand &MO, const MachineRegisterInfo *MRI)
constexpr bool isPowerOf2_32(uint32_t Value)
Return true if the argument is a power of two > 0.
Type * toTypedPointer(Type *Ty)
LLVM_ABI raw_ostream & dbgs()
dbgs() - This returns a reference to a raw_ostream for debugging messages.
LLVM_ABI void report_fatal_error(Error Err, bool gen_crash_diag=true)
constexpr bool isGenericCastablePtr(SPIRV::StorageClass::StorageClass SC)
class LLVM_GSL_OWNER SmallVector
Forward declaration of SmallVector so that calculateSmallVectorDefaultInlinedElements can reference s...
MachineInstr * passCopy(MachineInstr *Def, const MachineRegisterInfo *MRI)
bool isa(const From &Val)
isa<X> - Return true if the parameter to the template is an instance of one of the template type argu...
std::optional< SPIRV::LinkageType::LinkageType > getSpirvLinkageTypeFor(const SPIRVSubtarget &ST, const GlobalValue &GV)
LLVM_ABI raw_fd_ostream & errs()
This returns a reference to a raw_ostream for standard error.
SPIRV::StorageClass::StorageClass addressSpaceToStorageClass(unsigned AddrSpace, const SPIRVSubtarget &STI)
AtomicOrdering
Atomic ordering for LLVM's memory model.
InstructionSelector * createSPIRVInstructionSelector(const SPIRVTargetMachine &TM, const SPIRVSubtarget &Subtarget, const RegisterBankInfo &RBI)
std::string getStringValueFromReg(Register Reg, MachineRegisterInfo &MRI)
int64_t foldImm(const MachineOperand &MO, const MachineRegisterInfo *MRI)
DWARFExpression::Operation Op
ArrayRef(const T &OneElt) -> ArrayRef< T >
MachineInstr * getDefInstrMaybeConstant(Register &ConstReg, const MachineRegisterInfo *MRI)
constexpr unsigned BitWidth
decltype(auto) cast(const From &Val)
cast<X> - Return the argument parameter cast to the specified type.
bool hasInitializer(const GlobalVariable *GV)
bool isSpvIntrinsic(const MachineInstr &MI, Intrinsic::ID IntrinsicID)
MachineInstr * getVRegDef(MachineRegisterInfo &MRI, Register Reg)
SPIRV::MemorySemantics::MemorySemantics getMemSemantics(AtomicOrdering Ord)
std::string getLinkStringForBuiltIn(SPIRV::BuiltIn::BuiltIn BuiltInValue)
LLVM_ABI bool isTriviallyDead(const MachineInstr &MI, const MachineRegisterInfo &MRI)
Check whether an instruction MI is dead: it only defines dead virtual registers, and doesn't have oth...
MCRegisterClass TargetRegisterClass