36#include "llvm/IR/IntrinsicsSPIRV.h"
42#define DEBUG_TYPE "spirv-isel"
49 std::vector<std::pair<SPIRV::InstructionSet::InstructionSet, uint32_t>>;
54 std::optional<Register> Bias;
55 std::optional<Register>
Offset;
56 std::optional<Register> MinLod;
57 std::optional<Register> GradX;
58 std::optional<Register> GradY;
59 std::optional<Register> Lod;
60 std::optional<Register> Compare;
67 bool IsScalar =
false;
70llvm::SPIRV::SelectionControl::SelectionControl
71getSelectionOperandForImm(
int Imm) {
73 return SPIRV::SelectionControl::Flatten;
75 return SPIRV::SelectionControl::DontFlatten;
77 return SPIRV::SelectionControl::None;
81#define GET_GLOBALISEL_PREDICATE_BITSET
82#include "SPIRVGenGlobalISel.inc"
83#undef GET_GLOBALISEL_PREDICATE_BITSET
110#define GET_GLOBALISEL_PREDICATES_DECL
111#include "SPIRVGenGlobalISel.inc"
112#undef GET_GLOBALISEL_PREDICATES_DECL
114#define GET_GLOBALISEL_TEMPORARIES_DECL
115#include "SPIRVGenGlobalISel.inc"
116#undef GET_GLOBALISEL_TEMPORARIES_DECL
140 unsigned BitSetOpcode)
const;
144 unsigned BitSetOpcode)
const;
148 unsigned BitSetOpcode,
bool SwapPrimarySide)
const;
155 unsigned Opcode)
const;
158 unsigned Opcode)
const;
180 unsigned NewOpcode,
unsigned NegateOpcode = 0)
const;
189 Register castPtrToPtrToInt(
Register Ptr, SPIRV::StorageClass::StorageClass SC,
193 bool selectAtomicPtrValue(
209 unsigned OpType)
const;
277 unsigned Opcode)
const;
281 unsigned Opcode)
const;
285 unsigned Opcode)
const;
289 unsigned Opcode)
const;
291 template <
bool Signed>
294 template <
bool Signed>
301 template <
typename PickOpcodeFn>
304 PickOpcodeFn &&PickOpcode)
const;
321 template <
typename PickOpcodeFn>
324 PickOpcodeFn &&PickOpcode)
const;
342 bool IsSigned)
const;
344 bool IsSigned,
unsigned Opcode)
const;
346 bool IsSigned)
const;
352 bool IsSigned)
const;
393 GL::GLSLExtInst GLInst,
bool setMIFlags =
true,
394 bool useMISrc =
true,
396 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
397 CL::OpenCLExtInst CLInst,
bool setMIFlags =
true,
398 bool useMISrc =
true,
400 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
401 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
402 bool setMIFlags =
true,
bool useMISrc =
true,
404 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
405 const ExtInstList &ExtInsts,
bool setMIFlags =
true,
406 bool useMISrc =
true,
409 bool selectLog10(
Register ResVReg, SPIRVTypeInst ResType,
410 MachineInstr &
I)
const;
412 bool selectFpowi(
Register ResVReg, SPIRVTypeInst ResType,
413 MachineInstr &
I)
const;
415 bool selectSaturate(
Register ResVReg, SPIRVTypeInst ResType,
416 MachineInstr &
I)
const;
418 bool selectWaveOpInst(
Register ResVReg, SPIRVTypeInst ResType,
419 MachineInstr &
I,
unsigned Opcode)
const;
421 bool selectBarrierInst(MachineInstr &
I,
unsigned Scope,
unsigned MemSem,
422 bool WithGroupSync)
const;
424 bool selectWaveActiveCountBits(
Register ResVReg, SPIRVTypeInst ResType,
425 MachineInstr &
I)
const;
427 bool selectWaveActiveAllEqual(
Register ResVReg, SPIRVTypeInst ResType,
428 MachineInstr &
I)
const;
432 bool selectHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
433 MachineInstr &
I)
const;
435 bool selectCounterHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
436 MachineInstr &
I)
const;
438 bool selectReadImageIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
439 MachineInstr &
I)
const;
440 bool selectGetDimensionsIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
441 MachineInstr &
I)
const;
442 bool selectGetDimensionsLevelsIntrinsic(
Register &ResVReg,
443 SPIRVTypeInst ResType,
444 MachineInstr &
I)
const;
445 bool selectGetDimensionsMSIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
446 MachineInstr &
I)
const;
449 std::optional<Register> LodReg = std::nullopt)
const;
450 bool selectSampleBasicIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
451 MachineInstr &
I)
const;
452 bool selectCalculateLodIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
453 MachineInstr &
I)
const;
454 bool selectSampleBiasIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
455 MachineInstr &
I)
const;
456 bool selectSampleGradIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
457 MachineInstr &
I)
const;
458 bool selectSampleLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
459 MachineInstr &
I)
const;
460 bool selectLoadLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
461 MachineInstr &
I)
const;
462 bool selectSampleCmpIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
463 MachineInstr &
I)
const;
464 bool selectSampleCmpLevelZeroIntrinsic(
Register &ResVReg,
465 SPIRVTypeInst ResType,
466 MachineInstr &
I)
const;
467 bool selectGatherIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
468 MachineInstr &
I)
const;
469 bool selectImageWriteIntrinsic(MachineInstr &
I)
const;
470 bool selectResourceGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
471 MachineInstr &
I)
const;
472 bool selectPushConstantGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
473 MachineInstr &
I)
const;
474 bool selectResourceNonUniformIndex(
Register &ResVReg, SPIRVTypeInst ResType,
475 MachineInstr &
I)
const;
476 bool selectModf(
Register ResVReg, SPIRVTypeInst ResType,
477 MachineInstr &
I)
const;
478 bool selectUpdateCounter(
Register &ResVReg, SPIRVTypeInst ResType,
479 MachineInstr &
I)
const;
480 bool selectFrexp(
Register ResVReg, SPIRVTypeInst ResType,
481 MachineInstr &
I)
const;
483 bool selectLdexp(
Register ResVReg, SPIRVTypeInst ResType,
484 MachineInstr &
I)
const;
485 bool selectSincos(
Register ResVReg, SPIRVTypeInst ResType,
486 MachineInstr &
I)
const;
487 bool selectExp10(
Register ResVReg, SPIRVTypeInst ResType,
488 MachineInstr &
I)
const;
489 bool selectDerivativeInst(
Register ResVReg, SPIRVTypeInst ResType,
490 MachineInstr &
I,
const unsigned DPdOpCode)
const;
492 Register buildI32Constant(uint32_t Val, MachineInstr &
I,
493 SPIRVTypeInst ResType =
nullptr)
const;
494 Register buildI32ConstantInEntryBlock(uint32_t Val, MachineInstr &
I,
495 SPIRVTypeInst ResType =
nullptr)
const;
497 Register buildZerosVal(SPIRVTypeInst ResType, MachineInstr &
I)
const;
498 bool isScalarOrVectorIntConstantZero(
Register Reg)
const;
499 Register buildZerosValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
501 MachineInstr &
I)
const;
502 Register buildOnesValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
504 bool wrapIntoSpecConstantOp(MachineInstr &
I,
507 Register getUcharPtrTypeReg(MachineInstr &
I,
508 SPIRV::StorageClass::StorageClass SC)
const;
509 MachineInstrBuilder buildSpecConstantOp(MachineInstr &
I,
Register Dest,
511 uint32_t Opcode)
const;
512 MachineInstrBuilder buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
513 SPIRVTypeInst SrcPtrTy)
const;
514 Register buildPointerToResource(SPIRVTypeInst ResType,
515 SPIRV::StorageClass::StorageClass SC,
516 uint32_t Set, uint32_t
Binding,
517 uint32_t ArraySize,
Register IndexReg,
519 MachineIRBuilder MIRBuilder)
const;
520 SPIRVTypeInst widenTypeToVec4(SPIRVTypeInst
Type, MachineInstr &
I)
const;
521 bool extractSubvector(
Register &ResVReg, SPIRVTypeInst ResType,
522 Register &ReadReg, MachineInstr &InsertionPoint)
const;
523 bool generateImageReadOrFetch(
Register &ResVReg, SPIRVTypeInst ResType,
526 const ImageOperands *ImOps =
nullptr)
const;
527 bool generateSampleImage(
Register ResVReg, SPIRVTypeInst ResType,
529 Register CoordinateReg,
const ImageOperands &ImOps,
532 bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
533 Register ResVReg, SPIRVTypeInst ResType,
534 MachineInstr &
I)
const;
535 bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
536 Register ResVReg, SPIRVTypeInst ResType,
537 MachineInstr &
I)
const;
538 bool loadHandleBeforePosition(
Register &HandleReg, SPIRVTypeInst ResType,
539 GIntrinsic &HandleDef, MachineInstr &Pos)
const;
540 void decorateUsesAsNonUniform(
Register &NonUniformReg)
const;
541 bool errorIfInstrOutsideShader(MachineInstr &
I)
const;
543 std::optional<SplitParts> splitEvenOddLanes(
Register PopCountReg,
544 unsigned ComponentCount,
546 SPIRVTypeInst I32Type)
const;
549 handle64BitOverflow(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
550 Register SrcReg,
unsigned int Opcode,
551 std::function<
bool(
Register, SPIRVTypeInst,
552 MachineInstr &,
Register,
unsigned)>
556bool sampledTypeIsSignedInteger(
const llvm::Type *HandleType) {
558 if (
TET->getTargetExtName() ==
"spirv.Image") {
561 assert(
TET->getTargetExtName() ==
"spirv.SignedImage");
562 return TET->getTypeParameter(0)->isIntegerTy();
566#define GET_GLOBALISEL_IMPL
567#include "SPIRVGenGlobalISel.inc"
568#undef GET_GLOBALISEL_IMPL
574 TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
577#include
"SPIRVGenGlobalISel.inc"
580#include
"SPIRVGenGlobalISel.inc"
592 InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
597 if (HasVRegsReset == &MF)
612 for (
const auto &
MBB : MF) {
613 for (
const auto &
MI :
MBB) {
616 if (
MI.getOpcode() != SPIRV::ASSIGN_TYPE)
620 LLT DstType = MRI.
getType(DstReg);
622 LLT SrcType = MRI.
getType(SrcReg);
623 if (DstType != SrcType)
628 if (DstRC != SrcRC && SrcRC)
640 while (!Stack.empty()) {
645 switch (
MI->getOpcode()) {
646 case TargetOpcode::G_INTRINSIC:
647 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
648 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS: {
651 if (IntrID != Intrinsic::spv_const_composite &&
652 IntrID != Intrinsic::spv_undef && IntrID != Intrinsic::spv_poison)
656 case TargetOpcode::G_BUILD_VECTOR:
657 case TargetOpcode::G_SPLAT_VECTOR:
659 i < OpDef->getNumOperands(); i++) {
664 Stack.push_back(OpNestedDef);
667 case TargetOpcode::G_CONSTANT:
668 case TargetOpcode::G_FCONSTANT:
669 case TargetOpcode::G_IMPLICIT_DEF:
670 case SPIRV::OpConstantTrue:
671 case SPIRV::OpConstantFalse:
672 case SPIRV::OpConstantI:
673 case SPIRV::OpConstantF:
674 case SPIRV::OpConstantComposite:
675 case SPIRV::OpConstantCompositeContinuedINTEL:
676 case SPIRV::OpConstantSampler:
677 case SPIRV::OpConstantNull:
679 case SPIRV::OpPoisonKHR:
680 case SPIRV::OpConstantFunctionPointerINTEL:
707 case Intrinsic::spv_all:
708 case Intrinsic::spv_alloca:
709 case Intrinsic::spv_any:
710 case Intrinsic::spv_bitcast:
711 case Intrinsic::spv_const_composite:
712 case Intrinsic::spv_degrees:
713 case Intrinsic::spv_distance:
714 case Intrinsic::spv_extractelt:
715 case Intrinsic::spv_extractv:
716 case Intrinsic::spv_faceforward:
717 case Intrinsic::spv_fdot:
718 case Intrinsic::spv_firstbitlow:
719 case Intrinsic::spv_firstbitshigh:
720 case Intrinsic::spv_firstbituhigh:
721 case Intrinsic::spv_frac:
722 case Intrinsic::spv_gep:
723 case Intrinsic::spv_global_offset:
724 case Intrinsic::spv_global_size:
725 case Intrinsic::spv_group_id:
726 case Intrinsic::spv_insertelt:
727 case Intrinsic::spv_insertv:
728 case Intrinsic::spv_isinf:
729 case Intrinsic::spv_isnan:
730 case Intrinsic::spv_isfinite:
731 case Intrinsic::spv_isnormal:
732 case Intrinsic::spv_lerp:
733 case Intrinsic::spv_length:
734 case Intrinsic::spv_normalize:
735 case Intrinsic::spv_num_subgroups:
736 case Intrinsic::spv_num_workgroups:
737 case Intrinsic::spv_ptrcast:
738 case Intrinsic::spv_radians:
739 case Intrinsic::spv_reflect:
740 case Intrinsic::spv_refract:
741 case Intrinsic::spv_resource_getbasepointer:
742 case Intrinsic::spv_resource_getpointer:
743 case Intrinsic::spv_resource_handlefrombinding:
744 case Intrinsic::spv_resource_handlefromimplicitbinding:
745 case Intrinsic::spv_resource_nonuniformindex:
746 case Intrinsic::spv_resource_sample:
747 case Intrinsic::spv_rsqrt:
748 case Intrinsic::spv_saturate:
749 case Intrinsic::spv_sdot:
750 case Intrinsic::spv_sign:
751 case Intrinsic::spv_smoothstep:
752 case Intrinsic::spv_subgroup_id:
753 case Intrinsic::spv_subgroup_local_invocation_id:
754 case Intrinsic::spv_subgroup_max_size:
755 case Intrinsic::spv_subgroup_size:
756 case Intrinsic::spv_thread_id:
757 case Intrinsic::spv_thread_id_in_group:
758 case Intrinsic::spv_udot:
759 case Intrinsic::spv_undef:
760 case Intrinsic::spv_value_md:
761 case Intrinsic::spv_workgroup_size:
773 case SPIRV::OpTypeVoid:
774 case SPIRV::OpTypeBool:
775 case SPIRV::OpTypeInt:
776 case SPIRV::OpTypeFloat:
777 case SPIRV::OpTypeVector:
778 case SPIRV::OpTypeMatrix:
779 case SPIRV::OpTypeImage:
780 case SPIRV::OpTypeSampler:
781 case SPIRV::OpTypeSampledImage:
782 case SPIRV::OpTypeArray:
783 case SPIRV::OpTypeRuntimeArray:
784 case SPIRV::OpTypeStruct:
785 case SPIRV::OpTypeOpaque:
786 case SPIRV::OpTypePointer:
787 case SPIRV::OpTypeFunction:
788 case SPIRV::OpTypeEvent:
789 case SPIRV::OpTypeDeviceEvent:
790 case SPIRV::OpTypeReserveId:
791 case SPIRV::OpTypeQueue:
792 case SPIRV::OpTypePipe:
793 case SPIRV::OpTypeForwardPointer:
794 case SPIRV::OpTypePipeStorage:
795 case SPIRV::OpTypeNamedBarrier:
796 case SPIRV::OpTypeAccelerationStructureNV:
797 case SPIRV::OpTypeCooperativeMatrixNV:
798 case SPIRV::OpTypeCooperativeMatrixKHR:
808 if (
MI.getNumDefs() == 0)
811 for (
const auto &MO :
MI.all_defs()) {
813 if (
Reg.isPhysical()) {
818 if (
UseMI.getOpcode() != SPIRV::OpName) {
825 if (
MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE ||
MI.isFakeUse() ||
826 MI.isLifetimeMarker()) {
829 <<
"Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
840 if (
MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
841 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
844 LLVM_DEBUG(
dbgs() <<
"Dead: Intrinsic with no real side effects.\n");
849 if (
MI.mayStore() ||
MI.isCall() ||
850 (
MI.mayLoad() &&
MI.hasOrderedMemoryRef()) ||
MI.isPosition() ||
851 MI.isDebugInstr() ||
MI.isTerminator() ||
MI.isJumpTableDebugInfo()) {
852 LLVM_DEBUG(
dbgs() <<
"Not dead: instruction has side effects.\n");
863 LLVM_DEBUG(
dbgs() <<
"Dead: known opcode with no side effects\n");
870void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &
MI)
const {
872 for (
const auto &MO :
MI.all_defs()) {
876 SmallVector<MachineInstr *, 4> UselessOpNames;
879 "There is still a use of the dead function.");
882 for (MachineInstr *OpNameMI : UselessOpNames) {
884 OpNameMI->eraseFromParent();
889void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &
MI)
const {
892 removeOpNamesForDeadMI(
MI);
893 MI.eraseFromParent();
896bool SPIRVInstructionSelector::select(MachineInstr &
I) {
897 resetVRegsType(*
I.getParent()->getParent());
899 assert(
I.getParent() &&
"Instruction should be in a basic block!");
900 assert(
I.getParent()->getParent() &&
"Instruction should be in a function!");
905 removeDeadInstruction(
I);
912 if (Opcode == SPIRV::ASSIGN_TYPE) {
913 Register DstReg =
I.getOperand(0).getReg();
914 Register SrcReg =
I.getOperand(1).getReg();
917 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
918 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
919 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
920 Register SelectDstReg =
Def->getOperand(0).getReg();
921 bool SuccessToSelectSelect [[maybe_unused]] = selectSelect(
923 assert(SuccessToSelectSelect);
925 Def->eraseFromParent();
932 bool Res = selectImpl(
I, *CoverageInfo);
934 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
935 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
939 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
951 }
else if (
I.getNumDefs() == 1) {
963 removeDeadInstruction(
I);
968 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
969 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
975 bool HasDefs =
I.getNumDefs() > 0;
978 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
979 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
980 if (spvSelect(ResVReg, ResType,
I)) {
982 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
993 case TargetOpcode::G_CONSTANT:
994 case TargetOpcode::G_FCONSTANT:
1001 MachineInstr &
I)
const {
1004 if (DstRC != SrcRC && SrcRC)
1006 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
1013bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
1014 SPIRVTypeInst ResType,
1015 MachineInstr &
I)
const {
1016 const unsigned Opcode =
I.getOpcode();
1018 return selectImpl(
I, *CoverageInfo);
1020 case TargetOpcode::G_CONSTANT:
1021 case TargetOpcode::G_FCONSTANT:
1022 return selectConst(ResVReg, ResType,
I);
1023 case TargetOpcode::G_GLOBAL_VALUE:
1024 return selectGlobalValue(ResVReg,
I);
1025 case TargetOpcode::G_IMPLICIT_DEF:
1026 return selectOpUndef(ResVReg, ResType,
I);
1027 case TargetOpcode::G_FREEZE:
1028 return selectFreeze(ResVReg, ResType,
I);
1030 case TargetOpcode::G_INTRINSIC:
1031 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
1032 case TargetOpcode::G_INTRINSIC_CONVERGENT:
1033 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
1034 return selectIntrinsic(ResVReg, ResType,
I);
1035 case TargetOpcode::G_BITREVERSE:
1036 return selectBitreverse(ResVReg, ResType,
I);
1038 case TargetOpcode::G_BUILD_VECTOR:
1039 return selectBuildVector(ResVReg, ResType,
I);
1040 case TargetOpcode::G_SPLAT_VECTOR:
1041 return selectSplatVector(ResVReg, ResType,
I);
1042 case TargetOpcode::G_CONCAT_VECTORS:
1043 return selectConcatVectors(ResVReg, ResType,
I);
1045 case TargetOpcode::G_SHUFFLE_VECTOR: {
1046 MachineBasicBlock &BB = *
I.getParent();
1047 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
1050 .
addUse(
I.getOperand(1).getReg())
1051 .
addUse(
I.getOperand(2).getReg());
1052 for (
auto V :
I.getOperand(3).getShuffleMask())
1057 case TargetOpcode::G_MEMMOVE:
1058 case TargetOpcode::G_MEMCPY:
1059 case TargetOpcode::G_MEMCPY_INLINE:
1060 case TargetOpcode::G_MEMSET:
1061 case TargetOpcode::G_MEMSET_INLINE:
1062 return selectMemOperation(ResVReg,
I);
1064 case TargetOpcode::G_ICMP:
1065 return selectICmp(ResVReg, ResType,
I);
1066 case TargetOpcode::G_FCMP:
1067 return selectFCmp(ResVReg, ResType,
I);
1069 case TargetOpcode::G_FRAME_INDEX:
1070 return selectFrameIndex(ResVReg, ResType,
I);
1072 case TargetOpcode::G_LOAD:
1073 return selectLoad(ResVReg, ResType,
I);
1074 case TargetOpcode::G_STORE:
1075 return selectStore(
I);
1077 case TargetOpcode::G_BR:
1078 return selectBranch(
I);
1079 case TargetOpcode::G_BRCOND:
1080 return selectBranchCond(
I);
1082 case TargetOpcode::G_PHI:
1083 return selectPhi(ResVReg,
I);
1085 case TargetOpcode::G_FPTOSI:
1086 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1087 case TargetOpcode::G_FPTOUI:
1088 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1090 case TargetOpcode::G_FPTOSI_SAT:
1091 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1092 case TargetOpcode::G_FPTOUI_SAT:
1093 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1095 case TargetOpcode::G_SITOFP:
1096 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
1097 case TargetOpcode::G_UITOFP:
1098 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
1100 case TargetOpcode::G_CTPOP:
1101 return selectPopCount(ResVReg, ResType,
I, SPIRV::OpBitCount);
1102 case TargetOpcode::G_SMIN:
1103 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
1104 case TargetOpcode::G_UMIN:
1105 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
1107 case TargetOpcode::G_SMAX:
1108 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
1109 case TargetOpcode::G_UMAX:
1110 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
1112 case TargetOpcode::G_SCMP:
1113 return selectSUCmp(ResVReg, ResType,
I,
true);
1114 case TargetOpcode::G_UCMP:
1115 return selectSUCmp(ResVReg, ResType,
I,
false);
1116 case TargetOpcode::G_LROUND:
1117 case TargetOpcode::G_LLROUND: {
1120 MRI->
setRegClass(regForLround, &SPIRV::iIDRegClass);
1122 regForLround, *(
I.getParent()->getParent()));
1124 CL::round, GL::Round,
false);
1126 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1133 case TargetOpcode::G_STRICT_FMA:
1134 case TargetOpcode::G_FMA: {
1137 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1140 .
addUse(
I.getOperand(1).getReg())
1141 .
addUse(
I.getOperand(2).getReg())
1142 .
addUse(
I.getOperand(3).getReg())
1147 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1150 case TargetOpcode::G_FLDEXP:
1151 case TargetOpcode::G_STRICT_FLDEXP:
1152 return selectLdexp(ResVReg, ResType,
I);
1154 case TargetOpcode::G_FPOW:
1155 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1156 case TargetOpcode::G_FPOWI:
1157 return selectFpowi(ResVReg, ResType,
I);
1159 case TargetOpcode::G_FEXP:
1160 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1161 case TargetOpcode::G_FEXP2:
1162 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1163 case TargetOpcode::G_FEXP10:
1164 return selectExp10(ResVReg, ResType,
I);
1166 case TargetOpcode::G_FMODF:
1167 return selectModf(ResVReg, ResType,
I);
1168 case TargetOpcode::G_FSINCOS:
1169 return selectSincos(ResVReg, ResType,
I);
1171 case TargetOpcode::G_FLOG:
1172 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1173 case TargetOpcode::G_FLOG2:
1174 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1175 case TargetOpcode::G_FLOG10:
1176 return selectLog10(ResVReg, ResType,
I);
1178 case TargetOpcode::G_FABS:
1179 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1180 case TargetOpcode::G_ABS:
1181 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1183 case TargetOpcode::G_FMINNUM:
1184 case TargetOpcode::G_FMINIMUM:
1185 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1186 case TargetOpcode::G_FMAXNUM:
1187 case TargetOpcode::G_FMAXIMUM:
1188 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1190 case TargetOpcode::G_FCOPYSIGN:
1191 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1193 case TargetOpcode::G_FCEIL:
1194 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1195 case TargetOpcode::G_FFLOOR:
1196 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1198 case TargetOpcode::G_FCOS:
1199 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1200 case TargetOpcode::G_FSIN:
1201 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1202 case TargetOpcode::G_FTAN:
1203 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1204 case TargetOpcode::G_FACOS:
1205 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1206 case TargetOpcode::G_FASIN:
1207 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1208 case TargetOpcode::G_FATAN:
1209 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1210 case TargetOpcode::G_FATAN2:
1211 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1212 case TargetOpcode::G_FCOSH:
1213 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1214 case TargetOpcode::G_FSINH:
1215 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1216 case TargetOpcode::G_FTANH:
1217 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1219 case TargetOpcode::G_STRICT_FSQRT:
1220 case TargetOpcode::G_FSQRT:
1221 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1223 case TargetOpcode::G_CTTZ:
1224 case TargetOpcode::G_CTTZ_ZERO_POISON:
1225 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1226 case TargetOpcode::G_CTLZ:
1227 case TargetOpcode::G_CTLZ_ZERO_POISON:
1228 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1230 case TargetOpcode::G_INTRINSIC_ROUND:
1231 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1232 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1233 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1234 case TargetOpcode::G_INTRINSIC_TRUNC:
1235 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1236 case TargetOpcode::G_FRINT:
1237 case TargetOpcode::G_FNEARBYINT:
1238 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1240 case TargetOpcode::G_SMULH:
1241 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1242 case TargetOpcode::G_UMULH:
1243 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1245 case TargetOpcode::G_SADDSAT:
1246 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1247 case TargetOpcode::G_UADDSAT:
1248 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1249 case TargetOpcode::G_SSUBSAT:
1250 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1251 case TargetOpcode::G_USUBSAT:
1252 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1254 case TargetOpcode::G_FFREXP:
1255 return selectFrexp(ResVReg, ResType,
I);
1257 case TargetOpcode::G_UADDO:
1258 return selectOverflowArith(ResVReg, ResType,
I,
1259 ResType->
getOpcode() == SPIRV::OpTypeVector
1260 ? SPIRV::OpIAddCarryV
1261 : SPIRV::OpIAddCarryS);
1262 case TargetOpcode::G_USUBO:
1263 return selectOverflowArith(ResVReg, ResType,
I,
1264 ResType->
getOpcode() == SPIRV::OpTypeVector
1265 ? SPIRV::OpISubBorrowV
1266 : SPIRV::OpISubBorrowS);
1267 case TargetOpcode::G_UMULO:
1268 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1269 case TargetOpcode::G_SMULO:
1270 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1272 case TargetOpcode::G_SEXT:
1273 return selectExt(ResVReg, ResType,
I,
true);
1274 case TargetOpcode::G_ANYEXT:
1275 case TargetOpcode::G_ZEXT:
1276 return selectExt(ResVReg, ResType,
I,
false);
1277 case TargetOpcode::G_TRUNC:
1278 return selectTrunc(ResVReg, ResType,
I);
1279 case TargetOpcode::G_FPTRUNC:
1280 case TargetOpcode::G_FPEXT:
1281 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1283 case TargetOpcode::G_PTRTOINT:
1284 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1285 case TargetOpcode::G_INTTOPTR:
1286 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1287 case TargetOpcode::G_BITCAST:
1288 return selectBitcast(ResVReg, ResType,
I);
1289 case TargetOpcode::G_ADDRSPACE_CAST:
1290 return selectAddrSpaceCast(ResVReg, ResType,
I);
1291 case TargetOpcode::G_PTRMASK:
1292 return selectPtrMask(ResVReg, ResType,
I);
1293 case TargetOpcode::G_PTR_ADD: {
1295 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1299 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1300 (*II).getOpcode() == TargetOpcode::COPY ||
1301 (*II).getOpcode() == SPIRV::OpVariable ||
1302 (*II).getOpcode() == SPIRV::OpUntypedVariableKHR) &&
1303 getImm(
I.getOperand(2), MRI));
1305 bool IsGVInit =
false;
1309 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1310 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1311 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1312 (*UseIt).getOpcode() == SPIRV::OpVariable ||
1313 (*UseIt).getOpcode() == SPIRV::OpUntypedVariableKHR) {
1325 const bool UseUntypedPointers =
1326 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1327 if (UseUntypedPointers) {
1328 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1331 .
addImm(
static_cast<uint32_t
>(
1332 SPIRV::Opcode::UntypedInBoundsPtrAccessChainKHR))
1335 .
addUse(
I.getOperand(2).getReg())
1342 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1354 return diagnoseUnsupported(
1355 I,
"incompatible result and operand types in a bitcast");
1357 MachineInstrBuilder MIB =
1358 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1365 : SPIRV::OpInBoundsPtrAccessChain))
1369 .
addUse(
I.getOperand(2).getReg())
1372 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1376 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1378 .
addUse(
I.getOperand(2).getReg())
1387 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1390 .
addImm(
static_cast<uint32_t
>(
1391 SPIRV::Opcode::InBoundsPtrAccessChain))
1394 .
addUse(
I.getOperand(2).getReg());
1399 case TargetOpcode::G_ATOMICRMW_OR:
1400 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1401 case TargetOpcode::G_ATOMICRMW_ADD:
1402 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1403 case TargetOpcode::G_ATOMICRMW_AND:
1404 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1405 case TargetOpcode::G_ATOMICRMW_MAX:
1406 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1407 case TargetOpcode::G_ATOMICRMW_MIN:
1408 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1409 case TargetOpcode::G_ATOMICRMW_SUB:
1410 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1411 case TargetOpcode::G_ATOMICRMW_XOR:
1412 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1413 case TargetOpcode::G_ATOMICRMW_UMAX:
1414 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1415 case TargetOpcode::G_ATOMICRMW_UMIN:
1416 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1417 case TargetOpcode::G_ATOMICRMW_XCHG:
1418 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1420 case TargetOpcode::G_ATOMICRMW_FADD:
1421 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1422 case TargetOpcode::G_ATOMICRMW_FSUB:
1424 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1425 ResType->
getOpcode() == SPIRV::OpTypeVector
1427 : SPIRV::OpFNegate);
1428 case TargetOpcode::G_ATOMICRMW_FMIN:
1429 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1430 case TargetOpcode::G_ATOMICRMW_FMAX:
1431 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1433 case TargetOpcode::G_FENCE:
1434 return selectFence(
I);
1436 case TargetOpcode::G_STACKSAVE:
1437 return selectStackSave(ResVReg, ResType,
I);
1438 case TargetOpcode::G_STACKRESTORE:
1439 return selectStackRestore(
I);
1441 case TargetOpcode::G_UNMERGE_VALUES:
1444 case TargetOpcode::G_TRAP:
1445 case TargetOpcode::G_UBSANTRAP:
1446 return selectTrap(
I);
1451 case TargetOpcode::DBG_LABEL:
1453 case TargetOpcode::G_DEBUGTRAP:
1454 return selectDebugTrap(ResVReg, ResType,
I);
1455 case TargetOpcode::G_PREFETCH:
1456 return selectPrefetch(
I);
1463bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1464 SPIRVTypeInst ResType,
1465 MachineInstr &
I)
const {
1466 unsigned Opcode = SPIRV::OpNop;
1473bool SPIRVInstructionSelector::selectPrefetch(MachineInstr &
I)
const {
1482 MachineIRBuilder MIRBuilder(
I);
1484 const SPIRVTypeInst PointerSizeType =
1492 Register AddrVal =
I.getOperand(0).getReg();
1495 return selectExtInst(ExtReg, GR.
getOpTypeVoid(MIRBuilder),
I, CL::prefetch,
1497 {AddrVal, ConstIntOne});
1502bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1503 SPIRVTypeInst ResType,
1505 GL::GLSLExtInst GLInst,
1506 bool setMIFlags,
bool useMISrc,
1509 SPIRV::InstructionSet::InstructionSet::GLSL_std_450))
1510 return diagnoseUnsupported(
1512 "this instruction is only supported with the GLSL extended instruction "
1514 return selectExtInst(ResVReg, ResType,
I,
1515 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1516 setMIFlags, useMISrc, SrcRegs);
1519bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1520 SPIRVTypeInst ResType,
1522 CL::OpenCLExtInst CLInst,
1523 bool setMIFlags,
bool useMISrc,
1525 return selectExtInst(ResVReg, ResType,
I,
1526 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1527 setMIFlags, useMISrc, SrcRegs);
1530bool SPIRVInstructionSelector::selectExtInst(
1531 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1532 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1534 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1535 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1536 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1540bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1541 SPIRVTypeInst ResType,
1544 bool setMIFlags,
bool useMISrc,
1547 for (
const auto &[InstructionSet, Opcode] : Insts) {
1551 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1554 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1559 const unsigned NumOps =
I.getNumOperands();
1562 I.getOperand(Index).getType() ==
1563 MachineOperand::MachineOperandType::MO_IntrinsicID)
1566 MIB.
add(
I.getOperand(Index));
1578bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1579 SPIRVTypeInst ResType,
1580 MachineInstr &
I)
const {
1581 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1582 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1583 for (
const auto &Ex : ExtInsts) {
1584 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1585 uint32_t Opcode = Ex.second;
1589 MachineIRBuilder MIRBuilder(
I);
1592 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1599 const bool IsUntyped =
1600 PointerType->getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1602 BuildMI(*It->getParent(), It, It->getDebugLoc(),
1603 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
1604 : SPIRV::OpVariable))
1607 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
1613 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1616 .
addImm(
static_cast<uint32_t
>(Ex.first))
1618 .
add(
I.getOperand(2))
1622 Register ExpResReg =
I.getOperand(1).getReg();
1624 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1634bool SPIRVInstructionSelector::selectLdexp(
Register ResVReg,
1635 SPIRVTypeInst ResType,
1636 MachineInstr &
I)
const {
1637 Register XReg =
I.getOperand(1).getReg();
1638 Register ExpReg =
I.getOperand(2).getReg();
1644 if (ResType->
getOpcode() == SPIRV::OpTypeVector &&
1645 ExpType->
getOpcode() != SPIRV::OpTypeVector) {
1647 SPIRVTypeInst ExpVecType =
1651 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1652 TII.get(SPIRV::OpCompositeConstruct))
1655 for (
unsigned J = 0; J < NumElts; ++J)
1661 return selectExtInst(ResVReg, ResType,
I, CL::ldexp, GL::Ldexp,
1662 true,
false, {XReg, ExpReg});
1665bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1666 SPIRVTypeInst ResType,
1667 MachineInstr &
I)
const {
1668 Register CosResVReg =
I.getOperand(1).getReg();
1669 unsigned SrcIdx =
I.getNumExplicitDefs();
1674 MachineIRBuilder MIRBuilder(
I);
1676 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1683 const bool IsUntyped =
1684 PointerType->getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1686 BuildMI(*It->getParent(), It, It->getDebugLoc(),
1687 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
1688 : SPIRV::OpVariable))
1691 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
1695 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1698 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1700 .
add(
I.getOperand(SrcIdx))
1704 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1712 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1715 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1717 .
add(
I.getOperand(SrcIdx))
1719 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1722 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1724 .
add(
I.getOperand(SrcIdx))
1731bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1732 SPIRVTypeInst ResType,
1735 unsigned Opcode)
const {
1736 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1746std::optional<SplitParts> SPIRVInstructionSelector::splitEvenOddLanes(
1747 Register PopCountReg,
unsigned ComponentCount, MachineInstr &
I,
1748 SPIRVTypeInst I32Type)
const {
1751 if (ComponentCount == 1) {
1754 Parts.IsScalar =
true;
1755 Parts.Type = I32Type;
1763 if (!selectOpWithSrcs(Parts.High, I32Type,
I, {PopCountReg, IdxOne},
1764 SPIRV::OpVectorExtractDynamic))
1765 return std::nullopt;
1767 if (!selectOpWithSrcs(Parts.Low, I32Type,
I, {PopCountReg, IdxZero},
1768 SPIRV::OpVectorExtractDynamic))
1769 return std::nullopt;
1773 MachineIRBuilder MIRBuilder(
I);
1774 Parts.IsScalar =
false;
1781 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1782 TII.get(SPIRV::OpVectorShuffle))
1787 for (
unsigned J = 1; J < ComponentCount * 2; J += 2)
1792 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1793 TII.get(SPIRV::OpVectorShuffle))
1798 for (
unsigned J = 0; J < ComponentCount * 2; J += 2)
1806bool SPIRVInstructionSelector::selectPopCount16(
Register ResVReg,
1807 SPIRVTypeInst ResType,
1810 unsigned Opcode)
const {
1811 Register OpReg =
I.getOperand(1).getReg();
1814 MachineIRBuilder MIRBuilder(
I);
1816 SPIRVTypeInst I32VectorType =
1819 bool IsVector = NumElems > 1;
1820 SPIRVTypeInst ExtType = IsVector ? I32VectorType : I32Type;
1823 if (!selectOpWithSrcs(ExtReg, ExtType,
I, {OpReg}, SPIRV::OpUConvert))
1827 if (!selectPopCount32(PopCountReg, ExtType,
I, ExtReg, Opcode))
1830 return selectOpWithSrcs(ResVReg, ResType,
I, {PopCountReg}, ExtOpcode);
1833bool SPIRVInstructionSelector::selectPopCount32(
Register ResVReg,
1834 SPIRVTypeInst ResType,
1837 unsigned Opcode)
const {
1838 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
1841bool SPIRVInstructionSelector::selectPopCount64(
Register ResVReg,
1842 SPIRVTypeInst ResType,
1845 unsigned Opcode)
const {
1847 if (ComponentCount > 2)
1848 return handle64BitOverflow(
1849 ResVReg, ResType,
I, SrcReg, Opcode,
1851 unsigned O) {
return this->selectPopCount64(R,
T,
I, S, O); });
1853 MachineIRBuilder MIRBuilder(
I);
1858 I32Type, 2 * ComponentCount, MIRBuilder,
false);
1862 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
1867 if (!selectPopCount32(Pop32, VecI32Type,
I, Vec32, Opcode))
1871 auto MaybeParts = splitEvenOddLanes(Pop32, ComponentCount,
I, I32Type);
1874 SplitParts &Parts = *MaybeParts;
1877 unsigned OpAdd = Parts.IsScalar ? SPIRV::OpIAddS : SPIRV::OpIAddV;
1879 if (!selectOpWithSrcs(Sum, Parts.Type,
I, {Parts.High, Parts.Low}, OpAdd))
1884 unsigned ConvOp = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
1885 return selectOpWithSrcs(ResVReg, ResType,
I, {Sum}, ConvOp);
1888bool SPIRVInstructionSelector::selectPopCount(
Register ResVReg,
1889 SPIRVTypeInst ResType,
1891 unsigned Opcode)
const {
1896 if (!STI.getTargetTriple().isVulkanOS())
1897 return selectUnOp(ResVReg, ResType,
I, Opcode);
1899 Register OpReg =
I.getOperand(1).getReg();
1902 : SPIRV::OpUConvert;
1906 return selectPopCount16(ResVReg, ResType,
I, ExtOpcode, Opcode);
1908 return selectPopCount32(ResVReg, ResType,
I, OpReg, Opcode);
1910 return selectPopCount64(ResVReg, ResType,
I, OpReg, Opcode);
1912 return diagnoseUnsupported(
I,
"unsupported operand bit width for popcount");
1916bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1917 SPIRVTypeInst ResType,
1919 unsigned Opcode)
const {
1921 Register SrcReg =
I.getOperand(1).getReg();
1926 unsigned DefOpCode = DefIt->getOpcode();
1927 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1930 if (
auto *VRD =
getVRegDef(*MRI, DefIt->getOperand(1).getReg()))
1931 DefOpCode = VRD->getOpcode();
1933 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1934 DefOpCode == TargetOpcode::G_CONSTANT ||
1935 DefOpCode == SPIRV::OpVariable ||
1936 DefOpCode == SPIRV::OpUntypedVariableKHR ||
1937 DefOpCode == SPIRV::OpConstantI) {
1943 uint32_t SpecOpcode = 0;
1945 case SPIRV::OpConvertPtrToU:
1946 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1948 case SPIRV::OpConvertUToPtr:
1949 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1954 TII.get(SPIRV::OpSpecConstantOp))
1964 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1968bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1969 SPIRVTypeInst ResType,
1970 MachineInstr &
I)
const {
1971 Register OpReg =
I.getOperand(1).getReg();
1972 SPIRVTypeInst OpType =
1975 return diagnoseUnsupported(
1976 I,
"incompatible result and operand types in a bitcast");
1977 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1988 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1989 if (
MemOp->isNonTemporal())
1990 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1992 if (!ST->isShader() &&
MemOp->getAlign().value())
1993 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1997 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1998 if (
auto *MD =
MemOp->getAAInfo().Scope) {
2002 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
2004 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
2008 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
2012 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
2014 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
2026 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
2028 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
2030 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
2034bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
2035 SPIRVTypeInst ResType,
2036 MachineInstr &
I)
const {
2038 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2043 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2044 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2046 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2048 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2052 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2056 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2057 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
2058 I.getDebugLoc(),
I);
2062 MachineIRBuilder MIRBuilder(
I);
2064 if (
I.getNumMemOperands()) {
2065 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2066 if (MemOp->isAtomic())
2067 return selectAtomicLoad(ResVReg, ResType,
I);
2070 auto MIB = MIRBuilder.
buildInstr(SPIRV::OpLoad)
2074 if (!
I.getNumMemOperands()) {
2075 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2077 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2086Register SPIRVInstructionSelector::createPtrSizedIntReg(
2087 MachineIRBuilder &MIRBuilder)
const {
2088 SPIRVTypeInst IntType =
2098SPIRVInstructionSelector::convertPtrToInt(
Register PtrVal,
2099 MachineIRBuilder &MIRBuilder)
const {
2100 SPIRVTypeInst IntType =
2102 Register IntReg = createPtrSizedIntReg(MIRBuilder);
2103 MIRBuilder.
buildInstr(SPIRV::OpConvertPtrToU)
2111Register SPIRVInstructionSelector::castPtrToPtrToInt(
2112 Register Ptr, SPIRV::StorageClass::StorageClass SC,
2113 MachineIRBuilder &MIRBuilder)
const {
2114 SPIRVTypeInst IntType =
2116 SPIRVTypeInst PtrType =
2130bool SPIRVInstructionSelector::selectAtomicPtrValue(
2131 Register ResVReg, SPIRVTypeInst ResType, MachineIRBuilder &MIRBuilder,
2132 function_ref<
Register(SPIRVTypeInst IntType)> EmitAtomic)
const {
2140 Register IntResult = EmitAtomic(IntType);
2142 MIRBuilder.
buildInstr(SPIRV::OpConvertUToPtr)
2150bool SPIRVInstructionSelector::selectAtomicLoad(
Register ResVReg,
2151 SPIRVTypeInst ResType,
2152 MachineInstr &
I)
const {
2153 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2156 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2159 return diagnoseUnsupported(
2160 I,
"Lowering to SPIR-V of atomic load is only "
2161 "allowed for integer, floating point or pointer types");
2163 assert(
I.getNumMemOperands());
2164 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2165 assert(MemOp.isAtomic());
2167 uint32_t
Scope =
static_cast<uint32_t
>(
2169 Register ScopeReg = buildI32Constant(Scope,
I);
2175 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2176 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2179 Register MemSemReg = buildI32Constant(Sem,
I);
2181 MachineIRBuilder MIRBuilder(
I);
2185 return diagnoseUnsupported(
2186 I,
"Lowering to SPIR-V of atomic load is only "
2187 "allowed for pointer types for physical addressing model");
2192 SPIRV::StorageClass::StorageClass SC =
2194 return selectAtomicPtrValue(
2195 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2196 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2197 Register IntResult = createPtrSizedIntReg(MIRBuilder);
2208 auto AtomicLoad = MIRBuilder.
buildInstr(SPIRV::OpAtomicLoad)
2219bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
2221 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2222 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2227 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2228 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2230 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2235 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2239 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2240 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2241 SPIRVTypeInst SampledType =
2243 SPIRVTypeInst StoreValCompType =
2245 if (StoreValCompType && StoreValCompType != SampledType) {
2248 SPIRVTypeInst PackedType = widenTypeToVec4(SampledType,
I);
2251 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
2256 StoreVal = PackedReg;
2259 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2260 TII.get(SPIRV::OpImageWrite))
2266 if (sampledTypeIsSignedInteger(LLVMHandleType))
2269 BMI.constrainAllUses(
TII,
TRI, RBI);
2274 if (
I.getNumMemOperands()) {
2275 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2276 if (MemOp->isAtomic())
2277 return selectAtomicStore(
I);
2284 if (PtrSC == SPIRV::StorageClass::UniformConstant ||
2285 PtrSC == SPIRV::StorageClass::Input ||
2286 PtrSC == SPIRV::StorageClass::PushConstant)
2287 return diagnoseUnsupported(
2288 I,
"store into a read-only SPIR-V storage class is not allowed");
2290 MachineIRBuilder MIRBuilder(
I);
2292 if (!
I.getNumMemOperands()) {
2293 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2295 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2304bool SPIRVInstructionSelector::selectAtomicStore(MachineInstr &
I)
const {
2305 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2308 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2309 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2314 if (!PointeeType && PtrType &&
2315 PtrType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR)
2318 return diagnoseUnsupported(
I,
2319 "Lowering to SPIR-V of atomic store is only "
2320 "allowed for integer or floating point types");
2322 assert(
I.getNumMemOperands());
2323 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2324 assert(MemOp.isAtomic());
2326 uint32_t
Scope =
static_cast<uint32_t
>(
2328 Register ScopeReg = buildI32Constant(Scope,
I);
2334 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2335 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2338 Register MemSemReg = buildI32Constant(Sem,
I);
2339 MachineIRBuilder MIRBuilder(
I);
2343 return diagnoseUnsupported(
2344 I,
"Lowering to SPIR-V of atomic store is only "
2345 "allowed for pointer types for physical addressing model");
2350 SPIRV::StorageClass::StorageClass SC =
2352 return selectAtomicPtrValue(
2353 Register(), SPIRVTypeInst(), MIRBuilder, [&](SPIRVTypeInst IntType) {
2355 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2368 return diagnoseUnsupported(
I,
2369 "Lowering to SPIR-V of atomic store is only "
2370 "allowed for integer or floating point types");
2372 auto AtomicStore = MIRBuilder.
buildInstr(SPIRV::OpAtomicStore)
2382bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
2383 SPIRVTypeInst ResType,
2384 MachineInstr &
I)
const {
2385 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
2393 const Register PtrsReg =
I.getOperand(2).getReg();
2394 const uint32_t Alignment =
I.getOperand(3).getImm();
2395 const Register MaskReg =
I.getOperand(4).getReg();
2396 const Register PassthruReg =
I.getOperand(5).getReg();
2397 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2401 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
2412bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
2413 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
2420 const Register ValuesReg =
I.getOperand(1).getReg();
2421 const Register PtrsReg =
I.getOperand(2).getReg();
2422 const uint32_t Alignment =
I.getOperand(3).getImm();
2423 const Register MaskReg =
I.getOperand(4).getReg();
2424 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2428 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
2437bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
2438 const Twine &
Msg)
const {
2439 const Function &
F =
I.getMF()->getFunction();
2440 F.getContext().diagnose(
2441 DiagnosticInfoUnsupported(
F,
Msg,
I.getDebugLoc(),
DS_Error));
2445bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
2446 SPIRVTypeInst ResType,
2447 MachineInstr &
I)
const {
2448 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2449 return diagnoseUnsupported(
2450 I,
"llvm.stacksave intrinsic: this instruction requires the following "
2451 "SPIR-V extension: SPV_INTEL_variable_length_array");
2453 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
2460bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
2461 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2462 return diagnoseUnsupported(
2464 "llvm.stackrestore intrinsic: this instruction requires the following "
2465 "SPIR-V extension: SPV_INTEL_variable_length_array");
2466 if (!
I.getOperand(0).isReg())
2469 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
2470 .
addUse(
I.getOperand(0).getReg())
2476SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
2477 MachineIRBuilder MIRBuilder(
I);
2478 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
2485 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
2489 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2490 Type *ArrTy = ArrayType::get(ValTy, Num);
2492 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
2495 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
2506 const bool IsUntyped = VarTy->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
2507 auto MIBVar =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2508 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
2509 : SPIRV::OpVariable))
2512 .
addImm(SPIRV::StorageClass::UniformConstant);
2525bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
2528 Register DstReg =
I.getOperand(0).getReg();
2532 return diagnoseUnsupported(
2533 I,
"OpCopyMemory requires operands to have the same type");
2534 uint64_t CopySize =
getIConstVal(
I.getOperand(2).getReg(), MRI);
2538 return diagnoseUnsupported(
2539 I,
"Unable to determine pointee type size for OpCopyMemory");
2540 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
2541 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
2542 return diagnoseUnsupported(
2543 I,
"OpCopyMemory requires the size to match the pointee type size");
2544 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemory))
2547 if (
I.getNumMemOperands()) {
2548 MachineIRBuilder MIRBuilder(
I);
2555bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &
I,
2558 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemorySized))
2559 .
addUse(
I.getOperand(0).getReg())
2561 .
addUse(
I.getOperand(2).getReg());
2562 if (
I.getNumMemOperands()) {
2563 MachineIRBuilder MIRBuilder(
I);
2570bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
2571 MachineInstr &
I)
const {
2573 Register SizeReg =
I.getOperand(2).getReg();
2575 SizeDef && SizeDef->
getOpcode() == TargetOpcode::G_CONSTANT &&
2579 Register SrcReg =
I.getOperand(1).getReg();
2580 if (
I.getOpcode() == TargetOpcode::G_MEMSET ||
2581 I.getOpcode() == TargetOpcode::G_MEMSET_INLINE) {
2582 Register VarReg = getOrCreateMemSetGlobal(
I);
2585 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2587 ValTy,
I, SPIRV::StorageClass::UniformConstant);
2589 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
2593 if (!selectCopyMemory(
I, SrcReg))
2596 if (!selectCopyMemorySized(
I, SrcReg))
2599 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
2600 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
2605bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
2606 SPIRVTypeInst ResType,
2609 unsigned NegateOpcode)
const {
2611 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2612 uint32_t
Scope =
static_cast<uint32_t
>(
2614 MemOp->getSyncScopeID()));
2615 Register ScopeReg = buildI32Constant(Scope,
I);
2617 Register Ptr =
I.getOperand(1).getReg();
2618 uint32_t ScSem =
static_cast<uint32_t
>(
2622 Register MemSemReg = buildI32Constant(
2626 Register ValueReg =
I.getOperand(2).getReg();
2627 if (NegateOpcode != 0) {
2630 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
2636 if (NewOpcode != SPIRV::OpAtomicExchange)
2637 return diagnoseUnsupported(
2638 I,
"Lowering to SPIR-V of this atomic operation is not "
2639 "allowed for pointer types");
2641 return diagnoseUnsupported(
2642 I,
"Lowering to SPIR-V of atomic exchange is only "
2643 "allowed for pointer types for physical addressing model");
2650 MachineIRBuilder MIRBuilder(
I);
2652 return selectAtomicPtrValue(
2653 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2655 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2656 Register ExchangeResReg = createPtrSizedIntReg(MIRBuilder);
2657 MIRBuilder.
buildInstr(SPIRV::OpAtomicExchange)
2665 return ExchangeResReg;
2669 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2680bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2681 unsigned ArgI =
I.getNumOperands() - 1;
2683 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2684 SPIRVTypeInst SrcType =
2686 if (!SrcType || SrcType->
getOpcode() != SPIRV::OpTypeVector)
2688 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2692 unsigned CurrentIndex = 0;
2693 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2694 Register ResVReg =
I.getOperand(i).getReg();
2697 LLT ResLLT = MRI->
getType(ResVReg);
2703 ResType = ScalarType;
2709 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
2712 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2718 for (
unsigned j = 0;
j < NumElements; ++
j) {
2719 MIB.
addImm(CurrentIndex + j);
2721 CurrentIndex += NumElements;
2725 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2737bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2740 Register MemSemReg = buildI32Constant(MemSem,
I);
2744 Register ScopeReg = buildI32Constant(Scope,
I);
2746 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2753bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2754 SPIRVTypeInst ResType,
2756 unsigned Opcode)
const {
2757 Type *ResTy =
nullptr;
2760 return diagnoseUnsupported(
2762 "Not enough info to select the arithmetic with overflow instruction");
2764 return diagnoseUnsupported(
I,
2765 "Expect struct type result for the arithmetic "
2766 "with overflow instruction");
2772 MachineIRBuilder MIRBuilder(
I);
2774 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2775 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2781 Register ZeroReg = buildZerosVal(ResType,
I);
2786 if (ResName.
size() > 0)
2794 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2795 MIB.
addUse(
I.getOperand(i).getReg());
2800 MRI->
setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2801 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2803 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2804 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2811 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2812 .
addDef(
I.getOperand(1).getReg())
2820bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2821 SPIRVTypeInst ResType,
2822 MachineInstr &
I)
const {
2824 "selectAtomicCmpXchg only handles the spv_cmpxchg intrinsic");
2825 Register Ptr =
I.getOperand(2).getReg();
2826 Register ScopeReg =
I.getOperand(5).getReg();
2827 Register MemSemEqReg =
I.getOperand(6).getReg();
2828 Register MemSemNeqReg =
I.getOperand(7).getReg();
2830 Register Val =
I.getOperand(4).getReg();
2834 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2853 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2860 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2872 case SPIRV::StorageClass::DeviceOnlyINTEL:
2873 case SPIRV::StorageClass::HostOnlyINTEL:
2882 bool IsGRef =
false;
2883 bool IsAllowedRefs =
2885 unsigned Opcode = It.getOpcode();
2886 if (Opcode == SPIRV::OpConstantComposite ||
2887 Opcode == SPIRV::OpSpecConstantComposite ||
2888 Opcode == SPIRV::OpVariable ||
2889 Opcode == SPIRV::OpUntypedVariableKHR ||
2890 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2891 return IsGRef = true;
2892 return Opcode == SPIRV::OpName;
2894 return IsAllowedRefs && IsGRef;
2897Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2898 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2900 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2904SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2906 uint32_t Opcode)
const {
2907 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2908 TII.get(SPIRV::OpSpecConstantOp))
2916SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2917 SPIRVTypeInst SrcPtrTy)
const {
2918 SPIRVTypeInst GenericPtrTy =
2922 SPIRV::StorageClass::Generic),
2926 MachineInstrBuilder MIB = buildSpecConstantOp(
2928 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2938bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2939 SPIRVTypeInst ResType,
2940 MachineInstr &
I)
const {
2944 Register SrcPtr =
I.getOperand(1).getReg();
2949 return BuildCOPY(ResVReg, SrcPtr,
I);
2959 unsigned SpecOpcode = [&]() ->
unsigned {
2960 if (SrcSC == SPIRV::StorageClass::CodeSectionINTEL)
2961 return static_cast<uint32_t
>(SPIRV::Opcode::Bitcast);
2963 return static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric);
2965 return static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr);
2973 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2975 .constrainAllUses(
TII,
TRI, RBI);
2977 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2979 buildSpecConstantOp(
2981 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2982 .constrainAllUses(
TII,
TRI, RBI);
2989 return BuildCOPY(ResVReg, SrcPtr,
I);
2991 if ((SrcSC == SPIRV::StorageClass::Function &&
2992 DstSC == SPIRV::StorageClass::Private) ||
2993 (DstSC == SPIRV::StorageClass::Function &&
2994 SrcSC == SPIRV::StorageClass::Private))
2995 return BuildCOPY(ResVReg, SrcPtr,
I);
2999 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
3002 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
3005 SPIRVTypeInst GenericPtrTy =
3024 return selectUnOp(ResVReg, ResType,
I,
3025 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
3027 return selectUnOp(ResVReg, ResType,
I,
3028 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
3030 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
3032 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
3042bool SPIRVInstructionSelector::selectPtrMask(
Register ResVReg,
3043 SPIRVTypeInst ResType,
3044 MachineInstr &
I)
const {
3046 return diagnoseUnsupported(
3047 I,
"G_PTRMASK is not supported with logical SPIR-V");
3052 Register PtrReg =
I.getOperand(1).getReg();
3053 Register MaskReg =
I.getOperand(2).getReg();
3072 ? SPIRV::OpBitwiseAndV
3073 : SPIRV::OpBitwiseAndS;
3096 return SPIRV::OpFOrdEqual;
3098 return SPIRV::OpFOrdGreaterThanEqual;
3100 return SPIRV::OpFOrdGreaterThan;
3102 return SPIRV::OpFOrdLessThanEqual;
3104 return SPIRV::OpFOrdLessThan;
3106 return SPIRV::OpFOrdNotEqual;
3108 return SPIRV::OpOrdered;
3110 return SPIRV::OpFUnordEqual;
3112 return SPIRV::OpFUnordGreaterThanEqual;
3114 return SPIRV::OpFUnordGreaterThan;
3116 return SPIRV::OpFUnordLessThanEqual;
3118 return SPIRV::OpFUnordLessThan;
3120 return SPIRV::OpFUnordNotEqual;
3122 return SPIRV::OpUnordered;
3132 return SPIRV::OpIEqual;
3134 return SPIRV::OpINotEqual;
3136 return SPIRV::OpSGreaterThanEqual;
3138 return SPIRV::OpSGreaterThan;
3140 return SPIRV::OpSLessThanEqual;
3142 return SPIRV::OpSLessThan;
3144 return SPIRV::OpUGreaterThanEqual;
3146 return SPIRV::OpUGreaterThan;
3148 return SPIRV::OpULessThanEqual;
3150 return SPIRV::OpULessThan;
3159 return SPIRV::OpPtrEqual;
3161 return SPIRV::OpPtrNotEqual;
3172 return SPIRV::OpLogicalEqual;
3174 return SPIRV::OpLogicalNotEqual;
3212bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
3213 SPIRVTypeInst ResType,
3215 unsigned OpAnyOrAll)
const {
3216 assert(
I.getNumOperands() == 3);
3217 assert(
I.getOperand(2).isReg());
3219 Register InputRegister =
I.getOperand(2).getReg();
3222 assert(InputType &&
"VReg has no type assigned");
3225 bool IsVectorTy = InputType->
getOpcode() == SPIRV::OpTypeVector;
3226 if (IsBoolTy && !IsVectorTy) {
3227 assert(ResVReg ==
I.getOperand(0).getReg());
3228 return BuildCOPY(ResVReg, InputRegister,
I);
3232 unsigned SpirvNotEqualId =
3233 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
3235 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
3240 IsBoolTy ? InputRegister
3248 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
3250 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
3267bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
3268 SPIRVTypeInst ResType,
3269 MachineInstr &
I)
const {
3270 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
3273bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
3274 SPIRVTypeInst ResType,
3275 MachineInstr &
I)
const {
3276 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
3280bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
3281 SPIRVTypeInst ResType,
3282 MachineInstr &
I)
const {
3283 assert(
I.getNumOperands() == 4);
3284 assert(
I.getOperand(2).isReg());
3285 assert(
I.getOperand(3).isReg());
3287 [[maybe_unused]] SPIRVTypeInst VecType =
3292 "dot product requires a vector of at least 2 components");
3294 [[maybe_unused]] SPIRVTypeInst EltType =
3303 .
addUse(
I.getOperand(2).getReg())
3304 .
addUse(
I.getOperand(3).getReg())
3309bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
3310 SPIRVTypeInst ResType,
3313 assert(
I.getNumOperands() == 4);
3314 assert(
I.getOperand(2).isReg());
3315 assert(
I.getOperand(3).isReg());
3318 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3322 .
addUse(
I.getOperand(2).getReg())
3323 .
addUse(
I.getOperand(3).getReg())
3330bool SPIRVInstructionSelector::selectIntegerDotExpansion(
3331 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3332 assert(
I.getNumOperands() == 4);
3333 assert(
I.getOperand(2).isReg());
3334 assert(
I.getOperand(3).isReg());
3338 Register Vec0 =
I.getOperand(2).getReg();
3339 Register Vec1 =
I.getOperand(3).getReg();
3343 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
3352 "dot product requires a vector of at least 2 components");
3355 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3365 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3376 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3388bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
3389 SPIRVTypeInst ResType,
3390 MachineInstr &
I)
const {
3392 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
3395 .
addUse(
I.getOperand(2).getReg())
3400bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
3401 SPIRVTypeInst ResType,
3402 MachineInstr &
I)
const {
3404 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
3407 .
addUse(
I.getOperand(2).getReg())
3412bool SPIRVInstructionSelector::selectOpIsFinite(
Register ResVReg,
3413 SPIRVTypeInst ResType,
3414 MachineInstr &
I)
const {
3416 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsFinite))
3419 .
addUse(
I.getOperand(2).getReg())
3424bool SPIRVInstructionSelector::selectOpIsNormal(
Register ResVReg,
3425 SPIRVTypeInst ResType,
3426 MachineInstr &
I)
const {
3428 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNormal))
3431 .
addUse(
I.getOperand(2).getReg())
3436template <
bool Signed>
3437bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
3438 SPIRVTypeInst ResType,
3439 MachineInstr &
I)
const {
3440 assert(
I.getNumOperands() == 5);
3441 assert(
I.getOperand(2).isReg());
3442 assert(
I.getOperand(3).isReg());
3443 assert(
I.getOperand(4).isReg());
3446 Register Acc =
I.getOperand(2).getReg();
3450 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3452 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
3457 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
3460 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3472template <
bool Signed>
3473bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
3474 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3475 assert(
I.getNumOperands() == 5);
3476 assert(
I.getOperand(2).isReg());
3477 assert(
I.getOperand(3).isReg());
3478 assert(
I.getOperand(4).isReg());
3481 Register Acc =
I.getOperand(2).getReg();
3487 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
3491 for (
unsigned i = 0; i < 4; i++) {
3514 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
3534 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3549bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
3550 SPIRVTypeInst ResType,
3551 MachineInstr &
I)
const {
3552 assert(
I.getNumOperands() == 3);
3553 assert(
I.getOperand(2).isReg());
3555 Register VZero = buildZerosValF(ResType,
I);
3556 Register VOne = buildOnesValF(ResType,
I);
3558 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
3561 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3563 .
addUse(
I.getOperand(2).getReg())
3570bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
3571 SPIRVTypeInst ResType,
3572 MachineInstr &
I)
const {
3573 assert(
I.getNumOperands() == 3);
3574 assert(
I.getOperand(2).isReg());
3576 Register InputRegister =
I.getOperand(2).getReg();
3578 auto &
DL =
I.getDebugLoc();
3581 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3588 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
3590 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
3598 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3603 if (NeedsConversion) {
3604 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
3615bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
3616 SPIRVTypeInst ResType,
3618 unsigned Opcode)
const {
3622 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
3628 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
3629 BMI.addUse(
I.getOperand(J).getReg());
3636bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
3639 bool WithGroupSync)
const {
3641 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
3643 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
3645 assert(((Scope != SPIRV::Scope::Workgroup) ||
3646 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
3647 "Workgroup Scope must set WorkGroupMemory semantic "
3648 "in Barrier instruction");
3650 assert(((Scope != SPIRV::Scope::Device) ||
3651 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
3652 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
3653 "Device Scope must set UniformMemory and ImageMemory semantic "
3654 "in Barrier instruction");
3660 if (WithGroupSync) {
3661 Register ExecReg = buildI32Constant(SPIRV::Scope::Workgroup,
I);
3665 Register ScopeReg = buildI32Constant(Scope,
I);
3666 Register MemSemReg = buildI32Constant(MemSem,
I);
3668 MI.addUse(ScopeReg).addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
3672bool SPIRVInstructionSelector::selectWaveActiveCountBits(
3673 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3678 if (!selectWaveOpInst(BallotReg, BallotType,
I,
3679 SPIRV::OpGroupNonUniformBallot))
3684 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3689 .
addImm(SPIRV::GroupOperation::Reduce)
3696bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
3697 SPIRVTypeInst ResType,
3698 MachineInstr &
I)
const {
3703 Register InputReg =
I.getOperand(2).getReg();
3708 bool IsVector = NumElems > 1;
3721 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
3722 SPIRV::OpGroupNonUniformAllEqual);
3727 ElementResults.
reserve(NumElems);
3729 for (
unsigned Idx = 0; Idx < NumElems; ++Idx) {
3742 ElemInput = Extracted;
3748 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3759 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3770bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3771 SPIRVTypeInst ResType,
3772 MachineInstr &
I)
const {
3774 assert(
I.getNumOperands() == 3);
3776 auto Op =
I.getOperand(2);
3786 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3788 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3789 return diagnoseUnsupported(
I,
"WavePrefixBitCount requires boolean input");
3810 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3814 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3821bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3822 SPIRVTypeInst ResType,
3824 bool IsUnsigned)
const {
3825 return selectWaveReduce(
3826 ResVReg, ResType,
I, IsUnsigned,
3827 [&](
Register InputRegister,
bool IsUnsigned) {
3828 const bool IsFloatTy =
3830 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3831 : SPIRV::OpGroupNonUniformSMax;
3832 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3836bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3837 SPIRVTypeInst ResType,
3839 bool IsUnsigned)
const {
3840 return selectWaveReduce(
3841 ResVReg, ResType,
I, IsUnsigned,
3842 [&](
Register InputRegister,
bool IsUnsigned) {
3843 const bool IsFloatTy =
3845 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3846 : SPIRV::OpGroupNonUniformSMin;
3847 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3851bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3852 SPIRVTypeInst ResType,
3853 MachineInstr &
I)
const {
3854 return selectWaveReduce(ResVReg, ResType,
I,
false,
3855 [&](
Register InputRegister,
bool IsUnsigned) {
3857 InputRegister, SPIRV::OpTypeFloat);
3858 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3859 : SPIRV::OpGroupNonUniformIAdd;
3863bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3864 SPIRVTypeInst ResType,
3865 MachineInstr &
I)
const {
3866 return selectWaveReduce(ResVReg, ResType,
I,
false,
3867 [&](
Register InputRegister,
bool IsUnsigned) {
3869 InputRegister, SPIRV::OpTypeFloat);
3870 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3871 : SPIRV::OpGroupNonUniformIMul;
3875template <
typename PickOpcodeFn>
3876bool SPIRVInstructionSelector::selectWaveReduce(
3877 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3878 PickOpcodeFn &&PickOpcode)
const {
3879 assert(
I.getNumOperands() == 3);
3880 assert(
I.getOperand(2).isReg());
3882 Register InputRegister =
I.getOperand(2).getReg();
3886 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3889 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3895 .
addImm(SPIRV::GroupOperation::Reduce)
3896 .
addUse(
I.getOperand(2).getReg())
3901bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3902 SPIRVTypeInst ResType,
3904 unsigned Opcode)
const {
3905 return selectWaveReduce(
3906 ResVReg, ResType,
I,
false,
3907 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3910bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3911 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3912 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3913 [&](
Register InputRegister,
bool IsUnsigned) {
3915 InputRegister, SPIRV::OpTypeFloat);
3917 ? SPIRV::OpGroupNonUniformFAdd
3918 : SPIRV::OpGroupNonUniformIAdd;
3922bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3923 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3924 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3925 [&](
Register InputRegister,
bool IsUnsigned) {
3927 InputRegister, SPIRV::OpTypeFloat);
3929 ? SPIRV::OpGroupNonUniformFMul
3930 : SPIRV::OpGroupNonUniformIMul;
3934template <
typename PickOpcodeFn>
3935bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3936 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3937 PickOpcodeFn &&PickOpcode)
const {
3938 assert(
I.getNumOperands() == 3);
3939 assert(
I.getOperand(2).isReg());
3941 Register InputRegister =
I.getOperand(2).getReg();
3945 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3948 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3954 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3955 .
addUse(
I.getOperand(2).getReg())
3960bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
3961 SPIRVTypeInst ResType,
3964 assert(
I.getNumOperands() == 3);
3965 assert(
I.getOperand(2).isReg());
3967 Register InputRegister =
I.getOperand(2).getReg();
3973 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
3984bool SPIRVInstructionSelector::selectBitreverseViaI32(
Register ResVReg,
3985 SPIRVTypeInst ResType,
3992 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
3997 : SPIRV::OpUConvert;
4001 ShiftOp = SPIRV::OpShiftRightLogicalV;
4006 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4007 TII.get(SPIRV::OpConstantComposite))
4010 for (
unsigned It = 0; It <
N; ++It)
4014 ShiftConst = CompositeReg;
4019 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
4024 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
4029 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
4034 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
4037bool SPIRVInstructionSelector::handle64BitOverflow(
4039 unsigned int Opcode,
4046 "handle64BitOverflow should only be used for integer types");
4048 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
4050 MachineIRBuilder MIRBuilder(
I);
4052 SPIRVTypeInst I64x2Type =
4054 SPIRVTypeInst Vec2ResType =
4057 std::vector<Register> PartialRegs;
4059 unsigned CurrentComponent = 0;
4060 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
4064 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4065 TII.get(SPIRV::OpVectorShuffle))
4070 .
addImm(CurrentComponent)
4071 .
addImm(CurrentComponent + 1);
4081 PartialRegs.push_back(SubVecReg);
4084 if (CurrentComponent != ComponentCount) {
4090 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
4091 SPIRV::OpVectorExtractDynamic))
4100 PartialRegs.push_back(FinalElemResReg);
4104 return selectOpWithSrcs(ResVReg, ResType,
I, PartialRegs,
4105 SPIRV::OpCompositeConstruct);
4108bool SPIRVInstructionSelector::selectBitreverse64(
Register ResVReg,
4109 SPIRVTypeInst ResType,
4113 if (ComponentCount > 2)
4114 return handle64BitOverflow(
4115 ResVReg, ResType,
I, SrcReg, SPIRV::OpBitReverse,
4117 unsigned O) {
return this->selectBitreverse64(R,
T,
I, S); });
4119 MachineIRBuilder MIRBuilder(
I);
4123 I32Type, 2 * ComponentCount, MIRBuilder,
false);
4127 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
4132 if (!selectBitreverseNative(Reverse32, VecI32Type,
I, Vec32))
4139 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4140 TII.get(SPIRV::OpVectorShuffle))
4145 for (
unsigned J = 0; J < ComponentCount; ++J) {
4152 return selectOpWithSrcs(ResVReg, ResType,
I, {SwappedVec}, SPIRV::OpBitcast);
4155bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
4156 SPIRVTypeInst ResType,
4160 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
4168bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
4169 SPIRVTypeInst ResType,
4170 MachineInstr &
I)
const {
4171 Register OpReg =
I.getOperand(1).getReg();
4180 return selectBitreverseViaI32(ResVReg, ResType,
I, OpReg);
4182 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4184 return selectBitreverse64(ResVReg, ResType,
I, OpReg);
4186 return SPIRVInstructionSelector::diagnoseUnsupported(
4187 I,
"G_BITREVERSE only support 16,32,64 bits.");
4191 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4202 unsigned AndOp = SPIRV::OpBitwiseAndS;
4203 unsigned OrOp = SPIRV::OpBitwiseOrS;
4204 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
4205 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
4207 AndOp = SPIRV::OpBitwiseAndV;
4208 OrOp = SPIRV::OpBitwiseOrV;
4209 ShlOp = SPIRV::OpShiftLeftLogicalV;
4210 ShrOp = SPIRV::OpShiftRightLogicalV;
4216 const unsigned Shift) ->
Register {
4224 Register MaskReg = CreateConst(Mask);
4225 Register ShiftReg = CreateConst(Shift);
4232 if (!selectOpWithSrcs(
T1, ResType,
I, {Input, ShiftReg}, ShrOp) ||
4233 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
4234 !selectOpWithSrcs(T3, ResType,
I, {Input, MaskReg}, AndOp) ||
4235 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
4236 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
4244 uint64_t
Mask = ~0ull;
4245 while ((Shift >>= 1) > 0) {
4252 return BuildCOPY(ResVReg, Result,
I);
4255bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
4256 SPIRVTypeInst ResType,
4257 MachineInstr &
I)
const {
4258 assert(
I.getOperand(0).isReg() &&
I.getOperand(1).isReg() &&
4259 "G_FREEZE must define and use a register");
4260 Register OpReg =
I.getOperand(1).getReg();
4264 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
4277 if (MachineInstr *Def = MRI->
getVRegDef(OpReg)) {
4278 if (
Def->getOpcode() == TargetOpcode::COPY)
4281 switch (
Def->getOpcode()) {
4282 case SPIRV::ASSIGN_TYPE:
4283 if (MachineInstr *AssignToDef =
4285 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
4286 Reg =
Def->getOperand(2).getReg();
4289 case SPIRV::OpUndef:
4290 Reg =
Def->getOperand(1).getReg();
4293 unsigned DestOpCode;
4295 DestOpCode = SPIRV::OpConstantNull;
4296 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze of a "
4297 "static undef/poison lowered to OpConstantNull\n");
4299 DestOpCode = TargetOpcode::COPY;
4301 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze "
4302 "skipped, lowered as a copy of the operand\n");
4304 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
4305 .
addDef(
I.getOperand(0).getReg())
4313bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
4314 SPIRVTypeInst ResType,
4315 MachineInstr &
I)
const {
4317 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4319 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4323 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
4328 for (
unsigned i =
I.getNumExplicitDefs();
4329 i <
I.getNumExplicitOperands() && IsConst; ++i)
4333 if (!IsConst &&
N < 2)
4334 return diagnoseUnsupported(
4335 I,
"There must be at least two constituent operands in a vector");
4340 for (
unsigned i =
I.getNumExplicitDefs();
4341 i <
I.getNumExplicitOperands() && IsNullVector; ++i) {
4342 MachineInstr *
Def =
getDef(
I.getOperand(i), MRI);
4347 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4354 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4355 TII.get(IsConst ? SPIRV::OpConstantComposite
4356 : SPIRV::OpCompositeConstruct))
4359 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
4360 MIB.
addUse(
I.getOperand(i).getReg());
4365bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
4366 SPIRVTypeInst ResType,
4367 MachineInstr &
I)
const {
4369 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4371 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4376 unsigned OpIdx =
I.getNumExplicitDefs();
4377 if (!
I.getOperand(OpIdx).isReg())
4381 Register OpReg =
I.getOperand(OpIdx).getReg();
4384 if (!IsConst &&
N < 2)
4385 return diagnoseUnsupported(
4386 I,
"There must be at least two constituent operands in a vector");
4389 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4390 TII.get(IsConst ? SPIRV::OpConstantComposite
4391 : SPIRV::OpCompositeConstruct))
4394 for (
unsigned i = 0; i <
N; ++i)
4400bool SPIRVInstructionSelector::selectConcatVectors(
Register ResVReg,
4401 SPIRVTypeInst ResType,
4402 MachineInstr &
I)
const {
4406 if (ResType->
getOpcode() != SPIRV::OpTypeVector)
4408 "Cannot select G_CONCAT_VECTORS with a non-vector result");
4410 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4411 TII.get(SPIRV::OpCompositeConstruct))
4414 for (
unsigned OpIdx =
I.getNumExplicitDefs();
4415 OpIdx <
I.getNumExplicitOperands(); ++OpIdx)
4416 MIB.
addUse(
I.getOperand(OpIdx).getReg());
4421bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
4422 SPIRVTypeInst ResType,
4423 MachineInstr &
I)
const {
4428 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
4430 Opcode = SPIRV::OpDemoteToHelperInvocation;
4432 Opcode = SPIRV::OpKill;
4434 if (MachineInstr *NextI =
I.getNextNode()) {
4436 NextI->eraseFromParent();
4446bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
4447 SPIRVTypeInst ResType,
unsigned CmpOpc,
4448 MachineInstr &
I)
const {
4449 Register Cmp0 =
I.getOperand(2).getReg();
4450 Register Cmp1 =
I.getOperand(3).getReg();
4453 "CMP operands should have the same type");
4454 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
4464bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
4465 SPIRVTypeInst ResType,
4466 MachineInstr &
I)
const {
4467 auto Pred =
I.getOperand(1).getPredicate();
4470 Register CmpOperand =
I.getOperand(2).getReg();
4472 bool IsPtrCmp = CmpOperandType && CmpOperandType.
isPointer();
4477 Register Op1 =
I.getOperand(3).getReg();
4481 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
4486 I.getOperand(3).setReg(NewOp1);
4492 return selectCmp(ResVReg, ResType, CmpOpc,
I);
4496SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
4497 SPIRVTypeInst ResType)
const {
4499 SPIRVTypeInst SpvI32Ty =
4502 auto ConstInt = ConstantInt::get(LLVMTy, Val);
4509 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4512 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
4515 .
addImm(APInt(32, Val).getZExtValue());
4517 GR.
add(ConstInt,
MI);
4524Register SPIRVInstructionSelector::buildI32ConstantInEntryBlock(
4525 uint32_t Val, MachineInstr &
I, SPIRVTypeInst ResType)
const {
4527 SPIRVTypeInst SpvI32Ty =
4529 auto *ConstInt = ConstantInt::get(LLVMTy, Val);
4534 MachineBasicBlock &EntryBB = *InsertIt->getParent();
4535 MachineInstr *
MI =
nullptr;
4539 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantNull))
4543 uint64_t ImmVal = APInt(32, Val).getZExtValue();
4544 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantI))
4550 GR.
add(ConstInt,
MI);
4555bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
4556 SPIRVTypeInst ResType,
4557 MachineInstr &
I)
const {
4559 return selectCmp(ResVReg, ResType, CmpOp,
I);
4562bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
4563 SPIRVTypeInst ResType,
4564 MachineInstr &
I)
const {
4566 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
4573 if (ResType->
getOpcode() != SPIRV::OpTypeVector &&
4574 ResType->
getOpcode() != SPIRV::OpTypeFloat)
4577 MachineIRBuilder MIRBuilder(
I);
4584 APFloat ConstVal(3.3219280948873623);
4588 APFloat::rmNearestTiesToEven, &LosesInfo);
4592 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
4593 ? SPIRV::OpVectorTimesScalar
4596 if (!selectOpWithSrcs(ArgReg, ResType,
I,
4597 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
4599 if (!selectExtInst(ResVReg, ResType,
I,
4600 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
4610Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
4611 MachineInstr &
I)
const {
4614 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4619bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
4625 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
4633 if (
Def->getOpcode() == SPIRV::OpConstantNull)
4636 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
4637 Def->getOpcode() == SPIRV::OpConstantI)
4650 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
4651 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
4653 Intrinsic::spv_const_composite)) {
4654 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
4655 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
4656 if (!IsZero(
Def->getOperand(i).getReg()))
4665Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
4666 MachineInstr &
I)
const {
4670 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4675Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
4676 MachineInstr &
I)
const {
4680 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4686 SPIRVTypeInst ResType,
4687 MachineInstr &
I)
const {
4691 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4696bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
4697 SPIRVTypeInst ResType,
4698 MachineInstr &
I)
const {
4699 Register SelectFirstArg =
I.getOperand(2).getReg();
4700 Register SelectSecondArg =
I.getOperand(3).getReg();
4709 SPIRV::OpTypeVector;
4716 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
4717 }
else if (IsPtrTy) {
4718 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
4720 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
4723 assert(IsScalarBool &&
"OpSelect with a scalar result requires a scalar "
4724 "boolean condition");
4726 Opcode = SPIRV::OpSelectSFSCond;
4727 }
else if (IsPtrTy) {
4728 Opcode = SPIRV::OpSelectSPSCond;
4730 Opcode = SPIRV::OpSelectSISCond;
4733 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4736 .
addUse(
I.getOperand(1).getReg())
4745bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
4746 SPIRVTypeInst ResType,
4748 MachineInstr &InsertAt,
4749 bool IsSigned)
const {
4751 Register ZeroReg = buildZerosVal(ResType, InsertAt);
4752 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
4753 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
4755 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4767bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
4768 SPIRVTypeInst ResType,
4769 MachineInstr &
I,
bool IsSigned,
4770 unsigned Opcode)
const {
4771 Register SrcReg =
I.getOperand(1).getReg();
4777 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
4782 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I, IsSigned);
4784 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
4787bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
4788 SPIRVTypeInst ResType, MachineInstr &
I,
4789 bool IsSigned)
const {
4790 Register SrcReg =
I.getOperand(1).getReg();
4792 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
4796 if (ResType == SrcType)
4797 return BuildCOPY(ResVReg, SrcReg,
I);
4799 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4800 return selectUnOp(ResVReg, ResType,
I, Opcode);
4803bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
4804 SPIRVTypeInst ResType,
4806 bool IsSigned)
const {
4807 MachineIRBuilder MIRBuilder(
I);
4808 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
4820 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
4823 .
addUse(
I.getOperand(1).getReg())
4824 .
addUse(
I.getOperand(2).getReg())
4829 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
4832 .
addUse(
I.getOperand(1).getReg())
4833 .
addUse(
I.getOperand(2).getReg())
4841 unsigned SelectOpcode =
4842 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
4847 .
addUse(buildOnesVal(
true, ResType,
I))
4848 .
addUse(buildZerosVal(ResType,
I))
4855 .
addUse(buildOnesVal(
false, ResType,
I))
4860bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
4863 SPIRVTypeInst IntTy,
4864 SPIRVTypeInst BoolTy)
const {
4867 bool IsVectorTy = IntTy->
getOpcode() == SPIRV::OpTypeVector;
4868 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
4870 Register One = buildOnesVal(
false, IntTy,
I);
4878 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
4887bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
4888 SPIRVTypeInst ResType,
4889 MachineInstr &
I)
const {
4890 Register IntReg =
I.getOperand(1).getReg();
4893 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
4894 if (ArgType == ResType)
4895 return BuildCOPY(ResVReg, IntReg,
I);
4897 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4898 return selectUnOp(ResVReg, ResType,
I, Opcode);
4901bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
4902 SPIRVTypeInst ResType,
4903 MachineInstr &
I)
const {
4904 unsigned Opcode =
I.getOpcode();
4905 unsigned TpOpcode = ResType->
getOpcode();
4907 if (ResType.
isPointer() || TpOpcode == SPIRV::OpTypeEvent) {
4908 assert(Opcode == TargetOpcode::G_CONSTANT &&
4909 I.getOperand(1).getCImm()->isZero());
4910 MachineBasicBlock &DepMBB =
I.getMF()->front();
4913 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
4920 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
4923bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
4924 SPIRVTypeInst ResType,
4925 MachineInstr &
I)
const {
4926 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4933bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
4934 SPIRVTypeInst ResType,
4935 MachineInstr &
I)
const {
4937 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
4941 .
addUse(
I.getOperand(3).getReg())
4943 .
addUse(
I.getOperand(2).getReg());
4944 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
4950bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
4951 SPIRVTypeInst ResType,
4952 MachineInstr &
I)
const {
4953 Type *MaybeResTy =
nullptr;
4958 "Expected aggregate type for extractv instruction");
4960 SPIRV::AccessQualifier::ReadWrite,
false);
4964 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
4967 .
addUse(
I.getOperand(2).getReg());
4968 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
4974bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
4975 SPIRVTypeInst ResType,
4976 MachineInstr &
I)
const {
4977 if (
getImm(
I.getOperand(4), MRI))
4978 return selectInsertVal(ResVReg, ResType,
I);
4980 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
4983 .
addUse(
I.getOperand(2).getReg())
4984 .
addUse(
I.getOperand(3).getReg())
4985 .
addUse(
I.getOperand(4).getReg())
4990bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
4991 SPIRVTypeInst ResType,
4992 MachineInstr &
I)
const {
4993 if (
getImm(
I.getOperand(3), MRI))
4994 return selectExtractVal(ResVReg, ResType,
I);
4996 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
4999 .
addUse(
I.getOperand(2).getReg())
5000 .
addUse(
I.getOperand(3).getReg())
5005bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
5006 SPIRVTypeInst ResType,
5007 MachineInstr &
I)
const {
5008 const bool IsGEPInBounds =
I.getOperand(2).getImm();
5011 const bool UseUntypedPointers =
5012 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
5017 if (UseUntypedPointers) {
5019 Opcode = IsGEPInBounds ? SPIRV::OpUntypedInBoundsAccessChainKHR
5020 : SPIRV::OpUntypedAccessChainKHR;
5022 Opcode = IsGEPInBounds ? SPIRV::OpUntypedInBoundsPtrAccessChainKHR
5023 : SPIRV::OpUntypedPtrAccessChainKHR;
5032 IsGEPInBounds ? SPIRV::OpInBoundsAccessChain : SPIRV::OpAccessChain;
5034 Opcode = IsGEPInBounds ? SPIRV::OpInBoundsPtrAccessChain
5035 : SPIRV::OpPtrAccessChain;
5040 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5045 if (UseUntypedPointers) {
5060 while (Def &&
Def->getOpcode() == TargetOpcode::COPY &&
5061 Def->getOperand(1).isReg())
5063 if (Def &&
Def->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)
5064 if (
const auto *GVar =
5067 SPIRV::AccessQualifier::ReadWrite,
5071 return diagnoseUnsupported(
5072 I,
"could not deduce the base type of an untyped access chain");
5077 Res.addUse(BaseReg);
5079 const bool IsAccessChainOpcode =
5080 (Opcode == SPIRV::OpAccessChain ||
5081 Opcode == SPIRV::OpInBoundsAccessChain ||
5082 Opcode == SPIRV::OpUntypedAccessChainKHR ||
5083 Opcode == SPIRV::OpUntypedInBoundsAccessChainKHR);
5085 assert((!IsAccessChainOpcode || (
getImm(
I.getOperand(4), MRI) &&
5086 foldImm(
I.getOperand(4), MRI) == 0)) &&
5087 "Cannot translate GEP to OpAccessChain.");
5090 const unsigned StartingIndex = IsAccessChainOpcode ? 5 : 4;
5091 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
5092 Res.addUse(
I.getOperand(i).getReg());
5093 Res.constrainAllUses(
TII,
TRI, RBI);
5098bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
5100 unsigned Lim =
I.getNumExplicitOperands();
5101 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
5102 Register OpReg =
I.getOperand(i).getReg();
5103 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
5105 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
5106 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
5107 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
5120 SPIRVTypeInst WrapType = OpType;
5121 if (OpType->
getOpcode() == SPIRV::OpTypePointer &&
5123 SPIRV::StorageClass::CodeSectionINTEL) {
5125 SPIRV::StorageClass::Function,
I);
5132 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5133 TII.get(SPIRV::OpSpecConstantOp))
5136 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
5138 GR.
add(OpDefine, MIB);
5144bool SPIRVInstructionSelector::selectDerivativeInst(
5145 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
5146 const unsigned DPdOpCode)
const {
5149 if (!errorIfInstrOutsideShader(
I))
5155 Register SrcReg =
I.getOperand(2).getReg();
5160 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5163 .
addUse(
I.getOperand(2).getReg());
5165 MachineIRBuilder MIRBuilder(
I);
5168 if (componentCount != 1)
5176 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5181 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5186 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5194bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
5195 SPIRVTypeInst ResType,
5196 MachineInstr &
I)
const {
5200 case Intrinsic::spv_load:
5201 return selectLoad(ResVReg, ResType,
I);
5202 case Intrinsic::spv_atomic_load:
5203 return selectAtomicLoad(ResVReg, ResType,
I);
5204 case Intrinsic::spv_store:
5205 return selectStore(
I);
5206 case Intrinsic::spv_atomic_store:
5207 return selectAtomicStore(
I);
5208 case Intrinsic::spv_extractv:
5209 return selectExtractVal(ResVReg, ResType,
I);
5210 case Intrinsic::spv_insertv:
5211 return selectInsertVal(ResVReg, ResType,
I);
5212 case Intrinsic::spv_extractelt:
5213 return selectExtractElt(ResVReg, ResType,
I);
5214 case Intrinsic::spv_insertelt:
5215 return selectInsertElt(ResVReg, ResType,
I);
5216 case Intrinsic::spv_gep:
5217 return selectGEP(ResVReg, ResType,
I);
5218 case Intrinsic::spv_bitcast: {
5219 Register OpReg =
I.getOperand(2).getReg();
5220 SPIRVTypeInst OpType =
5224 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
5226 case Intrinsic::spv_unref_global:
5227 case Intrinsic::spv_init_global: {
5228 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
5233 Register GVarVReg =
MI->getOperand(0).getReg();
5234 if (!selectGlobalValue(GVarVReg, *
MI, Init))
5239 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
5241 MI->eraseFromParent();
5245 case Intrinsic::spv_undef: {
5246 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5252 case Intrinsic::spv_poison:
5253 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpPoisonKHR))
5258 case Intrinsic::spv_freeze:
5259 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
5262 .
addUse(
I.getOperand(2).getReg())
5265 case Intrinsic::spv_named_boolean_spec_constant: {
5266 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
5267 : SPIRV::OpSpecConstantFalse;
5269 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
5270 .
addDef(
I.getOperand(0).getReg())
5273 unsigned SpecId =
I.getOperand(2).getImm();
5275 SPIRV::Decoration::SpecId, {SpecId});
5279 case Intrinsic::spv_const_composite: {
5281 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
5287 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
5289 std::function<bool(
Register)> HasSpecConstOperand =
5299 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
5300 J < Def->getNumExplicitOperands(); ++J) {
5301 if (
Def->getOperand(J).isReg() &&
5302 HasSpecConstOperand(
Def->getOperand(J).getReg()))
5308 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
5309 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
5310 : SPIRV::OpConstantComposite;
5311 unsigned ContinuedOpc = HasSpecConst
5312 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
5313 : SPIRV::OpConstantCompositeContinuedINTEL;
5314 MachineIRBuilder MIR(
I);
5316 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
5318 for (
auto *Instr : Instructions) {
5319 Instr->setDebugLoc(
I.getDebugLoc());
5324 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5331 case Intrinsic::spv_assign_name: {
5332 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
5333 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
5334 for (
unsigned i =
I.getNumExplicitDefs() + 2;
5335 i <
I.getNumExplicitOperands(); ++i) {
5336 MIB.
addImm(
I.getOperand(i).getImm());
5341 case Intrinsic::spv_switch: {
5342 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
5343 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5344 if (
I.getOperand(i).isReg())
5345 MIB.
addReg(
I.getOperand(i).getReg());
5346 else if (
I.getOperand(i).isCImm())
5347 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
5348 else if (
I.getOperand(i).isMBB())
5349 MIB.
addMBB(
I.getOperand(i).getMBB());
5356 case Intrinsic::spv_loop_merge: {
5357 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
5358 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5359 if (
I.getOperand(i).isMBB())
5360 MIB.
addMBB(
I.getOperand(i).getMBB());
5367 case Intrinsic::spv_loop_control_intel: {
5369 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
5370 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
5375 case Intrinsic::spv_selection_merge: {
5377 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
5378 assert(
I.getOperand(1).isMBB() &&
5379 "operand 1 to spv_selection_merge must be a basic block");
5380 MIB.
addMBB(
I.getOperand(1).getMBB());
5381 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
5385 case Intrinsic::spv_cmpxchg:
5386 return selectAtomicCmpXchg(ResVReg, ResType,
I);
5387 case Intrinsic::spv_unreachable:
5388 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
5391 case Intrinsic::spv_abort:
5392 return selectAbort(
I);
5393 case Intrinsic::spv_alloca:
5394 return selectFrameIndex(ResVReg, ResType,
I);
5395 case Intrinsic::spv_alloca_array:
5396 return selectAllocaArray(ResVReg, ResType,
I);
5397 case Intrinsic::spv_assume:
5399 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
5400 .
addUse(
I.getOperand(1).getReg())
5405 case Intrinsic::spv_expect:
5407 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
5410 .
addUse(
I.getOperand(2).getReg())
5411 .
addUse(
I.getOperand(3).getReg())
5416 case Intrinsic::arithmetic_fence:
5417 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
5418 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
5421 .
addUse(
I.getOperand(2).getReg())
5425 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
5427 case Intrinsic::spv_thread_id:
5433 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
5435 case Intrinsic::spv_thread_id_in_group:
5441 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
5443 case Intrinsic::spv_group_id:
5449 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
5451 case Intrinsic::spv_flattened_thread_id_in_group:
5458 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
5460 case Intrinsic::spv_workgroup_size:
5461 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
5463 case Intrinsic::spv_global_size:
5464 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
5466 case Intrinsic::spv_global_offset:
5467 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
5469 case Intrinsic::spv_num_workgroups:
5470 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
5472 case Intrinsic::spv_subgroup_size:
5473 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
5475 case Intrinsic::spv_num_subgroups:
5476 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
5478 case Intrinsic::spv_subgroup_id:
5479 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
5480 case Intrinsic::spv_subgroup_local_invocation_id:
5481 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
5482 ResVReg, ResType,
I);
5483 case Intrinsic::spv_subgroup_max_size:
5484 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
5486 case Intrinsic::spv_fdot:
5487 return selectFloatDot(ResVReg, ResType,
I);
5488 case Intrinsic::spv_udot:
5489 case Intrinsic::spv_sdot:
5490 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5492 return selectIntegerDot(ResVReg, ResType,
I,
5493 IID == Intrinsic::spv_sdot);
5494 return selectIntegerDotExpansion(ResVReg, ResType,
I);
5495 case Intrinsic::spv_dot4add_i8packed:
5496 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5498 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
5499 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
5500 case Intrinsic::spv_dot4add_u8packed:
5501 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5503 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
5504 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
5505 case Intrinsic::spv_all:
5506 return selectAll(ResVReg, ResType,
I);
5507 case Intrinsic::spv_any:
5508 return selectAny(ResVReg, ResType,
I);
5509 case Intrinsic::spv_distance:
5510 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
5511 case Intrinsic::spv_lerp:
5512 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
5513 case Intrinsic::spv_length:
5514 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
5515 case Intrinsic::spv_degrees:
5516 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
5517 case Intrinsic::spv_faceforward:
5518 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
5519 case Intrinsic::spv_frac:
5520 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
5521 case Intrinsic::spv_isinf:
5522 return selectOpIsInf(ResVReg, ResType,
I);
5523 case Intrinsic::spv_isnan:
5524 return selectOpIsNan(ResVReg, ResType,
I);
5525 case Intrinsic::spv_isfinite:
5526 return selectOpIsFinite(ResVReg, ResType,
I);
5527 case Intrinsic::spv_isnormal:
5528 return selectOpIsNormal(ResVReg, ResType,
I);
5529 case Intrinsic::spv_normalize:
5530 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
5531 case Intrinsic::spv_refract:
5532 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
5533 case Intrinsic::spv_reflect:
5534 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
5535 case Intrinsic::spv_rsqrt:
5536 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
5537 case Intrinsic::spv_sign:
5538 return selectSign(ResVReg, ResType,
I);
5539 case Intrinsic::spv_smoothstep:
5540 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
5541 case Intrinsic::spv_firstbituhigh:
5542 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
5543 case Intrinsic::spv_firstbitshigh:
5544 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
5545 case Intrinsic::spv_firstbitlow:
5546 return selectFirstBitLow(ResVReg, ResType,
I);
5547 case Intrinsic::spv_all_memory_barrier:
5548 return selectBarrierInst(
I, SPIRV::Scope::Device,
5549 SPIRV::MemorySemantics::UniformMemory |
5550 SPIRV::MemorySemantics::ImageMemory |
5551 SPIRV::MemorySemantics::WorkgroupMemory,
5553 case Intrinsic::spv_all_memory_barrier_with_group_sync:
5554 return selectBarrierInst(
I, SPIRV::Scope::Device,
5555 SPIRV::MemorySemantics::UniformMemory |
5556 SPIRV::MemorySemantics::ImageMemory |
5557 SPIRV::MemorySemantics::WorkgroupMemory,
5559 case Intrinsic::spv_device_memory_barrier:
5560 return selectBarrierInst(
I, SPIRV::Scope::Device,
5561 SPIRV::MemorySemantics::UniformMemory |
5562 SPIRV::MemorySemantics::ImageMemory,
5564 case Intrinsic::spv_device_memory_barrier_with_group_sync:
5565 return selectBarrierInst(
I, SPIRV::Scope::Device,
5566 SPIRV::MemorySemantics::UniformMemory |
5567 SPIRV::MemorySemantics::ImageMemory,
5569 case Intrinsic::spv_group_memory_barrier:
5570 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5571 SPIRV::MemorySemantics::WorkgroupMemory,
5573 case Intrinsic::spv_group_memory_barrier_with_group_sync:
5574 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5575 SPIRV::MemorySemantics::WorkgroupMemory,
5577 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
5578 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
5579 SPIRV::StorageClass::StorageClass ResSC =
5582 return diagnoseUnsupported(
I,
"The target storage class is not castable "
5583 "from the Generic storage class");
5584 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
5592 case Intrinsic::spv_lifetime_start:
5593 case Intrinsic::spv_lifetime_end: {
5594 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
5595 : SPIRV::OpLifetimeStop;
5596 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
5597 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
5606 case Intrinsic::spv_saturate:
5607 return selectSaturate(ResVReg, ResType,
I);
5608 case Intrinsic::spv_nclamp:
5609 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
5610 case Intrinsic::spv_uclamp:
5611 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
5612 case Intrinsic::spv_sclamp:
5613 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
5614 case Intrinsic::spv_subgroup_prefix_bit_count:
5615 return selectWavePrefixBitCount(ResVReg, ResType,
I);
5616 case Intrinsic::spv_wave_active_countbits:
5617 return selectWaveActiveCountBits(ResVReg, ResType,
I);
5618 case Intrinsic::spv_wave_all_equal:
5619 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
5620 case Intrinsic::spv_wave_all:
5621 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
5622 case Intrinsic::spv_wave_any:
5623 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
5624 case Intrinsic::spv_subgroup_ballot:
5625 return selectWaveOpInst(ResVReg, ResType,
I,
5626 SPIRV::OpGroupNonUniformBallot);
5627 case Intrinsic::spv_wave_is_first_lane:
5628 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
5629 case Intrinsic::spv_wave_reduce_or:
5630 return selectWaveReduceOp(ResVReg, ResType,
I,
5631 SPIRV::OpGroupNonUniformBitwiseOr);
5632 case Intrinsic::spv_wave_reduce_xor:
5633 return selectWaveReduceOp(ResVReg, ResType,
I,
5634 SPIRV::OpGroupNonUniformBitwiseXor);
5635 case Intrinsic::spv_wave_reduce_and:
5636 return selectWaveReduceOp(ResVReg, ResType,
I,
5637 SPIRV::OpGroupNonUniformBitwiseAnd);
5638 case Intrinsic::spv_wave_reduce_umax:
5639 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
5640 case Intrinsic::spv_wave_reduce_max:
5641 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
5642 case Intrinsic::spv_wave_reduce_umin:
5643 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
5644 case Intrinsic::spv_wave_reduce_min:
5645 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
5646 case Intrinsic::spv_wave_reduce_sum:
5647 return selectWaveReduceSum(ResVReg, ResType,
I);
5648 case Intrinsic::spv_wave_product:
5649 return selectWaveReduceProduct(ResVReg, ResType,
I);
5650 case Intrinsic::spv_wave_readlane:
5651 return selectWaveOpInst(ResVReg, ResType,
I,
5652 SPIRV::OpGroupNonUniformShuffle);
5653 case Intrinsic::spv_wave_prefix_sum:
5654 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
5655 case Intrinsic::spv_wave_prefix_product:
5656 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
5657 case Intrinsic::spv_quad_read_across_x: {
5658 return selectQuadSwap(ResVReg, ResType,
I, 0);
5660 case Intrinsic::spv_quad_read_across_y: {
5661 return selectQuadSwap(ResVReg, ResType,
I, 1);
5663 case Intrinsic::spv_quad_read_across_diagonal: {
5664 return selectQuadSwap(ResVReg, ResType,
I, 2);
5666 case Intrinsic::spv_radians:
5667 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
5671 case Intrinsic::instrprof_increment:
5672 case Intrinsic::instrprof_increment_step:
5673 case Intrinsic::instrprof_value_profile:
5676 case Intrinsic::spv_value_md:
5678 case Intrinsic::spv_resource_handlefrombinding: {
5679 return selectHandleFromBinding(ResVReg, ResType,
I);
5681 case Intrinsic::spv_resource_counterhandlefrombinding:
5682 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
5683 case Intrinsic::spv_resource_updatecounter:
5684 return selectUpdateCounter(ResVReg, ResType,
I);
5685 case Intrinsic::spv_resource_store_typedbuffer: {
5686 return selectImageWriteIntrinsic(
I);
5688 case Intrinsic::spv_resource_load_typedbuffer: {
5689 return selectReadImageIntrinsic(ResVReg, ResType,
I);
5691 case Intrinsic::spv_resource_load_level: {
5692 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
5694 case Intrinsic::spv_resource_getdimensions_x:
5695 case Intrinsic::spv_resource_getdimensions_xy:
5696 case Intrinsic::spv_resource_getdimensions_xyz: {
5697 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
5699 case Intrinsic::spv_resource_getdimensions_levels_x:
5700 case Intrinsic::spv_resource_getdimensions_levels_xy:
5701 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
5702 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
5704 case Intrinsic::spv_resource_getdimensions_ms_xy:
5705 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
5706 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
5708 case Intrinsic::spv_resource_calculate_lod:
5709 case Intrinsic::spv_resource_calculate_lod_unclamped:
5710 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
5711 case Intrinsic::spv_resource_sample:
5712 case Intrinsic::spv_resource_sample_clamp:
5713 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
5714 case Intrinsic::spv_resource_samplebias:
5715 case Intrinsic::spv_resource_samplebias_clamp:
5716 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
5717 case Intrinsic::spv_resource_samplegrad:
5718 case Intrinsic::spv_resource_samplegrad_clamp:
5719 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
5720 case Intrinsic::spv_resource_samplelevel:
5721 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
5722 case Intrinsic::spv_resource_samplecmp:
5723 case Intrinsic::spv_resource_samplecmp_clamp:
5724 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
5725 case Intrinsic::spv_resource_samplecmplevelzero:
5726 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
5727 case Intrinsic::spv_resource_gather:
5728 case Intrinsic::spv_resource_gather_cmp:
5729 return selectGatherIntrinsic(ResVReg, ResType,
I);
5730 case Intrinsic::spv_resource_getbasepointer:
5731 case Intrinsic::spv_resource_getpointer: {
5732 return selectResourceGetPointer(ResVReg, ResType,
I);
5734 case Intrinsic::spv_pushconstant_getpointer: {
5735 return selectPushConstantGetPointer(ResVReg, ResType,
I);
5737 case Intrinsic::spv_discard: {
5738 return selectDiscard(ResVReg, ResType,
I);
5740 case Intrinsic::spv_resource_nonuniformindex: {
5741 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
5743 case Intrinsic::spv_unpackhalf2x16: {
5744 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
5746 case Intrinsic::spv_packhalf2x16: {
5747 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
5749 case Intrinsic::spv_ddx:
5750 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
5751 case Intrinsic::spv_ddy:
5752 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
5753 case Intrinsic::spv_ddx_coarse:
5754 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
5755 case Intrinsic::spv_ddy_coarse:
5756 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
5757 case Intrinsic::spv_ddx_fine:
5758 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
5759 case Intrinsic::spv_ddy_fine:
5760 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
5761 case Intrinsic::spv_fwidth:
5762 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
5763 case Intrinsic::spv_masked_gather:
5764 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5765 return selectMaskedGather(ResVReg, ResType,
I);
5766 return diagnoseUnsupported(
5767 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
5768 case Intrinsic::spv_masked_scatter:
5769 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5770 return selectMaskedScatter(
I);
5771 return diagnoseUnsupported(
5772 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
5773 case Intrinsic::returnaddress:
5774 case Intrinsic::frameaddress: {
5776 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5783 return diagnoseUnsupported(
I,
"intrinsic selection not implemented.");
5788bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
5789 SPIRVTypeInst ResType,
5790 MachineInstr &
I)
const {
5793 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
5800bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
5801 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5803 assert(Intr.getIntrinsicID() ==
5804 Intrinsic::spv_resource_counterhandlefrombinding);
5807 Register MainHandleReg = Intr.getOperand(2).getReg();
5809 assert(MainHandleDef->getIntrinsicID() ==
5810 Intrinsic::spv_resource_handlefrombinding);
5814 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
5815 Register IndexReg = MainHandleDef->getOperand(5).getReg();
5816 std::string CounterName =
5821 MachineIRBuilder MIRBuilder(
I);
5823 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
5825 ArraySize, IndexReg, CounterName, MIRBuilder);
5827 return BuildCOPY(ResVReg, CounterVarReg,
I);
5830bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
5831 SPIRVTypeInst ResType,
5832 MachineInstr &
I)
const {
5834 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
5836 Register CounterHandleReg = Intr.getOperand(2).getReg();
5837 Register IncrReg = Intr.getOperand(3).getReg();
5844 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
5845 assert(CounterVarPointeeType &&
5846 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
5847 "Counter variable must be a struct");
5849 SPIRV::StorageClass::StorageBuffer &&
5850 "Counter variable must be in the storage buffer storage class");
5852 "Counter variable must have exactly 1 member in the struct");
5853 const SPIRVTypeInst MemberType =
5856 "Counter variable struct must have a single i32 member");
5860 MachineIRBuilder MIRBuilder(
I);
5862 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
5865 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
5871 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5874 .
addUse(CounterHandleReg)
5881 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
5884 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
5887 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
5896 return BuildCOPY(ResVReg, AtomicRes,
I);
5904 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
5912bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
5913 SPIRVTypeInst ResType,
5914 MachineInstr &
I)
const {
5922 Register ImageReg =
I.getOperand(2).getReg();
5930 Register IdxReg =
I.getOperand(3).getReg();
5932 MachineInstr &Pos =
I;
5934 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
5938bool SPIRVInstructionSelector::generateSampleImage(
5941 DebugLoc Loc, MachineInstr &Pos)
const {
5952 if (!loadHandleBeforePosition(NewSamplerReg,
5958 MachineIRBuilder MIRBuilder(Pos);
5971 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
5972 ImOps.Lod.has_value();
5973 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
5974 : SPIRV::OpImageSampleImplicitLod;
5976 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
5977 : SPIRV::OpImageSampleDrefImplicitLod;
5986 MIB.
addUse(*ImOps.Compare);
5988 uint32_t ImageOperands = 0;
5990 ImageOperands |= SPIRV::ImageOperand::Bias;
5992 ImageOperands |= SPIRV::ImageOperand::Lod;
5993 if (ImOps.GradX && ImOps.GradY)
5994 ImageOperands |= SPIRV::ImageOperand::Grad;
5995 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
5997 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6000 "Non-constant offsets are not supported in sample instructions.");
6005 ImageOperands |= SPIRV::ImageOperand::MinLod;
6007 if (ImageOperands != 0) {
6008 MIB.
addImm(ImageOperands);
6009 if (ImageOperands & SPIRV::ImageOperand::Bias)
6011 if (ImageOperands & SPIRV::ImageOperand::Lod)
6013 if (ImageOperands & SPIRV::ImageOperand::Grad) {
6014 MIB.
addUse(*ImOps.GradX);
6015 MIB.
addUse(*ImOps.GradY);
6018 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6019 MIB.
addUse(*ImOps.Offset);
6020 if (ImageOperands & SPIRV::ImageOperand::MinLod)
6021 MIB.
addUse(*ImOps.MinLod);
6028bool SPIRVInstructionSelector::selectImageQuerySize(
6030 std::optional<Register> LodReg)
const {
6032 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
6035 "ImageReg is not an image type.");
6037 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6039 unsigned NumComponents = 0;
6041 case SPIRV::Dim::DIM_1D:
6042 case SPIRV::Dim::DIM_Buffer:
6043 NumComponents =
IsArray ? 2 : 1;
6045 case SPIRV::Dim::DIM_2D:
6046 case SPIRV::Dim::DIM_Cube:
6047 case SPIRV::Dim::DIM_Rect:
6048 NumComponents =
IsArray ? 3 : 2;
6050 case SPIRV::Dim::DIM_3D:
6054 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
6059 SPIRVTypeInst ResType =
6064 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6074bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
6075 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6076 Register ImageReg =
I.getOperand(2).getReg();
6083 return selectImageQuerySize(NewImageReg, ResVReg,
I);
6086bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
6087 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6088 Register ImageReg =
I.getOperand(2).getReg();
6097 Register LodReg =
I.getOperand(3).getReg();
6100 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
6102 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
6109 TII.get(SPIRV::OpImageQueryLevels))
6116 TII.get(SPIRV::OpCompositeConstruct))
6126bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
6127 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6128 Register ImageReg =
I.getOperand(2).getReg();
6139 "OpImageQuerySamples requires a multisampled image");
6141 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
6149 TII.get(SPIRV::OpImageQuerySamples))
6156 TII.get(SPIRV::OpCompositeConstruct))
6166bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
6167 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6168 Register ImageReg =
I.getOperand(2).getReg();
6169 Register SamplerReg =
I.getOperand(3).getReg();
6170 Register CoordinateReg =
I.getOperand(4).getReg();
6186 if (!loadHandleBeforePosition(
6191 MachineIRBuilder MIRBuilder(
I);
6197 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6207 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
6214 unsigned ExtractedIndex =
6216 Intrinsic::spv_resource_calculate_lod_unclamped
6220 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6221 TII.get(SPIRV::OpCompositeExtract))
6231bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
6232 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6233 Register ImageReg =
I.getOperand(2).getReg();
6234 Register SamplerReg =
I.getOperand(3).getReg();
6235 Register CoordinateReg =
I.getOperand(4).getReg();
6236 ImageOperands ImOps;
6237 if (
I.getNumOperands() > 5)
6238 ImOps.Offset =
I.getOperand(5).getReg();
6239 if (
I.getNumOperands() > 6)
6240 ImOps.MinLod =
I.getOperand(6).getReg();
6241 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6242 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6245bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
6246 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6247 Register ImageReg =
I.getOperand(2).getReg();
6248 Register SamplerReg =
I.getOperand(3).getReg();
6249 Register CoordinateReg =
I.getOperand(4).getReg();
6250 ImageOperands ImOps;
6251 ImOps.Bias =
I.getOperand(5).getReg();
6252 if (
I.getNumOperands() > 6)
6253 ImOps.Offset =
I.getOperand(6).getReg();
6254 if (
I.getNumOperands() > 7)
6255 ImOps.MinLod =
I.getOperand(7).getReg();
6256 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6257 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6260bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
6261 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6262 Register ImageReg =
I.getOperand(2).getReg();
6263 Register SamplerReg =
I.getOperand(3).getReg();
6264 Register CoordinateReg =
I.getOperand(4).getReg();
6265 ImageOperands ImOps;
6266 ImOps.GradX =
I.getOperand(5).getReg();
6267 ImOps.GradY =
I.getOperand(6).getReg();
6268 if (
I.getNumOperands() > 7)
6269 ImOps.Offset =
I.getOperand(7).getReg();
6270 if (
I.getNumOperands() > 8)
6271 ImOps.MinLod =
I.getOperand(8).getReg();
6272 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6273 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6276bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
6277 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6278 Register ImageReg =
I.getOperand(2).getReg();
6279 Register SamplerReg =
I.getOperand(3).getReg();
6280 Register CoordinateReg =
I.getOperand(4).getReg();
6281 ImageOperands ImOps;
6282 ImOps.Lod =
I.getOperand(5).getReg();
6283 if (
I.getNumOperands() > 6)
6284 ImOps.Offset =
I.getOperand(6).getReg();
6285 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6286 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6289bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
6290 SPIRVTypeInst ResType,
6291 MachineInstr &
I)
const {
6292 Register ImageReg =
I.getOperand(2).getReg();
6293 Register SamplerReg =
I.getOperand(3).getReg();
6294 Register CoordinateReg =
I.getOperand(4).getReg();
6295 ImageOperands ImOps;
6296 ImOps.Compare =
I.getOperand(5).getReg();
6297 if (
I.getNumOperands() > 6)
6298 ImOps.Offset =
I.getOperand(6).getReg();
6299 if (
I.getNumOperands() > 7)
6300 ImOps.MinLod =
I.getOperand(7).getReg();
6301 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6302 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6305bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
6306 SPIRVTypeInst ResType,
6307 MachineInstr &
I)
const {
6308 Register ImageReg =
I.getOperand(2).getReg();
6309 Register CoordinateReg =
I.getOperand(3).getReg();
6310 Register LodReg =
I.getOperand(4).getReg();
6312 ImageOperands ImOps;
6314 if (
I.getNumOperands() > 5)
6315 ImOps.Offset =
I.getOperand(5).getReg();
6327 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
6328 I.getDebugLoc(),
I, &ImOps);
6331bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
6332 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6333 Register ImageReg =
I.getOperand(2).getReg();
6334 Register SamplerReg =
I.getOperand(3).getReg();
6335 Register CoordinateReg =
I.getOperand(4).getReg();
6336 ImageOperands ImOps;
6337 ImOps.Compare =
I.getOperand(5).getReg();
6338 if (
I.getNumOperands() > 6)
6339 ImOps.Offset =
I.getOperand(6).getReg();
6342 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6343 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6346bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
6347 SPIRVTypeInst ResType,
6348 MachineInstr &
I)
const {
6349 Register ImageReg =
I.getOperand(2).getReg();
6350 Register SamplerReg =
I.getOperand(3).getReg();
6351 Register CoordinateReg =
I.getOperand(4).getReg();
6354 "ImageReg is not an image type.");
6359 ComponentOrCompareReg =
I.getOperand(5).getReg();
6360 OffsetReg =
I.getOperand(6).getReg();
6363 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
6367 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6368 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
6369 Dim != SPIRV::Dim::DIM_Rect) {
6371 "Gather operations are only supported for 2D, Cube, and Rect images.");
6378 if (!loadHandleBeforePosition(
6383 MachineIRBuilder MIRBuilder(
I);
6384 SPIRVTypeInst SampledImageType =
6389 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6397 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
6399 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
6401 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6406 .
addUse(ComponentOrCompareReg);
6408 uint32_t ImageOperands = 0;
6409 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
6410 if (Dim == SPIRV::Dim::DIM_Cube) {
6412 "Gather operations with offset are not supported for Cube images.");
6416 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6418 ImageOperands |= SPIRV::ImageOperand::Offset;
6422 if (ImageOperands != 0) {
6423 MIB.
addImm(ImageOperands);
6425 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6433bool SPIRVInstructionSelector::generateImageReadOrFetch(
6436 const ImageOperands *ImOps)
const {
6439 "ImageReg is not an image type.");
6441 bool IsSignedInteger =
6446 bool IsFetch = (SampledOp.getImm() == 1);
6448 auto AddOperands = [&](MachineInstrBuilder &MIB) {
6449 uint32_t ImageOperandsMask = 0;
6450 if (IsSignedInteger)
6451 ImageOperandsMask |= 0x1000;
6453 if (IsFetch && ImOps) {
6455 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
6456 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
6458 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
6460 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
6464 if (ImageOperandsMask != 0) {
6465 MIB.
addImm(ImageOperandsMask);
6466 if (IsFetch && ImOps) {
6469 if (ImOps->Offset &&
6470 (ImageOperandsMask &
6471 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
6472 MIB.
addUse(*ImOps->Offset);
6481 SPIRVTypeInst SampledType =
6484 SPIRVTypeInst ReadType =
6485 widenTypeToVec4(IsPacked ? SampledType : ResType, Pos);
6486 bool ReadTypeMatchesResult = ReadType == ResType;
6488 Register ReadReg = ReadTypeMatchesResult
6494 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
6500 BMI.constrainAllUses(
TII,
TRI, RBI);
6502 if (ReadTypeMatchesResult)
6515 if (ResultSize == 1) {
6524 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
6527bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
6528 SPIRVTypeInst ResType,
6529 MachineInstr &
I)
const {
6530 Register ResourcePtr =
I.getOperand(2).getReg();
6532 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
6541 MachineIRBuilder MIRBuilder(
I);
6546 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6552 if (
I.getNumExplicitOperands() > 3) {
6553 Register IndexReg =
I.getOperand(3).getReg();
6560bool SPIRVInstructionSelector::selectPushConstantGetPointer(
6561 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6566bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
6567 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6568 Register ObjReg =
I.getOperand(2).getReg();
6569 if (!BuildCOPY(ResVReg, ObjReg,
I))
6579 decorateUsesAsNonUniform(ResVReg);
6583void SPIRVInstructionSelector::decorateUsesAsNonUniform(
6586 {NonUniformReg,
nullptr}};
6587 llvm::SmallSet<Register, 8> Visited;
6588 while (WorkList.
size() > 0) {
6591 if (!Visited.
insert(CurrentReg).second)
6594 bool IsDecorated =
false;
6596 if (
Use.getOpcode() == SPIRV::OpDecorate &&
6597 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
6603 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
6605 if (ResultReg == CurrentReg)
6613 MachineInstr &InsertPt =
6616 SPIRV::Decoration::NonUniformEXT, {});
6621bool SPIRVInstructionSelector::extractSubvector(
6623 MachineInstr &InsertionPoint)
const {
6625 [[maybe_unused]] uint64_t InputSize =
6628 assert(InputSize > 1 &&
"The input must be a vector.");
6629 assert(ResultSize > 1 &&
"The result must be a vector.");
6630 assert(ResultSize < InputSize &&
6631 "Cannot extract more element than there are in the input.");
6635 for (uint64_t
I = 0;
I < ResultSize;
I++) {
6638 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6647 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
6649 TII.get(SPIRV::OpCompositeConstruct))
6653 for (
Register ComponentReg : ComponentRegisters)
6654 MIB.
addUse(ComponentReg);
6659bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
6660 MachineInstr &
I)
const {
6667 Register ImageReg =
I.getOperand(1).getReg();
6675 Register CoordinateReg =
I.getOperand(2).getReg();
6676 Register DataReg =
I.getOperand(3).getReg();
6679 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
6687Register SPIRVInstructionSelector::buildPointerToResource(
6688 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
6689 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
6690 StringRef Name, MachineIRBuilder MIRBuilder)
const {
6692 if (ArraySize == 1) {
6693 SPIRVTypeInst PtrType =
6696 "SpirvResType did not have an explicit layout.");
6701 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
6702 SPIRVTypeInst VarPointerType =
6705 VarPointerType, Set,
Binding, Name, MIRBuilder);
6707 SPIRVTypeInst ResPointerType =
6720bool SPIRVInstructionSelector::selectFirstBitSet16(
6721 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
6722 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
6724 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
6728 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
6731bool SPIRVInstructionSelector::selectFirstBitSet32(
6733 unsigned BitSetOpcode)
const {
6734 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6737 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6744bool SPIRVInstructionSelector::selectFirstBitSet64(
6746 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
6759 if (ComponentCount > 2) {
6760 auto Func = [
this, SwapPrimarySide](
Register ResVReg, SPIRVTypeInst ResType,
6762 unsigned Opcode) ->
bool {
6763 return this->selectFirstBitSet64(ResVReg, ResType,
I, SrcReg, Opcode,
6767 return handle64BitOverflow(ResVReg, ResType,
I, SrcReg, BitSetOpcode, Func);
6771 MachineIRBuilder MIRBuilder(
I);
6773 BaseType, 2 * ComponentCount, MIRBuilder,
false);
6777 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
6783 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
6790 bool IsScalarRes = ResType->
getOpcode() != SPIRV::OpTypeVector;
6793 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
6794 SPIRV::OpVectorExtractDynamic))
6796 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
6797 SPIRV::OpVectorExtractDynamic))
6801 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6802 TII.get(SPIRV::OpVectorShuffle))
6810 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
6816 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6817 TII.get(SPIRV::OpVectorShuffle))
6825 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
6845 SelectOp = SPIRV::OpSelectSISCond;
6846 AddOp = SPIRV::OpIAddS;
6854 SelectOp = SPIRV::OpSelectVIVCond;
6855 AddOp = SPIRV::OpIAddV;
6861 Register RegSecondaryOffset = Reg0;
6865 if (SwapPrimarySide) {
6866 PrimaryReg = LowReg;
6867 SecondaryReg = HighReg;
6868 RegPrimaryOffset = Reg0;
6869 RegSecondaryOffset = Reg32;
6874 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
6875 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
6880 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
6881 SPIRV::OpINotEqual))
6888 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
6889 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
6894 if (SwapPrimarySide) {
6896 if (!selectOpWithSrcs(RegAdd, ResType,
I,
6897 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
6908 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
6909 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
6914 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
6915 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
6918 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
6922bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
6923 SPIRVTypeInst ResType,
6925 bool IsSigned)
const {
6927 Register OpReg =
I.getOperand(2).getReg();
6930 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
6931 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
6935 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6937 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6939 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6942 return diagnoseUnsupported(
6944 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
6948bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
6949 SPIRVTypeInst ResType,
6950 MachineInstr &
I)
const {
6952 Register OpReg =
I.getOperand(2).getReg();
6957 unsigned ExtendOpcode = SPIRV::OpUConvert;
6958 unsigned BitSetOpcode = GL::FindILsb;
6962 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6964 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6966 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6969 return diagnoseUnsupported(
I,
6970 "spv_firstbitlow only supports 16,32,64 bits.");
6974bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
6975 SPIRVTypeInst ResType,
6976 MachineInstr &
I)
const {
6980 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
6983 .
addUse(
I.getOperand(2).getReg())
6986 unsigned Alignment =
I.getOperand(3).getImm();
7000 while (!Worklist.
empty()) {
7002 switch (
T->getOpcode()) {
7003 case SPIRV::OpTypeInt:
7004 case SPIRV::OpTypeFloat:
7005 case SPIRV::OpTypePointer:
7007 case SPIRV::OpTypeVector:
7008 case SPIRV::OpTypeMatrix:
7009 case SPIRV::OpTypeArray: {
7010 Register OperandReg =
T->getOperand(1).getReg();
7014 case SPIRV::OpTypeStruct:
7015 for (
unsigned Idx = 1,
E =
T->getNumOperands(); Idx <
E; ++Idx) {
7016 Register OperandReg =
T->getOperand(Idx).getReg();
7028bool SPIRVInstructionSelector::selectAbort(MachineInstr &
I)
const {
7029 assert(
I.getNumExplicitOperands() == 2);
7031 Register MsgReg =
I.getOperand(1).getReg();
7033 assert(MsgType &&
"Message argument of llvm.spv.abort has no SPIR-V type");
7036 return diagnoseUnsupported(
7038 "llvm.spv.abort message type must be a concrete SPIR-V type (numerical "
7039 "scalar, pointer, vector, matrix, or aggregate of such types)");
7042 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7049bool SPIRVInstructionSelector::selectTrap(MachineInstr &
I)
const {
7058 uint32_t MsgVal = ~0
u;
7059 if (
I.getOpcode() == TargetOpcode::G_UBSANTRAP)
7060 MsgVal =
static_cast<uint32_t
>(
I.getOperand(0).
getImm());
7063 Register MsgReg = buildI32ConstantInEntryBlock(MsgVal,
I, MsgType);
7066 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7073bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
7074 SPIRVTypeInst ResType,
7075 MachineInstr &
I)
const {
7082 bool UseUntypedPointers =
7083 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7085 UseUntypedPointers ? SPIRV::OpUntypedVariableKHR : SPIRV::OpVariable;
7087 auto MIB =
BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(Opcode))
7090 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7094 if (UseUntypedPointers) {
7098 return diagnoseUnsupported(
7099 I,
"could not deduce the data type of an untyped variable");
7105 unsigned Alignment =
I.getOperand(2).getImm();
7112bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
7117 const MachineInstr *PrevI =
I.getPrevNode();
7119 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
7123 .
addMBB(
I.getOperand(0).getMBB())
7128 .
addMBB(
I.getOperand(0).getMBB())
7133bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
7144 const MachineInstr *NextI =
I.getNextNode();
7146 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
7152 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
7154 .
addUse(
I.getOperand(0).getReg())
7155 .
addMBB(
I.getOperand(1).getMBB())
7161bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
7162 MachineInstr &
I)
const {
7164 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
7166 const unsigned NumOps =
I.getNumOperands();
7167 for (
unsigned i = 1; i <
NumOps; i += 2) {
7168 MIB.
addUse(
I.getOperand(i + 0).getReg());
7169 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
7175bool SPIRVInstructionSelector::selectGlobalValue(
7176 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
7178 MachineIRBuilder MIRBuilder(
I);
7179 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
7182 std::string GlobalIdent;
7184 unsigned &
ID = UnnamedGlobalIDs[GV];
7186 ID = UnnamedGlobalIDs.
size();
7187 GlobalIdent =
"__unnamed_" + Twine(ID).str();
7213 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
7220 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7225 MachineInstrBuilder MIB1 =
7226 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7229 MachineInstrBuilder MIB2 =
7231 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
7235 GR.
add(ConstVal, MIB2);
7243 MachineInstrBuilder MIB3 =
7244 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7247 GR.
add(ConstVal, MIB3);
7253 assert(NewReg != ResVReg);
7254 return BuildCOPY(ResVReg, NewReg,
I);
7264 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
7267 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
7273 SPIRVTypeInst ResType =
7277 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
7282 if (
GlobalVar->isExternallyInitialized() &&
7283 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
7284 constexpr unsigned ReadWriteINTEL = 3u;
7287 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
7293bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
7294 SPIRVTypeInst ResType,
7295 MachineInstr &
I)
const {
7297 return selectExtInst(ResVReg, ResType,
I, CL::log10);
7305 MachineIRBuilder MIRBuilder(
I);
7310 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7313 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
7315 .
add(
I.getOperand(1))
7320 ResType->
getOpcode() == SPIRV::OpTypeFloat);
7330 APFloat::rmNearestTiesToEven, &LosesInfo);
7334 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
7335 ? SPIRV::OpVectorTimesScalar
7346bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
7347 SPIRVTypeInst ResType,
7348 MachineInstr &
I)
const {
7351 return selectExtInst(ResVReg, ResType,
I, CL::pown);
7357 Register ExpReg =
I.getOperand(2).getReg();
7359 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
7360 SPIRV::OpConvertSToF))
7362 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
7369bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
7370 SPIRVTypeInst ResType,
7371 MachineInstr &
I)
const {
7387 MachineIRBuilder MIRBuilder(
I);
7388 SPIRVTypeInst FloatType =
7392 FloatType, MIRBuilder, SPIRV::StorageClass::Function);
7405 MachineBasicBlock &EntryBB =
I.getMF()->
front();
7406 const bool IsUntyped =
7407 PtrType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7409 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
7410 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
7411 : SPIRV::OpVariable))
7414 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7422 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7425 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
7428 .
add(
I.getOperand(
I.getNumExplicitDefs()))
7432 Register IntegralPartReg =
I.getOperand(1).getReg();
7435 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7445 assert(
false &&
"GLSL::Modf is deprecated.");
7456bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
7457 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7458 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7459 MachineIRBuilder MIRBuilder(
I);
7460 const SPIRVTypeInst Vec3Ty =
7463 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
7475 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7479 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7485 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7492 assert(
I.getOperand(2).isReg());
7493 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
7497 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
7508bool SPIRVInstructionSelector::loadBuiltinInputID(
7509 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7510 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7511 MachineIRBuilder MIRBuilder(
I);
7513 ResType, MIRBuilder, SPIRV::StorageClass::Input);
7528 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7532 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7541SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
7542 MachineInstr &
I)
const {
7543 MachineIRBuilder MIRBuilder(
I);
7544 if (
Type->getOpcode() != SPIRV::OpTypeVector)
7554bool SPIRVInstructionSelector::loadHandleBeforePosition(
7555 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
7556 MachineInstr &Pos)
const {
7559 Intrinsic::spv_resource_handlefrombinding);
7567 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
7568 MachineIRBuilder MIRBuilder(HandleDef);
7569 SPIRVTypeInst VarType = ResType;
7570 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
7572 if (IsStructuredBuffer) {
7577 if (ResType->
getOpcode() == SPIRV::OpTypeImage && ArraySize == 0)
7579 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
7582 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
7583 ArraySize, IndexReg, Name, MIRBuilder);
7587 uint32_t LoadOpcode =
7588 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
7598bool SPIRVInstructionSelector::errorIfInstrOutsideShader(
7599 MachineInstr &
I)
const {
7601 return diagnoseUnsupported(
7602 I,
"this instruction is only supported in shaders.");
7607InstructionSelector *
7611 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.
Represent a constant reference to an array (0 or more elements consecutively in memory),...
BlockFrequencyInfo pass uses BlockFrequencyInfoImpl implementation to estimate IR basic block frequen...
Predicate
This enumeration lists the possible predicates for CmpInst subclasses.
@ FCMP_OEQ
0 0 0 1 True if ordered and equal
@ ICMP_SLT
signed less than
@ ICMP_SLE
signed less or equal
@ FCMP_OLT
0 1 0 0 True if ordered and less than
@ FCMP_ULE
1 1 0 1 True if unordered, less than, or equal
@ FCMP_OGT
0 0 1 0 True if ordered and greater than
@ FCMP_OGE
0 0 1 1 True if ordered and greater than or equal
@ ICMP_UGE
unsigned greater or equal
@ ICMP_UGT
unsigned greater than
@ ICMP_SGT
signed greater than
@ FCMP_ULT
1 1 0 0 True if unordered or less than
@ FCMP_ONE
0 1 1 0 True if ordered and operands are unequal
@ FCMP_UEQ
1 0 0 1 True if unordered or equal
@ ICMP_ULT
unsigned less than
@ FCMP_UGT
1 0 1 0 True if unordered or greater than
@ FCMP_OLE
0 1 0 1 True if ordered and less than or equal
@ FCMP_ORD
0 1 1 1 True if ordered (no nans)
@ ICMP_SGE
signed greater or equal
@ FCMP_UNE
1 1 1 0 True if unordered or not equal
@ ICMP_ULE
unsigned less or equal
@ FCMP_UGE
1 0 1 1 True if unordered, greater than, or equal
@ FCMP_UNO
1 0 0 0 True if unordered: isnan(X) | isnan(Y)
static LLVM_ABI Constant * getNullValue(Type *Ty)
Constructor to create a '0' constant of arbitrary type.
LLVMContext & getContext() const
getContext - Return a reference to the LLVMContext associated with this function.
Represents a call to an intrinsic.
Intrinsic::ID getIntrinsicID() const
unsigned getAddressSpace() const
Module * getParent()
Get the module that this global value is contained inside of...
@ InternalLinkage
Rename collisions when linking (static functions).
static LLVM_ABI IntegerType * get(LLVMContext &C, unsigned NumBits)
This static method is the primary way of constructing an IntegerType.
static constexpr LLT scalar(unsigned SizeInBits)
Get a low-level scalar or aggregate "bag of bits".
constexpr bool isValid() const
constexpr uint16_t getNumElements() const
Returns the number of elements in a vector LLT.
constexpr bool isVector() const
static constexpr LLT pointer(unsigned AddressSpace, unsigned SizeInBits)
Get a low-level pointer in the given address space.
static constexpr LLT fixed_vector(unsigned NumElements, unsigned ScalarSizeInBits)
Get a low-level fixed-width vector of some number of elements and element width.
int getNumber() const
MachineBasicBlocks are uniquely numbered at the function level, unless they're not in a MachineFuncti...
LLVM_ABI iterator getFirstNonPHI()
Returns a pointer to the first instruction in this block that is not a PHINode instruction.
const MachineFunction * getParent() const
Return the MachineFunction containing this basic block.
MachineInstrBundleIterator< MachineInstr > iterator
MachineRegisterInfo & getRegInfo()
getRegInfo - Return information about the registers currently in use.
Function & getFunction()
Return the LLVM function that this machine code represents.
Helper class to build MachineInstr.
MachineBasicBlock::iterator getInsertPt()
Current insertion point for new instructions.
MachineInstrBuilder buildInstr(unsigned Opcode)
Build and insert <empty> = Opcode <empty>.
MachineFunction & getMF()
Getter for the function we currently build.
MachineRegisterInfo * getMRI()
Getter for MRI.
void constrainAllUses(const TargetInstrInfo &TII, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI) const
const MachineInstrBuilder & addUse(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a virtual register use operand.
const MachineInstrBuilder & addReg(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a new virtual register operand.
const MachineInstrBuilder & addImm(int64_t Val) const
Add a new immediate operand.
const MachineInstrBuilder & add(const MachineOperand &MO) const
const MachineInstrBuilder & addMBB(MachineBasicBlock *MBB, unsigned TargetFlags=0) const
const MachineInstrBuilder & addDef(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a virtual register definition operand.
const MachineInstrBuilder & setMIFlags(unsigned Flags) const
MachineInstr * getInstr() const
If conversion operators fail, use this method to get the MachineInstr explicitly.
Representation of each machine instruction.
unsigned getOpcode() const
Returns the opcode of this MachineInstr.
const MachineBasicBlock * getParent() const
unsigned getNumOperands() const
Retuns the total number of operands.
LLVM_ABI unsigned getNumExplicitOperands() const
Returns the number of non-implicit operands.
LLVM_ABI unsigned getNumExplicitDefs() const
Returns the number of non-implicit definitions.
LLVM_ABI void emitGenericError(const Twine &ErrMsg) const
LLVM_ABI const MachineFunction * getMF() const
Return the function that contains the basic block that this instruction belongs to.
const DebugLoc & getDebugLoc() const
Returns the debug location id of this MachineInstr.
const MachineOperand & getOperand(unsigned i) const
A description of a memory reference used in the backend.
@ MOVolatile
The memory access is volatile.
@ MONonTemporal
The memory access is non-temporal.
bool isReg() const
isReg - Tests if this is a MO_Register operand.
MachineBasicBlock * getMBB() const
Register getReg() const
getReg - Returns the register number.
MachineRegisterInfo - Keep track of information for virtual and physical registers,...
defusechain_instr_iterator< true, false, false, true > use_instr_iterator
use_instr_iterator/use_instr_begin/use_instr_end - Walk all uses of the specified register,...
const TargetRegisterClass * getRegClass(Register Reg) const
Return the register class of the specified virtual register.
LLVM_ABI 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.
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 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.
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.
Register createVirtualRegister(SPIRVTypeInst SpvType, SPIRVGlobalRegistry *GR, MachineRegisterInfo *MRI, const MachineFunction &MF)
unsigned getArrayComponentCount(const MachineRegisterInfo *MRI, const MachineInstr *ResType)
void buildOpDecorate(Register Reg, MachineIRBuilder &MIRBuilder, SPIRV::Decoration::Decoration Dec, ArrayRef< uint32_t > DecArgs, StringRef StrImm)
LLVM_ABI bool isNullOrNullSplat(const MachineInstr &MI, const MachineRegisterInfo &MRI, bool AllowUndefs=false)
Return true if the value is a constant 0 integer or a splatted vector of a constant 0 integer (with n...
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)
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