34#include "llvm/IR/IntrinsicsSPIRV.h"
39#define DEBUG_TYPE "spirv-isel"
46 std::vector<std::pair<SPIRV::InstructionSet::InstructionSet, uint32_t>>;
51 std::optional<Register> Bias;
52 std::optional<Register>
Offset;
53 std::optional<Register> MinLod;
54 std::optional<Register> GradX;
55 std::optional<Register> GradY;
56 std::optional<Register> Lod;
57 std::optional<Register> Compare;
64 bool IsScalar =
false;
67llvm::SPIRV::SelectionControl::SelectionControl
68getSelectionOperandForImm(
int Imm) {
70 return SPIRV::SelectionControl::Flatten;
72 return SPIRV::SelectionControl::DontFlatten;
74 return SPIRV::SelectionControl::None;
78#define GET_GLOBALISEL_PREDICATE_BITSET
79#include "SPIRVGenGlobalISel.inc"
80#undef GET_GLOBALISEL_PREDICATE_BITSET
107#define GET_GLOBALISEL_PREDICATES_DECL
108#include "SPIRVGenGlobalISel.inc"
109#undef GET_GLOBALISEL_PREDICATES_DECL
111#define GET_GLOBALISEL_TEMPORARIES_DECL
112#include "SPIRVGenGlobalISel.inc"
113#undef GET_GLOBALISEL_TEMPORARIES_DECL
137 unsigned BitSetOpcode)
const;
141 unsigned BitSetOpcode)
const;
145 unsigned BitSetOpcode,
bool SwapPrimarySide)
const;
149 unsigned BitSetOpcode,
150 bool SwapPrimarySide)
const;
157 unsigned Opcode)
const;
160 unsigned Opcode)
const;
179 unsigned NewOpcode,
unsigned NegateOpcode = 0)
const;
190 unsigned OpType)
const;
250 unsigned Opcode)
const;
254 unsigned Opcode)
const;
258 unsigned Opcode)
const;
262 unsigned Opcode)
const;
264 template <
bool Signed>
267 template <
bool Signed>
274 template <
typename PickOpcodeFn>
277 PickOpcodeFn &&PickOpcode)
const;
294 template <
typename PickOpcodeFn>
297 PickOpcodeFn &&PickOpcode)
const;
315 bool IsSigned)
const;
317 bool IsSigned,
unsigned Opcode)
const;
319 bool IsSigned)
const;
325 bool IsSigned)
const;
364 GL::GLSLExtInst GLInst,
bool setMIFlags =
true,
365 bool useMISrc =
true,
367 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
368 CL::OpenCLExtInst CLInst,
bool setMIFlags =
true,
369 bool useMISrc =
true,
371 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
372 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
373 bool setMIFlags =
true,
bool useMISrc =
true,
375 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
376 const ExtInstList &ExtInsts,
bool setMIFlags =
true,
377 bool useMISrc =
true,
380 bool selectLog10(
Register ResVReg, SPIRVTypeInst ResType,
381 MachineInstr &
I)
const;
383 bool selectFpowi(
Register ResVReg, SPIRVTypeInst ResType,
384 MachineInstr &
I)
const;
386 bool selectSaturate(
Register ResVReg, SPIRVTypeInst ResType,
387 MachineInstr &
I)
const;
389 bool selectWaveOpInst(
Register ResVReg, SPIRVTypeInst ResType,
390 MachineInstr &
I,
unsigned Opcode)
const;
392 bool selectBarrierInst(MachineInstr &
I,
unsigned Scope,
unsigned MemSem,
393 bool WithGroupSync)
const;
395 bool selectWaveActiveCountBits(
Register ResVReg, SPIRVTypeInst ResType,
396 MachineInstr &
I)
const;
398 bool selectWaveActiveAllEqual(
Register ResVReg, SPIRVTypeInst ResType,
399 MachineInstr &
I)
const;
403 bool selectHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
404 MachineInstr &
I)
const;
406 bool selectCounterHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
407 MachineInstr &
I)
const;
409 bool selectReadImageIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
410 MachineInstr &
I)
const;
411 bool selectGetDimensionsIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
412 MachineInstr &
I)
const;
413 bool selectGetDimensionsLevelsIntrinsic(
Register &ResVReg,
414 SPIRVTypeInst ResType,
415 MachineInstr &
I)
const;
416 bool selectGetDimensionsMSIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
417 MachineInstr &
I)
const;
420 std::optional<Register> LodReg = std::nullopt)
const;
421 bool selectSampleBasicIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
422 MachineInstr &
I)
const;
423 bool selectCalculateLodIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
424 MachineInstr &
I)
const;
425 bool selectSampleBiasIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
426 MachineInstr &
I)
const;
427 bool selectSampleGradIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
428 MachineInstr &
I)
const;
429 bool selectSampleLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
430 MachineInstr &
I)
const;
431 bool selectLoadLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
432 MachineInstr &
I)
const;
433 bool selectSampleCmpIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
434 MachineInstr &
I)
const;
435 bool selectSampleCmpLevelZeroIntrinsic(
Register &ResVReg,
436 SPIRVTypeInst ResType,
437 MachineInstr &
I)
const;
438 bool selectGatherIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
439 MachineInstr &
I)
const;
440 bool selectImageWriteIntrinsic(MachineInstr &
I)
const;
441 bool selectResourceGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
442 MachineInstr &
I)
const;
443 bool selectPushConstantGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
444 MachineInstr &
I)
const;
445 bool selectResourceNonUniformIndex(
Register &ResVReg, SPIRVTypeInst ResType,
446 MachineInstr &
I)
const;
447 bool selectModf(
Register ResVReg, SPIRVTypeInst ResType,
448 MachineInstr &
I)
const;
449 bool selectUpdateCounter(
Register &ResVReg, SPIRVTypeInst ResType,
450 MachineInstr &
I)
const;
451 bool selectFrexp(
Register ResVReg, SPIRVTypeInst ResType,
452 MachineInstr &
I)
const;
453 bool selectSincos(
Register ResVReg, SPIRVTypeInst ResType,
454 MachineInstr &
I)
const;
455 bool selectExp10(
Register ResVReg, SPIRVTypeInst ResType,
456 MachineInstr &
I)
const;
457 bool selectDerivativeInst(
Register ResVReg, SPIRVTypeInst ResType,
458 MachineInstr &
I,
const unsigned DPdOpCode)
const;
460 Register buildI32Constant(uint32_t Val, MachineInstr &
I,
461 SPIRVTypeInst ResType =
nullptr)
const;
463 Register buildZerosVal(SPIRVTypeInst ResType, MachineInstr &
I)
const;
464 bool isScalarOrVectorIntConstantZero(
Register Reg)
const;
465 Register buildZerosValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
467 MachineInstr &
I)
const;
468 Register buildOnesValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
470 bool wrapIntoSpecConstantOp(MachineInstr &
I,
473 Register getUcharPtrTypeReg(MachineInstr &
I,
474 SPIRV::StorageClass::StorageClass SC)
const;
475 MachineInstrBuilder buildSpecConstantOp(MachineInstr &
I,
Register Dest,
477 uint32_t Opcode)
const;
478 MachineInstrBuilder buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
479 SPIRVTypeInst SrcPtrTy)
const;
480 Register buildPointerToResource(SPIRVTypeInst ResType,
481 SPIRV::StorageClass::StorageClass SC,
482 uint32_t Set, uint32_t
Binding,
483 uint32_t ArraySize,
Register IndexReg,
485 MachineIRBuilder MIRBuilder)
const;
486 SPIRVTypeInst widenTypeToVec4(SPIRVTypeInst
Type, MachineInstr &
I)
const;
487 bool extractSubvector(
Register &ResVReg, SPIRVTypeInst ResType,
488 Register &ReadReg, MachineInstr &InsertionPoint)
const;
489 bool generateImageReadOrFetch(
Register &ResVReg, SPIRVTypeInst ResType,
492 const ImageOperands *ImOps =
nullptr)
const;
493 bool generateSampleImage(
Register ResVReg, SPIRVTypeInst ResType,
495 Register CoordinateReg,
const ImageOperands &ImOps,
498 bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
499 Register ResVReg, SPIRVTypeInst ResType,
500 MachineInstr &
I)
const;
501 bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
502 Register ResVReg, SPIRVTypeInst ResType,
503 MachineInstr &
I)
const;
504 bool loadHandleBeforePosition(
Register &HandleReg, SPIRVTypeInst ResType,
505 GIntrinsic &HandleDef, MachineInstr &Pos)
const;
506 void decorateUsesAsNonUniform(
Register &NonUniformReg)
const;
507 void errorIfInstrOutsideShader(MachineInstr &
I)
const;
509 std::optional<SplitParts> splitEvenOddLanes(
Register PopCountReg,
510 unsigned ComponentCount,
512 SPIRVTypeInst I32Type)
const;
515bool sampledTypeIsSignedInteger(
const llvm::Type *HandleType) {
517 if (
TET->getTargetExtName() ==
"spirv.Image") {
520 assert(
TET->getTargetExtName() ==
"spirv.SignedImage");
521 return TET->getTypeParameter(0)->isIntegerTy();
525#define GET_GLOBALISEL_IMPL
526#include "SPIRVGenGlobalISel.inc"
527#undef GET_GLOBALISEL_IMPL
533 TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
536#include
"SPIRVGenGlobalISel.inc"
539#include
"SPIRVGenGlobalISel.inc"
551 InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
555void SPIRVInstructionSelector::resetVRegsType(MachineFunction &MF) {
556 if (HasVRegsReset == &MF)
571 for (
const auto &
MBB : MF) {
572 for (
const auto &
MI :
MBB) {
575 if (
MI.getOpcode() != SPIRV::ASSIGN_TYPE)
579 LLT DstType = MRI.
getType(DstReg);
581 LLT SrcType = MRI.
getType(SrcReg);
582 if (DstType != SrcType)
587 if (DstRC != SrcRC && SrcRC)
599 while (!Stack.empty()) {
604 switch (
MI->getOpcode()) {
605 case TargetOpcode::G_INTRINSIC:
606 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
607 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS: {
610 if (IntrID != Intrinsic::spv_const_composite &&
611 IntrID != Intrinsic::spv_undef)
615 case TargetOpcode::G_BUILD_VECTOR:
616 case TargetOpcode::G_SPLAT_VECTOR:
618 i < OpDef->getNumOperands(); i++) {
623 Stack.push_back(OpNestedDef);
626 case TargetOpcode::G_CONSTANT:
627 case TargetOpcode::G_FCONSTANT:
628 case TargetOpcode::G_IMPLICIT_DEF:
629 case SPIRV::OpConstantTrue:
630 case SPIRV::OpConstantFalse:
631 case SPIRV::OpConstantI:
632 case SPIRV::OpConstantF:
633 case SPIRV::OpConstantComposite:
634 case SPIRV::OpConstantCompositeContinuedINTEL:
635 case SPIRV::OpConstantSampler:
636 case SPIRV::OpConstantNull:
638 case SPIRV::OpConstantFunctionPointerINTEL:
665 case Intrinsic::spv_all:
666 case Intrinsic::spv_alloca:
667 case Intrinsic::spv_any:
668 case Intrinsic::spv_bitcast:
669 case Intrinsic::spv_const_composite:
670 case Intrinsic::spv_cross:
671 case Intrinsic::spv_degrees:
672 case Intrinsic::spv_distance:
673 case Intrinsic::spv_extractelt:
674 case Intrinsic::spv_extractv:
675 case Intrinsic::spv_faceforward:
676 case Intrinsic::spv_fdot:
677 case Intrinsic::spv_firstbitlow:
678 case Intrinsic::spv_firstbitshigh:
679 case Intrinsic::spv_firstbituhigh:
680 case Intrinsic::spv_frac:
681 case Intrinsic::spv_gep:
682 case Intrinsic::spv_global_offset:
683 case Intrinsic::spv_global_size:
684 case Intrinsic::spv_group_id:
685 case Intrinsic::spv_insertelt:
686 case Intrinsic::spv_insertv:
687 case Intrinsic::spv_isinf:
688 case Intrinsic::spv_isnan:
689 case Intrinsic::spv_lerp:
690 case Intrinsic::spv_length:
691 case Intrinsic::spv_normalize:
692 case Intrinsic::spv_num_subgroups:
693 case Intrinsic::spv_num_workgroups:
694 case Intrinsic::spv_ptrcast:
695 case Intrinsic::spv_radians:
696 case Intrinsic::spv_reflect:
697 case Intrinsic::spv_refract:
698 case Intrinsic::spv_resource_getpointer:
699 case Intrinsic::spv_resource_handlefrombinding:
700 case Intrinsic::spv_resource_handlefromimplicitbinding:
701 case Intrinsic::spv_resource_nonuniformindex:
702 case Intrinsic::spv_resource_sample:
703 case Intrinsic::spv_rsqrt:
704 case Intrinsic::spv_saturate:
705 case Intrinsic::spv_sdot:
706 case Intrinsic::spv_sign:
707 case Intrinsic::spv_smoothstep:
708 case Intrinsic::spv_step:
709 case Intrinsic::spv_subgroup_id:
710 case Intrinsic::spv_subgroup_local_invocation_id:
711 case Intrinsic::spv_subgroup_max_size:
712 case Intrinsic::spv_subgroup_size:
713 case Intrinsic::spv_thread_id:
714 case Intrinsic::spv_thread_id_in_group:
715 case Intrinsic::spv_udot:
716 case Intrinsic::spv_undef:
717 case Intrinsic::spv_value_md:
718 case Intrinsic::spv_workgroup_size:
730 case SPIRV::OpTypeVoid:
731 case SPIRV::OpTypeBool:
732 case SPIRV::OpTypeInt:
733 case SPIRV::OpTypeFloat:
734 case SPIRV::OpTypeVector:
735 case SPIRV::OpTypeMatrix:
736 case SPIRV::OpTypeImage:
737 case SPIRV::OpTypeSampler:
738 case SPIRV::OpTypeSampledImage:
739 case SPIRV::OpTypeArray:
740 case SPIRV::OpTypeRuntimeArray:
741 case SPIRV::OpTypeStruct:
742 case SPIRV::OpTypeOpaque:
743 case SPIRV::OpTypePointer:
744 case SPIRV::OpTypeFunction:
745 case SPIRV::OpTypeEvent:
746 case SPIRV::OpTypeDeviceEvent:
747 case SPIRV::OpTypeReserveId:
748 case SPIRV::OpTypeQueue:
749 case SPIRV::OpTypePipe:
750 case SPIRV::OpTypeForwardPointer:
751 case SPIRV::OpTypePipeStorage:
752 case SPIRV::OpTypeNamedBarrier:
753 case SPIRV::OpTypeAccelerationStructureNV:
754 case SPIRV::OpTypeCooperativeMatrixNV:
755 case SPIRV::OpTypeCooperativeMatrixKHR:
765 if (
MI.getNumDefs() == 0)
768 for (
const auto &MO :
MI.all_defs()) {
770 if (
Reg.isPhysical()) {
775 if (
UseMI.getOpcode() != SPIRV::OpName) {
782 if (
MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE ||
MI.isFakeUse() ||
783 MI.isLifetimeMarker()) {
786 <<
"Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
797 if (
MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
798 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
801 LLVM_DEBUG(
dbgs() <<
"Dead: Intrinsic with no real side effects.\n");
806 if (
MI.mayStore() ||
MI.isCall() ||
807 (
MI.mayLoad() &&
MI.hasOrderedMemoryRef()) ||
MI.isPosition() ||
808 MI.isDebugInstr() ||
MI.isTerminator() ||
MI.isJumpTableDebugInfo()) {
809 LLVM_DEBUG(
dbgs() <<
"Not dead: instruction has side effects.\n");
820 LLVM_DEBUG(
dbgs() <<
"Dead: known opcode with no side effects\n");
827void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &
MI)
const {
829 for (
const auto &MO :
MI.all_defs()) {
833 SmallVector<MachineInstr *, 4> UselessOpNames;
836 "There is still a use of the dead function.");
839 for (MachineInstr *OpNameMI : UselessOpNames) {
841 OpNameMI->eraseFromParent();
846void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &
MI)
const {
849 removeOpNamesForDeadMI(
MI);
850 MI.eraseFromParent();
853bool SPIRVInstructionSelector::select(MachineInstr &
I) {
854 resetVRegsType(*
I.getParent()->getParent());
856 assert(
I.getParent() &&
"Instruction should be in a basic block!");
857 assert(
I.getParent()->getParent() &&
"Instruction should be in a function!");
862 removeDeadInstruction(
I);
869 if (Opcode == SPIRV::ASSIGN_TYPE) {
870 Register DstReg =
I.getOperand(0).getReg();
871 Register SrcReg =
I.getOperand(1).getReg();
874 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
875 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
876 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
877 Register SelectDstReg =
Def->getOperand(0).getReg();
878 bool SuccessToSelectSelect [[maybe_unused]] = selectSelect(
880 assert(SuccessToSelectSelect);
882 Def->eraseFromParent();
889 bool Res = selectImpl(
I, *CoverageInfo);
891 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
892 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
896 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
908 }
else if (
I.getNumDefs() == 1) {
920 removeDeadInstruction(
I);
925 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
926 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
932 bool HasDefs =
I.getNumDefs() > 0;
935 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
936 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
937 if (spvSelect(ResVReg, ResType,
I)) {
939 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
950 case TargetOpcode::G_CONSTANT:
951 case TargetOpcode::G_FCONSTANT:
958 MachineInstr &
I)
const {
961 if (DstRC != SrcRC && SrcRC)
963 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
970bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
971 SPIRVTypeInst ResType,
972 MachineInstr &
I)
const {
973 const unsigned Opcode =
I.getOpcode();
975 return selectImpl(
I, *CoverageInfo);
977 case TargetOpcode::G_CONSTANT:
978 case TargetOpcode::G_FCONSTANT:
979 return selectConst(ResVReg, ResType,
I);
980 case TargetOpcode::G_GLOBAL_VALUE:
981 return selectGlobalValue(ResVReg,
I);
982 case TargetOpcode::G_IMPLICIT_DEF:
983 return selectOpUndef(ResVReg, ResType,
I);
984 case TargetOpcode::G_FREEZE:
985 return selectFreeze(ResVReg, ResType,
I);
987 case TargetOpcode::G_INTRINSIC:
988 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
989 case TargetOpcode::G_INTRINSIC_CONVERGENT:
990 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
991 return selectIntrinsic(ResVReg, ResType,
I);
992 case TargetOpcode::G_BITREVERSE:
993 return selectBitreverse(ResVReg, ResType,
I);
995 case TargetOpcode::G_BUILD_VECTOR:
996 return selectBuildVector(ResVReg, ResType,
I);
997 case TargetOpcode::G_SPLAT_VECTOR:
998 return selectSplatVector(ResVReg, ResType,
I);
1000 case TargetOpcode::G_SHUFFLE_VECTOR: {
1001 MachineBasicBlock &BB = *
I.getParent();
1002 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
1005 .
addUse(
I.getOperand(1).getReg())
1006 .
addUse(
I.getOperand(2).getReg());
1007 for (
auto V :
I.getOperand(3).getShuffleMask())
1012 case TargetOpcode::G_MEMMOVE:
1013 case TargetOpcode::G_MEMCPY:
1014 case TargetOpcode::G_MEMSET:
1015 return selectMemOperation(ResVReg,
I);
1017 case TargetOpcode::G_ICMP:
1018 return selectICmp(ResVReg, ResType,
I);
1019 case TargetOpcode::G_FCMP:
1020 return selectFCmp(ResVReg, ResType,
I);
1022 case TargetOpcode::G_FRAME_INDEX:
1023 return selectFrameIndex(ResVReg, ResType,
I);
1025 case TargetOpcode::G_LOAD:
1026 return selectLoad(ResVReg, ResType,
I);
1027 case TargetOpcode::G_STORE:
1028 return selectStore(
I);
1030 case TargetOpcode::G_BR:
1031 return selectBranch(
I);
1032 case TargetOpcode::G_BRCOND:
1033 return selectBranchCond(
I);
1035 case TargetOpcode::G_PHI:
1036 return selectPhi(ResVReg,
I);
1038 case TargetOpcode::G_FPTOSI:
1039 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1040 case TargetOpcode::G_FPTOUI:
1041 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1043 case TargetOpcode::G_FPTOSI_SAT:
1044 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1045 case TargetOpcode::G_FPTOUI_SAT:
1046 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1048 case TargetOpcode::G_SITOFP:
1049 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
1050 case TargetOpcode::G_UITOFP:
1051 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
1053 case TargetOpcode::G_CTPOP:
1054 return selectPopCount(ResVReg, ResType,
I, SPIRV::OpBitCount);
1055 case TargetOpcode::G_SMIN:
1056 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
1057 case TargetOpcode::G_UMIN:
1058 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
1060 case TargetOpcode::G_SMAX:
1061 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
1062 case TargetOpcode::G_UMAX:
1063 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
1065 case TargetOpcode::G_SCMP:
1066 return selectSUCmp(ResVReg, ResType,
I,
true);
1067 case TargetOpcode::G_UCMP:
1068 return selectSUCmp(ResVReg, ResType,
I,
false);
1069 case TargetOpcode::G_LROUND:
1070 case TargetOpcode::G_LLROUND: {
1073 MRI->
setRegClass(regForLround, &SPIRV::iIDRegClass);
1075 regForLround, *(
I.getParent()->getParent()));
1077 CL::round, GL::Round,
false);
1079 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1086 case TargetOpcode::G_STRICT_FMA:
1087 case TargetOpcode::G_FMA: {
1090 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1093 .
addUse(
I.getOperand(1).getReg())
1094 .
addUse(
I.getOperand(2).getReg())
1095 .
addUse(
I.getOperand(3).getReg())
1100 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1103 case TargetOpcode::G_STRICT_FLDEXP:
1104 return selectExtInst(ResVReg, ResType,
I, CL::ldexp);
1106 case TargetOpcode::G_FPOW:
1107 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1108 case TargetOpcode::G_FPOWI:
1109 return selectFpowi(ResVReg, ResType,
I);
1111 case TargetOpcode::G_FEXP:
1112 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1113 case TargetOpcode::G_FEXP2:
1114 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1115 case TargetOpcode::G_FEXP10:
1116 return selectExp10(ResVReg, ResType,
I);
1118 case TargetOpcode::G_FMODF:
1119 return selectModf(ResVReg, ResType,
I);
1120 case TargetOpcode::G_FSINCOS:
1121 return selectSincos(ResVReg, ResType,
I);
1123 case TargetOpcode::G_FLOG:
1124 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1125 case TargetOpcode::G_FLOG2:
1126 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1127 case TargetOpcode::G_FLOG10:
1128 return selectLog10(ResVReg, ResType,
I);
1130 case TargetOpcode::G_FABS:
1131 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1132 case TargetOpcode::G_ABS:
1133 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1135 case TargetOpcode::G_FMINNUM:
1136 case TargetOpcode::G_FMINIMUM:
1137 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1138 case TargetOpcode::G_FMAXNUM:
1139 case TargetOpcode::G_FMAXIMUM:
1140 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1142 case TargetOpcode::G_FCOPYSIGN:
1143 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1145 case TargetOpcode::G_FCEIL:
1146 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1147 case TargetOpcode::G_FFLOOR:
1148 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1150 case TargetOpcode::G_FCOS:
1151 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1152 case TargetOpcode::G_FSIN:
1153 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1154 case TargetOpcode::G_FTAN:
1155 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1156 case TargetOpcode::G_FACOS:
1157 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1158 case TargetOpcode::G_FASIN:
1159 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1160 case TargetOpcode::G_FATAN:
1161 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1162 case TargetOpcode::G_FATAN2:
1163 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1164 case TargetOpcode::G_FCOSH:
1165 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1166 case TargetOpcode::G_FSINH:
1167 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1168 case TargetOpcode::G_FTANH:
1169 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1171 case TargetOpcode::G_STRICT_FSQRT:
1172 case TargetOpcode::G_FSQRT:
1173 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1175 case TargetOpcode::G_CTTZ:
1176 case TargetOpcode::G_CTTZ_ZERO_UNDEF:
1177 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1178 case TargetOpcode::G_CTLZ:
1179 case TargetOpcode::G_CTLZ_ZERO_UNDEF:
1180 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1182 case TargetOpcode::G_INTRINSIC_ROUND:
1183 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1184 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1185 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1186 case TargetOpcode::G_INTRINSIC_TRUNC:
1187 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1188 case TargetOpcode::G_FRINT:
1189 case TargetOpcode::G_FNEARBYINT:
1190 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1192 case TargetOpcode::G_SMULH:
1193 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1194 case TargetOpcode::G_UMULH:
1195 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1197 case TargetOpcode::G_SADDSAT:
1198 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1199 case TargetOpcode::G_UADDSAT:
1200 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1201 case TargetOpcode::G_SSUBSAT:
1202 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1203 case TargetOpcode::G_USUBSAT:
1204 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1206 case TargetOpcode::G_FFREXP:
1207 return selectFrexp(ResVReg, ResType,
I);
1209 case TargetOpcode::G_UADDO:
1210 return selectOverflowArith(ResVReg, ResType,
I,
1211 ResType->
getOpcode() == SPIRV::OpTypeVector
1212 ? SPIRV::OpIAddCarryV
1213 : SPIRV::OpIAddCarryS);
1214 case TargetOpcode::G_USUBO:
1215 return selectOverflowArith(ResVReg, ResType,
I,
1216 ResType->
getOpcode() == SPIRV::OpTypeVector
1217 ? SPIRV::OpISubBorrowV
1218 : SPIRV::OpISubBorrowS);
1219 case TargetOpcode::G_UMULO:
1220 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1221 case TargetOpcode::G_SMULO:
1222 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1224 case TargetOpcode::G_SEXT:
1225 return selectExt(ResVReg, ResType,
I,
true);
1226 case TargetOpcode::G_ANYEXT:
1227 case TargetOpcode::G_ZEXT:
1228 return selectExt(ResVReg, ResType,
I,
false);
1229 case TargetOpcode::G_TRUNC:
1230 return selectTrunc(ResVReg, ResType,
I);
1231 case TargetOpcode::G_FPTRUNC:
1232 case TargetOpcode::G_FPEXT:
1233 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1235 case TargetOpcode::G_PTRTOINT:
1236 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1237 case TargetOpcode::G_INTTOPTR:
1238 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1239 case TargetOpcode::G_BITCAST:
1240 return selectBitcast(ResVReg, ResType,
I);
1241 case TargetOpcode::G_ADDRSPACE_CAST:
1242 return selectAddrSpaceCast(ResVReg, ResType,
I);
1243 case TargetOpcode::G_PTR_ADD: {
1245 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1249 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1250 (*II).getOpcode() == TargetOpcode::COPY ||
1251 (*II).getOpcode() == SPIRV::OpVariable) &&
1252 getImm(
I.getOperand(2), MRI));
1254 bool IsGVInit =
false;
1258 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1259 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1260 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1261 (*UseIt).getOpcode() == SPIRV::OpVariable) {
1271 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1284 "incompatible result and operand types in a bitcast");
1286 MachineInstrBuilder MIB =
1287 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1294 : SPIRV::OpInBoundsPtrAccessChain))
1298 .
addUse(
I.getOperand(2).getReg())
1301 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1305 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1307 .
addUse(
I.getOperand(2).getReg())
1316 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1319 .
addImm(
static_cast<uint32_t
>(
1320 SPIRV::Opcode::InBoundsPtrAccessChain))
1323 .
addUse(
I.getOperand(2).getReg());
1328 case TargetOpcode::G_ATOMICRMW_OR:
1329 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1330 case TargetOpcode::G_ATOMICRMW_ADD:
1331 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1332 case TargetOpcode::G_ATOMICRMW_AND:
1333 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1334 case TargetOpcode::G_ATOMICRMW_MAX:
1335 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1336 case TargetOpcode::G_ATOMICRMW_MIN:
1337 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1338 case TargetOpcode::G_ATOMICRMW_SUB:
1339 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1340 case TargetOpcode::G_ATOMICRMW_XOR:
1341 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1342 case TargetOpcode::G_ATOMICRMW_UMAX:
1343 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1344 case TargetOpcode::G_ATOMICRMW_UMIN:
1345 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1346 case TargetOpcode::G_ATOMICRMW_XCHG:
1347 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1348 case TargetOpcode::G_ATOMIC_CMPXCHG:
1349 return selectAtomicCmpXchg(ResVReg, ResType,
I);
1351 case TargetOpcode::G_ATOMICRMW_FADD:
1352 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1353 case TargetOpcode::G_ATOMICRMW_FSUB:
1355 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1356 ResType->
getOpcode() == SPIRV::OpTypeVector
1358 : SPIRV::OpFNegate);
1359 case TargetOpcode::G_ATOMICRMW_FMIN:
1360 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1361 case TargetOpcode::G_ATOMICRMW_FMAX:
1362 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1364 case TargetOpcode::G_FENCE:
1365 return selectFence(
I);
1367 case TargetOpcode::G_STACKSAVE:
1368 return selectStackSave(ResVReg, ResType,
I);
1369 case TargetOpcode::G_STACKRESTORE:
1370 return selectStackRestore(
I);
1372 case TargetOpcode::G_UNMERGE_VALUES:
1378 case TargetOpcode::G_TRAP:
1379 case TargetOpcode::G_UBSANTRAP:
1380 case TargetOpcode::DBG_LABEL:
1382 case TargetOpcode::G_DEBUGTRAP:
1383 return selectDebugTrap(ResVReg, ResType,
I);
1390bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1391 SPIRVTypeInst ResType,
1392 MachineInstr &
I)
const {
1393 unsigned Opcode = SPIRV::OpNop;
1400bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1401 SPIRVTypeInst ResType,
1403 GL::GLSLExtInst GLInst,
1404 bool setMIFlags,
bool useMISrc,
1407 SPIRV::InstructionSet::InstructionSet::GLSL_std_450)) {
1408 std::string DiagMsg;
1409 raw_string_ostream OS(DiagMsg);
1410 I.print(OS,
true,
false,
false,
false);
1411 DiagMsg +=
" is only supported with the GLSL extended instruction set.\n";
1414 return selectExtInst(ResVReg, ResType,
I,
1415 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1416 setMIFlags, useMISrc, SrcRegs);
1419bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1420 SPIRVTypeInst ResType,
1422 CL::OpenCLExtInst CLInst,
1423 bool setMIFlags,
bool useMISrc,
1425 return selectExtInst(ResVReg, ResType,
I,
1426 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1427 setMIFlags, useMISrc, SrcRegs);
1430bool SPIRVInstructionSelector::selectExtInst(
1431 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1432 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1434 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1435 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1436 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1440bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1441 SPIRVTypeInst ResType,
1444 bool setMIFlags,
bool useMISrc,
1447 for (
const auto &[InstructionSet, Opcode] : Insts) {
1451 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1454 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1459 const unsigned NumOps =
I.getNumOperands();
1462 I.getOperand(Index).getType() ==
1463 MachineOperand::MachineOperandType::MO_IntrinsicID)
1466 MIB.
add(
I.getOperand(Index));
1478bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1479 SPIRVTypeInst ResType,
1480 MachineInstr &
I)
const {
1481 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1482 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1483 for (
const auto &Ex : ExtInsts) {
1484 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1485 uint32_t Opcode = Ex.second;
1489 MachineIRBuilder MIRBuilder(
I);
1492 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1497 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1500 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1503 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1506 .
addImm(
static_cast<uint32_t
>(Ex.first))
1508 .
add(
I.getOperand(2))
1512 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1513 .
addDef(
I.getOperand(1).getReg())
1522bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1523 SPIRVTypeInst ResType,
1524 MachineInstr &
I)
const {
1525 Register CosResVReg =
I.getOperand(1).getReg();
1526 unsigned SrcIdx =
I.getNumExplicitDefs();
1531 MachineIRBuilder MIRBuilder(
I);
1533 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1538 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1541 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1543 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1546 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1548 .
add(
I.getOperand(SrcIdx))
1551 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1559 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1562 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1564 .
add(
I.getOperand(SrcIdx))
1566 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1569 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1571 .
add(
I.getOperand(SrcIdx))
1578bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1579 SPIRVTypeInst ResType,
1581 std::vector<Register> Srcs,
1582 unsigned Opcode)
const {
1583 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1593std::optional<SplitParts> SPIRVInstructionSelector::splitEvenOddLanes(
1594 Register PopCountReg,
unsigned ComponentCount, MachineInstr &
I,
1595 SPIRVTypeInst I32Type)
const {
1598 if (ComponentCount == 1) {
1601 Parts.IsScalar =
true;
1602 Parts.Type = I32Type;
1610 if (!selectOpWithSrcs(Parts.High, I32Type,
I, {PopCountReg, IdxOne},
1611 SPIRV::OpVectorExtractDynamic))
1612 return std::nullopt;
1614 if (!selectOpWithSrcs(Parts.Low, I32Type,
I, {PopCountReg, IdxZero},
1615 SPIRV::OpVectorExtractDynamic))
1616 return std::nullopt;
1620 MachineIRBuilder MIRBuilder(
I);
1621 Parts.IsScalar =
false;
1628 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1629 TII.get(SPIRV::OpVectorShuffle))
1634 for (
unsigned J = 1; J < ComponentCount * 2; J += 2)
1639 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1640 TII.get(SPIRV::OpVectorShuffle))
1645 for (
unsigned J = 0; J < ComponentCount * 2; J += 2)
1653bool SPIRVInstructionSelector::selectPopCount16(
Register ResVReg,
1654 SPIRVTypeInst ResType,
1657 unsigned Opcode)
const {
1658 Register OpReg =
I.getOperand(1).getReg();
1661 MachineIRBuilder MIRBuilder(
I);
1663 SPIRVTypeInst I32VectorType =
1666 bool IsVector = NumElems > 1;
1667 SPIRVTypeInst ExtType = IsVector ? I32VectorType : I32Type;
1670 if (!selectOpWithSrcs(ExtReg, ExtType,
I, {OpReg}, SPIRV::OpUConvert))
1674 if (!selectPopCount32(PopCountReg, ExtType,
I, ExtReg, Opcode))
1677 return selectOpWithSrcs(ResVReg, ResType,
I, {PopCountReg}, ExtOpcode);
1680bool SPIRVInstructionSelector::selectPopCount32(
Register ResVReg,
1681 SPIRVTypeInst ResType,
1684 unsigned Opcode)
const {
1685 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
1688bool SPIRVInstructionSelector::selectPopCount64Overflow(
1690 unsigned int Opcode)
const {
1693 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
1695 MachineIRBuilder MIRBuilder(
I);
1698 SPIRVTypeInst I64x2Type =
1700 SPIRVTypeInst Vec2ResType =
1703 std::vector<Register> PartialRegs;
1705 unsigned CurrentComponent = 0;
1706 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
1710 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1711 TII.get(SPIRV::OpVectorShuffle))
1716 .
addImm(CurrentComponent)
1717 .
addImm(CurrentComponent + 1);
1724 if (!selectPopCount64(SubVecReg, Vec2ResType,
I, PopCountResult, Opcode))
1727 PartialRegs.push_back(SubVecReg);
1731 if (CurrentComponent != ComponentCount) {
1737 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
1738 SPIRV::OpVectorExtractDynamic))
1744 if (!selectPopCount64(FinalElemResReg,
BaseType,
I, FinalElemReg, Opcode))
1747 PartialRegs.push_back(FinalElemResReg);
1752 return selectOpWithSrcs(ResVReg, ResType,
I, std::move(PartialRegs),
1753 SPIRV::OpCompositeConstruct);
1756bool SPIRVInstructionSelector::selectPopCount64(
Register ResVReg,
1757 SPIRVTypeInst ResType,
1760 unsigned Opcode)
const {
1762 if (ComponentCount > 2)
1763 return selectPopCount64Overflow(ResVReg, ResType,
I, SrcReg, Opcode);
1765 MachineIRBuilder MIRBuilder(
I);
1770 I32Type, 2 * ComponentCount, MIRBuilder,
false);
1774 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
1779 if (!selectPopCount32(Pop32, VecI32Type,
I, Vec32, Opcode))
1783 auto MaybeParts = splitEvenOddLanes(Pop32, ComponentCount,
I, I32Type);
1786 SplitParts &Parts = *MaybeParts;
1789 unsigned OpAdd = Parts.IsScalar ? SPIRV::OpIAddS : SPIRV::OpIAddV;
1791 if (!selectOpWithSrcs(Sum, Parts.Type,
I, {Parts.High, Parts.Low}, OpAdd))
1796 unsigned ConvOp = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
1797 return selectOpWithSrcs(ResVReg, ResType,
I, {Sum}, ConvOp);
1800bool SPIRVInstructionSelector::selectPopCount(
Register ResVReg,
1801 SPIRVTypeInst ResType,
1803 unsigned Opcode)
const {
1808 if (!STI.getTargetTriple().isVulkanOS())
1809 return selectUnOp(ResVReg, ResType,
I, Opcode);
1811 Register OpReg =
I.getOperand(1).getReg();
1814 : SPIRV::OpUConvert;
1818 return selectPopCount16(ResVReg, ResType,
I, ExtOpcode, Opcode);
1820 return selectPopCount32(ResVReg, ResType,
I, OpReg, Opcode);
1822 return selectPopCount64(ResVReg, ResType,
I, OpReg, Opcode);
1828bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1829 SPIRVTypeInst ResType,
1831 unsigned Opcode)
const {
1833 Register SrcReg =
I.getOperand(1).getReg();
1838 unsigned DefOpCode = DefIt->getOpcode();
1839 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1842 if (
auto *VRD =
getVRegDef(*MRI, DefIt->getOperand(1).getReg()))
1843 DefOpCode = VRD->getOpcode();
1845 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1846 DefOpCode == TargetOpcode::G_CONSTANT ||
1847 DefOpCode == SPIRV::OpVariable || DefOpCode == SPIRV::OpConstantI) {
1853 uint32_t SpecOpcode = 0;
1855 case SPIRV::OpConvertPtrToU:
1856 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1858 case SPIRV::OpConvertUToPtr:
1859 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1864 TII.get(SPIRV::OpSpecConstantOp))
1874 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1878bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1879 SPIRVTypeInst ResType,
1880 MachineInstr &
I)
const {
1881 Register OpReg =
I.getOperand(1).getReg();
1882 SPIRVTypeInst OpType =
1886 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1896 if (
MemOp->isVolatile())
1897 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1898 if (
MemOp->isNonTemporal())
1899 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1901 if (!ST->isShader() &&
MemOp->getAlign().value())
1902 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1906 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1907 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1911 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1913 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1917 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1921 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1923 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1935 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1937 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1939 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
1943bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
1944 SPIRVTypeInst ResType,
1945 MachineInstr &
I)
const {
1947 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1952 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
1953 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1955 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1959 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1963 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1964 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
1965 I.getDebugLoc(),
I);
1969 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1973 if (!
I.getNumMemOperands()) {
1974 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1976 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1979 MachineIRBuilder MIRBuilder(
I);
1986bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
1988 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
1989 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1994 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer) {
1995 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2000 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2004 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2005 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2006 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2007 TII.get(SPIRV::OpImageWrite))
2013 if (sampledTypeIsSignedInteger(LLVMHandleType))
2016 BMI.constrainAllUses(
TII,
TRI, RBI);
2022 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpStore))
2025 if (!
I.getNumMemOperands()) {
2026 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2028 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2031 MachineIRBuilder MIRBuilder(
I);
2038bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
2039 SPIRVTypeInst ResType,
2040 MachineInstr &
I)
const {
2041 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
2049 const Register PtrsReg =
I.getOperand(2).getReg();
2050 const uint32_t Alignment =
I.getOperand(3).getImm();
2051 const Register MaskReg =
I.getOperand(4).getReg();
2052 const Register PassthruReg =
I.getOperand(5).getReg();
2053 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2057 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
2068bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
2069 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
2076 const Register ValuesReg =
I.getOperand(1).getReg();
2077 const Register PtrsReg =
I.getOperand(2).getReg();
2078 const uint32_t Alignment =
I.getOperand(3).getImm();
2079 const Register MaskReg =
I.getOperand(4).getReg();
2080 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2084 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
2093bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
2094 const Twine &Msg)
const {
2095 const Function &
F =
I.getMF()->getFunction();
2096 F.getContext().diagnose(
2097 DiagnosticInfoUnsupported(
F, Msg,
I.getDebugLoc(),
DS_Error));
2101bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
2102 SPIRVTypeInst ResType,
2103 MachineInstr &
I)
const {
2104 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2106 "llvm.stacksave intrinsic: this instruction requires the following "
2107 "SPIR-V extension: SPV_INTEL_variable_length_array",
2110 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
2117bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
2118 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2120 "llvm.stackrestore intrinsic: this instruction requires the following "
2121 "SPIR-V extension: SPV_INTEL_variable_length_array",
2123 if (!
I.getOperand(0).isReg())
2126 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
2127 .
addUse(
I.getOperand(0).getReg())
2133SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
2134 MachineIRBuilder MIRBuilder(
I);
2135 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
2142 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
2146 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2147 Type *ArrTy = ArrayType::get(ValTy, Num);
2149 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
2152 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
2159 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
2162 .
addImm(SPIRV::StorageClass::UniformConstant)
2173bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
2176 Register DstReg =
I.getOperand(0).getReg();
2181 uint64_t CopySize =
getIConstVal(
I.getOperand(2).getReg(), MRI);
2186 "Unable to determine pointee type size for OpCopyMemory");
2187 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
2188 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
2190 "OpCopyMemory requires the size to match the pointee type size");
2191 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemory))
2194 if (
I.getNumMemOperands()) {
2195 MachineIRBuilder MIRBuilder(
I);
2202bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &
I,
2205 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemorySized))
2206 .
addUse(
I.getOperand(0).getReg())
2208 .
addUse(
I.getOperand(2).getReg());
2209 if (
I.getNumMemOperands()) {
2210 MachineIRBuilder MIRBuilder(
I);
2217bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
2218 MachineInstr &
I)
const {
2219 Register SrcReg =
I.getOperand(1).getReg();
2220 if (
I.getOpcode() == TargetOpcode::G_MEMSET) {
2221 Register VarReg = getOrCreateMemSetGlobal(
I);
2224 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2226 ValTy,
I, SPIRV::StorageClass::UniformConstant);
2228 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
2232 if (!selectCopyMemory(
I, SrcReg))
2235 if (!selectCopyMemorySized(
I, SrcReg))
2238 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
2239 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
2244bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
2245 SPIRVTypeInst ResType,
2248 unsigned NegateOpcode)
const {
2250 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2253 Register ScopeReg = buildI32Constant(Scope,
I);
2255 Register Ptr =
I.getOperand(1).getReg();
2261 Register MemSemReg = buildI32Constant(MemSem ,
I);
2263 Register ValueReg =
I.getOperand(2).getReg();
2264 if (NegateOpcode != 0) {
2267 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
2272 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2283bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2284 unsigned ArgI =
I.getNumOperands() - 1;
2286 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2287 SPIRVTypeInst SrcType =
2289 if (!SrcType || SrcType->
getOpcode() != SPIRV::OpTypeVector)
2291 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2293 SPIRVTypeInst ScalarType =
2296 unsigned CurrentIndex = 0;
2297 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2298 Register ResVReg =
I.getOperand(i).getReg();
2301 LLT ResLLT = MRI->
getType(ResVReg);
2307 ResType = ScalarType;
2313 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
2316 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2322 for (
unsigned j = 0;
j < NumElements; ++
j) {
2323 MIB.
addImm(CurrentIndex + j);
2325 CurrentIndex += NumElements;
2329 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2341bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2344 Register MemSemReg = buildI32Constant(MemSem,
I);
2346 uint32_t
Scope =
static_cast<uint32_t
>(
2348 Register ScopeReg = buildI32Constant(Scope,
I);
2350 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2357bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2358 SPIRVTypeInst ResType,
2360 unsigned Opcode)
const {
2361 Type *ResTy =
nullptr;
2365 "Not enough info to select the arithmetic with overflow instruction");
2368 "with overflow instruction");
2374 MachineIRBuilder MIRBuilder(
I);
2376 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2377 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2383 Register ZeroReg = buildZerosVal(ResType,
I);
2388 if (ResName.
size() > 0)
2393 BuildMI(BB, MIRBuilder.getInsertPt(),
I.getDebugLoc(),
TII.get(Opcode))
2396 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2397 MIB.
addUse(
I.getOperand(i).getReg());
2402 MRI->
setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2403 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2405 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2406 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2413 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2414 .
addDef(
I.getOperand(1).getReg())
2422bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2423 SPIRVTypeInst ResType,
2424 MachineInstr &
I)
const {
2428 Register Ptr =
I.getOperand(2).getReg();
2431 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2434 ScopeReg = buildI32Constant(Scope,
I);
2436 unsigned ScSem =
static_cast<uint32_t
>(
2439 unsigned MemSemEq =
static_cast<uint32_t
>(
getMemSemantics(AO)) | ScSem;
2440 Register MemSemEqReg = buildI32Constant(MemSemEq,
I);
2442 unsigned MemSemNeq =
static_cast<uint32_t
>(
getMemSemantics(FO)) | ScSem;
2443 if (MemSemEq == MemSemNeq)
2444 MemSemNeqReg = MemSemEqReg;
2446 MemSemNeqReg = buildI32Constant(MemSemEq,
I);
2449 ScopeReg =
I.getOperand(5).getReg();
2450 MemSemEqReg =
I.getOperand(6).getReg();
2451 MemSemNeqReg =
I.getOperand(7).getReg();
2455 Register Val =
I.getOperand(4).getReg();
2459 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2478 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2485 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2497 case SPIRV::StorageClass::DeviceOnlyINTEL:
2498 case SPIRV::StorageClass::HostOnlyINTEL:
2507 bool IsGRef =
false;
2508 bool IsAllowedRefs =
2510 unsigned Opcode = It.getOpcode();
2511 if (Opcode == SPIRV::OpConstantComposite ||
2512 Opcode == SPIRV::OpSpecConstantComposite ||
2513 Opcode == SPIRV::OpVariable ||
2514 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2515 return IsGRef = true;
2516 return Opcode == SPIRV::OpName;
2518 return IsAllowedRefs && IsGRef;
2521Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2522 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2524 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2528SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2530 uint32_t Opcode)
const {
2531 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2532 TII.get(SPIRV::OpSpecConstantOp))
2540SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2541 SPIRVTypeInst SrcPtrTy)
const {
2542 SPIRVTypeInst GenericPtrTy =
2546 SPIRV::StorageClass::Generic),
2548 MachineFunction *MF =
I.getParent()->getParent();
2550 MachineInstrBuilder MIB = buildSpecConstantOp(
2552 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2562bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2563 SPIRVTypeInst ResType,
2564 MachineInstr &
I)
const {
2568 Register SrcPtr =
I.getOperand(1).getReg();
2572 if (SrcPtrTy->
getOpcode() != SPIRV::OpTypePointer ||
2573 ResType->
getOpcode() != SPIRV::OpTypePointer)
2574 return BuildCOPY(ResVReg, SrcPtr,
I);
2584 unsigned SpecOpcode =
2586 ?
static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric)
2589 ? static_cast<uint32_t>(
SPIRV::Opcode::GenericCastToPtr)
2596 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2598 .constrainAllUses(
TII,
TRI, RBI);
2600 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2602 buildSpecConstantOp(
2604 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2605 .constrainAllUses(
TII,
TRI, RBI);
2612 return BuildCOPY(ResVReg, SrcPtr,
I);
2614 if ((SrcSC == SPIRV::StorageClass::Function &&
2615 DstSC == SPIRV::StorageClass::Private) ||
2616 (DstSC == SPIRV::StorageClass::Function &&
2617 SrcSC == SPIRV::StorageClass::Private))
2618 return BuildCOPY(ResVReg, SrcPtr,
I);
2622 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2625 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2628 SPIRVTypeInst GenericPtrTy =
2647 return selectUnOp(ResVReg, ResType,
I,
2648 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
2650 return selectUnOp(ResVReg, ResType,
I,
2651 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
2653 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2655 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2665 return SPIRV::OpFOrdEqual;
2667 return SPIRV::OpFOrdGreaterThanEqual;
2669 return SPIRV::OpFOrdGreaterThan;
2671 return SPIRV::OpFOrdLessThanEqual;
2673 return SPIRV::OpFOrdLessThan;
2675 return SPIRV::OpFOrdNotEqual;
2677 return SPIRV::OpOrdered;
2679 return SPIRV::OpFUnordEqual;
2681 return SPIRV::OpFUnordGreaterThanEqual;
2683 return SPIRV::OpFUnordGreaterThan;
2685 return SPIRV::OpFUnordLessThanEqual;
2687 return SPIRV::OpFUnordLessThan;
2689 return SPIRV::OpFUnordNotEqual;
2691 return SPIRV::OpUnordered;
2701 return SPIRV::OpIEqual;
2703 return SPIRV::OpINotEqual;
2705 return SPIRV::OpSGreaterThanEqual;
2707 return SPIRV::OpSGreaterThan;
2709 return SPIRV::OpSLessThanEqual;
2711 return SPIRV::OpSLessThan;
2713 return SPIRV::OpUGreaterThanEqual;
2715 return SPIRV::OpUGreaterThan;
2717 return SPIRV::OpULessThanEqual;
2719 return SPIRV::OpULessThan;
2728 return SPIRV::OpPtrEqual;
2730 return SPIRV::OpPtrNotEqual;
2741 return SPIRV::OpLogicalEqual;
2743 return SPIRV::OpLogicalNotEqual;
2777bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
2778 SPIRVTypeInst ResType,
2780 unsigned OpAnyOrAll)
const {
2781 assert(
I.getNumOperands() == 3);
2782 assert(
I.getOperand(2).isReg());
2784 Register InputRegister =
I.getOperand(2).getReg();
2787 assert(InputType &&
"VReg has no type assigned");
2790 bool IsVectorTy = InputType->
getOpcode() == SPIRV::OpTypeVector;
2791 if (IsBoolTy && !IsVectorTy) {
2792 assert(ResVReg ==
I.getOperand(0).getReg());
2793 return BuildCOPY(ResVReg, InputRegister,
I);
2797 unsigned SpirvNotEqualId =
2798 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
2800 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
2805 IsBoolTy ? InputRegister
2813 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
2815 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
2832bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
2833 SPIRVTypeInst ResType,
2834 MachineInstr &
I)
const {
2835 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
2838bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
2839 SPIRVTypeInst ResType,
2840 MachineInstr &
I)
const {
2841 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
2845bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
2846 SPIRVTypeInst ResType,
2847 MachineInstr &
I)
const {
2848 assert(
I.getNumOperands() == 4);
2849 assert(
I.getOperand(2).isReg());
2850 assert(
I.getOperand(3).isReg());
2852 [[maybe_unused]] SPIRVTypeInst VecType =
2857 "dot product requires a vector of at least 2 components");
2859 [[maybe_unused]] SPIRVTypeInst EltType =
2868 .
addUse(
I.getOperand(2).getReg())
2869 .
addUse(
I.getOperand(3).getReg())
2874bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
2875 SPIRVTypeInst ResType,
2878 assert(
I.getNumOperands() == 4);
2879 assert(
I.getOperand(2).isReg());
2880 assert(
I.getOperand(3).isReg());
2883 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
2887 .
addUse(
I.getOperand(2).getReg())
2888 .
addUse(
I.getOperand(3).getReg())
2895bool SPIRVInstructionSelector::selectIntegerDotExpansion(
2896 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2897 assert(
I.getNumOperands() == 4);
2898 assert(
I.getOperand(2).isReg());
2899 assert(
I.getOperand(3).isReg());
2903 Register Vec0 =
I.getOperand(2).getReg();
2904 Register Vec1 =
I.getOperand(3).getReg();
2908 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
2917 "dot product requires a vector of at least 2 components");
2920 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2930 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2941 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2953bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
2954 SPIRVTypeInst ResType,
2955 MachineInstr &
I)
const {
2957 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
2960 .
addUse(
I.getOperand(2).getReg())
2965bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
2966 SPIRVTypeInst ResType,
2967 MachineInstr &
I)
const {
2969 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
2972 .
addUse(
I.getOperand(2).getReg())
2977template <
bool Signed>
2978bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
2979 SPIRVTypeInst ResType,
2980 MachineInstr &
I)
const {
2981 assert(
I.getNumOperands() == 5);
2982 assert(
I.getOperand(2).isReg());
2983 assert(
I.getOperand(3).isReg());
2984 assert(
I.getOperand(4).isReg());
2987 Register Acc =
I.getOperand(2).getReg();
2991 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
2993 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
2998 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
3001 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3013template <
bool Signed>
3014bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
3015 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3016 assert(
I.getNumOperands() == 5);
3017 assert(
I.getOperand(2).isReg());
3018 assert(
I.getOperand(3).isReg());
3019 assert(
I.getOperand(4).isReg());
3022 Register Acc =
I.getOperand(2).getReg();
3028 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
3032 for (
unsigned i = 0; i < 4; i++) {
3055 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
3075 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3090bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
3091 SPIRVTypeInst ResType,
3092 MachineInstr &
I)
const {
3093 assert(
I.getNumOperands() == 3);
3094 assert(
I.getOperand(2).isReg());
3096 Register VZero = buildZerosValF(ResType,
I);
3097 Register VOne = buildOnesValF(ResType,
I);
3099 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
3102 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3104 .
addUse(
I.getOperand(2).getReg())
3111bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
3112 SPIRVTypeInst ResType,
3113 MachineInstr &
I)
const {
3114 assert(
I.getNumOperands() == 3);
3115 assert(
I.getOperand(2).isReg());
3117 Register InputRegister =
I.getOperand(2).getReg();
3119 auto &
DL =
I.getDebugLoc();
3129 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
3131 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
3139 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3144 if (NeedsConversion) {
3145 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
3156bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
3157 SPIRVTypeInst ResType,
3159 unsigned Opcode)
const {
3163 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
3169 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
3170 BMI.addUse(
I.getOperand(J).getReg());
3177bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
3180 bool WithGroupSync)
const {
3182 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
3184 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
3186 assert(((Scope != SPIRV::Scope::Workgroup) ||
3187 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
3188 "Workgroup Scope must set WorkGroupMemory semantic "
3189 "in Barrier instruction");
3191 assert(((Scope != SPIRV::Scope::Device) ||
3192 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
3193 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
3194 "Device Scope must set UniformMemory and ImageMemory semantic "
3195 "in Barrier instruction");
3197 Register MemSemReg = buildI32Constant(MemSem,
I);
3198 Register ScopeReg = buildI32Constant(Scope,
I);
3204 if (WithGroupSync) {
3205 MI.addUse(ScopeReg);
3208 MI.addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
3212bool SPIRVInstructionSelector::selectWaveActiveCountBits(
3213 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3218 if (!selectWaveOpInst(BallotReg, BallotType,
I,
3219 SPIRV::OpGroupNonUniformBallot))
3224 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3229 .
addImm(SPIRV::GroupOperation::Reduce)
3238 if (
Type->getOpcode() != SPIRV::OpTypeVector)
3242 return Type->getOperand(2).getImm();
3245bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
3246 SPIRVTypeInst ResType,
3247 MachineInstr &
I)
const {
3252 Register InputReg =
I.getOperand(2).getReg();
3257 bool IsVector = NumElems > 1;
3260 SPIRVTypeInst ElemInputType = InputType;
3261 SPIRVTypeInst ElemBoolType = ResType;
3274 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
3275 SPIRV::OpGroupNonUniformAllEqual);
3280 ElementResults.
reserve(NumElems);
3282 for (
unsigned Idx = 0; Idx < NumElems; ++Idx) {
3295 ElemInput = Extracted;
3301 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3312 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3323bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3324 SPIRVTypeInst ResType,
3325 MachineInstr &
I)
const {
3327 assert(
I.getNumOperands() == 3);
3329 auto Op =
I.getOperand(2);
3341 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3363 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3367 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3374bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3375 SPIRVTypeInst ResType,
3377 bool IsUnsigned)
const {
3378 return selectWaveReduce(
3379 ResVReg, ResType,
I, IsUnsigned,
3380 [&](
Register InputRegister,
bool IsUnsigned) {
3381 const bool IsFloatTy =
3383 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3384 : SPIRV::OpGroupNonUniformSMax;
3385 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3389bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3390 SPIRVTypeInst ResType,
3392 bool IsUnsigned)
const {
3393 return selectWaveReduce(
3394 ResVReg, ResType,
I, IsUnsigned,
3395 [&](
Register InputRegister,
bool IsUnsigned) {
3396 const bool IsFloatTy =
3398 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3399 : SPIRV::OpGroupNonUniformSMin;
3400 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3404bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3405 SPIRVTypeInst ResType,
3406 MachineInstr &
I)
const {
3407 return selectWaveReduce(ResVReg, ResType,
I,
false,
3408 [&](
Register InputRegister,
bool IsUnsigned) {
3410 InputRegister, SPIRV::OpTypeFloat);
3411 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3412 : SPIRV::OpGroupNonUniformIAdd;
3416bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3417 SPIRVTypeInst ResType,
3418 MachineInstr &
I)
const {
3419 return selectWaveReduce(ResVReg, ResType,
I,
false,
3420 [&](
Register InputRegister,
bool IsUnsigned) {
3422 InputRegister, SPIRV::OpTypeFloat);
3423 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3424 : SPIRV::OpGroupNonUniformIMul;
3428template <
typename PickOpcodeFn>
3429bool SPIRVInstructionSelector::selectWaveReduce(
3430 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3431 PickOpcodeFn &&PickOpcode)
const {
3432 assert(
I.getNumOperands() == 3);
3433 assert(
I.getOperand(2).isReg());
3435 Register InputRegister =
I.getOperand(2).getReg();
3442 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3448 .
addImm(SPIRV::GroupOperation::Reduce)
3449 .
addUse(
I.getOperand(2).getReg())
3454bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3455 SPIRVTypeInst ResType,
3457 unsigned Opcode)
const {
3458 return selectWaveReduce(
3459 ResVReg, ResType,
I,
false,
3460 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3463bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3464 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3465 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3466 [&](
Register InputRegister,
bool IsUnsigned) {
3468 InputRegister, SPIRV::OpTypeFloat);
3470 ? SPIRV::OpGroupNonUniformFAdd
3471 : SPIRV::OpGroupNonUniformIAdd;
3475bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3476 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3477 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3478 [&](
Register InputRegister,
bool IsUnsigned) {
3480 InputRegister, SPIRV::OpTypeFloat);
3482 ? SPIRV::OpGroupNonUniformFMul
3483 : SPIRV::OpGroupNonUniformIMul;
3487template <
typename PickOpcodeFn>
3488bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3489 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3490 PickOpcodeFn &&PickOpcode)
const {
3491 assert(
I.getNumOperands() == 3);
3492 assert(
I.getOperand(2).isReg());
3494 Register InputRegister =
I.getOperand(2).getReg();
3501 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3507 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3508 .
addUse(
I.getOperand(2).getReg())
3513bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
3514 SPIRVTypeInst ResType,
3517 assert(
I.getNumOperands() == 3);
3518 assert(
I.getOperand(2).isReg());
3520 Register InputRegister =
I.getOperand(2).getReg();
3526 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
3537bool SPIRVInstructionSelector::selectBitreverse16(
Register ResVReg,
3538 SPIRVTypeInst ResType,
3543 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
3548 : SPIRV::OpUConvert;
3552 ShiftOp = SPIRV::OpShiftRightLogicalV;
3557 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3558 TII.get(SPIRV::OpConstantComposite))
3561 for (
unsigned It = 0; It <
N; ++It)
3565 ShiftConst = CompositeReg;
3570 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
3575 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
3580 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
3585 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
3588bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
3589 SPIRVTypeInst ResType,
3593 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
3601bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
3602 SPIRVTypeInst ResType,
3603 MachineInstr &
I)
const {
3604 Register OpReg =
I.getOperand(1).getReg();
3611 return selectBitreverse16(ResVReg, ResType,
I, OpReg);
3613 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
3618 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
3629 unsigned AndOp = SPIRV::OpBitwiseAndS;
3630 unsigned OrOp = SPIRV::OpBitwiseOrS;
3631 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
3632 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
3634 AndOp = SPIRV::OpBitwiseAndV;
3635 OrOp = SPIRV::OpBitwiseOrV;
3636 ShlOp = SPIRV::OpShiftLeftLogicalV;
3637 ShrOp = SPIRV::OpShiftRightLogicalV;
3643 const unsigned Shift) ->
Register {
3651 Register MaskReg = CreateConst(Mask);
3652 Register ShiftReg = CreateConst(Shift);
3659 if (!selectOpWithSrcs(
T1, ResType,
I, {Input, ShiftReg}, ShrOp) ||
3660 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
3661 !selectOpWithSrcs(T3, ResType,
I, {Input, MaskReg}, AndOp) ||
3662 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
3663 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
3671 uint64_t
Mask = ~0ull;
3672 while ((Shift >>= 1) > 0) {
3679 return BuildCOPY(ResVReg, Result,
I);
3682bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
3683 SPIRVTypeInst ResType,
3684 MachineInstr &
I)
const {
3690 if (!
I.getOperand(0).isReg() || !
I.getOperand(1).isReg())
3692 Register OpReg =
I.getOperand(1).getReg();
3693 if (MachineInstr *Def = MRI->
getVRegDef(OpReg)) {
3694 if (
Def->getOpcode() == TargetOpcode::COPY)
3697 switch (
Def->getOpcode()) {
3698 case SPIRV::ASSIGN_TYPE:
3699 if (MachineInstr *AssignToDef =
3701 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
3702 Reg =
Def->getOperand(2).getReg();
3705 case SPIRV::OpUndef:
3706 Reg =
Def->getOperand(1).getReg();
3709 unsigned DestOpCode;
3711 DestOpCode = SPIRV::OpConstantNull;
3713 DestOpCode = TargetOpcode::COPY;
3716 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
3717 .
addDef(
I.getOperand(0).getReg())
3725bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
3726 SPIRVTypeInst ResType,
3727 MachineInstr &
I)
const {
3729 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3731 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
3735 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
3740 for (
unsigned i =
I.getNumExplicitDefs();
3741 i <
I.getNumExplicitOperands() && IsConst; ++i)
3745 if (!IsConst &&
N < 2)
3747 "There must be at least two constituent operands in a vector");
3750 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3751 TII.get(IsConst ? SPIRV::OpConstantComposite
3752 : SPIRV::OpCompositeConstruct))
3755 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
3756 MIB.
addUse(
I.getOperand(i).getReg());
3761bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
3762 SPIRVTypeInst ResType,
3763 MachineInstr &
I)
const {
3765 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3767 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
3773 if (!
I.getOperand(
OpIdx).isReg())
3780 if (!IsConst &&
N < 2)
3782 "There must be at least two constituent operands in a vector");
3785 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3786 TII.get(IsConst ? SPIRV::OpConstantComposite
3787 : SPIRV::OpCompositeConstruct))
3790 for (
unsigned i = 0; i <
N; ++i)
3796bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
3797 SPIRVTypeInst ResType,
3798 MachineInstr &
I)
const {
3803 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
3805 Opcode = SPIRV::OpDemoteToHelperInvocation;
3807 Opcode = SPIRV::OpKill;
3809 if (MachineInstr *NextI =
I.getNextNode()) {
3811 NextI->eraseFromParent();
3821bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
3822 SPIRVTypeInst ResType,
unsigned CmpOpc,
3823 MachineInstr &
I)
const {
3824 Register Cmp0 =
I.getOperand(2).getReg();
3825 Register Cmp1 =
I.getOperand(3).getReg();
3828 "CMP operands should have the same type");
3829 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
3839bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
3840 SPIRVTypeInst ResType,
3841 MachineInstr &
I)
const {
3842 auto Pred =
I.getOperand(1).getPredicate();
3845 Register CmpOperand =
I.getOperand(2).getReg();
3852 return selectCmp(ResVReg, ResType, CmpOpc,
I);
3856SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
3857 SPIRVTypeInst ResType)
const {
3859 SPIRVTypeInst SpvI32Ty =
3862 auto ConstInt = ConstantInt::get(LLVMTy, Val);
3869 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
3872 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
3875 .
addImm(APInt(32, Val).getZExtValue());
3877 GR.
add(ConstInt,
MI);
3882bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
3883 SPIRVTypeInst ResType,
3884 MachineInstr &
I)
const {
3886 return selectCmp(ResVReg, ResType, CmpOp,
I);
3889bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
3890 SPIRVTypeInst ResType,
3891 MachineInstr &
I)
const {
3893 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
3900 if (ResType->
getOpcode() != SPIRV::OpTypeVector &&
3901 ResType->
getOpcode() != SPIRV::OpTypeFloat)
3904 MachineIRBuilder MIRBuilder(
I);
3906 SPIRVTypeInst SpirvScalarType = ResType->
getOpcode() == SPIRV::OpTypeVector
3912 "only float operands supported by GLSL extended math");
3915 MIRBuilder, SpirvScalarType);
3917 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
3918 ? SPIRV::OpVectorTimesScalar
3921 if (!selectOpWithSrcs(ArgReg, ResType,
I,
3922 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
3924 if (!selectExtInst(ResVReg, ResType,
I,
3925 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
3935Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
3936 MachineInstr &
I)
const {
3939 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3944bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
3950 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
3958 if (
Def->getOpcode() == SPIRV::OpConstantNull)
3961 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
3962 Def->getOpcode() == SPIRV::OpConstantI)
3975 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
3976 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
3978 Intrinsic::spv_const_composite)) {
3979 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
3980 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
3981 if (!IsZero(
Def->getOperand(i).getReg()))
3990Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
3991 MachineInstr &
I)
const {
3995 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4000Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
4001 MachineInstr &
I)
const {
4005 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4011 SPIRVTypeInst ResType,
4012 MachineInstr &
I)
const {
4016 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4021bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
4022 SPIRVTypeInst ResType,
4023 MachineInstr &
I)
const {
4024 Register SelectFirstArg =
I.getOperand(2).getReg();
4025 Register SelectSecondArg =
I.getOperand(3).getReg();
4034 SPIRV::OpTypeVector;
4041 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
4042 }
else if (IsPtrTy) {
4043 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
4045 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
4049 Opcode = IsScalarBool ? SPIRV::OpSelectSFSCond : SPIRV::OpSelectVFVCond;
4050 }
else if (IsPtrTy) {
4051 Opcode = IsScalarBool ? SPIRV::OpSelectSPSCond : SPIRV::OpSelectVPVCond;
4053 Opcode = IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4056 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4059 .
addUse(
I.getOperand(1).getReg())
4068bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
4069 SPIRVTypeInst ResType,
4071 MachineInstr &InsertAt,
4072 bool IsSigned)
const {
4074 Register ZeroReg = buildZerosVal(ResType, InsertAt);
4075 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
4076 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
4078 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4090bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
4091 SPIRVTypeInst ResType,
4092 MachineInstr &
I,
bool IsSigned,
4093 unsigned Opcode)
const {
4094 Register SrcReg =
I.getOperand(1).getReg();
4100 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
4105 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I,
false);
4107 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
4110bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
4111 SPIRVTypeInst ResType, MachineInstr &
I,
4112 bool IsSigned)
const {
4113 Register SrcReg =
I.getOperand(1).getReg();
4115 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
4119 if (ResType == SrcType)
4120 return BuildCOPY(ResVReg, SrcReg,
I);
4122 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4123 return selectUnOp(ResVReg, ResType,
I, Opcode);
4126bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
4127 SPIRVTypeInst ResType,
4129 bool IsSigned)
const {
4130 MachineIRBuilder MIRBuilder(
I);
4131 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
4146 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
4149 .
addUse(
I.getOperand(1).getReg())
4150 .
addUse(
I.getOperand(2).getReg())
4156 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
4159 .
addUse(
I.getOperand(1).getReg())
4160 .
addUse(
I.getOperand(2).getReg())
4168 unsigned SelectOpcode =
4169 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
4174 .
addUse(buildOnesVal(
true, ResType,
I))
4175 .
addUse(buildZerosVal(ResType,
I))
4182 .
addUse(buildOnesVal(
false, ResType,
I))
4187bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
4190 SPIRVTypeInst IntTy,
4191 SPIRVTypeInst BoolTy)
const {
4194 bool IsVectorTy = IntTy->
getOpcode() == SPIRV::OpTypeVector;
4195 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
4197 Register One = buildOnesVal(
false, IntTy,
I);
4205 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
4214bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
4215 SPIRVTypeInst ResType,
4216 MachineInstr &
I)
const {
4217 Register IntReg =
I.getOperand(1).getReg();
4220 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
4221 if (ArgType == ResType)
4222 return BuildCOPY(ResVReg, IntReg,
I);
4224 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4225 return selectUnOp(ResVReg, ResType,
I, Opcode);
4228bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
4229 SPIRVTypeInst ResType,
4230 MachineInstr &
I)
const {
4231 unsigned Opcode =
I.getOpcode();
4232 unsigned TpOpcode = ResType->
getOpcode();
4234 if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {
4235 assert(Opcode == TargetOpcode::G_CONSTANT &&
4236 I.getOperand(1).getCImm()->isZero());
4237 MachineBasicBlock &DepMBB =
I.getMF()->front();
4240 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
4247 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
4250bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
4251 SPIRVTypeInst ResType,
4252 MachineInstr &
I)
const {
4253 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4260bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
4261 SPIRVTypeInst ResType,
4262 MachineInstr &
I)
const {
4264 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
4268 .
addUse(
I.getOperand(3).getReg())
4270 .
addUse(
I.getOperand(2).getReg());
4271 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
4277bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
4278 SPIRVTypeInst ResType,
4279 MachineInstr &
I)
const {
4280 Type *MaybeResTy =
nullptr;
4285 "Expected aggregate type for extractv instruction");
4287 SPIRV::AccessQualifier::ReadWrite,
false);
4291 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
4294 .
addUse(
I.getOperand(2).getReg());
4295 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
4301bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
4302 SPIRVTypeInst ResType,
4303 MachineInstr &
I)
const {
4304 if (
getImm(
I.getOperand(4), MRI))
4305 return selectInsertVal(ResVReg, ResType,
I);
4307 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
4310 .
addUse(
I.getOperand(2).getReg())
4311 .
addUse(
I.getOperand(3).getReg())
4312 .
addUse(
I.getOperand(4).getReg())
4317bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
4318 SPIRVTypeInst ResType,
4319 MachineInstr &
I)
const {
4320 if (
getImm(
I.getOperand(3), MRI))
4321 return selectExtractVal(ResVReg, ResType,
I);
4323 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
4326 .
addUse(
I.getOperand(2).getReg())
4327 .
addUse(
I.getOperand(3).getReg())
4332bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
4333 SPIRVTypeInst ResType,
4334 MachineInstr &
I)
const {
4335 const bool IsGEPInBounds =
I.getOperand(2).getImm();
4341 ? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
4342 : SPIRV::OpAccessChain)
4343 : (IsGEPInBounds ?
SPIRV::OpInBoundsPtrAccessChain
4344 :
SPIRV::OpPtrAccessChain);
4346 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4350 .
addUse(
I.getOperand(3).getReg());
4352 (Opcode == SPIRV::OpPtrAccessChain ||
4353 Opcode == SPIRV::OpInBoundsPtrAccessChain ||
4354 (
getImm(
I.getOperand(4), MRI) &&
foldImm(
I.getOperand(4), MRI) == 0)) &&
4355 "Cannot translate GEP to OpAccessChain. First index must be 0.");
4358 const unsigned StartingIndex =
4359 (Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)
4362 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
4363 Res.addUse(
I.getOperand(i).getReg());
4364 Res.constrainAllUses(
TII,
TRI, RBI);
4369bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
4371 unsigned Lim =
I.getNumExplicitOperands();
4372 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
4373 Register OpReg =
I.getOperand(i).getReg();
4374 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
4376 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
4377 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
4378 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
4385 MachineFunction *MF =
I.getMF();
4397 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4398 TII.get(SPIRV::OpSpecConstantOp))
4401 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
4403 GR.
add(OpDefine, MIB);
4409bool SPIRVInstructionSelector::selectDerivativeInst(
4410 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
4411 const unsigned DPdOpCode)
const {
4414 errorIfInstrOutsideShader(
I);
4419 Register SrcReg =
I.getOperand(2).getReg();
4424 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4427 .
addUse(
I.getOperand(2).getReg());
4429 MachineIRBuilder MIRBuilder(
I);
4432 if (componentCount != 1)
4436 const TargetRegisterClass *RegClass = GR.
getRegClass(SrcType);
4440 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4445 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4450 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4458bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
4459 SPIRVTypeInst ResType,
4460 MachineInstr &
I)
const {
4464 case Intrinsic::spv_load:
4465 return selectLoad(ResVReg, ResType,
I);
4466 case Intrinsic::spv_store:
4467 return selectStore(
I);
4468 case Intrinsic::spv_extractv:
4469 return selectExtractVal(ResVReg, ResType,
I);
4470 case Intrinsic::spv_insertv:
4471 return selectInsertVal(ResVReg, ResType,
I);
4472 case Intrinsic::spv_extractelt:
4473 return selectExtractElt(ResVReg, ResType,
I);
4474 case Intrinsic::spv_insertelt:
4475 return selectInsertElt(ResVReg, ResType,
I);
4476 case Intrinsic::spv_gep:
4477 return selectGEP(ResVReg, ResType,
I);
4478 case Intrinsic::spv_bitcast: {
4479 Register OpReg =
I.getOperand(2).getReg();
4480 SPIRVTypeInst OpType =
4484 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
4486 case Intrinsic::spv_unref_global:
4487 case Intrinsic::spv_init_global: {
4488 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
4493 Register GVarVReg =
MI->getOperand(0).getReg();
4494 if (!selectGlobalValue(GVarVReg, *
MI, Init))
4499 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
4501 MI->eraseFromParent();
4505 case Intrinsic::spv_undef: {
4506 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4512 case Intrinsic::spv_named_boolean_spec_constant: {
4513 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
4514 : SPIRV::OpSpecConstantFalse;
4516 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
4517 .
addDef(
I.getOperand(0).getReg())
4520 unsigned SpecId =
I.getOperand(2).getImm();
4522 SPIRV::Decoration::SpecId, {SpecId});
4526 case Intrinsic::spv_const_composite: {
4528 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
4534 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
4536 std::function<bool(
Register)> HasSpecConstOperand =
4546 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
4547 J < Def->getNumExplicitOperands(); ++J) {
4548 if (
Def->getOperand(J).isReg() &&
4549 HasSpecConstOperand(
Def->getOperand(J).getReg()))
4555 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
4556 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
4557 : SPIRV::OpConstantComposite;
4558 unsigned ContinuedOpc = HasSpecConst
4559 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
4560 : SPIRV::OpConstantCompositeContinuedINTEL;
4561 MachineIRBuilder MIR(
I);
4563 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
4565 for (
auto *Instr : Instructions) {
4566 Instr->setDebugLoc(
I.getDebugLoc());
4571 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4578 case Intrinsic::spv_assign_name: {
4579 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
4580 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
4581 for (
unsigned i =
I.getNumExplicitDefs() + 2;
4582 i <
I.getNumExplicitOperands(); ++i) {
4583 MIB.
addImm(
I.getOperand(i).getImm());
4588 case Intrinsic::spv_switch: {
4589 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
4590 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
4591 if (
I.getOperand(i).isReg())
4592 MIB.
addReg(
I.getOperand(i).getReg());
4593 else if (
I.getOperand(i).isCImm())
4594 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
4595 else if (
I.getOperand(i).isMBB())
4596 MIB.
addMBB(
I.getOperand(i).getMBB());
4603 case Intrinsic::spv_loop_merge: {
4604 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
4605 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
4606 if (
I.getOperand(i).isMBB())
4607 MIB.
addMBB(
I.getOperand(i).getMBB());
4614 case Intrinsic::spv_loop_control_intel: {
4616 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
4617 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
4622 case Intrinsic::spv_selection_merge: {
4624 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
4625 assert(
I.getOperand(1).isMBB() &&
4626 "operand 1 to spv_selection_merge must be a basic block");
4627 MIB.
addMBB(
I.getOperand(1).getMBB());
4628 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
4632 case Intrinsic::spv_cmpxchg:
4633 return selectAtomicCmpXchg(ResVReg, ResType,
I);
4634 case Intrinsic::spv_unreachable:
4635 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
4638 case Intrinsic::spv_alloca:
4639 return selectFrameIndex(ResVReg, ResType,
I);
4640 case Intrinsic::spv_alloca_array:
4641 return selectAllocaArray(ResVReg, ResType,
I);
4642 case Intrinsic::spv_assume:
4644 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
4645 .
addUse(
I.getOperand(1).getReg())
4650 case Intrinsic::spv_expect:
4652 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
4655 .
addUse(
I.getOperand(2).getReg())
4656 .
addUse(
I.getOperand(3).getReg())
4661 case Intrinsic::arithmetic_fence:
4662 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
4663 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
4666 .
addUse(
I.getOperand(2).getReg())
4670 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
4672 case Intrinsic::spv_thread_id:
4678 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
4680 case Intrinsic::spv_thread_id_in_group:
4686 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
4688 case Intrinsic::spv_group_id:
4694 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
4696 case Intrinsic::spv_flattened_thread_id_in_group:
4703 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
4705 case Intrinsic::spv_workgroup_size:
4706 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
4708 case Intrinsic::spv_global_size:
4709 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
4711 case Intrinsic::spv_global_offset:
4712 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
4714 case Intrinsic::spv_num_workgroups:
4715 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
4717 case Intrinsic::spv_subgroup_size:
4718 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
4720 case Intrinsic::spv_num_subgroups:
4721 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
4723 case Intrinsic::spv_subgroup_id:
4724 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
4725 case Intrinsic::spv_subgroup_local_invocation_id:
4726 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
4727 ResVReg, ResType,
I);
4728 case Intrinsic::spv_subgroup_max_size:
4729 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
4731 case Intrinsic::spv_fdot:
4732 return selectFloatDot(ResVReg, ResType,
I);
4733 case Intrinsic::spv_udot:
4734 case Intrinsic::spv_sdot:
4735 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4737 return selectIntegerDot(ResVReg, ResType,
I,
4738 IID == Intrinsic::spv_sdot);
4739 return selectIntegerDotExpansion(ResVReg, ResType,
I);
4740 case Intrinsic::spv_dot4add_i8packed:
4741 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4743 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
4744 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
4745 case Intrinsic::spv_dot4add_u8packed:
4746 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4748 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
4749 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
4750 case Intrinsic::spv_all:
4751 return selectAll(ResVReg, ResType,
I);
4752 case Intrinsic::spv_any:
4753 return selectAny(ResVReg, ResType,
I);
4754 case Intrinsic::spv_cross:
4755 return selectExtInst(ResVReg, ResType,
I, CL::cross, GL::Cross);
4756 case Intrinsic::spv_distance:
4757 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
4758 case Intrinsic::spv_lerp:
4759 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
4760 case Intrinsic::spv_length:
4761 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
4762 case Intrinsic::spv_degrees:
4763 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
4764 case Intrinsic::spv_faceforward:
4765 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
4766 case Intrinsic::spv_frac:
4767 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
4768 case Intrinsic::spv_isinf:
4769 return selectOpIsInf(ResVReg, ResType,
I);
4770 case Intrinsic::spv_isnan:
4771 return selectOpIsNan(ResVReg, ResType,
I);
4772 case Intrinsic::spv_normalize:
4773 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
4774 case Intrinsic::spv_refract:
4775 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
4776 case Intrinsic::spv_reflect:
4777 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
4778 case Intrinsic::spv_rsqrt:
4779 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
4780 case Intrinsic::spv_sign:
4781 return selectSign(ResVReg, ResType,
I);
4782 case Intrinsic::spv_smoothstep:
4783 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
4784 case Intrinsic::spv_firstbituhigh:
4785 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
4786 case Intrinsic::spv_firstbitshigh:
4787 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
4788 case Intrinsic::spv_firstbitlow:
4789 return selectFirstBitLow(ResVReg, ResType,
I);
4790 case Intrinsic::spv_group_memory_barrier:
4791 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
4792 SPIRV::MemorySemantics::WorkgroupMemory,
4794 case Intrinsic::spv_group_memory_barrier_with_group_sync:
4795 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
4796 SPIRV::MemorySemantics::WorkgroupMemory,
4798 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
4799 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
4800 SPIRV::StorageClass::StorageClass ResSC =
4804 "Generic storage class");
4805 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
4813 case Intrinsic::spv_lifetime_start:
4814 case Intrinsic::spv_lifetime_end: {
4815 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
4816 : SPIRV::OpLifetimeStop;
4817 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
4818 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
4827 case Intrinsic::spv_saturate:
4828 return selectSaturate(ResVReg, ResType,
I);
4829 case Intrinsic::spv_nclamp:
4830 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
4831 case Intrinsic::spv_uclamp:
4832 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
4833 case Intrinsic::spv_sclamp:
4834 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
4835 case Intrinsic::spv_subgroup_prefix_bit_count:
4836 return selectWavePrefixBitCount(ResVReg, ResType,
I);
4837 case Intrinsic::spv_wave_active_countbits:
4838 return selectWaveActiveCountBits(ResVReg, ResType,
I);
4839 case Intrinsic::spv_wave_all_equal:
4840 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
4841 case Intrinsic::spv_wave_all:
4842 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
4843 case Intrinsic::spv_wave_any:
4844 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
4845 case Intrinsic::spv_subgroup_ballot:
4846 return selectWaveOpInst(ResVReg, ResType,
I,
4847 SPIRV::OpGroupNonUniformBallot);
4848 case Intrinsic::spv_wave_is_first_lane:
4849 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
4850 case Intrinsic::spv_wave_reduce_or:
4851 return selectWaveReduceOp(ResVReg, ResType,
I,
4852 SPIRV::OpGroupNonUniformBitwiseOr);
4853 case Intrinsic::spv_wave_reduce_xor:
4854 return selectWaveReduceOp(ResVReg, ResType,
I,
4855 SPIRV::OpGroupNonUniformBitwiseXor);
4856 case Intrinsic::spv_wave_reduce_and:
4857 return selectWaveReduceOp(ResVReg, ResType,
I,
4858 SPIRV::OpGroupNonUniformBitwiseAnd);
4859 case Intrinsic::spv_wave_reduce_umax:
4860 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
4861 case Intrinsic::spv_wave_reduce_max:
4862 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
4863 case Intrinsic::spv_wave_reduce_umin:
4864 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
4865 case Intrinsic::spv_wave_reduce_min:
4866 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
4867 case Intrinsic::spv_wave_reduce_sum:
4868 return selectWaveReduceSum(ResVReg, ResType,
I);
4869 case Intrinsic::spv_wave_product:
4870 return selectWaveReduceProduct(ResVReg, ResType,
I);
4871 case Intrinsic::spv_wave_readlane:
4872 return selectWaveOpInst(ResVReg, ResType,
I,
4873 SPIRV::OpGroupNonUniformShuffle);
4874 case Intrinsic::spv_wave_prefix_sum:
4875 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
4876 case Intrinsic::spv_wave_prefix_product:
4877 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
4878 case Intrinsic::spv_quad_read_across_x: {
4879 return selectQuadSwap(ResVReg, ResType,
I, 0);
4881 case Intrinsic::spv_quad_read_across_y: {
4882 return selectQuadSwap(ResVReg, ResType,
I, 1);
4884 case Intrinsic::spv_step:
4885 return selectExtInst(ResVReg, ResType,
I, CL::step, GL::Step);
4886 case Intrinsic::spv_radians:
4887 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
4891 case Intrinsic::instrprof_increment:
4892 case Intrinsic::instrprof_increment_step:
4893 case Intrinsic::instrprof_value_profile:
4896 case Intrinsic::spv_value_md:
4898 case Intrinsic::spv_resource_handlefrombinding: {
4899 return selectHandleFromBinding(ResVReg, ResType,
I);
4901 case Intrinsic::spv_resource_counterhandlefrombinding:
4902 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
4903 case Intrinsic::spv_resource_updatecounter:
4904 return selectUpdateCounter(ResVReg, ResType,
I);
4905 case Intrinsic::spv_resource_store_typedbuffer: {
4906 return selectImageWriteIntrinsic(
I);
4908 case Intrinsic::spv_resource_load_typedbuffer: {
4909 return selectReadImageIntrinsic(ResVReg, ResType,
I);
4911 case Intrinsic::spv_resource_load_level: {
4912 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
4914 case Intrinsic::spv_resource_getdimensions_x:
4915 case Intrinsic::spv_resource_getdimensions_xy:
4916 case Intrinsic::spv_resource_getdimensions_xyz: {
4917 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
4919 case Intrinsic::spv_resource_getdimensions_levels_x:
4920 case Intrinsic::spv_resource_getdimensions_levels_xy:
4921 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
4922 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
4924 case Intrinsic::spv_resource_getdimensions_ms_xy:
4925 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
4926 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
4928 case Intrinsic::spv_resource_calculate_lod:
4929 case Intrinsic::spv_resource_calculate_lod_unclamped:
4930 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
4931 case Intrinsic::spv_resource_sample:
4932 case Intrinsic::spv_resource_sample_clamp:
4933 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
4934 case Intrinsic::spv_resource_samplebias:
4935 case Intrinsic::spv_resource_samplebias_clamp:
4936 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
4937 case Intrinsic::spv_resource_samplegrad:
4938 case Intrinsic::spv_resource_samplegrad_clamp:
4939 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
4940 case Intrinsic::spv_resource_samplelevel:
4941 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
4942 case Intrinsic::spv_resource_samplecmp:
4943 case Intrinsic::spv_resource_samplecmp_clamp:
4944 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
4945 case Intrinsic::spv_resource_samplecmplevelzero:
4946 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
4947 case Intrinsic::spv_resource_gather:
4948 case Intrinsic::spv_resource_gather_cmp:
4949 return selectGatherIntrinsic(ResVReg, ResType,
I);
4950 case Intrinsic::spv_resource_getpointer: {
4951 return selectResourceGetPointer(ResVReg, ResType,
I);
4953 case Intrinsic::spv_pushconstant_getpointer: {
4954 return selectPushConstantGetPointer(ResVReg, ResType,
I);
4956 case Intrinsic::spv_discard: {
4957 return selectDiscard(ResVReg, ResType,
I);
4959 case Intrinsic::spv_resource_nonuniformindex: {
4960 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
4962 case Intrinsic::spv_unpackhalf2x16: {
4963 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
4965 case Intrinsic::spv_packhalf2x16: {
4966 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
4968 case Intrinsic::spv_ddx:
4969 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
4970 case Intrinsic::spv_ddy:
4971 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
4972 case Intrinsic::spv_ddx_coarse:
4973 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
4974 case Intrinsic::spv_ddy_coarse:
4975 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
4976 case Intrinsic::spv_ddx_fine:
4977 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
4978 case Intrinsic::spv_ddy_fine:
4979 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
4980 case Intrinsic::spv_fwidth:
4981 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
4982 case Intrinsic::spv_masked_gather:
4983 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
4984 return selectMaskedGather(ResVReg, ResType,
I);
4985 return diagnoseUnsupported(
4986 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
4987 case Intrinsic::spv_masked_scatter:
4988 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
4989 return selectMaskedScatter(
I);
4990 return diagnoseUnsupported(
4991 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
4993 std::string DiagMsg;
4994 raw_string_ostream OS(DiagMsg);
4996 DiagMsg =
"Intrinsic selection not implemented: " + DiagMsg;
5003bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
5004 SPIRVTypeInst ResType,
5005 MachineInstr &
I)
const {
5008 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
5015bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
5016 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5018 assert(Intr.getIntrinsicID() ==
5019 Intrinsic::spv_resource_counterhandlefrombinding);
5022 Register MainHandleReg = Intr.getOperand(2).getReg();
5024 assert(MainHandleDef->getIntrinsicID() ==
5025 Intrinsic::spv_resource_handlefrombinding);
5029 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
5030 Register IndexReg = MainHandleDef->getOperand(5).getReg();
5031 std::string CounterName =
5036 MachineIRBuilder MIRBuilder(
I);
5038 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
5040 ArraySize, IndexReg, CounterName, MIRBuilder);
5042 return BuildCOPY(ResVReg, CounterVarReg,
I);
5045bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
5046 SPIRVTypeInst ResType,
5047 MachineInstr &
I)
const {
5049 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
5051 Register CounterHandleReg = Intr.getOperand(2).getReg();
5052 Register IncrReg = Intr.getOperand(3).getReg();
5059 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
5060 assert(CounterVarPointeeType &&
5061 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
5062 "Counter variable must be a struct");
5064 SPIRV::StorageClass::StorageBuffer &&
5065 "Counter variable must be in the storage buffer storage class");
5067 "Counter variable must have exactly 1 member in the struct");
5068 const SPIRVTypeInst MemberType =
5071 "Counter variable struct must have a single i32 member");
5075 MachineIRBuilder MIRBuilder(
I);
5077 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
5080 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
5086 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5089 .
addUse(CounterHandleReg)
5096 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
5099 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
5102 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
5111 return BuildCOPY(ResVReg, AtomicRes,
I);
5119 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
5127bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
5128 SPIRVTypeInst ResType,
5129 MachineInstr &
I)
const {
5137 Register ImageReg =
I.getOperand(2).getReg();
5145 Register IdxReg =
I.getOperand(3).getReg();
5147 MachineInstr &Pos =
I;
5149 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
5153bool SPIRVInstructionSelector::generateSampleImage(
5156 DebugLoc Loc, MachineInstr &Pos)
const {
5167 if (!loadHandleBeforePosition(NewSamplerReg,
5173 MachineIRBuilder MIRBuilder(Pos);
5186 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
5187 ImOps.Lod.has_value();
5188 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
5189 : SPIRV::OpImageSampleImplicitLod;
5191 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
5192 : SPIRV::OpImageSampleDrefImplicitLod;
5201 MIB.
addUse(*ImOps.Compare);
5203 uint32_t ImageOperands = 0;
5205 ImageOperands |= SPIRV::ImageOperand::Bias;
5207 ImageOperands |= SPIRV::ImageOperand::Lod;
5208 if (ImOps.GradX && ImOps.GradY)
5209 ImageOperands |= SPIRV::ImageOperand::Grad;
5210 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
5212 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
5215 "Non-constant offsets are not supported in sample instructions.");
5219 ImageOperands |= SPIRV::ImageOperand::MinLod;
5221 if (ImageOperands != 0) {
5222 MIB.
addImm(ImageOperands);
5223 if (ImageOperands & SPIRV::ImageOperand::Bias)
5225 if (ImageOperands & SPIRV::ImageOperand::Lod)
5227 if (ImageOperands & SPIRV::ImageOperand::Grad) {
5228 MIB.
addUse(*ImOps.GradX);
5229 MIB.
addUse(*ImOps.GradY);
5232 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
5233 MIB.
addUse(*ImOps.Offset);
5234 if (ImageOperands & SPIRV::ImageOperand::MinLod)
5235 MIB.
addUse(*ImOps.MinLod);
5242bool SPIRVInstructionSelector::selectImageQuerySize(
5244 std::optional<Register> LodReg)
const {
5246 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
5249 "ImageReg is not an image type.");
5251 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
5253 unsigned NumComponents = 0;
5255 case SPIRV::Dim::DIM_1D:
5256 case SPIRV::Dim::DIM_Buffer:
5257 NumComponents =
IsArray ? 2 : 1;
5259 case SPIRV::Dim::DIM_2D:
5260 case SPIRV::Dim::DIM_Cube:
5261 case SPIRV::Dim::DIM_Rect:
5262 NumComponents =
IsArray ? 3 : 2;
5264 case SPIRV::Dim::DIM_3D:
5268 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
5273 SPIRVTypeInst ResType =
5278 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5288bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
5289 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5290 Register ImageReg =
I.getOperand(2).getReg();
5297 return selectImageQuerySize(NewImageReg, ResVReg,
I);
5300bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
5301 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5302 Register ImageReg =
I.getOperand(2).getReg();
5311 Register LodReg =
I.getOperand(3).getReg();
5314 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
5316 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
5323 TII.get(SPIRV::OpImageQueryLevels))
5330 TII.get(SPIRV::OpCompositeConstruct))
5340bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
5341 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5342 Register ImageReg =
I.getOperand(2).getReg();
5353 "OpImageQuerySamples requires a multisampled image");
5355 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
5363 TII.get(SPIRV::OpImageQuerySamples))
5370 TII.get(SPIRV::OpCompositeConstruct))
5380bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
5381 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5382 Register ImageReg =
I.getOperand(2).getReg();
5383 Register SamplerReg =
I.getOperand(3).getReg();
5384 Register CoordinateReg =
I.getOperand(4).getReg();
5400 if (!loadHandleBeforePosition(
5405 MachineIRBuilder MIRBuilder(
I);
5411 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
5421 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
5428 unsigned ExtractedIndex =
5430 Intrinsic::spv_resource_calculate_lod_unclamped
5434 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5435 TII.get(SPIRV::OpCompositeExtract))
5445bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
5446 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5447 Register ImageReg =
I.getOperand(2).getReg();
5448 Register SamplerReg =
I.getOperand(3).getReg();
5449 Register CoordinateReg =
I.getOperand(4).getReg();
5450 ImageOperands ImOps;
5451 if (
I.getNumOperands() > 5)
5452 ImOps.Offset =
I.getOperand(5).getReg();
5453 if (
I.getNumOperands() > 6)
5454 ImOps.MinLod =
I.getOperand(6).getReg();
5455 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5456 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5459bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
5460 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5461 Register ImageReg =
I.getOperand(2).getReg();
5462 Register SamplerReg =
I.getOperand(3).getReg();
5463 Register CoordinateReg =
I.getOperand(4).getReg();
5464 ImageOperands ImOps;
5465 ImOps.Bias =
I.getOperand(5).getReg();
5466 if (
I.getNumOperands() > 6)
5467 ImOps.Offset =
I.getOperand(6).getReg();
5468 if (
I.getNumOperands() > 7)
5469 ImOps.MinLod =
I.getOperand(7).getReg();
5470 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5471 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5474bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
5475 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5476 Register ImageReg =
I.getOperand(2).getReg();
5477 Register SamplerReg =
I.getOperand(3).getReg();
5478 Register CoordinateReg =
I.getOperand(4).getReg();
5479 ImageOperands ImOps;
5480 ImOps.GradX =
I.getOperand(5).getReg();
5481 ImOps.GradY =
I.getOperand(6).getReg();
5482 if (
I.getNumOperands() > 7)
5483 ImOps.Offset =
I.getOperand(7).getReg();
5484 if (
I.getNumOperands() > 8)
5485 ImOps.MinLod =
I.getOperand(8).getReg();
5486 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5487 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5490bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
5491 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5492 Register ImageReg =
I.getOperand(2).getReg();
5493 Register SamplerReg =
I.getOperand(3).getReg();
5494 Register CoordinateReg =
I.getOperand(4).getReg();
5495 ImageOperands ImOps;
5496 ImOps.Lod =
I.getOperand(5).getReg();
5497 if (
I.getNumOperands() > 6)
5498 ImOps.Offset =
I.getOperand(6).getReg();
5499 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5500 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5503bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
5504 SPIRVTypeInst ResType,
5505 MachineInstr &
I)
const {
5506 Register ImageReg =
I.getOperand(2).getReg();
5507 Register SamplerReg =
I.getOperand(3).getReg();
5508 Register CoordinateReg =
I.getOperand(4).getReg();
5509 ImageOperands ImOps;
5510 ImOps.Compare =
I.getOperand(5).getReg();
5511 if (
I.getNumOperands() > 6)
5512 ImOps.Offset =
I.getOperand(6).getReg();
5513 if (
I.getNumOperands() > 7)
5514 ImOps.MinLod =
I.getOperand(7).getReg();
5515 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5516 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5519bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
5520 SPIRVTypeInst ResType,
5521 MachineInstr &
I)
const {
5522 Register ImageReg =
I.getOperand(2).getReg();
5523 Register CoordinateReg =
I.getOperand(3).getReg();
5524 Register LodReg =
I.getOperand(4).getReg();
5526 ImageOperands ImOps;
5528 if (
I.getNumOperands() > 5)
5529 ImOps.Offset =
I.getOperand(5).getReg();
5541 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
5542 I.getDebugLoc(),
I, &ImOps);
5545bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
5546 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5547 Register ImageReg =
I.getOperand(2).getReg();
5548 Register SamplerReg =
I.getOperand(3).getReg();
5549 Register CoordinateReg =
I.getOperand(4).getReg();
5550 ImageOperands ImOps;
5551 ImOps.Compare =
I.getOperand(5).getReg();
5552 if (
I.getNumOperands() > 6)
5553 ImOps.Offset =
I.getOperand(6).getReg();
5556 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5557 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5560bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
5561 SPIRVTypeInst ResType,
5562 MachineInstr &
I)
const {
5563 Register ImageReg =
I.getOperand(2).getReg();
5564 Register SamplerReg =
I.getOperand(3).getReg();
5565 Register CoordinateReg =
I.getOperand(4).getReg();
5568 "ImageReg is not an image type.");
5573 ComponentOrCompareReg =
I.getOperand(5).getReg();
5574 OffsetReg =
I.getOperand(6).getReg();
5577 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
5581 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
5582 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
5583 Dim != SPIRV::Dim::DIM_Rect) {
5585 "Gather operations are only supported for 2D, Cube, and Rect images.");
5592 if (!loadHandleBeforePosition(
5597 MachineIRBuilder MIRBuilder(
I);
5598 SPIRVTypeInst SampledImageType =
5603 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
5611 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
5613 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
5615 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5620 .
addUse(ComponentOrCompareReg);
5622 uint32_t ImageOperands = 0;
5623 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
5624 if (Dim == SPIRV::Dim::DIM_Cube) {
5626 "Gather operations with offset are not supported for Cube images.");
5630 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
5632 ImageOperands |= SPIRV::ImageOperand::Offset;
5636 if (ImageOperands != 0) {
5637 MIB.
addImm(ImageOperands);
5639 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
5647bool SPIRVInstructionSelector::generateImageReadOrFetch(
5650 const ImageOperands *ImOps)
const {
5653 "ImageReg is not an image type.");
5655 bool IsSignedInteger =
5660 bool IsFetch = (SampledOp.getImm() == 1);
5662 auto AddOperands = [&](MachineInstrBuilder &MIB) {
5663 uint32_t ImageOperandsMask = 0;
5664 if (IsSignedInteger)
5665 ImageOperandsMask |= 0x1000;
5667 if (IsFetch && ImOps) {
5669 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
5670 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
5672 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
5674 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
5678 if (ImageOperandsMask != 0) {
5679 MIB.
addImm(ImageOperandsMask);
5680 if (IsFetch && ImOps) {
5683 if (ImOps->Offset &&
5684 (ImageOperandsMask &
5685 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
5686 MIB.
addUse(*ImOps->Offset);
5692 if (ResultSize == 4) {
5695 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
5702 BMI.constrainAllUses(
TII,
TRI, RBI);
5706 SPIRVTypeInst ReadType = widenTypeToVec4(ResType, Pos);
5710 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
5716 BMI.constrainAllUses(
TII,
TRI, RBI);
5718 if (ResultSize == 1) {
5727 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
5730bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
5731 SPIRVTypeInst ResType,
5732 MachineInstr &
I)
const {
5733 Register ResourcePtr =
I.getOperand(2).getReg();
5735 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
5744 MachineIRBuilder MIRBuilder(
I);
5746 Register IndexReg =
I.getOperand(3).getReg();
5749 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5759bool SPIRVInstructionSelector::selectPushConstantGetPointer(
5760 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5765bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
5766 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5767 Register ObjReg =
I.getOperand(2).getReg();
5768 if (!BuildCOPY(ResVReg, ObjReg,
I))
5778 decorateUsesAsNonUniform(ResVReg);
5782void SPIRVInstructionSelector::decorateUsesAsNonUniform(
5785 while (WorkList.
size() > 0) {
5789 bool IsDecorated =
false;
5791 if (
Use.getOpcode() == SPIRV::OpDecorate &&
5792 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
5798 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
5800 if (ResultReg == CurrentReg)
5808 SPIRV::Decoration::NonUniformEXT, {});
5813bool SPIRVInstructionSelector::extractSubvector(
5815 MachineInstr &InsertionPoint)
const {
5817 [[maybe_unused]] uint64_t InputSize =
5820 assert(InputSize > 1 &&
"The input must be a vector.");
5821 assert(ResultSize > 1 &&
"The result must be a vector.");
5822 assert(ResultSize < InputSize &&
5823 "Cannot extract more element than there are in the input.");
5826 const TargetRegisterClass *ScalarRegClass = GR.
getRegClass(ScalarType);
5827 for (uint64_t
I = 0;
I < ResultSize;
I++) {
5830 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
5839 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
5841 TII.get(SPIRV::OpCompositeConstruct))
5845 for (
Register ComponentReg : ComponentRegisters)
5846 MIB.
addUse(ComponentReg);
5851bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
5852 MachineInstr &
I)
const {
5859 Register ImageReg =
I.getOperand(1).getReg();
5867 Register CoordinateReg =
I.getOperand(2).getReg();
5868 Register DataReg =
I.getOperand(3).getReg();
5871 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
5879Register SPIRVInstructionSelector::buildPointerToResource(
5880 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
5881 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
5882 StringRef Name, MachineIRBuilder MIRBuilder)
const {
5884 if (ArraySize == 1) {
5885 SPIRVTypeInst PtrType =
5888 "SpirvResType did not have an explicit layout.");
5893 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
5894 SPIRVTypeInst VarPointerType =
5897 VarPointerType, Set,
Binding, Name, MIRBuilder);
5899 SPIRVTypeInst ResPointerType =
5912bool SPIRVInstructionSelector::selectFirstBitSet16(
5913 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
5914 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
5916 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
5920 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
5923bool SPIRVInstructionSelector::selectFirstBitSet32(
5925 unsigned BitSetOpcode)
const {
5926 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
5929 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
5936bool SPIRVInstructionSelector::selectFirstBitSet64Overflow(
5938 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
5945 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
5947 MachineIRBuilder MIRBuilder(
I);
5950 SPIRVTypeInst I64x2Type =
5952 SPIRVTypeInst Vec2ResType =
5955 std::vector<Register> PartialRegs;
5958 unsigned CurrentComponent = 0;
5959 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
5965 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5966 TII.get(SPIRV::OpVectorShuffle))
5971 .
addImm(CurrentComponent)
5972 .
addImm(CurrentComponent + 1);
5979 if (!selectFirstBitSet64(SubVecBitSetReg, Vec2ResType,
I, BitSetResult,
5980 BitSetOpcode, SwapPrimarySide))
5983 PartialRegs.push_back(SubVecBitSetReg);
5987 if (CurrentComponent != ComponentCount) {
5993 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
5994 SPIRV::OpVectorExtractDynamic))
6000 if (!selectFirstBitSet64(FinalElemBitSetReg,
BaseType,
I, FinalElemReg,
6001 BitSetOpcode, SwapPrimarySide))
6004 PartialRegs.push_back(FinalElemBitSetReg);
6009 return selectOpWithSrcs(ResVReg, ResType,
I, std::move(PartialRegs),
6010 SPIRV::OpCompositeConstruct);
6013bool SPIRVInstructionSelector::selectFirstBitSet64(
6015 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
6028 if (ComponentCount > 2) {
6029 return selectFirstBitSet64Overflow(ResVReg, ResType,
I, SrcReg,
6030 BitSetOpcode, SwapPrimarySide);
6034 MachineIRBuilder MIRBuilder(
I);
6036 BaseType, 2 * ComponentCount, MIRBuilder,
false);
6040 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
6046 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
6053 bool IsScalarRes = ResType->
getOpcode() != SPIRV::OpTypeVector;
6056 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
6057 SPIRV::OpVectorExtractDynamic))
6059 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
6060 SPIRV::OpVectorExtractDynamic))
6064 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6065 TII.get(SPIRV::OpVectorShuffle))
6073 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
6079 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6080 TII.get(SPIRV::OpVectorShuffle))
6088 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
6108 SelectOp = SPIRV::OpSelectSISCond;
6109 AddOp = SPIRV::OpIAddS;
6117 SelectOp = SPIRV::OpSelectVIVCond;
6118 AddOp = SPIRV::OpIAddV;
6124 Register RegSecondaryOffset = Reg0;
6128 if (SwapPrimarySide) {
6129 PrimaryReg = LowReg;
6130 SecondaryReg = HighReg;
6131 RegPrimaryOffset = Reg0;
6132 RegSecondaryOffset = Reg32;
6137 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
6138 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
6143 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
6144 SPIRV::OpINotEqual))
6151 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
6152 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
6157 if (SwapPrimarySide) {
6159 if (!selectOpWithSrcs(RegAdd, ResType,
I,
6160 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
6171 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
6172 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
6177 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
6178 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
6181 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
6185bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
6186 SPIRVTypeInst ResType,
6188 bool IsSigned)
const {
6190 Register OpReg =
I.getOperand(2).getReg();
6193 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
6194 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
6198 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6200 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6202 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6206 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
6210bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
6211 SPIRVTypeInst ResType,
6212 MachineInstr &
I)
const {
6214 Register OpReg =
I.getOperand(2).getReg();
6219 unsigned ExtendOpcode = SPIRV::OpUConvert;
6220 unsigned BitSetOpcode = GL::FindILsb;
6224 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6226 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6228 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6235bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
6236 SPIRVTypeInst ResType,
6237 MachineInstr &
I)
const {
6241 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
6244 .
addUse(
I.getOperand(2).getReg())
6247 unsigned Alignment =
I.getOperand(3).getImm();
6253bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
6254 SPIRVTypeInst ResType,
6255 MachineInstr &
I)
const {
6259 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
6262 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
6265 unsigned Alignment =
I.getOperand(2).getImm();
6272bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
6277 const MachineInstr *PrevI =
I.getPrevNode();
6279 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
6283 .
addMBB(
I.getOperand(0).getMBB())
6288 .
addMBB(
I.getOperand(0).getMBB())
6293bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
6304 const MachineInstr *NextI =
I.getNextNode();
6306 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
6312 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
6314 .
addUse(
I.getOperand(0).getReg())
6315 .
addMBB(
I.getOperand(1).getMBB())
6321bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
6322 MachineInstr &
I)
const {
6324 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
6326 const unsigned NumOps =
I.getNumOperands();
6327 for (
unsigned i = 1; i <
NumOps; i += 2) {
6328 MIB.
addUse(
I.getOperand(i + 0).getReg());
6329 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
6335bool SPIRVInstructionSelector::selectGlobalValue(
6336 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
6338 MachineIRBuilder MIRBuilder(
I);
6339 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
6342 std::string GlobalIdent;
6344 unsigned &
ID = UnnamedGlobalIDs[GV];
6346 ID = UnnamedGlobalIDs.
size();
6347 GlobalIdent =
"__unnamed_" + Twine(
ID).str();
6373 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
6380 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
6385 MachineInstrBuilder MIB1 =
6386 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
6389 MachineInstrBuilder MIB2 =
6391 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
6395 GR.
add(ConstVal, MIB2);
6403 MachineInstrBuilder MIB3 =
6404 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
6407 GR.
add(ConstVal, MIB3);
6411 assert(NewReg != ResVReg);
6412 return BuildCOPY(ResVReg, NewReg,
I);
6422 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
6425 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
6431 SPIRVTypeInst ResType =
6435 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
6440 if (
GlobalVar->isExternallyInitialized() &&
6441 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
6442 constexpr unsigned ReadWriteINTEL = 3u;
6445 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
6451bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
6452 SPIRVTypeInst ResType,
6453 MachineInstr &
I)
const {
6455 return selectExtInst(ResVReg, ResType,
I, CL::log10);
6463 MachineIRBuilder MIRBuilder(
I);
6468 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6471 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6473 .
add(
I.getOperand(1))
6478 ResType->
getOpcode() == SPIRV::OpTypeFloat);
6480 SPIRVTypeInst SpirvScalarType = ResType->
getOpcode() == SPIRV::OpTypeVector
6488 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
6489 ? SPIRV::OpVectorTimesScalar
6500bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
6501 SPIRVTypeInst ResType,
6502 MachineInstr &
I)
const {
6505 return selectExtInst(ResVReg, ResType,
I, CL::pown);
6511 Register ExpReg =
I.getOperand(2).getReg();
6513 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
6514 SPIRV::OpConvertSToF))
6516 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
6523bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
6524 SPIRVTypeInst ResType,
6525 MachineInstr &
I)
const {
6541 MachineIRBuilder MIRBuilder(
I);
6544 ResType, MIRBuilder, SPIRV::StorageClass::Function);
6556 MachineBasicBlock &EntryBB =
I.getMF()->front();
6560 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
6563 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
6569 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6572 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
6575 .
add(
I.getOperand(
I.getNumExplicitDefs()))
6579 Register IntegralPartReg =
I.getOperand(1).getReg();
6580 if (IntegralPartReg.
isValid()) {
6582 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6593 assert(
false &&
"GLSL::Modf is deprecated.");
6604bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
6605 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
6606 SPIRVTypeInst ResType, MachineInstr &
I)
const {
6607 MachineIRBuilder MIRBuilder(
I);
6608 const SPIRVTypeInst Vec3Ty =
6611 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
6623 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
6627 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
6633 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6640 assert(
I.getOperand(2).isReg());
6641 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
6645 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6656bool SPIRVInstructionSelector::loadBuiltinInputID(
6657 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
6658 SPIRVTypeInst ResType, MachineInstr &
I)
const {
6659 MachineIRBuilder MIRBuilder(
I);
6661 ResType, MIRBuilder, SPIRV::StorageClass::Input);
6676 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
6680 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6689SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
6690 MachineInstr &
I)
const {
6691 MachineIRBuilder MIRBuilder(
I);
6692 if (
Type->getOpcode() != SPIRV::OpTypeVector)
6696 if (VectorSize == 4)
6704bool SPIRVInstructionSelector::loadHandleBeforePosition(
6705 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
6706 MachineInstr &Pos)
const {
6709 Intrinsic::spv_resource_handlefrombinding);
6717 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
6718 MachineIRBuilder MIRBuilder(HandleDef);
6719 SPIRVTypeInst VarType = ResType;
6720 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
6722 if (IsStructuredBuffer) {
6727 if (ResType->
getOpcode() == SPIRV::OpTypeImage && ArraySize == 0)
6729 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
6732 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
6733 ArraySize, IndexReg, Name, MIRBuilder);
6737 uint32_t LoadOpcode =
6738 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
6748void SPIRVInstructionSelector::errorIfInstrOutsideShader(
6749 MachineInstr &
I)
const {
6751 std::string DiagMsg;
6752 raw_string_ostream OS(DiagMsg);
6753 I.print(OS,
true,
false,
false,
false);
6754 DiagMsg +=
" is only supported in shaders.\n";
6760InstructionSelector *
6764 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)
unsigned getVectorSizeOrOne(SPIRVTypeInst Type)
bool isDead(const MachineInstr &MI, const MachineRegisterInfo &MRI)
BaseType
A given derived pointer can have multiple base pointers through phi/selects.
static TableGen::Emitter::Opt Y("gen-skeleton-entry", EmitSkeleton, "Generate example skeleton entry")
static const fltSemantics & IEEEsingle()
static const fltSemantics & IEEEdouble()
static const fltSemantics & IEEEhalf()
static APFloat getOne(const fltSemantics &Sem, bool Negative=false)
Factory for Positive and Negative One.
static APFloat getZero(const fltSemantics &Sem, bool Negative=false)
Factory for Positive and Negative Zero.
static APInt getAllOnes(unsigned numBits)
Return an APInt of a specified width with all bits set.
ArrayRef - Represent a constant reference to an array (0 or more elements consecutively in memory),...
BlockFrequencyInfo pass uses BlockFrequencyInfoImpl implementation to estimate IR basic block frequen...
Predicate
This enumeration lists the possible predicates for CmpInst subclasses.
@ FCMP_OEQ
0 0 0 1 True if ordered and equal
@ ICMP_SLT
signed less than
@ ICMP_SLE
signed less or equal
@ FCMP_OLT
0 1 0 0 True if ordered and less than
@ FCMP_ULE
1 1 0 1 True if unordered, less than, or equal
@ FCMP_OGT
0 0 1 0 True if ordered and greater than
@ FCMP_OGE
0 0 1 1 True if ordered and greater than or equal
@ ICMP_UGE
unsigned greater or equal
@ ICMP_UGT
unsigned greater than
@ ICMP_SGT
signed greater than
@ FCMP_ULT
1 1 0 0 True if unordered or less than
@ FCMP_ONE
0 1 1 0 True if ordered and operands are unequal
@ FCMP_UEQ
1 0 0 1 True if unordered or equal
@ ICMP_ULT
unsigned less than
@ FCMP_UGT
1 0 1 0 True if unordered or greater than
@ FCMP_OLE
0 1 0 1 True if ordered and less than or equal
@ FCMP_ORD
0 1 1 1 True if ordered (no nans)
@ ICMP_SGE
signed greater or equal
@ FCMP_UNE
1 1 1 0 True if unordered or not equal
@ ICMP_ULE
unsigned less or equal
@ FCMP_UGE
1 0 1 1 True if unordered, greater than, or equal
@ FCMP_UNO
1 0 0 0 True if unordered: isnan(X) | isnan(Y)
static LLVM_ABI Constant * getNullValue(Type *Ty)
Constructor to create a '0' constant of arbitrary type.
LLVMContext & getContext() const
getContext - Return a reference to the LLVMContext associated with this function.
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 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
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.
int64_t getIConstValSext(Register ConstReg, const MachineRegisterInfo *MRI)
MachineInstrBuilder BuildMI(MachineFunction &MF, const MIMetadata &MIMD, const MCInstrDesc &MCID)
Builder interface. Specify how to create the initial instruction itself.
bool isTypeFoldingSupported(unsigned Opcode)
decltype(auto) dyn_cast(const From &Val)
dyn_cast<X> - Return the argument parameter cast to the specified type.
FunctionAddr VTableAddr uintptr_t uintptr_t Int32Ty
void addNumImm(const APInt &Imm, MachineInstrBuilder &MIB)
LLVM_ABI void salvageDebugInfo(const MachineRegisterInfo &MRI, MachineInstr &MI)
Assuming the instruction MI is going to be deleted, attempt to salvage debug users of MI by writing t...
LLVM_ABI void constrainSelectedInstRegOperands(MachineInstr &I, const TargetInstrInfo &TII, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI)
Mutate the newly-selected instruction I to constrain its (possibly generic) virtual register operands...
bool isPreISelGenericOpcode(unsigned Opcode)
Check whether the given Opcode is a generic opcode that is not supposed to appear after ISel.
Register createVirtualRegister(SPIRVTypeInst SpvType, SPIRVGlobalRegistry *GR, MachineRegisterInfo *MRI, const MachineFunction &MF)
unsigned getArrayComponentCount(const MachineRegisterInfo *MRI, const MachineInstr *ResType)
uint64_t getIConstVal(Register ConstReg, const MachineRegisterInfo *MRI)
SmallVector< MachineInstr *, 4 > createContinuedInstructions(MachineIRBuilder &MIRBuilder, unsigned Opcode, unsigned MinWC, unsigned ContinuedOpcode, ArrayRef< Register > Args, Register ReturnRegister, Register TypeID)
SPIRV::MemorySemantics::MemorySemantics getMemSemanticsForStorageClass(SPIRV::StorageClass::StorageClass SC)
constexpr unsigned storageClassToAddressSpace(SPIRV::StorageClass::StorageClass SC)
MachineBasicBlock::iterator getFirstValidInstructionInsertPoint(MachineBasicBlock &BB)
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)
MachineBasicBlock::iterator getOpVariableMBBIt(MachineInstr &I)
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...