34#include "llvm/IR/IntrinsicsSPIRV.h"
40#define DEBUG_TYPE "spirv-isel"
47 std::vector<std::pair<SPIRV::InstructionSet::InstructionSet, uint32_t>>;
52 std::optional<Register> Bias;
53 std::optional<Register>
Offset;
54 std::optional<Register> MinLod;
55 std::optional<Register> GradX;
56 std::optional<Register> GradY;
57 std::optional<Register> Lod;
58 std::optional<Register> Compare;
65 bool IsScalar =
false;
68llvm::SPIRV::SelectionControl::SelectionControl
69getSelectionOperandForImm(
int Imm) {
71 return SPIRV::SelectionControl::Flatten;
73 return SPIRV::SelectionControl::DontFlatten;
75 return SPIRV::SelectionControl::None;
79#define GET_GLOBALISEL_PREDICATE_BITSET
80#include "SPIRVGenGlobalISel.inc"
81#undef GET_GLOBALISEL_PREDICATE_BITSET
108#define GET_GLOBALISEL_PREDICATES_DECL
109#include "SPIRVGenGlobalISel.inc"
110#undef GET_GLOBALISEL_PREDICATES_DECL
112#define GET_GLOBALISEL_TEMPORARIES_DECL
113#include "SPIRVGenGlobalISel.inc"
114#undef GET_GLOBALISEL_TEMPORARIES_DECL
138 unsigned BitSetOpcode)
const;
142 unsigned BitSetOpcode)
const;
146 unsigned BitSetOpcode,
bool SwapPrimarySide)
const;
153 unsigned Opcode)
const;
156 unsigned Opcode)
const;
178 unsigned NewOpcode,
unsigned NegateOpcode = 0)
const;
189 unsigned OpType)
const;
252 unsigned Opcode)
const;
256 unsigned Opcode)
const;
260 unsigned Opcode)
const;
262 template <
bool Signed>
265 template <
bool Signed>
272 template <
typename PickOpcodeFn>
275 PickOpcodeFn &&PickOpcode)
const;
292 template <
typename PickOpcodeFn>
295 PickOpcodeFn &&PickOpcode)
const;
313 bool IsSigned)
const;
315 bool IsSigned,
unsigned Opcode)
const;
317 bool IsSigned)
const;
323 bool IsSigned)
const;
362 GL::GLSLExtInst GLInst,
bool setMIFlags =
true,
363 bool useMISrc =
true,
365 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
366 CL::OpenCLExtInst CLInst,
bool setMIFlags =
true,
367 bool useMISrc =
true,
369 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
370 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
371 bool setMIFlags =
true,
bool useMISrc =
true,
373 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
374 const ExtInstList &ExtInsts,
bool setMIFlags =
true,
375 bool useMISrc =
true,
378 bool selectLog10(
Register ResVReg, SPIRVTypeInst ResType,
379 MachineInstr &
I)
const;
381 bool selectFpowi(
Register ResVReg, SPIRVTypeInst ResType,
382 MachineInstr &
I)
const;
384 bool selectSaturate(
Register ResVReg, SPIRVTypeInst ResType,
385 MachineInstr &
I)
const;
387 bool selectWaveOpInst(
Register ResVReg, SPIRVTypeInst ResType,
388 MachineInstr &
I,
unsigned Opcode)
const;
390 bool selectBarrierInst(MachineInstr &
I,
unsigned Scope,
unsigned MemSem,
391 bool WithGroupSync)
const;
393 bool selectWaveActiveCountBits(
Register ResVReg, SPIRVTypeInst ResType,
394 MachineInstr &
I)
const;
396 bool selectWaveActiveAllEqual(
Register ResVReg, SPIRVTypeInst ResType,
397 MachineInstr &
I)
const;
401 bool selectHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
402 MachineInstr &
I)
const;
404 bool selectCounterHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
405 MachineInstr &
I)
const;
407 bool selectReadImageIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
408 MachineInstr &
I)
const;
409 bool selectGetDimensionsIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
410 MachineInstr &
I)
const;
411 bool selectGetDimensionsLevelsIntrinsic(
Register &ResVReg,
412 SPIRVTypeInst ResType,
413 MachineInstr &
I)
const;
414 bool selectGetDimensionsMSIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
415 MachineInstr &
I)
const;
418 std::optional<Register> LodReg = std::nullopt)
const;
419 bool selectSampleBasicIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
420 MachineInstr &
I)
const;
421 bool selectCalculateLodIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
422 MachineInstr &
I)
const;
423 bool selectSampleBiasIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
424 MachineInstr &
I)
const;
425 bool selectSampleGradIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
426 MachineInstr &
I)
const;
427 bool selectSampleLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
428 MachineInstr &
I)
const;
429 bool selectLoadLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
430 MachineInstr &
I)
const;
431 bool selectSampleCmpIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
432 MachineInstr &
I)
const;
433 bool selectSampleCmpLevelZeroIntrinsic(
Register &ResVReg,
434 SPIRVTypeInst ResType,
435 MachineInstr &
I)
const;
436 bool selectGatherIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
437 MachineInstr &
I)
const;
438 bool selectImageWriteIntrinsic(MachineInstr &
I)
const;
439 bool selectResourceGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
440 MachineInstr &
I)
const;
441 bool selectPushConstantGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
442 MachineInstr &
I)
const;
443 bool selectResourceNonUniformIndex(
Register &ResVReg, SPIRVTypeInst ResType,
444 MachineInstr &
I)
const;
445 bool selectModf(
Register ResVReg, SPIRVTypeInst ResType,
446 MachineInstr &
I)
const;
447 bool selectUpdateCounter(
Register &ResVReg, SPIRVTypeInst ResType,
448 MachineInstr &
I)
const;
449 bool selectFrexp(
Register ResVReg, SPIRVTypeInst ResType,
450 MachineInstr &
I)
const;
451 bool selectSincos(
Register ResVReg, SPIRVTypeInst ResType,
452 MachineInstr &
I)
const;
453 bool selectExp10(
Register ResVReg, SPIRVTypeInst ResType,
454 MachineInstr &
I)
const;
455 bool selectDerivativeInst(
Register ResVReg, SPIRVTypeInst ResType,
456 MachineInstr &
I,
const unsigned DPdOpCode)
const;
458 Register buildI32Constant(uint32_t Val, MachineInstr &
I,
459 SPIRVTypeInst ResType =
nullptr)
const;
461 Register buildZerosVal(SPIRVTypeInst ResType, MachineInstr &
I)
const;
462 bool isScalarOrVectorIntConstantZero(
Register Reg)
const;
463 Register buildZerosValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
465 MachineInstr &
I)
const;
466 Register buildOnesValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
468 bool wrapIntoSpecConstantOp(MachineInstr &
I,
471 Register getUcharPtrTypeReg(MachineInstr &
I,
472 SPIRV::StorageClass::StorageClass SC)
const;
473 MachineInstrBuilder buildSpecConstantOp(MachineInstr &
I,
Register Dest,
475 uint32_t Opcode)
const;
476 MachineInstrBuilder buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
477 SPIRVTypeInst SrcPtrTy)
const;
478 Register buildPointerToResource(SPIRVTypeInst ResType,
479 SPIRV::StorageClass::StorageClass SC,
480 uint32_t Set, uint32_t
Binding,
481 uint32_t ArraySize,
Register IndexReg,
483 MachineIRBuilder MIRBuilder)
const;
484 SPIRVTypeInst widenTypeToVec4(SPIRVTypeInst
Type, MachineInstr &
I)
const;
485 bool extractSubvector(
Register &ResVReg, SPIRVTypeInst ResType,
486 Register &ReadReg, MachineInstr &InsertionPoint)
const;
487 bool generateImageReadOrFetch(
Register &ResVReg, SPIRVTypeInst ResType,
490 const ImageOperands *ImOps =
nullptr)
const;
491 bool generateSampleImage(
Register ResVReg, SPIRVTypeInst ResType,
493 Register CoordinateReg,
const ImageOperands &ImOps,
496 bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
497 Register ResVReg, SPIRVTypeInst ResType,
498 MachineInstr &
I)
const;
499 bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
500 Register ResVReg, SPIRVTypeInst ResType,
501 MachineInstr &
I)
const;
502 bool loadHandleBeforePosition(
Register &HandleReg, SPIRVTypeInst ResType,
503 GIntrinsic &HandleDef, MachineInstr &Pos)
const;
504 void decorateUsesAsNonUniform(
Register &NonUniformReg)
const;
505 void errorIfInstrOutsideShader(MachineInstr &
I)
const;
507 std::optional<SplitParts> splitEvenOddLanes(
Register PopCountReg,
508 unsigned ComponentCount,
510 SPIRVTypeInst I32Type)
const;
513 handle64BitOverflow(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
514 Register SrcReg,
unsigned int Opcode,
515 std::function<
bool(
Register, SPIRVTypeInst,
516 MachineInstr &,
Register,
unsigned)>
520bool sampledTypeIsSignedInteger(
const llvm::Type *HandleType) {
522 if (
TET->getTargetExtName() ==
"spirv.Image") {
525 assert(
TET->getTargetExtName() ==
"spirv.SignedImage");
526 return TET->getTypeParameter(0)->isIntegerTy();
530#define GET_GLOBALISEL_IMPL
531#include "SPIRVGenGlobalISel.inc"
532#undef GET_GLOBALISEL_IMPL
538 TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
541#include
"SPIRVGenGlobalISel.inc"
544#include
"SPIRVGenGlobalISel.inc"
556 InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
560void SPIRVInstructionSelector::resetVRegsType(MachineFunction &MF) {
561 if (HasVRegsReset == &MF)
576 for (
const auto &
MBB : MF) {
577 for (
const auto &
MI :
MBB) {
580 if (
MI.getOpcode() != SPIRV::ASSIGN_TYPE)
584 LLT DstType = MRI.
getType(DstReg);
586 LLT SrcType = MRI.
getType(SrcReg);
587 if (DstType != SrcType)
592 if (DstRC != SrcRC && SrcRC)
604 while (!Stack.empty()) {
609 switch (
MI->getOpcode()) {
610 case TargetOpcode::G_INTRINSIC:
611 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
612 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS: {
615 if (IntrID != Intrinsic::spv_const_composite &&
616 IntrID != Intrinsic::spv_undef)
620 case TargetOpcode::G_BUILD_VECTOR:
621 case TargetOpcode::G_SPLAT_VECTOR:
623 i < OpDef->getNumOperands(); i++) {
628 Stack.push_back(OpNestedDef);
631 case TargetOpcode::G_CONSTANT:
632 case TargetOpcode::G_FCONSTANT:
633 case TargetOpcode::G_IMPLICIT_DEF:
634 case SPIRV::OpConstantTrue:
635 case SPIRV::OpConstantFalse:
636 case SPIRV::OpConstantI:
637 case SPIRV::OpConstantF:
638 case SPIRV::OpConstantComposite:
639 case SPIRV::OpConstantCompositeContinuedINTEL:
640 case SPIRV::OpConstantSampler:
641 case SPIRV::OpConstantNull:
643 case SPIRV::OpConstantFunctionPointerINTEL:
670 case Intrinsic::spv_all:
671 case Intrinsic::spv_alloca:
672 case Intrinsic::spv_any:
673 case Intrinsic::spv_bitcast:
674 case Intrinsic::spv_const_composite:
675 case Intrinsic::spv_cross:
676 case Intrinsic::spv_degrees:
677 case Intrinsic::spv_distance:
678 case Intrinsic::spv_extractelt:
679 case Intrinsic::spv_extractv:
680 case Intrinsic::spv_faceforward:
681 case Intrinsic::spv_fdot:
682 case Intrinsic::spv_firstbitlow:
683 case Intrinsic::spv_firstbitshigh:
684 case Intrinsic::spv_firstbituhigh:
685 case Intrinsic::spv_frac:
686 case Intrinsic::spv_gep:
687 case Intrinsic::spv_global_offset:
688 case Intrinsic::spv_global_size:
689 case Intrinsic::spv_group_id:
690 case Intrinsic::spv_insertelt:
691 case Intrinsic::spv_insertv:
692 case Intrinsic::spv_isinf:
693 case Intrinsic::spv_isnan:
694 case Intrinsic::spv_lerp:
695 case Intrinsic::spv_length:
696 case Intrinsic::spv_normalize:
697 case Intrinsic::spv_num_subgroups:
698 case Intrinsic::spv_num_workgroups:
699 case Intrinsic::spv_ptrcast:
700 case Intrinsic::spv_radians:
701 case Intrinsic::spv_reflect:
702 case Intrinsic::spv_refract:
703 case Intrinsic::spv_resource_getpointer:
704 case Intrinsic::spv_resource_handlefrombinding:
705 case Intrinsic::spv_resource_handlefromimplicitbinding:
706 case Intrinsic::spv_resource_nonuniformindex:
707 case Intrinsic::spv_resource_sample:
708 case Intrinsic::spv_rsqrt:
709 case Intrinsic::spv_saturate:
710 case Intrinsic::spv_sdot:
711 case Intrinsic::spv_sign:
712 case Intrinsic::spv_smoothstep:
713 case Intrinsic::spv_step:
714 case Intrinsic::spv_subgroup_id:
715 case Intrinsic::spv_subgroup_local_invocation_id:
716 case Intrinsic::spv_subgroup_max_size:
717 case Intrinsic::spv_subgroup_size:
718 case Intrinsic::spv_thread_id:
719 case Intrinsic::spv_thread_id_in_group:
720 case Intrinsic::spv_udot:
721 case Intrinsic::spv_undef:
722 case Intrinsic::spv_value_md:
723 case Intrinsic::spv_workgroup_size:
735 case SPIRV::OpTypeVoid:
736 case SPIRV::OpTypeBool:
737 case SPIRV::OpTypeInt:
738 case SPIRV::OpTypeFloat:
739 case SPIRV::OpTypeVector:
740 case SPIRV::OpTypeMatrix:
741 case SPIRV::OpTypeImage:
742 case SPIRV::OpTypeSampler:
743 case SPIRV::OpTypeSampledImage:
744 case SPIRV::OpTypeArray:
745 case SPIRV::OpTypeRuntimeArray:
746 case SPIRV::OpTypeStruct:
747 case SPIRV::OpTypeOpaque:
748 case SPIRV::OpTypePointer:
749 case SPIRV::OpTypeFunction:
750 case SPIRV::OpTypeEvent:
751 case SPIRV::OpTypeDeviceEvent:
752 case SPIRV::OpTypeReserveId:
753 case SPIRV::OpTypeQueue:
754 case SPIRV::OpTypePipe:
755 case SPIRV::OpTypeForwardPointer:
756 case SPIRV::OpTypePipeStorage:
757 case SPIRV::OpTypeNamedBarrier:
758 case SPIRV::OpTypeAccelerationStructureNV:
759 case SPIRV::OpTypeCooperativeMatrixNV:
760 case SPIRV::OpTypeCooperativeMatrixKHR:
770 if (
MI.getNumDefs() == 0)
773 for (
const auto &MO :
MI.all_defs()) {
775 if (
Reg.isPhysical()) {
780 if (
UseMI.getOpcode() != SPIRV::OpName) {
787 if (
MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE ||
MI.isFakeUse() ||
788 MI.isLifetimeMarker()) {
791 <<
"Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
802 if (
MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
803 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
806 LLVM_DEBUG(
dbgs() <<
"Dead: Intrinsic with no real side effects.\n");
811 if (
MI.mayStore() ||
MI.isCall() ||
812 (
MI.mayLoad() &&
MI.hasOrderedMemoryRef()) ||
MI.isPosition() ||
813 MI.isDebugInstr() ||
MI.isTerminator() ||
MI.isJumpTableDebugInfo()) {
814 LLVM_DEBUG(
dbgs() <<
"Not dead: instruction has side effects.\n");
825 LLVM_DEBUG(
dbgs() <<
"Dead: known opcode with no side effects\n");
832void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &
MI)
const {
834 for (
const auto &MO :
MI.all_defs()) {
838 SmallVector<MachineInstr *, 4> UselessOpNames;
841 "There is still a use of the dead function.");
844 for (MachineInstr *OpNameMI : UselessOpNames) {
846 OpNameMI->eraseFromParent();
851void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &
MI)
const {
854 removeOpNamesForDeadMI(
MI);
855 MI.eraseFromParent();
858bool SPIRVInstructionSelector::select(MachineInstr &
I) {
859 resetVRegsType(*
I.getParent()->getParent());
861 assert(
I.getParent() &&
"Instruction should be in a basic block!");
862 assert(
I.getParent()->getParent() &&
"Instruction should be in a function!");
867 removeDeadInstruction(
I);
874 if (Opcode == SPIRV::ASSIGN_TYPE) {
875 Register DstReg =
I.getOperand(0).getReg();
876 Register SrcReg =
I.getOperand(1).getReg();
879 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
880 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
881 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
882 Register SelectDstReg =
Def->getOperand(0).getReg();
883 bool SuccessToSelectSelect [[maybe_unused]] = selectSelect(
885 assert(SuccessToSelectSelect);
887 Def->eraseFromParent();
894 bool Res = selectImpl(
I, *CoverageInfo);
896 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
897 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
901 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
913 }
else if (
I.getNumDefs() == 1) {
925 removeDeadInstruction(
I);
930 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
931 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
937 bool HasDefs =
I.getNumDefs() > 0;
940 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
941 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
942 if (spvSelect(ResVReg, ResType,
I)) {
944 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
955 case TargetOpcode::G_CONSTANT:
956 case TargetOpcode::G_FCONSTANT:
963 MachineInstr &
I)
const {
966 if (DstRC != SrcRC && SrcRC)
968 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
975bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
976 SPIRVTypeInst ResType,
977 MachineInstr &
I)
const {
978 const unsigned Opcode =
I.getOpcode();
980 return selectImpl(
I, *CoverageInfo);
982 case TargetOpcode::G_CONSTANT:
983 case TargetOpcode::G_FCONSTANT:
984 return selectConst(ResVReg, ResType,
I);
985 case TargetOpcode::G_GLOBAL_VALUE:
986 return selectGlobalValue(ResVReg,
I);
987 case TargetOpcode::G_IMPLICIT_DEF:
988 return selectOpUndef(ResVReg, ResType,
I);
989 case TargetOpcode::G_FREEZE:
990 return selectFreeze(ResVReg, ResType,
I);
992 case TargetOpcode::G_INTRINSIC:
993 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
994 case TargetOpcode::G_INTRINSIC_CONVERGENT:
995 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
996 return selectIntrinsic(ResVReg, ResType,
I);
997 case TargetOpcode::G_BITREVERSE:
998 return selectBitreverse(ResVReg, ResType,
I);
1000 case TargetOpcode::G_BUILD_VECTOR:
1001 return selectBuildVector(ResVReg, ResType,
I);
1002 case TargetOpcode::G_SPLAT_VECTOR:
1003 return selectSplatVector(ResVReg, ResType,
I);
1005 case TargetOpcode::G_SHUFFLE_VECTOR: {
1006 MachineBasicBlock &BB = *
I.getParent();
1007 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
1010 .
addUse(
I.getOperand(1).getReg())
1011 .
addUse(
I.getOperand(2).getReg());
1012 for (
auto V :
I.getOperand(3).getShuffleMask())
1017 case TargetOpcode::G_MEMMOVE:
1018 case TargetOpcode::G_MEMCPY:
1019 case TargetOpcode::G_MEMSET:
1020 return selectMemOperation(ResVReg,
I);
1022 case TargetOpcode::G_ICMP:
1023 return selectICmp(ResVReg, ResType,
I);
1024 case TargetOpcode::G_FCMP:
1025 return selectFCmp(ResVReg, ResType,
I);
1027 case TargetOpcode::G_FRAME_INDEX:
1028 return selectFrameIndex(ResVReg, ResType,
I);
1030 case TargetOpcode::G_LOAD:
1031 return selectLoad(ResVReg, ResType,
I);
1032 case TargetOpcode::G_STORE:
1033 return selectStore(
I);
1035 case TargetOpcode::G_BR:
1036 return selectBranch(
I);
1037 case TargetOpcode::G_BRCOND:
1038 return selectBranchCond(
I);
1040 case TargetOpcode::G_PHI:
1041 return selectPhi(ResVReg,
I);
1043 case TargetOpcode::G_FPTOSI:
1044 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1045 case TargetOpcode::G_FPTOUI:
1046 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1048 case TargetOpcode::G_FPTOSI_SAT:
1049 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1050 case TargetOpcode::G_FPTOUI_SAT:
1051 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1053 case TargetOpcode::G_SITOFP:
1054 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
1055 case TargetOpcode::G_UITOFP:
1056 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
1058 case TargetOpcode::G_CTPOP:
1059 return selectPopCount(ResVReg, ResType,
I, SPIRV::OpBitCount);
1060 case TargetOpcode::G_SMIN:
1061 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
1062 case TargetOpcode::G_UMIN:
1063 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
1065 case TargetOpcode::G_SMAX:
1066 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
1067 case TargetOpcode::G_UMAX:
1068 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
1070 case TargetOpcode::G_SCMP:
1071 return selectSUCmp(ResVReg, ResType,
I,
true);
1072 case TargetOpcode::G_UCMP:
1073 return selectSUCmp(ResVReg, ResType,
I,
false);
1074 case TargetOpcode::G_LROUND:
1075 case TargetOpcode::G_LLROUND: {
1078 MRI->
setRegClass(regForLround, &SPIRV::iIDRegClass);
1080 regForLround, *(
I.getParent()->getParent()));
1082 CL::round, GL::Round,
false);
1084 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1091 case TargetOpcode::G_STRICT_FMA:
1092 case TargetOpcode::G_FMA: {
1095 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1098 .
addUse(
I.getOperand(1).getReg())
1099 .
addUse(
I.getOperand(2).getReg())
1100 .
addUse(
I.getOperand(3).getReg())
1105 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1108 case TargetOpcode::G_STRICT_FLDEXP:
1109 return selectExtInst(ResVReg, ResType,
I, CL::ldexp);
1111 case TargetOpcode::G_FPOW:
1112 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1113 case TargetOpcode::G_FPOWI:
1114 return selectFpowi(ResVReg, ResType,
I);
1116 case TargetOpcode::G_FEXP:
1117 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1118 case TargetOpcode::G_FEXP2:
1119 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1120 case TargetOpcode::G_FEXP10:
1121 return selectExp10(ResVReg, ResType,
I);
1123 case TargetOpcode::G_FMODF:
1124 return selectModf(ResVReg, ResType,
I);
1125 case TargetOpcode::G_FSINCOS:
1126 return selectSincos(ResVReg, ResType,
I);
1128 case TargetOpcode::G_FLOG:
1129 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1130 case TargetOpcode::G_FLOG2:
1131 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1132 case TargetOpcode::G_FLOG10:
1133 return selectLog10(ResVReg, ResType,
I);
1135 case TargetOpcode::G_FABS:
1136 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1137 case TargetOpcode::G_ABS:
1138 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1140 case TargetOpcode::G_FMINNUM:
1141 case TargetOpcode::G_FMINIMUM:
1142 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1143 case TargetOpcode::G_FMAXNUM:
1144 case TargetOpcode::G_FMAXIMUM:
1145 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1147 case TargetOpcode::G_FCOPYSIGN:
1148 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1150 case TargetOpcode::G_FCEIL:
1151 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1152 case TargetOpcode::G_FFLOOR:
1153 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1155 case TargetOpcode::G_FCOS:
1156 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1157 case TargetOpcode::G_FSIN:
1158 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1159 case TargetOpcode::G_FTAN:
1160 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1161 case TargetOpcode::G_FACOS:
1162 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1163 case TargetOpcode::G_FASIN:
1164 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1165 case TargetOpcode::G_FATAN:
1166 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1167 case TargetOpcode::G_FATAN2:
1168 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1169 case TargetOpcode::G_FCOSH:
1170 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1171 case TargetOpcode::G_FSINH:
1172 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1173 case TargetOpcode::G_FTANH:
1174 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1176 case TargetOpcode::G_STRICT_FSQRT:
1177 case TargetOpcode::G_FSQRT:
1178 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1180 case TargetOpcode::G_CTTZ:
1181 case TargetOpcode::G_CTTZ_ZERO_UNDEF:
1182 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1183 case TargetOpcode::G_CTLZ:
1184 case TargetOpcode::G_CTLZ_ZERO_UNDEF:
1185 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1187 case TargetOpcode::G_INTRINSIC_ROUND:
1188 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1189 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1190 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1191 case TargetOpcode::G_INTRINSIC_TRUNC:
1192 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1193 case TargetOpcode::G_FRINT:
1194 case TargetOpcode::G_FNEARBYINT:
1195 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1197 case TargetOpcode::G_SMULH:
1198 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1199 case TargetOpcode::G_UMULH:
1200 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1202 case TargetOpcode::G_SADDSAT:
1203 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1204 case TargetOpcode::G_UADDSAT:
1205 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1206 case TargetOpcode::G_SSUBSAT:
1207 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1208 case TargetOpcode::G_USUBSAT:
1209 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1211 case TargetOpcode::G_FFREXP:
1212 return selectFrexp(ResVReg, ResType,
I);
1214 case TargetOpcode::G_UADDO:
1215 return selectOverflowArith(ResVReg, ResType,
I,
1216 ResType->
getOpcode() == SPIRV::OpTypeVector
1217 ? SPIRV::OpIAddCarryV
1218 : SPIRV::OpIAddCarryS);
1219 case TargetOpcode::G_USUBO:
1220 return selectOverflowArith(ResVReg, ResType,
I,
1221 ResType->
getOpcode() == SPIRV::OpTypeVector
1222 ? SPIRV::OpISubBorrowV
1223 : SPIRV::OpISubBorrowS);
1224 case TargetOpcode::G_UMULO:
1225 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1226 case TargetOpcode::G_SMULO:
1227 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1229 case TargetOpcode::G_SEXT:
1230 return selectExt(ResVReg, ResType,
I,
true);
1231 case TargetOpcode::G_ANYEXT:
1232 case TargetOpcode::G_ZEXT:
1233 return selectExt(ResVReg, ResType,
I,
false);
1234 case TargetOpcode::G_TRUNC:
1235 return selectTrunc(ResVReg, ResType,
I);
1236 case TargetOpcode::G_FPTRUNC:
1237 case TargetOpcode::G_FPEXT:
1238 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1240 case TargetOpcode::G_PTRTOINT:
1241 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1242 case TargetOpcode::G_INTTOPTR:
1243 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1244 case TargetOpcode::G_BITCAST:
1245 return selectBitcast(ResVReg, ResType,
I);
1246 case TargetOpcode::G_ADDRSPACE_CAST:
1247 return selectAddrSpaceCast(ResVReg, ResType,
I);
1248 case TargetOpcode::G_PTR_ADD: {
1250 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1254 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1255 (*II).getOpcode() == TargetOpcode::COPY ||
1256 (*II).getOpcode() == SPIRV::OpVariable) &&
1257 getImm(
I.getOperand(2), MRI));
1259 bool IsGVInit =
false;
1263 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1264 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1265 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1266 (*UseIt).getOpcode() == SPIRV::OpVariable) {
1276 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1289 "incompatible result and operand types in a bitcast");
1291 MachineInstrBuilder MIB =
1292 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1299 : SPIRV::OpInBoundsPtrAccessChain))
1303 .
addUse(
I.getOperand(2).getReg())
1306 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1310 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1312 .
addUse(
I.getOperand(2).getReg())
1321 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1324 .
addImm(
static_cast<uint32_t
>(
1325 SPIRV::Opcode::InBoundsPtrAccessChain))
1328 .
addUse(
I.getOperand(2).getReg());
1333 case TargetOpcode::G_ATOMICRMW_OR:
1334 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1335 case TargetOpcode::G_ATOMICRMW_ADD:
1336 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1337 case TargetOpcode::G_ATOMICRMW_AND:
1338 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1339 case TargetOpcode::G_ATOMICRMW_MAX:
1340 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1341 case TargetOpcode::G_ATOMICRMW_MIN:
1342 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1343 case TargetOpcode::G_ATOMICRMW_SUB:
1344 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1345 case TargetOpcode::G_ATOMICRMW_XOR:
1346 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1347 case TargetOpcode::G_ATOMICRMW_UMAX:
1348 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1349 case TargetOpcode::G_ATOMICRMW_UMIN:
1350 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1351 case TargetOpcode::G_ATOMICRMW_XCHG:
1352 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1354 case TargetOpcode::G_ATOMICRMW_FADD:
1355 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1356 case TargetOpcode::G_ATOMICRMW_FSUB:
1358 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1359 ResType->
getOpcode() == SPIRV::OpTypeVector
1361 : SPIRV::OpFNegate);
1362 case TargetOpcode::G_ATOMICRMW_FMIN:
1363 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1364 case TargetOpcode::G_ATOMICRMW_FMAX:
1365 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1367 case TargetOpcode::G_FENCE:
1368 return selectFence(
I);
1370 case TargetOpcode::G_STACKSAVE:
1371 return selectStackSave(ResVReg, ResType,
I);
1372 case TargetOpcode::G_STACKRESTORE:
1373 return selectStackRestore(
I);
1375 case TargetOpcode::G_UNMERGE_VALUES:
1381 case TargetOpcode::G_TRAP:
1382 case TargetOpcode::G_UBSANTRAP:
1383 case TargetOpcode::DBG_LABEL:
1385 case TargetOpcode::G_DEBUGTRAP:
1386 return selectDebugTrap(ResVReg, ResType,
I);
1393bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1394 SPIRVTypeInst ResType,
1395 MachineInstr &
I)
const {
1396 unsigned Opcode = SPIRV::OpNop;
1403bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1404 SPIRVTypeInst ResType,
1406 GL::GLSLExtInst GLInst,
1407 bool setMIFlags,
bool useMISrc,
1410 SPIRV::InstructionSet::InstructionSet::GLSL_std_450)) {
1411 std::string DiagMsg;
1412 raw_string_ostream OS(DiagMsg);
1413 I.print(OS,
true,
false,
false,
false);
1414 DiagMsg +=
" is only supported with the GLSL extended instruction set.\n";
1417 return selectExtInst(ResVReg, ResType,
I,
1418 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1419 setMIFlags, useMISrc, SrcRegs);
1422bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1423 SPIRVTypeInst ResType,
1425 CL::OpenCLExtInst CLInst,
1426 bool setMIFlags,
bool useMISrc,
1428 return selectExtInst(ResVReg, ResType,
I,
1429 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1430 setMIFlags, useMISrc, SrcRegs);
1433bool SPIRVInstructionSelector::selectExtInst(
1434 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1435 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1437 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1438 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1439 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1443bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1444 SPIRVTypeInst ResType,
1447 bool setMIFlags,
bool useMISrc,
1450 for (
const auto &[InstructionSet, Opcode] : Insts) {
1454 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1457 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1462 const unsigned NumOps =
I.getNumOperands();
1465 I.getOperand(Index).getType() ==
1466 MachineOperand::MachineOperandType::MO_IntrinsicID)
1469 MIB.
add(
I.getOperand(Index));
1481bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1482 SPIRVTypeInst ResType,
1483 MachineInstr &
I)
const {
1484 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1485 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1486 for (
const auto &Ex : ExtInsts) {
1487 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1488 uint32_t Opcode = Ex.second;
1492 MachineIRBuilder MIRBuilder(
I);
1495 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1500 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1503 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1506 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1509 .
addImm(
static_cast<uint32_t
>(Ex.first))
1511 .
add(
I.getOperand(2))
1515 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1516 .
addDef(
I.getOperand(1).getReg())
1525bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1526 SPIRVTypeInst ResType,
1527 MachineInstr &
I)
const {
1528 Register CosResVReg =
I.getOperand(1).getReg();
1529 unsigned SrcIdx =
I.getNumExplicitDefs();
1534 MachineIRBuilder MIRBuilder(
I);
1536 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1541 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1544 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1546 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1549 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1551 .
add(
I.getOperand(SrcIdx))
1554 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1562 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1565 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1567 .
add(
I.getOperand(SrcIdx))
1569 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1572 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1574 .
add(
I.getOperand(SrcIdx))
1581bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1582 SPIRVTypeInst ResType,
1584 std::vector<Register> Srcs,
1585 unsigned Opcode)
const {
1586 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1596std::optional<SplitParts> SPIRVInstructionSelector::splitEvenOddLanes(
1597 Register PopCountReg,
unsigned ComponentCount, MachineInstr &
I,
1598 SPIRVTypeInst I32Type)
const {
1601 if (ComponentCount == 1) {
1604 Parts.IsScalar =
true;
1605 Parts.Type = I32Type;
1613 if (!selectOpWithSrcs(Parts.High, I32Type,
I, {PopCountReg, IdxOne},
1614 SPIRV::OpVectorExtractDynamic))
1615 return std::nullopt;
1617 if (!selectOpWithSrcs(Parts.Low, I32Type,
I, {PopCountReg, IdxZero},
1618 SPIRV::OpVectorExtractDynamic))
1619 return std::nullopt;
1623 MachineIRBuilder MIRBuilder(
I);
1624 Parts.IsScalar =
false;
1631 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1632 TII.get(SPIRV::OpVectorShuffle))
1637 for (
unsigned J = 1; J < ComponentCount * 2; J += 2)
1642 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1643 TII.get(SPIRV::OpVectorShuffle))
1648 for (
unsigned J = 0; J < ComponentCount * 2; J += 2)
1656bool SPIRVInstructionSelector::selectPopCount16(
Register ResVReg,
1657 SPIRVTypeInst ResType,
1660 unsigned Opcode)
const {
1661 Register OpReg =
I.getOperand(1).getReg();
1664 MachineIRBuilder MIRBuilder(
I);
1666 SPIRVTypeInst I32VectorType =
1669 bool IsVector = NumElems > 1;
1670 SPIRVTypeInst ExtType = IsVector ? I32VectorType : I32Type;
1673 if (!selectOpWithSrcs(ExtReg, ExtType,
I, {OpReg}, SPIRV::OpUConvert))
1677 if (!selectPopCount32(PopCountReg, ExtType,
I, ExtReg, Opcode))
1680 return selectOpWithSrcs(ResVReg, ResType,
I, {PopCountReg}, ExtOpcode);
1683bool SPIRVInstructionSelector::selectPopCount32(
Register ResVReg,
1684 SPIRVTypeInst ResType,
1687 unsigned Opcode)
const {
1688 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
1691bool SPIRVInstructionSelector::selectPopCount64(
Register ResVReg,
1692 SPIRVTypeInst ResType,
1695 unsigned Opcode)
const {
1697 if (ComponentCount > 2)
1698 return handle64BitOverflow(
1699 ResVReg, ResType,
I, SrcReg, Opcode,
1701 unsigned O) {
return this->selectPopCount64(R,
T,
I, S, O); });
1703 MachineIRBuilder MIRBuilder(
I);
1708 I32Type, 2 * ComponentCount, MIRBuilder,
false);
1712 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
1717 if (!selectPopCount32(Pop32, VecI32Type,
I, Vec32, Opcode))
1721 auto MaybeParts = splitEvenOddLanes(Pop32, ComponentCount,
I, I32Type);
1724 SplitParts &Parts = *MaybeParts;
1727 unsigned OpAdd = Parts.IsScalar ? SPIRV::OpIAddS : SPIRV::OpIAddV;
1729 if (!selectOpWithSrcs(Sum, Parts.Type,
I, {Parts.High, Parts.Low}, OpAdd))
1734 unsigned ConvOp = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
1735 return selectOpWithSrcs(ResVReg, ResType,
I, {Sum}, ConvOp);
1738bool SPIRVInstructionSelector::selectPopCount(
Register ResVReg,
1739 SPIRVTypeInst ResType,
1741 unsigned Opcode)
const {
1746 if (!STI.getTargetTriple().isVulkanOS())
1747 return selectUnOp(ResVReg, ResType,
I, Opcode);
1749 Register OpReg =
I.getOperand(1).getReg();
1752 : SPIRV::OpUConvert;
1756 return selectPopCount16(ResVReg, ResType,
I, ExtOpcode, Opcode);
1758 return selectPopCount32(ResVReg, ResType,
I, OpReg, Opcode);
1760 return selectPopCount64(ResVReg, ResType,
I, OpReg, Opcode);
1766bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1767 SPIRVTypeInst ResType,
1769 unsigned Opcode)
const {
1771 Register SrcReg =
I.getOperand(1).getReg();
1776 unsigned DefOpCode = DefIt->getOpcode();
1777 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1780 if (
auto *VRD =
getVRegDef(*MRI, DefIt->getOperand(1).getReg()))
1781 DefOpCode = VRD->getOpcode();
1783 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1784 DefOpCode == TargetOpcode::G_CONSTANT ||
1785 DefOpCode == SPIRV::OpVariable || DefOpCode == SPIRV::OpConstantI) {
1791 uint32_t SpecOpcode = 0;
1793 case SPIRV::OpConvertPtrToU:
1794 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1796 case SPIRV::OpConvertUToPtr:
1797 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1802 TII.get(SPIRV::OpSpecConstantOp))
1812 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1816bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1817 SPIRVTypeInst ResType,
1818 MachineInstr &
I)
const {
1819 Register OpReg =
I.getOperand(1).getReg();
1820 SPIRVTypeInst OpType =
1824 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1834 if (
MemOp->isVolatile())
1835 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1836 if (
MemOp->isNonTemporal())
1837 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1839 if (!ST->isShader() &&
MemOp->getAlign().value())
1840 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1844 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1845 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1849 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1851 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1855 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1859 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1861 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1873 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1875 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1877 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
1881bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
1882 SPIRVTypeInst ResType,
1883 MachineInstr &
I)
const {
1885 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1890 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
1891 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1893 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1897 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1901 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1902 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
1903 I.getDebugLoc(),
I);
1907 MachineIRBuilder MIRBuilder(
I);
1909 if (
I.getNumMemOperands()) {
1910 const MachineMemOperand *MemOp = *
I.memoperands_begin();
1911 if (MemOp->isAtomic())
1912 return selectAtomicLoad(ResVReg, ResType,
I);
1915 auto MIB = MIRBuilder.buildInstr(SPIRV::OpLoad)
1919 if (!
I.getNumMemOperands()) {
1920 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1922 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1931bool SPIRVInstructionSelector::selectAtomicLoad(
Register ResVReg,
1932 SPIRVTypeInst ResType,
1933 MachineInstr &
I)
const {
1934 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
1937 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1940 return diagnoseUnsupported(
I,
1941 "Lowering to SPIR-V of atomic load is only "
1942 "allowed for integer or floating point types");
1944 assert(
I.getNumMemOperands());
1945 const MachineMemOperand &MemOp = **
I.memoperands_begin();
1946 assert(MemOp.isAtomic());
1949 if (MemOp.isVolatile())
1950 return diagnoseUnsupported(
I,
"Lowering to SPIR-V of atomic load of "
1951 "volatile memory is not supported");
1955 Register ScopeReg = buildI32Constant(Scope,
I);
1963 MachineIRBuilder MIRBuilder(
I);
1964 auto AtomicLoad = MIRBuilder.buildInstr(SPIRV::OpAtomicLoad)
1970 AtomicLoad.constrainAllUses(
TII,
TRI, RBI);
1974bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
1976 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
1977 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1982 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
1983 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1988 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1992 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1993 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1994 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1995 TII.get(SPIRV::OpImageWrite))
2001 if (sampledTypeIsSignedInteger(LLVMHandleType))
2004 BMI.constrainAllUses(
TII,
TRI, RBI);
2009 if (
I.getNumMemOperands()) {
2010 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2011 if (MemOp->isAtomic())
2012 return selectAtomicStore(
I);
2015 MachineIRBuilder MIRBuilder(
I);
2016 auto MIB = MIRBuilder.buildInstr(SPIRV::OpStore).
addUse(Ptr).
addUse(StoreVal);
2017 if (!
I.getNumMemOperands()) {
2018 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2020 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2029bool SPIRVInstructionSelector::selectAtomicStore(MachineInstr &
I)
const {
2030 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2033 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2034 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2039 return diagnoseUnsupported(
I,
2040 "Lowering to SPIR-V of atomic store is only "
2041 "allowed for integer or floating point types");
2043 assert(
I.getNumMemOperands());
2044 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2045 assert(MemOp.isAtomic());
2049 if (MemOp.isVolatile())
2050 return diagnoseUnsupported(
I,
"Lowering to SPIR-V of atomic store of "
2051 "volatile memory is not supported");
2055 Register ScopeReg = buildI32Constant(Scope,
I);
2063 MachineIRBuilder MIRBuilder(
I);
2064 auto AtomicStore = MIRBuilder.buildInstr(SPIRV::OpAtomicStore)
2069 AtomicStore.constrainAllUses(
TII,
TRI, RBI);
2073bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
2074 SPIRVTypeInst ResType,
2075 MachineInstr &
I)
const {
2076 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
2084 const Register PtrsReg =
I.getOperand(2).getReg();
2085 const uint32_t Alignment =
I.getOperand(3).getImm();
2086 const Register MaskReg =
I.getOperand(4).getReg();
2087 const Register PassthruReg =
I.getOperand(5).getReg();
2088 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2092 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
2103bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
2104 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
2111 const Register ValuesReg =
I.getOperand(1).getReg();
2112 const Register PtrsReg =
I.getOperand(2).getReg();
2113 const uint32_t Alignment =
I.getOperand(3).getImm();
2114 const Register MaskReg =
I.getOperand(4).getReg();
2115 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2119 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
2128bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
2129 const Twine &Msg)
const {
2130 const Function &
F =
I.getMF()->getFunction();
2131 F.getContext().diagnose(
2132 DiagnosticInfoUnsupported(
F, Msg,
I.getDebugLoc(),
DS_Error));
2136bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
2137 SPIRVTypeInst ResType,
2138 MachineInstr &
I)
const {
2139 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2141 "llvm.stacksave intrinsic: this instruction requires the following "
2142 "SPIR-V extension: SPV_INTEL_variable_length_array",
2145 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
2152bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
2153 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2155 "llvm.stackrestore intrinsic: this instruction requires the following "
2156 "SPIR-V extension: SPV_INTEL_variable_length_array",
2158 if (!
I.getOperand(0).isReg())
2161 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
2162 .
addUse(
I.getOperand(0).getReg())
2168SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
2169 MachineIRBuilder MIRBuilder(
I);
2170 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
2177 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
2181 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2182 Type *ArrTy = ArrayType::get(ValTy, Num);
2184 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
2187 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
2194 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
2197 .
addImm(SPIRV::StorageClass::UniformConstant)
2208bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
2211 Register DstReg =
I.getOperand(0).getReg();
2216 uint64_t CopySize =
getIConstVal(
I.getOperand(2).getReg(), MRI);
2221 "Unable to determine pointee type size for OpCopyMemory");
2222 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
2223 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
2225 "OpCopyMemory requires the size to match the pointee type size");
2226 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemory))
2229 if (
I.getNumMemOperands()) {
2230 MachineIRBuilder MIRBuilder(
I);
2237bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &
I,
2240 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemorySized))
2241 .
addUse(
I.getOperand(0).getReg())
2243 .
addUse(
I.getOperand(2).getReg());
2244 if (
I.getNumMemOperands()) {
2245 MachineIRBuilder MIRBuilder(
I);
2252bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
2253 MachineInstr &
I)
const {
2254 Register SrcReg =
I.getOperand(1).getReg();
2255 if (
I.getOpcode() == TargetOpcode::G_MEMSET) {
2256 Register VarReg = getOrCreateMemSetGlobal(
I);
2259 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2261 ValTy,
I, SPIRV::StorageClass::UniformConstant);
2263 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
2267 if (!selectCopyMemory(
I, SrcReg))
2270 if (!selectCopyMemorySized(
I, SrcReg))
2273 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
2274 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
2279bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
2280 SPIRVTypeInst ResType,
2283 unsigned NegateOpcode)
const {
2285 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2288 Register ScopeReg = buildI32Constant(Scope,
I);
2290 Register Ptr =
I.getOperand(1).getReg();
2291 uint32_t ScSem =
static_cast<uint32_t
>(
2295 Register MemSemReg = buildI32Constant(MemSem,
I);
2297 Register ValueReg =
I.getOperand(2).getReg();
2298 if (NegateOpcode != 0) {
2301 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
2306 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2317bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2318 unsigned ArgI =
I.getNumOperands() - 1;
2320 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2321 SPIRVTypeInst SrcType =
2323 if (!SrcType || SrcType->
getOpcode() != SPIRV::OpTypeVector)
2325 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2329 unsigned CurrentIndex = 0;
2330 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2331 Register ResVReg =
I.getOperand(i).getReg();
2334 LLT ResLLT = MRI->
getType(ResVReg);
2340 ResType = ScalarType;
2346 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
2349 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2355 for (
unsigned j = 0;
j < NumElements; ++
j) {
2356 MIB.
addImm(CurrentIndex + j);
2358 CurrentIndex += NumElements;
2362 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2374bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2377 Register MemSemReg = buildI32Constant(MemSem,
I);
2379 uint32_t
Scope =
static_cast<uint32_t
>(
2381 Register ScopeReg = buildI32Constant(Scope,
I);
2383 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2390bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2391 SPIRVTypeInst ResType,
2393 unsigned Opcode)
const {
2394 Type *ResTy =
nullptr;
2398 "Not enough info to select the arithmetic with overflow instruction");
2401 "with overflow instruction");
2407 MachineIRBuilder MIRBuilder(
I);
2409 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2410 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2416 Register ZeroReg = buildZerosVal(ResType,
I);
2421 if (ResName.
size() > 0)
2426 BuildMI(BB, MIRBuilder.getInsertPt(),
I.getDebugLoc(),
TII.get(Opcode))
2429 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2430 MIB.
addUse(
I.getOperand(i).getReg());
2435 MRI->
setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2436 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2438 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2439 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2446 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2447 .
addDef(
I.getOperand(1).getReg())
2455bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2456 SPIRVTypeInst ResType,
2457 MachineInstr &
I)
const {
2459 "selectAtomicCmpXchg only handles the spv_cmpxchg intrinsic");
2460 Register Ptr =
I.getOperand(2).getReg();
2461 Register ScopeReg =
I.getOperand(5).getReg();
2462 Register MemSemEqReg =
I.getOperand(6).getReg();
2463 Register MemSemNeqReg =
I.getOperand(7).getReg();
2465 Register Val =
I.getOperand(4).getReg();
2469 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2488 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2495 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2507 case SPIRV::StorageClass::DeviceOnlyINTEL:
2508 case SPIRV::StorageClass::HostOnlyINTEL:
2517 bool IsGRef =
false;
2518 bool IsAllowedRefs =
2520 unsigned Opcode = It.getOpcode();
2521 if (Opcode == SPIRV::OpConstantComposite ||
2522 Opcode == SPIRV::OpSpecConstantComposite ||
2523 Opcode == SPIRV::OpVariable ||
2524 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2525 return IsGRef = true;
2526 return Opcode == SPIRV::OpName;
2528 return IsAllowedRefs && IsGRef;
2531Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2532 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2534 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2538SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2540 uint32_t Opcode)
const {
2541 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2542 TII.get(SPIRV::OpSpecConstantOp))
2550SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2551 SPIRVTypeInst SrcPtrTy)
const {
2552 SPIRVTypeInst GenericPtrTy =
2556 SPIRV::StorageClass::Generic),
2558 MachineFunction *MF =
I.getParent()->getParent();
2560 MachineInstrBuilder MIB = buildSpecConstantOp(
2562 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2572bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2573 SPIRVTypeInst ResType,
2574 MachineInstr &
I)
const {
2578 Register SrcPtr =
I.getOperand(1).getReg();
2582 if (SrcPtrTy->
getOpcode() != SPIRV::OpTypePointer ||
2583 ResType->
getOpcode() != SPIRV::OpTypePointer)
2584 return BuildCOPY(ResVReg, SrcPtr,
I);
2594 unsigned SpecOpcode =
2596 ?
static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric)
2599 ? static_cast<uint32_t>(
SPIRV::Opcode::GenericCastToPtr)
2606 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2608 .constrainAllUses(
TII,
TRI, RBI);
2610 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2612 buildSpecConstantOp(
2614 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2615 .constrainAllUses(
TII,
TRI, RBI);
2622 return BuildCOPY(ResVReg, SrcPtr,
I);
2624 if ((SrcSC == SPIRV::StorageClass::Function &&
2625 DstSC == SPIRV::StorageClass::Private) ||
2626 (DstSC == SPIRV::StorageClass::Function &&
2627 SrcSC == SPIRV::StorageClass::Private))
2628 return BuildCOPY(ResVReg, SrcPtr,
I);
2632 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2635 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2638 SPIRVTypeInst GenericPtrTy =
2657 return selectUnOp(ResVReg, ResType,
I,
2658 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
2660 return selectUnOp(ResVReg, ResType,
I,
2661 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
2663 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2665 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2675 return SPIRV::OpFOrdEqual;
2677 return SPIRV::OpFOrdGreaterThanEqual;
2679 return SPIRV::OpFOrdGreaterThan;
2681 return SPIRV::OpFOrdLessThanEqual;
2683 return SPIRV::OpFOrdLessThan;
2685 return SPIRV::OpFOrdNotEqual;
2687 return SPIRV::OpOrdered;
2689 return SPIRV::OpFUnordEqual;
2691 return SPIRV::OpFUnordGreaterThanEqual;
2693 return SPIRV::OpFUnordGreaterThan;
2695 return SPIRV::OpFUnordLessThanEqual;
2697 return SPIRV::OpFUnordLessThan;
2699 return SPIRV::OpFUnordNotEqual;
2701 return SPIRV::OpUnordered;
2711 return SPIRV::OpIEqual;
2713 return SPIRV::OpINotEqual;
2715 return SPIRV::OpSGreaterThanEqual;
2717 return SPIRV::OpSGreaterThan;
2719 return SPIRV::OpSLessThanEqual;
2721 return SPIRV::OpSLessThan;
2723 return SPIRV::OpUGreaterThanEqual;
2725 return SPIRV::OpUGreaterThan;
2727 return SPIRV::OpULessThanEqual;
2729 return SPIRV::OpULessThan;
2738 return SPIRV::OpPtrEqual;
2740 return SPIRV::OpPtrNotEqual;
2751 return SPIRV::OpLogicalEqual;
2753 return SPIRV::OpLogicalNotEqual;
2787bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
2788 SPIRVTypeInst ResType,
2790 unsigned OpAnyOrAll)
const {
2791 assert(
I.getNumOperands() == 3);
2792 assert(
I.getOperand(2).isReg());
2794 Register InputRegister =
I.getOperand(2).getReg();
2797 assert(InputType &&
"VReg has no type assigned");
2800 bool IsVectorTy = InputType->
getOpcode() == SPIRV::OpTypeVector;
2801 if (IsBoolTy && !IsVectorTy) {
2802 assert(ResVReg ==
I.getOperand(0).getReg());
2803 return BuildCOPY(ResVReg, InputRegister,
I);
2807 unsigned SpirvNotEqualId =
2808 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
2810 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
2815 IsBoolTy ? InputRegister
2823 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
2825 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
2842bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
2843 SPIRVTypeInst ResType,
2844 MachineInstr &
I)
const {
2845 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
2848bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
2849 SPIRVTypeInst ResType,
2850 MachineInstr &
I)
const {
2851 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
2855bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
2856 SPIRVTypeInst ResType,
2857 MachineInstr &
I)
const {
2858 assert(
I.getNumOperands() == 4);
2859 assert(
I.getOperand(2).isReg());
2860 assert(
I.getOperand(3).isReg());
2862 [[maybe_unused]] SPIRVTypeInst VecType =
2867 "dot product requires a vector of at least 2 components");
2869 [[maybe_unused]] SPIRVTypeInst EltType =
2878 .
addUse(
I.getOperand(2).getReg())
2879 .
addUse(
I.getOperand(3).getReg())
2884bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
2885 SPIRVTypeInst ResType,
2888 assert(
I.getNumOperands() == 4);
2889 assert(
I.getOperand(2).isReg());
2890 assert(
I.getOperand(3).isReg());
2893 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
2897 .
addUse(
I.getOperand(2).getReg())
2898 .
addUse(
I.getOperand(3).getReg())
2905bool SPIRVInstructionSelector::selectIntegerDotExpansion(
2906 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2907 assert(
I.getNumOperands() == 4);
2908 assert(
I.getOperand(2).isReg());
2909 assert(
I.getOperand(3).isReg());
2913 Register Vec0 =
I.getOperand(2).getReg();
2914 Register Vec1 =
I.getOperand(3).getReg();
2918 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
2927 "dot product requires a vector of at least 2 components");
2930 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2940 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2951 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2963bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
2964 SPIRVTypeInst ResType,
2965 MachineInstr &
I)
const {
2967 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
2970 .
addUse(
I.getOperand(2).getReg())
2975bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
2976 SPIRVTypeInst ResType,
2977 MachineInstr &
I)
const {
2979 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
2982 .
addUse(
I.getOperand(2).getReg())
2987template <
bool Signed>
2988bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
2989 SPIRVTypeInst ResType,
2990 MachineInstr &
I)
const {
2991 assert(
I.getNumOperands() == 5);
2992 assert(
I.getOperand(2).isReg());
2993 assert(
I.getOperand(3).isReg());
2994 assert(
I.getOperand(4).isReg());
2997 Register Acc =
I.getOperand(2).getReg();
3001 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3003 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
3008 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
3011 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3023template <
bool Signed>
3024bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
3025 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3026 assert(
I.getNumOperands() == 5);
3027 assert(
I.getOperand(2).isReg());
3028 assert(
I.getOperand(3).isReg());
3029 assert(
I.getOperand(4).isReg());
3032 Register Acc =
I.getOperand(2).getReg();
3038 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
3042 for (
unsigned i = 0; i < 4; i++) {
3065 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
3085 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3100bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
3101 SPIRVTypeInst ResType,
3102 MachineInstr &
I)
const {
3103 assert(
I.getNumOperands() == 3);
3104 assert(
I.getOperand(2).isReg());
3106 Register VZero = buildZerosValF(ResType,
I);
3107 Register VOne = buildOnesValF(ResType,
I);
3109 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
3112 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3114 .
addUse(
I.getOperand(2).getReg())
3121bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
3122 SPIRVTypeInst ResType,
3123 MachineInstr &
I)
const {
3124 assert(
I.getNumOperands() == 3);
3125 assert(
I.getOperand(2).isReg());
3127 Register InputRegister =
I.getOperand(2).getReg();
3129 auto &
DL =
I.getDebugLoc();
3139 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
3141 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
3149 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3154 if (NeedsConversion) {
3155 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
3166bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
3167 SPIRVTypeInst ResType,
3169 unsigned Opcode)
const {
3173 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
3179 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
3180 BMI.addUse(
I.getOperand(J).getReg());
3187bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
3190 bool WithGroupSync)
const {
3192 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
3194 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
3196 assert(((Scope != SPIRV::Scope::Workgroup) ||
3197 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
3198 "Workgroup Scope must set WorkGroupMemory semantic "
3199 "in Barrier instruction");
3201 assert(((Scope != SPIRV::Scope::Device) ||
3202 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
3203 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
3204 "Device Scope must set UniformMemory and ImageMemory semantic "
3205 "in Barrier instruction");
3211 if (WithGroupSync) {
3212 Register ExecReg = buildI32Constant(SPIRV::Scope::Workgroup,
I);
3216 Register ScopeReg = buildI32Constant(Scope,
I);
3217 Register MemSemReg = buildI32Constant(MemSem,
I);
3219 MI.addUse(ScopeReg).addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
3223bool SPIRVInstructionSelector::selectWaveActiveCountBits(
3224 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3229 if (!selectWaveOpInst(BallotReg, BallotType,
I,
3230 SPIRV::OpGroupNonUniformBallot))
3235 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3240 .
addImm(SPIRV::GroupOperation::Reduce)
3247bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
3248 SPIRVTypeInst ResType,
3249 MachineInstr &
I)
const {
3254 Register InputReg =
I.getOperand(2).getReg();
3259 bool IsVector = NumElems > 1;
3272 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
3273 SPIRV::OpGroupNonUniformAllEqual);
3278 ElementResults.
reserve(NumElems);
3280 for (
unsigned Idx = 0; Idx < NumElems; ++Idx) {
3293 ElemInput = Extracted;
3299 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3310 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3321bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3322 SPIRVTypeInst ResType,
3323 MachineInstr &
I)
const {
3325 assert(
I.getNumOperands() == 3);
3327 auto Op =
I.getOperand(2);
3339 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3361 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3365 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3372bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3373 SPIRVTypeInst ResType,
3375 bool IsUnsigned)
const {
3376 return selectWaveReduce(
3377 ResVReg, ResType,
I, IsUnsigned,
3378 [&](
Register InputRegister,
bool IsUnsigned) {
3379 const bool IsFloatTy =
3381 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3382 : SPIRV::OpGroupNonUniformSMax;
3383 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3387bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3388 SPIRVTypeInst ResType,
3390 bool IsUnsigned)
const {
3391 return selectWaveReduce(
3392 ResVReg, ResType,
I, IsUnsigned,
3393 [&](
Register InputRegister,
bool IsUnsigned) {
3394 const bool IsFloatTy =
3396 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3397 : SPIRV::OpGroupNonUniformSMin;
3398 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3402bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3403 SPIRVTypeInst ResType,
3404 MachineInstr &
I)
const {
3405 return selectWaveReduce(ResVReg, ResType,
I,
false,
3406 [&](
Register InputRegister,
bool IsUnsigned) {
3408 InputRegister, SPIRV::OpTypeFloat);
3409 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3410 : SPIRV::OpGroupNonUniformIAdd;
3414bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3415 SPIRVTypeInst ResType,
3416 MachineInstr &
I)
const {
3417 return selectWaveReduce(ResVReg, ResType,
I,
false,
3418 [&](
Register InputRegister,
bool IsUnsigned) {
3420 InputRegister, SPIRV::OpTypeFloat);
3421 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3422 : SPIRV::OpGroupNonUniformIMul;
3426template <
typename PickOpcodeFn>
3427bool SPIRVInstructionSelector::selectWaveReduce(
3428 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3429 PickOpcodeFn &&PickOpcode)
const {
3430 assert(
I.getNumOperands() == 3);
3431 assert(
I.getOperand(2).isReg());
3433 Register InputRegister =
I.getOperand(2).getReg();
3440 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3446 .
addImm(SPIRV::GroupOperation::Reduce)
3447 .
addUse(
I.getOperand(2).getReg())
3452bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3453 SPIRVTypeInst ResType,
3455 unsigned Opcode)
const {
3456 return selectWaveReduce(
3457 ResVReg, ResType,
I,
false,
3458 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3461bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3462 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3463 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3464 [&](
Register InputRegister,
bool IsUnsigned) {
3466 InputRegister, SPIRV::OpTypeFloat);
3468 ? SPIRV::OpGroupNonUniformFAdd
3469 : SPIRV::OpGroupNonUniformIAdd;
3473bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3474 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3475 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3476 [&](
Register InputRegister,
bool IsUnsigned) {
3478 InputRegister, SPIRV::OpTypeFloat);
3480 ? SPIRV::OpGroupNonUniformFMul
3481 : SPIRV::OpGroupNonUniformIMul;
3485template <
typename PickOpcodeFn>
3486bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3487 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3488 PickOpcodeFn &&PickOpcode)
const {
3489 assert(
I.getNumOperands() == 3);
3490 assert(
I.getOperand(2).isReg());
3492 Register InputRegister =
I.getOperand(2).getReg();
3499 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3505 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3506 .
addUse(
I.getOperand(2).getReg())
3511bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
3512 SPIRVTypeInst ResType,
3515 assert(
I.getNumOperands() == 3);
3516 assert(
I.getOperand(2).isReg());
3518 Register InputRegister =
I.getOperand(2).getReg();
3524 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
3535bool SPIRVInstructionSelector::selectBitreverse16(
Register ResVReg,
3536 SPIRVTypeInst ResType,
3541 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
3546 : SPIRV::OpUConvert;
3550 ShiftOp = SPIRV::OpShiftRightLogicalV;
3555 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3556 TII.get(SPIRV::OpConstantComposite))
3559 for (
unsigned It = 0; It <
N; ++It)
3563 ShiftConst = CompositeReg;
3568 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
3573 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
3578 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
3583 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
3586bool SPIRVInstructionSelector::handle64BitOverflow(
3588 unsigned int Opcode,
3595 "handle64BitOverflow should only be used for integer types");
3597 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
3599 MachineIRBuilder MIRBuilder(
I);
3601 SPIRVTypeInst I64x2Type =
3603 SPIRVTypeInst Vec2ResType =
3606 std::vector<Register> PartialRegs;
3608 unsigned CurrentComponent = 0;
3609 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
3613 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3614 TII.get(SPIRV::OpVectorShuffle))
3619 .
addImm(CurrentComponent)
3620 .
addImm(CurrentComponent + 1);
3630 PartialRegs.push_back(SubVecReg);
3633 if (CurrentComponent != ComponentCount) {
3639 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
3640 SPIRV::OpVectorExtractDynamic))
3649 PartialRegs.push_back(FinalElemResReg);
3653 return selectOpWithSrcs(ResVReg, ResType,
I, std::move(PartialRegs),
3654 SPIRV::OpCompositeConstruct);
3657bool SPIRVInstructionSelector::selectBitreverse64(
Register ResVReg,
3658 SPIRVTypeInst ResType,
3662 if (ComponentCount > 2)
3663 return handle64BitOverflow(
3664 ResVReg, ResType,
I, SrcReg, SPIRV::OpBitReverse,
3666 unsigned O) {
return this->selectBitreverse64(R,
T,
I, S); });
3668 MachineIRBuilder MIRBuilder(
I);
3673 I32Type, 2 * ComponentCount, MIRBuilder,
false);
3677 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
3682 if (!selectBitreverseNative(Reverse32, VecI32Type,
I, Vec32))
3686 auto MaybeParts = splitEvenOddLanes(Reverse32, ComponentCount,
I, I32Type);
3689 SplitParts &Parts = *MaybeParts;
3696 if (!selectOpWithSrcs(SwappedVec, VecI32Type,
I, {Parts.High, Parts.Low},
3697 SPIRV::OpCompositeConstruct))
3701 return selectOpWithSrcs(ResVReg, ResType,
I, {SwappedVec}, SPIRV::OpBitcast);
3704bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
3705 SPIRVTypeInst ResType,
3709 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
3717bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
3718 SPIRVTypeInst ResType,
3719 MachineInstr &
I)
const {
3720 Register OpReg =
I.getOperand(1).getReg();
3728 return selectBitreverse16(ResVReg, ResType,
I, OpReg);
3730 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
3732 return selectBitreverse64(ResVReg, ResType,
I, OpReg);
3739 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
3750 unsigned AndOp = SPIRV::OpBitwiseAndS;
3751 unsigned OrOp = SPIRV::OpBitwiseOrS;
3752 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
3753 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
3755 AndOp = SPIRV::OpBitwiseAndV;
3756 OrOp = SPIRV::OpBitwiseOrV;
3757 ShlOp = SPIRV::OpShiftLeftLogicalV;
3758 ShrOp = SPIRV::OpShiftRightLogicalV;
3764 const unsigned Shift) ->
Register {
3772 Register MaskReg = CreateConst(Mask);
3773 Register ShiftReg = CreateConst(Shift);
3780 if (!selectOpWithSrcs(
T1, ResType,
I, {Input, ShiftReg}, ShrOp) ||
3781 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
3782 !selectOpWithSrcs(T3, ResType,
I, {Input, MaskReg}, AndOp) ||
3783 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
3784 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
3792 uint64_t
Mask = ~0ull;
3793 while ((Shift >>= 1) > 0) {
3800 return BuildCOPY(ResVReg, Result,
I);
3803bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
3804 SPIRVTypeInst ResType,
3805 MachineInstr &
I)
const {
3811 if (!
I.getOperand(0).isReg() || !
I.getOperand(1).isReg())
3813 Register OpReg =
I.getOperand(1).getReg();
3814 if (MachineInstr *Def = MRI->
getVRegDef(OpReg)) {
3815 if (
Def->getOpcode() == TargetOpcode::COPY)
3818 switch (
Def->getOpcode()) {
3819 case SPIRV::ASSIGN_TYPE:
3820 if (MachineInstr *AssignToDef =
3822 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
3823 Reg =
Def->getOperand(2).getReg();
3826 case SPIRV::OpUndef:
3827 Reg =
Def->getOperand(1).getReg();
3830 unsigned DestOpCode;
3832 DestOpCode = SPIRV::OpConstantNull;
3834 DestOpCode = TargetOpcode::COPY;
3837 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
3838 .
addDef(
I.getOperand(0).getReg())
3846bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
3847 SPIRVTypeInst ResType,
3848 MachineInstr &
I)
const {
3850 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3852 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
3856 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
3861 for (
unsigned i =
I.getNumExplicitDefs();
3862 i <
I.getNumExplicitOperands() && IsConst; ++i)
3866 if (!IsConst &&
N < 2)
3868 "There must be at least two constituent operands in a vector");
3871 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3872 TII.get(IsConst ? SPIRV::OpConstantComposite
3873 : SPIRV::OpCompositeConstruct))
3876 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
3877 MIB.
addUse(
I.getOperand(i).getReg());
3882bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
3883 SPIRVTypeInst ResType,
3884 MachineInstr &
I)
const {
3886 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3888 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
3894 if (!
I.getOperand(
OpIdx).isReg())
3901 if (!IsConst &&
N < 2)
3903 "There must be at least two constituent operands in a vector");
3906 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3907 TII.get(IsConst ? SPIRV::OpConstantComposite
3908 : SPIRV::OpCompositeConstruct))
3911 for (
unsigned i = 0; i <
N; ++i)
3917bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
3918 SPIRVTypeInst ResType,
3919 MachineInstr &
I)
const {
3924 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
3926 Opcode = SPIRV::OpDemoteToHelperInvocation;
3928 Opcode = SPIRV::OpKill;
3930 if (MachineInstr *NextI =
I.getNextNode()) {
3932 NextI->eraseFromParent();
3942bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
3943 SPIRVTypeInst ResType,
unsigned CmpOpc,
3944 MachineInstr &
I)
const {
3945 Register Cmp0 =
I.getOperand(2).getReg();
3946 Register Cmp1 =
I.getOperand(3).getReg();
3949 "CMP operands should have the same type");
3950 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
3960bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
3961 SPIRVTypeInst ResType,
3962 MachineInstr &
I)
const {
3963 auto Pred =
I.getOperand(1).getPredicate();
3966 Register CmpOperand =
I.getOperand(2).getReg();
3973 return selectCmp(ResVReg, ResType, CmpOpc,
I);
3977SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
3978 SPIRVTypeInst ResType)
const {
3980 SPIRVTypeInst SpvI32Ty =
3983 auto ConstInt = ConstantInt::get(LLVMTy, Val);
3990 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
3993 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
3996 .
addImm(APInt(32, Val).getZExtValue());
3998 GR.
add(ConstInt,
MI);
4003bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
4004 SPIRVTypeInst ResType,
4005 MachineInstr &
I)
const {
4007 return selectCmp(ResVReg, ResType, CmpOp,
I);
4010bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
4011 SPIRVTypeInst ResType,
4012 MachineInstr &
I)
const {
4014 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
4021 if (ResType->
getOpcode() != SPIRV::OpTypeVector &&
4022 ResType->
getOpcode() != SPIRV::OpTypeFloat)
4025 MachineIRBuilder MIRBuilder(
I);
4032 APFloat ConstVal(3.3219280948873623);
4036 APFloat::rmNearestTiesToEven, &LosesInfo);
4040 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
4041 ? SPIRV::OpVectorTimesScalar
4044 if (!selectOpWithSrcs(ArgReg, ResType,
I,
4045 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
4047 if (!selectExtInst(ResVReg, ResType,
I,
4048 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
4058Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
4059 MachineInstr &
I)
const {
4062 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4067bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
4073 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
4081 if (
Def->getOpcode() == SPIRV::OpConstantNull)
4084 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
4085 Def->getOpcode() == SPIRV::OpConstantI)
4098 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
4099 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
4101 Intrinsic::spv_const_composite)) {
4102 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
4103 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
4104 if (!IsZero(
Def->getOperand(i).getReg()))
4113Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
4114 MachineInstr &
I)
const {
4118 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4123Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
4124 MachineInstr &
I)
const {
4128 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4134 SPIRVTypeInst ResType,
4135 MachineInstr &
I)
const {
4139 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4144bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
4145 SPIRVTypeInst ResType,
4146 MachineInstr &
I)
const {
4147 Register SelectFirstArg =
I.getOperand(2).getReg();
4148 Register SelectSecondArg =
I.getOperand(3).getReg();
4157 SPIRV::OpTypeVector;
4164 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
4165 }
else if (IsPtrTy) {
4166 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
4168 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
4172 Opcode = IsScalarBool ? SPIRV::OpSelectSFSCond : SPIRV::OpSelectVFVCond;
4173 }
else if (IsPtrTy) {
4174 Opcode = IsScalarBool ? SPIRV::OpSelectSPSCond : SPIRV::OpSelectVPVCond;
4176 Opcode = IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4179 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4182 .
addUse(
I.getOperand(1).getReg())
4191bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
4192 SPIRVTypeInst ResType,
4194 MachineInstr &InsertAt,
4195 bool IsSigned)
const {
4197 Register ZeroReg = buildZerosVal(ResType, InsertAt);
4198 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
4199 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
4201 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4213bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
4214 SPIRVTypeInst ResType,
4215 MachineInstr &
I,
bool IsSigned,
4216 unsigned Opcode)
const {
4217 Register SrcReg =
I.getOperand(1).getReg();
4223 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
4228 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I,
false);
4230 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
4233bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
4234 SPIRVTypeInst ResType, MachineInstr &
I,
4235 bool IsSigned)
const {
4236 Register SrcReg =
I.getOperand(1).getReg();
4238 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
4242 if (ResType == SrcType)
4243 return BuildCOPY(ResVReg, SrcReg,
I);
4245 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4246 return selectUnOp(ResVReg, ResType,
I, Opcode);
4249bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
4250 SPIRVTypeInst ResType,
4252 bool IsSigned)
const {
4253 MachineIRBuilder MIRBuilder(
I);
4254 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
4269 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
4272 .
addUse(
I.getOperand(1).getReg())
4273 .
addUse(
I.getOperand(2).getReg())
4279 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
4282 .
addUse(
I.getOperand(1).getReg())
4283 .
addUse(
I.getOperand(2).getReg())
4291 unsigned SelectOpcode =
4292 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
4297 .
addUse(buildOnesVal(
true, ResType,
I))
4298 .
addUse(buildZerosVal(ResType,
I))
4305 .
addUse(buildOnesVal(
false, ResType,
I))
4310bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
4313 SPIRVTypeInst IntTy,
4314 SPIRVTypeInst BoolTy)
const {
4317 bool IsVectorTy = IntTy->
getOpcode() == SPIRV::OpTypeVector;
4318 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
4320 Register One = buildOnesVal(
false, IntTy,
I);
4328 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
4337bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
4338 SPIRVTypeInst ResType,
4339 MachineInstr &
I)
const {
4340 Register IntReg =
I.getOperand(1).getReg();
4343 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
4344 if (ArgType == ResType)
4345 return BuildCOPY(ResVReg, IntReg,
I);
4347 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4348 return selectUnOp(ResVReg, ResType,
I, Opcode);
4351bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
4352 SPIRVTypeInst ResType,
4353 MachineInstr &
I)
const {
4354 unsigned Opcode =
I.getOpcode();
4355 unsigned TpOpcode = ResType->
getOpcode();
4357 if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {
4358 assert(Opcode == TargetOpcode::G_CONSTANT &&
4359 I.getOperand(1).getCImm()->isZero());
4360 MachineBasicBlock &DepMBB =
I.getMF()->front();
4363 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
4370 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
4373bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
4374 SPIRVTypeInst ResType,
4375 MachineInstr &
I)
const {
4376 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4383bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
4384 SPIRVTypeInst ResType,
4385 MachineInstr &
I)
const {
4387 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
4391 .
addUse(
I.getOperand(3).getReg())
4393 .
addUse(
I.getOperand(2).getReg());
4394 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
4400bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
4401 SPIRVTypeInst ResType,
4402 MachineInstr &
I)
const {
4403 Type *MaybeResTy =
nullptr;
4408 "Expected aggregate type for extractv instruction");
4410 SPIRV::AccessQualifier::ReadWrite,
false);
4414 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
4417 .
addUse(
I.getOperand(2).getReg());
4418 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
4424bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
4425 SPIRVTypeInst ResType,
4426 MachineInstr &
I)
const {
4427 if (
getImm(
I.getOperand(4), MRI))
4428 return selectInsertVal(ResVReg, ResType,
I);
4430 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
4433 .
addUse(
I.getOperand(2).getReg())
4434 .
addUse(
I.getOperand(3).getReg())
4435 .
addUse(
I.getOperand(4).getReg())
4440bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
4441 SPIRVTypeInst ResType,
4442 MachineInstr &
I)
const {
4443 if (
getImm(
I.getOperand(3), MRI))
4444 return selectExtractVal(ResVReg, ResType,
I);
4446 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
4449 .
addUse(
I.getOperand(2).getReg())
4450 .
addUse(
I.getOperand(3).getReg())
4455bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
4456 SPIRVTypeInst ResType,
4457 MachineInstr &
I)
const {
4458 const bool IsGEPInBounds =
I.getOperand(2).getImm();
4464 ? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
4465 : SPIRV::OpAccessChain)
4466 : (IsGEPInBounds ?
SPIRV::OpInBoundsPtrAccessChain
4467 :
SPIRV::OpPtrAccessChain);
4469 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4473 .
addUse(
I.getOperand(3).getReg());
4475 (Opcode == SPIRV::OpPtrAccessChain ||
4476 Opcode == SPIRV::OpInBoundsPtrAccessChain ||
4477 (
getImm(
I.getOperand(4), MRI) &&
foldImm(
I.getOperand(4), MRI) == 0)) &&
4478 "Cannot translate GEP to OpAccessChain. First index must be 0.");
4481 const unsigned StartingIndex =
4482 (Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)
4485 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
4486 Res.addUse(
I.getOperand(i).getReg());
4487 Res.constrainAllUses(
TII,
TRI, RBI);
4492bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
4494 unsigned Lim =
I.getNumExplicitOperands();
4495 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
4496 Register OpReg =
I.getOperand(i).getReg();
4497 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
4499 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
4500 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
4501 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
4508 MachineFunction *MF =
I.getMF();
4520 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4521 TII.get(SPIRV::OpSpecConstantOp))
4524 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
4526 GR.
add(OpDefine, MIB);
4532bool SPIRVInstructionSelector::selectDerivativeInst(
4533 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
4534 const unsigned DPdOpCode)
const {
4537 errorIfInstrOutsideShader(
I);
4542 Register SrcReg =
I.getOperand(2).getReg();
4547 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4550 .
addUse(
I.getOperand(2).getReg());
4552 MachineIRBuilder MIRBuilder(
I);
4555 if (componentCount != 1)
4559 const TargetRegisterClass *RegClass = GR.
getRegClass(SrcType);
4563 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4568 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4573 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4581bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
4582 SPIRVTypeInst ResType,
4583 MachineInstr &
I)
const {
4587 case Intrinsic::spv_load:
4588 return selectLoad(ResVReg, ResType,
I);
4589 case Intrinsic::spv_atomic_load:
4590 return selectAtomicLoad(ResVReg, ResType,
I);
4591 case Intrinsic::spv_store:
4592 return selectStore(
I);
4593 case Intrinsic::spv_atomic_store:
4594 return selectAtomicStore(
I);
4595 case Intrinsic::spv_extractv:
4596 return selectExtractVal(ResVReg, ResType,
I);
4597 case Intrinsic::spv_insertv:
4598 return selectInsertVal(ResVReg, ResType,
I);
4599 case Intrinsic::spv_extractelt:
4600 return selectExtractElt(ResVReg, ResType,
I);
4601 case Intrinsic::spv_insertelt:
4602 return selectInsertElt(ResVReg, ResType,
I);
4603 case Intrinsic::spv_gep:
4604 return selectGEP(ResVReg, ResType,
I);
4605 case Intrinsic::spv_bitcast: {
4606 Register OpReg =
I.getOperand(2).getReg();
4607 SPIRVTypeInst OpType =
4611 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
4613 case Intrinsic::spv_unref_global:
4614 case Intrinsic::spv_init_global: {
4615 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
4620 Register GVarVReg =
MI->getOperand(0).getReg();
4621 if (!selectGlobalValue(GVarVReg, *
MI, Init))
4626 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
4628 MI->eraseFromParent();
4632 case Intrinsic::spv_undef: {
4633 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4639 case Intrinsic::spv_named_boolean_spec_constant: {
4640 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
4641 : SPIRV::OpSpecConstantFalse;
4643 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
4644 .
addDef(
I.getOperand(0).getReg())
4647 unsigned SpecId =
I.getOperand(2).getImm();
4649 SPIRV::Decoration::SpecId, {SpecId});
4653 case Intrinsic::spv_const_composite: {
4655 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
4661 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
4663 std::function<bool(
Register)> HasSpecConstOperand =
4673 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
4674 J < Def->getNumExplicitOperands(); ++J) {
4675 if (
Def->getOperand(J).isReg() &&
4676 HasSpecConstOperand(
Def->getOperand(J).getReg()))
4682 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
4683 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
4684 : SPIRV::OpConstantComposite;
4685 unsigned ContinuedOpc = HasSpecConst
4686 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
4687 : SPIRV::OpConstantCompositeContinuedINTEL;
4688 MachineIRBuilder MIR(
I);
4690 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
4692 for (
auto *Instr : Instructions) {
4693 Instr->setDebugLoc(
I.getDebugLoc());
4698 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4705 case Intrinsic::spv_assign_name: {
4706 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
4707 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
4708 for (
unsigned i =
I.getNumExplicitDefs() + 2;
4709 i <
I.getNumExplicitOperands(); ++i) {
4710 MIB.
addImm(
I.getOperand(i).getImm());
4715 case Intrinsic::spv_switch: {
4716 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
4717 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
4718 if (
I.getOperand(i).isReg())
4719 MIB.
addReg(
I.getOperand(i).getReg());
4720 else if (
I.getOperand(i).isCImm())
4721 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
4722 else if (
I.getOperand(i).isMBB())
4723 MIB.
addMBB(
I.getOperand(i).getMBB());
4730 case Intrinsic::spv_loop_merge: {
4731 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
4732 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
4733 if (
I.getOperand(i).isMBB())
4734 MIB.
addMBB(
I.getOperand(i).getMBB());
4741 case Intrinsic::spv_loop_control_intel: {
4743 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
4744 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
4749 case Intrinsic::spv_selection_merge: {
4751 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
4752 assert(
I.getOperand(1).isMBB() &&
4753 "operand 1 to spv_selection_merge must be a basic block");
4754 MIB.
addMBB(
I.getOperand(1).getMBB());
4755 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
4759 case Intrinsic::spv_cmpxchg:
4760 return selectAtomicCmpXchg(ResVReg, ResType,
I);
4761 case Intrinsic::spv_unreachable:
4762 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
4765 case Intrinsic::spv_alloca:
4766 return selectFrameIndex(ResVReg, ResType,
I);
4767 case Intrinsic::spv_alloca_array:
4768 return selectAllocaArray(ResVReg, ResType,
I);
4769 case Intrinsic::spv_assume:
4771 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
4772 .
addUse(
I.getOperand(1).getReg())
4777 case Intrinsic::spv_expect:
4779 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
4782 .
addUse(
I.getOperand(2).getReg())
4783 .
addUse(
I.getOperand(3).getReg())
4788 case Intrinsic::arithmetic_fence:
4789 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
4790 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
4793 .
addUse(
I.getOperand(2).getReg())
4797 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
4799 case Intrinsic::spv_thread_id:
4805 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
4807 case Intrinsic::spv_thread_id_in_group:
4813 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
4815 case Intrinsic::spv_group_id:
4821 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
4823 case Intrinsic::spv_flattened_thread_id_in_group:
4830 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
4832 case Intrinsic::spv_workgroup_size:
4833 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
4835 case Intrinsic::spv_global_size:
4836 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
4838 case Intrinsic::spv_global_offset:
4839 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
4841 case Intrinsic::spv_num_workgroups:
4842 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
4844 case Intrinsic::spv_subgroup_size:
4845 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
4847 case Intrinsic::spv_num_subgroups:
4848 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
4850 case Intrinsic::spv_subgroup_id:
4851 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
4852 case Intrinsic::spv_subgroup_local_invocation_id:
4853 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
4854 ResVReg, ResType,
I);
4855 case Intrinsic::spv_subgroup_max_size:
4856 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
4858 case Intrinsic::spv_fdot:
4859 return selectFloatDot(ResVReg, ResType,
I);
4860 case Intrinsic::spv_udot:
4861 case Intrinsic::spv_sdot:
4862 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4864 return selectIntegerDot(ResVReg, ResType,
I,
4865 IID == Intrinsic::spv_sdot);
4866 return selectIntegerDotExpansion(ResVReg, ResType,
I);
4867 case Intrinsic::spv_dot4add_i8packed:
4868 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4870 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
4871 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
4872 case Intrinsic::spv_dot4add_u8packed:
4873 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4875 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
4876 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
4877 case Intrinsic::spv_all:
4878 return selectAll(ResVReg, ResType,
I);
4879 case Intrinsic::spv_any:
4880 return selectAny(ResVReg, ResType,
I);
4881 case Intrinsic::spv_cross:
4882 return selectExtInst(ResVReg, ResType,
I, CL::cross, GL::Cross);
4883 case Intrinsic::spv_distance:
4884 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
4885 case Intrinsic::spv_lerp:
4886 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
4887 case Intrinsic::spv_length:
4888 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
4889 case Intrinsic::spv_degrees:
4890 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
4891 case Intrinsic::spv_faceforward:
4892 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
4893 case Intrinsic::spv_frac:
4894 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
4895 case Intrinsic::spv_isinf:
4896 return selectOpIsInf(ResVReg, ResType,
I);
4897 case Intrinsic::spv_isnan:
4898 return selectOpIsNan(ResVReg, ResType,
I);
4899 case Intrinsic::spv_normalize:
4900 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
4901 case Intrinsic::spv_refract:
4902 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
4903 case Intrinsic::spv_reflect:
4904 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
4905 case Intrinsic::spv_rsqrt:
4906 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
4907 case Intrinsic::spv_sign:
4908 return selectSign(ResVReg, ResType,
I);
4909 case Intrinsic::spv_smoothstep:
4910 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
4911 case Intrinsic::spv_firstbituhigh:
4912 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
4913 case Intrinsic::spv_firstbitshigh:
4914 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
4915 case Intrinsic::spv_firstbitlow:
4916 return selectFirstBitLow(ResVReg, ResType,
I);
4917 case Intrinsic::spv_all_memory_barrier:
4918 return selectBarrierInst(
I, SPIRV::Scope::Device,
4919 SPIRV::MemorySemantics::UniformMemory |
4920 SPIRV::MemorySemantics::ImageMemory |
4921 SPIRV::MemorySemantics::WorkgroupMemory,
4923 case Intrinsic::spv_all_memory_barrier_with_group_sync:
4924 return selectBarrierInst(
I, SPIRV::Scope::Device,
4925 SPIRV::MemorySemantics::UniformMemory |
4926 SPIRV::MemorySemantics::ImageMemory |
4927 SPIRV::MemorySemantics::WorkgroupMemory,
4929 case Intrinsic::spv_device_memory_barrier:
4930 return selectBarrierInst(
I, SPIRV::Scope::Device,
4931 SPIRV::MemorySemantics::UniformMemory |
4932 SPIRV::MemorySemantics::ImageMemory,
4934 case Intrinsic::spv_device_memory_barrier_with_group_sync:
4935 return selectBarrierInst(
I, SPIRV::Scope::Device,
4936 SPIRV::MemorySemantics::UniformMemory |
4937 SPIRV::MemorySemantics::ImageMemory,
4939 case Intrinsic::spv_group_memory_barrier:
4940 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
4941 SPIRV::MemorySemantics::WorkgroupMemory,
4943 case Intrinsic::spv_group_memory_barrier_with_group_sync:
4944 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
4945 SPIRV::MemorySemantics::WorkgroupMemory,
4947 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
4948 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
4949 SPIRV::StorageClass::StorageClass ResSC =
4953 "Generic storage class");
4954 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
4962 case Intrinsic::spv_lifetime_start:
4963 case Intrinsic::spv_lifetime_end: {
4964 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
4965 : SPIRV::OpLifetimeStop;
4966 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
4967 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
4976 case Intrinsic::spv_saturate:
4977 return selectSaturate(ResVReg, ResType,
I);
4978 case Intrinsic::spv_nclamp:
4979 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
4980 case Intrinsic::spv_uclamp:
4981 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
4982 case Intrinsic::spv_sclamp:
4983 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
4984 case Intrinsic::spv_subgroup_prefix_bit_count:
4985 return selectWavePrefixBitCount(ResVReg, ResType,
I);
4986 case Intrinsic::spv_wave_active_countbits:
4987 return selectWaveActiveCountBits(ResVReg, ResType,
I);
4988 case Intrinsic::spv_wave_all_equal:
4989 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
4990 case Intrinsic::spv_wave_all:
4991 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
4992 case Intrinsic::spv_wave_any:
4993 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
4994 case Intrinsic::spv_subgroup_ballot:
4995 return selectWaveOpInst(ResVReg, ResType,
I,
4996 SPIRV::OpGroupNonUniformBallot);
4997 case Intrinsic::spv_wave_is_first_lane:
4998 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
4999 case Intrinsic::spv_wave_reduce_or:
5000 return selectWaveReduceOp(ResVReg, ResType,
I,
5001 SPIRV::OpGroupNonUniformBitwiseOr);
5002 case Intrinsic::spv_wave_reduce_xor:
5003 return selectWaveReduceOp(ResVReg, ResType,
I,
5004 SPIRV::OpGroupNonUniformBitwiseXor);
5005 case Intrinsic::spv_wave_reduce_and:
5006 return selectWaveReduceOp(ResVReg, ResType,
I,
5007 SPIRV::OpGroupNonUniformBitwiseAnd);
5008 case Intrinsic::spv_wave_reduce_umax:
5009 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
5010 case Intrinsic::spv_wave_reduce_max:
5011 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
5012 case Intrinsic::spv_wave_reduce_umin:
5013 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
5014 case Intrinsic::spv_wave_reduce_min:
5015 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
5016 case Intrinsic::spv_wave_reduce_sum:
5017 return selectWaveReduceSum(ResVReg, ResType,
I);
5018 case Intrinsic::spv_wave_product:
5019 return selectWaveReduceProduct(ResVReg, ResType,
I);
5020 case Intrinsic::spv_wave_readlane:
5021 return selectWaveOpInst(ResVReg, ResType,
I,
5022 SPIRV::OpGroupNonUniformShuffle);
5023 case Intrinsic::spv_wave_prefix_sum:
5024 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
5025 case Intrinsic::spv_wave_prefix_product:
5026 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
5027 case Intrinsic::spv_quad_read_across_x: {
5028 return selectQuadSwap(ResVReg, ResType,
I, 0);
5030 case Intrinsic::spv_quad_read_across_y: {
5031 return selectQuadSwap(ResVReg, ResType,
I, 1);
5033 case Intrinsic::spv_step:
5034 return selectExtInst(ResVReg, ResType,
I, CL::step, GL::Step);
5035 case Intrinsic::spv_radians:
5036 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
5040 case Intrinsic::instrprof_increment:
5041 case Intrinsic::instrprof_increment_step:
5042 case Intrinsic::instrprof_value_profile:
5045 case Intrinsic::spv_value_md:
5047 case Intrinsic::spv_resource_handlefrombinding: {
5048 return selectHandleFromBinding(ResVReg, ResType,
I);
5050 case Intrinsic::spv_resource_counterhandlefrombinding:
5051 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
5052 case Intrinsic::spv_resource_updatecounter:
5053 return selectUpdateCounter(ResVReg, ResType,
I);
5054 case Intrinsic::spv_resource_store_typedbuffer: {
5055 return selectImageWriteIntrinsic(
I);
5057 case Intrinsic::spv_resource_load_typedbuffer: {
5058 return selectReadImageIntrinsic(ResVReg, ResType,
I);
5060 case Intrinsic::spv_resource_load_level: {
5061 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
5063 case Intrinsic::spv_resource_getdimensions_x:
5064 case Intrinsic::spv_resource_getdimensions_xy:
5065 case Intrinsic::spv_resource_getdimensions_xyz: {
5066 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
5068 case Intrinsic::spv_resource_getdimensions_levels_x:
5069 case Intrinsic::spv_resource_getdimensions_levels_xy:
5070 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
5071 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
5073 case Intrinsic::spv_resource_getdimensions_ms_xy:
5074 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
5075 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
5077 case Intrinsic::spv_resource_calculate_lod:
5078 case Intrinsic::spv_resource_calculate_lod_unclamped:
5079 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
5080 case Intrinsic::spv_resource_sample:
5081 case Intrinsic::spv_resource_sample_clamp:
5082 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
5083 case Intrinsic::spv_resource_samplebias:
5084 case Intrinsic::spv_resource_samplebias_clamp:
5085 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
5086 case Intrinsic::spv_resource_samplegrad:
5087 case Intrinsic::spv_resource_samplegrad_clamp:
5088 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
5089 case Intrinsic::spv_resource_samplelevel:
5090 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
5091 case Intrinsic::spv_resource_samplecmp:
5092 case Intrinsic::spv_resource_samplecmp_clamp:
5093 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
5094 case Intrinsic::spv_resource_samplecmplevelzero:
5095 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
5096 case Intrinsic::spv_resource_gather:
5097 case Intrinsic::spv_resource_gather_cmp:
5098 return selectGatherIntrinsic(ResVReg, ResType,
I);
5099 case Intrinsic::spv_resource_getpointer: {
5100 return selectResourceGetPointer(ResVReg, ResType,
I);
5102 case Intrinsic::spv_pushconstant_getpointer: {
5103 return selectPushConstantGetPointer(ResVReg, ResType,
I);
5105 case Intrinsic::spv_discard: {
5106 return selectDiscard(ResVReg, ResType,
I);
5108 case Intrinsic::spv_resource_nonuniformindex: {
5109 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
5111 case Intrinsic::spv_unpackhalf2x16: {
5112 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
5114 case Intrinsic::spv_packhalf2x16: {
5115 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
5117 case Intrinsic::spv_ddx:
5118 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
5119 case Intrinsic::spv_ddy:
5120 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
5121 case Intrinsic::spv_ddx_coarse:
5122 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
5123 case Intrinsic::spv_ddy_coarse:
5124 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
5125 case Intrinsic::spv_ddx_fine:
5126 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
5127 case Intrinsic::spv_ddy_fine:
5128 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
5129 case Intrinsic::spv_fwidth:
5130 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
5131 case Intrinsic::spv_masked_gather:
5132 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5133 return selectMaskedGather(ResVReg, ResType,
I);
5134 return diagnoseUnsupported(
5135 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
5136 case Intrinsic::spv_masked_scatter:
5137 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5138 return selectMaskedScatter(
I);
5139 return diagnoseUnsupported(
5140 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
5142 std::string DiagMsg;
5143 raw_string_ostream OS(DiagMsg);
5145 DiagMsg =
"Intrinsic selection not implemented: " + DiagMsg;
5152bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
5153 SPIRVTypeInst ResType,
5154 MachineInstr &
I)
const {
5157 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
5164bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
5165 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5167 assert(Intr.getIntrinsicID() ==
5168 Intrinsic::spv_resource_counterhandlefrombinding);
5171 Register MainHandleReg = Intr.getOperand(2).getReg();
5173 assert(MainHandleDef->getIntrinsicID() ==
5174 Intrinsic::spv_resource_handlefrombinding);
5178 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
5179 Register IndexReg = MainHandleDef->getOperand(5).getReg();
5180 std::string CounterName =
5185 MachineIRBuilder MIRBuilder(
I);
5187 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
5189 ArraySize, IndexReg, CounterName, MIRBuilder);
5191 return BuildCOPY(ResVReg, CounterVarReg,
I);
5194bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
5195 SPIRVTypeInst ResType,
5196 MachineInstr &
I)
const {
5198 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
5200 Register CounterHandleReg = Intr.getOperand(2).getReg();
5201 Register IncrReg = Intr.getOperand(3).getReg();
5208 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
5209 assert(CounterVarPointeeType &&
5210 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
5211 "Counter variable must be a struct");
5213 SPIRV::StorageClass::StorageBuffer &&
5214 "Counter variable must be in the storage buffer storage class");
5216 "Counter variable must have exactly 1 member in the struct");
5217 const SPIRVTypeInst MemberType =
5220 "Counter variable struct must have a single i32 member");
5224 MachineIRBuilder MIRBuilder(
I);
5226 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
5229 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
5235 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5238 .
addUse(CounterHandleReg)
5245 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
5248 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
5251 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
5260 return BuildCOPY(ResVReg, AtomicRes,
I);
5268 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
5276bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
5277 SPIRVTypeInst ResType,
5278 MachineInstr &
I)
const {
5286 Register ImageReg =
I.getOperand(2).getReg();
5294 Register IdxReg =
I.getOperand(3).getReg();
5296 MachineInstr &Pos =
I;
5298 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
5302bool SPIRVInstructionSelector::generateSampleImage(
5305 DebugLoc Loc, MachineInstr &Pos)
const {
5316 if (!loadHandleBeforePosition(NewSamplerReg,
5322 MachineIRBuilder MIRBuilder(Pos);
5335 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
5336 ImOps.Lod.has_value();
5337 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
5338 : SPIRV::OpImageSampleImplicitLod;
5340 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
5341 : SPIRV::OpImageSampleDrefImplicitLod;
5350 MIB.
addUse(*ImOps.Compare);
5352 uint32_t ImageOperands = 0;
5354 ImageOperands |= SPIRV::ImageOperand::Bias;
5356 ImageOperands |= SPIRV::ImageOperand::Lod;
5357 if (ImOps.GradX && ImOps.GradY)
5358 ImageOperands |= SPIRV::ImageOperand::Grad;
5359 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
5361 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
5364 "Non-constant offsets are not supported in sample instructions.");
5368 ImageOperands |= SPIRV::ImageOperand::MinLod;
5370 if (ImageOperands != 0) {
5371 MIB.
addImm(ImageOperands);
5372 if (ImageOperands & SPIRV::ImageOperand::Bias)
5374 if (ImageOperands & SPIRV::ImageOperand::Lod)
5376 if (ImageOperands & SPIRV::ImageOperand::Grad) {
5377 MIB.
addUse(*ImOps.GradX);
5378 MIB.
addUse(*ImOps.GradY);
5381 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
5382 MIB.
addUse(*ImOps.Offset);
5383 if (ImageOperands & SPIRV::ImageOperand::MinLod)
5384 MIB.
addUse(*ImOps.MinLod);
5391bool SPIRVInstructionSelector::selectImageQuerySize(
5393 std::optional<Register> LodReg)
const {
5395 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
5398 "ImageReg is not an image type.");
5400 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
5402 unsigned NumComponents = 0;
5404 case SPIRV::Dim::DIM_1D:
5405 case SPIRV::Dim::DIM_Buffer:
5406 NumComponents =
IsArray ? 2 : 1;
5408 case SPIRV::Dim::DIM_2D:
5409 case SPIRV::Dim::DIM_Cube:
5410 case SPIRV::Dim::DIM_Rect:
5411 NumComponents =
IsArray ? 3 : 2;
5413 case SPIRV::Dim::DIM_3D:
5417 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
5422 SPIRVTypeInst ResType =
5427 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5437bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
5438 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5439 Register ImageReg =
I.getOperand(2).getReg();
5446 return selectImageQuerySize(NewImageReg, ResVReg,
I);
5449bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
5450 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5451 Register ImageReg =
I.getOperand(2).getReg();
5460 Register LodReg =
I.getOperand(3).getReg();
5463 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
5465 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
5472 TII.get(SPIRV::OpImageQueryLevels))
5479 TII.get(SPIRV::OpCompositeConstruct))
5489bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
5490 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5491 Register ImageReg =
I.getOperand(2).getReg();
5502 "OpImageQuerySamples requires a multisampled image");
5504 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
5512 TII.get(SPIRV::OpImageQuerySamples))
5519 TII.get(SPIRV::OpCompositeConstruct))
5529bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
5530 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5531 Register ImageReg =
I.getOperand(2).getReg();
5532 Register SamplerReg =
I.getOperand(3).getReg();
5533 Register CoordinateReg =
I.getOperand(4).getReg();
5549 if (!loadHandleBeforePosition(
5554 MachineIRBuilder MIRBuilder(
I);
5560 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
5570 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
5577 unsigned ExtractedIndex =
5579 Intrinsic::spv_resource_calculate_lod_unclamped
5583 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5584 TII.get(SPIRV::OpCompositeExtract))
5594bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
5595 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5596 Register ImageReg =
I.getOperand(2).getReg();
5597 Register SamplerReg =
I.getOperand(3).getReg();
5598 Register CoordinateReg =
I.getOperand(4).getReg();
5599 ImageOperands ImOps;
5600 if (
I.getNumOperands() > 5)
5601 ImOps.Offset =
I.getOperand(5).getReg();
5602 if (
I.getNumOperands() > 6)
5603 ImOps.MinLod =
I.getOperand(6).getReg();
5604 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5605 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5608bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
5609 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5610 Register ImageReg =
I.getOperand(2).getReg();
5611 Register SamplerReg =
I.getOperand(3).getReg();
5612 Register CoordinateReg =
I.getOperand(4).getReg();
5613 ImageOperands ImOps;
5614 ImOps.Bias =
I.getOperand(5).getReg();
5615 if (
I.getNumOperands() > 6)
5616 ImOps.Offset =
I.getOperand(6).getReg();
5617 if (
I.getNumOperands() > 7)
5618 ImOps.MinLod =
I.getOperand(7).getReg();
5619 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5620 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5623bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
5624 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5625 Register ImageReg =
I.getOperand(2).getReg();
5626 Register SamplerReg =
I.getOperand(3).getReg();
5627 Register CoordinateReg =
I.getOperand(4).getReg();
5628 ImageOperands ImOps;
5629 ImOps.GradX =
I.getOperand(5).getReg();
5630 ImOps.GradY =
I.getOperand(6).getReg();
5631 if (
I.getNumOperands() > 7)
5632 ImOps.Offset =
I.getOperand(7).getReg();
5633 if (
I.getNumOperands() > 8)
5634 ImOps.MinLod =
I.getOperand(8).getReg();
5635 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5636 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5639bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
5640 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5641 Register ImageReg =
I.getOperand(2).getReg();
5642 Register SamplerReg =
I.getOperand(3).getReg();
5643 Register CoordinateReg =
I.getOperand(4).getReg();
5644 ImageOperands ImOps;
5645 ImOps.Lod =
I.getOperand(5).getReg();
5646 if (
I.getNumOperands() > 6)
5647 ImOps.Offset =
I.getOperand(6).getReg();
5648 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5649 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5652bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
5653 SPIRVTypeInst ResType,
5654 MachineInstr &
I)
const {
5655 Register ImageReg =
I.getOperand(2).getReg();
5656 Register SamplerReg =
I.getOperand(3).getReg();
5657 Register CoordinateReg =
I.getOperand(4).getReg();
5658 ImageOperands ImOps;
5659 ImOps.Compare =
I.getOperand(5).getReg();
5660 if (
I.getNumOperands() > 6)
5661 ImOps.Offset =
I.getOperand(6).getReg();
5662 if (
I.getNumOperands() > 7)
5663 ImOps.MinLod =
I.getOperand(7).getReg();
5664 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5665 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5668bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
5669 SPIRVTypeInst ResType,
5670 MachineInstr &
I)
const {
5671 Register ImageReg =
I.getOperand(2).getReg();
5672 Register CoordinateReg =
I.getOperand(3).getReg();
5673 Register LodReg =
I.getOperand(4).getReg();
5675 ImageOperands ImOps;
5677 if (
I.getNumOperands() > 5)
5678 ImOps.Offset =
I.getOperand(5).getReg();
5690 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
5691 I.getDebugLoc(),
I, &ImOps);
5694bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
5695 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5696 Register ImageReg =
I.getOperand(2).getReg();
5697 Register SamplerReg =
I.getOperand(3).getReg();
5698 Register CoordinateReg =
I.getOperand(4).getReg();
5699 ImageOperands ImOps;
5700 ImOps.Compare =
I.getOperand(5).getReg();
5701 if (
I.getNumOperands() > 6)
5702 ImOps.Offset =
I.getOperand(6).getReg();
5705 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5706 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5709bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
5710 SPIRVTypeInst ResType,
5711 MachineInstr &
I)
const {
5712 Register ImageReg =
I.getOperand(2).getReg();
5713 Register SamplerReg =
I.getOperand(3).getReg();
5714 Register CoordinateReg =
I.getOperand(4).getReg();
5717 "ImageReg is not an image type.");
5722 ComponentOrCompareReg =
I.getOperand(5).getReg();
5723 OffsetReg =
I.getOperand(6).getReg();
5726 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
5730 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
5731 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
5732 Dim != SPIRV::Dim::DIM_Rect) {
5734 "Gather operations are only supported for 2D, Cube, and Rect images.");
5741 if (!loadHandleBeforePosition(
5746 MachineIRBuilder MIRBuilder(
I);
5747 SPIRVTypeInst SampledImageType =
5752 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
5760 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
5762 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
5764 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5769 .
addUse(ComponentOrCompareReg);
5771 uint32_t ImageOperands = 0;
5772 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
5773 if (Dim == SPIRV::Dim::DIM_Cube) {
5775 "Gather operations with offset are not supported for Cube images.");
5779 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
5781 ImageOperands |= SPIRV::ImageOperand::Offset;
5785 if (ImageOperands != 0) {
5786 MIB.
addImm(ImageOperands);
5788 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
5796bool SPIRVInstructionSelector::generateImageReadOrFetch(
5799 const ImageOperands *ImOps)
const {
5802 "ImageReg is not an image type.");
5804 bool IsSignedInteger =
5809 bool IsFetch = (SampledOp.getImm() == 1);
5811 auto AddOperands = [&](MachineInstrBuilder &MIB) {
5812 uint32_t ImageOperandsMask = 0;
5813 if (IsSignedInteger)
5814 ImageOperandsMask |= 0x1000;
5816 if (IsFetch && ImOps) {
5818 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
5819 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
5821 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
5823 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
5827 if (ImageOperandsMask != 0) {
5828 MIB.
addImm(ImageOperandsMask);
5829 if (IsFetch && ImOps) {
5832 if (ImOps->Offset &&
5833 (ImageOperandsMask &
5834 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
5835 MIB.
addUse(*ImOps->Offset);
5841 if (ResultSize == 4) {
5844 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
5851 BMI.constrainAllUses(
TII,
TRI, RBI);
5855 SPIRVTypeInst ReadType = widenTypeToVec4(ResType, Pos);
5859 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
5865 BMI.constrainAllUses(
TII,
TRI, RBI);
5867 if (ResultSize == 1) {
5876 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
5879bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
5880 SPIRVTypeInst ResType,
5881 MachineInstr &
I)
const {
5882 Register ResourcePtr =
I.getOperand(2).getReg();
5884 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
5893 MachineIRBuilder MIRBuilder(
I);
5895 Register IndexReg =
I.getOperand(3).getReg();
5898 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5908bool SPIRVInstructionSelector::selectPushConstantGetPointer(
5909 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5914bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
5915 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5916 Register ObjReg =
I.getOperand(2).getReg();
5917 if (!BuildCOPY(ResVReg, ObjReg,
I))
5927 decorateUsesAsNonUniform(ResVReg);
5931void SPIRVInstructionSelector::decorateUsesAsNonUniform(
5934 while (WorkList.
size() > 0) {
5938 bool IsDecorated =
false;
5940 if (
Use.getOpcode() == SPIRV::OpDecorate &&
5941 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
5947 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
5949 if (ResultReg == CurrentReg)
5957 SPIRV::Decoration::NonUniformEXT, {});
5962bool SPIRVInstructionSelector::extractSubvector(
5964 MachineInstr &InsertionPoint)
const {
5966 [[maybe_unused]] uint64_t InputSize =
5969 assert(InputSize > 1 &&
"The input must be a vector.");
5970 assert(ResultSize > 1 &&
"The result must be a vector.");
5971 assert(ResultSize < InputSize &&
5972 "Cannot extract more element than there are in the input.");
5975 const TargetRegisterClass *ScalarRegClass = GR.
getRegClass(ScalarType);
5976 for (uint64_t
I = 0;
I < ResultSize;
I++) {
5979 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
5988 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
5990 TII.get(SPIRV::OpCompositeConstruct))
5994 for (
Register ComponentReg : ComponentRegisters)
5995 MIB.
addUse(ComponentReg);
6000bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
6001 MachineInstr &
I)
const {
6008 Register ImageReg =
I.getOperand(1).getReg();
6016 Register CoordinateReg =
I.getOperand(2).getReg();
6017 Register DataReg =
I.getOperand(3).getReg();
6020 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
6028Register SPIRVInstructionSelector::buildPointerToResource(
6029 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
6030 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
6031 StringRef Name, MachineIRBuilder MIRBuilder)
const {
6033 if (ArraySize == 1) {
6034 SPIRVTypeInst PtrType =
6037 "SpirvResType did not have an explicit layout.");
6042 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
6043 SPIRVTypeInst VarPointerType =
6046 VarPointerType, Set,
Binding, Name, MIRBuilder);
6048 SPIRVTypeInst ResPointerType =
6061bool SPIRVInstructionSelector::selectFirstBitSet16(
6062 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
6063 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
6065 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
6069 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
6072bool SPIRVInstructionSelector::selectFirstBitSet32(
6074 unsigned BitSetOpcode)
const {
6075 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6078 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6085bool SPIRVInstructionSelector::selectFirstBitSet64(
6087 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
6100 if (ComponentCount > 2) {
6101 auto Func = [
this, SwapPrimarySide](
Register ResVReg, SPIRVTypeInst ResType,
6103 unsigned Opcode) ->
bool {
6104 return this->selectFirstBitSet64(ResVReg, ResType,
I, SrcReg, Opcode,
6108 return handle64BitOverflow(ResVReg, ResType,
I, SrcReg, BitSetOpcode, Func);
6112 MachineIRBuilder MIRBuilder(
I);
6114 BaseType, 2 * ComponentCount, MIRBuilder,
false);
6118 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
6124 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
6131 bool IsScalarRes = ResType->
getOpcode() != SPIRV::OpTypeVector;
6134 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
6135 SPIRV::OpVectorExtractDynamic))
6137 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
6138 SPIRV::OpVectorExtractDynamic))
6142 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6143 TII.get(SPIRV::OpVectorShuffle))
6151 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
6157 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6158 TII.get(SPIRV::OpVectorShuffle))
6166 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
6186 SelectOp = SPIRV::OpSelectSISCond;
6187 AddOp = SPIRV::OpIAddS;
6195 SelectOp = SPIRV::OpSelectVIVCond;
6196 AddOp = SPIRV::OpIAddV;
6202 Register RegSecondaryOffset = Reg0;
6206 if (SwapPrimarySide) {
6207 PrimaryReg = LowReg;
6208 SecondaryReg = HighReg;
6209 RegPrimaryOffset = Reg0;
6210 RegSecondaryOffset = Reg32;
6215 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
6216 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
6221 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
6222 SPIRV::OpINotEqual))
6229 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
6230 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
6235 if (SwapPrimarySide) {
6237 if (!selectOpWithSrcs(RegAdd, ResType,
I,
6238 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
6249 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
6250 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
6255 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
6256 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
6259 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
6263bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
6264 SPIRVTypeInst ResType,
6266 bool IsSigned)
const {
6268 Register OpReg =
I.getOperand(2).getReg();
6271 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
6272 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
6276 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6278 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6280 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6284 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
6288bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
6289 SPIRVTypeInst ResType,
6290 MachineInstr &
I)
const {
6292 Register OpReg =
I.getOperand(2).getReg();
6297 unsigned ExtendOpcode = SPIRV::OpUConvert;
6298 unsigned BitSetOpcode = GL::FindILsb;
6302 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6304 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6306 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6313bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
6314 SPIRVTypeInst ResType,
6315 MachineInstr &
I)
const {
6319 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
6322 .
addUse(
I.getOperand(2).getReg())
6325 unsigned Alignment =
I.getOperand(3).getImm();
6331bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
6332 SPIRVTypeInst ResType,
6333 MachineInstr &
I)
const {
6337 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
6340 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
6343 unsigned Alignment =
I.getOperand(2).getImm();
6350bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
6355 const MachineInstr *PrevI =
I.getPrevNode();
6357 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
6361 .
addMBB(
I.getOperand(0).getMBB())
6366 .
addMBB(
I.getOperand(0).getMBB())
6371bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
6382 const MachineInstr *NextI =
I.getNextNode();
6384 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
6390 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
6392 .
addUse(
I.getOperand(0).getReg())
6393 .
addMBB(
I.getOperand(1).getMBB())
6399bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
6400 MachineInstr &
I)
const {
6402 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
6404 const unsigned NumOps =
I.getNumOperands();
6405 for (
unsigned i = 1; i <
NumOps; i += 2) {
6406 MIB.
addUse(
I.getOperand(i + 0).getReg());
6407 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
6413bool SPIRVInstructionSelector::selectGlobalValue(
6414 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
6416 MachineIRBuilder MIRBuilder(
I);
6417 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
6420 std::string GlobalIdent;
6422 unsigned &
ID = UnnamedGlobalIDs[GV];
6424 ID = UnnamedGlobalIDs.
size();
6425 GlobalIdent =
"__unnamed_" + Twine(
ID).str();
6451 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
6458 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
6463 MachineInstrBuilder MIB1 =
6464 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
6467 MachineInstrBuilder MIB2 =
6469 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
6473 GR.
add(ConstVal, MIB2);
6481 MachineInstrBuilder MIB3 =
6482 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
6485 GR.
add(ConstVal, MIB3);
6489 assert(NewReg != ResVReg);
6490 return BuildCOPY(ResVReg, NewReg,
I);
6500 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
6503 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
6509 SPIRVTypeInst ResType =
6513 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
6518 if (
GlobalVar->isExternallyInitialized() &&
6519 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
6520 constexpr unsigned ReadWriteINTEL = 3u;
6523 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
6529bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
6530 SPIRVTypeInst ResType,
6531 MachineInstr &
I)
const {
6533 return selectExtInst(ResVReg, ResType,
I, CL::log10);
6541 MachineIRBuilder MIRBuilder(
I);
6546 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6549 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6551 .
add(
I.getOperand(1))
6556 ResType->
getOpcode() == SPIRV::OpTypeFloat);
6566 APFloat::rmNearestTiesToEven, &LosesInfo);
6570 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
6571 ? SPIRV::OpVectorTimesScalar
6582bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
6583 SPIRVTypeInst ResType,
6584 MachineInstr &
I)
const {
6587 return selectExtInst(ResVReg, ResType,
I, CL::pown);
6593 Register ExpReg =
I.getOperand(2).getReg();
6595 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
6596 SPIRV::OpConvertSToF))
6598 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
6605bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
6606 SPIRVTypeInst ResType,
6607 MachineInstr &
I)
const {
6623 MachineIRBuilder MIRBuilder(
I);
6626 ResType, MIRBuilder, SPIRV::StorageClass::Function);
6639 MachineBasicBlock &EntryBB =
I.getMF()->front();
6641 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
6644 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
6650 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6653 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
6656 .
add(
I.getOperand(
I.getNumExplicitDefs()))
6660 Register IntegralPartReg =
I.getOperand(1).getReg();
6661 if (IntegralPartReg.
isValid()) {
6663 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6674 assert(
false &&
"GLSL::Modf is deprecated.");
6685bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
6686 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
6687 SPIRVTypeInst ResType, MachineInstr &
I)
const {
6688 MachineIRBuilder MIRBuilder(
I);
6689 const SPIRVTypeInst Vec3Ty =
6692 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
6704 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
6708 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
6714 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6721 assert(
I.getOperand(2).isReg());
6722 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
6726 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6737bool SPIRVInstructionSelector::loadBuiltinInputID(
6738 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
6739 SPIRVTypeInst ResType, MachineInstr &
I)
const {
6740 MachineIRBuilder MIRBuilder(
I);
6742 ResType, MIRBuilder, SPIRV::StorageClass::Input);
6757 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
6761 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6770SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
6771 MachineInstr &
I)
const {
6772 MachineIRBuilder MIRBuilder(
I);
6773 if (
Type->getOpcode() != SPIRV::OpTypeVector)
6783bool SPIRVInstructionSelector::loadHandleBeforePosition(
6784 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
6785 MachineInstr &Pos)
const {
6788 Intrinsic::spv_resource_handlefrombinding);
6796 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
6797 MachineIRBuilder MIRBuilder(HandleDef);
6798 SPIRVTypeInst VarType = ResType;
6799 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
6801 if (IsStructuredBuffer) {
6806 if (ResType->
getOpcode() == SPIRV::OpTypeImage && ArraySize == 0)
6808 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
6811 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
6812 ArraySize, IndexReg, Name, MIRBuilder);
6816 uint32_t LoadOpcode =
6817 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
6827void SPIRVInstructionSelector::errorIfInstrOutsideShader(
6828 MachineInstr &
I)
const {
6830 std::string DiagMsg;
6831 raw_string_ostream OS(DiagMsg);
6832 I.print(OS,
true,
false,
false,
false);
6833 DiagMsg +=
" is only supported in shaders.\n";
6839InstructionSelector *
6843 return new SPIRVInstructionSelector(TM, Subtarget, RBI);
MachineInstrBuilder & UseMI
#define GET_GLOBALISEL_PREDICATES_INIT
#define GET_GLOBALISEL_TEMPORARIES_INIT
assert(UImm &&(UImm !=~static_cast< T >(0)) &&"Invalid immediate!")
This file declares a class to represent arbitrary precision floating point values and provide a varie...
static bool selectUnmergeValues(MachineInstrBuilder &MIB, const ARMBaseInstrInfo &TII, MachineRegisterInfo &MRI, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI)
MachineBasicBlock MachineBasicBlock::iterator DebugLoc DL
static uint8_t SwapBits(uint8_t Val)
static GCRegistry::Add< CoreCLRGC > E("coreclr", "CoreCLR-compatible GC")
DXIL Resource Implicit Binding
Declares convenience wrapper classes for interpreting MachineInstr instances as specific generic oper...
const HexagonInstrInfo * TII
LLVMTypeRef LLVMIntType(unsigned NumBits)
const size_t AbstractManglingParser< Derived, Alloc >::NumOps
Loop::LoopBounds::Direction Direction
Register const TargetRegisterInfo * TRI
Promote Memory to Register
MachineInstr unsigned OpIdx
uint64_t IntrinsicInst * II
static StringRef getName(Value *V)
static unsigned getFCmpOpcode(CmpInst::Predicate Pred, unsigned Size)
static APFloat getOneFP(const Type *LLVMFloatTy)
static bool isUSMStorageClass(SPIRV::StorageClass::StorageClass SC)
static bool isASCastInGVar(MachineRegisterInfo *MRI, Register ResVReg)
static bool mayApplyGenericSelection(unsigned Opcode)
static APFloat getZeroFP(const Type *LLVMFloatTy)
std::vector< std::pair< SPIRV::InstructionSet::InstructionSet, uint32_t > > ExtInstList
static bool intrinsicHasSideEffects(Intrinsic::ID ID)
static unsigned getBoolCmpOpcode(unsigned PredNum)
static unsigned getICmpOpcode(unsigned PredNum)
static bool isOpcodeWithNoSideEffects(unsigned Opcode)
static void addMemoryOperands(MachineMemOperand *MemOp, MachineInstrBuilder &MIB, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry &GR)
static bool isConstReg(MachineRegisterInfo *MRI, MachineInstr *OpDef)
static unsigned getPtrCmpOpcode(unsigned Pred)
bool isDead(const MachineInstr &MI, const MachineRegisterInfo &MRI)
BaseType
A given derived pointer can have multiple base pointers through phi/selects.
static TableGen::Emitter::Opt Y("gen-skeleton-entry", EmitSkeleton, "Generate example skeleton entry")
static ManagedStatic< cl::opt< FnT >, OptCreatorT > CallbackFunction
static const fltSemantics & IEEEsingle()
static const fltSemantics & IEEEdouble()
static const fltSemantics & IEEEhalf()
const fltSemantics & getSemantics() const
static APFloat getOne(const fltSemantics &Sem, bool Negative=false)
Factory for Positive and Negative One.
static APFloat getZero(const fltSemantics &Sem, bool Negative=false)
Factory for Positive and Negative Zero.
static APInt getAllOnes(unsigned numBits)
Return an APInt of a specified width with all bits set.
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.
Represents a call to an intrinsic.
Intrinsic::ID getIntrinsicID() const
unsigned getAddressSpace() const
Module * getParent()
Get the module that this global value is contained inside of...
@ InternalLinkage
Rename collisions when linking (static functions).
static LLVM_ABI IntegerType * get(LLVMContext &C, unsigned NumBits)
This static method is the primary way of constructing an IntegerType.
static constexpr LLT scalar(unsigned SizeInBits)
Get a low-level scalar or aggregate "bag of bits".
constexpr bool isValid() const
constexpr uint16_t getNumElements() const
Returns the number of elements in a vector LLT.
constexpr bool isVector() const
static constexpr LLT pointer(unsigned AddressSpace, unsigned SizeInBits)
Get a low-level pointer in the given address space.
static constexpr LLT fixed_vector(unsigned NumElements, unsigned ScalarSizeInBits)
Get a low-level fixed-width vector of some number of elements and element width.
int getNumber() const
MachineBasicBlocks are uniquely numbered at the function level, unless they're not in a MachineFuncti...
LLVM_ABI iterator getFirstNonPHI()
Returns a pointer to the first instruction in this block that is not a PHINode instruction.
const MachineFunction * getParent() const
Return the MachineFunction containing this basic block.
MachineInstrBundleIterator< MachineInstr > iterator
MachineRegisterInfo & getRegInfo()
getRegInfo - Return information about the registers currently in use.
Function & getFunction()
Return the LLVM function that this machine code represents.
Helper class to build MachineInstr.
MachineBasicBlock::iterator getInsertPt()
Current insertion point for new instructions.
MachineInstrBuilder buildInstr(unsigned Opcode)
Build and insert <empty> = Opcode <empty>.
MachineFunction & getMF()
Getter for the function we currently build.
MachineRegisterInfo * getMRI()
Getter for MRI.
void constrainAllUses(const TargetInstrInfo &TII, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI) const
const MachineInstrBuilder & addUse(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a virtual register use operand.
const MachineInstrBuilder & addReg(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a new virtual register operand.
const MachineInstrBuilder & addImm(int64_t Val) const
Add a new immediate operand.
const MachineInstrBuilder & add(const MachineOperand &MO) const
const MachineInstrBuilder & addMBB(MachineBasicBlock *MBB, unsigned TargetFlags=0) const
const MachineInstrBuilder & addDef(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a virtual register definition operand.
const MachineInstrBuilder & setMIFlags(unsigned Flags) const
MachineInstr * getInstr() const
If conversion operators fail, use this method to get the MachineInstr explicitly.
Representation of each machine instruction.
unsigned getOpcode() const
Returns the opcode of this MachineInstr.
const MachineBasicBlock * getParent() const
unsigned getNumOperands() const
Retuns the total number of operands.
LLVM_ABI unsigned getNumExplicitOperands() const
Returns the number of non-implicit operands.
LLVM_ABI unsigned getNumExplicitDefs() const
Returns the number of non-implicit definitions.
LLVM_ABI void emitGenericError(const Twine &ErrMsg) const
LLVM_ABI const MachineFunction * getMF() const
Return the function that contains the basic block that this instruction belongs to.
const DebugLoc & getDebugLoc() const
Returns the debug location id of this MachineInstr.
const MachineOperand & getOperand(unsigned i) const
A description of a memory reference used in the backend.
@ MOVolatile
The memory access is volatile.
@ MONonTemporal
The memory access is non-temporal.
bool isReg() const
isReg - Tests if this is a MO_Register operand.
MachineBasicBlock * getMBB() const
Register getReg() const
getReg - Returns the register number.
MachineRegisterInfo - Keep track of information for virtual and physical registers,...
defusechain_instr_iterator< true, false, false, true > use_instr_iterator
use_instr_iterator/use_instr_begin/use_instr_end - Walk all uses of the specified register,...
const TargetRegisterClass * getRegClass(Register Reg) const
Return the register class of the specified virtual register.
LLVM_ABI MachineInstr * getVRegDef(Register Reg) const
getVRegDef - Return the machine instr that defines the specified virtual register or null if none is ...
use_instr_iterator use_instr_begin(Register RegNo) const
static def_instr_iterator def_instr_end()
defusechain_instr_iterator< false, true, false, true > def_instr_iterator
def_instr_iterator/def_instr_begin/def_instr_end - Walk all defs of the specified register,...
LLVM_ABI Register createVirtualRegister(const TargetRegisterClass *RegClass, StringRef Name="")
createVirtualRegister - Create and return a new virtual register in the function with the specified r...
def_instr_iterator def_instr_begin(Register RegNo) const
LLT getType(Register Reg) const
Get the low-level type of Reg or LLT{} if Reg is not a generic (target independent) virtual register.
static use_instr_iterator use_instr_end()
iterator_range< use_instr_nodbg_iterator > use_nodbg_instructions(Register Reg) const
LLVM_ABI void setType(Register VReg, LLT Ty)
Set the low-level type of VReg to Ty.
const MachineFunction & getMF() const
LLVM_ABI void setRegClass(Register Reg, const TargetRegisterClass *RC)
setRegClass - Set the register class of the specified virtual register.
LLVM_ABI Register createGenericVirtualRegister(LLT Ty, StringRef Name="")
Create and return a new generic virtual register with low-level type Ty.
const TargetRegisterClass * getRegClassOrNull(Register Reg) const
Return the register class of Reg, or null if Reg has not been assigned a register class yet.
iterator_range< use_instr_iterator > use_instructions(Register Reg) const
unsigned getNumVirtRegs() const
getNumVirtRegs - Return the number of virtual registers created.
LLVM_ABI void replaceRegWith(Register FromReg, Register ToReg)
replaceRegWith - Replace all instances of FromReg with ToReg in the machine function.
Analysis providing profile information.
Holds all the information related to register banks.
Wrapper class representing virtual and physical registers.
constexpr bool isValid() const
constexpr bool isPhysical() const
Return true if the specified register number is in the physical register namespace.
bool isScalarOrVectorSigned(SPIRVTypeInst Type) const
SPIRVTypeInst getOrCreateOpTypeSampledImage(SPIRVTypeInst ImageType, MachineIRBuilder &MIRBuilder)
void assignSPIRVTypeToVReg(SPIRVTypeInst Type, Register VReg, const MachineFunction &MF)
const TargetRegisterClass * getRegClass(SPIRVTypeInst SpvType) const
MachineInstr * getOrAddMemAliasingINTELInst(MachineIRBuilder &MIRBuilder, const MDNode *AliasingListMD)
bool isAggregateType(SPIRVTypeInst Type) const
unsigned getScalarOrVectorBitWidth(SPIRVTypeInst Type) const
SPIRVTypeInst getOrCreateSPIRVIntegerType(unsigned BitWidth, MachineIRBuilder &MIRBuilder)
SPIRVTypeInst getOrCreateSPIRVVectorType(SPIRVTypeInst BaseType, unsigned NumElements, MachineIRBuilder &MIRBuilder, bool EmitIR)
Register buildGlobalVariable(Register Reg, SPIRVTypeInst BaseType, StringRef Name, const GlobalValue *GV, SPIRV::StorageClass::StorageClass Storage, const MachineInstr *Init, bool IsConst, const std::optional< SPIRV::LinkageType::LinkageType > &LinkageType, MachineIRBuilder &MIRBuilder, bool IsInstSelector)
SPIRVTypeInst getResultType(Register VReg, MachineFunction *MF=nullptr)
unsigned getScalarOrVectorComponentCount(Register VReg) const
const Type * getTypeForSPIRVType(SPIRVTypeInst Ty) const
bool isBitcastCompatible(SPIRVTypeInst Type1, SPIRVTypeInst Type2) const
unsigned getPointerSize() const
Register getOrCreateConstFP(APFloat Val, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
LLT getRegType(SPIRVTypeInst SpvType) const
void invalidateMachineInstr(MachineInstr *MI)
SPIRVTypeInst getOrCreateSPIRVBoolType(MachineIRBuilder &MIRBuilder, bool EmitIR)
SPIRVTypeInst getOrCreateSPIRVPointerType(const Type *BaseType, MachineIRBuilder &MIRBuilder, SPIRV::StorageClass::StorageClass SC)
bool isScalarOfType(Register VReg, unsigned TypeOpcode) const
Register getSPIRVTypeID(SPIRVTypeInst SpirvType) const
Register getOrCreateConstInt(uint64_t Val, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
Register getOrCreateConstIntArray(uint64_t Val, size_t Num, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII)
bool findValueAttrs(const MachineInstr *Key, Type *&Ty, StringRef &Name)
SPIRVTypeInst retrieveScalarOrVectorIntType(SPIRVTypeInst Type) const
Register getOrCreateGlobalVariableWithBinding(SPIRVTypeInst VarType, uint32_t Set, uint32_t Binding, StringRef Name, MachineIRBuilder &MIRBuilder)
SPIRVTypeInst changePointerStorageClass(SPIRVTypeInst PtrType, SPIRV::StorageClass::StorageClass SC, MachineInstr &I)
Register getOrCreateConstVector(uint64_t Val, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
Register buildConstantFP(APFloat Val, MachineIRBuilder &MIRBuilder, SPIRVTypeInst SpvType=nullptr)
void addGlobalObject(const Value *V, const MachineFunction *MF, Register R)
SPIRVTypeInst getScalarOrVectorComponentType(SPIRVTypeInst Type) const
void recordFunctionPointer(const MachineOperand *MO, const Function *F)
SPIRVTypeInst getOrCreateSPIRVFloatType(unsigned BitWidth, MachineInstr &I, const SPIRVInstrInfo &TII)
SPIRVTypeInst getPointeeType(SPIRVTypeInst PtrType)
SPIRVTypeInst getOrCreateSPIRVType(const Type *Type, MachineInstr &I, SPIRV::AccessQualifier::AccessQualifier AQ, bool EmitIR)
bool isScalarOrVectorOfType(Register VReg, unsigned TypeOpcode) const
MachineFunction * setCurrentFunc(MachineFunction &MF)
Register getOrCreateConstNullPtr(MachineIRBuilder &MIRBuilder, SPIRVTypeInst SpvType)
SPIRVTypeInst getSPIRVTypeForVReg(Register VReg, const MachineFunction *MF=nullptr) const
Type * getDeducedGlobalValueType(const GlobalValue *Global)
Register getOrCreateUndef(MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII)
SPIRV::StorageClass::StorageClass getPointerStorageClass(Register VReg) const
bool erase(const MachineInstr *MI)
bool add(SPIRV::IRHandle Handle, const MachineInstr *MI)
Register find(SPIRV::IRHandle Handle, const MachineFunction *MF)
bool isPhysicalSPIRV() const
bool isAtLeastSPIRVVer(VersionTuple VerToCompareTo) const
bool canUseExtInstSet(SPIRV::InstructionSet::InstructionSet E) const
bool isLogicalSPIRV() const
bool canUseExtension(SPIRV::Extension::Extension E) const
bool isTypeIntOrFloat() const
bool erase(PtrType Ptr)
Remove pointer from the set.
std::pair< iterator, bool > insert(PtrType Ptr)
Inserts Ptr if and only if there is no element in the container equal to Ptr.
bool contains(ConstPtrType Ptr) const
SmallPtrSet - This class implements a set which is optimized for holding SmallSize or less elements.
reference emplace_back(ArgTypes &&... Args)
void reserve(size_type N)
void push_back(const T &Elt)
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
constexpr size_t size() const
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.
self_iterator getIterator()
NodeTy * getNextNode()
Get the next node, or nullptr for the list tail.
#define llvm_unreachable(msg)
Marks that the current location is not supposed to be reachable.
constexpr char IsConst[]
Key for Kernel::Arg::Metadata::mIsConst.
constexpr std::underlying_type_t< E > Mask()
Get a bitmask with 1s in all places up to the high-order bit of E's largest value.
unsigned ID
LLVM IR allows to use arbitrary numbers as calling convention identifiers.
Scope
Defines the scope in which this symbol should be visible: Default – Visible in the public interface o...
NodeAddr< DefNode * > Def
NodeAddr< InstrNode * > Instr
NodeAddr< UseNode * > Use
NodeAddr< FuncNode * > Func
BaseReg
Stack frame base register. Bit 0 of FREInfo.Info.
This is an optimization pass for GlobalISel generic memory operations.
void buildOpName(Register Target, const StringRef &Name, MachineIRBuilder &MIRBuilder)
@ Low
Lower the current thread's priority such that it does not affect foreground tasks significantly.
FunctionAddr VTableAddr Value
bool all_of(R &&range, UnaryPredicate P)
Provide wrappers to std::all_of which take ranges instead of having to pass begin/end explicitly.
MachineBasicBlock::iterator getOpVariableMBBIt(MachineFunction &MF)
int64_t getIConstValSext(Register ConstReg, const MachineRegisterInfo *MRI)
MachineInstrBuilder BuildMI(MachineFunction &MF, const MIMetadata &MIMD, const MCInstrDesc &MCID)
Builder interface. Specify how to create the initial instruction itself.
bool isTypeFoldingSupported(unsigned Opcode)
decltype(auto) dyn_cast(const From &Val)
dyn_cast<X> - Return the argument parameter cast to the specified type.
FunctionAddr VTableAddr uintptr_t uintptr_t Int32Ty
void addNumImm(const APInt &Imm, MachineInstrBuilder &MIB)
LLVM_ABI void salvageDebugInfo(const MachineRegisterInfo &MRI, MachineInstr &MI)
Assuming the instruction MI is going to be deleted, attempt to salvage debug users of MI by writing t...
LLVM_ABI void constrainSelectedInstRegOperands(MachineInstr &I, const TargetInstrInfo &TII, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI)
Mutate the newly-selected instruction I to constrain its (possibly generic) virtual register operands...
bool isPreISelGenericOpcode(unsigned Opcode)
Check whether the given Opcode is a generic opcode that is not supposed to appear after ISel.
Register createVirtualRegister(SPIRVTypeInst SpvType, SPIRVGlobalRegistry *GR, MachineRegisterInfo *MRI, const MachineFunction &MF)
unsigned getArrayComponentCount(const MachineRegisterInfo *MRI, const MachineInstr *ResType)
uint64_t getIConstVal(Register ConstReg, const MachineRegisterInfo *MRI)
SmallVector< MachineInstr *, 4 > createContinuedInstructions(MachineIRBuilder &MIRBuilder, unsigned Opcode, unsigned MinWC, unsigned ContinuedOpcode, ArrayRef< Register > Args, Register ReturnRegister, Register TypeID)
SPIRV::MemorySemantics::MemorySemantics getMemSemanticsForStorageClass(SPIRV::StorageClass::StorageClass SC)
constexpr unsigned storageClassToAddressSpace(SPIRV::StorageClass::StorageClass SC)
bool any_of(R &&range, UnaryPredicate P)
Provide wrappers to std::any_of which take ranges instead of having to pass begin/end explicitly.
void buildOpDecorate(Register Reg, MachineIRBuilder &MIRBuilder, SPIRV::Decoration::Decoration Dec, const std::vector< uint32_t > &DecArgs, StringRef StrImm)
MachineInstr * getImm(const MachineOperand &MO, const MachineRegisterInfo *MRI)
constexpr bool isPowerOf2_32(uint32_t Value)
Return true if the argument is a power of two > 0.
Type * toTypedPointer(Type *Ty)
LLVM_ABI raw_ostream & dbgs()
dbgs() - This returns a reference to a raw_ostream for debugging messages.
LLVM_ABI void report_fatal_error(Error Err, bool gen_crash_diag=true)
constexpr bool isGenericCastablePtr(SPIRV::StorageClass::StorageClass SC)
class LLVM_GSL_OWNER SmallVector
Forward declaration of SmallVector so that calculateSmallVectorDefaultInlinedElements can reference s...
MachineInstr * passCopy(MachineInstr *Def, const MachineRegisterInfo *MRI)
bool isa(const From &Val)
isa<X> - Return true if the parameter to the template is an instance of one of the template type argu...
std::optional< SPIRV::LinkageType::LinkageType > getSpirvLinkageTypeFor(const SPIRVSubtarget &ST, const GlobalValue &GV)
LLVM_ABI raw_fd_ostream & errs()
This returns a reference to a raw_ostream for standard error.
SPIRV::StorageClass::StorageClass addressSpaceToStorageClass(unsigned AddrSpace, const SPIRVSubtarget &STI)
AtomicOrdering
Atomic ordering for LLVM's memory model.
SPIRV::Scope::Scope getMemScope(LLVMContext &Ctx, SyncScope::ID Id)
InstructionSelector * createSPIRVInstructionSelector(const SPIRVTargetMachine &TM, const SPIRVSubtarget &Subtarget, const RegisterBankInfo &RBI)
std::string getStringValueFromReg(Register Reg, MachineRegisterInfo &MRI)
int64_t foldImm(const MachineOperand &MO, const MachineRegisterInfo *MRI)
DWARFExpression::Operation Op
ArrayRef(const T &OneElt) -> ArrayRef< T >
MachineInstr * getDefInstrMaybeConstant(Register &ConstReg, const MachineRegisterInfo *MRI)
constexpr unsigned BitWidth
decltype(auto) cast(const From &Val)
cast<X> - Return the argument parameter cast to the specified type.
bool hasInitializer(const GlobalVariable *GV)
bool isSpvIntrinsic(const MachineInstr &MI, Intrinsic::ID IntrinsicID)
void addStringImm(const StringRef &Str, MCInst &Inst)
MachineInstr * getVRegDef(MachineRegisterInfo &MRI, Register Reg)
SPIRV::MemorySemantics::MemorySemantics getMemSemantics(AtomicOrdering Ord)
std::string getLinkStringForBuiltIn(SPIRV::BuiltIn::BuiltIn BuiltInValue)
LLVM_ABI bool isTriviallyDead(const MachineInstr &MI, const MachineRegisterInfo &MRI)
Check whether an instruction MI is dead: it only defines dead virtual registers, and doesn't have oth...