33#include "llvm/IR/IntrinsicsSPIRV.h"
37#define DEBUG_TYPE "spirv-isel"
44 std::vector<std::pair<SPIRV::InstructionSet::InstructionSet, uint32_t>>;
49 std::optional<Register> Bias;
50 std::optional<Register>
Offset;
51 std::optional<Register> MinLod;
52 std::optional<Register> GradX;
53 std::optional<Register> GradY;
54 std::optional<Register> Lod;
55 std::optional<Register> Compare;
58llvm::SPIRV::SelectionControl::SelectionControl
59getSelectionOperandForImm(
int Imm) {
61 return SPIRV::SelectionControl::Flatten;
63 return SPIRV::SelectionControl::DontFlatten;
65 return SPIRV::SelectionControl::None;
69#define GET_GLOBALISEL_PREDICATE_BITSET
70#include "SPIRVGenGlobalISel.inc"
71#undef GET_GLOBALISEL_PREDICATE_BITSET
98#define GET_GLOBALISEL_PREDICATES_DECL
99#include "SPIRVGenGlobalISel.inc"
100#undef GET_GLOBALISEL_PREDICATES_DECL
102#define GET_GLOBALISEL_TEMPORARIES_DECL
103#include "SPIRVGenGlobalISel.inc"
104#undef GET_GLOBALISEL_TEMPORARIES_DECL
128 unsigned BitSetOpcode)
const;
132 unsigned BitSetOpcode)
const;
136 unsigned BitSetOpcode,
bool SwapPrimarySide)
const;
140 unsigned BitSetOpcode,
141 bool SwapPrimarySide)
const;
148 unsigned Opcode)
const;
151 unsigned Opcode)
const;
170 unsigned NewOpcode,
unsigned NegateOpcode = 0)
const;
181 unsigned OpType)
const;
236 template <
bool Signed>
239 template <
bool Signed>
246 template <
typename PickOpcodeFn>
249 PickOpcodeFn &&PickOpcode)
const;
266 template <
typename PickOpcodeFn>
269 PickOpcodeFn &&PickOpcode)
const;
284 bool IsSigned)
const;
286 bool IsSigned,
unsigned Opcode)
const;
288 bool IsSigned)
const;
294 bool IsSigned)
const;
333 GL::GLSLExtInst GLInst,
bool setMIFlags =
true,
334 bool useMISrc =
true,
336 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
337 CL::OpenCLExtInst CLInst,
bool setMIFlags =
true,
338 bool useMISrc =
true,
340 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
341 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
342 bool setMIFlags =
true,
bool useMISrc =
true,
344 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
345 const ExtInstList &ExtInsts,
bool setMIFlags =
true,
346 bool useMISrc =
true,
349 bool selectLog10(
Register ResVReg, SPIRVTypeInst ResType,
350 MachineInstr &
I)
const;
352 bool selectFpowi(
Register ResVReg, SPIRVTypeInst ResType,
353 MachineInstr &
I)
const;
355 bool selectSaturate(
Register ResVReg, SPIRVTypeInst ResType,
356 MachineInstr &
I)
const;
358 bool selectWaveOpInst(
Register ResVReg, SPIRVTypeInst ResType,
359 MachineInstr &
I,
unsigned Opcode)
const;
361 bool selectWaveActiveCountBits(
Register ResVReg, SPIRVTypeInst ResType,
362 MachineInstr &
I)
const;
364 bool selectWaveActiveAllEqual(
Register ResVReg, SPIRVTypeInst ResType,
365 MachineInstr &
I)
const;
369 bool selectHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
370 MachineInstr &
I)
const;
372 bool selectCounterHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
373 MachineInstr &
I)
const;
375 bool selectReadImageIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
376 MachineInstr &
I)
const;
377 bool selectSampleBasicIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
378 MachineInstr &
I)
const;
379 bool selectSampleBiasIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
380 MachineInstr &
I)
const;
381 bool selectSampleGradIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
382 MachineInstr &
I)
const;
383 bool selectSampleLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
384 MachineInstr &
I)
const;
385 bool selectLoadLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
386 MachineInstr &
I)
const;
387 bool selectSampleCmpIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
388 MachineInstr &
I)
const;
389 bool selectSampleCmpLevelZeroIntrinsic(
Register &ResVReg,
390 SPIRVTypeInst ResType,
391 MachineInstr &
I)
const;
392 bool selectGatherIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
393 MachineInstr &
I)
const;
394 bool selectImageWriteIntrinsic(MachineInstr &
I)
const;
395 bool selectResourceGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
396 MachineInstr &
I)
const;
397 bool selectPushConstantGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
398 MachineInstr &
I)
const;
399 bool selectResourceNonUniformIndex(
Register &ResVReg, SPIRVTypeInst ResType,
400 MachineInstr &
I)
const;
401 bool selectModf(
Register ResVReg, SPIRVTypeInst ResType,
402 MachineInstr &
I)
const;
403 bool selectUpdateCounter(
Register &ResVReg, SPIRVTypeInst ResType,
404 MachineInstr &
I)
const;
405 bool selectFrexp(
Register ResVReg, SPIRVTypeInst ResType,
406 MachineInstr &
I)
const;
407 bool selectSincos(
Register ResVReg, SPIRVTypeInst ResType,
408 MachineInstr &
I)
const;
409 bool selectExp10(
Register ResVReg, SPIRVTypeInst ResType,
410 MachineInstr &
I)
const;
411 bool selectDerivativeInst(
Register ResVReg, SPIRVTypeInst ResType,
412 MachineInstr &
I,
const unsigned DPdOpCode)
const;
414 Register buildI32Constant(uint32_t Val, MachineInstr &
I,
415 SPIRVTypeInst ResType =
nullptr)
const;
417 Register buildZerosVal(SPIRVTypeInst ResType, MachineInstr &
I)
const;
418 bool isScalarOrVectorIntConstantZero(
Register Reg)
const;
419 Register buildZerosValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
421 MachineInstr &
I)
const;
422 Register buildOnesValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
424 bool wrapIntoSpecConstantOp(MachineInstr &
I,
427 Register getUcharPtrTypeReg(MachineInstr &
I,
428 SPIRV::StorageClass::StorageClass SC)
const;
429 MachineInstrBuilder buildSpecConstantOp(MachineInstr &
I,
Register Dest,
431 uint32_t Opcode)
const;
432 MachineInstrBuilder buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
433 SPIRVTypeInst SrcPtrTy)
const;
434 Register buildPointerToResource(SPIRVTypeInst ResType,
435 SPIRV::StorageClass::StorageClass SC,
436 uint32_t Set, uint32_t
Binding,
437 uint32_t ArraySize,
Register IndexReg,
439 MachineIRBuilder MIRBuilder)
const;
440 SPIRVTypeInst widenTypeToVec4(SPIRVTypeInst
Type, MachineInstr &
I)
const;
441 bool extractSubvector(
Register &ResVReg, SPIRVTypeInst ResType,
442 Register &ReadReg, MachineInstr &InsertionPoint)
const;
443 bool generateImageReadOrFetch(
Register &ResVReg, SPIRVTypeInst ResType,
446 const ImageOperands *ImOps =
nullptr)
const;
447 bool generateSampleImage(
Register ResVReg, SPIRVTypeInst ResType,
449 Register CoordinateReg,
const ImageOperands &ImOps,
452 bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
453 Register ResVReg, SPIRVTypeInst ResType,
454 MachineInstr &
I)
const;
455 bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
456 Register ResVReg, SPIRVTypeInst ResType,
457 MachineInstr &
I)
const;
458 bool loadHandleBeforePosition(
Register &HandleReg, SPIRVTypeInst ResType,
459 GIntrinsic &HandleDef, MachineInstr &Pos)
const;
460 void decorateUsesAsNonUniform(
Register &NonUniformReg)
const;
461 void errorIfInstrOutsideShader(MachineInstr &
I)
const;
464bool sampledTypeIsSignedInteger(
const llvm::Type *HandleType) {
466 if (
TET->getTargetExtName() ==
"spirv.Image") {
469 assert(
TET->getTargetExtName() ==
"spirv.SignedImage");
470 return TET->getTypeParameter(0)->isIntegerTy();
474#define GET_GLOBALISEL_IMPL
475#include "SPIRVGenGlobalISel.inc"
476#undef GET_GLOBALISEL_IMPL
482 TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
485#include
"SPIRVGenGlobalISel.inc"
488#include
"SPIRVGenGlobalISel.inc"
500 InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
504void SPIRVInstructionSelector::resetVRegsType(MachineFunction &MF) {
505 if (HasVRegsReset == &MF)
520 for (
const auto &
MBB : MF) {
521 for (
const auto &
MI :
MBB) {
524 if (
MI.getOpcode() != SPIRV::ASSIGN_TYPE)
528 LLT DstType = MRI.
getType(DstReg);
530 LLT SrcType = MRI.
getType(SrcReg);
531 if (DstType != SrcType)
536 if (DstRC != SrcRC && SrcRC)
548 while (!Stack.empty()) {
553 switch (
MI->getOpcode()) {
554 case TargetOpcode::G_INTRINSIC:
555 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
556 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
558 Intrinsic::spv_const_composite)
561 case TargetOpcode::G_BUILD_VECTOR:
562 case TargetOpcode::G_SPLAT_VECTOR:
564 i < OpDef->getNumOperands(); i++) {
569 Stack.push_back(OpNestedDef);
572 case TargetOpcode::G_CONSTANT:
573 case TargetOpcode::G_FCONSTANT:
574 case TargetOpcode::G_IMPLICIT_DEF:
575 case SPIRV::OpConstantTrue:
576 case SPIRV::OpConstantFalse:
577 case SPIRV::OpConstantI:
578 case SPIRV::OpConstantF:
579 case SPIRV::OpConstantComposite:
580 case SPIRV::OpConstantCompositeContinuedINTEL:
581 case SPIRV::OpConstantSampler:
582 case SPIRV::OpConstantNull:
584 case SPIRV::OpConstantFunctionPointerINTEL:
611 case Intrinsic::spv_all:
612 case Intrinsic::spv_alloca:
613 case Intrinsic::spv_any:
614 case Intrinsic::spv_bitcast:
615 case Intrinsic::spv_const_composite:
616 case Intrinsic::spv_cross:
617 case Intrinsic::spv_degrees:
618 case Intrinsic::spv_distance:
619 case Intrinsic::spv_extractelt:
620 case Intrinsic::spv_extractv:
621 case Intrinsic::spv_faceforward:
622 case Intrinsic::spv_fdot:
623 case Intrinsic::spv_firstbitlow:
624 case Intrinsic::spv_firstbitshigh:
625 case Intrinsic::spv_firstbituhigh:
626 case Intrinsic::spv_frac:
627 case Intrinsic::spv_gep:
628 case Intrinsic::spv_global_offset:
629 case Intrinsic::spv_global_size:
630 case Intrinsic::spv_group_id:
631 case Intrinsic::spv_insertelt:
632 case Intrinsic::spv_insertv:
633 case Intrinsic::spv_isinf:
634 case Intrinsic::spv_isnan:
635 case Intrinsic::spv_lerp:
636 case Intrinsic::spv_length:
637 case Intrinsic::spv_normalize:
638 case Intrinsic::spv_num_subgroups:
639 case Intrinsic::spv_num_workgroups:
640 case Intrinsic::spv_ptrcast:
641 case Intrinsic::spv_radians:
642 case Intrinsic::spv_reflect:
643 case Intrinsic::spv_refract:
644 case Intrinsic::spv_resource_getpointer:
645 case Intrinsic::spv_resource_handlefrombinding:
646 case Intrinsic::spv_resource_handlefromimplicitbinding:
647 case Intrinsic::spv_resource_nonuniformindex:
648 case Intrinsic::spv_resource_sample:
649 case Intrinsic::spv_rsqrt:
650 case Intrinsic::spv_saturate:
651 case Intrinsic::spv_sdot:
652 case Intrinsic::spv_sign:
653 case Intrinsic::spv_smoothstep:
654 case Intrinsic::spv_step:
655 case Intrinsic::spv_subgroup_id:
656 case Intrinsic::spv_subgroup_local_invocation_id:
657 case Intrinsic::spv_subgroup_max_size:
658 case Intrinsic::spv_subgroup_size:
659 case Intrinsic::spv_thread_id:
660 case Intrinsic::spv_thread_id_in_group:
661 case Intrinsic::spv_udot:
662 case Intrinsic::spv_undef:
663 case Intrinsic::spv_value_md:
664 case Intrinsic::spv_workgroup_size:
676 case SPIRV::OpTypeVoid:
677 case SPIRV::OpTypeBool:
678 case SPIRV::OpTypeInt:
679 case SPIRV::OpTypeFloat:
680 case SPIRV::OpTypeVector:
681 case SPIRV::OpTypeMatrix:
682 case SPIRV::OpTypeImage:
683 case SPIRV::OpTypeSampler:
684 case SPIRV::OpTypeSampledImage:
685 case SPIRV::OpTypeArray:
686 case SPIRV::OpTypeRuntimeArray:
687 case SPIRV::OpTypeStruct:
688 case SPIRV::OpTypeOpaque:
689 case SPIRV::OpTypePointer:
690 case SPIRV::OpTypeFunction:
691 case SPIRV::OpTypeEvent:
692 case SPIRV::OpTypeDeviceEvent:
693 case SPIRV::OpTypeReserveId:
694 case SPIRV::OpTypeQueue:
695 case SPIRV::OpTypePipe:
696 case SPIRV::OpTypeForwardPointer:
697 case SPIRV::OpTypePipeStorage:
698 case SPIRV::OpTypeNamedBarrier:
699 case SPIRV::OpTypeAccelerationStructureNV:
700 case SPIRV::OpTypeCooperativeMatrixNV:
701 case SPIRV::OpTypeCooperativeMatrixKHR:
711 if (
MI.getNumDefs() == 0)
714 for (
const auto &MO :
MI.all_defs()) {
716 if (
Reg.isPhysical()) {
721 if (
UseMI.getOpcode() != SPIRV::OpName) {
728 if (
MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE ||
MI.isFakeUse() ||
729 MI.isLifetimeMarker()) {
732 <<
"Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
743 if (
MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
744 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
747 LLVM_DEBUG(
dbgs() <<
"Dead: Intrinsic with no real side effects.\n");
752 if (
MI.mayStore() ||
MI.isCall() ||
753 (
MI.mayLoad() &&
MI.hasOrderedMemoryRef()) ||
MI.isPosition() ||
754 MI.isDebugInstr() ||
MI.isTerminator() ||
MI.isJumpTableDebugInfo()) {
755 LLVM_DEBUG(
dbgs() <<
"Not dead: instruction has side effects.\n");
766 LLVM_DEBUG(
dbgs() <<
"Dead: known opcode with no side effects\n");
773void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &
MI)
const {
775 for (
const auto &MO :
MI.all_defs()) {
779 SmallVector<MachineInstr *, 4> UselessOpNames;
782 "There is still a use of the dead function.");
785 for (MachineInstr *OpNameMI : UselessOpNames) {
787 OpNameMI->eraseFromParent();
792void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &
MI)
const {
795 removeOpNamesForDeadMI(
MI);
796 MI.eraseFromParent();
799bool SPIRVInstructionSelector::select(MachineInstr &
I) {
800 resetVRegsType(*
I.getParent()->getParent());
802 assert(
I.getParent() &&
"Instruction should be in a basic block!");
803 assert(
I.getParent()->getParent() &&
"Instruction should be in a function!");
808 removeDeadInstruction(
I);
815 if (Opcode == SPIRV::ASSIGN_TYPE) {
816 Register DstReg =
I.getOperand(0).getReg();
817 Register SrcReg =
I.getOperand(1).getReg();
820 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
821 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
822 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
823 Register SelectDstReg =
Def->getOperand(0).getReg();
824 bool SuccessToSelectSelect [[maybe_unused]] = selectSelect(
826 assert(SuccessToSelectSelect);
828 Def->eraseFromParent();
835 bool Res = selectImpl(
I, *CoverageInfo);
837 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
838 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
842 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
854 }
else if (
I.getNumDefs() == 1) {
866 removeDeadInstruction(
I);
871 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
872 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
878 bool HasDefs =
I.getNumDefs() > 0;
881 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
882 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
883 if (spvSelect(ResVReg, ResType,
I)) {
885 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
896 case TargetOpcode::G_CONSTANT:
897 case TargetOpcode::G_FCONSTANT:
899 case TargetOpcode::G_SADDO:
900 case TargetOpcode::G_SSUBO:
907 MachineInstr &
I)
const {
910 if (DstRC != SrcRC && SrcRC)
912 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
919bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
920 SPIRVTypeInst ResType,
921 MachineInstr &
I)
const {
922 const unsigned Opcode =
I.getOpcode();
924 return selectImpl(
I, *CoverageInfo);
926 case TargetOpcode::G_CONSTANT:
927 case TargetOpcode::G_FCONSTANT:
928 return selectConst(ResVReg, ResType,
I);
929 case TargetOpcode::G_GLOBAL_VALUE:
930 return selectGlobalValue(ResVReg,
I);
931 case TargetOpcode::G_IMPLICIT_DEF:
932 return selectOpUndef(ResVReg, ResType,
I);
933 case TargetOpcode::G_FREEZE:
934 return selectFreeze(ResVReg, ResType,
I);
936 case TargetOpcode::G_INTRINSIC:
937 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
938 case TargetOpcode::G_INTRINSIC_CONVERGENT:
939 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
940 return selectIntrinsic(ResVReg, ResType,
I);
941 case TargetOpcode::G_BITREVERSE:
942 return selectBitreverse(ResVReg, ResType,
I);
944 case TargetOpcode::G_BUILD_VECTOR:
945 return selectBuildVector(ResVReg, ResType,
I);
946 case TargetOpcode::G_SPLAT_VECTOR:
947 return selectSplatVector(ResVReg, ResType,
I);
949 case TargetOpcode::G_SHUFFLE_VECTOR: {
950 MachineBasicBlock &BB = *
I.getParent();
951 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
954 .
addUse(
I.getOperand(1).getReg())
955 .
addUse(
I.getOperand(2).getReg());
956 for (
auto V :
I.getOperand(3).getShuffleMask())
961 case TargetOpcode::G_MEMMOVE:
962 case TargetOpcode::G_MEMCPY:
963 case TargetOpcode::G_MEMSET:
964 return selectMemOperation(ResVReg,
I);
966 case TargetOpcode::G_ICMP:
967 return selectICmp(ResVReg, ResType,
I);
968 case TargetOpcode::G_FCMP:
969 return selectFCmp(ResVReg, ResType,
I);
971 case TargetOpcode::G_FRAME_INDEX:
972 return selectFrameIndex(ResVReg, ResType,
I);
974 case TargetOpcode::G_LOAD:
975 return selectLoad(ResVReg, ResType,
I);
976 case TargetOpcode::G_STORE:
977 return selectStore(
I);
979 case TargetOpcode::G_BR:
980 return selectBranch(
I);
981 case TargetOpcode::G_BRCOND:
982 return selectBranchCond(
I);
984 case TargetOpcode::G_PHI:
985 return selectPhi(ResVReg,
I);
987 case TargetOpcode::G_FPTOSI:
988 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
989 case TargetOpcode::G_FPTOUI:
990 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
992 case TargetOpcode::G_FPTOSI_SAT:
993 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
994 case TargetOpcode::G_FPTOUI_SAT:
995 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
997 case TargetOpcode::G_SITOFP:
998 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
999 case TargetOpcode::G_UITOFP:
1000 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
1002 case TargetOpcode::G_CTPOP:
1003 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitCount);
1004 case TargetOpcode::G_SMIN:
1005 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
1006 case TargetOpcode::G_UMIN:
1007 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
1009 case TargetOpcode::G_SMAX:
1010 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
1011 case TargetOpcode::G_UMAX:
1012 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
1014 case TargetOpcode::G_SCMP:
1015 return selectSUCmp(ResVReg, ResType,
I,
true);
1016 case TargetOpcode::G_UCMP:
1017 return selectSUCmp(ResVReg, ResType,
I,
false);
1018 case TargetOpcode::G_LROUND:
1019 case TargetOpcode::G_LLROUND: {
1022 MRI->
setRegClass(regForLround, &SPIRV::iIDRegClass);
1024 regForLround, *(
I.getParent()->getParent()));
1026 CL::round, GL::Round,
false);
1028 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1035 case TargetOpcode::G_STRICT_FMA:
1036 case TargetOpcode::G_FMA: {
1039 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1042 .
addUse(
I.getOperand(1).getReg())
1043 .
addUse(
I.getOperand(2).getReg())
1044 .
addUse(
I.getOperand(3).getReg())
1049 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1052 case TargetOpcode::G_STRICT_FLDEXP:
1053 return selectExtInst(ResVReg, ResType,
I, CL::ldexp);
1055 case TargetOpcode::G_FPOW:
1056 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1057 case TargetOpcode::G_FPOWI:
1058 return selectFpowi(ResVReg, ResType,
I);
1060 case TargetOpcode::G_FEXP:
1061 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1062 case TargetOpcode::G_FEXP2:
1063 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1064 case TargetOpcode::G_FEXP10:
1065 return selectExp10(ResVReg, ResType,
I);
1067 case TargetOpcode::G_FMODF:
1068 return selectModf(ResVReg, ResType,
I);
1069 case TargetOpcode::G_FSINCOS:
1070 return selectSincos(ResVReg, ResType,
I);
1072 case TargetOpcode::G_FLOG:
1073 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1074 case TargetOpcode::G_FLOG2:
1075 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1076 case TargetOpcode::G_FLOG10:
1077 return selectLog10(ResVReg, ResType,
I);
1079 case TargetOpcode::G_FABS:
1080 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1081 case TargetOpcode::G_ABS:
1082 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1084 case TargetOpcode::G_FMINNUM:
1085 case TargetOpcode::G_FMINIMUM:
1086 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1087 case TargetOpcode::G_FMAXNUM:
1088 case TargetOpcode::G_FMAXIMUM:
1089 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1091 case TargetOpcode::G_FCOPYSIGN:
1092 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1094 case TargetOpcode::G_FCEIL:
1095 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1096 case TargetOpcode::G_FFLOOR:
1097 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1099 case TargetOpcode::G_FCOS:
1100 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1101 case TargetOpcode::G_FSIN:
1102 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1103 case TargetOpcode::G_FTAN:
1104 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1105 case TargetOpcode::G_FACOS:
1106 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1107 case TargetOpcode::G_FASIN:
1108 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1109 case TargetOpcode::G_FATAN:
1110 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1111 case TargetOpcode::G_FATAN2:
1112 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1113 case TargetOpcode::G_FCOSH:
1114 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1115 case TargetOpcode::G_FSINH:
1116 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1117 case TargetOpcode::G_FTANH:
1118 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1120 case TargetOpcode::G_STRICT_FSQRT:
1121 case TargetOpcode::G_FSQRT:
1122 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1124 case TargetOpcode::G_CTTZ:
1125 case TargetOpcode::G_CTTZ_ZERO_UNDEF:
1126 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1127 case TargetOpcode::G_CTLZ:
1128 case TargetOpcode::G_CTLZ_ZERO_UNDEF:
1129 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1131 case TargetOpcode::G_INTRINSIC_ROUND:
1132 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1133 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1134 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1135 case TargetOpcode::G_INTRINSIC_TRUNC:
1136 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1137 case TargetOpcode::G_FRINT:
1138 case TargetOpcode::G_FNEARBYINT:
1139 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1141 case TargetOpcode::G_SMULH:
1142 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1143 case TargetOpcode::G_UMULH:
1144 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1146 case TargetOpcode::G_SADDSAT:
1147 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1148 case TargetOpcode::G_UADDSAT:
1149 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1150 case TargetOpcode::G_SSUBSAT:
1151 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1152 case TargetOpcode::G_USUBSAT:
1153 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1155 case TargetOpcode::G_FFREXP:
1156 return selectFrexp(ResVReg, ResType,
I);
1158 case TargetOpcode::G_UADDO:
1159 return selectOverflowArith(ResVReg, ResType,
I,
1160 ResType->
getOpcode() == SPIRV::OpTypeVector
1161 ? SPIRV::OpIAddCarryV
1162 : SPIRV::OpIAddCarryS);
1163 case TargetOpcode::G_USUBO:
1164 return selectOverflowArith(ResVReg, ResType,
I,
1165 ResType->
getOpcode() == SPIRV::OpTypeVector
1166 ? SPIRV::OpISubBorrowV
1167 : SPIRV::OpISubBorrowS);
1168 case TargetOpcode::G_UMULO:
1169 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1170 case TargetOpcode::G_SMULO:
1171 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1173 case TargetOpcode::G_SEXT:
1174 return selectExt(ResVReg, ResType,
I,
true);
1175 case TargetOpcode::G_ANYEXT:
1176 case TargetOpcode::G_ZEXT:
1177 return selectExt(ResVReg, ResType,
I,
false);
1178 case TargetOpcode::G_TRUNC:
1179 return selectTrunc(ResVReg, ResType,
I);
1180 case TargetOpcode::G_FPTRUNC:
1181 case TargetOpcode::G_FPEXT:
1182 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1184 case TargetOpcode::G_PTRTOINT:
1185 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1186 case TargetOpcode::G_INTTOPTR:
1187 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1188 case TargetOpcode::G_BITCAST:
1189 return selectBitcast(ResVReg, ResType,
I);
1190 case TargetOpcode::G_ADDRSPACE_CAST:
1191 return selectAddrSpaceCast(ResVReg, ResType,
I);
1192 case TargetOpcode::G_PTR_ADD: {
1194 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1198 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1199 (*II).getOpcode() == TargetOpcode::COPY ||
1200 (*II).getOpcode() == SPIRV::OpVariable) &&
1201 getImm(
I.getOperand(2), MRI));
1203 bool IsGVInit =
false;
1207 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1208 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1209 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1210 (*UseIt).getOpcode() == SPIRV::OpVariable) {
1220 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1233 "incompatible result and operand types in a bitcast");
1235 MachineInstrBuilder MIB =
1236 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1243 : SPIRV::OpInBoundsPtrAccessChain))
1247 .
addUse(
I.getOperand(2).getReg())
1250 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1254 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1256 .
addUse(
I.getOperand(2).getReg())
1265 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1268 .
addImm(
static_cast<uint32_t
>(
1269 SPIRV::Opcode::InBoundsPtrAccessChain))
1272 .
addUse(
I.getOperand(2).getReg());
1277 case TargetOpcode::G_ATOMICRMW_OR:
1278 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1279 case TargetOpcode::G_ATOMICRMW_ADD:
1280 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1281 case TargetOpcode::G_ATOMICRMW_AND:
1282 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1283 case TargetOpcode::G_ATOMICRMW_MAX:
1284 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1285 case TargetOpcode::G_ATOMICRMW_MIN:
1286 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1287 case TargetOpcode::G_ATOMICRMW_SUB:
1288 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1289 case TargetOpcode::G_ATOMICRMW_XOR:
1290 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1291 case TargetOpcode::G_ATOMICRMW_UMAX:
1292 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1293 case TargetOpcode::G_ATOMICRMW_UMIN:
1294 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1295 case TargetOpcode::G_ATOMICRMW_XCHG:
1296 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1297 case TargetOpcode::G_ATOMIC_CMPXCHG:
1298 return selectAtomicCmpXchg(ResVReg, ResType,
I);
1300 case TargetOpcode::G_ATOMICRMW_FADD:
1301 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1302 case TargetOpcode::G_ATOMICRMW_FSUB:
1304 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1305 ResType->
getOpcode() == SPIRV::OpTypeVector
1307 : SPIRV::OpFNegate);
1308 case TargetOpcode::G_ATOMICRMW_FMIN:
1309 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1310 case TargetOpcode::G_ATOMICRMW_FMAX:
1311 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1313 case TargetOpcode::G_FENCE:
1314 return selectFence(
I);
1316 case TargetOpcode::G_STACKSAVE:
1317 return selectStackSave(ResVReg, ResType,
I);
1318 case TargetOpcode::G_STACKRESTORE:
1319 return selectStackRestore(
I);
1321 case TargetOpcode::G_UNMERGE_VALUES:
1327 case TargetOpcode::G_TRAP:
1328 case TargetOpcode::G_UBSANTRAP:
1329 case TargetOpcode::DBG_LABEL:
1331 case TargetOpcode::G_DEBUGTRAP:
1332 return selectDebugTrap(ResVReg, ResType,
I);
1339bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1340 SPIRVTypeInst ResType,
1341 MachineInstr &
I)
const {
1342 unsigned Opcode = SPIRV::OpNop;
1349bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1350 SPIRVTypeInst ResType,
1352 GL::GLSLExtInst GLInst,
1353 bool setMIFlags,
bool useMISrc,
1356 SPIRV::InstructionSet::InstructionSet::GLSL_std_450)) {
1357 std::string DiagMsg;
1358 raw_string_ostream OS(DiagMsg);
1359 I.print(OS,
true,
false,
false,
false);
1360 DiagMsg +=
" is only supported with the GLSL extended instruction set.\n";
1363 return selectExtInst(ResVReg, ResType,
I,
1364 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1365 setMIFlags, useMISrc, SrcRegs);
1368bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1369 SPIRVTypeInst ResType,
1371 CL::OpenCLExtInst CLInst,
1372 bool setMIFlags,
bool useMISrc,
1374 return selectExtInst(ResVReg, ResType,
I,
1375 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1376 setMIFlags, useMISrc, SrcRegs);
1379bool SPIRVInstructionSelector::selectExtInst(
1380 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1381 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1383 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1384 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1385 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1389bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1390 SPIRVTypeInst ResType,
1393 bool setMIFlags,
bool useMISrc,
1396 for (
const auto &[InstructionSet, Opcode] : Insts) {
1400 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1403 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1408 const unsigned NumOps =
I.getNumOperands();
1411 I.getOperand(Index).getType() ==
1412 MachineOperand::MachineOperandType::MO_IntrinsicID)
1415 MIB.
add(
I.getOperand(Index));
1427bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1428 SPIRVTypeInst ResType,
1429 MachineInstr &
I)
const {
1430 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1431 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1432 for (
const auto &Ex : ExtInsts) {
1433 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1434 uint32_t Opcode = Ex.second;
1438 MachineIRBuilder MIRBuilder(
I);
1441 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1446 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1449 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1452 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1455 .
addImm(
static_cast<uint32_t
>(Ex.first))
1457 .
add(
I.getOperand(2))
1461 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1462 .
addDef(
I.getOperand(1).getReg())
1471bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1472 SPIRVTypeInst ResType,
1473 MachineInstr &
I)
const {
1474 Register CosResVReg =
I.getOperand(1).getReg();
1475 unsigned SrcIdx =
I.getNumExplicitDefs();
1480 MachineIRBuilder MIRBuilder(
I);
1482 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1487 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1490 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1492 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1495 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1497 .
add(
I.getOperand(SrcIdx))
1500 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1508 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1511 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1513 .
add(
I.getOperand(SrcIdx))
1515 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1518 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1520 .
add(
I.getOperand(SrcIdx))
1527bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1528 SPIRVTypeInst ResType,
1530 std::vector<Register> Srcs,
1531 unsigned Opcode)
const {
1532 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1542bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1543 SPIRVTypeInst ResType,
1545 unsigned Opcode)
const {
1547 Register SrcReg =
I.getOperand(1).getReg();
1552 unsigned DefOpCode = DefIt->getOpcode();
1553 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1556 if (
auto *VRD =
getVRegDef(*MRI, DefIt->getOperand(1).getReg()))
1557 DefOpCode = VRD->getOpcode();
1559 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1560 DefOpCode == TargetOpcode::G_CONSTANT ||
1561 DefOpCode == SPIRV::OpVariable || DefOpCode == SPIRV::OpConstantI) {
1567 uint32_t SpecOpcode = 0;
1569 case SPIRV::OpConvertPtrToU:
1570 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1572 case SPIRV::OpConvertUToPtr:
1573 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1578 TII.get(SPIRV::OpSpecConstantOp))
1588 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1592bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1593 SPIRVTypeInst ResType,
1594 MachineInstr &
I)
const {
1595 Register OpReg =
I.getOperand(1).getReg();
1596 SPIRVTypeInst OpType =
1600 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1610 if (
MemOp->isVolatile())
1611 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1612 if (
MemOp->isNonTemporal())
1613 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1615 if (!ST->isShader() &&
MemOp->getAlign().value())
1616 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1620 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1621 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1625 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1627 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1631 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1635 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1637 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1649 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1651 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1653 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
1657bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
1658 SPIRVTypeInst ResType,
1659 MachineInstr &
I)
const {
1661 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1666 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
1667 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1669 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1673 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1677 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1678 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
1679 I.getDebugLoc(),
I);
1683 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1687 if (!
I.getNumMemOperands()) {
1688 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1690 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1693 MachineIRBuilder MIRBuilder(
I);
1700bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
1702 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
1703 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1708 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
1709 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1714 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1718 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1719 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1720 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1721 TII.get(SPIRV::OpImageWrite))
1727 if (sampledTypeIsSignedInteger(LLVMHandleType))
1730 BMI.constrainAllUses(
TII,
TRI, RBI);
1736 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpStore))
1739 if (!
I.getNumMemOperands()) {
1740 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1742 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1745 MachineIRBuilder MIRBuilder(
I);
1752bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
1753 SPIRVTypeInst ResType,
1754 MachineInstr &
I)
const {
1755 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
1763 Register PtrsReg =
I.getOperand(2).getReg();
1764 uint32_t Alignment =
I.getOperand(3).getImm();
1765 Register MaskReg =
I.getOperand(4).getReg();
1766 Register PassthruReg =
I.getOperand(5).getReg();
1767 Register AlignmentReg = buildI32Constant(Alignment,
I);
1771 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
1782bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
1783 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
1790 Register ValuesReg =
I.getOperand(1).getReg();
1791 Register PtrsReg =
I.getOperand(2).getReg();
1792 uint32_t Alignment =
I.getOperand(3).getImm();
1793 Register MaskReg =
I.getOperand(4).getReg();
1794 Register AlignmentReg = buildI32Constant(Alignment,
I);
1798 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
1807bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
1808 const Twine &Msg)
const {
1809 const Function &
F =
I.getMF()->getFunction();
1810 F.getContext().diagnose(
1811 DiagnosticInfoUnsupported(
F, Msg,
I.getDebugLoc(),
DS_Error));
1815bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
1816 SPIRVTypeInst ResType,
1817 MachineInstr &
I)
const {
1818 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
1820 "llvm.stacksave intrinsic: this instruction requires the following "
1821 "SPIR-V extension: SPV_INTEL_variable_length_array",
1824 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
1831bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
1832 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
1834 "llvm.stackrestore intrinsic: this instruction requires the following "
1835 "SPIR-V extension: SPV_INTEL_variable_length_array",
1837 if (!
I.getOperand(0).isReg())
1840 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
1841 .
addUse(
I.getOperand(0).getReg())
1847SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
1848 MachineIRBuilder MIRBuilder(
I);
1849 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1856 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
1860 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
1861 Type *ArrTy = ArrayType::get(ValTy, Num);
1863 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
1866 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
1873 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
1876 .
addImm(SPIRV::StorageClass::UniformConstant)
1887bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
1890 Register DstReg =
I.getOperand(0).getReg();
1895 uint64_t CopySize =
getIConstVal(
I.getOperand(2).getReg(), MRI);
1900 "Unable to determine pointee type size for OpCopyMemory");
1901 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
1902 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
1904 "OpCopyMemory requires the size to match the pointee type size");
1905 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemory))
1908 if (
I.getNumMemOperands()) {
1909 MachineIRBuilder MIRBuilder(
I);
1916bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &
I,
1919 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemorySized))
1920 .
addUse(
I.getOperand(0).getReg())
1922 .
addUse(
I.getOperand(2).getReg());
1923 if (
I.getNumMemOperands()) {
1924 MachineIRBuilder MIRBuilder(
I);
1931bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
1932 MachineInstr &
I)
const {
1933 Register SrcReg =
I.getOperand(1).getReg();
1934 if (
I.getOpcode() == TargetOpcode::G_MEMSET) {
1935 Register VarReg = getOrCreateMemSetGlobal(
I);
1938 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
1940 ValTy,
I, SPIRV::StorageClass::UniformConstant);
1942 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
1946 if (!selectCopyMemory(
I, SrcReg))
1949 if (!selectCopyMemorySized(
I, SrcReg))
1952 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
1953 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
1958bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
1959 SPIRVTypeInst ResType,
1962 unsigned NegateOpcode)
const {
1964 const MachineMemOperand *MemOp = *
I.memoperands_begin();
1967 Register ScopeReg = buildI32Constant(Scope,
I);
1969 Register Ptr =
I.getOperand(1).getReg();
1975 Register MemSemReg = buildI32Constant(MemSem ,
I);
1977 Register ValueReg =
I.getOperand(2).getReg();
1978 if (NegateOpcode != 0) {
1981 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
1986 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
1997bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
1998 unsigned ArgI =
I.getNumOperands() - 1;
2000 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2001 SPIRVTypeInst SrcType =
2003 if (!SrcType || SrcType->
getOpcode() != SPIRV::OpTypeVector)
2005 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2007 SPIRVTypeInst ScalarType =
2010 unsigned CurrentIndex = 0;
2011 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2012 Register ResVReg =
I.getOperand(i).getReg();
2015 LLT ResLLT = MRI->
getType(ResVReg);
2021 ResType = ScalarType;
2027 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
2030 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2036 for (
unsigned j = 0;
j < NumElements; ++
j) {
2037 MIB.
addImm(CurrentIndex + j);
2039 CurrentIndex += NumElements;
2043 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2055bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2058 Register MemSemReg = buildI32Constant(MemSem,
I);
2060 uint32_t
Scope =
static_cast<uint32_t
>(
2062 Register ScopeReg = buildI32Constant(Scope,
I);
2064 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2071bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2072 SPIRVTypeInst ResType,
2074 unsigned Opcode)
const {
2075 Type *ResTy =
nullptr;
2079 "Not enough info to select the arithmetic with overflow instruction");
2082 "with overflow instruction");
2088 MachineIRBuilder MIRBuilder(
I);
2090 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2091 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2097 Register ZeroReg = buildZerosVal(ResType,
I);
2102 if (ResName.
size() > 0)
2107 BuildMI(BB, MIRBuilder.getInsertPt(),
I.getDebugLoc(),
TII.get(Opcode))
2110 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2111 MIB.
addUse(
I.getOperand(i).getReg());
2116 MRI->
setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2117 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2119 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2120 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2127 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2128 .
addDef(
I.getOperand(1).getReg())
2136bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2137 SPIRVTypeInst ResType,
2138 MachineInstr &
I)
const {
2142 Register Ptr =
I.getOperand(2).getReg();
2145 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2148 ScopeReg = buildI32Constant(Scope,
I);
2150 unsigned ScSem =
static_cast<uint32_t
>(
2153 unsigned MemSemEq =
static_cast<uint32_t
>(
getMemSemantics(AO)) | ScSem;
2154 Register MemSemEqReg = buildI32Constant(MemSemEq,
I);
2156 unsigned MemSemNeq =
static_cast<uint32_t
>(
getMemSemantics(FO)) | ScSem;
2157 if (MemSemEq == MemSemNeq)
2158 MemSemNeqReg = MemSemEqReg;
2160 MemSemNeqReg = buildI32Constant(MemSemEq,
I);
2163 ScopeReg =
I.getOperand(5).getReg();
2164 MemSemEqReg =
I.getOperand(6).getReg();
2165 MemSemNeqReg =
I.getOperand(7).getReg();
2169 Register Val =
I.getOperand(4).getReg();
2173 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2192 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2199 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2211 case SPIRV::StorageClass::DeviceOnlyINTEL:
2212 case SPIRV::StorageClass::HostOnlyINTEL:
2221 bool IsGRef =
false;
2222 bool IsAllowedRefs =
2224 unsigned Opcode = It.getOpcode();
2225 if (Opcode == SPIRV::OpConstantComposite ||
2226 Opcode == SPIRV::OpVariable ||
2227 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2228 return IsGRef = true;
2229 return Opcode == SPIRV::OpName;
2231 return IsAllowedRefs && IsGRef;
2234Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2235 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2237 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2241SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2243 uint32_t Opcode)
const {
2244 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2245 TII.get(SPIRV::OpSpecConstantOp))
2253SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2254 SPIRVTypeInst SrcPtrTy)
const {
2255 SPIRVTypeInst GenericPtrTy =
2259 SPIRV::StorageClass::Generic),
2261 MachineFunction *MF =
I.getParent()->getParent();
2263 MachineInstrBuilder MIB = buildSpecConstantOp(
2265 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2275bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2276 SPIRVTypeInst ResType,
2277 MachineInstr &
I)
const {
2281 Register SrcPtr =
I.getOperand(1).getReg();
2285 if (SrcPtrTy->
getOpcode() != SPIRV::OpTypePointer ||
2286 ResType->
getOpcode() != SPIRV::OpTypePointer)
2287 return BuildCOPY(ResVReg, SrcPtr,
I);
2297 unsigned SpecOpcode =
2299 ?
static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric)
2302 ? static_cast<uint32_t>(
SPIRV::Opcode::GenericCastToPtr)
2309 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2311 .constrainAllUses(
TII,
TRI, RBI);
2313 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2315 buildSpecConstantOp(
2317 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2318 .constrainAllUses(
TII,
TRI, RBI);
2325 return BuildCOPY(ResVReg, SrcPtr,
I);
2327 if ((SrcSC == SPIRV::StorageClass::Function &&
2328 DstSC == SPIRV::StorageClass::Private) ||
2329 (DstSC == SPIRV::StorageClass::Function &&
2330 SrcSC == SPIRV::StorageClass::Private))
2331 return BuildCOPY(ResVReg, SrcPtr,
I);
2335 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2338 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2341 SPIRVTypeInst GenericPtrTy =
2360 return selectUnOp(ResVReg, ResType,
I,
2361 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
2363 return selectUnOp(ResVReg, ResType,
I,
2364 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
2366 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2368 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2378 return SPIRV::OpFOrdEqual;
2380 return SPIRV::OpFOrdGreaterThanEqual;
2382 return SPIRV::OpFOrdGreaterThan;
2384 return SPIRV::OpFOrdLessThanEqual;
2386 return SPIRV::OpFOrdLessThan;
2388 return SPIRV::OpFOrdNotEqual;
2390 return SPIRV::OpOrdered;
2392 return SPIRV::OpFUnordEqual;
2394 return SPIRV::OpFUnordGreaterThanEqual;
2396 return SPIRV::OpFUnordGreaterThan;
2398 return SPIRV::OpFUnordLessThanEqual;
2400 return SPIRV::OpFUnordLessThan;
2402 return SPIRV::OpFUnordNotEqual;
2404 return SPIRV::OpUnordered;
2414 return SPIRV::OpIEqual;
2416 return SPIRV::OpINotEqual;
2418 return SPIRV::OpSGreaterThanEqual;
2420 return SPIRV::OpSGreaterThan;
2422 return SPIRV::OpSLessThanEqual;
2424 return SPIRV::OpSLessThan;
2426 return SPIRV::OpUGreaterThanEqual;
2428 return SPIRV::OpUGreaterThan;
2430 return SPIRV::OpULessThanEqual;
2432 return SPIRV::OpULessThan;
2441 return SPIRV::OpPtrEqual;
2443 return SPIRV::OpPtrNotEqual;
2454 return SPIRV::OpLogicalEqual;
2456 return SPIRV::OpLogicalNotEqual;
2490bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
2491 SPIRVTypeInst ResType,
2493 unsigned OpAnyOrAll)
const {
2494 assert(
I.getNumOperands() == 3);
2495 assert(
I.getOperand(2).isReg());
2497 Register InputRegister =
I.getOperand(2).getReg();
2500 assert(InputType &&
"VReg has no type assigned");
2503 bool IsVectorTy = InputType->
getOpcode() == SPIRV::OpTypeVector;
2504 if (IsBoolTy && !IsVectorTy) {
2505 assert(ResVReg ==
I.getOperand(0).getReg());
2506 return BuildCOPY(ResVReg, InputRegister,
I);
2510 unsigned SpirvNotEqualId =
2511 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
2513 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
2518 IsBoolTy ? InputRegister
2526 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
2528 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
2545bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
2546 SPIRVTypeInst ResType,
2547 MachineInstr &
I)
const {
2548 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
2551bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
2552 SPIRVTypeInst ResType,
2553 MachineInstr &
I)
const {
2554 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
2558bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
2559 SPIRVTypeInst ResType,
2560 MachineInstr &
I)
const {
2561 assert(
I.getNumOperands() == 4);
2562 assert(
I.getOperand(2).isReg());
2563 assert(
I.getOperand(3).isReg());
2565 [[maybe_unused]] SPIRVTypeInst VecType =
2570 "dot product requires a vector of at least 2 components");
2572 [[maybe_unused]] SPIRVTypeInst EltType =
2581 .
addUse(
I.getOperand(2).getReg())
2582 .
addUse(
I.getOperand(3).getReg())
2587bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
2588 SPIRVTypeInst ResType,
2591 assert(
I.getNumOperands() == 4);
2592 assert(
I.getOperand(2).isReg());
2593 assert(
I.getOperand(3).isReg());
2596 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
2600 .
addUse(
I.getOperand(2).getReg())
2601 .
addUse(
I.getOperand(3).getReg())
2608bool SPIRVInstructionSelector::selectIntegerDotExpansion(
2609 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2610 assert(
I.getNumOperands() == 4);
2611 assert(
I.getOperand(2).isReg());
2612 assert(
I.getOperand(3).isReg());
2616 Register Vec0 =
I.getOperand(2).getReg();
2617 Register Vec1 =
I.getOperand(3).getReg();
2621 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
2630 "dot product requires a vector of at least 2 components");
2633 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2643 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2654 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2666bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
2667 SPIRVTypeInst ResType,
2668 MachineInstr &
I)
const {
2670 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
2673 .
addUse(
I.getOperand(2).getReg())
2678bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
2679 SPIRVTypeInst ResType,
2680 MachineInstr &
I)
const {
2682 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
2685 .
addUse(
I.getOperand(2).getReg())
2690template <
bool Signed>
2691bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
2692 SPIRVTypeInst ResType,
2693 MachineInstr &
I)
const {
2694 assert(
I.getNumOperands() == 5);
2695 assert(
I.getOperand(2).isReg());
2696 assert(
I.getOperand(3).isReg());
2697 assert(
I.getOperand(4).isReg());
2700 Register Acc =
I.getOperand(2).getReg();
2704 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
2706 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
2711 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
2714 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2726template <
bool Signed>
2727bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
2728 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2729 assert(
I.getNumOperands() == 5);
2730 assert(
I.getOperand(2).isReg());
2731 assert(
I.getOperand(3).isReg());
2732 assert(
I.getOperand(4).isReg());
2735 Register Acc =
I.getOperand(2).getReg();
2741 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
2745 for (
unsigned i = 0; i < 4; i++) {
2768 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
2788 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2803bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
2804 SPIRVTypeInst ResType,
2805 MachineInstr &
I)
const {
2806 assert(
I.getNumOperands() == 3);
2807 assert(
I.getOperand(2).isReg());
2809 Register VZero = buildZerosValF(ResType,
I);
2810 Register VOne = buildOnesValF(ResType,
I);
2812 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
2815 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
2817 .
addUse(
I.getOperand(2).getReg())
2824bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
2825 SPIRVTypeInst ResType,
2826 MachineInstr &
I)
const {
2827 assert(
I.getNumOperands() == 3);
2828 assert(
I.getOperand(2).isReg());
2830 Register InputRegister =
I.getOperand(2).getReg();
2832 auto &
DL =
I.getDebugLoc();
2842 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
2844 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
2852 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
2857 if (NeedsConversion) {
2858 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
2869bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
2870 SPIRVTypeInst ResType,
2872 unsigned Opcode)
const {
2876 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
2882 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
2883 BMI.addUse(
I.getOperand(J).getReg());
2890bool SPIRVInstructionSelector::selectWaveActiveCountBits(
2891 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2896 if (!selectWaveOpInst(BallotReg, BallotType,
I,
2897 SPIRV::OpGroupNonUniformBallot))
2902 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
2907 .
addImm(SPIRV::GroupOperation::Reduce)
2916 if (
Type->getOpcode() != SPIRV::OpTypeVector)
2920 return Type->getOperand(2).getImm();
2923bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
2924 SPIRVTypeInst ResType,
2925 MachineInstr &
I)
const {
2930 Register InputReg =
I.getOperand(2).getReg();
2935 bool IsVector = NumElems > 1;
2938 SPIRVTypeInst ElemInputType = InputType;
2939 SPIRVTypeInst ElemBoolType = ResType;
2952 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
2953 SPIRV::OpGroupNonUniformAllEqual);
2958 ElementResults.
reserve(NumElems);
2960 for (
unsigned Idx = 0; Idx < NumElems; ++Idx) {
2973 ElemInput = Extracted;
2979 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
2990 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3001bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3002 SPIRVTypeInst ResType,
3003 MachineInstr &
I)
const {
3005 assert(
I.getNumOperands() == 3);
3007 auto Op =
I.getOperand(2);
3019 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3041 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3045 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3052bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3053 SPIRVTypeInst ResType,
3055 bool IsUnsigned)
const {
3056 return selectWaveReduce(
3057 ResVReg, ResType,
I, IsUnsigned,
3058 [&](
Register InputRegister,
bool IsUnsigned) {
3059 const bool IsFloatTy =
3061 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3062 : SPIRV::OpGroupNonUniformSMax;
3063 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3067bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3068 SPIRVTypeInst ResType,
3070 bool IsUnsigned)
const {
3071 return selectWaveReduce(
3072 ResVReg, ResType,
I, IsUnsigned,
3073 [&](
Register InputRegister,
bool IsUnsigned) {
3074 const bool IsFloatTy =
3076 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3077 : SPIRV::OpGroupNonUniformSMin;
3078 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3082bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3083 SPIRVTypeInst ResType,
3084 MachineInstr &
I)
const {
3085 return selectWaveReduce(ResVReg, ResType,
I,
false,
3086 [&](
Register InputRegister,
bool IsUnsigned) {
3088 InputRegister, SPIRV::OpTypeFloat);
3089 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3090 : SPIRV::OpGroupNonUniformIAdd;
3094bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3095 SPIRVTypeInst ResType,
3096 MachineInstr &
I)
const {
3097 return selectWaveReduce(ResVReg, ResType,
I,
false,
3098 [&](
Register InputRegister,
bool IsUnsigned) {
3100 InputRegister, SPIRV::OpTypeFloat);
3101 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3102 : SPIRV::OpGroupNonUniformIMul;
3106template <
typename PickOpcodeFn>
3107bool SPIRVInstructionSelector::selectWaveReduce(
3108 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3109 PickOpcodeFn &&PickOpcode)
const {
3110 assert(
I.getNumOperands() == 3);
3111 assert(
I.getOperand(2).isReg());
3113 Register InputRegister =
I.getOperand(2).getReg();
3120 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3126 .
addImm(SPIRV::GroupOperation::Reduce)
3127 .
addUse(
I.getOperand(2).getReg())
3132bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3133 SPIRVTypeInst ResType,
3135 unsigned Opcode)
const {
3136 return selectWaveReduce(
3137 ResVReg, ResType,
I,
false,
3138 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3141bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3142 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3143 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3144 [&](
Register InputRegister,
bool IsUnsigned) {
3146 InputRegister, SPIRV::OpTypeFloat);
3148 ? SPIRV::OpGroupNonUniformFAdd
3149 : SPIRV::OpGroupNonUniformIAdd;
3153bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3154 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3155 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3156 [&](
Register InputRegister,
bool IsUnsigned) {
3158 InputRegister, SPIRV::OpTypeFloat);
3160 ? SPIRV::OpGroupNonUniformFMul
3161 : SPIRV::OpGroupNonUniformIMul;
3165template <
typename PickOpcodeFn>
3166bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3167 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3168 PickOpcodeFn &&PickOpcode)
const {
3169 assert(
I.getNumOperands() == 3);
3170 assert(
I.getOperand(2).isReg());
3172 Register InputRegister =
I.getOperand(2).getReg();
3179 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3185 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3186 .
addUse(
I.getOperand(2).getReg())
3191bool SPIRVInstructionSelector::selectBitreverse16(
Register ResVReg,
3192 SPIRVTypeInst ResType,
3197 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
3202 : SPIRV::OpUConvert;
3206 ShiftOp = SPIRV::OpShiftRightLogicalV;
3211 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3212 TII.get(SPIRV::OpConstantComposite))
3215 for (
unsigned It = 0; It <
N; ++It)
3219 ShiftConst = CompositeReg;
3224 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
3229 if (!selectBitreverse32(BitrevReg, Int32Type,
I, ExtReg))
3234 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
3239 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
3242bool SPIRVInstructionSelector::selectBitreverse32(
Register ResVReg,
3243 SPIRVTypeInst ResType,
3247 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
3255bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
3256 SPIRVTypeInst ResType,
3257 MachineInstr &
I)
const {
3258 Register OpReg =
I.getOperand(1).getReg();
3262 return selectBitreverse16(ResVReg, ResType,
I, OpReg);
3264 return selectBitreverse32(ResVReg, ResType,
I, OpReg);
3268bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
3269 SPIRVTypeInst ResType,
3270 MachineInstr &
I)
const {
3276 if (!
I.getOperand(0).isReg() || !
I.getOperand(1).isReg())
3278 Register OpReg =
I.getOperand(1).getReg();
3279 if (MachineInstr *Def = MRI->
getVRegDef(OpReg)) {
3280 if (
Def->getOpcode() == TargetOpcode::COPY)
3283 switch (
Def->getOpcode()) {
3284 case SPIRV::ASSIGN_TYPE:
3285 if (MachineInstr *AssignToDef =
3287 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
3288 Reg =
Def->getOperand(2).getReg();
3291 case SPIRV::OpUndef:
3292 Reg =
Def->getOperand(1).getReg();
3295 unsigned DestOpCode;
3297 DestOpCode = SPIRV::OpConstantNull;
3299 DestOpCode = TargetOpcode::COPY;
3302 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
3303 .
addDef(
I.getOperand(0).getReg())
3311bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
3312 SPIRVTypeInst ResType,
3313 MachineInstr &
I)
const {
3315 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3317 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
3321 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
3326 for (
unsigned i =
I.getNumExplicitDefs();
3327 i <
I.getNumExplicitOperands() && IsConst; ++i)
3331 if (!IsConst &&
N < 2)
3333 "There must be at least two constituent operands in a vector");
3336 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3337 TII.get(IsConst ? SPIRV::OpConstantComposite
3338 : SPIRV::OpCompositeConstruct))
3341 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
3342 MIB.
addUse(
I.getOperand(i).getReg());
3347bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
3348 SPIRVTypeInst ResType,
3349 MachineInstr &
I)
const {
3351 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3353 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
3359 if (!
I.getOperand(
OpIdx).isReg())
3366 if (!IsConst &&
N < 2)
3368 "There must be at least two constituent operands in a vector");
3371 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3372 TII.get(IsConst ? SPIRV::OpConstantComposite
3373 : SPIRV::OpCompositeConstruct))
3376 for (
unsigned i = 0; i <
N; ++i)
3382bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
3383 SPIRVTypeInst ResType,
3384 MachineInstr &
I)
const {
3389 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
3391 Opcode = SPIRV::OpDemoteToHelperInvocation;
3393 Opcode = SPIRV::OpKill;
3395 if (MachineInstr *NextI =
I.getNextNode()) {
3397 NextI->eraseFromParent();
3407bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
3408 SPIRVTypeInst ResType,
unsigned CmpOpc,
3409 MachineInstr &
I)
const {
3410 Register Cmp0 =
I.getOperand(2).getReg();
3411 Register Cmp1 =
I.getOperand(3).getReg();
3414 "CMP operands should have the same type");
3415 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
3425bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
3426 SPIRVTypeInst ResType,
3427 MachineInstr &
I)
const {
3428 auto Pred =
I.getOperand(1).getPredicate();
3431 Register CmpOperand =
I.getOperand(2).getReg();
3438 return selectCmp(ResVReg, ResType, CmpOpc,
I);
3442SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
3443 SPIRVTypeInst ResType)
const {
3445 SPIRVTypeInst SpvI32Ty =
3448 auto ConstInt = ConstantInt::get(LLVMTy, Val);
3455 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
3458 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
3461 .
addImm(APInt(32, Val).getZExtValue());
3463 GR.
add(ConstInt,
MI);
3468bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
3469 SPIRVTypeInst ResType,
3470 MachineInstr &
I)
const {
3472 return selectCmp(ResVReg, ResType, CmpOp,
I);
3475bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
3476 SPIRVTypeInst ResType,
3477 MachineInstr &
I)
const {
3479 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
3486 if (ResType->
getOpcode() != SPIRV::OpTypeVector &&
3487 ResType->
getOpcode() != SPIRV::OpTypeFloat)
3490 MachineIRBuilder MIRBuilder(
I);
3492 SPIRVTypeInst SpirvScalarType = ResType->
getOpcode() == SPIRV::OpTypeVector
3498 "only float operands supported by GLSL extended math");
3501 MIRBuilder, SpirvScalarType);
3503 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
3504 ? SPIRV::OpVectorTimesScalar
3507 if (!selectOpWithSrcs(ArgReg, ResType,
I,
3508 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
3510 if (!selectExtInst(ResVReg, ResType,
I,
3511 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
3521Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
3522 MachineInstr &
I)
const {
3525 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3530bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
3536 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
3544 if (
Def->getOpcode() == SPIRV::OpConstantNull)
3547 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
3548 Def->getOpcode() == SPIRV::OpConstantI)
3561 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
3562 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
3564 Intrinsic::spv_const_composite)) {
3565 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
3566 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
3567 if (!IsZero(
Def->getOperand(i).getReg()))
3576Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
3577 MachineInstr &
I)
const {
3581 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3586Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
3587 MachineInstr &
I)
const {
3591 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3597 SPIRVTypeInst ResType,
3598 MachineInstr &
I)
const {
3602 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3607bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
3608 SPIRVTypeInst ResType,
3609 MachineInstr &
I)
const {
3610 Register SelectFirstArg =
I.getOperand(2).getReg();
3611 Register SelectSecondArg =
I.getOperand(3).getReg();
3620 SPIRV::OpTypeVector;
3627 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
3628 }
else if (IsPtrTy) {
3629 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
3631 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
3635 Opcode = IsScalarBool ? SPIRV::OpSelectSFSCond : SPIRV::OpSelectVFVCond;
3636 }
else if (IsPtrTy) {
3637 Opcode = IsScalarBool ? SPIRV::OpSelectSPSCond : SPIRV::OpSelectVPVCond;
3639 Opcode = IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
3642 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
3645 .
addUse(
I.getOperand(1).getReg())
3654bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
3655 SPIRVTypeInst ResType,
3657 MachineInstr &InsertAt,
3658 bool IsSigned)
const {
3660 Register ZeroReg = buildZerosVal(ResType, InsertAt);
3661 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
3662 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
3664 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
3676bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
3677 SPIRVTypeInst ResType,
3678 MachineInstr &
I,
bool IsSigned,
3679 unsigned Opcode)
const {
3680 Register SrcReg =
I.getOperand(1).getReg();
3686 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
3691 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I,
false);
3693 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
3696bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
3697 SPIRVTypeInst ResType, MachineInstr &
I,
3698 bool IsSigned)
const {
3699 Register SrcReg =
I.getOperand(1).getReg();
3701 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
3705 if (ResType == SrcType)
3706 return BuildCOPY(ResVReg, SrcReg,
I);
3708 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
3709 return selectUnOp(ResVReg, ResType,
I, Opcode);
3712bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
3713 SPIRVTypeInst ResType,
3715 bool IsSigned)
const {
3716 MachineIRBuilder MIRBuilder(
I);
3717 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
3732 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
3735 .
addUse(
I.getOperand(1).getReg())
3736 .
addUse(
I.getOperand(2).getReg())
3742 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
3745 .
addUse(
I.getOperand(1).getReg())
3746 .
addUse(
I.getOperand(2).getReg())
3754 unsigned SelectOpcode =
3755 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
3760 .
addUse(buildOnesVal(
true, ResType,
I))
3761 .
addUse(buildZerosVal(ResType,
I))
3768 .
addUse(buildOnesVal(
false, ResType,
I))
3773bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
3776 SPIRVTypeInst IntTy,
3777 SPIRVTypeInst BoolTy)
const {
3780 bool IsVectorTy = IntTy->
getOpcode() == SPIRV::OpTypeVector;
3781 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
3783 Register One = buildOnesVal(
false, IntTy,
I);
3791 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
3800bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
3801 SPIRVTypeInst ResType,
3802 MachineInstr &
I)
const {
3803 Register IntReg =
I.getOperand(1).getReg();
3806 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
3807 if (ArgType == ResType)
3808 return BuildCOPY(ResVReg, IntReg,
I);
3810 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
3811 return selectUnOp(ResVReg, ResType,
I, Opcode);
3814bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
3815 SPIRVTypeInst ResType,
3816 MachineInstr &
I)
const {
3817 unsigned Opcode =
I.getOpcode();
3818 unsigned TpOpcode = ResType->
getOpcode();
3820 if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {
3821 assert(Opcode == TargetOpcode::G_CONSTANT &&
3822 I.getOperand(1).getCImm()->isZero());
3823 MachineBasicBlock &DepMBB =
I.getMF()->front();
3826 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
3833 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
3836bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
3837 SPIRVTypeInst ResType,
3838 MachineInstr &
I)
const {
3839 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
3846bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
3847 SPIRVTypeInst ResType,
3848 MachineInstr &
I)
const {
3850 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
3854 .
addUse(
I.getOperand(3).getReg())
3856 .
addUse(
I.getOperand(2).getReg());
3857 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
3863bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
3864 SPIRVTypeInst ResType,
3865 MachineInstr &
I)
const {
3866 Type *MaybeResTy =
nullptr;
3871 "Expected aggregate type for extractv instruction");
3873 SPIRV::AccessQualifier::ReadWrite,
false);
3877 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3880 .
addUse(
I.getOperand(2).getReg());
3881 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
3887bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
3888 SPIRVTypeInst ResType,
3889 MachineInstr &
I)
const {
3890 if (
getImm(
I.getOperand(4), MRI))
3891 return selectInsertVal(ResVReg, ResType,
I);
3893 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
3896 .
addUse(
I.getOperand(2).getReg())
3897 .
addUse(
I.getOperand(3).getReg())
3898 .
addUse(
I.getOperand(4).getReg())
3903bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
3904 SPIRVTypeInst ResType,
3905 MachineInstr &
I)
const {
3906 if (
getImm(
I.getOperand(3), MRI))
3907 return selectExtractVal(ResVReg, ResType,
I);
3909 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
3912 .
addUse(
I.getOperand(2).getReg())
3913 .
addUse(
I.getOperand(3).getReg())
3918bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
3919 SPIRVTypeInst ResType,
3920 MachineInstr &
I)
const {
3921 const bool IsGEPInBounds =
I.getOperand(2).getImm();
3927 ? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
3928 : SPIRV::OpAccessChain)
3929 : (IsGEPInBounds ?
SPIRV::OpInBoundsPtrAccessChain
3930 :
SPIRV::OpPtrAccessChain);
3932 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
3936 .
addUse(
I.getOperand(3).getReg());
3938 (Opcode == SPIRV::OpPtrAccessChain ||
3939 Opcode == SPIRV::OpInBoundsPtrAccessChain ||
3940 (
getImm(
I.getOperand(4), MRI) &&
foldImm(
I.getOperand(4), MRI) == 0)) &&
3941 "Cannot translate GEP to OpAccessChain. First index must be 0.");
3944 const unsigned StartingIndex =
3945 (Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)
3948 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
3949 Res.addUse(
I.getOperand(i).getReg());
3950 Res.constrainAllUses(
TII,
TRI, RBI);
3955bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
3957 unsigned Lim =
I.getNumExplicitOperands();
3958 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
3959 Register OpReg =
I.getOperand(i).getReg();
3960 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
3962 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
3963 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
3964 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
3971 MachineFunction *MF =
I.getMF();
3983 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3984 TII.get(SPIRV::OpSpecConstantOp))
3987 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
3989 GR.
add(OpDefine, MIB);
3995bool SPIRVInstructionSelector::selectDerivativeInst(
3996 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
3997 const unsigned DPdOpCode)
const {
4000 errorIfInstrOutsideShader(
I);
4005 Register SrcReg =
I.getOperand(2).getReg();
4010 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4013 .
addUse(
I.getOperand(2).getReg());
4015 MachineIRBuilder MIRBuilder(
I);
4018 if (componentCount != 1)
4022 const TargetRegisterClass *RegClass = GR.
getRegClass(SrcType);
4026 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4031 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4036 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4044bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
4045 SPIRVTypeInst ResType,
4046 MachineInstr &
I)
const {
4050 case Intrinsic::spv_load:
4051 return selectLoad(ResVReg, ResType,
I);
4052 case Intrinsic::spv_store:
4053 return selectStore(
I);
4054 case Intrinsic::spv_extractv:
4055 return selectExtractVal(ResVReg, ResType,
I);
4056 case Intrinsic::spv_insertv:
4057 return selectInsertVal(ResVReg, ResType,
I);
4058 case Intrinsic::spv_extractelt:
4059 return selectExtractElt(ResVReg, ResType,
I);
4060 case Intrinsic::spv_insertelt:
4061 return selectInsertElt(ResVReg, ResType,
I);
4062 case Intrinsic::spv_gep:
4063 return selectGEP(ResVReg, ResType,
I);
4064 case Intrinsic::spv_bitcast: {
4065 Register OpReg =
I.getOperand(2).getReg();
4066 SPIRVTypeInst OpType =
4070 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
4072 case Intrinsic::spv_unref_global:
4073 case Intrinsic::spv_init_global: {
4074 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
4079 Register GVarVReg =
MI->getOperand(0).getReg();
4080 if (!selectGlobalValue(GVarVReg, *
MI, Init))
4085 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
4087 MI->eraseFromParent();
4091 case Intrinsic::spv_undef: {
4092 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4098 case Intrinsic::spv_const_composite: {
4100 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
4106 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
4108 MachineIRBuilder MIR(
I);
4110 MIR, SPIRV::OpConstantComposite, 3,
4111 SPIRV::OpConstantCompositeContinuedINTEL, CompositeArgs, ResVReg,
4113 for (
auto *Instr : Instructions) {
4114 Instr->setDebugLoc(
I.getDebugLoc());
4119 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4126 case Intrinsic::spv_assign_name: {
4127 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
4128 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
4129 for (
unsigned i =
I.getNumExplicitDefs() + 2;
4130 i <
I.getNumExplicitOperands(); ++i) {
4131 MIB.
addImm(
I.getOperand(i).getImm());
4136 case Intrinsic::spv_switch: {
4137 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
4138 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
4139 if (
I.getOperand(i).isReg())
4140 MIB.
addReg(
I.getOperand(i).getReg());
4141 else if (
I.getOperand(i).isCImm())
4142 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
4143 else if (
I.getOperand(i).isMBB())
4144 MIB.
addMBB(
I.getOperand(i).getMBB());
4151 case Intrinsic::spv_loop_merge: {
4152 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
4153 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
4154 if (
I.getOperand(i).isMBB())
4155 MIB.
addMBB(
I.getOperand(i).getMBB());
4162 case Intrinsic::spv_loop_control_intel: {
4164 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
4165 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
4170 case Intrinsic::spv_selection_merge: {
4172 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
4173 assert(
I.getOperand(1).isMBB() &&
4174 "operand 1 to spv_selection_merge must be a basic block");
4175 MIB.
addMBB(
I.getOperand(1).getMBB());
4176 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
4180 case Intrinsic::spv_cmpxchg:
4181 return selectAtomicCmpXchg(ResVReg, ResType,
I);
4182 case Intrinsic::spv_unreachable:
4183 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
4186 case Intrinsic::spv_alloca:
4187 return selectFrameIndex(ResVReg, ResType,
I);
4188 case Intrinsic::spv_alloca_array:
4189 return selectAllocaArray(ResVReg, ResType,
I);
4190 case Intrinsic::spv_assume:
4192 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
4193 .
addUse(
I.getOperand(1).getReg())
4198 case Intrinsic::spv_expect:
4200 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
4203 .
addUse(
I.getOperand(2).getReg())
4204 .
addUse(
I.getOperand(3).getReg())
4209 case Intrinsic::arithmetic_fence:
4210 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
4211 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
4214 .
addUse(
I.getOperand(2).getReg())
4218 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
4220 case Intrinsic::spv_thread_id:
4226 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
4228 case Intrinsic::spv_thread_id_in_group:
4234 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
4236 case Intrinsic::spv_group_id:
4242 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
4244 case Intrinsic::spv_flattened_thread_id_in_group:
4251 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
4253 case Intrinsic::spv_workgroup_size:
4254 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
4256 case Intrinsic::spv_global_size:
4257 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
4259 case Intrinsic::spv_global_offset:
4260 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
4262 case Intrinsic::spv_num_workgroups:
4263 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
4265 case Intrinsic::spv_subgroup_size:
4266 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
4268 case Intrinsic::spv_num_subgroups:
4269 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
4271 case Intrinsic::spv_subgroup_id:
4272 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
4273 case Intrinsic::spv_subgroup_local_invocation_id:
4274 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
4275 ResVReg, ResType,
I);
4276 case Intrinsic::spv_subgroup_max_size:
4277 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
4279 case Intrinsic::spv_fdot:
4280 return selectFloatDot(ResVReg, ResType,
I);
4281 case Intrinsic::spv_udot:
4282 case Intrinsic::spv_sdot:
4283 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4285 return selectIntegerDot(ResVReg, ResType,
I,
4286 IID == Intrinsic::spv_sdot);
4287 return selectIntegerDotExpansion(ResVReg, ResType,
I);
4288 case Intrinsic::spv_dot4add_i8packed:
4289 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4291 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
4292 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
4293 case Intrinsic::spv_dot4add_u8packed:
4294 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4296 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
4297 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
4298 case Intrinsic::spv_all:
4299 return selectAll(ResVReg, ResType,
I);
4300 case Intrinsic::spv_any:
4301 return selectAny(ResVReg, ResType,
I);
4302 case Intrinsic::spv_cross:
4303 return selectExtInst(ResVReg, ResType,
I, CL::cross, GL::Cross);
4304 case Intrinsic::spv_distance:
4305 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
4306 case Intrinsic::spv_lerp:
4307 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
4308 case Intrinsic::spv_length:
4309 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
4310 case Intrinsic::spv_degrees:
4311 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
4312 case Intrinsic::spv_faceforward:
4313 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
4314 case Intrinsic::spv_frac:
4315 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
4316 case Intrinsic::spv_isinf:
4317 return selectOpIsInf(ResVReg, ResType,
I);
4318 case Intrinsic::spv_isnan:
4319 return selectOpIsNan(ResVReg, ResType,
I);
4320 case Intrinsic::spv_normalize:
4321 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
4322 case Intrinsic::spv_refract:
4323 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
4324 case Intrinsic::spv_reflect:
4325 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
4326 case Intrinsic::spv_rsqrt:
4327 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
4328 case Intrinsic::spv_sign:
4329 return selectSign(ResVReg, ResType,
I);
4330 case Intrinsic::spv_smoothstep:
4331 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
4332 case Intrinsic::spv_firstbituhigh:
4333 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
4334 case Intrinsic::spv_firstbitshigh:
4335 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
4336 case Intrinsic::spv_firstbitlow:
4337 return selectFirstBitLow(ResVReg, ResType,
I);
4338 case Intrinsic::spv_group_memory_barrier_with_group_sync: {
4340 buildI32Constant(SPIRV::MemorySemantics::SequentiallyConsistent,
I);
4341 Register ScopeReg = buildI32Constant(SPIRV::Scope::Workgroup,
I);
4343 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpControlBarrier))
4350 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
4351 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
4352 SPIRV::StorageClass::StorageClass ResSC =
4356 "Generic storage class");
4357 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
4365 case Intrinsic::spv_lifetime_start:
4366 case Intrinsic::spv_lifetime_end: {
4367 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
4368 : SPIRV::OpLifetimeStop;
4369 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
4370 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
4379 case Intrinsic::spv_saturate:
4380 return selectSaturate(ResVReg, ResType,
I);
4381 case Intrinsic::spv_nclamp:
4382 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
4383 case Intrinsic::spv_uclamp:
4384 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
4385 case Intrinsic::spv_sclamp:
4386 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
4387 case Intrinsic::spv_subgroup_prefix_bit_count:
4388 return selectWavePrefixBitCount(ResVReg, ResType,
I);
4389 case Intrinsic::spv_wave_active_countbits:
4390 return selectWaveActiveCountBits(ResVReg, ResType,
I);
4391 case Intrinsic::spv_wave_all_equal:
4392 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
4393 case Intrinsic::spv_wave_all:
4394 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
4395 case Intrinsic::spv_wave_any:
4396 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
4397 case Intrinsic::spv_subgroup_ballot:
4398 return selectWaveOpInst(ResVReg, ResType,
I,
4399 SPIRV::OpGroupNonUniformBallot);
4400 case Intrinsic::spv_wave_is_first_lane:
4401 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
4402 case Intrinsic::spv_wave_reduce_or:
4403 return selectWaveReduceOp(ResVReg, ResType,
I,
4404 SPIRV::OpGroupNonUniformBitwiseOr);
4405 case Intrinsic::spv_wave_reduce_xor:
4406 return selectWaveReduceOp(ResVReg, ResType,
I,
4407 SPIRV::OpGroupNonUniformBitwiseXor);
4408 case Intrinsic::spv_wave_reduce_umax:
4409 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
4410 case Intrinsic::spv_wave_reduce_max:
4411 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
4412 case Intrinsic::spv_wave_reduce_umin:
4413 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
4414 case Intrinsic::spv_wave_reduce_min:
4415 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
4416 case Intrinsic::spv_wave_reduce_sum:
4417 return selectWaveReduceSum(ResVReg, ResType,
I);
4418 case Intrinsic::spv_wave_product:
4419 return selectWaveReduceProduct(ResVReg, ResType,
I);
4420 case Intrinsic::spv_wave_readlane:
4421 return selectWaveOpInst(ResVReg, ResType,
I,
4422 SPIRV::OpGroupNonUniformShuffle);
4423 case Intrinsic::spv_wave_prefix_sum:
4424 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
4425 case Intrinsic::spv_wave_prefix_product:
4426 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
4427 case Intrinsic::spv_step:
4428 return selectExtInst(ResVReg, ResType,
I, CL::step, GL::Step);
4429 case Intrinsic::spv_radians:
4430 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
4434 case Intrinsic::instrprof_increment:
4435 case Intrinsic::instrprof_increment_step:
4436 case Intrinsic::instrprof_value_profile:
4439 case Intrinsic::spv_value_md:
4441 case Intrinsic::spv_resource_handlefrombinding: {
4442 return selectHandleFromBinding(ResVReg, ResType,
I);
4444 case Intrinsic::spv_resource_counterhandlefrombinding:
4445 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
4446 case Intrinsic::spv_resource_updatecounter:
4447 return selectUpdateCounter(ResVReg, ResType,
I);
4448 case Intrinsic::spv_resource_store_typedbuffer: {
4449 return selectImageWriteIntrinsic(
I);
4451 case Intrinsic::spv_resource_load_typedbuffer: {
4452 return selectReadImageIntrinsic(ResVReg, ResType,
I);
4454 case Intrinsic::spv_resource_load_level: {
4455 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
4457 case Intrinsic::spv_resource_sample:
4458 case Intrinsic::spv_resource_sample_clamp:
4459 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
4460 case Intrinsic::spv_resource_samplebias:
4461 case Intrinsic::spv_resource_samplebias_clamp:
4462 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
4463 case Intrinsic::spv_resource_samplegrad:
4464 case Intrinsic::spv_resource_samplegrad_clamp:
4465 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
4466 case Intrinsic::spv_resource_samplelevel:
4467 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
4468 case Intrinsic::spv_resource_samplecmp:
4469 case Intrinsic::spv_resource_samplecmp_clamp:
4470 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
4471 case Intrinsic::spv_resource_samplecmplevelzero:
4472 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
4473 case Intrinsic::spv_resource_gather:
4474 case Intrinsic::spv_resource_gather_cmp:
4475 return selectGatherIntrinsic(ResVReg, ResType,
I);
4476 case Intrinsic::spv_resource_getpointer: {
4477 return selectResourceGetPointer(ResVReg, ResType,
I);
4479 case Intrinsic::spv_pushconstant_getpointer: {
4480 return selectPushConstantGetPointer(ResVReg, ResType,
I);
4482 case Intrinsic::spv_discard: {
4483 return selectDiscard(ResVReg, ResType,
I);
4485 case Intrinsic::spv_resource_nonuniformindex: {
4486 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
4488 case Intrinsic::spv_unpackhalf2x16: {
4489 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
4491 case Intrinsic::spv_packhalf2x16: {
4492 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
4494 case Intrinsic::spv_ddx:
4495 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
4496 case Intrinsic::spv_ddy:
4497 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
4498 case Intrinsic::spv_ddx_coarse:
4499 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
4500 case Intrinsic::spv_ddy_coarse:
4501 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
4502 case Intrinsic::spv_ddx_fine:
4503 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
4504 case Intrinsic::spv_ddy_fine:
4505 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
4506 case Intrinsic::spv_fwidth:
4507 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
4508 case Intrinsic::spv_masked_gather:
4509 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
4510 return selectMaskedGather(ResVReg, ResType,
I);
4511 return diagnoseUnsupported(
4512 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
4513 case Intrinsic::spv_masked_scatter:
4514 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
4515 return selectMaskedScatter(
I);
4516 return diagnoseUnsupported(
4517 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
4519 std::string DiagMsg;
4520 raw_string_ostream OS(DiagMsg);
4522 DiagMsg =
"Intrinsic selection not implemented: " + DiagMsg;
4529bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
4530 SPIRVTypeInst ResType,
4531 MachineInstr &
I)
const {
4534 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
4541bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
4542 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4544 assert(Intr.getIntrinsicID() ==
4545 Intrinsic::spv_resource_counterhandlefrombinding);
4548 Register MainHandleReg = Intr.getOperand(2).getReg();
4550 assert(MainHandleDef->getIntrinsicID() ==
4551 Intrinsic::spv_resource_handlefrombinding);
4555 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
4556 Register IndexReg = MainHandleDef->getOperand(5).getReg();
4557 std::string CounterName =
4562 MachineIRBuilder MIRBuilder(
I);
4564 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
4566 ArraySize, IndexReg, CounterName, MIRBuilder);
4568 return BuildCOPY(ResVReg, CounterVarReg,
I);
4571bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
4572 SPIRVTypeInst ResType,
4573 MachineInstr &
I)
const {
4575 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
4577 Register CounterHandleReg = Intr.getOperand(2).getReg();
4578 Register IncrReg = Intr.getOperand(3).getReg();
4585 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
4586 assert(CounterVarPointeeType &&
4587 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
4588 "Counter variable must be a struct");
4590 SPIRV::StorageClass::StorageBuffer &&
4591 "Counter variable must be in the storage buffer storage class");
4593 "Counter variable must have exactly 1 member in the struct");
4594 const SPIRVTypeInst MemberType =
4597 "Counter variable struct must have a single i32 member");
4601 MachineIRBuilder MIRBuilder(
I);
4603 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
4606 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
4612 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
4615 .
addUse(CounterHandleReg)
4622 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
4625 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
4628 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
4637 return BuildCOPY(ResVReg, AtomicRes,
I);
4645 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
4653bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
4654 SPIRVTypeInst ResType,
4655 MachineInstr &
I)
const {
4663 Register ImageReg =
I.getOperand(2).getReg();
4671 Register IdxReg =
I.getOperand(3).getReg();
4673 MachineInstr &Pos =
I;
4675 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
4679bool SPIRVInstructionSelector::generateSampleImage(
4682 DebugLoc Loc, MachineInstr &Pos)
const {
4693 if (!loadHandleBeforePosition(NewSamplerReg,
4699 MachineIRBuilder MIRBuilder(Pos);
4712 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
4713 ImOps.Lod.has_value();
4714 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
4715 : SPIRV::OpImageSampleImplicitLod;
4717 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
4718 : SPIRV::OpImageSampleDrefImplicitLod;
4727 MIB.
addUse(*ImOps.Compare);
4729 uint32_t ImageOperands = 0;
4731 ImageOperands |= SPIRV::ImageOperand::Bias;
4733 ImageOperands |= SPIRV::ImageOperand::Lod;
4734 if (ImOps.GradX && ImOps.GradY)
4735 ImageOperands |= SPIRV::ImageOperand::Grad;
4736 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
4738 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
4741 "Non-constant offsets are not supported in sample instructions.");
4745 ImageOperands |= SPIRV::ImageOperand::MinLod;
4747 if (ImageOperands != 0) {
4748 MIB.
addImm(ImageOperands);
4749 if (ImageOperands & SPIRV::ImageOperand::Bias)
4751 if (ImageOperands & SPIRV::ImageOperand::Lod)
4753 if (ImageOperands & SPIRV::ImageOperand::Grad) {
4754 MIB.
addUse(*ImOps.GradX);
4755 MIB.
addUse(*ImOps.GradY);
4758 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
4759 MIB.
addUse(*ImOps.Offset);
4760 if (ImageOperands & SPIRV::ImageOperand::MinLod)
4761 MIB.
addUse(*ImOps.MinLod);
4768bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
4769 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4770 Register ImageReg =
I.getOperand(2).getReg();
4771 Register SamplerReg =
I.getOperand(3).getReg();
4772 Register CoordinateReg =
I.getOperand(4).getReg();
4773 ImageOperands ImOps;
4774 if (
I.getNumOperands() > 5)
4775 ImOps.Offset =
I.getOperand(5).getReg();
4776 if (
I.getNumOperands() > 6)
4777 ImOps.MinLod =
I.getOperand(6).getReg();
4778 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4779 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4782bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
4783 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4784 Register ImageReg =
I.getOperand(2).getReg();
4785 Register SamplerReg =
I.getOperand(3).getReg();
4786 Register CoordinateReg =
I.getOperand(4).getReg();
4787 ImageOperands ImOps;
4788 ImOps.Bias =
I.getOperand(5).getReg();
4789 if (
I.getNumOperands() > 6)
4790 ImOps.Offset =
I.getOperand(6).getReg();
4791 if (
I.getNumOperands() > 7)
4792 ImOps.MinLod =
I.getOperand(7).getReg();
4793 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4794 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4797bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
4798 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4799 Register ImageReg =
I.getOperand(2).getReg();
4800 Register SamplerReg =
I.getOperand(3).getReg();
4801 Register CoordinateReg =
I.getOperand(4).getReg();
4802 ImageOperands ImOps;
4803 ImOps.GradX =
I.getOperand(5).getReg();
4804 ImOps.GradY =
I.getOperand(6).getReg();
4805 if (
I.getNumOperands() > 7)
4806 ImOps.Offset =
I.getOperand(7).getReg();
4807 if (
I.getNumOperands() > 8)
4808 ImOps.MinLod =
I.getOperand(8).getReg();
4809 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4810 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4813bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
4814 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4815 Register ImageReg =
I.getOperand(2).getReg();
4816 Register SamplerReg =
I.getOperand(3).getReg();
4817 Register CoordinateReg =
I.getOperand(4).getReg();
4818 ImageOperands ImOps;
4819 ImOps.Lod =
I.getOperand(5).getReg();
4820 if (
I.getNumOperands() > 6)
4821 ImOps.Offset =
I.getOperand(6).getReg();
4822 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4823 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4826bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
4827 SPIRVTypeInst ResType,
4828 MachineInstr &
I)
const {
4829 Register ImageReg =
I.getOperand(2).getReg();
4830 Register SamplerReg =
I.getOperand(3).getReg();
4831 Register CoordinateReg =
I.getOperand(4).getReg();
4832 ImageOperands ImOps;
4833 ImOps.Compare =
I.getOperand(5).getReg();
4834 if (
I.getNumOperands() > 6)
4835 ImOps.Offset =
I.getOperand(6).getReg();
4836 if (
I.getNumOperands() > 7)
4837 ImOps.MinLod =
I.getOperand(7).getReg();
4838 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4839 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4842bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
4843 SPIRVTypeInst ResType,
4844 MachineInstr &
I)
const {
4845 Register ImageReg =
I.getOperand(2).getReg();
4846 Register CoordinateReg =
I.getOperand(3).getReg();
4847 Register LodReg =
I.getOperand(4).getReg();
4849 ImageOperands ImOps;
4851 if (
I.getNumOperands() > 5)
4852 ImOps.Offset =
I.getOperand(5).getReg();
4864 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
4865 I.getDebugLoc(),
I, &ImOps);
4868bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
4869 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4870 Register ImageReg =
I.getOperand(2).getReg();
4871 Register SamplerReg =
I.getOperand(3).getReg();
4872 Register CoordinateReg =
I.getOperand(4).getReg();
4873 ImageOperands ImOps;
4874 ImOps.Compare =
I.getOperand(5).getReg();
4875 if (
I.getNumOperands() > 6)
4876 ImOps.Offset =
I.getOperand(6).getReg();
4879 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
4880 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
4883bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
4884 SPIRVTypeInst ResType,
4885 MachineInstr &
I)
const {
4886 Register ImageReg =
I.getOperand(2).getReg();
4887 Register SamplerReg =
I.getOperand(3).getReg();
4888 Register CoordinateReg =
I.getOperand(4).getReg();
4891 "ImageReg is not an image type.");
4896 ComponentOrCompareReg =
I.getOperand(5).getReg();
4897 OffsetReg =
I.getOperand(6).getReg();
4900 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
4904 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
4905 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
4906 Dim != SPIRV::Dim::DIM_Rect) {
4908 "Gather operations are only supported for 2D, Cube, and Rect images.");
4915 if (!loadHandleBeforePosition(
4920 MachineIRBuilder MIRBuilder(
I);
4921 SPIRVTypeInst SampledImageType =
4926 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
4934 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
4936 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
4938 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4943 .
addUse(ComponentOrCompareReg);
4945 uint32_t ImageOperands = 0;
4946 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
4947 if (Dim == SPIRV::Dim::DIM_Cube) {
4949 "Gather operations with offset are not supported for Cube images.");
4953 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
4955 ImageOperands |= SPIRV::ImageOperand::Offset;
4959 if (ImageOperands != 0) {
4960 MIB.
addImm(ImageOperands);
4962 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
4970bool SPIRVInstructionSelector::generateImageReadOrFetch(
4973 const ImageOperands *ImOps)
const {
4976 "ImageReg is not an image type.");
4978 bool IsSignedInteger =
4983 bool IsFetch = (SampledOp.getImm() == 1);
4985 auto AddOperands = [&](MachineInstrBuilder &MIB) {
4986 uint32_t ImageOperandsMask = 0;
4987 if (IsSignedInteger)
4988 ImageOperandsMask |= 0x1000;
4990 if (IsFetch && ImOps) {
4992 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
4993 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
4995 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
4997 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
5001 if (ImageOperandsMask != 0) {
5002 MIB.
addImm(ImageOperandsMask);
5003 if (IsFetch && ImOps) {
5006 if (ImOps->Offset &&
5007 (ImageOperandsMask &
5008 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
5009 MIB.
addUse(*ImOps->Offset);
5015 if (ResultSize == 4) {
5018 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
5025 BMI.constrainAllUses(
TII,
TRI, RBI);
5029 SPIRVTypeInst ReadType = widenTypeToVec4(ResType, Pos);
5033 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
5039 BMI.constrainAllUses(
TII,
TRI, RBI);
5041 if (ResultSize == 1) {
5050 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
5053bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
5054 SPIRVTypeInst ResType,
5055 MachineInstr &
I)
const {
5056 Register ResourcePtr =
I.getOperand(2).getReg();
5058 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
5067 MachineIRBuilder MIRBuilder(
I);
5069 Register IndexReg =
I.getOperand(3).getReg();
5072 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5082bool SPIRVInstructionSelector::selectPushConstantGetPointer(
5083 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5088bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
5089 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5090 Register ObjReg =
I.getOperand(2).getReg();
5091 if (!BuildCOPY(ResVReg, ObjReg,
I))
5101 decorateUsesAsNonUniform(ResVReg);
5105void SPIRVInstructionSelector::decorateUsesAsNonUniform(
5108 while (WorkList.
size() > 0) {
5112 bool IsDecorated =
false;
5114 if (
Use.getOpcode() == SPIRV::OpDecorate &&
5115 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
5121 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
5123 if (ResultReg == CurrentReg)
5131 SPIRV::Decoration::NonUniformEXT, {});
5136bool SPIRVInstructionSelector::extractSubvector(
5138 MachineInstr &InsertionPoint)
const {
5140 [[maybe_unused]] uint64_t InputSize =
5143 assert(InputSize > 1 &&
"The input must be a vector.");
5144 assert(ResultSize > 1 &&
"The result must be a vector.");
5145 assert(ResultSize < InputSize &&
5146 "Cannot extract more element than there are in the input.");
5149 const TargetRegisterClass *ScalarRegClass = GR.
getRegClass(ScalarType);
5150 for (uint64_t
I = 0;
I < ResultSize;
I++) {
5153 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
5162 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
5164 TII.get(SPIRV::OpCompositeConstruct))
5168 for (
Register ComponentReg : ComponentRegisters)
5169 MIB.
addUse(ComponentReg);
5174bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
5175 MachineInstr &
I)
const {
5182 Register ImageReg =
I.getOperand(1).getReg();
5190 Register CoordinateReg =
I.getOperand(2).getReg();
5191 Register DataReg =
I.getOperand(3).getReg();
5194 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
5202Register SPIRVInstructionSelector::buildPointerToResource(
5203 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
5204 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
5205 StringRef Name, MachineIRBuilder MIRBuilder)
const {
5207 if (ArraySize == 1) {
5208 SPIRVTypeInst PtrType =
5211 "SpirvResType did not have an explicit layout.");
5216 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
5217 SPIRVTypeInst VarPointerType =
5220 VarPointerType, Set,
Binding, Name, MIRBuilder);
5222 SPIRVTypeInst ResPointerType =
5235bool SPIRVInstructionSelector::selectFirstBitSet16(
5236 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
5237 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
5239 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
5243 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
5246bool SPIRVInstructionSelector::selectFirstBitSet32(
5248 unsigned BitSetOpcode)
const {
5249 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
5252 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
5259bool SPIRVInstructionSelector::selectFirstBitSet64Overflow(
5261 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
5268 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
5270 MachineIRBuilder MIRBuilder(
I);
5273 SPIRVTypeInst I64x2Type =
5275 SPIRVTypeInst Vec2ResType =
5278 std::vector<Register> PartialRegs;
5281 unsigned CurrentComponent = 0;
5282 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
5288 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5289 TII.get(SPIRV::OpVectorShuffle))
5294 .
addImm(CurrentComponent)
5295 .
addImm(CurrentComponent + 1);
5302 if (!selectFirstBitSet64(SubVecBitSetReg, Vec2ResType,
I, BitSetResult,
5303 BitSetOpcode, SwapPrimarySide))
5306 PartialRegs.push_back(SubVecBitSetReg);
5310 if (CurrentComponent != ComponentCount) {
5316 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
5317 SPIRV::OpVectorExtractDynamic))
5323 if (!selectFirstBitSet64(FinalElemBitSetReg,
BaseType,
I, FinalElemReg,
5324 BitSetOpcode, SwapPrimarySide))
5327 PartialRegs.push_back(FinalElemBitSetReg);
5332 return selectOpWithSrcs(ResVReg, ResType,
I, std::move(PartialRegs),
5333 SPIRV::OpCompositeConstruct);
5336bool SPIRVInstructionSelector::selectFirstBitSet64(
5338 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
5351 if (ComponentCount > 2) {
5352 return selectFirstBitSet64Overflow(ResVReg, ResType,
I, SrcReg,
5353 BitSetOpcode, SwapPrimarySide);
5357 MachineIRBuilder MIRBuilder(
I);
5359 BaseType, 2 * ComponentCount, MIRBuilder,
false);
5363 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
5369 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
5376 bool IsScalarRes = ResType->
getOpcode() != SPIRV::OpTypeVector;
5379 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntZero},
5380 SPIRV::OpVectorExtractDynamic))
5382 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntOne},
5383 SPIRV::OpVectorExtractDynamic))
5387 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5388 TII.get(SPIRV::OpVectorShuffle))
5396 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
5402 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5403 TII.get(SPIRV::OpVectorShuffle))
5411 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
5431 SelectOp = SPIRV::OpSelectSISCond;
5432 AddOp = SPIRV::OpIAddS;
5440 SelectOp = SPIRV::OpSelectVIVCond;
5441 AddOp = SPIRV::OpIAddV;
5451 if (SwapPrimarySide) {
5452 PrimaryReg = LowReg;
5453 SecondaryReg = HighReg;
5454 PrimaryShiftReg = Reg0;
5455 SecondaryShiftReg = Reg32;
5460 if (!selectOpWithSrcs(BReg, BoolType,
I, {PrimaryReg, NegOneReg},
5466 if (!selectOpWithSrcs(TmpReg, ResType,
I, {
BReg, SecondaryReg, PrimaryReg},
5472 if (!selectOpWithSrcs(ValReg, ResType,
I,
5473 {
BReg, SecondaryShiftReg, PrimaryShiftReg}, SelectOp))
5476 return selectOpWithSrcs(ResVReg, ResType,
I, {ValReg, TmpReg}, AddOp);
5479bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
5480 SPIRVTypeInst ResType,
5482 bool IsSigned)
const {
5484 Register OpReg =
I.getOperand(2).getReg();
5487 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
5488 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
5492 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
5494 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
5496 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
5500 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
5504bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
5505 SPIRVTypeInst ResType,
5506 MachineInstr &
I)
const {
5508 Register OpReg =
I.getOperand(2).getReg();
5513 unsigned ExtendOpcode = SPIRV::OpUConvert;
5514 unsigned BitSetOpcode = GL::FindILsb;
5518 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
5520 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
5522 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
5529bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
5530 SPIRVTypeInst ResType,
5531 MachineInstr &
I)
const {
5535 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
5538 .
addUse(
I.getOperand(2).getReg())
5541 unsigned Alignment =
I.getOperand(3).getImm();
5547bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
5548 SPIRVTypeInst ResType,
5549 MachineInstr &
I)
const {
5553 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
5556 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
5559 unsigned Alignment =
I.getOperand(2).getImm();
5566bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
5571 const MachineInstr *PrevI =
I.getPrevNode();
5573 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
5577 .
addMBB(
I.getOperand(0).getMBB())
5582 .
addMBB(
I.getOperand(0).getMBB())
5587bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
5598 const MachineInstr *NextI =
I.getNextNode();
5600 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
5606 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
5608 .
addUse(
I.getOperand(0).getReg())
5609 .
addMBB(
I.getOperand(1).getMBB())
5615bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
5616 MachineInstr &
I)
const {
5618 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
5620 const unsigned NumOps =
I.getNumOperands();
5621 for (
unsigned i = 1; i <
NumOps; i += 2) {
5622 MIB.
addUse(
I.getOperand(i + 0).getReg());
5623 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
5629bool SPIRVInstructionSelector::selectGlobalValue(
5630 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
5632 MachineIRBuilder MIRBuilder(
I);
5633 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
5636 std::string GlobalIdent;
5638 unsigned &
ID = UnnamedGlobalIDs[GV];
5640 ID = UnnamedGlobalIDs.
size();
5641 GlobalIdent =
"__unnamed_" + Twine(
ID).str();
5667 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
5674 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
5679 MachineInstrBuilder MIB1 =
5680 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5683 MachineInstrBuilder MIB2 =
5685 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
5689 GR.
add(ConstVal, MIB2);
5697 MachineInstrBuilder MIB3 =
5698 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5701 GR.
add(ConstVal, MIB3);
5705 assert(NewReg != ResVReg);
5706 return BuildCOPY(ResVReg, NewReg,
I);
5716 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
5722 SPIRVTypeInst ResType =
5726 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
5731 if (
GlobalVar->isExternallyInitialized() &&
5732 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
5733 constexpr unsigned ReadWriteINTEL = 3u;
5736 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
5742bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
5743 SPIRVTypeInst ResType,
5744 MachineInstr &
I)
const {
5746 return selectExtInst(ResVReg, ResType,
I, CL::log10);
5754 MachineIRBuilder MIRBuilder(
I);
5759 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
5762 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
5764 .
add(
I.getOperand(1))
5769 ResType->
getOpcode() == SPIRV::OpTypeFloat);
5771 SPIRVTypeInst SpirvScalarType = ResType->
getOpcode() == SPIRV::OpTypeVector
5779 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
5780 ? SPIRV::OpVectorTimesScalar
5791bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
5792 SPIRVTypeInst ResType,
5793 MachineInstr &
I)
const {
5796 return selectExtInst(ResVReg, ResType,
I, CL::pown);
5802 Register ExpReg =
I.getOperand(2).getReg();
5804 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
5805 SPIRV::OpConvertSToF))
5807 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
5814bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
5815 SPIRVTypeInst ResType,
5816 MachineInstr &
I)
const {
5832 MachineIRBuilder MIRBuilder(
I);
5835 ResType, MIRBuilder, SPIRV::StorageClass::Function);
5847 MachineBasicBlock &EntryBB =
I.getMF()->front();
5851 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
5854 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
5860 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
5863 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
5866 .
add(
I.getOperand(
I.getNumExplicitDefs()))
5870 Register IntegralPartReg =
I.getOperand(1).getReg();
5871 if (IntegralPartReg.
isValid()) {
5873 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
5884 assert(
false &&
"GLSL::Modf is deprecated.");
5895bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
5896 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
5897 SPIRVTypeInst ResType, MachineInstr &
I)
const {
5898 MachineIRBuilder MIRBuilder(
I);
5899 const SPIRVTypeInst Vec3Ty =
5902 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
5914 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
5918 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
5924 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
5931 assert(
I.getOperand(2).isReg());
5932 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
5936 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
5947bool SPIRVInstructionSelector::loadBuiltinInputID(
5948 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
5949 SPIRVTypeInst ResType, MachineInstr &
I)
const {
5950 MachineIRBuilder MIRBuilder(
I);
5952 ResType, MIRBuilder, SPIRV::StorageClass::Input);
5967 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
5971 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
5980SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
5981 MachineInstr &
I)
const {
5982 MachineIRBuilder MIRBuilder(
I);
5983 if (
Type->getOpcode() != SPIRV::OpTypeVector)
5987 if (VectorSize == 4)
5995bool SPIRVInstructionSelector::loadHandleBeforePosition(
5996 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
5997 MachineInstr &Pos)
const {
6000 Intrinsic::spv_resource_handlefrombinding);
6008 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
6009 MachineIRBuilder MIRBuilder(HandleDef);
6010 SPIRVTypeInst VarType = ResType;
6011 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
6013 if (IsStructuredBuffer) {
6019 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
6020 ArraySize, IndexReg, Name, MIRBuilder);
6024 uint32_t LoadOpcode =
6025 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
6035void SPIRVInstructionSelector::errorIfInstrOutsideShader(
6036 MachineInstr &
I)
const {
6038 std::string DiagMsg;
6039 raw_string_ostream OS(DiagMsg);
6040 I.print(OS,
true,
false,
false,
false);
6041 DiagMsg +=
" is only supported in shaders.\n";
6047InstructionSelector *
6051 return new SPIRVInstructionSelector(TM, Subtarget, RBI);
MachineInstrBuilder & UseMI
#define GET_GLOBALISEL_PREDICATES_INIT
#define GET_GLOBALISEL_TEMPORARIES_INIT
assert(UImm &&(UImm !=~static_cast< T >(0)) &&"Invalid immediate!")
This file declares a class to represent arbitrary precision floating point values and provide a varie...
static bool selectUnmergeValues(MachineInstrBuilder &MIB, const ARMBaseInstrInfo &TII, MachineRegisterInfo &MRI, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI)
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
static 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
Register const TargetRegisterInfo * TRI
Promote Memory to Register
MachineInstr unsigned OpIdx
uint64_t IntrinsicInst * II
static StringRef getName(Value *V)
static unsigned getFCmpOpcode(CmpInst::Predicate Pred, unsigned Size)
static 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)
unsigned getVectorSizeOrOne(SPIRVTypeInst Type)
bool isDead(const MachineInstr &MI, const MachineRegisterInfo &MRI)
BaseType
A given derived pointer can have multiple base pointers through phi/selects.
static TableGen::Emitter::Opt Y("gen-skeleton-entry", EmitSkeleton, "Generate example skeleton entry")
static const fltSemantics & IEEEsingle()
static const fltSemantics & IEEEdouble()
static const fltSemantics & IEEEhalf()
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.
ArrayRef - 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.
Intrinsic::ID getIntrinsicID() const
unsigned getAddressSpace() const
Module * getParent()
Get the module that this global value is contained inside of...
@ InternalLinkage
Rename collisions when linking (static functions).
static LLVM_ABI IntegerType * get(LLVMContext &C, unsigned NumBits)
This static method is the primary way of constructing an IntegerType.
static constexpr LLT scalar(unsigned SizeInBits)
Get a low-level scalar or aggregate "bag of bits".
constexpr bool isValid() const
constexpr uint16_t getNumElements() const
Returns the number of elements in a vector LLT.
constexpr bool isVector() const
static constexpr LLT pointer(unsigned AddressSpace, unsigned SizeInBits)
Get a low-level pointer in the given address space.
static constexpr LLT fixed_vector(unsigned NumElements, unsigned ScalarSizeInBits)
Get a low-level fixed-width vector of some number of elements and element width.
int getNumber() const
MachineBasicBlocks are uniquely numbered at the function level, unless they're not in a MachineFuncti...
LLVM_ABI iterator getFirstNonPHI()
Returns a pointer to the first instruction in this block that is not a PHINode instruction.
const MachineFunction * getParent() const
Return the MachineFunction containing this basic block.
MachineInstrBundleIterator< MachineInstr > iterator
MachineRegisterInfo & getRegInfo()
getRegInfo - Return information about the registers currently in use.
Function & getFunction()
Return the LLVM function that this machine code represents.
Helper class to build MachineInstr.
MachineBasicBlock::iterator getInsertPt()
Current insertion point for new instructions.
MachineInstrBuilder buildInstr(unsigned Opcode)
Build and insert <empty> = Opcode <empty>.
MachineFunction & getMF()
Getter for the function we currently build.
MachineRegisterInfo * getMRI()
Getter for MRI.
void constrainAllUses(const TargetInstrInfo &TII, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI) const
const MachineInstrBuilder & addUse(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a virtual register use operand.
const MachineInstrBuilder & addReg(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a new virtual register operand.
const MachineInstrBuilder & addImm(int64_t Val) const
Add a new immediate operand.
const MachineInstrBuilder & add(const MachineOperand &MO) const
const MachineInstrBuilder & addMBB(MachineBasicBlock *MBB, unsigned TargetFlags=0) const
const MachineInstrBuilder & addDef(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a virtual register definition operand.
const MachineInstrBuilder & setMIFlags(unsigned Flags) const
MachineInstr * getInstr() const
If conversion operators fail, use this method to get the MachineInstr explicitly.
Representation of each machine instruction.
unsigned getOpcode() const
Returns the opcode of this MachineInstr.
const MachineBasicBlock * getParent() const
unsigned getNumOperands() const
Retuns the total number of operands.
LLVM_ABI unsigned getNumExplicitOperands() const
Returns the number of non-implicit operands.
LLVM_ABI unsigned getNumExplicitDefs() const
Returns the number of non-implicit definitions.
LLVM_ABI void emitGenericError(const Twine &ErrMsg) const
LLVM_ABI const MachineFunction * getMF() const
Return the function that contains the basic block that this instruction belongs to.
const DebugLoc & getDebugLoc() const
Returns the debug location id of this MachineInstr.
const MachineOperand & getOperand(unsigned i) const
A description of a memory reference used in the backend.
@ MOVolatile
The memory access is volatile.
@ MONonTemporal
The memory access is non-temporal.
bool isReg() const
isReg - Tests if this is a MO_Register operand.
MachineBasicBlock * getMBB() const
Register getReg() const
getReg - Returns the register number.
MachineRegisterInfo - Keep track of information for virtual and physical registers,...
defusechain_instr_iterator< true, false, false, true > use_instr_iterator
use_instr_iterator/use_instr_begin/use_instr_end - Walk all uses of the specified register,...
const TargetRegisterClass * getRegClass(Register Reg) const
Return the register class of the specified virtual register.
LLVM_ABI MachineInstr * getVRegDef(Register Reg) const
getVRegDef - Return the machine instr that defines the specified virtual register or null if none is ...
use_instr_iterator use_instr_begin(Register RegNo) const
static def_instr_iterator def_instr_end()
defusechain_instr_iterator< false, true, false, true > def_instr_iterator
def_instr_iterator/def_instr_begin/def_instr_end - Walk all defs of the specified register,...
LLVM_ABI Register createVirtualRegister(const TargetRegisterClass *RegClass, StringRef Name="")
createVirtualRegister - Create and return a new virtual register in the function with the specified r...
def_instr_iterator def_instr_begin(Register RegNo) const
LLT getType(Register Reg) const
Get the low-level type of Reg or LLT{} if Reg is not a generic (target independent) virtual register.
static use_instr_iterator use_instr_end()
iterator_range< use_instr_nodbg_iterator > use_nodbg_instructions(Register Reg) const
LLVM_ABI void setType(Register VReg, LLT Ty)
Set the low-level type of VReg to Ty.
const MachineFunction & getMF() const
LLVM_ABI void setRegClass(Register Reg, const TargetRegisterClass *RC)
setRegClass - Set the register class of the specified virtual register.
LLVM_ABI Register createGenericVirtualRegister(LLT Ty, StringRef Name="")
Create and return a new generic virtual register with low-level type Ty.
const TargetRegisterClass * getRegClassOrNull(Register Reg) const
Return the register class of Reg, or null if Reg has not been assigned a register class yet.
iterator_range< use_instr_iterator > use_instructions(Register Reg) const
unsigned getNumVirtRegs() const
getNumVirtRegs - Return the number of virtual registers created.
LLVM_ABI void replaceRegWith(Register FromReg, Register ToReg)
replaceRegWith - Replace all instances of FromReg with ToReg in the machine function.
Analysis providing profile information.
Holds all the information related to register banks.
Wrapper class representing virtual and physical registers.
constexpr bool isValid() const
constexpr bool isPhysical() const
Return true if the specified register number is in the physical register namespace.
bool isScalarOrVectorSigned(SPIRVTypeInst Type) const
SPIRVTypeInst getOrCreateOpTypeSampledImage(SPIRVTypeInst ImageType, MachineIRBuilder &MIRBuilder)
void assignSPIRVTypeToVReg(SPIRVTypeInst Type, Register VReg, const MachineFunction &MF)
const TargetRegisterClass * getRegClass(SPIRVTypeInst SpvType) const
MachineInstr * getOrAddMemAliasingINTELInst(MachineIRBuilder &MIRBuilder, const MDNode *AliasingListMD)
bool isAggregateType(SPIRVTypeInst Type) const
unsigned getScalarOrVectorBitWidth(SPIRVTypeInst Type) const
SPIRVTypeInst getOrCreateSPIRVIntegerType(unsigned BitWidth, MachineIRBuilder &MIRBuilder)
SPIRVTypeInst getOrCreateSPIRVVectorType(SPIRVTypeInst BaseType, unsigned NumElements, MachineIRBuilder &MIRBuilder, bool EmitIR)
Register buildGlobalVariable(Register Reg, SPIRVTypeInst BaseType, StringRef Name, const GlobalValue *GV, SPIRV::StorageClass::StorageClass Storage, const MachineInstr *Init, bool IsConst, const std::optional< SPIRV::LinkageType::LinkageType > &LinkageType, MachineIRBuilder &MIRBuilder, bool IsInstSelector)
SPIRVTypeInst getResultType(Register VReg, MachineFunction *MF=nullptr)
unsigned getScalarOrVectorComponentCount(Register VReg) const
const Type * getTypeForSPIRVType(SPIRVTypeInst Ty) const
bool isBitcastCompatible(SPIRVTypeInst Type1, SPIRVTypeInst Type2) const
unsigned getPointerSize() const
Register getOrCreateConstFP(APFloat Val, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
LLT getRegType(SPIRVTypeInst SpvType) const
void invalidateMachineInstr(MachineInstr *MI)
SPIRVTypeInst getOrCreateSPIRVBoolType(MachineIRBuilder &MIRBuilder, bool EmitIR)
SPIRVTypeInst getOrCreateSPIRVPointerType(const Type *BaseType, MachineIRBuilder &MIRBuilder, SPIRV::StorageClass::StorageClass SC)
bool isScalarOfType(Register VReg, unsigned TypeOpcode) const
Register getSPIRVTypeID(SPIRVTypeInst SpirvType) const
Register getOrCreateConstInt(uint64_t Val, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
Register getOrCreateConstIntArray(uint64_t Val, size_t Num, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII)
bool findValueAttrs(const MachineInstr *Key, Type *&Ty, StringRef &Name)
SPIRVTypeInst retrieveScalarOrVectorIntType(SPIRVTypeInst Type) const
Register getOrCreateGlobalVariableWithBinding(SPIRVTypeInst VarType, uint32_t Set, uint32_t Binding, StringRef Name, MachineIRBuilder &MIRBuilder)
SPIRVTypeInst changePointerStorageClass(SPIRVTypeInst PtrType, SPIRV::StorageClass::StorageClass SC, MachineInstr &I)
Register getOrCreateConstVector(uint64_t Val, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
Register buildConstantFP(APFloat Val, MachineIRBuilder &MIRBuilder, SPIRVTypeInst SpvType=nullptr)
void addGlobalObject(const Value *V, const MachineFunction *MF, Register R)
SPIRVTypeInst getScalarOrVectorComponentType(SPIRVTypeInst Type) const
void recordFunctionPointer(const MachineOperand *MO, const Function *F)
SPIRVTypeInst getOrCreateSPIRVFloatType(unsigned BitWidth, MachineInstr &I, const SPIRVInstrInfo &TII)
SPIRVTypeInst getPointeeType(SPIRVTypeInst PtrType)
SPIRVTypeInst getOrCreateSPIRVType(const Type *Type, MachineInstr &I, SPIRV::AccessQualifier::AccessQualifier AQ, bool EmitIR)
bool isScalarOrVectorOfType(Register VReg, unsigned TypeOpcode) const
MachineFunction * setCurrentFunc(MachineFunction &MF)
Register getOrCreateConstNullPtr(MachineIRBuilder &MIRBuilder, SPIRVTypeInst SpvType)
SPIRVTypeInst getSPIRVTypeForVReg(Register VReg, const MachineFunction *MF=nullptr) const
Type * getDeducedGlobalValueType(const GlobalValue *Global)
Register getOrCreateUndef(MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII)
SPIRV::StorageClass::StorageClass getPointerStorageClass(Register VReg) const
bool erase(const MachineInstr *MI)
bool add(SPIRV::IRHandle Handle, const MachineInstr *MI)
Register find(SPIRV::IRHandle Handle, const MachineFunction *MF)
bool isPhysicalSPIRV() const
bool isAtLeastSPIRVVer(VersionTuple VerToCompareTo) const
bool canUseExtInstSet(SPIRV::InstructionSet::InstructionSet E) const
bool isLogicalSPIRV() const
bool canUseExtension(SPIRV::Extension::Extension E) const
bool erase(PtrType Ptr)
Remove pointer from the set.
std::pair< iterator, bool > insert(PtrType Ptr)
Inserts Ptr if and only if there is no element in the container equal to Ptr.
bool contains(ConstPtrType Ptr) const
SmallPtrSet - This class implements a set which is optimized for holding SmallSize or less elements.
reference emplace_back(ArgTypes &&... Args)
void reserve(size_type N)
void push_back(const T &Elt)
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
constexpr size_t size() const
size - Get the string size.
static LLVM_ABI StructType * get(LLVMContext &Context, ArrayRef< Type * > Elements, bool isPacked=false)
This static method is the primary way to create a literal StructType.
Twine - A lightweight data structure for efficiently representing the concatenation of temporary valu...
The instances of the Type class are immutable: once they are created, they are never changed.
@ HalfTyID
16-bit floating point type
@ FloatTyID
32-bit floating point type
@ DoubleTyID
64-bit floating point type
Type * getScalarType() const
If this is a vector type, return the element type, otherwise return 'this'.
bool isStructTy() const
True if this is an instance of StructType.
bool isAggregateType() const
Return true if the type is an aggregate type.
TypeID getTypeID() const
Return the type id for the type.
Value * getOperand(unsigned i) const
LLVM_ABI StringRef getName() const
Return a constant reference to the value's name.
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.
unsigned ID
LLVM IR allows to use arbitrary numbers as calling convention identifiers.
Scope
Defines the scope in which this symbol should be visible: Default – Visible in the public interface o...
NodeAddr< DefNode * > Def
NodeAddr< InstrNode * > Instr
NodeAddr< UseNode * > Use
BaseReg
Stack frame base register. Bit 0 of FREInfo.Info.
This is an optimization pass for GlobalISel generic memory operations.
void buildOpName(Register Target, const StringRef &Name, MachineIRBuilder &MIRBuilder)
bool all_of(R &&range, UnaryPredicate P)
Provide wrappers to std::all_of which take ranges instead of having to pass begin/end explicitly.
int64_t getIConstValSext(Register ConstReg, const MachineRegisterInfo *MRI)
MachineInstrBuilder BuildMI(MachineFunction &MF, const MIMetadata &MIMD, const MCInstrDesc &MCID)
Builder interface. Specify how to create the initial instruction itself.
bool isTypeFoldingSupported(unsigned Opcode)
decltype(auto) dyn_cast(const From &Val)
dyn_cast<X> - Return the argument parameter cast to the specified type.
FunctionAddr VTableAddr uintptr_t uintptr_t Int32Ty
void addNumImm(const APInt &Imm, MachineInstrBuilder &MIB)
LLVM_ABI void salvageDebugInfo(const MachineRegisterInfo &MRI, MachineInstr &MI)
Assuming the instruction MI is going to be deleted, attempt to salvage debug users of MI by writing t...
LLVM_ABI void constrainSelectedInstRegOperands(MachineInstr &I, const TargetInstrInfo &TII, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI)
Mutate the newly-selected instruction I to constrain its (possibly generic) virtual register operands...
bool isPreISelGenericOpcode(unsigned Opcode)
Check whether the given Opcode is a generic opcode that is not supposed to appear after ISel.
Register createVirtualRegister(SPIRVTypeInst SpvType, SPIRVGlobalRegistry *GR, MachineRegisterInfo *MRI, const MachineFunction &MF)
unsigned getArrayComponentCount(const MachineRegisterInfo *MRI, const MachineInstr *ResType)
uint64_t getIConstVal(Register ConstReg, const MachineRegisterInfo *MRI)
SmallVector< MachineInstr *, 4 > createContinuedInstructions(MachineIRBuilder &MIRBuilder, unsigned Opcode, unsigned MinWC, unsigned ContinuedOpcode, ArrayRef< Register > Args, Register ReturnRegister, Register TypeID)
SPIRV::MemorySemantics::MemorySemantics getMemSemanticsForStorageClass(SPIRV::StorageClass::StorageClass SC)
constexpr unsigned storageClassToAddressSpace(SPIRV::StorageClass::StorageClass SC)
MachineBasicBlock::iterator getFirstValidInstructionInsertPoint(MachineBasicBlock &BB)
void buildOpDecorate(Register Reg, MachineIRBuilder &MIRBuilder, SPIRV::Decoration::Decoration Dec, const std::vector< uint32_t > &DecArgs, StringRef StrImm)
MachineBasicBlock::iterator getOpVariableMBBIt(MachineInstr &I)
MachineInstr * getImm(const MachineOperand &MO, const MachineRegisterInfo *MRI)
Type * toTypedPointer(Type *Ty)
LLVM_ABI raw_ostream & dbgs()
dbgs() - This returns a reference to a raw_ostream for debugging messages.
LLVM_ABI void report_fatal_error(Error Err, bool gen_crash_diag=true)
constexpr bool isGenericCastablePtr(SPIRV::StorageClass::StorageClass SC)
class LLVM_GSL_OWNER SmallVector
Forward declaration of SmallVector so that calculateSmallVectorDefaultInlinedElements can reference s...
MachineInstr * passCopy(MachineInstr *Def, const MachineRegisterInfo *MRI)
bool isa(const From &Val)
isa<X> - Return true if the parameter to the template is an instance of one of the template type argu...
std::optional< SPIRV::LinkageType::LinkageType > getSpirvLinkageTypeFor(const SPIRVSubtarget &ST, const GlobalValue &GV)
LLVM_ABI raw_fd_ostream & errs()
This returns a reference to a raw_ostream for standard error.
SPIRV::StorageClass::StorageClass addressSpaceToStorageClass(unsigned AddrSpace, const SPIRVSubtarget &STI)
AtomicOrdering
Atomic ordering for LLVM's memory model.
SPIRV::Scope::Scope getMemScope(LLVMContext &Ctx, SyncScope::ID Id)
InstructionSelector * createSPIRVInstructionSelector(const SPIRVTargetMachine &TM, const SPIRVSubtarget &Subtarget, const RegisterBankInfo &RBI)
std::string getStringValueFromReg(Register Reg, MachineRegisterInfo &MRI)
int64_t foldImm(const MachineOperand &MO, const MachineRegisterInfo *MRI)
DWARFExpression::Operation Op
ArrayRef(const T &OneElt) -> ArrayRef< T >
MachineInstr * getDefInstrMaybeConstant(Register &ConstReg, const MachineRegisterInfo *MRI)
constexpr unsigned BitWidth
decltype(auto) cast(const From &Val)
cast<X> - Return the argument parameter cast to the specified type.
bool hasInitializer(const GlobalVariable *GV)
void addStringImm(const StringRef &Str, MCInst &Inst)
MachineInstr * getVRegDef(MachineRegisterInfo &MRI, Register Reg)
SPIRV::MemorySemantics::MemorySemantics getMemSemantics(AtomicOrdering Ord)
std::string getLinkStringForBuiltIn(SPIRV::BuiltIn::BuiltIn BuiltInValue)
LLVM_ABI bool isTriviallyDead(const MachineInstr &MI, const MachineRegisterInfo &MRI)
Check whether an instruction MI is dead: it only defines dead virtual registers, and doesn't have oth...