34#include "llvm/IR/IntrinsicsSPIRV.h"
40#define DEBUG_TYPE "spirv-isel"
47 std::vector<std::pair<SPIRV::InstructionSet::InstructionSet, uint32_t>>;
52 std::optional<Register> Bias;
53 std::optional<Register>
Offset;
54 std::optional<Register> MinLod;
55 std::optional<Register> GradX;
56 std::optional<Register> GradY;
57 std::optional<Register> Lod;
58 std::optional<Register> Compare;
65 bool IsScalar =
false;
68llvm::SPIRV::SelectionControl::SelectionControl
69getSelectionOperandForImm(
int Imm) {
71 return SPIRV::SelectionControl::Flatten;
73 return SPIRV::SelectionControl::DontFlatten;
75 return SPIRV::SelectionControl::None;
79#define GET_GLOBALISEL_PREDICATE_BITSET
80#include "SPIRVGenGlobalISel.inc"
81#undef GET_GLOBALISEL_PREDICATE_BITSET
108#define GET_GLOBALISEL_PREDICATES_DECL
109#include "SPIRVGenGlobalISel.inc"
110#undef GET_GLOBALISEL_PREDICATES_DECL
112#define GET_GLOBALISEL_TEMPORARIES_DECL
113#include "SPIRVGenGlobalISel.inc"
114#undef GET_GLOBALISEL_TEMPORARIES_DECL
138 unsigned BitSetOpcode)
const;
142 unsigned BitSetOpcode)
const;
146 unsigned BitSetOpcode,
bool SwapPrimarySide)
const;
153 unsigned Opcode)
const;
156 unsigned Opcode)
const;
178 unsigned NewOpcode,
unsigned NegateOpcode = 0)
const;
189 unsigned OpType)
const;
254 unsigned Opcode)
const;
258 unsigned Opcode)
const;
262 unsigned Opcode)
const;
266 unsigned Opcode)
const;
268 template <
bool Signed>
271 template <
bool Signed>
278 template <
typename PickOpcodeFn>
281 PickOpcodeFn &&PickOpcode)
const;
298 template <
typename PickOpcodeFn>
301 PickOpcodeFn &&PickOpcode)
const;
319 bool IsSigned)
const;
321 bool IsSigned,
unsigned Opcode)
const;
323 bool IsSigned)
const;
329 bool IsSigned)
const;
370 GL::GLSLExtInst GLInst,
bool setMIFlags =
true,
371 bool useMISrc =
true,
373 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
374 CL::OpenCLExtInst CLInst,
bool setMIFlags =
true,
375 bool useMISrc =
true,
377 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
378 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
379 bool setMIFlags =
true,
bool useMISrc =
true,
381 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
382 const ExtInstList &ExtInsts,
bool setMIFlags =
true,
383 bool useMISrc =
true,
386 bool selectLog10(
Register ResVReg, SPIRVTypeInst ResType,
387 MachineInstr &
I)
const;
389 bool selectFpowi(
Register ResVReg, SPIRVTypeInst ResType,
390 MachineInstr &
I)
const;
392 bool selectSaturate(
Register ResVReg, SPIRVTypeInst ResType,
393 MachineInstr &
I)
const;
395 bool selectWaveOpInst(
Register ResVReg, SPIRVTypeInst ResType,
396 MachineInstr &
I,
unsigned Opcode)
const;
398 bool selectBarrierInst(MachineInstr &
I,
unsigned Scope,
unsigned MemSem,
399 bool WithGroupSync)
const;
401 bool selectWaveActiveCountBits(
Register ResVReg, SPIRVTypeInst ResType,
402 MachineInstr &
I)
const;
404 bool selectWaveActiveAllEqual(
Register ResVReg, SPIRVTypeInst ResType,
405 MachineInstr &
I)
const;
409 bool selectHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
410 MachineInstr &
I)
const;
412 bool selectCounterHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
413 MachineInstr &
I)
const;
415 bool selectReadImageIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
416 MachineInstr &
I)
const;
417 bool selectGetDimensionsIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
418 MachineInstr &
I)
const;
419 bool selectGetDimensionsLevelsIntrinsic(
Register &ResVReg,
420 SPIRVTypeInst ResType,
421 MachineInstr &
I)
const;
422 bool selectGetDimensionsMSIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
423 MachineInstr &
I)
const;
426 std::optional<Register> LodReg = std::nullopt)
const;
427 bool selectSampleBasicIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
428 MachineInstr &
I)
const;
429 bool selectCalculateLodIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
430 MachineInstr &
I)
const;
431 bool selectSampleBiasIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
432 MachineInstr &
I)
const;
433 bool selectSampleGradIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
434 MachineInstr &
I)
const;
435 bool selectSampleLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
436 MachineInstr &
I)
const;
437 bool selectLoadLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
438 MachineInstr &
I)
const;
439 bool selectSampleCmpIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
440 MachineInstr &
I)
const;
441 bool selectSampleCmpLevelZeroIntrinsic(
Register &ResVReg,
442 SPIRVTypeInst ResType,
443 MachineInstr &
I)
const;
444 bool selectGatherIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
445 MachineInstr &
I)
const;
446 bool selectImageWriteIntrinsic(MachineInstr &
I)
const;
447 bool selectResourceGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
448 MachineInstr &
I)
const;
449 bool selectPushConstantGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
450 MachineInstr &
I)
const;
451 bool selectResourceNonUniformIndex(
Register &ResVReg, SPIRVTypeInst ResType,
452 MachineInstr &
I)
const;
453 bool selectModf(
Register ResVReg, SPIRVTypeInst ResType,
454 MachineInstr &
I)
const;
455 bool selectUpdateCounter(
Register &ResVReg, SPIRVTypeInst ResType,
456 MachineInstr &
I)
const;
457 bool selectFrexp(
Register ResVReg, SPIRVTypeInst ResType,
458 MachineInstr &
I)
const;
459 bool selectSincos(
Register ResVReg, SPIRVTypeInst ResType,
460 MachineInstr &
I)
const;
461 bool selectExp10(
Register ResVReg, SPIRVTypeInst ResType,
462 MachineInstr &
I)
const;
463 bool selectDerivativeInst(
Register ResVReg, SPIRVTypeInst ResType,
464 MachineInstr &
I,
const unsigned DPdOpCode)
const;
466 Register buildI32Constant(uint32_t Val, MachineInstr &
I,
467 SPIRVTypeInst ResType =
nullptr)
const;
468 Register buildI32ConstantInEntryBlock(uint32_t Val, MachineInstr &
I,
469 SPIRVTypeInst ResType =
nullptr)
const;
471 Register buildZerosVal(SPIRVTypeInst ResType, MachineInstr &
I)
const;
472 bool isScalarOrVectorIntConstantZero(
Register Reg)
const;
473 Register buildZerosValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
475 MachineInstr &
I)
const;
476 Register buildOnesValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
478 bool wrapIntoSpecConstantOp(MachineInstr &
I,
481 Register getUcharPtrTypeReg(MachineInstr &
I,
482 SPIRV::StorageClass::StorageClass SC)
const;
483 MachineInstrBuilder buildSpecConstantOp(MachineInstr &
I,
Register Dest,
485 uint32_t Opcode)
const;
486 MachineInstrBuilder buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
487 SPIRVTypeInst SrcPtrTy)
const;
488 Register buildPointerToResource(SPIRVTypeInst ResType,
489 SPIRV::StorageClass::StorageClass SC,
490 uint32_t Set, uint32_t
Binding,
491 uint32_t ArraySize,
Register IndexReg,
493 MachineIRBuilder MIRBuilder)
const;
494 SPIRVTypeInst widenTypeToVec4(SPIRVTypeInst
Type, MachineInstr &
I)
const;
495 bool extractSubvector(
Register &ResVReg, SPIRVTypeInst ResType,
496 Register &ReadReg, MachineInstr &InsertionPoint)
const;
497 bool generateImageReadOrFetch(
Register &ResVReg, SPIRVTypeInst ResType,
500 const ImageOperands *ImOps =
nullptr)
const;
501 bool generateSampleImage(
Register ResVReg, SPIRVTypeInst ResType,
503 Register CoordinateReg,
const ImageOperands &ImOps,
506 bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
507 Register ResVReg, SPIRVTypeInst ResType,
508 MachineInstr &
I)
const;
509 bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
510 Register ResVReg, SPIRVTypeInst ResType,
511 MachineInstr &
I)
const;
512 bool loadHandleBeforePosition(
Register &HandleReg, SPIRVTypeInst ResType,
513 GIntrinsic &HandleDef, MachineInstr &Pos)
const;
514 void decorateUsesAsNonUniform(
Register &NonUniformReg)
const;
515 bool errorIfInstrOutsideShader(MachineInstr &
I)
const;
517 std::optional<SplitParts> splitEvenOddLanes(
Register PopCountReg,
518 unsigned ComponentCount,
520 SPIRVTypeInst I32Type)
const;
523 handle64BitOverflow(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
524 Register SrcReg,
unsigned int Opcode,
525 std::function<
bool(
Register, SPIRVTypeInst,
526 MachineInstr &,
Register,
unsigned)>
530bool sampledTypeIsSignedInteger(
const llvm::Type *HandleType) {
532 if (
TET->getTargetExtName() ==
"spirv.Image") {
535 assert(
TET->getTargetExtName() ==
"spirv.SignedImage");
536 return TET->getTypeParameter(0)->isIntegerTy();
540#define GET_GLOBALISEL_IMPL
541#include "SPIRVGenGlobalISel.inc"
542#undef GET_GLOBALISEL_IMPL
548 TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
551#include
"SPIRVGenGlobalISel.inc"
554#include
"SPIRVGenGlobalISel.inc"
566 InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
570void SPIRVInstructionSelector::resetVRegsType(MachineFunction &MF) {
571 if (HasVRegsReset == &MF)
586 for (
const auto &
MBB : MF) {
587 for (
const auto &
MI :
MBB) {
590 if (
MI.getOpcode() != SPIRV::ASSIGN_TYPE)
594 LLT DstType = MRI.
getType(DstReg);
596 LLT SrcType = MRI.
getType(SrcReg);
597 if (DstType != SrcType)
602 if (DstRC != SrcRC && SrcRC)
614 while (!Stack.empty()) {
619 switch (
MI->getOpcode()) {
620 case TargetOpcode::G_INTRINSIC:
621 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
622 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS: {
625 if (IntrID != Intrinsic::spv_const_composite &&
626 IntrID != Intrinsic::spv_undef)
630 case TargetOpcode::G_BUILD_VECTOR:
631 case TargetOpcode::G_SPLAT_VECTOR:
633 i < OpDef->getNumOperands(); i++) {
638 Stack.push_back(OpNestedDef);
641 case TargetOpcode::G_CONSTANT:
642 case TargetOpcode::G_FCONSTANT:
643 case TargetOpcode::G_IMPLICIT_DEF:
644 case SPIRV::OpConstantTrue:
645 case SPIRV::OpConstantFalse:
646 case SPIRV::OpConstantI:
647 case SPIRV::OpConstantF:
648 case SPIRV::OpConstantComposite:
649 case SPIRV::OpConstantCompositeContinuedINTEL:
650 case SPIRV::OpConstantSampler:
651 case SPIRV::OpConstantNull:
653 case SPIRV::OpConstantFunctionPointerINTEL:
680 case Intrinsic::spv_all:
681 case Intrinsic::spv_alloca:
682 case Intrinsic::spv_any:
683 case Intrinsic::spv_bitcast:
684 case Intrinsic::spv_const_composite:
685 case Intrinsic::spv_cross:
686 case Intrinsic::spv_degrees:
687 case Intrinsic::spv_distance:
688 case Intrinsic::spv_extractelt:
689 case Intrinsic::spv_extractv:
690 case Intrinsic::spv_faceforward:
691 case Intrinsic::spv_fdot:
692 case Intrinsic::spv_firstbitlow:
693 case Intrinsic::spv_firstbitshigh:
694 case Intrinsic::spv_firstbituhigh:
695 case Intrinsic::spv_frac:
696 case Intrinsic::spv_gep:
697 case Intrinsic::spv_global_offset:
698 case Intrinsic::spv_global_size:
699 case Intrinsic::spv_group_id:
700 case Intrinsic::spv_insertelt:
701 case Intrinsic::spv_insertv:
702 case Intrinsic::spv_isinf:
703 case Intrinsic::spv_isnan:
704 case Intrinsic::spv_isfinite:
705 case Intrinsic::spv_isnormal:
706 case Intrinsic::spv_lerp:
707 case Intrinsic::spv_length:
708 case Intrinsic::spv_normalize:
709 case Intrinsic::spv_num_subgroups:
710 case Intrinsic::spv_num_workgroups:
711 case Intrinsic::spv_ptrcast:
712 case Intrinsic::spv_radians:
713 case Intrinsic::spv_reflect:
714 case Intrinsic::spv_refract:
715 case Intrinsic::spv_resource_getbasepointer:
716 case Intrinsic::spv_resource_getpointer:
717 case Intrinsic::spv_resource_handlefrombinding:
718 case Intrinsic::spv_resource_handlefromimplicitbinding:
719 case Intrinsic::spv_resource_nonuniformindex:
720 case Intrinsic::spv_resource_sample:
721 case Intrinsic::spv_rsqrt:
722 case Intrinsic::spv_saturate:
723 case Intrinsic::spv_sdot:
724 case Intrinsic::spv_sign:
725 case Intrinsic::spv_smoothstep:
726 case Intrinsic::spv_step:
727 case Intrinsic::spv_subgroup_id:
728 case Intrinsic::spv_subgroup_local_invocation_id:
729 case Intrinsic::spv_subgroup_max_size:
730 case Intrinsic::spv_subgroup_size:
731 case Intrinsic::spv_thread_id:
732 case Intrinsic::spv_thread_id_in_group:
733 case Intrinsic::spv_udot:
734 case Intrinsic::spv_undef:
735 case Intrinsic::spv_value_md:
736 case Intrinsic::spv_workgroup_size:
748 case SPIRV::OpTypeVoid:
749 case SPIRV::OpTypeBool:
750 case SPIRV::OpTypeInt:
751 case SPIRV::OpTypeFloat:
752 case SPIRV::OpTypeVector:
753 case SPIRV::OpTypeMatrix:
754 case SPIRV::OpTypeImage:
755 case SPIRV::OpTypeSampler:
756 case SPIRV::OpTypeSampledImage:
757 case SPIRV::OpTypeArray:
758 case SPIRV::OpTypeRuntimeArray:
759 case SPIRV::OpTypeStruct:
760 case SPIRV::OpTypeOpaque:
761 case SPIRV::OpTypePointer:
762 case SPIRV::OpTypeFunction:
763 case SPIRV::OpTypeEvent:
764 case SPIRV::OpTypeDeviceEvent:
765 case SPIRV::OpTypeReserveId:
766 case SPIRV::OpTypeQueue:
767 case SPIRV::OpTypePipe:
768 case SPIRV::OpTypeForwardPointer:
769 case SPIRV::OpTypePipeStorage:
770 case SPIRV::OpTypeNamedBarrier:
771 case SPIRV::OpTypeAccelerationStructureNV:
772 case SPIRV::OpTypeCooperativeMatrixNV:
773 case SPIRV::OpTypeCooperativeMatrixKHR:
783 if (
MI.getNumDefs() == 0)
786 for (
const auto &MO :
MI.all_defs()) {
788 if (
Reg.isPhysical()) {
793 if (
UseMI.getOpcode() != SPIRV::OpName) {
800 if (
MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE ||
MI.isFakeUse() ||
801 MI.isLifetimeMarker()) {
804 <<
"Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
815 if (
MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
816 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
819 LLVM_DEBUG(
dbgs() <<
"Dead: Intrinsic with no real side effects.\n");
824 if (
MI.mayStore() ||
MI.isCall() ||
825 (
MI.mayLoad() &&
MI.hasOrderedMemoryRef()) ||
MI.isPosition() ||
826 MI.isDebugInstr() ||
MI.isTerminator() ||
MI.isJumpTableDebugInfo()) {
827 LLVM_DEBUG(
dbgs() <<
"Not dead: instruction has side effects.\n");
838 LLVM_DEBUG(
dbgs() <<
"Dead: known opcode with no side effects\n");
845void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &
MI)
const {
847 for (
const auto &MO :
MI.all_defs()) {
851 SmallVector<MachineInstr *, 4> UselessOpNames;
854 "There is still a use of the dead function.");
857 for (MachineInstr *OpNameMI : UselessOpNames) {
859 OpNameMI->eraseFromParent();
864void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &
MI)
const {
867 removeOpNamesForDeadMI(
MI);
868 MI.eraseFromParent();
871bool SPIRVInstructionSelector::select(MachineInstr &
I) {
872 resetVRegsType(*
I.getParent()->getParent());
874 assert(
I.getParent() &&
"Instruction should be in a basic block!");
875 assert(
I.getParent()->getParent() &&
"Instruction should be in a function!");
880 removeDeadInstruction(
I);
887 if (Opcode == SPIRV::ASSIGN_TYPE) {
888 Register DstReg =
I.getOperand(0).getReg();
889 Register SrcReg =
I.getOperand(1).getReg();
892 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
893 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
894 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
895 Register SelectDstReg =
Def->getOperand(0).getReg();
896 bool SuccessToSelectSelect [[maybe_unused]] = selectSelect(
898 assert(SuccessToSelectSelect);
900 Def->eraseFromParent();
907 bool Res = selectImpl(
I, *CoverageInfo);
909 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
910 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
914 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
926 }
else if (
I.getNumDefs() == 1) {
938 removeDeadInstruction(
I);
943 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
944 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
950 bool HasDefs =
I.getNumDefs() > 0;
953 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
954 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
955 if (spvSelect(ResVReg, ResType,
I)) {
957 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
968 case TargetOpcode::G_CONSTANT:
969 case TargetOpcode::G_FCONSTANT:
976 MachineInstr &
I)
const {
979 if (DstRC != SrcRC && SrcRC)
981 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
988bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
989 SPIRVTypeInst ResType,
990 MachineInstr &
I)
const {
991 const unsigned Opcode =
I.getOpcode();
993 return selectImpl(
I, *CoverageInfo);
995 case TargetOpcode::G_CONSTANT:
996 case TargetOpcode::G_FCONSTANT:
997 return selectConst(ResVReg, ResType,
I);
998 case TargetOpcode::G_GLOBAL_VALUE:
999 return selectGlobalValue(ResVReg,
I);
1000 case TargetOpcode::G_IMPLICIT_DEF:
1001 return selectOpUndef(ResVReg, ResType,
I);
1002 case TargetOpcode::G_FREEZE:
1003 return selectFreeze(ResVReg, ResType,
I);
1005 case TargetOpcode::G_INTRINSIC:
1006 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
1007 case TargetOpcode::G_INTRINSIC_CONVERGENT:
1008 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
1009 return selectIntrinsic(ResVReg, ResType,
I);
1010 case TargetOpcode::G_BITREVERSE:
1011 return selectBitreverse(ResVReg, ResType,
I);
1013 case TargetOpcode::G_BUILD_VECTOR:
1014 return selectBuildVector(ResVReg, ResType,
I);
1015 case TargetOpcode::G_SPLAT_VECTOR:
1016 return selectSplatVector(ResVReg, ResType,
I);
1018 case TargetOpcode::G_SHUFFLE_VECTOR: {
1019 MachineBasicBlock &BB = *
I.getParent();
1020 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
1023 .
addUse(
I.getOperand(1).getReg())
1024 .
addUse(
I.getOperand(2).getReg());
1025 for (
auto V :
I.getOperand(3).getShuffleMask())
1030 case TargetOpcode::G_MEMMOVE:
1031 case TargetOpcode::G_MEMCPY:
1032 case TargetOpcode::G_MEMSET:
1033 return selectMemOperation(ResVReg,
I);
1035 case TargetOpcode::G_ICMP:
1036 return selectICmp(ResVReg, ResType,
I);
1037 case TargetOpcode::G_FCMP:
1038 return selectFCmp(ResVReg, ResType,
I);
1040 case TargetOpcode::G_FRAME_INDEX:
1041 return selectFrameIndex(ResVReg, ResType,
I);
1043 case TargetOpcode::G_LOAD:
1044 return selectLoad(ResVReg, ResType,
I);
1045 case TargetOpcode::G_STORE:
1046 return selectStore(
I);
1048 case TargetOpcode::G_BR:
1049 return selectBranch(
I);
1050 case TargetOpcode::G_BRCOND:
1051 return selectBranchCond(
I);
1053 case TargetOpcode::G_PHI:
1054 return selectPhi(ResVReg,
I);
1056 case TargetOpcode::G_FPTOSI:
1057 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1058 case TargetOpcode::G_FPTOUI:
1059 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1061 case TargetOpcode::G_FPTOSI_SAT:
1062 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1063 case TargetOpcode::G_FPTOUI_SAT:
1064 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1066 case TargetOpcode::G_SITOFP:
1067 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
1068 case TargetOpcode::G_UITOFP:
1069 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
1071 case TargetOpcode::G_CTPOP:
1072 return selectPopCount(ResVReg, ResType,
I, SPIRV::OpBitCount);
1073 case TargetOpcode::G_SMIN:
1074 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
1075 case TargetOpcode::G_UMIN:
1076 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
1078 case TargetOpcode::G_SMAX:
1079 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
1080 case TargetOpcode::G_UMAX:
1081 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
1083 case TargetOpcode::G_SCMP:
1084 return selectSUCmp(ResVReg, ResType,
I,
true);
1085 case TargetOpcode::G_UCMP:
1086 return selectSUCmp(ResVReg, ResType,
I,
false);
1087 case TargetOpcode::G_LROUND:
1088 case TargetOpcode::G_LLROUND: {
1091 MRI->
setRegClass(regForLround, &SPIRV::iIDRegClass);
1093 regForLround, *(
I.getParent()->getParent()));
1095 CL::round, GL::Round,
false);
1097 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1104 case TargetOpcode::G_STRICT_FMA:
1105 case TargetOpcode::G_FMA: {
1108 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1111 .
addUse(
I.getOperand(1).getReg())
1112 .
addUse(
I.getOperand(2).getReg())
1113 .
addUse(
I.getOperand(3).getReg())
1118 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1121 case TargetOpcode::G_STRICT_FLDEXP:
1122 return selectExtInst(ResVReg, ResType,
I, CL::ldexp);
1124 case TargetOpcode::G_FPOW:
1125 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1126 case TargetOpcode::G_FPOWI:
1127 return selectFpowi(ResVReg, ResType,
I);
1129 case TargetOpcode::G_FEXP:
1130 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1131 case TargetOpcode::G_FEXP2:
1132 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1133 case TargetOpcode::G_FEXP10:
1134 return selectExp10(ResVReg, ResType,
I);
1136 case TargetOpcode::G_FMODF:
1137 return selectModf(ResVReg, ResType,
I);
1138 case TargetOpcode::G_FSINCOS:
1139 return selectSincos(ResVReg, ResType,
I);
1141 case TargetOpcode::G_FLOG:
1142 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1143 case TargetOpcode::G_FLOG2:
1144 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1145 case TargetOpcode::G_FLOG10:
1146 return selectLog10(ResVReg, ResType,
I);
1148 case TargetOpcode::G_FABS:
1149 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1150 case TargetOpcode::G_ABS:
1151 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1153 case TargetOpcode::G_FMINNUM:
1154 case TargetOpcode::G_FMINIMUM:
1155 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1156 case TargetOpcode::G_FMAXNUM:
1157 case TargetOpcode::G_FMAXIMUM:
1158 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1160 case TargetOpcode::G_FCOPYSIGN:
1161 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1163 case TargetOpcode::G_FCEIL:
1164 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1165 case TargetOpcode::G_FFLOOR:
1166 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1168 case TargetOpcode::G_FCOS:
1169 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1170 case TargetOpcode::G_FSIN:
1171 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1172 case TargetOpcode::G_FTAN:
1173 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1174 case TargetOpcode::G_FACOS:
1175 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1176 case TargetOpcode::G_FASIN:
1177 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1178 case TargetOpcode::G_FATAN:
1179 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1180 case TargetOpcode::G_FATAN2:
1181 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1182 case TargetOpcode::G_FCOSH:
1183 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1184 case TargetOpcode::G_FSINH:
1185 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1186 case TargetOpcode::G_FTANH:
1187 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1189 case TargetOpcode::G_STRICT_FSQRT:
1190 case TargetOpcode::G_FSQRT:
1191 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1193 case TargetOpcode::G_CTTZ:
1194 case TargetOpcode::G_CTTZ_ZERO_POISON:
1195 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1196 case TargetOpcode::G_CTLZ:
1197 case TargetOpcode::G_CTLZ_ZERO_POISON:
1198 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1200 case TargetOpcode::G_INTRINSIC_ROUND:
1201 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1202 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1203 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1204 case TargetOpcode::G_INTRINSIC_TRUNC:
1205 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1206 case TargetOpcode::G_FRINT:
1207 case TargetOpcode::G_FNEARBYINT:
1208 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1210 case TargetOpcode::G_SMULH:
1211 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1212 case TargetOpcode::G_UMULH:
1213 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1215 case TargetOpcode::G_SADDSAT:
1216 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1217 case TargetOpcode::G_UADDSAT:
1218 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1219 case TargetOpcode::G_SSUBSAT:
1220 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1221 case TargetOpcode::G_USUBSAT:
1222 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1224 case TargetOpcode::G_FFREXP:
1225 return selectFrexp(ResVReg, ResType,
I);
1227 case TargetOpcode::G_UADDO:
1228 return selectOverflowArith(ResVReg, ResType,
I,
1229 ResType->
getOpcode() == SPIRV::OpTypeVector
1230 ? SPIRV::OpIAddCarryV
1231 : SPIRV::OpIAddCarryS);
1232 case TargetOpcode::G_USUBO:
1233 return selectOverflowArith(ResVReg, ResType,
I,
1234 ResType->
getOpcode() == SPIRV::OpTypeVector
1235 ? SPIRV::OpISubBorrowV
1236 : SPIRV::OpISubBorrowS);
1237 case TargetOpcode::G_UMULO:
1238 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1239 case TargetOpcode::G_SMULO:
1240 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1242 case TargetOpcode::G_SEXT:
1243 return selectExt(ResVReg, ResType,
I,
true);
1244 case TargetOpcode::G_ANYEXT:
1245 case TargetOpcode::G_ZEXT:
1246 return selectExt(ResVReg, ResType,
I,
false);
1247 case TargetOpcode::G_TRUNC:
1248 return selectTrunc(ResVReg, ResType,
I);
1249 case TargetOpcode::G_FPTRUNC:
1250 case TargetOpcode::G_FPEXT:
1251 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1253 case TargetOpcode::G_PTRTOINT:
1254 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1255 case TargetOpcode::G_INTTOPTR:
1256 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1257 case TargetOpcode::G_BITCAST:
1258 return selectBitcast(ResVReg, ResType,
I);
1259 case TargetOpcode::G_ADDRSPACE_CAST:
1260 return selectAddrSpaceCast(ResVReg, ResType,
I);
1261 case TargetOpcode::G_PTR_ADD: {
1263 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1267 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1268 (*II).getOpcode() == TargetOpcode::COPY ||
1269 (*II).getOpcode() == SPIRV::OpVariable) &&
1270 getImm(
I.getOperand(2), MRI));
1272 bool IsGVInit =
false;
1276 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1277 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1278 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1279 (*UseIt).getOpcode() == SPIRV::OpVariable) {
1289 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1301 return diagnoseUnsupported(
1302 I,
"incompatible result and operand types in a bitcast");
1304 MachineInstrBuilder MIB =
1305 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1312 : SPIRV::OpInBoundsPtrAccessChain))
1316 .
addUse(
I.getOperand(2).getReg())
1319 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1323 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1325 .
addUse(
I.getOperand(2).getReg())
1334 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1337 .
addImm(
static_cast<uint32_t
>(
1338 SPIRV::Opcode::InBoundsPtrAccessChain))
1341 .
addUse(
I.getOperand(2).getReg());
1346 case TargetOpcode::G_ATOMICRMW_OR:
1347 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1348 case TargetOpcode::G_ATOMICRMW_ADD:
1349 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1350 case TargetOpcode::G_ATOMICRMW_AND:
1351 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1352 case TargetOpcode::G_ATOMICRMW_MAX:
1353 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1354 case TargetOpcode::G_ATOMICRMW_MIN:
1355 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1356 case TargetOpcode::G_ATOMICRMW_SUB:
1357 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1358 case TargetOpcode::G_ATOMICRMW_XOR:
1359 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1360 case TargetOpcode::G_ATOMICRMW_UMAX:
1361 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1362 case TargetOpcode::G_ATOMICRMW_UMIN:
1363 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1364 case TargetOpcode::G_ATOMICRMW_XCHG:
1365 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1367 case TargetOpcode::G_ATOMICRMW_FADD:
1368 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1369 case TargetOpcode::G_ATOMICRMW_FSUB:
1371 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1372 ResType->
getOpcode() == SPIRV::OpTypeVector
1374 : SPIRV::OpFNegate);
1375 case TargetOpcode::G_ATOMICRMW_FMIN:
1376 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1377 case TargetOpcode::G_ATOMICRMW_FMAX:
1378 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1380 case TargetOpcode::G_FENCE:
1381 return selectFence(
I);
1383 case TargetOpcode::G_STACKSAVE:
1384 return selectStackSave(ResVReg, ResType,
I);
1385 case TargetOpcode::G_STACKRESTORE:
1386 return selectStackRestore(
I);
1388 case TargetOpcode::G_UNMERGE_VALUES:
1391 case TargetOpcode::G_TRAP:
1392 case TargetOpcode::G_UBSANTRAP:
1393 return selectTrap(
I);
1398 case TargetOpcode::DBG_LABEL:
1400 case TargetOpcode::G_DEBUGTRAP:
1401 return selectDebugTrap(ResVReg, ResType,
I);
1408bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1409 SPIRVTypeInst ResType,
1410 MachineInstr &
I)
const {
1411 unsigned Opcode = SPIRV::OpNop;
1418bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1419 SPIRVTypeInst ResType,
1421 GL::GLSLExtInst GLInst,
1422 bool setMIFlags,
bool useMISrc,
1425 SPIRV::InstructionSet::InstructionSet::GLSL_std_450))
1426 return diagnoseUnsupported(
1428 "this instruction is only supported with the GLSL extended instruction "
1430 return selectExtInst(ResVReg, ResType,
I,
1431 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1432 setMIFlags, useMISrc, SrcRegs);
1435bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1436 SPIRVTypeInst ResType,
1438 CL::OpenCLExtInst CLInst,
1439 bool setMIFlags,
bool useMISrc,
1441 return selectExtInst(ResVReg, ResType,
I,
1442 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1443 setMIFlags, useMISrc, SrcRegs);
1446bool SPIRVInstructionSelector::selectExtInst(
1447 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1448 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1450 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1451 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1452 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1456bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1457 SPIRVTypeInst ResType,
1460 bool setMIFlags,
bool useMISrc,
1463 for (
const auto &[InstructionSet, Opcode] : Insts) {
1467 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1470 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1475 const unsigned NumOps =
I.getNumOperands();
1478 I.getOperand(Index).getType() ==
1479 MachineOperand::MachineOperandType::MO_IntrinsicID)
1482 MIB.
add(
I.getOperand(Index));
1494bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1495 SPIRVTypeInst ResType,
1496 MachineInstr &
I)
const {
1497 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1498 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1499 for (
const auto &Ex : ExtInsts) {
1500 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1501 uint32_t Opcode = Ex.second;
1505 MachineIRBuilder MIRBuilder(
I);
1508 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1513 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1516 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1519 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1522 .
addImm(
static_cast<uint32_t
>(Ex.first))
1524 .
add(
I.getOperand(2))
1528 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1529 .
addDef(
I.getOperand(1).getReg())
1538bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1539 SPIRVTypeInst ResType,
1540 MachineInstr &
I)
const {
1541 Register CosResVReg =
I.getOperand(1).getReg();
1542 unsigned SrcIdx =
I.getNumExplicitDefs();
1547 MachineIRBuilder MIRBuilder(
I);
1549 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1554 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1557 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1559 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1562 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1564 .
add(
I.getOperand(SrcIdx))
1567 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1575 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1578 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1580 .
add(
I.getOperand(SrcIdx))
1582 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1585 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1587 .
add(
I.getOperand(SrcIdx))
1594bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1595 SPIRVTypeInst ResType,
1597 std::vector<Register> Srcs,
1598 unsigned Opcode)
const {
1599 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1609std::optional<SplitParts> SPIRVInstructionSelector::splitEvenOddLanes(
1610 Register PopCountReg,
unsigned ComponentCount, MachineInstr &
I,
1611 SPIRVTypeInst I32Type)
const {
1614 if (ComponentCount == 1) {
1617 Parts.IsScalar =
true;
1618 Parts.Type = I32Type;
1626 if (!selectOpWithSrcs(Parts.High, I32Type,
I, {PopCountReg, IdxOne},
1627 SPIRV::OpVectorExtractDynamic))
1628 return std::nullopt;
1630 if (!selectOpWithSrcs(Parts.Low, I32Type,
I, {PopCountReg, IdxZero},
1631 SPIRV::OpVectorExtractDynamic))
1632 return std::nullopt;
1636 MachineIRBuilder MIRBuilder(
I);
1637 Parts.IsScalar =
false;
1644 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1645 TII.get(SPIRV::OpVectorShuffle))
1650 for (
unsigned J = 1; J < ComponentCount * 2; J += 2)
1655 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1656 TII.get(SPIRV::OpVectorShuffle))
1661 for (
unsigned J = 0; J < ComponentCount * 2; J += 2)
1669bool SPIRVInstructionSelector::selectPopCount16(
Register ResVReg,
1670 SPIRVTypeInst ResType,
1673 unsigned Opcode)
const {
1674 Register OpReg =
I.getOperand(1).getReg();
1677 MachineIRBuilder MIRBuilder(
I);
1679 SPIRVTypeInst I32VectorType =
1682 bool IsVector = NumElems > 1;
1683 SPIRVTypeInst ExtType = IsVector ? I32VectorType : I32Type;
1686 if (!selectOpWithSrcs(ExtReg, ExtType,
I, {OpReg}, SPIRV::OpUConvert))
1690 if (!selectPopCount32(PopCountReg, ExtType,
I, ExtReg, Opcode))
1693 return selectOpWithSrcs(ResVReg, ResType,
I, {PopCountReg}, ExtOpcode);
1696bool SPIRVInstructionSelector::selectPopCount32(
Register ResVReg,
1697 SPIRVTypeInst ResType,
1700 unsigned Opcode)
const {
1701 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
1704bool SPIRVInstructionSelector::selectPopCount64(
Register ResVReg,
1705 SPIRVTypeInst ResType,
1708 unsigned Opcode)
const {
1710 if (ComponentCount > 2)
1711 return handle64BitOverflow(
1712 ResVReg, ResType,
I, SrcReg, Opcode,
1714 unsigned O) {
return this->selectPopCount64(R,
T,
I, S, O); });
1716 MachineIRBuilder MIRBuilder(
I);
1721 I32Type, 2 * ComponentCount, MIRBuilder,
false);
1725 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
1730 if (!selectPopCount32(Pop32, VecI32Type,
I, Vec32, Opcode))
1734 auto MaybeParts = splitEvenOddLanes(Pop32, ComponentCount,
I, I32Type);
1737 SplitParts &Parts = *MaybeParts;
1740 unsigned OpAdd = Parts.IsScalar ? SPIRV::OpIAddS : SPIRV::OpIAddV;
1742 if (!selectOpWithSrcs(Sum, Parts.Type,
I, {Parts.High, Parts.Low}, OpAdd))
1747 unsigned ConvOp = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
1748 return selectOpWithSrcs(ResVReg, ResType,
I, {Sum}, ConvOp);
1751bool SPIRVInstructionSelector::selectPopCount(
Register ResVReg,
1752 SPIRVTypeInst ResType,
1754 unsigned Opcode)
const {
1759 if (!STI.getTargetTriple().isVulkanOS())
1760 return selectUnOp(ResVReg, ResType,
I, Opcode);
1762 Register OpReg =
I.getOperand(1).getReg();
1765 : SPIRV::OpUConvert;
1769 return selectPopCount16(ResVReg, ResType,
I, ExtOpcode, Opcode);
1771 return selectPopCount32(ResVReg, ResType,
I, OpReg, Opcode);
1773 return selectPopCount64(ResVReg, ResType,
I, OpReg, Opcode);
1775 return diagnoseUnsupported(
I,
"unsupported operand bit width for popcount");
1779bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1780 SPIRVTypeInst ResType,
1782 unsigned Opcode)
const {
1784 Register SrcReg =
I.getOperand(1).getReg();
1789 unsigned DefOpCode = DefIt->getOpcode();
1790 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1793 if (
auto *VRD =
getVRegDef(*MRI, DefIt->getOperand(1).getReg()))
1794 DefOpCode = VRD->getOpcode();
1796 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1797 DefOpCode == TargetOpcode::G_CONSTANT ||
1798 DefOpCode == SPIRV::OpVariable || DefOpCode == SPIRV::OpConstantI) {
1804 uint32_t SpecOpcode = 0;
1806 case SPIRV::OpConvertPtrToU:
1807 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1809 case SPIRV::OpConvertUToPtr:
1810 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1815 TII.get(SPIRV::OpSpecConstantOp))
1825 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1829bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1830 SPIRVTypeInst ResType,
1831 MachineInstr &
I)
const {
1832 Register OpReg =
I.getOperand(1).getReg();
1833 SPIRVTypeInst OpType =
1836 return diagnoseUnsupported(
1837 I,
"incompatible result and operand types in a bitcast");
1838 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1848 if (
MemOp->isVolatile())
1849 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1850 if (
MemOp->isNonTemporal())
1851 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1853 if (!ST->isShader() &&
MemOp->getAlign().value())
1854 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1858 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1859 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1863 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1865 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1869 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1873 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1875 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1887 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1889 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1891 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
1895bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
1896 SPIRVTypeInst ResType,
1897 MachineInstr &
I)
const {
1899 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1904 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
1905 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
1907 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1909 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1913 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1917 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1918 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
1919 I.getDebugLoc(),
I);
1923 MachineIRBuilder MIRBuilder(
I);
1925 if (
I.getNumMemOperands()) {
1926 const MachineMemOperand *MemOp = *
I.memoperands_begin();
1927 if (MemOp->isAtomic())
1928 return selectAtomicLoad(ResVReg, ResType,
I);
1931 auto MIB = MIRBuilder.buildInstr(SPIRV::OpLoad)
1935 if (!
I.getNumMemOperands()) {
1936 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1938 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1947bool SPIRVInstructionSelector::selectAtomicLoad(
Register ResVReg,
1948 SPIRVTypeInst ResType,
1949 MachineInstr &
I)
const {
1950 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
1953 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1956 return diagnoseUnsupported(
I,
1957 "Lowering to SPIR-V of atomic load is only "
1958 "allowed for integer or floating point types");
1960 assert(
I.getNumMemOperands());
1961 const MachineMemOperand &MemOp = **
I.memoperands_begin();
1962 assert(MemOp.isAtomic());
1966 Register ScopeReg = buildI32Constant(Scope,
I);
1972 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
1973 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
1976 MachineIRBuilder MIRBuilder(
I);
1977 auto AtomicLoad = MIRBuilder.buildInstr(SPIRV::OpAtomicLoad)
1983 AtomicLoad.constrainAllUses(
TII,
TRI, RBI);
1987bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
1989 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
1990 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1995 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
1996 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
1998 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2003 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2007 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2008 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2009 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2010 TII.get(SPIRV::OpImageWrite))
2016 if (sampledTypeIsSignedInteger(LLVMHandleType))
2019 BMI.constrainAllUses(
TII,
TRI, RBI);
2024 if (
I.getNumMemOperands()) {
2025 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2026 if (MemOp->isAtomic())
2027 return selectAtomicStore(
I);
2030 MachineIRBuilder MIRBuilder(
I);
2031 auto MIB = MIRBuilder.buildInstr(SPIRV::OpStore).
addUse(Ptr).
addUse(StoreVal);
2032 if (!
I.getNumMemOperands()) {
2033 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2035 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2044bool SPIRVInstructionSelector::selectAtomicStore(MachineInstr &
I)
const {
2045 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2048 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2049 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2054 return diagnoseUnsupported(
I,
2055 "Lowering to SPIR-V of atomic store is only "
2056 "allowed for integer or floating point types");
2058 assert(
I.getNumMemOperands());
2059 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2060 assert(MemOp.isAtomic());
2064 Register ScopeReg = buildI32Constant(Scope,
I);
2070 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2071 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2074 MachineIRBuilder MIRBuilder(
I);
2075 auto AtomicStore = MIRBuilder.buildInstr(SPIRV::OpAtomicStore)
2080 AtomicStore.constrainAllUses(
TII,
TRI, RBI);
2084bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
2085 SPIRVTypeInst ResType,
2086 MachineInstr &
I)
const {
2087 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
2095 const Register PtrsReg =
I.getOperand(2).getReg();
2096 const uint32_t Alignment =
I.getOperand(3).getImm();
2097 const Register MaskReg =
I.getOperand(4).getReg();
2098 const Register PassthruReg =
I.getOperand(5).getReg();
2099 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2103 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
2114bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
2115 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
2122 const Register ValuesReg =
I.getOperand(1).getReg();
2123 const Register PtrsReg =
I.getOperand(2).getReg();
2124 const uint32_t Alignment =
I.getOperand(3).getImm();
2125 const Register MaskReg =
I.getOperand(4).getReg();
2126 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2130 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
2139bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
2140 const Twine &Msg)
const {
2141 const Function &
F =
I.getMF()->getFunction();
2142 F.getContext().diagnose(
2143 DiagnosticInfoUnsupported(
F, Msg,
I.getDebugLoc(),
DS_Error));
2147bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
2148 SPIRVTypeInst ResType,
2149 MachineInstr &
I)
const {
2150 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2151 return diagnoseUnsupported(
2152 I,
"llvm.stacksave intrinsic: this instruction requires the following "
2153 "SPIR-V extension: SPV_INTEL_variable_length_array");
2155 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
2162bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
2163 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2164 return diagnoseUnsupported(
2166 "llvm.stackrestore intrinsic: this instruction requires the following "
2167 "SPIR-V extension: SPV_INTEL_variable_length_array");
2168 if (!
I.getOperand(0).isReg())
2171 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
2172 .
addUse(
I.getOperand(0).getReg())
2178SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
2179 MachineIRBuilder MIRBuilder(
I);
2180 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
2187 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
2191 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2192 Type *ArrTy = ArrayType::get(ValTy, Num);
2194 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
2197 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
2204 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
2207 .
addImm(SPIRV::StorageClass::UniformConstant)
2218bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
2221 Register DstReg =
I.getOperand(0).getReg();
2225 return diagnoseUnsupported(
2226 I,
"OpCopyMemory requires operands to have the same type");
2227 uint64_t CopySize =
getIConstVal(
I.getOperand(2).getReg(), MRI);
2231 return diagnoseUnsupported(
2232 I,
"Unable to determine pointee type size for OpCopyMemory");
2233 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
2234 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
2235 return diagnoseUnsupported(
2236 I,
"OpCopyMemory requires the size to match the pointee type size");
2237 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemory))
2240 if (
I.getNumMemOperands()) {
2241 MachineIRBuilder MIRBuilder(
I);
2248bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &
I,
2251 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemorySized))
2252 .
addUse(
I.getOperand(0).getReg())
2254 .
addUse(
I.getOperand(2).getReg());
2255 if (
I.getNumMemOperands()) {
2256 MachineIRBuilder MIRBuilder(
I);
2263bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
2264 MachineInstr &
I)
const {
2265 Register SrcReg =
I.getOperand(1).getReg();
2266 if (
I.getOpcode() == TargetOpcode::G_MEMSET) {
2267 Register VarReg = getOrCreateMemSetGlobal(
I);
2270 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2272 ValTy,
I, SPIRV::StorageClass::UniformConstant);
2274 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
2278 if (!selectCopyMemory(
I, SrcReg))
2281 if (!selectCopyMemorySized(
I, SrcReg))
2284 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
2285 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
2290bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
2291 SPIRVTypeInst ResType,
2294 unsigned NegateOpcode)
const {
2296 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2299 Register ScopeReg = buildI32Constant(Scope,
I);
2301 Register Ptr =
I.getOperand(1).getReg();
2302 uint32_t ScSem =
static_cast<uint32_t
>(
2306 Register MemSemReg = buildI32Constant(MemSem,
I);
2308 Register ValueReg =
I.getOperand(2).getReg();
2309 if (NegateOpcode != 0) {
2312 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
2317 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2328bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2329 unsigned ArgI =
I.getNumOperands() - 1;
2331 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2332 SPIRVTypeInst SrcType =
2334 if (!SrcType || SrcType->
getOpcode() != SPIRV::OpTypeVector)
2336 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2340 unsigned CurrentIndex = 0;
2341 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2342 Register ResVReg =
I.getOperand(i).getReg();
2345 LLT ResLLT = MRI->
getType(ResVReg);
2351 ResType = ScalarType;
2357 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
2360 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2366 for (
unsigned j = 0;
j < NumElements; ++
j) {
2367 MIB.
addImm(CurrentIndex + j);
2369 CurrentIndex += NumElements;
2373 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2385bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2388 Register MemSemReg = buildI32Constant(MemSem,
I);
2390 uint32_t
Scope =
static_cast<uint32_t
>(
2392 Register ScopeReg = buildI32Constant(Scope,
I);
2394 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2401bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2402 SPIRVTypeInst ResType,
2404 unsigned Opcode)
const {
2405 Type *ResTy =
nullptr;
2408 return diagnoseUnsupported(
2410 "Not enough info to select the arithmetic with overflow instruction");
2412 return diagnoseUnsupported(
I,
2413 "Expect struct type result for the arithmetic "
2414 "with overflow instruction");
2420 MachineIRBuilder MIRBuilder(
I);
2422 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2423 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2429 Register ZeroReg = buildZerosVal(ResType,
I);
2434 if (ResName.
size() > 0)
2439 BuildMI(BB, MIRBuilder.getInsertPt(),
I.getDebugLoc(),
TII.get(Opcode))
2442 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2443 MIB.
addUse(
I.getOperand(i).getReg());
2448 MRI->
setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2449 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2451 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2452 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2459 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2460 .
addDef(
I.getOperand(1).getReg())
2468bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2469 SPIRVTypeInst ResType,
2470 MachineInstr &
I)
const {
2472 "selectAtomicCmpXchg only handles the spv_cmpxchg intrinsic");
2473 Register Ptr =
I.getOperand(2).getReg();
2474 Register ScopeReg =
I.getOperand(5).getReg();
2475 Register MemSemEqReg =
I.getOperand(6).getReg();
2476 Register MemSemNeqReg =
I.getOperand(7).getReg();
2478 Register Val =
I.getOperand(4).getReg();
2482 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2501 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2508 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2520 case SPIRV::StorageClass::DeviceOnlyINTEL:
2521 case SPIRV::StorageClass::HostOnlyINTEL:
2530 bool IsGRef =
false;
2531 bool IsAllowedRefs =
2533 unsigned Opcode = It.getOpcode();
2534 if (Opcode == SPIRV::OpConstantComposite ||
2535 Opcode == SPIRV::OpSpecConstantComposite ||
2536 Opcode == SPIRV::OpVariable ||
2537 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2538 return IsGRef = true;
2539 return Opcode == SPIRV::OpName;
2541 return IsAllowedRefs && IsGRef;
2544Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2545 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2547 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2551SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2553 uint32_t Opcode)
const {
2554 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2555 TII.get(SPIRV::OpSpecConstantOp))
2563SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2564 SPIRVTypeInst SrcPtrTy)
const {
2565 SPIRVTypeInst GenericPtrTy =
2569 SPIRV::StorageClass::Generic),
2571 MachineFunction *MF =
I.getParent()->getParent();
2573 MachineInstrBuilder MIB = buildSpecConstantOp(
2575 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2585bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2586 SPIRVTypeInst ResType,
2587 MachineInstr &
I)
const {
2591 Register SrcPtr =
I.getOperand(1).getReg();
2595 if (SrcPtrTy->
getOpcode() != SPIRV::OpTypePointer ||
2596 ResType->
getOpcode() != SPIRV::OpTypePointer)
2597 return BuildCOPY(ResVReg, SrcPtr,
I);
2607 unsigned SpecOpcode =
2609 ?
static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric)
2612 ? static_cast<uint32_t>(
SPIRV::Opcode::GenericCastToPtr)
2619 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2621 .constrainAllUses(
TII,
TRI, RBI);
2623 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2625 buildSpecConstantOp(
2627 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2628 .constrainAllUses(
TII,
TRI, RBI);
2635 return BuildCOPY(ResVReg, SrcPtr,
I);
2637 if ((SrcSC == SPIRV::StorageClass::Function &&
2638 DstSC == SPIRV::StorageClass::Private) ||
2639 (DstSC == SPIRV::StorageClass::Function &&
2640 SrcSC == SPIRV::StorageClass::Private))
2641 return BuildCOPY(ResVReg, SrcPtr,
I);
2645 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2648 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2651 SPIRVTypeInst GenericPtrTy =
2670 return selectUnOp(ResVReg, ResType,
I,
2671 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
2673 return selectUnOp(ResVReg, ResType,
I,
2674 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
2676 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2678 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2688 return SPIRV::OpFOrdEqual;
2690 return SPIRV::OpFOrdGreaterThanEqual;
2692 return SPIRV::OpFOrdGreaterThan;
2694 return SPIRV::OpFOrdLessThanEqual;
2696 return SPIRV::OpFOrdLessThan;
2698 return SPIRV::OpFOrdNotEqual;
2700 return SPIRV::OpOrdered;
2702 return SPIRV::OpFUnordEqual;
2704 return SPIRV::OpFUnordGreaterThanEqual;
2706 return SPIRV::OpFUnordGreaterThan;
2708 return SPIRV::OpFUnordLessThanEqual;
2710 return SPIRV::OpFUnordLessThan;
2712 return SPIRV::OpFUnordNotEqual;
2714 return SPIRV::OpUnordered;
2724 return SPIRV::OpIEqual;
2726 return SPIRV::OpINotEqual;
2728 return SPIRV::OpSGreaterThanEqual;
2730 return SPIRV::OpSGreaterThan;
2732 return SPIRV::OpSLessThanEqual;
2734 return SPIRV::OpSLessThan;
2736 return SPIRV::OpUGreaterThanEqual;
2738 return SPIRV::OpUGreaterThan;
2740 return SPIRV::OpULessThanEqual;
2742 return SPIRV::OpULessThan;
2751 return SPIRV::OpPtrEqual;
2753 return SPIRV::OpPtrNotEqual;
2764 return SPIRV::OpLogicalEqual;
2766 return SPIRV::OpLogicalNotEqual;
2800bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
2801 SPIRVTypeInst ResType,
2803 unsigned OpAnyOrAll)
const {
2804 assert(
I.getNumOperands() == 3);
2805 assert(
I.getOperand(2).isReg());
2807 Register InputRegister =
I.getOperand(2).getReg();
2810 assert(InputType &&
"VReg has no type assigned");
2813 bool IsVectorTy = InputType->
getOpcode() == SPIRV::OpTypeVector;
2814 if (IsBoolTy && !IsVectorTy) {
2815 assert(ResVReg ==
I.getOperand(0).getReg());
2816 return BuildCOPY(ResVReg, InputRegister,
I);
2820 unsigned SpirvNotEqualId =
2821 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
2823 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
2828 IsBoolTy ? InputRegister
2836 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
2838 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
2855bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
2856 SPIRVTypeInst ResType,
2857 MachineInstr &
I)
const {
2858 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
2861bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
2862 SPIRVTypeInst ResType,
2863 MachineInstr &
I)
const {
2864 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
2868bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
2869 SPIRVTypeInst ResType,
2870 MachineInstr &
I)
const {
2871 assert(
I.getNumOperands() == 4);
2872 assert(
I.getOperand(2).isReg());
2873 assert(
I.getOperand(3).isReg());
2875 [[maybe_unused]] SPIRVTypeInst VecType =
2880 "dot product requires a vector of at least 2 components");
2882 [[maybe_unused]] SPIRVTypeInst EltType =
2891 .
addUse(
I.getOperand(2).getReg())
2892 .
addUse(
I.getOperand(3).getReg())
2897bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
2898 SPIRVTypeInst ResType,
2901 assert(
I.getNumOperands() == 4);
2902 assert(
I.getOperand(2).isReg());
2903 assert(
I.getOperand(3).isReg());
2906 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
2910 .
addUse(
I.getOperand(2).getReg())
2911 .
addUse(
I.getOperand(3).getReg())
2918bool SPIRVInstructionSelector::selectIntegerDotExpansion(
2919 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2920 assert(
I.getNumOperands() == 4);
2921 assert(
I.getOperand(2).isReg());
2922 assert(
I.getOperand(3).isReg());
2926 Register Vec0 =
I.getOperand(2).getReg();
2927 Register Vec1 =
I.getOperand(3).getReg();
2931 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
2940 "dot product requires a vector of at least 2 components");
2943 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2953 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2964 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2976bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
2977 SPIRVTypeInst ResType,
2978 MachineInstr &
I)
const {
2980 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
2983 .
addUse(
I.getOperand(2).getReg())
2988bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
2989 SPIRVTypeInst ResType,
2990 MachineInstr &
I)
const {
2992 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
2995 .
addUse(
I.getOperand(2).getReg())
3000bool SPIRVInstructionSelector::selectOpIsFinite(
Register ResVReg,
3001 SPIRVTypeInst ResType,
3002 MachineInstr &
I)
const {
3004 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsFinite))
3007 .
addUse(
I.getOperand(2).getReg())
3012bool SPIRVInstructionSelector::selectOpIsNormal(
Register ResVReg,
3013 SPIRVTypeInst ResType,
3014 MachineInstr &
I)
const {
3016 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNormal))
3019 .
addUse(
I.getOperand(2).getReg())
3024template <
bool Signed>
3025bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
3026 SPIRVTypeInst ResType,
3027 MachineInstr &
I)
const {
3028 assert(
I.getNumOperands() == 5);
3029 assert(
I.getOperand(2).isReg());
3030 assert(
I.getOperand(3).isReg());
3031 assert(
I.getOperand(4).isReg());
3034 Register Acc =
I.getOperand(2).getReg();
3038 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3040 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
3045 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
3048 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3060template <
bool Signed>
3061bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
3062 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3063 assert(
I.getNumOperands() == 5);
3064 assert(
I.getOperand(2).isReg());
3065 assert(
I.getOperand(3).isReg());
3066 assert(
I.getOperand(4).isReg());
3069 Register Acc =
I.getOperand(2).getReg();
3075 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
3079 for (
unsigned i = 0; i < 4; i++) {
3102 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
3122 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3137bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
3138 SPIRVTypeInst ResType,
3139 MachineInstr &
I)
const {
3140 assert(
I.getNumOperands() == 3);
3141 assert(
I.getOperand(2).isReg());
3143 Register VZero = buildZerosValF(ResType,
I);
3144 Register VOne = buildOnesValF(ResType,
I);
3146 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
3149 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3151 .
addUse(
I.getOperand(2).getReg())
3158bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
3159 SPIRVTypeInst ResType,
3160 MachineInstr &
I)
const {
3161 assert(
I.getNumOperands() == 3);
3162 assert(
I.getOperand(2).isReg());
3164 Register InputRegister =
I.getOperand(2).getReg();
3166 auto &
DL =
I.getDebugLoc();
3169 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3176 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
3178 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
3186 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3191 if (NeedsConversion) {
3192 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
3203bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
3204 SPIRVTypeInst ResType,
3206 unsigned Opcode)
const {
3210 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
3216 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
3217 BMI.addUse(
I.getOperand(J).getReg());
3224bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
3227 bool WithGroupSync)
const {
3229 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
3231 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
3233 assert(((Scope != SPIRV::Scope::Workgroup) ||
3234 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
3235 "Workgroup Scope must set WorkGroupMemory semantic "
3236 "in Barrier instruction");
3238 assert(((Scope != SPIRV::Scope::Device) ||
3239 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
3240 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
3241 "Device Scope must set UniformMemory and ImageMemory semantic "
3242 "in Barrier instruction");
3248 if (WithGroupSync) {
3249 Register ExecReg = buildI32Constant(SPIRV::Scope::Workgroup,
I);
3253 Register ScopeReg = buildI32Constant(Scope,
I);
3254 Register MemSemReg = buildI32Constant(MemSem,
I);
3256 MI.addUse(ScopeReg).addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
3260bool SPIRVInstructionSelector::selectWaveActiveCountBits(
3261 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3266 if (!selectWaveOpInst(BallotReg, BallotType,
I,
3267 SPIRV::OpGroupNonUniformBallot))
3272 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3277 .
addImm(SPIRV::GroupOperation::Reduce)
3284bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
3285 SPIRVTypeInst ResType,
3286 MachineInstr &
I)
const {
3291 Register InputReg =
I.getOperand(2).getReg();
3296 bool IsVector = NumElems > 1;
3309 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
3310 SPIRV::OpGroupNonUniformAllEqual);
3315 ElementResults.
reserve(NumElems);
3317 for (
unsigned Idx = 0; Idx < NumElems; ++Idx) {
3330 ElemInput = Extracted;
3336 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3347 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3358bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3359 SPIRVTypeInst ResType,
3360 MachineInstr &
I)
const {
3362 assert(
I.getNumOperands() == 3);
3364 auto Op =
I.getOperand(2);
3374 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3376 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3377 return diagnoseUnsupported(
I,
"WavePrefixBitCount requires boolean input");
3398 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3402 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3409bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3410 SPIRVTypeInst ResType,
3412 bool IsUnsigned)
const {
3413 return selectWaveReduce(
3414 ResVReg, ResType,
I, IsUnsigned,
3415 [&](
Register InputRegister,
bool IsUnsigned) {
3416 const bool IsFloatTy =
3418 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3419 : SPIRV::OpGroupNonUniformSMax;
3420 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3424bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3425 SPIRVTypeInst ResType,
3427 bool IsUnsigned)
const {
3428 return selectWaveReduce(
3429 ResVReg, ResType,
I, IsUnsigned,
3430 [&](
Register InputRegister,
bool IsUnsigned) {
3431 const bool IsFloatTy =
3433 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3434 : SPIRV::OpGroupNonUniformSMin;
3435 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3439bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3440 SPIRVTypeInst ResType,
3441 MachineInstr &
I)
const {
3442 return selectWaveReduce(ResVReg, ResType,
I,
false,
3443 [&](
Register InputRegister,
bool IsUnsigned) {
3445 InputRegister, SPIRV::OpTypeFloat);
3446 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3447 : SPIRV::OpGroupNonUniformIAdd;
3451bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3452 SPIRVTypeInst ResType,
3453 MachineInstr &
I)
const {
3454 return selectWaveReduce(ResVReg, ResType,
I,
false,
3455 [&](
Register InputRegister,
bool IsUnsigned) {
3457 InputRegister, SPIRV::OpTypeFloat);
3458 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3459 : SPIRV::OpGroupNonUniformIMul;
3463template <
typename PickOpcodeFn>
3464bool SPIRVInstructionSelector::selectWaveReduce(
3465 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3466 PickOpcodeFn &&PickOpcode)
const {
3467 assert(
I.getNumOperands() == 3);
3468 assert(
I.getOperand(2).isReg());
3470 Register InputRegister =
I.getOperand(2).getReg();
3474 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3477 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3483 .
addImm(SPIRV::GroupOperation::Reduce)
3484 .
addUse(
I.getOperand(2).getReg())
3489bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3490 SPIRVTypeInst ResType,
3492 unsigned Opcode)
const {
3493 return selectWaveReduce(
3494 ResVReg, ResType,
I,
false,
3495 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3498bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3499 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3500 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3501 [&](
Register InputRegister,
bool IsUnsigned) {
3503 InputRegister, SPIRV::OpTypeFloat);
3505 ? SPIRV::OpGroupNonUniformFAdd
3506 : SPIRV::OpGroupNonUniformIAdd;
3510bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3511 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3512 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3513 [&](
Register InputRegister,
bool IsUnsigned) {
3515 InputRegister, SPIRV::OpTypeFloat);
3517 ? SPIRV::OpGroupNonUniformFMul
3518 : SPIRV::OpGroupNonUniformIMul;
3522template <
typename PickOpcodeFn>
3523bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3524 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3525 PickOpcodeFn &&PickOpcode)
const {
3526 assert(
I.getNumOperands() == 3);
3527 assert(
I.getOperand(2).isReg());
3529 Register InputRegister =
I.getOperand(2).getReg();
3533 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3536 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3542 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3543 .
addUse(
I.getOperand(2).getReg())
3548bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
3549 SPIRVTypeInst ResType,
3552 assert(
I.getNumOperands() == 3);
3553 assert(
I.getOperand(2).isReg());
3555 Register InputRegister =
I.getOperand(2).getReg();
3561 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
3572bool SPIRVInstructionSelector::selectBitreverse16(
Register ResVReg,
3573 SPIRVTypeInst ResType,
3578 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
3583 : SPIRV::OpUConvert;
3587 ShiftOp = SPIRV::OpShiftRightLogicalV;
3592 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3593 TII.get(SPIRV::OpConstantComposite))
3596 for (
unsigned It = 0; It <
N; ++It)
3600 ShiftConst = CompositeReg;
3605 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
3610 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
3615 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
3620 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
3623bool SPIRVInstructionSelector::handle64BitOverflow(
3625 unsigned int Opcode,
3632 "handle64BitOverflow should only be used for integer types");
3634 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
3636 MachineIRBuilder MIRBuilder(
I);
3638 SPIRVTypeInst I64x2Type =
3640 SPIRVTypeInst Vec2ResType =
3643 std::vector<Register> PartialRegs;
3645 unsigned CurrentComponent = 0;
3646 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
3650 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3651 TII.get(SPIRV::OpVectorShuffle))
3656 .
addImm(CurrentComponent)
3657 .
addImm(CurrentComponent + 1);
3667 PartialRegs.push_back(SubVecReg);
3670 if (CurrentComponent != ComponentCount) {
3676 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
3677 SPIRV::OpVectorExtractDynamic))
3686 PartialRegs.push_back(FinalElemResReg);
3690 return selectOpWithSrcs(ResVReg, ResType,
I, std::move(PartialRegs),
3691 SPIRV::OpCompositeConstruct);
3694bool SPIRVInstructionSelector::selectBitreverse64(
Register ResVReg,
3695 SPIRVTypeInst ResType,
3699 if (ComponentCount > 2)
3700 return handle64BitOverflow(
3701 ResVReg, ResType,
I, SrcReg, SPIRV::OpBitReverse,
3703 unsigned O) {
return this->selectBitreverse64(R,
T,
I, S); });
3705 MachineIRBuilder MIRBuilder(
I);
3709 I32Type, 2 * ComponentCount, MIRBuilder,
false);
3713 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
3718 if (!selectBitreverseNative(Reverse32, VecI32Type,
I, Vec32))
3725 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3726 TII.get(SPIRV::OpVectorShuffle))
3731 for (
unsigned J = 0; J < ComponentCount; ++J) {
3738 return selectOpWithSrcs(ResVReg, ResType,
I, {SwappedVec}, SPIRV::OpBitcast);
3741bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
3742 SPIRVTypeInst ResType,
3746 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
3754bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
3755 SPIRVTypeInst ResType,
3756 MachineInstr &
I)
const {
3757 Register OpReg =
I.getOperand(1).getReg();
3765 return selectBitreverse16(ResVReg, ResType,
I, OpReg);
3767 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
3769 return selectBitreverse64(ResVReg, ResType,
I, OpReg);
3771 return SPIRVInstructionSelector::diagnoseUnsupported(
3772 I,
"G_BITREVERSE only support 16,32,64 bits.");
3776 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
3787 unsigned AndOp = SPIRV::OpBitwiseAndS;
3788 unsigned OrOp = SPIRV::OpBitwiseOrS;
3789 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
3790 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
3792 AndOp = SPIRV::OpBitwiseAndV;
3793 OrOp = SPIRV::OpBitwiseOrV;
3794 ShlOp = SPIRV::OpShiftLeftLogicalV;
3795 ShrOp = SPIRV::OpShiftRightLogicalV;
3801 const unsigned Shift) ->
Register {
3809 Register MaskReg = CreateConst(Mask);
3810 Register ShiftReg = CreateConst(Shift);
3817 if (!selectOpWithSrcs(
T1, ResType,
I, {Input, ShiftReg}, ShrOp) ||
3818 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
3819 !selectOpWithSrcs(T3, ResType,
I, {Input, MaskReg}, AndOp) ||
3820 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
3821 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
3829 uint64_t
Mask = ~0ull;
3830 while ((Shift >>= 1) > 0) {
3837 return BuildCOPY(ResVReg, Result,
I);
3840bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
3841 SPIRVTypeInst ResType,
3842 MachineInstr &
I)
const {
3848 if (!
I.getOperand(0).isReg() || !
I.getOperand(1).isReg())
3850 Register OpReg =
I.getOperand(1).getReg();
3851 if (MachineInstr *Def = MRI->
getVRegDef(OpReg)) {
3852 if (
Def->getOpcode() == TargetOpcode::COPY)
3855 switch (
Def->getOpcode()) {
3856 case SPIRV::ASSIGN_TYPE:
3857 if (MachineInstr *AssignToDef =
3859 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
3860 Reg =
Def->getOperand(2).getReg();
3863 case SPIRV::OpUndef:
3864 Reg =
Def->getOperand(1).getReg();
3867 unsigned DestOpCode;
3869 DestOpCode = SPIRV::OpConstantNull;
3871 DestOpCode = TargetOpcode::COPY;
3874 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
3875 .
addDef(
I.getOperand(0).getReg())
3883bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
3884 SPIRVTypeInst ResType,
3885 MachineInstr &
I)
const {
3887 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3889 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
3893 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
3898 for (
unsigned i =
I.getNumExplicitDefs();
3899 i <
I.getNumExplicitOperands() && IsConst; ++i)
3903 if (!IsConst &&
N < 2)
3904 return diagnoseUnsupported(
3905 I,
"There must be at least two constituent operands in a vector");
3908 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3909 TII.get(IsConst ? SPIRV::OpConstantComposite
3910 : SPIRV::OpCompositeConstruct))
3913 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
3914 MIB.
addUse(
I.getOperand(i).getReg());
3919bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
3920 SPIRVTypeInst ResType,
3921 MachineInstr &
I)
const {
3923 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3925 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
3931 if (!
I.getOperand(
OpIdx).isReg())
3938 if (!IsConst &&
N < 2)
3939 return diagnoseUnsupported(
3940 I,
"There must be at least two constituent operands in a vector");
3943 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3944 TII.get(IsConst ? SPIRV::OpConstantComposite
3945 : SPIRV::OpCompositeConstruct))
3948 for (
unsigned i = 0; i <
N; ++i)
3954bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
3955 SPIRVTypeInst ResType,
3956 MachineInstr &
I)
const {
3961 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
3963 Opcode = SPIRV::OpDemoteToHelperInvocation;
3965 Opcode = SPIRV::OpKill;
3967 if (MachineInstr *NextI =
I.getNextNode()) {
3969 NextI->eraseFromParent();
3979bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
3980 SPIRVTypeInst ResType,
unsigned CmpOpc,
3981 MachineInstr &
I)
const {
3982 Register Cmp0 =
I.getOperand(2).getReg();
3983 Register Cmp1 =
I.getOperand(3).getReg();
3986 "CMP operands should have the same type");
3987 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
3997bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
3998 SPIRVTypeInst ResType,
3999 MachineInstr &
I)
const {
4000 auto Pred =
I.getOperand(1).getPredicate();
4003 Register CmpOperand =
I.getOperand(2).getReg();
4008 Register Op1 =
I.getOperand(3).getReg();
4012 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
4017 I.getOperand(3).setReg(NewOp1);
4023 return selectCmp(ResVReg, ResType, CmpOpc,
I);
4027SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
4028 SPIRVTypeInst ResType)
const {
4030 SPIRVTypeInst SpvI32Ty =
4033 auto ConstInt = ConstantInt::get(LLVMTy, Val);
4040 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4043 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
4046 .
addImm(APInt(32, Val).getZExtValue());
4048 GR.
add(ConstInt,
MI);
4055Register SPIRVInstructionSelector::buildI32ConstantInEntryBlock(
4056 uint32_t Val, MachineInstr &
I, SPIRVTypeInst ResType)
const {
4058 SPIRVTypeInst SpvI32Ty =
4060 auto *ConstInt = ConstantInt::get(LLVMTy, Val);
4065 MachineBasicBlock &EntryBB = *InsertIt->getParent();
4066 MachineInstr *
MI =
nullptr;
4070 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantNull))
4074 uint64_t ImmVal = APInt(32, Val).getZExtValue();
4075 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantI))
4081 GR.
add(ConstInt,
MI);
4086bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
4087 SPIRVTypeInst ResType,
4088 MachineInstr &
I)
const {
4090 return selectCmp(ResVReg, ResType, CmpOp,
I);
4093bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
4094 SPIRVTypeInst ResType,
4095 MachineInstr &
I)
const {
4097 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
4104 if (ResType->
getOpcode() != SPIRV::OpTypeVector &&
4105 ResType->
getOpcode() != SPIRV::OpTypeFloat)
4108 MachineIRBuilder MIRBuilder(
I);
4115 APFloat ConstVal(3.3219280948873623);
4119 APFloat::rmNearestTiesToEven, &LosesInfo);
4123 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
4124 ? SPIRV::OpVectorTimesScalar
4127 if (!selectOpWithSrcs(ArgReg, ResType,
I,
4128 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
4130 if (!selectExtInst(ResVReg, ResType,
I,
4131 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
4141Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
4142 MachineInstr &
I)
const {
4145 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4150bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
4156 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
4164 if (
Def->getOpcode() == SPIRV::OpConstantNull)
4167 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
4168 Def->getOpcode() == SPIRV::OpConstantI)
4181 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
4182 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
4184 Intrinsic::spv_const_composite)) {
4185 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
4186 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
4187 if (!IsZero(
Def->getOperand(i).getReg()))
4196Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
4197 MachineInstr &
I)
const {
4201 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4206Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
4207 MachineInstr &
I)
const {
4211 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4217 SPIRVTypeInst ResType,
4218 MachineInstr &
I)
const {
4222 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4227bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
4228 SPIRVTypeInst ResType,
4229 MachineInstr &
I)
const {
4230 Register SelectFirstArg =
I.getOperand(2).getReg();
4231 Register SelectSecondArg =
I.getOperand(3).getReg();
4240 SPIRV::OpTypeVector;
4247 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
4248 }
else if (IsPtrTy) {
4249 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
4251 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
4254 assert(IsScalarBool &&
"OpSelect with a scalar result requires a scalar "
4255 "boolean condition");
4257 Opcode = SPIRV::OpSelectSFSCond;
4258 }
else if (IsPtrTy) {
4259 Opcode = SPIRV::OpSelectSPSCond;
4261 Opcode = SPIRV::OpSelectSISCond;
4264 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4267 .
addUse(
I.getOperand(1).getReg())
4276bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
4277 SPIRVTypeInst ResType,
4279 MachineInstr &InsertAt,
4280 bool IsSigned)
const {
4282 Register ZeroReg = buildZerosVal(ResType, InsertAt);
4283 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
4284 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
4286 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4298bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
4299 SPIRVTypeInst ResType,
4300 MachineInstr &
I,
bool IsSigned,
4301 unsigned Opcode)
const {
4302 Register SrcReg =
I.getOperand(1).getReg();
4308 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
4313 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I,
false);
4315 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
4318bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
4319 SPIRVTypeInst ResType, MachineInstr &
I,
4320 bool IsSigned)
const {
4321 Register SrcReg =
I.getOperand(1).getReg();
4323 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
4327 if (ResType == SrcType)
4328 return BuildCOPY(ResVReg, SrcReg,
I);
4330 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4331 return selectUnOp(ResVReg, ResType,
I, Opcode);
4334bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
4335 SPIRVTypeInst ResType,
4337 bool IsSigned)
const {
4338 MachineIRBuilder MIRBuilder(
I);
4339 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
4351 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
4354 .
addUse(
I.getOperand(1).getReg())
4355 .
addUse(
I.getOperand(2).getReg())
4360 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
4363 .
addUse(
I.getOperand(1).getReg())
4364 .
addUse(
I.getOperand(2).getReg())
4372 unsigned SelectOpcode =
4373 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
4378 .
addUse(buildOnesVal(
true, ResType,
I))
4379 .
addUse(buildZerosVal(ResType,
I))
4386 .
addUse(buildOnesVal(
false, ResType,
I))
4391bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
4394 SPIRVTypeInst IntTy,
4395 SPIRVTypeInst BoolTy)
const {
4398 bool IsVectorTy = IntTy->
getOpcode() == SPIRV::OpTypeVector;
4399 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
4401 Register One = buildOnesVal(
false, IntTy,
I);
4409 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
4418bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
4419 SPIRVTypeInst ResType,
4420 MachineInstr &
I)
const {
4421 Register IntReg =
I.getOperand(1).getReg();
4424 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
4425 if (ArgType == ResType)
4426 return BuildCOPY(ResVReg, IntReg,
I);
4428 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4429 return selectUnOp(ResVReg, ResType,
I, Opcode);
4432bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
4433 SPIRVTypeInst ResType,
4434 MachineInstr &
I)
const {
4435 unsigned Opcode =
I.getOpcode();
4436 unsigned TpOpcode = ResType->
getOpcode();
4438 if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {
4439 assert(Opcode == TargetOpcode::G_CONSTANT &&
4440 I.getOperand(1).getCImm()->isZero());
4441 MachineBasicBlock &DepMBB =
I.getMF()->front();
4444 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
4451 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
4454bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
4455 SPIRVTypeInst ResType,
4456 MachineInstr &
I)
const {
4457 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4464bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
4465 SPIRVTypeInst ResType,
4466 MachineInstr &
I)
const {
4468 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
4472 .
addUse(
I.getOperand(3).getReg())
4474 .
addUse(
I.getOperand(2).getReg());
4475 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
4481bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
4482 SPIRVTypeInst ResType,
4483 MachineInstr &
I)
const {
4484 Type *MaybeResTy =
nullptr;
4489 "Expected aggregate type for extractv instruction");
4491 SPIRV::AccessQualifier::ReadWrite,
false);
4495 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
4498 .
addUse(
I.getOperand(2).getReg());
4499 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
4505bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
4506 SPIRVTypeInst ResType,
4507 MachineInstr &
I)
const {
4508 if (
getImm(
I.getOperand(4), MRI))
4509 return selectInsertVal(ResVReg, ResType,
I);
4511 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
4514 .
addUse(
I.getOperand(2).getReg())
4515 .
addUse(
I.getOperand(3).getReg())
4516 .
addUse(
I.getOperand(4).getReg())
4521bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
4522 SPIRVTypeInst ResType,
4523 MachineInstr &
I)
const {
4524 if (
getImm(
I.getOperand(3), MRI))
4525 return selectExtractVal(ResVReg, ResType,
I);
4527 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
4530 .
addUse(
I.getOperand(2).getReg())
4531 .
addUse(
I.getOperand(3).getReg())
4536bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
4537 SPIRVTypeInst ResType,
4538 MachineInstr &
I)
const {
4539 const bool IsGEPInBounds =
I.getOperand(2).getImm();
4545 ? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
4546 : SPIRV::OpAccessChain)
4547 : (IsGEPInBounds ?
SPIRV::OpInBoundsPtrAccessChain
4548 :
SPIRV::OpPtrAccessChain);
4550 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4554 .
addUse(
I.getOperand(3).getReg());
4556 (Opcode == SPIRV::OpPtrAccessChain ||
4557 Opcode == SPIRV::OpInBoundsPtrAccessChain ||
4558 (
getImm(
I.getOperand(4), MRI) &&
foldImm(
I.getOperand(4), MRI) == 0)) &&
4559 "Cannot translate GEP to OpAccessChain. First index must be 0.");
4562 const unsigned StartingIndex =
4563 (Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)
4566 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
4567 Res.addUse(
I.getOperand(i).getReg());
4568 Res.constrainAllUses(
TII,
TRI, RBI);
4573bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
4575 unsigned Lim =
I.getNumExplicitOperands();
4576 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
4577 Register OpReg =
I.getOperand(i).getReg();
4578 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
4580 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
4581 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
4582 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
4589 MachineFunction *MF =
I.getMF();
4601 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4602 TII.get(SPIRV::OpSpecConstantOp))
4605 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
4607 GR.
add(OpDefine, MIB);
4613bool SPIRVInstructionSelector::selectDerivativeInst(
4614 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
4615 const unsigned DPdOpCode)
const {
4618 if (!errorIfInstrOutsideShader(
I))
4624 Register SrcReg =
I.getOperand(2).getReg();
4629 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4632 .
addUse(
I.getOperand(2).getReg());
4634 MachineIRBuilder MIRBuilder(
I);
4637 if (componentCount != 1)
4641 const TargetRegisterClass *RegClass = GR.
getRegClass(SrcType);
4645 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4650 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4655 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4663bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
4664 SPIRVTypeInst ResType,
4665 MachineInstr &
I)
const {
4669 case Intrinsic::spv_load:
4670 return selectLoad(ResVReg, ResType,
I);
4671 case Intrinsic::spv_atomic_load:
4672 return selectAtomicLoad(ResVReg, ResType,
I);
4673 case Intrinsic::spv_store:
4674 return selectStore(
I);
4675 case Intrinsic::spv_atomic_store:
4676 return selectAtomicStore(
I);
4677 case Intrinsic::spv_extractv:
4678 return selectExtractVal(ResVReg, ResType,
I);
4679 case Intrinsic::spv_insertv:
4680 return selectInsertVal(ResVReg, ResType,
I);
4681 case Intrinsic::spv_extractelt:
4682 return selectExtractElt(ResVReg, ResType,
I);
4683 case Intrinsic::spv_insertelt:
4684 return selectInsertElt(ResVReg, ResType,
I);
4685 case Intrinsic::spv_gep:
4686 return selectGEP(ResVReg, ResType,
I);
4687 case Intrinsic::spv_bitcast: {
4688 Register OpReg =
I.getOperand(2).getReg();
4689 SPIRVTypeInst OpType =
4693 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
4695 case Intrinsic::spv_unref_global:
4696 case Intrinsic::spv_init_global: {
4697 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
4702 Register GVarVReg =
MI->getOperand(0).getReg();
4703 if (!selectGlobalValue(GVarVReg, *
MI, Init))
4708 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
4710 MI->eraseFromParent();
4714 case Intrinsic::spv_undef: {
4715 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4721 case Intrinsic::spv_named_boolean_spec_constant: {
4722 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
4723 : SPIRV::OpSpecConstantFalse;
4725 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
4726 .
addDef(
I.getOperand(0).getReg())
4729 unsigned SpecId =
I.getOperand(2).getImm();
4731 SPIRV::Decoration::SpecId, {SpecId});
4735 case Intrinsic::spv_const_composite: {
4737 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
4743 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
4745 std::function<bool(
Register)> HasSpecConstOperand =
4755 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
4756 J < Def->getNumExplicitOperands(); ++J) {
4757 if (
Def->getOperand(J).isReg() &&
4758 HasSpecConstOperand(
Def->getOperand(J).getReg()))
4764 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
4765 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
4766 : SPIRV::OpConstantComposite;
4767 unsigned ContinuedOpc = HasSpecConst
4768 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
4769 : SPIRV::OpConstantCompositeContinuedINTEL;
4770 MachineIRBuilder MIR(
I);
4772 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
4774 for (
auto *Instr : Instructions) {
4775 Instr->setDebugLoc(
I.getDebugLoc());
4780 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4787 case Intrinsic::spv_assign_name: {
4788 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
4789 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
4790 for (
unsigned i =
I.getNumExplicitDefs() + 2;
4791 i <
I.getNumExplicitOperands(); ++i) {
4792 MIB.
addImm(
I.getOperand(i).getImm());
4797 case Intrinsic::spv_switch: {
4798 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
4799 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
4800 if (
I.getOperand(i).isReg())
4801 MIB.
addReg(
I.getOperand(i).getReg());
4802 else if (
I.getOperand(i).isCImm())
4803 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
4804 else if (
I.getOperand(i).isMBB())
4805 MIB.
addMBB(
I.getOperand(i).getMBB());
4812 case Intrinsic::spv_loop_merge: {
4813 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
4814 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
4815 if (
I.getOperand(i).isMBB())
4816 MIB.
addMBB(
I.getOperand(i).getMBB());
4823 case Intrinsic::spv_loop_control_intel: {
4825 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
4826 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
4831 case Intrinsic::spv_selection_merge: {
4833 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
4834 assert(
I.getOperand(1).isMBB() &&
4835 "operand 1 to spv_selection_merge must be a basic block");
4836 MIB.
addMBB(
I.getOperand(1).getMBB());
4837 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
4841 case Intrinsic::spv_cmpxchg:
4842 return selectAtomicCmpXchg(ResVReg, ResType,
I);
4843 case Intrinsic::spv_unreachable:
4844 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
4847 case Intrinsic::spv_abort:
4848 return selectAbort(
I);
4849 case Intrinsic::spv_alloca:
4850 return selectFrameIndex(ResVReg, ResType,
I);
4851 case Intrinsic::spv_alloca_array:
4852 return selectAllocaArray(ResVReg, ResType,
I);
4853 case Intrinsic::spv_assume:
4855 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
4856 .
addUse(
I.getOperand(1).getReg())
4861 case Intrinsic::spv_expect:
4863 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
4866 .
addUse(
I.getOperand(2).getReg())
4867 .
addUse(
I.getOperand(3).getReg())
4872 case Intrinsic::arithmetic_fence:
4873 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
4874 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
4877 .
addUse(
I.getOperand(2).getReg())
4881 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
4883 case Intrinsic::spv_thread_id:
4889 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
4891 case Intrinsic::spv_thread_id_in_group:
4897 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
4899 case Intrinsic::spv_group_id:
4905 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
4907 case Intrinsic::spv_flattened_thread_id_in_group:
4914 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
4916 case Intrinsic::spv_workgroup_size:
4917 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
4919 case Intrinsic::spv_global_size:
4920 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
4922 case Intrinsic::spv_global_offset:
4923 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
4925 case Intrinsic::spv_num_workgroups:
4926 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
4928 case Intrinsic::spv_subgroup_size:
4929 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
4931 case Intrinsic::spv_num_subgroups:
4932 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
4934 case Intrinsic::spv_subgroup_id:
4935 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
4936 case Intrinsic::spv_subgroup_local_invocation_id:
4937 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
4938 ResVReg, ResType,
I);
4939 case Intrinsic::spv_subgroup_max_size:
4940 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
4942 case Intrinsic::spv_fdot:
4943 return selectFloatDot(ResVReg, ResType,
I);
4944 case Intrinsic::spv_udot:
4945 case Intrinsic::spv_sdot:
4946 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4948 return selectIntegerDot(ResVReg, ResType,
I,
4949 IID == Intrinsic::spv_sdot);
4950 return selectIntegerDotExpansion(ResVReg, ResType,
I);
4951 case Intrinsic::spv_dot4add_i8packed:
4952 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4954 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
4955 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
4956 case Intrinsic::spv_dot4add_u8packed:
4957 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4959 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
4960 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
4961 case Intrinsic::spv_all:
4962 return selectAll(ResVReg, ResType,
I);
4963 case Intrinsic::spv_any:
4964 return selectAny(ResVReg, ResType,
I);
4965 case Intrinsic::spv_cross:
4966 return selectExtInst(ResVReg, ResType,
I, CL::cross, GL::Cross);
4967 case Intrinsic::spv_distance:
4968 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
4969 case Intrinsic::spv_lerp:
4970 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
4971 case Intrinsic::spv_length:
4972 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
4973 case Intrinsic::spv_degrees:
4974 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
4975 case Intrinsic::spv_faceforward:
4976 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
4977 case Intrinsic::spv_frac:
4978 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
4979 case Intrinsic::spv_isinf:
4980 return selectOpIsInf(ResVReg, ResType,
I);
4981 case Intrinsic::spv_isnan:
4982 return selectOpIsNan(ResVReg, ResType,
I);
4983 case Intrinsic::spv_isfinite:
4984 return selectOpIsFinite(ResVReg, ResType,
I);
4985 case Intrinsic::spv_isnormal:
4986 return selectOpIsNormal(ResVReg, ResType,
I);
4987 case Intrinsic::spv_normalize:
4988 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
4989 case Intrinsic::spv_refract:
4990 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
4991 case Intrinsic::spv_reflect:
4992 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
4993 case Intrinsic::spv_rsqrt:
4994 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
4995 case Intrinsic::spv_sign:
4996 return selectSign(ResVReg, ResType,
I);
4997 case Intrinsic::spv_smoothstep:
4998 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
4999 case Intrinsic::spv_firstbituhigh:
5000 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
5001 case Intrinsic::spv_firstbitshigh:
5002 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
5003 case Intrinsic::spv_firstbitlow:
5004 return selectFirstBitLow(ResVReg, ResType,
I);
5005 case Intrinsic::spv_all_memory_barrier:
5006 return selectBarrierInst(
I, SPIRV::Scope::Device,
5007 SPIRV::MemorySemantics::UniformMemory |
5008 SPIRV::MemorySemantics::ImageMemory |
5009 SPIRV::MemorySemantics::WorkgroupMemory,
5011 case Intrinsic::spv_all_memory_barrier_with_group_sync:
5012 return selectBarrierInst(
I, SPIRV::Scope::Device,
5013 SPIRV::MemorySemantics::UniformMemory |
5014 SPIRV::MemorySemantics::ImageMemory |
5015 SPIRV::MemorySemantics::WorkgroupMemory,
5017 case Intrinsic::spv_device_memory_barrier:
5018 return selectBarrierInst(
I, SPIRV::Scope::Device,
5019 SPIRV::MemorySemantics::UniformMemory |
5020 SPIRV::MemorySemantics::ImageMemory,
5022 case Intrinsic::spv_device_memory_barrier_with_group_sync:
5023 return selectBarrierInst(
I, SPIRV::Scope::Device,
5024 SPIRV::MemorySemantics::UniformMemory |
5025 SPIRV::MemorySemantics::ImageMemory,
5027 case Intrinsic::spv_group_memory_barrier:
5028 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5029 SPIRV::MemorySemantics::WorkgroupMemory,
5031 case Intrinsic::spv_group_memory_barrier_with_group_sync:
5032 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5033 SPIRV::MemorySemantics::WorkgroupMemory,
5035 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
5036 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
5037 SPIRV::StorageClass::StorageClass ResSC =
5040 return diagnoseUnsupported(
I,
"The target storage class is not castable "
5041 "from the Generic storage class");
5042 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
5050 case Intrinsic::spv_lifetime_start:
5051 case Intrinsic::spv_lifetime_end: {
5052 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
5053 : SPIRV::OpLifetimeStop;
5054 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
5055 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
5064 case Intrinsic::spv_saturate:
5065 return selectSaturate(ResVReg, ResType,
I);
5066 case Intrinsic::spv_nclamp:
5067 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
5068 case Intrinsic::spv_uclamp:
5069 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
5070 case Intrinsic::spv_sclamp:
5071 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
5072 case Intrinsic::spv_subgroup_prefix_bit_count:
5073 return selectWavePrefixBitCount(ResVReg, ResType,
I);
5074 case Intrinsic::spv_wave_active_countbits:
5075 return selectWaveActiveCountBits(ResVReg, ResType,
I);
5076 case Intrinsic::spv_wave_all_equal:
5077 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
5078 case Intrinsic::spv_wave_all:
5079 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
5080 case Intrinsic::spv_wave_any:
5081 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
5082 case Intrinsic::spv_subgroup_ballot:
5083 return selectWaveOpInst(ResVReg, ResType,
I,
5084 SPIRV::OpGroupNonUniformBallot);
5085 case Intrinsic::spv_wave_is_first_lane:
5086 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
5087 case Intrinsic::spv_wave_reduce_or:
5088 return selectWaveReduceOp(ResVReg, ResType,
I,
5089 SPIRV::OpGroupNonUniformBitwiseOr);
5090 case Intrinsic::spv_wave_reduce_xor:
5091 return selectWaveReduceOp(ResVReg, ResType,
I,
5092 SPIRV::OpGroupNonUniformBitwiseXor);
5093 case Intrinsic::spv_wave_reduce_and:
5094 return selectWaveReduceOp(ResVReg, ResType,
I,
5095 SPIRV::OpGroupNonUniformBitwiseAnd);
5096 case Intrinsic::spv_wave_reduce_umax:
5097 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
5098 case Intrinsic::spv_wave_reduce_max:
5099 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
5100 case Intrinsic::spv_wave_reduce_umin:
5101 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
5102 case Intrinsic::spv_wave_reduce_min:
5103 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
5104 case Intrinsic::spv_wave_reduce_sum:
5105 return selectWaveReduceSum(ResVReg, ResType,
I);
5106 case Intrinsic::spv_wave_product:
5107 return selectWaveReduceProduct(ResVReg, ResType,
I);
5108 case Intrinsic::spv_wave_readlane:
5109 return selectWaveOpInst(ResVReg, ResType,
I,
5110 SPIRV::OpGroupNonUniformShuffle);
5111 case Intrinsic::spv_wave_prefix_sum:
5112 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
5113 case Intrinsic::spv_wave_prefix_product:
5114 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
5115 case Intrinsic::spv_quad_read_across_x: {
5116 return selectQuadSwap(ResVReg, ResType,
I, 0);
5118 case Intrinsic::spv_quad_read_across_y: {
5119 return selectQuadSwap(ResVReg, ResType,
I, 1);
5121 case Intrinsic::spv_step:
5122 return selectExtInst(ResVReg, ResType,
I, CL::step, GL::Step);
5123 case Intrinsic::spv_radians:
5124 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
5128 case Intrinsic::instrprof_increment:
5129 case Intrinsic::instrprof_increment_step:
5130 case Intrinsic::instrprof_value_profile:
5133 case Intrinsic::spv_value_md:
5135 case Intrinsic::spv_resource_handlefrombinding: {
5136 return selectHandleFromBinding(ResVReg, ResType,
I);
5138 case Intrinsic::spv_resource_counterhandlefrombinding:
5139 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
5140 case Intrinsic::spv_resource_updatecounter:
5141 return selectUpdateCounter(ResVReg, ResType,
I);
5142 case Intrinsic::spv_resource_store_typedbuffer: {
5143 return selectImageWriteIntrinsic(
I);
5145 case Intrinsic::spv_resource_load_typedbuffer: {
5146 return selectReadImageIntrinsic(ResVReg, ResType,
I);
5148 case Intrinsic::spv_resource_load_level: {
5149 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
5151 case Intrinsic::spv_resource_getdimensions_x:
5152 case Intrinsic::spv_resource_getdimensions_xy:
5153 case Intrinsic::spv_resource_getdimensions_xyz: {
5154 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
5156 case Intrinsic::spv_resource_getdimensions_levels_x:
5157 case Intrinsic::spv_resource_getdimensions_levels_xy:
5158 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
5159 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
5161 case Intrinsic::spv_resource_getdimensions_ms_xy:
5162 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
5163 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
5165 case Intrinsic::spv_resource_calculate_lod:
5166 case Intrinsic::spv_resource_calculate_lod_unclamped:
5167 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
5168 case Intrinsic::spv_resource_sample:
5169 case Intrinsic::spv_resource_sample_clamp:
5170 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
5171 case Intrinsic::spv_resource_samplebias:
5172 case Intrinsic::spv_resource_samplebias_clamp:
5173 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
5174 case Intrinsic::spv_resource_samplegrad:
5175 case Intrinsic::spv_resource_samplegrad_clamp:
5176 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
5177 case Intrinsic::spv_resource_samplelevel:
5178 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
5179 case Intrinsic::spv_resource_samplecmp:
5180 case Intrinsic::spv_resource_samplecmp_clamp:
5181 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
5182 case Intrinsic::spv_resource_samplecmplevelzero:
5183 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
5184 case Intrinsic::spv_resource_gather:
5185 case Intrinsic::spv_resource_gather_cmp:
5186 return selectGatherIntrinsic(ResVReg, ResType,
I);
5187 case Intrinsic::spv_resource_getbasepointer:
5188 case Intrinsic::spv_resource_getpointer: {
5189 return selectResourceGetPointer(ResVReg, ResType,
I);
5191 case Intrinsic::spv_pushconstant_getpointer: {
5192 return selectPushConstantGetPointer(ResVReg, ResType,
I);
5194 case Intrinsic::spv_discard: {
5195 return selectDiscard(ResVReg, ResType,
I);
5197 case Intrinsic::spv_resource_nonuniformindex: {
5198 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
5200 case Intrinsic::spv_unpackhalf2x16: {
5201 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
5203 case Intrinsic::spv_packhalf2x16: {
5204 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
5206 case Intrinsic::spv_ddx:
5207 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
5208 case Intrinsic::spv_ddy:
5209 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
5210 case Intrinsic::spv_ddx_coarse:
5211 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
5212 case Intrinsic::spv_ddy_coarse:
5213 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
5214 case Intrinsic::spv_ddx_fine:
5215 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
5216 case Intrinsic::spv_ddy_fine:
5217 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
5218 case Intrinsic::spv_fwidth:
5219 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
5220 case Intrinsic::spv_masked_gather:
5221 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5222 return selectMaskedGather(ResVReg, ResType,
I);
5223 return diagnoseUnsupported(
5224 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
5225 case Intrinsic::spv_masked_scatter:
5226 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5227 return selectMaskedScatter(
I);
5228 return diagnoseUnsupported(
5229 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
5230 case Intrinsic::returnaddress:
5231 case Intrinsic::frameaddress: {
5233 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5240 return diagnoseUnsupported(
I,
"intrinsic selection not implemented.");
5245bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
5246 SPIRVTypeInst ResType,
5247 MachineInstr &
I)
const {
5250 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
5257bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
5258 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5260 assert(Intr.getIntrinsicID() ==
5261 Intrinsic::spv_resource_counterhandlefrombinding);
5264 Register MainHandleReg = Intr.getOperand(2).getReg();
5266 assert(MainHandleDef->getIntrinsicID() ==
5267 Intrinsic::spv_resource_handlefrombinding);
5271 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
5272 Register IndexReg = MainHandleDef->getOperand(5).getReg();
5273 std::string CounterName =
5278 MachineIRBuilder MIRBuilder(
I);
5280 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
5282 ArraySize, IndexReg, CounterName, MIRBuilder);
5284 return BuildCOPY(ResVReg, CounterVarReg,
I);
5287bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
5288 SPIRVTypeInst ResType,
5289 MachineInstr &
I)
const {
5291 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
5293 Register CounterHandleReg = Intr.getOperand(2).getReg();
5294 Register IncrReg = Intr.getOperand(3).getReg();
5301 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
5302 assert(CounterVarPointeeType &&
5303 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
5304 "Counter variable must be a struct");
5306 SPIRV::StorageClass::StorageBuffer &&
5307 "Counter variable must be in the storage buffer storage class");
5309 "Counter variable must have exactly 1 member in the struct");
5310 const SPIRVTypeInst MemberType =
5313 "Counter variable struct must have a single i32 member");
5317 MachineIRBuilder MIRBuilder(
I);
5319 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
5322 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
5328 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5331 .
addUse(CounterHandleReg)
5338 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
5341 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
5344 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
5353 return BuildCOPY(ResVReg, AtomicRes,
I);
5361 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
5369bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
5370 SPIRVTypeInst ResType,
5371 MachineInstr &
I)
const {
5379 Register ImageReg =
I.getOperand(2).getReg();
5387 Register IdxReg =
I.getOperand(3).getReg();
5389 MachineInstr &Pos =
I;
5391 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
5395bool SPIRVInstructionSelector::generateSampleImage(
5398 DebugLoc Loc, MachineInstr &Pos)
const {
5409 if (!loadHandleBeforePosition(NewSamplerReg,
5415 MachineIRBuilder MIRBuilder(Pos);
5428 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
5429 ImOps.Lod.has_value();
5430 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
5431 : SPIRV::OpImageSampleImplicitLod;
5433 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
5434 : SPIRV::OpImageSampleDrefImplicitLod;
5443 MIB.
addUse(*ImOps.Compare);
5445 uint32_t ImageOperands = 0;
5447 ImageOperands |= SPIRV::ImageOperand::Bias;
5449 ImageOperands |= SPIRV::ImageOperand::Lod;
5450 if (ImOps.GradX && ImOps.GradY)
5451 ImageOperands |= SPIRV::ImageOperand::Grad;
5452 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
5454 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
5457 "Non-constant offsets are not supported in sample instructions.");
5461 ImageOperands |= SPIRV::ImageOperand::MinLod;
5463 if (ImageOperands != 0) {
5464 MIB.
addImm(ImageOperands);
5465 if (ImageOperands & SPIRV::ImageOperand::Bias)
5467 if (ImageOperands & SPIRV::ImageOperand::Lod)
5469 if (ImageOperands & SPIRV::ImageOperand::Grad) {
5470 MIB.
addUse(*ImOps.GradX);
5471 MIB.
addUse(*ImOps.GradY);
5474 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
5475 MIB.
addUse(*ImOps.Offset);
5476 if (ImageOperands & SPIRV::ImageOperand::MinLod)
5477 MIB.
addUse(*ImOps.MinLod);
5484bool SPIRVInstructionSelector::selectImageQuerySize(
5486 std::optional<Register> LodReg)
const {
5488 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
5491 "ImageReg is not an image type.");
5493 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
5495 unsigned NumComponents = 0;
5497 case SPIRV::Dim::DIM_1D:
5498 case SPIRV::Dim::DIM_Buffer:
5499 NumComponents =
IsArray ? 2 : 1;
5501 case SPIRV::Dim::DIM_2D:
5502 case SPIRV::Dim::DIM_Cube:
5503 case SPIRV::Dim::DIM_Rect:
5504 NumComponents =
IsArray ? 3 : 2;
5506 case SPIRV::Dim::DIM_3D:
5510 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
5515 SPIRVTypeInst ResType =
5520 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5530bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
5531 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5532 Register ImageReg =
I.getOperand(2).getReg();
5539 return selectImageQuerySize(NewImageReg, ResVReg,
I);
5542bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
5543 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5544 Register ImageReg =
I.getOperand(2).getReg();
5553 Register LodReg =
I.getOperand(3).getReg();
5556 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
5558 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
5565 TII.get(SPIRV::OpImageQueryLevels))
5572 TII.get(SPIRV::OpCompositeConstruct))
5582bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
5583 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5584 Register ImageReg =
I.getOperand(2).getReg();
5595 "OpImageQuerySamples requires a multisampled image");
5597 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
5605 TII.get(SPIRV::OpImageQuerySamples))
5612 TII.get(SPIRV::OpCompositeConstruct))
5622bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
5623 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5624 Register ImageReg =
I.getOperand(2).getReg();
5625 Register SamplerReg =
I.getOperand(3).getReg();
5626 Register CoordinateReg =
I.getOperand(4).getReg();
5642 if (!loadHandleBeforePosition(
5647 MachineIRBuilder MIRBuilder(
I);
5653 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
5663 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
5670 unsigned ExtractedIndex =
5672 Intrinsic::spv_resource_calculate_lod_unclamped
5676 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5677 TII.get(SPIRV::OpCompositeExtract))
5687bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
5688 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5689 Register ImageReg =
I.getOperand(2).getReg();
5690 Register SamplerReg =
I.getOperand(3).getReg();
5691 Register CoordinateReg =
I.getOperand(4).getReg();
5692 ImageOperands ImOps;
5693 if (
I.getNumOperands() > 5)
5694 ImOps.Offset =
I.getOperand(5).getReg();
5695 if (
I.getNumOperands() > 6)
5696 ImOps.MinLod =
I.getOperand(6).getReg();
5697 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5698 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5701bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
5702 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5703 Register ImageReg =
I.getOperand(2).getReg();
5704 Register SamplerReg =
I.getOperand(3).getReg();
5705 Register CoordinateReg =
I.getOperand(4).getReg();
5706 ImageOperands ImOps;
5707 ImOps.Bias =
I.getOperand(5).getReg();
5708 if (
I.getNumOperands() > 6)
5709 ImOps.Offset =
I.getOperand(6).getReg();
5710 if (
I.getNumOperands() > 7)
5711 ImOps.MinLod =
I.getOperand(7).getReg();
5712 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5713 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5716bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
5717 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5718 Register ImageReg =
I.getOperand(2).getReg();
5719 Register SamplerReg =
I.getOperand(3).getReg();
5720 Register CoordinateReg =
I.getOperand(4).getReg();
5721 ImageOperands ImOps;
5722 ImOps.GradX =
I.getOperand(5).getReg();
5723 ImOps.GradY =
I.getOperand(6).getReg();
5724 if (
I.getNumOperands() > 7)
5725 ImOps.Offset =
I.getOperand(7).getReg();
5726 if (
I.getNumOperands() > 8)
5727 ImOps.MinLod =
I.getOperand(8).getReg();
5728 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5729 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5732bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
5733 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5734 Register ImageReg =
I.getOperand(2).getReg();
5735 Register SamplerReg =
I.getOperand(3).getReg();
5736 Register CoordinateReg =
I.getOperand(4).getReg();
5737 ImageOperands ImOps;
5738 ImOps.Lod =
I.getOperand(5).getReg();
5739 if (
I.getNumOperands() > 6)
5740 ImOps.Offset =
I.getOperand(6).getReg();
5741 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5742 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5745bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
5746 SPIRVTypeInst ResType,
5747 MachineInstr &
I)
const {
5748 Register ImageReg =
I.getOperand(2).getReg();
5749 Register SamplerReg =
I.getOperand(3).getReg();
5750 Register CoordinateReg =
I.getOperand(4).getReg();
5751 ImageOperands ImOps;
5752 ImOps.Compare =
I.getOperand(5).getReg();
5753 if (
I.getNumOperands() > 6)
5754 ImOps.Offset =
I.getOperand(6).getReg();
5755 if (
I.getNumOperands() > 7)
5756 ImOps.MinLod =
I.getOperand(7).getReg();
5757 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5758 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5761bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
5762 SPIRVTypeInst ResType,
5763 MachineInstr &
I)
const {
5764 Register ImageReg =
I.getOperand(2).getReg();
5765 Register CoordinateReg =
I.getOperand(3).getReg();
5766 Register LodReg =
I.getOperand(4).getReg();
5768 ImageOperands ImOps;
5770 if (
I.getNumOperands() > 5)
5771 ImOps.Offset =
I.getOperand(5).getReg();
5783 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
5784 I.getDebugLoc(),
I, &ImOps);
5787bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
5788 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5789 Register ImageReg =
I.getOperand(2).getReg();
5790 Register SamplerReg =
I.getOperand(3).getReg();
5791 Register CoordinateReg =
I.getOperand(4).getReg();
5792 ImageOperands ImOps;
5793 ImOps.Compare =
I.getOperand(5).getReg();
5794 if (
I.getNumOperands() > 6)
5795 ImOps.Offset =
I.getOperand(6).getReg();
5798 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5799 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5802bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
5803 SPIRVTypeInst ResType,
5804 MachineInstr &
I)
const {
5805 Register ImageReg =
I.getOperand(2).getReg();
5806 Register SamplerReg =
I.getOperand(3).getReg();
5807 Register CoordinateReg =
I.getOperand(4).getReg();
5810 "ImageReg is not an image type.");
5815 ComponentOrCompareReg =
I.getOperand(5).getReg();
5816 OffsetReg =
I.getOperand(6).getReg();
5819 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
5823 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
5824 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
5825 Dim != SPIRV::Dim::DIM_Rect) {
5827 "Gather operations are only supported for 2D, Cube, and Rect images.");
5834 if (!loadHandleBeforePosition(
5839 MachineIRBuilder MIRBuilder(
I);
5840 SPIRVTypeInst SampledImageType =
5845 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
5853 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
5855 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
5857 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5862 .
addUse(ComponentOrCompareReg);
5864 uint32_t ImageOperands = 0;
5865 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
5866 if (Dim == SPIRV::Dim::DIM_Cube) {
5868 "Gather operations with offset are not supported for Cube images.");
5872 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
5874 ImageOperands |= SPIRV::ImageOperand::Offset;
5878 if (ImageOperands != 0) {
5879 MIB.
addImm(ImageOperands);
5881 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
5889bool SPIRVInstructionSelector::generateImageReadOrFetch(
5892 const ImageOperands *ImOps)
const {
5895 "ImageReg is not an image type.");
5897 bool IsSignedInteger =
5902 bool IsFetch = (SampledOp.getImm() == 1);
5904 auto AddOperands = [&](MachineInstrBuilder &MIB) {
5905 uint32_t ImageOperandsMask = 0;
5906 if (IsSignedInteger)
5907 ImageOperandsMask |= 0x1000;
5909 if (IsFetch && ImOps) {
5911 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
5912 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
5914 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
5916 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
5920 if (ImageOperandsMask != 0) {
5921 MIB.
addImm(ImageOperandsMask);
5922 if (IsFetch && ImOps) {
5925 if (ImOps->Offset &&
5926 (ImageOperandsMask &
5927 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
5928 MIB.
addUse(*ImOps->Offset);
5934 if (ResultSize == 4) {
5937 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
5944 BMI.constrainAllUses(
TII,
TRI, RBI);
5948 SPIRVTypeInst ReadType = widenTypeToVec4(ResType, Pos);
5952 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
5958 BMI.constrainAllUses(
TII,
TRI, RBI);
5960 if (ResultSize == 1) {
5969 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
5972bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
5973 SPIRVTypeInst ResType,
5974 MachineInstr &
I)
const {
5975 Register ResourcePtr =
I.getOperand(2).getReg();
5977 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
5986 MachineIRBuilder MIRBuilder(
I);
5991 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5997 if (
I.getNumExplicitOperands() > 3) {
5998 Register IndexReg =
I.getOperand(3).getReg();
6005bool SPIRVInstructionSelector::selectPushConstantGetPointer(
6006 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6011bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
6012 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6013 Register ObjReg =
I.getOperand(2).getReg();
6014 if (!BuildCOPY(ResVReg, ObjReg,
I))
6024 decorateUsesAsNonUniform(ResVReg);
6028void SPIRVInstructionSelector::decorateUsesAsNonUniform(
6031 while (WorkList.
size() > 0) {
6035 bool IsDecorated =
false;
6037 if (
Use.getOpcode() == SPIRV::OpDecorate &&
6038 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
6044 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
6046 if (ResultReg == CurrentReg)
6054 SPIRV::Decoration::NonUniformEXT, {});
6059bool SPIRVInstructionSelector::extractSubvector(
6061 MachineInstr &InsertionPoint)
const {
6063 [[maybe_unused]] uint64_t InputSize =
6066 assert(InputSize > 1 &&
"The input must be a vector.");
6067 assert(ResultSize > 1 &&
"The result must be a vector.");
6068 assert(ResultSize < InputSize &&
6069 "Cannot extract more element than there are in the input.");
6072 const TargetRegisterClass *ScalarRegClass = GR.
getRegClass(ScalarType);
6073 for (uint64_t
I = 0;
I < ResultSize;
I++) {
6076 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6085 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
6087 TII.get(SPIRV::OpCompositeConstruct))
6091 for (
Register ComponentReg : ComponentRegisters)
6092 MIB.
addUse(ComponentReg);
6097bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
6098 MachineInstr &
I)
const {
6105 Register ImageReg =
I.getOperand(1).getReg();
6113 Register CoordinateReg =
I.getOperand(2).getReg();
6114 Register DataReg =
I.getOperand(3).getReg();
6117 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
6125Register SPIRVInstructionSelector::buildPointerToResource(
6126 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
6127 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
6128 StringRef Name, MachineIRBuilder MIRBuilder)
const {
6130 if (ArraySize == 1) {
6131 SPIRVTypeInst PtrType =
6134 "SpirvResType did not have an explicit layout.");
6139 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
6140 SPIRVTypeInst VarPointerType =
6143 VarPointerType, Set,
Binding, Name, MIRBuilder);
6145 SPIRVTypeInst ResPointerType =
6158bool SPIRVInstructionSelector::selectFirstBitSet16(
6159 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
6160 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
6162 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
6166 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
6169bool SPIRVInstructionSelector::selectFirstBitSet32(
6171 unsigned BitSetOpcode)
const {
6172 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6175 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6182bool SPIRVInstructionSelector::selectFirstBitSet64(
6184 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
6197 if (ComponentCount > 2) {
6198 auto Func = [
this, SwapPrimarySide](
Register ResVReg, SPIRVTypeInst ResType,
6200 unsigned Opcode) ->
bool {
6201 return this->selectFirstBitSet64(ResVReg, ResType,
I, SrcReg, Opcode,
6205 return handle64BitOverflow(ResVReg, ResType,
I, SrcReg, BitSetOpcode, Func);
6209 MachineIRBuilder MIRBuilder(
I);
6211 BaseType, 2 * ComponentCount, MIRBuilder,
false);
6215 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
6221 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
6228 bool IsScalarRes = ResType->
getOpcode() != SPIRV::OpTypeVector;
6231 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
6232 SPIRV::OpVectorExtractDynamic))
6234 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
6235 SPIRV::OpVectorExtractDynamic))
6239 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6240 TII.get(SPIRV::OpVectorShuffle))
6248 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
6254 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6255 TII.get(SPIRV::OpVectorShuffle))
6263 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
6283 SelectOp = SPIRV::OpSelectSISCond;
6284 AddOp = SPIRV::OpIAddS;
6292 SelectOp = SPIRV::OpSelectVIVCond;
6293 AddOp = SPIRV::OpIAddV;
6299 Register RegSecondaryOffset = Reg0;
6303 if (SwapPrimarySide) {
6304 PrimaryReg = LowReg;
6305 SecondaryReg = HighReg;
6306 RegPrimaryOffset = Reg0;
6307 RegSecondaryOffset = Reg32;
6312 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
6313 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
6318 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
6319 SPIRV::OpINotEqual))
6326 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
6327 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
6332 if (SwapPrimarySide) {
6334 if (!selectOpWithSrcs(RegAdd, ResType,
I,
6335 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
6346 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
6347 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
6352 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
6353 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
6356 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
6360bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
6361 SPIRVTypeInst ResType,
6363 bool IsSigned)
const {
6365 Register OpReg =
I.getOperand(2).getReg();
6368 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
6369 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
6373 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6375 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6377 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6380 return diagnoseUnsupported(
6382 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
6386bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
6387 SPIRVTypeInst ResType,
6388 MachineInstr &
I)
const {
6390 Register OpReg =
I.getOperand(2).getReg();
6395 unsigned ExtendOpcode = SPIRV::OpUConvert;
6396 unsigned BitSetOpcode = GL::FindILsb;
6400 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6402 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6404 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6407 return diagnoseUnsupported(
I,
6408 "spv_firstbitlow only supports 16,32,64 bits.");
6412bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
6413 SPIRVTypeInst ResType,
6414 MachineInstr &
I)
const {
6418 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
6421 .
addUse(
I.getOperand(2).getReg())
6424 unsigned Alignment =
I.getOperand(3).getImm();
6438 while (!Worklist.
empty()) {
6440 switch (
T->getOpcode()) {
6441 case SPIRV::OpTypeInt:
6442 case SPIRV::OpTypeFloat:
6443 case SPIRV::OpTypePointer:
6445 case SPIRV::OpTypeVector:
6446 case SPIRV::OpTypeMatrix:
6447 case SPIRV::OpTypeArray: {
6448 Register OperandReg =
T->getOperand(1).getReg();
6452 case SPIRV::OpTypeStruct:
6453 for (
unsigned Idx = 1,
E =
T->getNumOperands(); Idx <
E; ++Idx) {
6454 Register OperandReg =
T->getOperand(Idx).getReg();
6466bool SPIRVInstructionSelector::selectAbort(MachineInstr &
I)
const {
6467 assert(
I.getNumExplicitOperands() == 2);
6469 Register MsgReg =
I.getOperand(1).getReg();
6471 assert(MsgType &&
"Message argument of llvm.spv.abort has no SPIR-V type");
6474 return diagnoseUnsupported(
6476 "llvm.spv.abort message type must be a concrete SPIR-V type (numerical "
6477 "scalar, pointer, vector, matrix, or aggregate of such types)");
6480 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
6487bool SPIRVInstructionSelector::selectTrap(MachineInstr &
I)
const {
6496 uint32_t MsgVal = ~0
u;
6497 if (
I.getOpcode() == TargetOpcode::G_UBSANTRAP)
6498 MsgVal =
static_cast<uint32_t
>(
I.getOperand(0).
getImm());
6501 Register MsgReg = buildI32ConstantInEntryBlock(MsgVal,
I, MsgType);
6504 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
6511bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
6512 SPIRVTypeInst ResType,
6513 MachineInstr &
I)
const {
6517 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
6520 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
6523 unsigned Alignment =
I.getOperand(2).getImm();
6530bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
6535 const MachineInstr *PrevI =
I.getPrevNode();
6537 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
6541 .
addMBB(
I.getOperand(0).getMBB())
6546 .
addMBB(
I.getOperand(0).getMBB())
6551bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
6562 const MachineInstr *NextI =
I.getNextNode();
6564 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
6570 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
6572 .
addUse(
I.getOperand(0).getReg())
6573 .
addMBB(
I.getOperand(1).getMBB())
6579bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
6580 MachineInstr &
I)
const {
6582 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
6584 const unsigned NumOps =
I.getNumOperands();
6585 for (
unsigned i = 1; i <
NumOps; i += 2) {
6586 MIB.
addUse(
I.getOperand(i + 0).getReg());
6587 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
6593bool SPIRVInstructionSelector::selectGlobalValue(
6594 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
6596 MachineIRBuilder MIRBuilder(
I);
6597 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
6600 std::string GlobalIdent;
6602 unsigned &
ID = UnnamedGlobalIDs[GV];
6604 ID = UnnamedGlobalIDs.
size();
6605 GlobalIdent =
"__unnamed_" + Twine(
ID).str();
6631 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
6638 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
6643 MachineInstrBuilder MIB1 =
6644 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
6647 MachineInstrBuilder MIB2 =
6649 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
6653 GR.
add(ConstVal, MIB2);
6661 MachineInstrBuilder MIB3 =
6662 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
6665 GR.
add(ConstVal, MIB3);
6669 assert(NewReg != ResVReg);
6670 return BuildCOPY(ResVReg, NewReg,
I);
6680 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
6683 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
6689 SPIRVTypeInst ResType =
6693 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
6698 if (
GlobalVar->isExternallyInitialized() &&
6699 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
6700 constexpr unsigned ReadWriteINTEL = 3u;
6703 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
6709bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
6710 SPIRVTypeInst ResType,
6711 MachineInstr &
I)
const {
6713 return selectExtInst(ResVReg, ResType,
I, CL::log10);
6721 MachineIRBuilder MIRBuilder(
I);
6726 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6729 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6731 .
add(
I.getOperand(1))
6736 ResType->
getOpcode() == SPIRV::OpTypeFloat);
6746 APFloat::rmNearestTiesToEven, &LosesInfo);
6750 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
6751 ? SPIRV::OpVectorTimesScalar
6762bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
6763 SPIRVTypeInst ResType,
6764 MachineInstr &
I)
const {
6767 return selectExtInst(ResVReg, ResType,
I, CL::pown);
6773 Register ExpReg =
I.getOperand(2).getReg();
6775 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
6776 SPIRV::OpConvertSToF))
6778 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
6785bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
6786 SPIRVTypeInst ResType,
6787 MachineInstr &
I)
const {
6803 MachineIRBuilder MIRBuilder(
I);
6806 ResType, MIRBuilder, SPIRV::StorageClass::Function);
6819 MachineBasicBlock &EntryBB =
I.getMF()->
front();
6821 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
6824 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
6830 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6833 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
6836 .
add(
I.getOperand(
I.getNumExplicitDefs()))
6840 Register IntegralPartReg =
I.getOperand(1).getReg();
6841 if (IntegralPartReg.
isValid()) {
6843 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6854 assert(
false &&
"GLSL::Modf is deprecated.");
6865bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
6866 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
6867 SPIRVTypeInst ResType, MachineInstr &
I)
const {
6868 MachineIRBuilder MIRBuilder(
I);
6869 const SPIRVTypeInst Vec3Ty =
6872 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
6884 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
6888 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
6894 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6901 assert(
I.getOperand(2).isReg());
6902 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
6906 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6917bool SPIRVInstructionSelector::loadBuiltinInputID(
6918 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
6919 SPIRVTypeInst ResType, MachineInstr &
I)
const {
6920 MachineIRBuilder MIRBuilder(
I);
6922 ResType, MIRBuilder, SPIRV::StorageClass::Input);
6937 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
6941 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6950SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
6951 MachineInstr &
I)
const {
6952 MachineIRBuilder MIRBuilder(
I);
6953 if (
Type->getOpcode() != SPIRV::OpTypeVector)
6963bool SPIRVInstructionSelector::loadHandleBeforePosition(
6964 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
6965 MachineInstr &Pos)
const {
6968 Intrinsic::spv_resource_handlefrombinding);
6976 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
6977 MachineIRBuilder MIRBuilder(HandleDef);
6978 SPIRVTypeInst VarType = ResType;
6979 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
6981 if (IsStructuredBuffer) {
6986 if (ResType->
getOpcode() == SPIRV::OpTypeImage && ArraySize == 0)
6988 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
6991 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
6992 ArraySize, IndexReg, Name, MIRBuilder);
6996 uint32_t LoadOpcode =
6997 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
7007bool SPIRVInstructionSelector::errorIfInstrOutsideShader(
7008 MachineInstr &
I)
const {
7010 return diagnoseUnsupported(
7011 I,
"this instruction is only supported in shaders.");
7016InstructionSelector *
7020 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 bool isConcreteSPIRVType(SPIRVTypeInst Ty, const SPIRVGlobalRegistry &GR)
static APFloat getOneFP(const Type *LLVMFloatTy)
static bool isUSMStorageClass(SPIRV::StorageClass::StorageClass SC)
static bool isASCastInGVar(MachineRegisterInfo *MRI, Register ResVReg)
static bool mayApplyGenericSelection(unsigned Opcode)
static APFloat getZeroFP(const Type *LLVMFloatTy)
std::vector< std::pair< SPIRV::InstructionSet::InstructionSet, uint32_t > > ExtInstList
static bool intrinsicHasSideEffects(Intrinsic::ID ID)
static unsigned getBoolCmpOpcode(unsigned PredNum)
static unsigned getICmpOpcode(unsigned PredNum)
static bool isOpcodeWithNoSideEffects(unsigned Opcode)
static void addMemoryOperands(MachineMemOperand *MemOp, MachineInstrBuilder &MIB, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry &GR)
static bool isConstReg(MachineRegisterInfo *MRI, MachineInstr *OpDef)
static unsigned getPtrCmpOpcode(unsigned Pred)
bool isDead(const MachineInstr &MI, const MachineRegisterInfo &MRI)
BaseType
A given derived pointer can have multiple base pointers through phi/selects.
static TableGen::Emitter::Opt Y("gen-skeleton-entry", EmitSkeleton, "Generate example skeleton entry")
static ManagedStatic< cl::opt< FnT >, OptCreatorT > CallbackFunction
static const fltSemantics & IEEEsingle()
static const fltSemantics & IEEEdouble()
static const fltSemantics & IEEEhalf()
const fltSemantics & getSemantics() const
static APFloat getOne(const fltSemantics &Sem, bool Negative=false)
Factory for Positive and Negative One.
static APFloat getZero(const fltSemantics &Sem, bool Negative=false)
Factory for Positive and Negative Zero.
static APInt getAllOnes(unsigned numBits)
Return an APInt of a specified width with all bits set.
Represent a constant reference to an array (0 or more elements consecutively in memory),...
BlockFrequencyInfo pass uses BlockFrequencyInfoImpl implementation to estimate IR basic block frequen...
Predicate
This enumeration lists the possible predicates for CmpInst subclasses.
@ FCMP_OEQ
0 0 0 1 True if ordered and equal
@ ICMP_SLT
signed less than
@ ICMP_SLE
signed less or equal
@ FCMP_OLT
0 1 0 0 True if ordered and less than
@ FCMP_ULE
1 1 0 1 True if unordered, less than, or equal
@ FCMP_OGT
0 0 1 0 True if ordered and greater than
@ FCMP_OGE
0 0 1 1 True if ordered and greater than or equal
@ ICMP_UGE
unsigned greater or equal
@ ICMP_UGT
unsigned greater than
@ ICMP_SGT
signed greater than
@ FCMP_ULT
1 1 0 0 True if unordered or less than
@ FCMP_ONE
0 1 1 0 True if ordered and operands are unequal
@ FCMP_UEQ
1 0 0 1 True if unordered or equal
@ ICMP_ULT
unsigned less than
@ FCMP_UGT
1 0 1 0 True if unordered or greater than
@ FCMP_OLE
0 1 0 1 True if ordered and less than or equal
@ FCMP_ORD
0 1 1 1 True if ordered (no nans)
@ ICMP_SGE
signed greater or equal
@ FCMP_UNE
1 1 1 0 True if unordered or not equal
@ ICMP_ULE
unsigned less or equal
@ FCMP_UGE
1 0 1 1 True if unordered, greater than, or equal
@ FCMP_UNO
1 0 0 0 True if unordered: isnan(X) | isnan(Y)
static LLVM_ABI Constant * getNullValue(Type *Ty)
Constructor to create a '0' constant of arbitrary type.
LLVMContext & getContext() const
getContext - Return a reference to the LLVMContext associated with this function.
Represents a call to an intrinsic.
Intrinsic::ID getIntrinsicID() const
unsigned getAddressSpace() const
Module * getParent()
Get the module that this global value is contained inside of...
@ InternalLinkage
Rename collisions when linking (static functions).
static LLVM_ABI IntegerType * get(LLVMContext &C, unsigned NumBits)
This static method is the primary way of constructing an IntegerType.
static constexpr LLT scalar(unsigned SizeInBits)
Get a low-level scalar or aggregate "bag of bits".
constexpr bool isValid() const
constexpr uint16_t getNumElements() const
Returns the number of elements in a vector LLT.
constexpr bool isVector() const
static constexpr LLT pointer(unsigned AddressSpace, unsigned SizeInBits)
Get a low-level pointer in the given address space.
static constexpr LLT fixed_vector(unsigned NumElements, unsigned ScalarSizeInBits)
Get a low-level fixed-width vector of some number of elements and element width.
int getNumber() const
MachineBasicBlocks are uniquely numbered at the function level, unless they're not in a MachineFuncti...
LLVM_ABI iterator getFirstNonPHI()
Returns a pointer to the first instruction in this block that is not a PHINode instruction.
const MachineFunction * getParent() const
Return the MachineFunction containing this basic block.
MachineInstrBundleIterator< MachineInstr > iterator
MachineRegisterInfo & getRegInfo()
getRegInfo - Return information about the registers currently in use.
Function & getFunction()
Return the LLVM function that this machine code represents.
Helper class to build MachineInstr.
MachineBasicBlock::iterator getInsertPt()
Current insertion point for new instructions.
MachineInstrBuilder buildInstr(unsigned Opcode)
Build and insert <empty> = Opcode <empty>.
MachineFunction & getMF()
Getter for the function we currently build.
MachineRegisterInfo * getMRI()
Getter for MRI.
void constrainAllUses(const TargetInstrInfo &TII, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI) const
const MachineInstrBuilder & addUse(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a virtual register use operand.
const MachineInstrBuilder & addReg(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a new virtual register operand.
const MachineInstrBuilder & addImm(int64_t Val) const
Add a new immediate operand.
const MachineInstrBuilder & add(const MachineOperand &MO) const
const MachineInstrBuilder & addMBB(MachineBasicBlock *MBB, unsigned TargetFlags=0) const
const MachineInstrBuilder & addDef(Register RegNo, RegState Flags={}, unsigned SubReg=0) const
Add a virtual register definition operand.
const MachineInstrBuilder & setMIFlags(unsigned Flags) const
MachineInstr * getInstr() const
If conversion operators fail, use this method to get the MachineInstr explicitly.
Representation of each machine instruction.
unsigned getOpcode() const
Returns the opcode of this MachineInstr.
const MachineBasicBlock * getParent() const
unsigned getNumOperands() const
Retuns the total number of operands.
LLVM_ABI unsigned getNumExplicitOperands() const
Returns the number of non-implicit operands.
LLVM_ABI unsigned getNumExplicitDefs() const
Returns the number of non-implicit definitions.
LLVM_ABI void emitGenericError(const Twine &ErrMsg) const
LLVM_ABI const MachineFunction * getMF() const
Return the function that contains the basic block that this instruction belongs to.
const DebugLoc & getDebugLoc() const
Returns the debug location id of this MachineInstr.
const MachineOperand & getOperand(unsigned i) const
A description of a memory reference used in the backend.
@ MOVolatile
The memory access is volatile.
@ MONonTemporal
The memory access is non-temporal.
bool isReg() const
isReg - Tests if this is a MO_Register operand.
MachineBasicBlock * getMBB() const
Register getReg() const
getReg - Returns the register number.
MachineRegisterInfo - Keep track of information for virtual and physical registers,...
defusechain_instr_iterator< true, false, false, true > use_instr_iterator
use_instr_iterator/use_instr_begin/use_instr_end - Walk all uses of the specified register,...
const TargetRegisterClass * getRegClass(Register Reg) const
Return the register class of the specified virtual register.
LLVM_ABI MachineInstr * getVRegDef(Register Reg) const
getVRegDef - Return the machine instr that defines the specified virtual register or null if none is ...
use_instr_iterator use_instr_begin(Register RegNo) const
static def_instr_iterator def_instr_end()
defusechain_instr_iterator< false, true, false, true > def_instr_iterator
def_instr_iterator/def_instr_begin/def_instr_end - Walk all defs of the specified register,...
LLVM_ABI Register createVirtualRegister(const TargetRegisterClass *RegClass, StringRef Name="")
createVirtualRegister - Create and return a new virtual register in the function with the specified r...
def_instr_iterator def_instr_begin(Register RegNo) const
LLT getType(Register Reg) const
Get the low-level type of Reg or LLT{} if Reg is not a generic (target independent) virtual register.
static use_instr_iterator use_instr_end()
iterator_range< use_instr_nodbg_iterator > use_nodbg_instructions(Register Reg) const
LLVM_ABI void setType(Register VReg, LLT Ty)
Set the low-level type of VReg to Ty.
const MachineFunction & getMF() const
LLVM_ABI void setRegClass(Register Reg, const TargetRegisterClass *RC)
setRegClass - Set the register class of the specified virtual register.
LLVM_ABI Register createGenericVirtualRegister(LLT Ty, StringRef Name="")
Create and return a new generic virtual register with low-level type Ty.
const TargetRegisterClass * getRegClassOrNull(Register Reg) const
Return the register class of Reg, or null if Reg has not been assigned a register class yet.
iterator_range< use_instr_iterator > use_instructions(Register Reg) const
unsigned getNumVirtRegs() const
getNumVirtRegs - Return the number of virtual registers created.
LLVM_ABI void replaceRegWith(Register FromReg, Register ToReg)
replaceRegWith - Replace all instances of FromReg with ToReg in the machine function.
Analysis providing profile information.
Holds all the information related to register banks.
Wrapper class representing virtual and physical registers.
constexpr bool isValid() const
constexpr bool isPhysical() const
Return true if the specified register number is in the physical register namespace.
bool isScalarOrVectorSigned(SPIRVTypeInst Type) const
SPIRVTypeInst getOrCreateOpTypeSampledImage(SPIRVTypeInst ImageType, MachineIRBuilder &MIRBuilder)
void assignSPIRVTypeToVReg(SPIRVTypeInst Type, Register VReg, const MachineFunction &MF)
const TargetRegisterClass * getRegClass(SPIRVTypeInst SpvType) const
MachineInstr * getOrAddMemAliasingINTELInst(MachineIRBuilder &MIRBuilder, const MDNode *AliasingListMD)
bool isAggregateType(SPIRVTypeInst Type) const
unsigned getScalarOrVectorBitWidth(SPIRVTypeInst Type) const
SPIRVTypeInst getOrCreateSPIRVIntegerType(unsigned BitWidth, MachineIRBuilder &MIRBuilder)
SPIRVTypeInst getOrCreateSPIRVVectorType(SPIRVTypeInst BaseType, unsigned NumElements, MachineIRBuilder &MIRBuilder, bool EmitIR)
Register buildGlobalVariable(Register Reg, SPIRVTypeInst BaseType, StringRef Name, const GlobalValue *GV, SPIRV::StorageClass::StorageClass Storage, const MachineInstr *Init, bool IsConst, const std::optional< SPIRV::LinkageType::LinkageType > &LinkageType, MachineIRBuilder &MIRBuilder, bool IsInstSelector)
SPIRVTypeInst getResultType(Register VReg, MachineFunction *MF=nullptr)
unsigned getScalarOrVectorComponentCount(Register VReg) const
const Type * getTypeForSPIRVType(SPIRVTypeInst Ty) const
bool isBitcastCompatible(SPIRVTypeInst Type1, SPIRVTypeInst Type2) const
unsigned getPointerSize() const
Register getOrCreateConstFP(APFloat Val, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
LLT getRegType(SPIRVTypeInst SpvType) const
void invalidateMachineInstr(MachineInstr *MI)
SPIRVTypeInst getOrCreateSPIRVBoolType(MachineIRBuilder &MIRBuilder, bool EmitIR)
SPIRVTypeInst getOrCreateSPIRVPointerType(const Type *BaseType, MachineIRBuilder &MIRBuilder, SPIRV::StorageClass::StorageClass SC)
bool isScalarOfType(Register VReg, unsigned TypeOpcode) const
Register getSPIRVTypeID(SPIRVTypeInst SpirvType) const
Register getOrCreateConstInt(uint64_t Val, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
Register getOrCreateConstIntArray(uint64_t Val, size_t Num, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII)
bool findValueAttrs(const MachineInstr *Key, Type *&Ty, StringRef &Name)
SPIRVTypeInst retrieveScalarOrVectorIntType(SPIRVTypeInst Type) const
Register getOrCreateGlobalVariableWithBinding(SPIRVTypeInst VarType, uint32_t Set, uint32_t Binding, StringRef Name, MachineIRBuilder &MIRBuilder)
SPIRVTypeInst changePointerStorageClass(SPIRVTypeInst PtrType, SPIRV::StorageClass::StorageClass SC, MachineInstr &I)
Register getOrCreateConstVector(uint64_t Val, MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII, bool ZeroAsNull=true)
Register buildConstantFP(APFloat Val, MachineIRBuilder &MIRBuilder, SPIRVTypeInst SpvType=nullptr)
void addGlobalObject(const Value *V, const MachineFunction *MF, Register R)
SPIRVTypeInst getScalarOrVectorComponentType(SPIRVTypeInst Type) const
void recordFunctionPointer(const MachineOperand *MO, const Function *F)
SPIRVTypeInst getOrCreateSPIRVFloatType(unsigned BitWidth, MachineInstr &I, const SPIRVInstrInfo &TII)
SPIRVTypeInst getPointeeType(SPIRVTypeInst PtrType)
SPIRVTypeInst getOrCreateSPIRVType(const Type *Type, MachineInstr &I, SPIRV::AccessQualifier::AccessQualifier AQ, bool EmitIR)
bool isScalarOrVectorOfType(Register VReg, unsigned TypeOpcode) const
MachineFunction * setCurrentFunc(MachineFunction &MF)
Register getOrCreateConstNullPtr(MachineIRBuilder &MIRBuilder, SPIRVTypeInst SpvType)
SPIRVTypeInst getSPIRVTypeForVReg(Register VReg, const MachineFunction *MF=nullptr) const
Type * getDeducedGlobalValueType(const GlobalValue *Global)
Register getOrCreateUndef(MachineInstr &I, SPIRVTypeInst SpvType, const SPIRVInstrInfo &TII)
SPIRV::StorageClass::StorageClass getPointerStorageClass(Register VReg) const
bool erase(const MachineInstr *MI)
bool add(SPIRV::IRHandle Handle, const MachineInstr *MI)
Register find(SPIRV::IRHandle Handle, const MachineFunction *MF)
bool isPhysicalSPIRV() const
bool isAtLeastSPIRVVer(VersionTuple VerToCompareTo) const
bool canUseExtInstSet(SPIRV::InstructionSet::InstructionSet E) const
bool isLogicalSPIRV() const
bool canUseExtension(SPIRV::Extension::Extension E) const
bool isTypeIntOrFloat() const
bool erase(PtrType Ptr)
Remove pointer from the set.
std::pair< iterator, bool > insert(PtrType Ptr)
Inserts Ptr if and only if there is no element in the container equal to Ptr.
bool contains(ConstPtrType Ptr) const
SmallPtrSet - This class implements a set which is optimized for holding SmallSize or less elements.
reference emplace_back(ArgTypes &&... Args)
void reserve(size_type N)
void push_back(const T &Elt)
This is a 'vector' (really, a variable-sized array), optimized for the case when the array is small.
constexpr size_t size() const
Get the string size.
static LLVM_ABI StructType * get(LLVMContext &Context, ArrayRef< Type * > Elements, bool isPacked=false)
This static method is the primary way to create a literal StructType.
Twine - A lightweight data structure for efficiently representing the concatenation of temporary valu...
The instances of the Type class are immutable: once they are created, they are never changed.
@ HalfTyID
16-bit floating point type
@ FloatTyID
32-bit floating point type
@ DoubleTyID
64-bit floating point type
Type * getScalarType() const
If this is a vector type, return the element type, otherwise return 'this'.
bool isStructTy() const
True if this is an instance of StructType.
bool isAggregateType() const
Return true if the type is an aggregate type.
TypeID getTypeID() const
Return the type id for the type.
Value * getOperand(unsigned i) const
LLVM_ABI StringRef getName() const
Return a constant reference to the value's name.
self_iterator getIterator()
NodeTy * getNextNode()
Get the next node, or nullptr for the list tail.
#define llvm_unreachable(msg)
Marks that the current location is not supposed to be reachable.
constexpr char IsConst[]
Key for Kernel::Arg::Metadata::mIsConst.
constexpr std::underlying_type_t< E > Mask()
Get a bitmask with 1s in all places up to the high-order bit of E's largest value.
unsigned ID
LLVM IR allows to use arbitrary numbers as calling convention identifiers.
Scope
Defines the scope in which this symbol should be visible: Default – Visible in the public interface o...
NodeAddr< DefNode * > Def
NodeAddr< InstrNode * > Instr
NodeAddr< UseNode * > Use
NodeAddr< FuncNode * > Func
BaseReg
Stack frame base register. Bit 0 of FREInfo.Info.
This is an optimization pass for GlobalISel generic memory operations.
void buildOpName(Register Target, const StringRef &Name, MachineIRBuilder &MIRBuilder)
@ Low
Lower the current thread's priority such that it does not affect foreground tasks significantly.
FunctionAddr VTableAddr Value
bool all_of(R &&range, UnaryPredicate P)
Provide wrappers to std::all_of which take ranges instead of having to pass begin/end explicitly.
MachineBasicBlock::iterator getOpVariableMBBIt(MachineFunction &MF)
int64_t getIConstValSext(Register ConstReg, const MachineRegisterInfo *MRI)
MachineInstrBuilder BuildMI(MachineFunction &MF, const MIMetadata &MIMD, const MCInstrDesc &MCID)
Builder interface. Specify how to create the initial instruction itself.
bool isTypeFoldingSupported(unsigned Opcode)
decltype(auto) dyn_cast(const From &Val)
dyn_cast<X> - Return the argument parameter cast to the specified type.
FunctionAddr VTableAddr uintptr_t uintptr_t Int32Ty
void addNumImm(const APInt &Imm, MachineInstrBuilder &MIB)
LLVM_ABI void salvageDebugInfo(const MachineRegisterInfo &MRI, MachineInstr &MI)
Assuming the instruction MI is going to be deleted, attempt to salvage debug users of MI by writing t...
LLVM_ABI void constrainSelectedInstRegOperands(MachineInstr &I, const TargetInstrInfo &TII, const TargetRegisterInfo &TRI, const RegisterBankInfo &RBI)
Mutate the newly-selected instruction I to constrain its (possibly generic) virtual register operands...
bool isPreISelGenericOpcode(unsigned Opcode)
Check whether the given Opcode is a generic opcode that is not supposed to appear after ISel.
Register createVirtualRegister(SPIRVTypeInst SpvType, SPIRVGlobalRegistry *GR, MachineRegisterInfo *MRI, const MachineFunction &MF)
unsigned getArrayComponentCount(const MachineRegisterInfo *MRI, const MachineInstr *ResType)
uint64_t getIConstVal(Register ConstReg, const MachineRegisterInfo *MRI)
SmallVector< MachineInstr *, 4 > createContinuedInstructions(MachineIRBuilder &MIRBuilder, unsigned Opcode, unsigned MinWC, unsigned ContinuedOpcode, ArrayRef< Register > Args, Register ReturnRegister, Register TypeID)
SPIRV::MemorySemantics::MemorySemantics getMemSemanticsForStorageClass(SPIRV::StorageClass::StorageClass SC)
constexpr unsigned storageClassToAddressSpace(SPIRV::StorageClass::StorageClass SC)
bool any_of(R &&range, UnaryPredicate P)
Provide wrappers to std::any_of which take ranges instead of having to pass begin/end explicitly.
void buildOpDecorate(Register Reg, MachineIRBuilder &MIRBuilder, SPIRV::Decoration::Decoration Dec, const std::vector< uint32_t > &DecArgs, StringRef StrImm)
MachineInstr * getImm(const MachineOperand &MO, const MachineRegisterInfo *MRI)
constexpr bool isPowerOf2_32(uint32_t Value)
Return true if the argument is a power of two > 0.
Type * toTypedPointer(Type *Ty)
LLVM_ABI raw_ostream & dbgs()
dbgs() - This returns a reference to a raw_ostream for debugging messages.
LLVM_ABI void report_fatal_error(Error Err, bool gen_crash_diag=true)
constexpr bool isGenericCastablePtr(SPIRV::StorageClass::StorageClass SC)
class LLVM_GSL_OWNER SmallVector
Forward declaration of SmallVector so that calculateSmallVectorDefaultInlinedElements can reference s...
MachineInstr * passCopy(MachineInstr *Def, const MachineRegisterInfo *MRI)
bool isa(const From &Val)
isa<X> - Return true if the parameter to the template is an instance of one of the template type argu...
std::optional< SPIRV::LinkageType::LinkageType > getSpirvLinkageTypeFor(const SPIRVSubtarget &ST, const GlobalValue &GV)
LLVM_ABI raw_fd_ostream & errs()
This returns a reference to a raw_ostream for standard error.
SPIRV::StorageClass::StorageClass addressSpaceToStorageClass(unsigned AddrSpace, const SPIRVSubtarget &STI)
AtomicOrdering
Atomic ordering for LLVM's memory model.
SPIRV::Scope::Scope getMemScope(LLVMContext &Ctx, SyncScope::ID Id)
InstructionSelector * createSPIRVInstructionSelector(const SPIRVTargetMachine &TM, const SPIRVSubtarget &Subtarget, const RegisterBankInfo &RBI)
std::string getStringValueFromReg(Register Reg, MachineRegisterInfo &MRI)
int64_t foldImm(const MachineOperand &MO, const MachineRegisterInfo *MRI)
DWARFExpression::Operation Op
ArrayRef(const T &OneElt) -> ArrayRef< T >
MachineInstr * getDefInstrMaybeConstant(Register &ConstReg, const MachineRegisterInfo *MRI)
constexpr unsigned BitWidth
decltype(auto) cast(const From &Val)
cast<X> - Return the argument parameter cast to the specified type.
bool hasInitializer(const GlobalVariable *GV)
bool isSpvIntrinsic(const MachineInstr &MI, Intrinsic::ID IntrinsicID)
void addStringImm(const StringRef &Str, MCInst &Inst)
MachineInstr * getVRegDef(MachineRegisterInfo &MRI, Register Reg)
SPIRV::MemorySemantics::MemorySemantics getMemSemantics(AtomicOrdering Ord)
std::string getLinkStringForBuiltIn(SPIRV::BuiltIn::BuiltIn BuiltInValue)
LLVM_ABI bool isTriviallyDead(const MachineInstr &MI, const MachineRegisterInfo &MRI)
Check whether an instruction MI is dead: it only defines dead virtual registers, and doesn't have oth...