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 && IntrID != Intrinsic::spv_poison)
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::OpPoisonKHR:
654 case SPIRV::OpConstantFunctionPointerINTEL:
681 case Intrinsic::spv_all:
682 case Intrinsic::spv_alloca:
683 case Intrinsic::spv_any:
684 case Intrinsic::spv_bitcast:
685 case Intrinsic::spv_const_composite:
686 case Intrinsic::spv_cross:
687 case Intrinsic::spv_degrees:
688 case Intrinsic::spv_distance:
689 case Intrinsic::spv_extractelt:
690 case Intrinsic::spv_extractv:
691 case Intrinsic::spv_faceforward:
692 case Intrinsic::spv_fdot:
693 case Intrinsic::spv_firstbitlow:
694 case Intrinsic::spv_firstbitshigh:
695 case Intrinsic::spv_firstbituhigh:
696 case Intrinsic::spv_frac:
697 case Intrinsic::spv_gep:
698 case Intrinsic::spv_global_offset:
699 case Intrinsic::spv_global_size:
700 case Intrinsic::spv_group_id:
701 case Intrinsic::spv_insertelt:
702 case Intrinsic::spv_insertv:
703 case Intrinsic::spv_isinf:
704 case Intrinsic::spv_isnan:
705 case Intrinsic::spv_isfinite:
706 case Intrinsic::spv_isnormal:
707 case Intrinsic::spv_lerp:
708 case Intrinsic::spv_length:
709 case Intrinsic::spv_normalize:
710 case Intrinsic::spv_num_subgroups:
711 case Intrinsic::spv_num_workgroups:
712 case Intrinsic::spv_ptrcast:
713 case Intrinsic::spv_radians:
714 case Intrinsic::spv_reflect:
715 case Intrinsic::spv_refract:
716 case Intrinsic::spv_resource_getbasepointer:
717 case Intrinsic::spv_resource_getpointer:
718 case Intrinsic::spv_resource_handlefrombinding:
719 case Intrinsic::spv_resource_handlefromimplicitbinding:
720 case Intrinsic::spv_resource_nonuniformindex:
721 case Intrinsic::spv_resource_sample:
722 case Intrinsic::spv_rsqrt:
723 case Intrinsic::spv_saturate:
724 case Intrinsic::spv_sdot:
725 case Intrinsic::spv_sign:
726 case Intrinsic::spv_smoothstep:
727 case Intrinsic::spv_step:
728 case Intrinsic::spv_subgroup_id:
729 case Intrinsic::spv_subgroup_local_invocation_id:
730 case Intrinsic::spv_subgroup_max_size:
731 case Intrinsic::spv_subgroup_size:
732 case Intrinsic::spv_thread_id:
733 case Intrinsic::spv_thread_id_in_group:
734 case Intrinsic::spv_udot:
735 case Intrinsic::spv_undef:
736 case Intrinsic::spv_value_md:
737 case Intrinsic::spv_workgroup_size:
749 case SPIRV::OpTypeVoid:
750 case SPIRV::OpTypeBool:
751 case SPIRV::OpTypeInt:
752 case SPIRV::OpTypeFloat:
753 case SPIRV::OpTypeVector:
754 case SPIRV::OpTypeMatrix:
755 case SPIRV::OpTypeImage:
756 case SPIRV::OpTypeSampler:
757 case SPIRV::OpTypeSampledImage:
758 case SPIRV::OpTypeArray:
759 case SPIRV::OpTypeRuntimeArray:
760 case SPIRV::OpTypeStruct:
761 case SPIRV::OpTypeOpaque:
762 case SPIRV::OpTypePointer:
763 case SPIRV::OpTypeFunction:
764 case SPIRV::OpTypeEvent:
765 case SPIRV::OpTypeDeviceEvent:
766 case SPIRV::OpTypeReserveId:
767 case SPIRV::OpTypeQueue:
768 case SPIRV::OpTypePipe:
769 case SPIRV::OpTypeForwardPointer:
770 case SPIRV::OpTypePipeStorage:
771 case SPIRV::OpTypeNamedBarrier:
772 case SPIRV::OpTypeAccelerationStructureNV:
773 case SPIRV::OpTypeCooperativeMatrixNV:
774 case SPIRV::OpTypeCooperativeMatrixKHR:
784 if (
MI.getNumDefs() == 0)
787 for (
const auto &MO :
MI.all_defs()) {
789 if (
Reg.isPhysical()) {
794 if (
UseMI.getOpcode() != SPIRV::OpName) {
801 if (
MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE ||
MI.isFakeUse() ||
802 MI.isLifetimeMarker()) {
805 <<
"Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
816 if (
MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
817 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
820 LLVM_DEBUG(
dbgs() <<
"Dead: Intrinsic with no real side effects.\n");
825 if (
MI.mayStore() ||
MI.isCall() ||
826 (
MI.mayLoad() &&
MI.hasOrderedMemoryRef()) ||
MI.isPosition() ||
827 MI.isDebugInstr() ||
MI.isTerminator() ||
MI.isJumpTableDebugInfo()) {
828 LLVM_DEBUG(
dbgs() <<
"Not dead: instruction has side effects.\n");
839 LLVM_DEBUG(
dbgs() <<
"Dead: known opcode with no side effects\n");
846void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &
MI)
const {
848 for (
const auto &MO :
MI.all_defs()) {
852 SmallVector<MachineInstr *, 4> UselessOpNames;
855 "There is still a use of the dead function.");
858 for (MachineInstr *OpNameMI : UselessOpNames) {
860 OpNameMI->eraseFromParent();
865void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &
MI)
const {
868 removeOpNamesForDeadMI(
MI);
869 MI.eraseFromParent();
872bool SPIRVInstructionSelector::select(MachineInstr &
I) {
873 resetVRegsType(*
I.getParent()->getParent());
875 assert(
I.getParent() &&
"Instruction should be in a basic block!");
876 assert(
I.getParent()->getParent() &&
"Instruction should be in a function!");
881 removeDeadInstruction(
I);
888 if (Opcode == SPIRV::ASSIGN_TYPE) {
889 Register DstReg =
I.getOperand(0).getReg();
890 Register SrcReg =
I.getOperand(1).getReg();
893 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
894 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
895 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
896 Register SelectDstReg =
Def->getOperand(0).getReg();
897 bool SuccessToSelectSelect [[maybe_unused]] = selectSelect(
899 assert(SuccessToSelectSelect);
901 Def->eraseFromParent();
908 bool Res = selectImpl(
I, *CoverageInfo);
910 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
911 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
915 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
927 }
else if (
I.getNumDefs() == 1) {
939 removeDeadInstruction(
I);
944 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
945 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
951 bool HasDefs =
I.getNumDefs() > 0;
954 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
955 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
956 if (spvSelect(ResVReg, ResType,
I)) {
958 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
969 case TargetOpcode::G_CONSTANT:
970 case TargetOpcode::G_FCONSTANT:
977 MachineInstr &
I)
const {
980 if (DstRC != SrcRC && SrcRC)
982 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
989bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
990 SPIRVTypeInst ResType,
991 MachineInstr &
I)
const {
992 const unsigned Opcode =
I.getOpcode();
994 return selectImpl(
I, *CoverageInfo);
996 case TargetOpcode::G_CONSTANT:
997 case TargetOpcode::G_FCONSTANT:
998 return selectConst(ResVReg, ResType,
I);
999 case TargetOpcode::G_GLOBAL_VALUE:
1000 return selectGlobalValue(ResVReg,
I);
1001 case TargetOpcode::G_IMPLICIT_DEF:
1002 return selectOpUndef(ResVReg, ResType,
I);
1003 case TargetOpcode::G_FREEZE:
1004 return selectFreeze(ResVReg, ResType,
I);
1006 case TargetOpcode::G_INTRINSIC:
1007 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
1008 case TargetOpcode::G_INTRINSIC_CONVERGENT:
1009 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
1010 return selectIntrinsic(ResVReg, ResType,
I);
1011 case TargetOpcode::G_BITREVERSE:
1012 return selectBitreverse(ResVReg, ResType,
I);
1014 case TargetOpcode::G_BUILD_VECTOR:
1015 return selectBuildVector(ResVReg, ResType,
I);
1016 case TargetOpcode::G_SPLAT_VECTOR:
1017 return selectSplatVector(ResVReg, ResType,
I);
1019 case TargetOpcode::G_SHUFFLE_VECTOR: {
1020 MachineBasicBlock &BB = *
I.getParent();
1021 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
1024 .
addUse(
I.getOperand(1).getReg())
1025 .
addUse(
I.getOperand(2).getReg());
1026 for (
auto V :
I.getOperand(3).getShuffleMask())
1031 case TargetOpcode::G_MEMMOVE:
1032 case TargetOpcode::G_MEMCPY:
1033 case TargetOpcode::G_MEMSET:
1034 return selectMemOperation(ResVReg,
I);
1036 case TargetOpcode::G_ICMP:
1037 return selectICmp(ResVReg, ResType,
I);
1038 case TargetOpcode::G_FCMP:
1039 return selectFCmp(ResVReg, ResType,
I);
1041 case TargetOpcode::G_FRAME_INDEX:
1042 return selectFrameIndex(ResVReg, ResType,
I);
1044 case TargetOpcode::G_LOAD:
1045 return selectLoad(ResVReg, ResType,
I);
1046 case TargetOpcode::G_STORE:
1047 return selectStore(
I);
1049 case TargetOpcode::G_BR:
1050 return selectBranch(
I);
1051 case TargetOpcode::G_BRCOND:
1052 return selectBranchCond(
I);
1054 case TargetOpcode::G_PHI:
1055 return selectPhi(ResVReg,
I);
1057 case TargetOpcode::G_FPTOSI:
1058 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1059 case TargetOpcode::G_FPTOUI:
1060 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1062 case TargetOpcode::G_FPTOSI_SAT:
1063 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1064 case TargetOpcode::G_FPTOUI_SAT:
1065 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1067 case TargetOpcode::G_SITOFP:
1068 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
1069 case TargetOpcode::G_UITOFP:
1070 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
1072 case TargetOpcode::G_CTPOP:
1073 return selectPopCount(ResVReg, ResType,
I, SPIRV::OpBitCount);
1074 case TargetOpcode::G_SMIN:
1075 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
1076 case TargetOpcode::G_UMIN:
1077 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
1079 case TargetOpcode::G_SMAX:
1080 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
1081 case TargetOpcode::G_UMAX:
1082 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
1084 case TargetOpcode::G_SCMP:
1085 return selectSUCmp(ResVReg, ResType,
I,
true);
1086 case TargetOpcode::G_UCMP:
1087 return selectSUCmp(ResVReg, ResType,
I,
false);
1088 case TargetOpcode::G_LROUND:
1089 case TargetOpcode::G_LLROUND: {
1092 MRI->
setRegClass(regForLround, &SPIRV::iIDRegClass);
1094 regForLround, *(
I.getParent()->getParent()));
1096 CL::round, GL::Round,
false);
1098 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1105 case TargetOpcode::G_STRICT_FMA:
1106 case TargetOpcode::G_FMA: {
1109 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1112 .
addUse(
I.getOperand(1).getReg())
1113 .
addUse(
I.getOperand(2).getReg())
1114 .
addUse(
I.getOperand(3).getReg())
1119 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1122 case TargetOpcode::G_STRICT_FLDEXP:
1123 return selectExtInst(ResVReg, ResType,
I, CL::ldexp);
1125 case TargetOpcode::G_FPOW:
1126 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1127 case TargetOpcode::G_FPOWI:
1128 return selectFpowi(ResVReg, ResType,
I);
1130 case TargetOpcode::G_FEXP:
1131 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1132 case TargetOpcode::G_FEXP2:
1133 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1134 case TargetOpcode::G_FEXP10:
1135 return selectExp10(ResVReg, ResType,
I);
1137 case TargetOpcode::G_FMODF:
1138 return selectModf(ResVReg, ResType,
I);
1139 case TargetOpcode::G_FSINCOS:
1140 return selectSincos(ResVReg, ResType,
I);
1142 case TargetOpcode::G_FLOG:
1143 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1144 case TargetOpcode::G_FLOG2:
1145 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1146 case TargetOpcode::G_FLOG10:
1147 return selectLog10(ResVReg, ResType,
I);
1149 case TargetOpcode::G_FABS:
1150 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1151 case TargetOpcode::G_ABS:
1152 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1154 case TargetOpcode::G_FMINNUM:
1155 case TargetOpcode::G_FMINIMUM:
1156 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1157 case TargetOpcode::G_FMAXNUM:
1158 case TargetOpcode::G_FMAXIMUM:
1159 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1161 case TargetOpcode::G_FCOPYSIGN:
1162 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1164 case TargetOpcode::G_FCEIL:
1165 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1166 case TargetOpcode::G_FFLOOR:
1167 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1169 case TargetOpcode::G_FCOS:
1170 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1171 case TargetOpcode::G_FSIN:
1172 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1173 case TargetOpcode::G_FTAN:
1174 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1175 case TargetOpcode::G_FACOS:
1176 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1177 case TargetOpcode::G_FASIN:
1178 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1179 case TargetOpcode::G_FATAN:
1180 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1181 case TargetOpcode::G_FATAN2:
1182 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1183 case TargetOpcode::G_FCOSH:
1184 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1185 case TargetOpcode::G_FSINH:
1186 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1187 case TargetOpcode::G_FTANH:
1188 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1190 case TargetOpcode::G_STRICT_FSQRT:
1191 case TargetOpcode::G_FSQRT:
1192 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1194 case TargetOpcode::G_CTTZ:
1195 case TargetOpcode::G_CTTZ_ZERO_POISON:
1196 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1197 case TargetOpcode::G_CTLZ:
1198 case TargetOpcode::G_CTLZ_ZERO_POISON:
1199 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1201 case TargetOpcode::G_INTRINSIC_ROUND:
1202 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1203 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1204 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1205 case TargetOpcode::G_INTRINSIC_TRUNC:
1206 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1207 case TargetOpcode::G_FRINT:
1208 case TargetOpcode::G_FNEARBYINT:
1209 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1211 case TargetOpcode::G_SMULH:
1212 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1213 case TargetOpcode::G_UMULH:
1214 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1216 case TargetOpcode::G_SADDSAT:
1217 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1218 case TargetOpcode::G_UADDSAT:
1219 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1220 case TargetOpcode::G_SSUBSAT:
1221 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1222 case TargetOpcode::G_USUBSAT:
1223 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1225 case TargetOpcode::G_FFREXP:
1226 return selectFrexp(ResVReg, ResType,
I);
1228 case TargetOpcode::G_UADDO:
1229 return selectOverflowArith(ResVReg, ResType,
I,
1230 ResType->
getOpcode() == SPIRV::OpTypeVector
1231 ? SPIRV::OpIAddCarryV
1232 : SPIRV::OpIAddCarryS);
1233 case TargetOpcode::G_USUBO:
1234 return selectOverflowArith(ResVReg, ResType,
I,
1235 ResType->
getOpcode() == SPIRV::OpTypeVector
1236 ? SPIRV::OpISubBorrowV
1237 : SPIRV::OpISubBorrowS);
1238 case TargetOpcode::G_UMULO:
1239 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1240 case TargetOpcode::G_SMULO:
1241 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1243 case TargetOpcode::G_SEXT:
1244 return selectExt(ResVReg, ResType,
I,
true);
1245 case TargetOpcode::G_ANYEXT:
1246 case TargetOpcode::G_ZEXT:
1247 return selectExt(ResVReg, ResType,
I,
false);
1248 case TargetOpcode::G_TRUNC:
1249 return selectTrunc(ResVReg, ResType,
I);
1250 case TargetOpcode::G_FPTRUNC:
1251 case TargetOpcode::G_FPEXT:
1252 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1254 case TargetOpcode::G_PTRTOINT:
1255 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1256 case TargetOpcode::G_INTTOPTR:
1257 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1258 case TargetOpcode::G_BITCAST:
1259 return selectBitcast(ResVReg, ResType,
I);
1260 case TargetOpcode::G_ADDRSPACE_CAST:
1261 return selectAddrSpaceCast(ResVReg, ResType,
I);
1262 case TargetOpcode::G_PTR_ADD: {
1264 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1268 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1269 (*II).getOpcode() == TargetOpcode::COPY ||
1270 (*II).getOpcode() == SPIRV::OpVariable) &&
1271 getImm(
I.getOperand(2), MRI));
1273 bool IsGVInit =
false;
1277 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1278 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1279 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1280 (*UseIt).getOpcode() == SPIRV::OpVariable) {
1290 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1302 return diagnoseUnsupported(
1303 I,
"incompatible result and operand types in a bitcast");
1305 MachineInstrBuilder MIB =
1306 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1313 : SPIRV::OpInBoundsPtrAccessChain))
1317 .
addUse(
I.getOperand(2).getReg())
1320 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1324 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1326 .
addUse(
I.getOperand(2).getReg())
1335 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1338 .
addImm(
static_cast<uint32_t
>(
1339 SPIRV::Opcode::InBoundsPtrAccessChain))
1342 .
addUse(
I.getOperand(2).getReg());
1347 case TargetOpcode::G_ATOMICRMW_OR:
1348 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1349 case TargetOpcode::G_ATOMICRMW_ADD:
1350 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1351 case TargetOpcode::G_ATOMICRMW_AND:
1352 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1353 case TargetOpcode::G_ATOMICRMW_MAX:
1354 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1355 case TargetOpcode::G_ATOMICRMW_MIN:
1356 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1357 case TargetOpcode::G_ATOMICRMW_SUB:
1358 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1359 case TargetOpcode::G_ATOMICRMW_XOR:
1360 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1361 case TargetOpcode::G_ATOMICRMW_UMAX:
1362 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1363 case TargetOpcode::G_ATOMICRMW_UMIN:
1364 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1365 case TargetOpcode::G_ATOMICRMW_XCHG:
1366 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1368 case TargetOpcode::G_ATOMICRMW_FADD:
1369 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1370 case TargetOpcode::G_ATOMICRMW_FSUB:
1372 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1373 ResType->
getOpcode() == SPIRV::OpTypeVector
1375 : SPIRV::OpFNegate);
1376 case TargetOpcode::G_ATOMICRMW_FMIN:
1377 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1378 case TargetOpcode::G_ATOMICRMW_FMAX:
1379 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1381 case TargetOpcode::G_FENCE:
1382 return selectFence(
I);
1384 case TargetOpcode::G_STACKSAVE:
1385 return selectStackSave(ResVReg, ResType,
I);
1386 case TargetOpcode::G_STACKRESTORE:
1387 return selectStackRestore(
I);
1389 case TargetOpcode::G_UNMERGE_VALUES:
1392 case TargetOpcode::G_TRAP:
1393 case TargetOpcode::G_UBSANTRAP:
1394 return selectTrap(
I);
1399 case TargetOpcode::DBG_LABEL:
1401 case TargetOpcode::G_DEBUGTRAP:
1402 return selectDebugTrap(ResVReg, ResType,
I);
1409bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1410 SPIRVTypeInst ResType,
1411 MachineInstr &
I)
const {
1412 unsigned Opcode = SPIRV::OpNop;
1419bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1420 SPIRVTypeInst ResType,
1422 GL::GLSLExtInst GLInst,
1423 bool setMIFlags,
bool useMISrc,
1426 SPIRV::InstructionSet::InstructionSet::GLSL_std_450))
1427 return diagnoseUnsupported(
1429 "this instruction is only supported with the GLSL extended instruction "
1431 return selectExtInst(ResVReg, ResType,
I,
1432 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1433 setMIFlags, useMISrc, SrcRegs);
1436bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1437 SPIRVTypeInst ResType,
1439 CL::OpenCLExtInst CLInst,
1440 bool setMIFlags,
bool useMISrc,
1442 return selectExtInst(ResVReg, ResType,
I,
1443 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1444 setMIFlags, useMISrc, SrcRegs);
1447bool SPIRVInstructionSelector::selectExtInst(
1448 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1449 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1451 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1452 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1453 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1457bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1458 SPIRVTypeInst ResType,
1461 bool setMIFlags,
bool useMISrc,
1464 for (
const auto &[InstructionSet, Opcode] : Insts) {
1468 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1471 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1476 const unsigned NumOps =
I.getNumOperands();
1479 I.getOperand(Index).getType() ==
1480 MachineOperand::MachineOperandType::MO_IntrinsicID)
1483 MIB.
add(
I.getOperand(Index));
1495bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1496 SPIRVTypeInst ResType,
1497 MachineInstr &
I)
const {
1498 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1499 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1500 for (
const auto &Ex : ExtInsts) {
1501 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1502 uint32_t Opcode = Ex.second;
1506 MachineIRBuilder MIRBuilder(
I);
1509 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1514 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1517 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1520 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1523 .
addImm(
static_cast<uint32_t
>(Ex.first))
1525 .
add(
I.getOperand(2))
1529 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1530 .
addDef(
I.getOperand(1).getReg())
1539bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1540 SPIRVTypeInst ResType,
1541 MachineInstr &
I)
const {
1542 Register CosResVReg =
I.getOperand(1).getReg();
1543 unsigned SrcIdx =
I.getNumExplicitDefs();
1548 MachineIRBuilder MIRBuilder(
I);
1550 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1555 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1558 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1560 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1563 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1565 .
add(
I.getOperand(SrcIdx))
1568 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1576 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1579 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1581 .
add(
I.getOperand(SrcIdx))
1583 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1586 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1588 .
add(
I.getOperand(SrcIdx))
1595bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1596 SPIRVTypeInst ResType,
1598 std::vector<Register> Srcs,
1599 unsigned Opcode)
const {
1600 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1610std::optional<SplitParts> SPIRVInstructionSelector::splitEvenOddLanes(
1611 Register PopCountReg,
unsigned ComponentCount, MachineInstr &
I,
1612 SPIRVTypeInst I32Type)
const {
1615 if (ComponentCount == 1) {
1618 Parts.IsScalar =
true;
1619 Parts.Type = I32Type;
1627 if (!selectOpWithSrcs(Parts.High, I32Type,
I, {PopCountReg, IdxOne},
1628 SPIRV::OpVectorExtractDynamic))
1629 return std::nullopt;
1631 if (!selectOpWithSrcs(Parts.Low, I32Type,
I, {PopCountReg, IdxZero},
1632 SPIRV::OpVectorExtractDynamic))
1633 return std::nullopt;
1637 MachineIRBuilder MIRBuilder(
I);
1638 Parts.IsScalar =
false;
1645 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1646 TII.get(SPIRV::OpVectorShuffle))
1651 for (
unsigned J = 1; J < ComponentCount * 2; J += 2)
1656 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1657 TII.get(SPIRV::OpVectorShuffle))
1662 for (
unsigned J = 0; J < ComponentCount * 2; J += 2)
1670bool SPIRVInstructionSelector::selectPopCount16(
Register ResVReg,
1671 SPIRVTypeInst ResType,
1674 unsigned Opcode)
const {
1675 Register OpReg =
I.getOperand(1).getReg();
1678 MachineIRBuilder MIRBuilder(
I);
1680 SPIRVTypeInst I32VectorType =
1683 bool IsVector = NumElems > 1;
1684 SPIRVTypeInst ExtType = IsVector ? I32VectorType : I32Type;
1687 if (!selectOpWithSrcs(ExtReg, ExtType,
I, {OpReg}, SPIRV::OpUConvert))
1691 if (!selectPopCount32(PopCountReg, ExtType,
I, ExtReg, Opcode))
1694 return selectOpWithSrcs(ResVReg, ResType,
I, {PopCountReg}, ExtOpcode);
1697bool SPIRVInstructionSelector::selectPopCount32(
Register ResVReg,
1698 SPIRVTypeInst ResType,
1701 unsigned Opcode)
const {
1702 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
1705bool SPIRVInstructionSelector::selectPopCount64(
Register ResVReg,
1706 SPIRVTypeInst ResType,
1709 unsigned Opcode)
const {
1711 if (ComponentCount > 2)
1712 return handle64BitOverflow(
1713 ResVReg, ResType,
I, SrcReg, Opcode,
1715 unsigned O) {
return this->selectPopCount64(R,
T,
I, S, O); });
1717 MachineIRBuilder MIRBuilder(
I);
1722 I32Type, 2 * ComponentCount, MIRBuilder,
false);
1726 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
1731 if (!selectPopCount32(Pop32, VecI32Type,
I, Vec32, Opcode))
1735 auto MaybeParts = splitEvenOddLanes(Pop32, ComponentCount,
I, I32Type);
1738 SplitParts &Parts = *MaybeParts;
1741 unsigned OpAdd = Parts.IsScalar ? SPIRV::OpIAddS : SPIRV::OpIAddV;
1743 if (!selectOpWithSrcs(Sum, Parts.Type,
I, {Parts.High, Parts.Low}, OpAdd))
1748 unsigned ConvOp = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
1749 return selectOpWithSrcs(ResVReg, ResType,
I, {Sum}, ConvOp);
1752bool SPIRVInstructionSelector::selectPopCount(
Register ResVReg,
1753 SPIRVTypeInst ResType,
1755 unsigned Opcode)
const {
1760 if (!STI.getTargetTriple().isVulkanOS())
1761 return selectUnOp(ResVReg, ResType,
I, Opcode);
1763 Register OpReg =
I.getOperand(1).getReg();
1766 : SPIRV::OpUConvert;
1770 return selectPopCount16(ResVReg, ResType,
I, ExtOpcode, Opcode);
1772 return selectPopCount32(ResVReg, ResType,
I, OpReg, Opcode);
1774 return selectPopCount64(ResVReg, ResType,
I, OpReg, Opcode);
1776 return diagnoseUnsupported(
I,
"unsupported operand bit width for popcount");
1780bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1781 SPIRVTypeInst ResType,
1783 unsigned Opcode)
const {
1785 Register SrcReg =
I.getOperand(1).getReg();
1790 unsigned DefOpCode = DefIt->getOpcode();
1791 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1794 if (
auto *VRD =
getVRegDef(*MRI, DefIt->getOperand(1).getReg()))
1795 DefOpCode = VRD->getOpcode();
1797 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1798 DefOpCode == TargetOpcode::G_CONSTANT ||
1799 DefOpCode == SPIRV::OpVariable || DefOpCode == SPIRV::OpConstantI) {
1805 uint32_t SpecOpcode = 0;
1807 case SPIRV::OpConvertPtrToU:
1808 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1810 case SPIRV::OpConvertUToPtr:
1811 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1816 TII.get(SPIRV::OpSpecConstantOp))
1826 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1830bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1831 SPIRVTypeInst ResType,
1832 MachineInstr &
I)
const {
1833 Register OpReg =
I.getOperand(1).getReg();
1834 SPIRVTypeInst OpType =
1837 return diagnoseUnsupported(
1838 I,
"incompatible result and operand types in a bitcast");
1839 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1849 if (
MemOp->isVolatile())
1850 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1851 if (
MemOp->isNonTemporal())
1852 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1854 if (!ST->isShader() &&
MemOp->getAlign().value())
1855 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1859 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1860 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1864 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1866 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1870 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1874 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1876 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1888 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1890 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1892 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
1896bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
1897 SPIRVTypeInst ResType,
1898 MachineInstr &
I)
const {
1900 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1905 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
1906 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
1908 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1910 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1914 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1918 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1919 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
1920 I.getDebugLoc(),
I);
1924 MachineIRBuilder MIRBuilder(
I);
1926 if (
I.getNumMemOperands()) {
1927 const MachineMemOperand *MemOp = *
I.memoperands_begin();
1928 if (MemOp->isAtomic())
1929 return selectAtomicLoad(ResVReg, ResType,
I);
1932 auto MIB = MIRBuilder.buildInstr(SPIRV::OpLoad)
1936 if (!
I.getNumMemOperands()) {
1937 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1939 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1948bool SPIRVInstructionSelector::selectAtomicLoad(
Register ResVReg,
1949 SPIRVTypeInst ResType,
1950 MachineInstr &
I)
const {
1951 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
1954 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1957 return diagnoseUnsupported(
I,
1958 "Lowering to SPIR-V of atomic load is only "
1959 "allowed for integer or floating point types");
1961 assert(
I.getNumMemOperands());
1962 const MachineMemOperand &MemOp = **
I.memoperands_begin();
1963 assert(MemOp.isAtomic());
1967 Register ScopeReg = buildI32Constant(Scope,
I);
1973 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
1974 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
1977 MachineIRBuilder MIRBuilder(
I);
1978 auto AtomicLoad = MIRBuilder.buildInstr(SPIRV::OpAtomicLoad)
1984 AtomicLoad.constrainAllUses(
TII,
TRI, RBI);
1988bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
1990 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
1991 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1996 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
1997 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
1999 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2004 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2008 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2009 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2010 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2011 TII.get(SPIRV::OpImageWrite))
2017 if (sampledTypeIsSignedInteger(LLVMHandleType))
2020 BMI.constrainAllUses(
TII,
TRI, RBI);
2025 if (
I.getNumMemOperands()) {
2026 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2027 if (MemOp->isAtomic())
2028 return selectAtomicStore(
I);
2031 MachineIRBuilder MIRBuilder(
I);
2032 auto MIB = MIRBuilder.buildInstr(SPIRV::OpStore).
addUse(Ptr).
addUse(StoreVal);
2033 if (!
I.getNumMemOperands()) {
2034 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2036 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2045bool SPIRVInstructionSelector::selectAtomicStore(MachineInstr &
I)
const {
2046 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2049 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2050 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2055 return diagnoseUnsupported(
I,
2056 "Lowering to SPIR-V of atomic store is only "
2057 "allowed for integer or floating point types");
2059 assert(
I.getNumMemOperands());
2060 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2061 assert(MemOp.isAtomic());
2065 Register ScopeReg = buildI32Constant(Scope,
I);
2071 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2072 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2075 MachineIRBuilder MIRBuilder(
I);
2076 auto AtomicStore = MIRBuilder.buildInstr(SPIRV::OpAtomicStore)
2081 AtomicStore.constrainAllUses(
TII,
TRI, RBI);
2085bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
2086 SPIRVTypeInst ResType,
2087 MachineInstr &
I)
const {
2088 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
2096 const Register PtrsReg =
I.getOperand(2).getReg();
2097 const uint32_t Alignment =
I.getOperand(3).getImm();
2098 const Register MaskReg =
I.getOperand(4).getReg();
2099 const Register PassthruReg =
I.getOperand(5).getReg();
2100 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2104 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
2115bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
2116 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
2123 const Register ValuesReg =
I.getOperand(1).getReg();
2124 const Register PtrsReg =
I.getOperand(2).getReg();
2125 const uint32_t Alignment =
I.getOperand(3).getImm();
2126 const Register MaskReg =
I.getOperand(4).getReg();
2127 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2131 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
2140bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
2141 const Twine &Msg)
const {
2142 const Function &
F =
I.getMF()->getFunction();
2143 F.getContext().diagnose(
2144 DiagnosticInfoUnsupported(
F, Msg,
I.getDebugLoc(),
DS_Error));
2148bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
2149 SPIRVTypeInst ResType,
2150 MachineInstr &
I)
const {
2151 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2152 return diagnoseUnsupported(
2153 I,
"llvm.stacksave intrinsic: this instruction requires the following "
2154 "SPIR-V extension: SPV_INTEL_variable_length_array");
2156 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
2163bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
2164 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2165 return diagnoseUnsupported(
2167 "llvm.stackrestore intrinsic: this instruction requires the following "
2168 "SPIR-V extension: SPV_INTEL_variable_length_array");
2169 if (!
I.getOperand(0).isReg())
2172 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
2173 .
addUse(
I.getOperand(0).getReg())
2179SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
2180 MachineIRBuilder MIRBuilder(
I);
2181 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
2188 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
2192 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2193 Type *ArrTy = ArrayType::get(ValTy, Num);
2195 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
2198 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
2205 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
2208 .
addImm(SPIRV::StorageClass::UniformConstant)
2219bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
2222 Register DstReg =
I.getOperand(0).getReg();
2226 return diagnoseUnsupported(
2227 I,
"OpCopyMemory requires operands to have the same type");
2228 uint64_t CopySize =
getIConstVal(
I.getOperand(2).getReg(), MRI);
2232 return diagnoseUnsupported(
2233 I,
"Unable to determine pointee type size for OpCopyMemory");
2234 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
2235 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
2236 return diagnoseUnsupported(
2237 I,
"OpCopyMemory requires the size to match the pointee type size");
2238 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemory))
2241 if (
I.getNumMemOperands()) {
2242 MachineIRBuilder MIRBuilder(
I);
2249bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &
I,
2252 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemorySized))
2253 .
addUse(
I.getOperand(0).getReg())
2255 .
addUse(
I.getOperand(2).getReg());
2256 if (
I.getNumMemOperands()) {
2257 MachineIRBuilder MIRBuilder(
I);
2264bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
2265 MachineInstr &
I)
const {
2266 Register SrcReg =
I.getOperand(1).getReg();
2267 if (
I.getOpcode() == TargetOpcode::G_MEMSET) {
2268 Register VarReg = getOrCreateMemSetGlobal(
I);
2271 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2273 ValTy,
I, SPIRV::StorageClass::UniformConstant);
2275 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
2279 if (!selectCopyMemory(
I, SrcReg))
2282 if (!selectCopyMemorySized(
I, SrcReg))
2285 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
2286 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
2291bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
2292 SPIRVTypeInst ResType,
2295 unsigned NegateOpcode)
const {
2297 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2300 Register ScopeReg = buildI32Constant(Scope,
I);
2302 Register Ptr =
I.getOperand(1).getReg();
2303 uint32_t ScSem =
static_cast<uint32_t
>(
2307 Register MemSemReg = buildI32Constant(MemSem,
I);
2309 Register ValueReg =
I.getOperand(2).getReg();
2310 if (NegateOpcode != 0) {
2313 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
2318 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2329bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2330 unsigned ArgI =
I.getNumOperands() - 1;
2332 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2333 SPIRVTypeInst SrcType =
2335 if (!SrcType || SrcType->
getOpcode() != SPIRV::OpTypeVector)
2337 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2341 unsigned CurrentIndex = 0;
2342 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2343 Register ResVReg =
I.getOperand(i).getReg();
2346 LLT ResLLT = MRI->
getType(ResVReg);
2352 ResType = ScalarType;
2358 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
2361 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2367 for (
unsigned j = 0;
j < NumElements; ++
j) {
2368 MIB.
addImm(CurrentIndex + j);
2370 CurrentIndex += NumElements;
2374 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2386bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2389 Register MemSemReg = buildI32Constant(MemSem,
I);
2391 uint32_t
Scope =
static_cast<uint32_t
>(
2393 Register ScopeReg = buildI32Constant(Scope,
I);
2395 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2402bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2403 SPIRVTypeInst ResType,
2405 unsigned Opcode)
const {
2406 Type *ResTy =
nullptr;
2409 return diagnoseUnsupported(
2411 "Not enough info to select the arithmetic with overflow instruction");
2413 return diagnoseUnsupported(
I,
2414 "Expect struct type result for the arithmetic "
2415 "with overflow instruction");
2421 MachineIRBuilder MIRBuilder(
I);
2423 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2424 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2430 Register ZeroReg = buildZerosVal(ResType,
I);
2435 if (ResName.
size() > 0)
2440 BuildMI(BB, MIRBuilder.getInsertPt(),
I.getDebugLoc(),
TII.get(Opcode))
2443 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2444 MIB.
addUse(
I.getOperand(i).getReg());
2449 MRI->
setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2450 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2452 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2453 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2460 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2461 .
addDef(
I.getOperand(1).getReg())
2469bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2470 SPIRVTypeInst ResType,
2471 MachineInstr &
I)
const {
2473 "selectAtomicCmpXchg only handles the spv_cmpxchg intrinsic");
2474 Register Ptr =
I.getOperand(2).getReg();
2475 Register ScopeReg =
I.getOperand(5).getReg();
2476 Register MemSemEqReg =
I.getOperand(6).getReg();
2477 Register MemSemNeqReg =
I.getOperand(7).getReg();
2479 Register Val =
I.getOperand(4).getReg();
2483 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2502 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2509 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2521 case SPIRV::StorageClass::DeviceOnlyINTEL:
2522 case SPIRV::StorageClass::HostOnlyINTEL:
2531 bool IsGRef =
false;
2532 bool IsAllowedRefs =
2534 unsigned Opcode = It.getOpcode();
2535 if (Opcode == SPIRV::OpConstantComposite ||
2536 Opcode == SPIRV::OpSpecConstantComposite ||
2537 Opcode == SPIRV::OpVariable ||
2538 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2539 return IsGRef = true;
2540 return Opcode == SPIRV::OpName;
2542 return IsAllowedRefs && IsGRef;
2545Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2546 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2548 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2552SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2554 uint32_t Opcode)
const {
2555 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2556 TII.get(SPIRV::OpSpecConstantOp))
2564SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2565 SPIRVTypeInst SrcPtrTy)
const {
2566 SPIRVTypeInst GenericPtrTy =
2570 SPIRV::StorageClass::Generic),
2572 MachineFunction *MF =
I.getParent()->getParent();
2574 MachineInstrBuilder MIB = buildSpecConstantOp(
2576 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2586bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2587 SPIRVTypeInst ResType,
2588 MachineInstr &
I)
const {
2592 Register SrcPtr =
I.getOperand(1).getReg();
2596 if (SrcPtrTy->
getOpcode() != SPIRV::OpTypePointer ||
2597 ResType->
getOpcode() != SPIRV::OpTypePointer)
2598 return BuildCOPY(ResVReg, SrcPtr,
I);
2608 unsigned SpecOpcode =
2610 ?
static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric)
2613 ? static_cast<uint32_t>(
SPIRV::Opcode::GenericCastToPtr)
2620 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2622 .constrainAllUses(
TII,
TRI, RBI);
2624 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2626 buildSpecConstantOp(
2628 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2629 .constrainAllUses(
TII,
TRI, RBI);
2636 return BuildCOPY(ResVReg, SrcPtr,
I);
2638 if ((SrcSC == SPIRV::StorageClass::Function &&
2639 DstSC == SPIRV::StorageClass::Private) ||
2640 (DstSC == SPIRV::StorageClass::Function &&
2641 SrcSC == SPIRV::StorageClass::Private))
2642 return BuildCOPY(ResVReg, SrcPtr,
I);
2646 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2649 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2652 SPIRVTypeInst GenericPtrTy =
2671 return selectUnOp(ResVReg, ResType,
I,
2672 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
2674 return selectUnOp(ResVReg, ResType,
I,
2675 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
2677 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2679 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2689 return SPIRV::OpFOrdEqual;
2691 return SPIRV::OpFOrdGreaterThanEqual;
2693 return SPIRV::OpFOrdGreaterThan;
2695 return SPIRV::OpFOrdLessThanEqual;
2697 return SPIRV::OpFOrdLessThan;
2699 return SPIRV::OpFOrdNotEqual;
2701 return SPIRV::OpOrdered;
2703 return SPIRV::OpFUnordEqual;
2705 return SPIRV::OpFUnordGreaterThanEqual;
2707 return SPIRV::OpFUnordGreaterThan;
2709 return SPIRV::OpFUnordLessThanEqual;
2711 return SPIRV::OpFUnordLessThan;
2713 return SPIRV::OpFUnordNotEqual;
2715 return SPIRV::OpUnordered;
2725 return SPIRV::OpIEqual;
2727 return SPIRV::OpINotEqual;
2729 return SPIRV::OpSGreaterThanEqual;
2731 return SPIRV::OpSGreaterThan;
2733 return SPIRV::OpSLessThanEqual;
2735 return SPIRV::OpSLessThan;
2737 return SPIRV::OpUGreaterThanEqual;
2739 return SPIRV::OpUGreaterThan;
2741 return SPIRV::OpULessThanEqual;
2743 return SPIRV::OpULessThan;
2752 return SPIRV::OpPtrEqual;
2754 return SPIRV::OpPtrNotEqual;
2765 return SPIRV::OpLogicalEqual;
2767 return SPIRV::OpLogicalNotEqual;
2801bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
2802 SPIRVTypeInst ResType,
2804 unsigned OpAnyOrAll)
const {
2805 assert(
I.getNumOperands() == 3);
2806 assert(
I.getOperand(2).isReg());
2808 Register InputRegister =
I.getOperand(2).getReg();
2811 assert(InputType &&
"VReg has no type assigned");
2814 bool IsVectorTy = InputType->
getOpcode() == SPIRV::OpTypeVector;
2815 if (IsBoolTy && !IsVectorTy) {
2816 assert(ResVReg ==
I.getOperand(0).getReg());
2817 return BuildCOPY(ResVReg, InputRegister,
I);
2821 unsigned SpirvNotEqualId =
2822 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
2824 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
2829 IsBoolTy ? InputRegister
2837 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
2839 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
2856bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
2857 SPIRVTypeInst ResType,
2858 MachineInstr &
I)
const {
2859 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
2862bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
2863 SPIRVTypeInst ResType,
2864 MachineInstr &
I)
const {
2865 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
2869bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
2870 SPIRVTypeInst ResType,
2871 MachineInstr &
I)
const {
2872 assert(
I.getNumOperands() == 4);
2873 assert(
I.getOperand(2).isReg());
2874 assert(
I.getOperand(3).isReg());
2876 [[maybe_unused]] SPIRVTypeInst VecType =
2881 "dot product requires a vector of at least 2 components");
2883 [[maybe_unused]] SPIRVTypeInst EltType =
2892 .
addUse(
I.getOperand(2).getReg())
2893 .
addUse(
I.getOperand(3).getReg())
2898bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
2899 SPIRVTypeInst ResType,
2902 assert(
I.getNumOperands() == 4);
2903 assert(
I.getOperand(2).isReg());
2904 assert(
I.getOperand(3).isReg());
2907 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
2911 .
addUse(
I.getOperand(2).getReg())
2912 .
addUse(
I.getOperand(3).getReg())
2919bool SPIRVInstructionSelector::selectIntegerDotExpansion(
2920 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
2921 assert(
I.getNumOperands() == 4);
2922 assert(
I.getOperand(2).isReg());
2923 assert(
I.getOperand(3).isReg());
2927 Register Vec0 =
I.getOperand(2).getReg();
2928 Register Vec1 =
I.getOperand(3).getReg();
2932 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
2941 "dot product requires a vector of at least 2 components");
2944 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2954 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2965 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
2977bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
2978 SPIRVTypeInst ResType,
2979 MachineInstr &
I)
const {
2981 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
2984 .
addUse(
I.getOperand(2).getReg())
2989bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
2990 SPIRVTypeInst ResType,
2991 MachineInstr &
I)
const {
2993 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
2996 .
addUse(
I.getOperand(2).getReg())
3001bool SPIRVInstructionSelector::selectOpIsFinite(
Register ResVReg,
3002 SPIRVTypeInst ResType,
3003 MachineInstr &
I)
const {
3005 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsFinite))
3008 .
addUse(
I.getOperand(2).getReg())
3013bool SPIRVInstructionSelector::selectOpIsNormal(
Register ResVReg,
3014 SPIRVTypeInst ResType,
3015 MachineInstr &
I)
const {
3017 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNormal))
3020 .
addUse(
I.getOperand(2).getReg())
3025template <
bool Signed>
3026bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
3027 SPIRVTypeInst ResType,
3028 MachineInstr &
I)
const {
3029 assert(
I.getNumOperands() == 5);
3030 assert(
I.getOperand(2).isReg());
3031 assert(
I.getOperand(3).isReg());
3032 assert(
I.getOperand(4).isReg());
3035 Register Acc =
I.getOperand(2).getReg();
3039 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3041 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
3046 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
3049 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3061template <
bool Signed>
3062bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
3063 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3064 assert(
I.getNumOperands() == 5);
3065 assert(
I.getOperand(2).isReg());
3066 assert(
I.getOperand(3).isReg());
3067 assert(
I.getOperand(4).isReg());
3070 Register Acc =
I.getOperand(2).getReg();
3076 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
3080 for (
unsigned i = 0; i < 4; i++) {
3103 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
3123 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3138bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
3139 SPIRVTypeInst ResType,
3140 MachineInstr &
I)
const {
3141 assert(
I.getNumOperands() == 3);
3142 assert(
I.getOperand(2).isReg());
3144 Register VZero = buildZerosValF(ResType,
I);
3145 Register VOne = buildOnesValF(ResType,
I);
3147 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
3150 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3152 .
addUse(
I.getOperand(2).getReg())
3159bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
3160 SPIRVTypeInst ResType,
3161 MachineInstr &
I)
const {
3162 assert(
I.getNumOperands() == 3);
3163 assert(
I.getOperand(2).isReg());
3165 Register InputRegister =
I.getOperand(2).getReg();
3167 auto &
DL =
I.getDebugLoc();
3170 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3177 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
3179 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
3187 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3192 if (NeedsConversion) {
3193 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
3204bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
3205 SPIRVTypeInst ResType,
3207 unsigned Opcode)
const {
3211 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
3217 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
3218 BMI.addUse(
I.getOperand(J).getReg());
3225bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
3228 bool WithGroupSync)
const {
3230 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
3232 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
3234 assert(((Scope != SPIRV::Scope::Workgroup) ||
3235 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
3236 "Workgroup Scope must set WorkGroupMemory semantic "
3237 "in Barrier instruction");
3239 assert(((Scope != SPIRV::Scope::Device) ||
3240 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
3241 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
3242 "Device Scope must set UniformMemory and ImageMemory semantic "
3243 "in Barrier instruction");
3249 if (WithGroupSync) {
3250 Register ExecReg = buildI32Constant(SPIRV::Scope::Workgroup,
I);
3254 Register ScopeReg = buildI32Constant(Scope,
I);
3255 Register MemSemReg = buildI32Constant(MemSem,
I);
3257 MI.addUse(ScopeReg).addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
3261bool SPIRVInstructionSelector::selectWaveActiveCountBits(
3262 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3267 if (!selectWaveOpInst(BallotReg, BallotType,
I,
3268 SPIRV::OpGroupNonUniformBallot))
3273 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3278 .
addImm(SPIRV::GroupOperation::Reduce)
3285bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
3286 SPIRVTypeInst ResType,
3287 MachineInstr &
I)
const {
3292 Register InputReg =
I.getOperand(2).getReg();
3297 bool IsVector = NumElems > 1;
3310 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
3311 SPIRV::OpGroupNonUniformAllEqual);
3316 ElementResults.
reserve(NumElems);
3318 for (
unsigned Idx = 0; Idx < NumElems; ++Idx) {
3331 ElemInput = Extracted;
3337 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3348 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3359bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3360 SPIRVTypeInst ResType,
3361 MachineInstr &
I)
const {
3363 assert(
I.getNumOperands() == 3);
3365 auto Op =
I.getOperand(2);
3375 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3377 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3378 return diagnoseUnsupported(
I,
"WavePrefixBitCount requires boolean input");
3399 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3403 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3410bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3411 SPIRVTypeInst ResType,
3413 bool IsUnsigned)
const {
3414 return selectWaveReduce(
3415 ResVReg, ResType,
I, IsUnsigned,
3416 [&](
Register InputRegister,
bool IsUnsigned) {
3417 const bool IsFloatTy =
3419 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3420 : SPIRV::OpGroupNonUniformSMax;
3421 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3425bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3426 SPIRVTypeInst ResType,
3428 bool IsUnsigned)
const {
3429 return selectWaveReduce(
3430 ResVReg, ResType,
I, IsUnsigned,
3431 [&](
Register InputRegister,
bool IsUnsigned) {
3432 const bool IsFloatTy =
3434 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3435 : SPIRV::OpGroupNonUniformSMin;
3436 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3440bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3441 SPIRVTypeInst ResType,
3442 MachineInstr &
I)
const {
3443 return selectWaveReduce(ResVReg, ResType,
I,
false,
3444 [&](
Register InputRegister,
bool IsUnsigned) {
3446 InputRegister, SPIRV::OpTypeFloat);
3447 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3448 : SPIRV::OpGroupNonUniformIAdd;
3452bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3453 SPIRVTypeInst ResType,
3454 MachineInstr &
I)
const {
3455 return selectWaveReduce(ResVReg, ResType,
I,
false,
3456 [&](
Register InputRegister,
bool IsUnsigned) {
3458 InputRegister, SPIRV::OpTypeFloat);
3459 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3460 : SPIRV::OpGroupNonUniformIMul;
3464template <
typename PickOpcodeFn>
3465bool SPIRVInstructionSelector::selectWaveReduce(
3466 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3467 PickOpcodeFn &&PickOpcode)
const {
3468 assert(
I.getNumOperands() == 3);
3469 assert(
I.getOperand(2).isReg());
3471 Register InputRegister =
I.getOperand(2).getReg();
3475 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3478 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3484 .
addImm(SPIRV::GroupOperation::Reduce)
3485 .
addUse(
I.getOperand(2).getReg())
3490bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3491 SPIRVTypeInst ResType,
3493 unsigned Opcode)
const {
3494 return selectWaveReduce(
3495 ResVReg, ResType,
I,
false,
3496 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3499bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3500 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3501 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3502 [&](
Register InputRegister,
bool IsUnsigned) {
3504 InputRegister, SPIRV::OpTypeFloat);
3506 ? SPIRV::OpGroupNonUniformFAdd
3507 : SPIRV::OpGroupNonUniformIAdd;
3511bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3512 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3513 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3514 [&](
Register InputRegister,
bool IsUnsigned) {
3516 InputRegister, SPIRV::OpTypeFloat);
3518 ? SPIRV::OpGroupNonUniformFMul
3519 : SPIRV::OpGroupNonUniformIMul;
3523template <
typename PickOpcodeFn>
3524bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3525 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3526 PickOpcodeFn &&PickOpcode)
const {
3527 assert(
I.getNumOperands() == 3);
3528 assert(
I.getOperand(2).isReg());
3530 Register InputRegister =
I.getOperand(2).getReg();
3534 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3537 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3543 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3544 .
addUse(
I.getOperand(2).getReg())
3549bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
3550 SPIRVTypeInst ResType,
3553 assert(
I.getNumOperands() == 3);
3554 assert(
I.getOperand(2).isReg());
3556 Register InputRegister =
I.getOperand(2).getReg();
3562 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
3573bool SPIRVInstructionSelector::selectBitreverse16(
Register ResVReg,
3574 SPIRVTypeInst ResType,
3579 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
3584 : SPIRV::OpUConvert;
3588 ShiftOp = SPIRV::OpShiftRightLogicalV;
3593 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3594 TII.get(SPIRV::OpConstantComposite))
3597 for (
unsigned It = 0; It <
N; ++It)
3601 ShiftConst = CompositeReg;
3606 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
3611 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
3616 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
3621 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
3624bool SPIRVInstructionSelector::handle64BitOverflow(
3626 unsigned int Opcode,
3633 "handle64BitOverflow should only be used for integer types");
3635 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
3637 MachineIRBuilder MIRBuilder(
I);
3639 SPIRVTypeInst I64x2Type =
3641 SPIRVTypeInst Vec2ResType =
3644 std::vector<Register> PartialRegs;
3646 unsigned CurrentComponent = 0;
3647 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
3651 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3652 TII.get(SPIRV::OpVectorShuffle))
3657 .
addImm(CurrentComponent)
3658 .
addImm(CurrentComponent + 1);
3668 PartialRegs.push_back(SubVecReg);
3671 if (CurrentComponent != ComponentCount) {
3677 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
3678 SPIRV::OpVectorExtractDynamic))
3687 PartialRegs.push_back(FinalElemResReg);
3691 return selectOpWithSrcs(ResVReg, ResType,
I, std::move(PartialRegs),
3692 SPIRV::OpCompositeConstruct);
3695bool SPIRVInstructionSelector::selectBitreverse64(
Register ResVReg,
3696 SPIRVTypeInst ResType,
3700 if (ComponentCount > 2)
3701 return handle64BitOverflow(
3702 ResVReg, ResType,
I, SrcReg, SPIRV::OpBitReverse,
3704 unsigned O) {
return this->selectBitreverse64(R,
T,
I, S); });
3706 MachineIRBuilder MIRBuilder(
I);
3710 I32Type, 2 * ComponentCount, MIRBuilder,
false);
3714 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
3719 if (!selectBitreverseNative(Reverse32, VecI32Type,
I, Vec32))
3726 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3727 TII.get(SPIRV::OpVectorShuffle))
3732 for (
unsigned J = 0; J < ComponentCount; ++J) {
3739 return selectOpWithSrcs(ResVReg, ResType,
I, {SwappedVec}, SPIRV::OpBitcast);
3742bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
3743 SPIRVTypeInst ResType,
3747 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
3755bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
3756 SPIRVTypeInst ResType,
3757 MachineInstr &
I)
const {
3758 Register OpReg =
I.getOperand(1).getReg();
3766 return selectBitreverse16(ResVReg, ResType,
I, OpReg);
3768 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
3770 return selectBitreverse64(ResVReg, ResType,
I, OpReg);
3772 return SPIRVInstructionSelector::diagnoseUnsupported(
3773 I,
"G_BITREVERSE only support 16,32,64 bits.");
3777 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
3788 unsigned AndOp = SPIRV::OpBitwiseAndS;
3789 unsigned OrOp = SPIRV::OpBitwiseOrS;
3790 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
3791 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
3793 AndOp = SPIRV::OpBitwiseAndV;
3794 OrOp = SPIRV::OpBitwiseOrV;
3795 ShlOp = SPIRV::OpShiftLeftLogicalV;
3796 ShrOp = SPIRV::OpShiftRightLogicalV;
3802 const unsigned Shift) ->
Register {
3810 Register MaskReg = CreateConst(Mask);
3811 Register ShiftReg = CreateConst(Shift);
3818 if (!selectOpWithSrcs(
T1, ResType,
I, {Input, ShiftReg}, ShrOp) ||
3819 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
3820 !selectOpWithSrcs(T3, ResType,
I, {Input, MaskReg}, AndOp) ||
3821 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
3822 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
3830 uint64_t
Mask = ~0ull;
3831 while ((Shift >>= 1) > 0) {
3838 return BuildCOPY(ResVReg, Result,
I);
3841bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
3842 SPIRVTypeInst ResType,
3843 MachineInstr &
I)
const {
3844 assert(
I.getOperand(0).isReg() &&
I.getOperand(1).isReg() &&
3845 "G_FREEZE must define and use a register");
3846 Register OpReg =
I.getOperand(1).getReg();
3850 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
3863 if (MachineInstr *Def = MRI->
getVRegDef(OpReg)) {
3864 if (
Def->getOpcode() == TargetOpcode::COPY)
3867 switch (
Def->getOpcode()) {
3868 case SPIRV::ASSIGN_TYPE:
3869 if (MachineInstr *AssignToDef =
3871 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
3872 Reg =
Def->getOperand(2).getReg();
3875 case SPIRV::OpUndef:
3876 Reg =
Def->getOperand(1).getReg();
3879 unsigned DestOpCode;
3881 DestOpCode = SPIRV::OpConstantNull;
3882 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze of a "
3883 "static undef/poison lowered to OpConstantNull\n");
3885 DestOpCode = TargetOpcode::COPY;
3887 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze "
3888 "skipped, lowered as a copy of the operand\n");
3890 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
3891 .
addDef(
I.getOperand(0).getReg())
3899bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
3900 SPIRVTypeInst ResType,
3901 MachineInstr &
I)
const {
3903 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3905 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
3909 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
3914 for (
unsigned i =
I.getNumExplicitDefs();
3915 i <
I.getNumExplicitOperands() && IsConst; ++i)
3919 if (!IsConst &&
N < 2)
3920 return diagnoseUnsupported(
3921 I,
"There must be at least two constituent operands in a vector");
3924 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3925 TII.get(IsConst ? SPIRV::OpConstantComposite
3926 : SPIRV::OpCompositeConstruct))
3929 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
3930 MIB.
addUse(
I.getOperand(i).getReg());
3935bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
3936 SPIRVTypeInst ResType,
3937 MachineInstr &
I)
const {
3939 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
3941 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
3947 if (!
I.getOperand(
OpIdx).isReg())
3954 if (!IsConst &&
N < 2)
3955 return diagnoseUnsupported(
3956 I,
"There must be at least two constituent operands in a vector");
3959 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3960 TII.get(IsConst ? SPIRV::OpConstantComposite
3961 : SPIRV::OpCompositeConstruct))
3964 for (
unsigned i = 0; i <
N; ++i)
3970bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
3971 SPIRVTypeInst ResType,
3972 MachineInstr &
I)
const {
3977 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
3979 Opcode = SPIRV::OpDemoteToHelperInvocation;
3981 Opcode = SPIRV::OpKill;
3983 if (MachineInstr *NextI =
I.getNextNode()) {
3985 NextI->eraseFromParent();
3995bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
3996 SPIRVTypeInst ResType,
unsigned CmpOpc,
3997 MachineInstr &
I)
const {
3998 Register Cmp0 =
I.getOperand(2).getReg();
3999 Register Cmp1 =
I.getOperand(3).getReg();
4002 "CMP operands should have the same type");
4003 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
4013bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
4014 SPIRVTypeInst ResType,
4015 MachineInstr &
I)
const {
4016 auto Pred =
I.getOperand(1).getPredicate();
4019 Register CmpOperand =
I.getOperand(2).getReg();
4024 Register Op1 =
I.getOperand(3).getReg();
4028 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
4033 I.getOperand(3).setReg(NewOp1);
4039 return selectCmp(ResVReg, ResType, CmpOpc,
I);
4043SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
4044 SPIRVTypeInst ResType)
const {
4046 SPIRVTypeInst SpvI32Ty =
4049 auto ConstInt = ConstantInt::get(LLVMTy, Val);
4056 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4059 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
4062 .
addImm(APInt(32, Val).getZExtValue());
4064 GR.
add(ConstInt,
MI);
4071Register SPIRVInstructionSelector::buildI32ConstantInEntryBlock(
4072 uint32_t Val, MachineInstr &
I, SPIRVTypeInst ResType)
const {
4074 SPIRVTypeInst SpvI32Ty =
4076 auto *ConstInt = ConstantInt::get(LLVMTy, Val);
4081 MachineBasicBlock &EntryBB = *InsertIt->getParent();
4082 MachineInstr *
MI =
nullptr;
4086 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantNull))
4090 uint64_t ImmVal = APInt(32, Val).getZExtValue();
4091 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantI))
4097 GR.
add(ConstInt,
MI);
4102bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
4103 SPIRVTypeInst ResType,
4104 MachineInstr &
I)
const {
4106 return selectCmp(ResVReg, ResType, CmpOp,
I);
4109bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
4110 SPIRVTypeInst ResType,
4111 MachineInstr &
I)
const {
4113 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
4120 if (ResType->
getOpcode() != SPIRV::OpTypeVector &&
4121 ResType->
getOpcode() != SPIRV::OpTypeFloat)
4124 MachineIRBuilder MIRBuilder(
I);
4131 APFloat ConstVal(3.3219280948873623);
4135 APFloat::rmNearestTiesToEven, &LosesInfo);
4139 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
4140 ? SPIRV::OpVectorTimesScalar
4143 if (!selectOpWithSrcs(ArgReg, ResType,
I,
4144 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
4146 if (!selectExtInst(ResVReg, ResType,
I,
4147 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
4157Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
4158 MachineInstr &
I)
const {
4161 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4166bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
4172 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
4180 if (
Def->getOpcode() == SPIRV::OpConstantNull)
4183 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
4184 Def->getOpcode() == SPIRV::OpConstantI)
4197 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
4198 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
4200 Intrinsic::spv_const_composite)) {
4201 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
4202 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
4203 if (!IsZero(
Def->getOperand(i).getReg()))
4212Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
4213 MachineInstr &
I)
const {
4217 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4222Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
4223 MachineInstr &
I)
const {
4227 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4233 SPIRVTypeInst ResType,
4234 MachineInstr &
I)
const {
4238 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4243bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
4244 SPIRVTypeInst ResType,
4245 MachineInstr &
I)
const {
4246 Register SelectFirstArg =
I.getOperand(2).getReg();
4247 Register SelectSecondArg =
I.getOperand(3).getReg();
4256 SPIRV::OpTypeVector;
4263 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
4264 }
else if (IsPtrTy) {
4265 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
4267 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
4270 assert(IsScalarBool &&
"OpSelect with a scalar result requires a scalar "
4271 "boolean condition");
4273 Opcode = SPIRV::OpSelectSFSCond;
4274 }
else if (IsPtrTy) {
4275 Opcode = SPIRV::OpSelectSPSCond;
4277 Opcode = SPIRV::OpSelectSISCond;
4280 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4283 .
addUse(
I.getOperand(1).getReg())
4292bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
4293 SPIRVTypeInst ResType,
4295 MachineInstr &InsertAt,
4296 bool IsSigned)
const {
4298 Register ZeroReg = buildZerosVal(ResType, InsertAt);
4299 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
4300 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
4302 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4314bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
4315 SPIRVTypeInst ResType,
4316 MachineInstr &
I,
bool IsSigned,
4317 unsigned Opcode)
const {
4318 Register SrcReg =
I.getOperand(1).getReg();
4324 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
4329 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I,
false);
4331 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
4334bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
4335 SPIRVTypeInst ResType, MachineInstr &
I,
4336 bool IsSigned)
const {
4337 Register SrcReg =
I.getOperand(1).getReg();
4339 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
4343 if (ResType == SrcType)
4344 return BuildCOPY(ResVReg, SrcReg,
I);
4346 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4347 return selectUnOp(ResVReg, ResType,
I, Opcode);
4350bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
4351 SPIRVTypeInst ResType,
4353 bool IsSigned)
const {
4354 MachineIRBuilder MIRBuilder(
I);
4355 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
4367 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
4370 .
addUse(
I.getOperand(1).getReg())
4371 .
addUse(
I.getOperand(2).getReg())
4376 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
4379 .
addUse(
I.getOperand(1).getReg())
4380 .
addUse(
I.getOperand(2).getReg())
4388 unsigned SelectOpcode =
4389 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
4394 .
addUse(buildOnesVal(
true, ResType,
I))
4395 .
addUse(buildZerosVal(ResType,
I))
4402 .
addUse(buildOnesVal(
false, ResType,
I))
4407bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
4410 SPIRVTypeInst IntTy,
4411 SPIRVTypeInst BoolTy)
const {
4414 bool IsVectorTy = IntTy->
getOpcode() == SPIRV::OpTypeVector;
4415 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
4417 Register One = buildOnesVal(
false, IntTy,
I);
4425 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
4434bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
4435 SPIRVTypeInst ResType,
4436 MachineInstr &
I)
const {
4437 Register IntReg =
I.getOperand(1).getReg();
4440 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
4441 if (ArgType == ResType)
4442 return BuildCOPY(ResVReg, IntReg,
I);
4444 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4445 return selectUnOp(ResVReg, ResType,
I, Opcode);
4448bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
4449 SPIRVTypeInst ResType,
4450 MachineInstr &
I)
const {
4451 unsigned Opcode =
I.getOpcode();
4452 unsigned TpOpcode = ResType->
getOpcode();
4454 if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {
4455 assert(Opcode == TargetOpcode::G_CONSTANT &&
4456 I.getOperand(1).getCImm()->isZero());
4457 MachineBasicBlock &DepMBB =
I.getMF()->front();
4460 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
4467 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
4470bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
4471 SPIRVTypeInst ResType,
4472 MachineInstr &
I)
const {
4473 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4480bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
4481 SPIRVTypeInst ResType,
4482 MachineInstr &
I)
const {
4484 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
4488 .
addUse(
I.getOperand(3).getReg())
4490 .
addUse(
I.getOperand(2).getReg());
4491 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
4497bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
4498 SPIRVTypeInst ResType,
4499 MachineInstr &
I)
const {
4500 Type *MaybeResTy =
nullptr;
4505 "Expected aggregate type for extractv instruction");
4507 SPIRV::AccessQualifier::ReadWrite,
false);
4511 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
4514 .
addUse(
I.getOperand(2).getReg());
4515 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
4521bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
4522 SPIRVTypeInst ResType,
4523 MachineInstr &
I)
const {
4524 if (
getImm(
I.getOperand(4), MRI))
4525 return selectInsertVal(ResVReg, ResType,
I);
4527 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
4530 .
addUse(
I.getOperand(2).getReg())
4531 .
addUse(
I.getOperand(3).getReg())
4532 .
addUse(
I.getOperand(4).getReg())
4537bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
4538 SPIRVTypeInst ResType,
4539 MachineInstr &
I)
const {
4540 if (
getImm(
I.getOperand(3), MRI))
4541 return selectExtractVal(ResVReg, ResType,
I);
4543 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
4546 .
addUse(
I.getOperand(2).getReg())
4547 .
addUse(
I.getOperand(3).getReg())
4552bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
4553 SPIRVTypeInst ResType,
4554 MachineInstr &
I)
const {
4555 const bool IsGEPInBounds =
I.getOperand(2).getImm();
4561 ? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
4562 : SPIRV::OpAccessChain)
4563 : (IsGEPInBounds ?
SPIRV::OpInBoundsPtrAccessChain
4564 :
SPIRV::OpPtrAccessChain);
4566 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4570 .
addUse(
I.getOperand(3).getReg());
4572 (Opcode == SPIRV::OpPtrAccessChain ||
4573 Opcode == SPIRV::OpInBoundsPtrAccessChain ||
4574 (
getImm(
I.getOperand(4), MRI) &&
foldImm(
I.getOperand(4), MRI) == 0)) &&
4575 "Cannot translate GEP to OpAccessChain. First index must be 0.");
4578 const unsigned StartingIndex =
4579 (Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)
4582 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
4583 Res.addUse(
I.getOperand(i).getReg());
4584 Res.constrainAllUses(
TII,
TRI, RBI);
4589bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
4591 unsigned Lim =
I.getNumExplicitOperands();
4592 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
4593 Register OpReg =
I.getOperand(i).getReg();
4594 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
4596 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
4597 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
4598 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
4605 MachineFunction *MF =
I.getMF();
4617 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4618 TII.get(SPIRV::OpSpecConstantOp))
4621 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
4623 GR.
add(OpDefine, MIB);
4629bool SPIRVInstructionSelector::selectDerivativeInst(
4630 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
4631 const unsigned DPdOpCode)
const {
4634 if (!errorIfInstrOutsideShader(
I))
4640 Register SrcReg =
I.getOperand(2).getReg();
4645 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4648 .
addUse(
I.getOperand(2).getReg());
4650 MachineIRBuilder MIRBuilder(
I);
4653 if (componentCount != 1)
4657 const TargetRegisterClass *RegClass = GR.
getRegClass(SrcType);
4661 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4666 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4671 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4679bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
4680 SPIRVTypeInst ResType,
4681 MachineInstr &
I)
const {
4685 case Intrinsic::spv_load:
4686 return selectLoad(ResVReg, ResType,
I);
4687 case Intrinsic::spv_atomic_load:
4688 return selectAtomicLoad(ResVReg, ResType,
I);
4689 case Intrinsic::spv_store:
4690 return selectStore(
I);
4691 case Intrinsic::spv_atomic_store:
4692 return selectAtomicStore(
I);
4693 case Intrinsic::spv_extractv:
4694 return selectExtractVal(ResVReg, ResType,
I);
4695 case Intrinsic::spv_insertv:
4696 return selectInsertVal(ResVReg, ResType,
I);
4697 case Intrinsic::spv_extractelt:
4698 return selectExtractElt(ResVReg, ResType,
I);
4699 case Intrinsic::spv_insertelt:
4700 return selectInsertElt(ResVReg, ResType,
I);
4701 case Intrinsic::spv_gep:
4702 return selectGEP(ResVReg, ResType,
I);
4703 case Intrinsic::spv_bitcast: {
4704 Register OpReg =
I.getOperand(2).getReg();
4705 SPIRVTypeInst OpType =
4709 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
4711 case Intrinsic::spv_unref_global:
4712 case Intrinsic::spv_init_global: {
4713 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
4718 Register GVarVReg =
MI->getOperand(0).getReg();
4719 if (!selectGlobalValue(GVarVReg, *
MI, Init))
4724 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
4726 MI->eraseFromParent();
4730 case Intrinsic::spv_undef: {
4731 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4737 case Intrinsic::spv_poison:
4738 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpPoisonKHR))
4743 case Intrinsic::spv_freeze:
4744 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
4747 .
addUse(
I.getOperand(2).getReg())
4750 case Intrinsic::spv_named_boolean_spec_constant: {
4751 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
4752 : SPIRV::OpSpecConstantFalse;
4754 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
4755 .
addDef(
I.getOperand(0).getReg())
4758 unsigned SpecId =
I.getOperand(2).getImm();
4760 SPIRV::Decoration::SpecId, {SpecId});
4764 case Intrinsic::spv_const_composite: {
4766 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
4772 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
4774 std::function<bool(
Register)> HasSpecConstOperand =
4784 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
4785 J < Def->getNumExplicitOperands(); ++J) {
4786 if (
Def->getOperand(J).isReg() &&
4787 HasSpecConstOperand(
Def->getOperand(J).getReg()))
4793 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
4794 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
4795 : SPIRV::OpConstantComposite;
4796 unsigned ContinuedOpc = HasSpecConst
4797 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
4798 : SPIRV::OpConstantCompositeContinuedINTEL;
4799 MachineIRBuilder MIR(
I);
4801 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
4803 for (
auto *Instr : Instructions) {
4804 Instr->setDebugLoc(
I.getDebugLoc());
4809 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4816 case Intrinsic::spv_assign_name: {
4817 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
4818 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
4819 for (
unsigned i =
I.getNumExplicitDefs() + 2;
4820 i <
I.getNumExplicitOperands(); ++i) {
4821 MIB.
addImm(
I.getOperand(i).getImm());
4826 case Intrinsic::spv_switch: {
4827 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
4828 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
4829 if (
I.getOperand(i).isReg())
4830 MIB.
addReg(
I.getOperand(i).getReg());
4831 else if (
I.getOperand(i).isCImm())
4832 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
4833 else if (
I.getOperand(i).isMBB())
4834 MIB.
addMBB(
I.getOperand(i).getMBB());
4841 case Intrinsic::spv_loop_merge: {
4842 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
4843 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
4844 if (
I.getOperand(i).isMBB())
4845 MIB.
addMBB(
I.getOperand(i).getMBB());
4852 case Intrinsic::spv_loop_control_intel: {
4854 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
4855 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
4860 case Intrinsic::spv_selection_merge: {
4862 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
4863 assert(
I.getOperand(1).isMBB() &&
4864 "operand 1 to spv_selection_merge must be a basic block");
4865 MIB.
addMBB(
I.getOperand(1).getMBB());
4866 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
4870 case Intrinsic::spv_cmpxchg:
4871 return selectAtomicCmpXchg(ResVReg, ResType,
I);
4872 case Intrinsic::spv_unreachable:
4873 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
4876 case Intrinsic::spv_abort:
4877 return selectAbort(
I);
4878 case Intrinsic::spv_alloca:
4879 return selectFrameIndex(ResVReg, ResType,
I);
4880 case Intrinsic::spv_alloca_array:
4881 return selectAllocaArray(ResVReg, ResType,
I);
4882 case Intrinsic::spv_assume:
4884 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
4885 .
addUse(
I.getOperand(1).getReg())
4890 case Intrinsic::spv_expect:
4892 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
4895 .
addUse(
I.getOperand(2).getReg())
4896 .
addUse(
I.getOperand(3).getReg())
4901 case Intrinsic::arithmetic_fence:
4902 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
4903 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
4906 .
addUse(
I.getOperand(2).getReg())
4910 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
4912 case Intrinsic::spv_thread_id:
4918 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
4920 case Intrinsic::spv_thread_id_in_group:
4926 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
4928 case Intrinsic::spv_group_id:
4934 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
4936 case Intrinsic::spv_flattened_thread_id_in_group:
4943 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
4945 case Intrinsic::spv_workgroup_size:
4946 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
4948 case Intrinsic::spv_global_size:
4949 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
4951 case Intrinsic::spv_global_offset:
4952 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
4954 case Intrinsic::spv_num_workgroups:
4955 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
4957 case Intrinsic::spv_subgroup_size:
4958 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
4960 case Intrinsic::spv_num_subgroups:
4961 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
4963 case Intrinsic::spv_subgroup_id:
4964 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
4965 case Intrinsic::spv_subgroup_local_invocation_id:
4966 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
4967 ResVReg, ResType,
I);
4968 case Intrinsic::spv_subgroup_max_size:
4969 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
4971 case Intrinsic::spv_fdot:
4972 return selectFloatDot(ResVReg, ResType,
I);
4973 case Intrinsic::spv_udot:
4974 case Intrinsic::spv_sdot:
4975 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4977 return selectIntegerDot(ResVReg, ResType,
I,
4978 IID == Intrinsic::spv_sdot);
4979 return selectIntegerDotExpansion(ResVReg, ResType,
I);
4980 case Intrinsic::spv_dot4add_i8packed:
4981 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4983 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
4984 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
4985 case Intrinsic::spv_dot4add_u8packed:
4986 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
4988 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
4989 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
4990 case Intrinsic::spv_all:
4991 return selectAll(ResVReg, ResType,
I);
4992 case Intrinsic::spv_any:
4993 return selectAny(ResVReg, ResType,
I);
4994 case Intrinsic::spv_cross:
4995 return selectExtInst(ResVReg, ResType,
I, CL::cross, GL::Cross);
4996 case Intrinsic::spv_distance:
4997 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
4998 case Intrinsic::spv_lerp:
4999 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
5000 case Intrinsic::spv_length:
5001 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
5002 case Intrinsic::spv_degrees:
5003 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
5004 case Intrinsic::spv_faceforward:
5005 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
5006 case Intrinsic::spv_frac:
5007 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
5008 case Intrinsic::spv_isinf:
5009 return selectOpIsInf(ResVReg, ResType,
I);
5010 case Intrinsic::spv_isnan:
5011 return selectOpIsNan(ResVReg, ResType,
I);
5012 case Intrinsic::spv_isfinite:
5013 return selectOpIsFinite(ResVReg, ResType,
I);
5014 case Intrinsic::spv_isnormal:
5015 return selectOpIsNormal(ResVReg, ResType,
I);
5016 case Intrinsic::spv_normalize:
5017 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
5018 case Intrinsic::spv_refract:
5019 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
5020 case Intrinsic::spv_reflect:
5021 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
5022 case Intrinsic::spv_rsqrt:
5023 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
5024 case Intrinsic::spv_sign:
5025 return selectSign(ResVReg, ResType,
I);
5026 case Intrinsic::spv_smoothstep:
5027 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
5028 case Intrinsic::spv_firstbituhigh:
5029 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
5030 case Intrinsic::spv_firstbitshigh:
5031 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
5032 case Intrinsic::spv_firstbitlow:
5033 return selectFirstBitLow(ResVReg, ResType,
I);
5034 case Intrinsic::spv_all_memory_barrier:
5035 return selectBarrierInst(
I, SPIRV::Scope::Device,
5036 SPIRV::MemorySemantics::UniformMemory |
5037 SPIRV::MemorySemantics::ImageMemory |
5038 SPIRV::MemorySemantics::WorkgroupMemory,
5040 case Intrinsic::spv_all_memory_barrier_with_group_sync:
5041 return selectBarrierInst(
I, SPIRV::Scope::Device,
5042 SPIRV::MemorySemantics::UniformMemory |
5043 SPIRV::MemorySemantics::ImageMemory |
5044 SPIRV::MemorySemantics::WorkgroupMemory,
5046 case Intrinsic::spv_device_memory_barrier:
5047 return selectBarrierInst(
I, SPIRV::Scope::Device,
5048 SPIRV::MemorySemantics::UniformMemory |
5049 SPIRV::MemorySemantics::ImageMemory,
5051 case Intrinsic::spv_device_memory_barrier_with_group_sync:
5052 return selectBarrierInst(
I, SPIRV::Scope::Device,
5053 SPIRV::MemorySemantics::UniformMemory |
5054 SPIRV::MemorySemantics::ImageMemory,
5056 case Intrinsic::spv_group_memory_barrier:
5057 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5058 SPIRV::MemorySemantics::WorkgroupMemory,
5060 case Intrinsic::spv_group_memory_barrier_with_group_sync:
5061 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5062 SPIRV::MemorySemantics::WorkgroupMemory,
5064 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
5065 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
5066 SPIRV::StorageClass::StorageClass ResSC =
5069 return diagnoseUnsupported(
I,
"The target storage class is not castable "
5070 "from the Generic storage class");
5071 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
5079 case Intrinsic::spv_lifetime_start:
5080 case Intrinsic::spv_lifetime_end: {
5081 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
5082 : SPIRV::OpLifetimeStop;
5083 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
5084 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
5093 case Intrinsic::spv_saturate:
5094 return selectSaturate(ResVReg, ResType,
I);
5095 case Intrinsic::spv_nclamp:
5096 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
5097 case Intrinsic::spv_uclamp:
5098 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
5099 case Intrinsic::spv_sclamp:
5100 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
5101 case Intrinsic::spv_subgroup_prefix_bit_count:
5102 return selectWavePrefixBitCount(ResVReg, ResType,
I);
5103 case Intrinsic::spv_wave_active_countbits:
5104 return selectWaveActiveCountBits(ResVReg, ResType,
I);
5105 case Intrinsic::spv_wave_all_equal:
5106 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
5107 case Intrinsic::spv_wave_all:
5108 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
5109 case Intrinsic::spv_wave_any:
5110 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
5111 case Intrinsic::spv_subgroup_ballot:
5112 return selectWaveOpInst(ResVReg, ResType,
I,
5113 SPIRV::OpGroupNonUniformBallot);
5114 case Intrinsic::spv_wave_is_first_lane:
5115 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
5116 case Intrinsic::spv_wave_reduce_or:
5117 return selectWaveReduceOp(ResVReg, ResType,
I,
5118 SPIRV::OpGroupNonUniformBitwiseOr);
5119 case Intrinsic::spv_wave_reduce_xor:
5120 return selectWaveReduceOp(ResVReg, ResType,
I,
5121 SPIRV::OpGroupNonUniformBitwiseXor);
5122 case Intrinsic::spv_wave_reduce_and:
5123 return selectWaveReduceOp(ResVReg, ResType,
I,
5124 SPIRV::OpGroupNonUniformBitwiseAnd);
5125 case Intrinsic::spv_wave_reduce_umax:
5126 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
5127 case Intrinsic::spv_wave_reduce_max:
5128 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
5129 case Intrinsic::spv_wave_reduce_umin:
5130 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
5131 case Intrinsic::spv_wave_reduce_min:
5132 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
5133 case Intrinsic::spv_wave_reduce_sum:
5134 return selectWaveReduceSum(ResVReg, ResType,
I);
5135 case Intrinsic::spv_wave_product:
5136 return selectWaveReduceProduct(ResVReg, ResType,
I);
5137 case Intrinsic::spv_wave_readlane:
5138 return selectWaveOpInst(ResVReg, ResType,
I,
5139 SPIRV::OpGroupNonUniformShuffle);
5140 case Intrinsic::spv_wave_prefix_sum:
5141 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
5142 case Intrinsic::spv_wave_prefix_product:
5143 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
5144 case Intrinsic::spv_quad_read_across_x: {
5145 return selectQuadSwap(ResVReg, ResType,
I, 0);
5147 case Intrinsic::spv_quad_read_across_y: {
5148 return selectQuadSwap(ResVReg, ResType,
I, 1);
5150 case Intrinsic::spv_step:
5151 return selectExtInst(ResVReg, ResType,
I, CL::step, GL::Step);
5152 case Intrinsic::spv_radians:
5153 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
5157 case Intrinsic::instrprof_increment:
5158 case Intrinsic::instrprof_increment_step:
5159 case Intrinsic::instrprof_value_profile:
5162 case Intrinsic::spv_value_md:
5164 case Intrinsic::spv_resource_handlefrombinding: {
5165 return selectHandleFromBinding(ResVReg, ResType,
I);
5167 case Intrinsic::spv_resource_counterhandlefrombinding:
5168 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
5169 case Intrinsic::spv_resource_updatecounter:
5170 return selectUpdateCounter(ResVReg, ResType,
I);
5171 case Intrinsic::spv_resource_store_typedbuffer: {
5172 return selectImageWriteIntrinsic(
I);
5174 case Intrinsic::spv_resource_load_typedbuffer: {
5175 return selectReadImageIntrinsic(ResVReg, ResType,
I);
5177 case Intrinsic::spv_resource_load_level: {
5178 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
5180 case Intrinsic::spv_resource_getdimensions_x:
5181 case Intrinsic::spv_resource_getdimensions_xy:
5182 case Intrinsic::spv_resource_getdimensions_xyz: {
5183 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
5185 case Intrinsic::spv_resource_getdimensions_levels_x:
5186 case Intrinsic::spv_resource_getdimensions_levels_xy:
5187 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
5188 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
5190 case Intrinsic::spv_resource_getdimensions_ms_xy:
5191 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
5192 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
5194 case Intrinsic::spv_resource_calculate_lod:
5195 case Intrinsic::spv_resource_calculate_lod_unclamped:
5196 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
5197 case Intrinsic::spv_resource_sample:
5198 case Intrinsic::spv_resource_sample_clamp:
5199 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
5200 case Intrinsic::spv_resource_samplebias:
5201 case Intrinsic::spv_resource_samplebias_clamp:
5202 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
5203 case Intrinsic::spv_resource_samplegrad:
5204 case Intrinsic::spv_resource_samplegrad_clamp:
5205 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
5206 case Intrinsic::spv_resource_samplelevel:
5207 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
5208 case Intrinsic::spv_resource_samplecmp:
5209 case Intrinsic::spv_resource_samplecmp_clamp:
5210 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
5211 case Intrinsic::spv_resource_samplecmplevelzero:
5212 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
5213 case Intrinsic::spv_resource_gather:
5214 case Intrinsic::spv_resource_gather_cmp:
5215 return selectGatherIntrinsic(ResVReg, ResType,
I);
5216 case Intrinsic::spv_resource_getbasepointer:
5217 case Intrinsic::spv_resource_getpointer: {
5218 return selectResourceGetPointer(ResVReg, ResType,
I);
5220 case Intrinsic::spv_pushconstant_getpointer: {
5221 return selectPushConstantGetPointer(ResVReg, ResType,
I);
5223 case Intrinsic::spv_discard: {
5224 return selectDiscard(ResVReg, ResType,
I);
5226 case Intrinsic::spv_resource_nonuniformindex: {
5227 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
5229 case Intrinsic::spv_unpackhalf2x16: {
5230 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
5232 case Intrinsic::spv_packhalf2x16: {
5233 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
5235 case Intrinsic::spv_ddx:
5236 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
5237 case Intrinsic::spv_ddy:
5238 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
5239 case Intrinsic::spv_ddx_coarse:
5240 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
5241 case Intrinsic::spv_ddy_coarse:
5242 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
5243 case Intrinsic::spv_ddx_fine:
5244 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
5245 case Intrinsic::spv_ddy_fine:
5246 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
5247 case Intrinsic::spv_fwidth:
5248 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
5249 case Intrinsic::spv_masked_gather:
5250 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5251 return selectMaskedGather(ResVReg, ResType,
I);
5252 return diagnoseUnsupported(
5253 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
5254 case Intrinsic::spv_masked_scatter:
5255 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5256 return selectMaskedScatter(
I);
5257 return diagnoseUnsupported(
5258 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
5259 case Intrinsic::returnaddress:
5260 case Intrinsic::frameaddress: {
5262 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5269 return diagnoseUnsupported(
I,
"intrinsic selection not implemented.");
5274bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
5275 SPIRVTypeInst ResType,
5276 MachineInstr &
I)
const {
5279 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
5286bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
5287 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5289 assert(Intr.getIntrinsicID() ==
5290 Intrinsic::spv_resource_counterhandlefrombinding);
5293 Register MainHandleReg = Intr.getOperand(2).getReg();
5295 assert(MainHandleDef->getIntrinsicID() ==
5296 Intrinsic::spv_resource_handlefrombinding);
5300 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
5301 Register IndexReg = MainHandleDef->getOperand(5).getReg();
5302 std::string CounterName =
5307 MachineIRBuilder MIRBuilder(
I);
5309 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
5311 ArraySize, IndexReg, CounterName, MIRBuilder);
5313 return BuildCOPY(ResVReg, CounterVarReg,
I);
5316bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
5317 SPIRVTypeInst ResType,
5318 MachineInstr &
I)
const {
5320 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
5322 Register CounterHandleReg = Intr.getOperand(2).getReg();
5323 Register IncrReg = Intr.getOperand(3).getReg();
5330 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
5331 assert(CounterVarPointeeType &&
5332 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
5333 "Counter variable must be a struct");
5335 SPIRV::StorageClass::StorageBuffer &&
5336 "Counter variable must be in the storage buffer storage class");
5338 "Counter variable must have exactly 1 member in the struct");
5339 const SPIRVTypeInst MemberType =
5342 "Counter variable struct must have a single i32 member");
5346 MachineIRBuilder MIRBuilder(
I);
5348 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
5351 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
5357 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5360 .
addUse(CounterHandleReg)
5367 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
5370 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
5373 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
5382 return BuildCOPY(ResVReg, AtomicRes,
I);
5390 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
5398bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
5399 SPIRVTypeInst ResType,
5400 MachineInstr &
I)
const {
5408 Register ImageReg =
I.getOperand(2).getReg();
5416 Register IdxReg =
I.getOperand(3).getReg();
5418 MachineInstr &Pos =
I;
5420 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
5424bool SPIRVInstructionSelector::generateSampleImage(
5427 DebugLoc Loc, MachineInstr &Pos)
const {
5438 if (!loadHandleBeforePosition(NewSamplerReg,
5444 MachineIRBuilder MIRBuilder(Pos);
5457 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
5458 ImOps.Lod.has_value();
5459 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
5460 : SPIRV::OpImageSampleImplicitLod;
5462 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
5463 : SPIRV::OpImageSampleDrefImplicitLod;
5472 MIB.
addUse(*ImOps.Compare);
5474 uint32_t ImageOperands = 0;
5476 ImageOperands |= SPIRV::ImageOperand::Bias;
5478 ImageOperands |= SPIRV::ImageOperand::Lod;
5479 if (ImOps.GradX && ImOps.GradY)
5480 ImageOperands |= SPIRV::ImageOperand::Grad;
5481 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
5483 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
5486 "Non-constant offsets are not supported in sample instructions.");
5490 ImageOperands |= SPIRV::ImageOperand::MinLod;
5492 if (ImageOperands != 0) {
5493 MIB.
addImm(ImageOperands);
5494 if (ImageOperands & SPIRV::ImageOperand::Bias)
5496 if (ImageOperands & SPIRV::ImageOperand::Lod)
5498 if (ImageOperands & SPIRV::ImageOperand::Grad) {
5499 MIB.
addUse(*ImOps.GradX);
5500 MIB.
addUse(*ImOps.GradY);
5503 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
5504 MIB.
addUse(*ImOps.Offset);
5505 if (ImageOperands & SPIRV::ImageOperand::MinLod)
5506 MIB.
addUse(*ImOps.MinLod);
5513bool SPIRVInstructionSelector::selectImageQuerySize(
5515 std::optional<Register> LodReg)
const {
5517 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
5520 "ImageReg is not an image type.");
5522 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
5524 unsigned NumComponents = 0;
5526 case SPIRV::Dim::DIM_1D:
5527 case SPIRV::Dim::DIM_Buffer:
5528 NumComponents =
IsArray ? 2 : 1;
5530 case SPIRV::Dim::DIM_2D:
5531 case SPIRV::Dim::DIM_Cube:
5532 case SPIRV::Dim::DIM_Rect:
5533 NumComponents =
IsArray ? 3 : 2;
5535 case SPIRV::Dim::DIM_3D:
5539 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
5544 SPIRVTypeInst ResType =
5549 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5559bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
5560 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5561 Register ImageReg =
I.getOperand(2).getReg();
5568 return selectImageQuerySize(NewImageReg, ResVReg,
I);
5571bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
5572 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5573 Register ImageReg =
I.getOperand(2).getReg();
5582 Register LodReg =
I.getOperand(3).getReg();
5585 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
5587 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
5594 TII.get(SPIRV::OpImageQueryLevels))
5601 TII.get(SPIRV::OpCompositeConstruct))
5611bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
5612 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5613 Register ImageReg =
I.getOperand(2).getReg();
5624 "OpImageQuerySamples requires a multisampled image");
5626 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
5634 TII.get(SPIRV::OpImageQuerySamples))
5641 TII.get(SPIRV::OpCompositeConstruct))
5651bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
5652 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5653 Register ImageReg =
I.getOperand(2).getReg();
5654 Register SamplerReg =
I.getOperand(3).getReg();
5655 Register CoordinateReg =
I.getOperand(4).getReg();
5671 if (!loadHandleBeforePosition(
5676 MachineIRBuilder MIRBuilder(
I);
5682 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
5692 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
5699 unsigned ExtractedIndex =
5701 Intrinsic::spv_resource_calculate_lod_unclamped
5705 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5706 TII.get(SPIRV::OpCompositeExtract))
5716bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
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 if (
I.getNumOperands() > 5)
5723 ImOps.Offset =
I.getOperand(5).getReg();
5724 if (
I.getNumOperands() > 6)
5725 ImOps.MinLod =
I.getOperand(6).getReg();
5726 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5727 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5730bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
5731 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5732 Register ImageReg =
I.getOperand(2).getReg();
5733 Register SamplerReg =
I.getOperand(3).getReg();
5734 Register CoordinateReg =
I.getOperand(4).getReg();
5735 ImageOperands ImOps;
5736 ImOps.Bias =
I.getOperand(5).getReg();
5737 if (
I.getNumOperands() > 6)
5738 ImOps.Offset =
I.getOperand(6).getReg();
5739 if (
I.getNumOperands() > 7)
5740 ImOps.MinLod =
I.getOperand(7).getReg();
5741 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5742 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5745bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
5746 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5747 Register ImageReg =
I.getOperand(2).getReg();
5748 Register SamplerReg =
I.getOperand(3).getReg();
5749 Register CoordinateReg =
I.getOperand(4).getReg();
5750 ImageOperands ImOps;
5751 ImOps.GradX =
I.getOperand(5).getReg();
5752 ImOps.GradY =
I.getOperand(6).getReg();
5753 if (
I.getNumOperands() > 7)
5754 ImOps.Offset =
I.getOperand(7).getReg();
5755 if (
I.getNumOperands() > 8)
5756 ImOps.MinLod =
I.getOperand(8).getReg();
5757 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5758 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5761bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
5762 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5763 Register ImageReg =
I.getOperand(2).getReg();
5764 Register SamplerReg =
I.getOperand(3).getReg();
5765 Register CoordinateReg =
I.getOperand(4).getReg();
5766 ImageOperands ImOps;
5767 ImOps.Lod =
I.getOperand(5).getReg();
5768 if (
I.getNumOperands() > 6)
5769 ImOps.Offset =
I.getOperand(6).getReg();
5770 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5771 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5774bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
5775 SPIRVTypeInst ResType,
5776 MachineInstr &
I)
const {
5777 Register ImageReg =
I.getOperand(2).getReg();
5778 Register SamplerReg =
I.getOperand(3).getReg();
5779 Register CoordinateReg =
I.getOperand(4).getReg();
5780 ImageOperands ImOps;
5781 ImOps.Compare =
I.getOperand(5).getReg();
5782 if (
I.getNumOperands() > 6)
5783 ImOps.Offset =
I.getOperand(6).getReg();
5784 if (
I.getNumOperands() > 7)
5785 ImOps.MinLod =
I.getOperand(7).getReg();
5786 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5787 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5790bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
5791 SPIRVTypeInst ResType,
5792 MachineInstr &
I)
const {
5793 Register ImageReg =
I.getOperand(2).getReg();
5794 Register CoordinateReg =
I.getOperand(3).getReg();
5795 Register LodReg =
I.getOperand(4).getReg();
5797 ImageOperands ImOps;
5799 if (
I.getNumOperands() > 5)
5800 ImOps.Offset =
I.getOperand(5).getReg();
5812 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
5813 I.getDebugLoc(),
I, &ImOps);
5816bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
5817 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5818 Register ImageReg =
I.getOperand(2).getReg();
5819 Register SamplerReg =
I.getOperand(3).getReg();
5820 Register CoordinateReg =
I.getOperand(4).getReg();
5821 ImageOperands ImOps;
5822 ImOps.Compare =
I.getOperand(5).getReg();
5823 if (
I.getNumOperands() > 6)
5824 ImOps.Offset =
I.getOperand(6).getReg();
5827 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5828 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5831bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
5832 SPIRVTypeInst ResType,
5833 MachineInstr &
I)
const {
5834 Register ImageReg =
I.getOperand(2).getReg();
5835 Register SamplerReg =
I.getOperand(3).getReg();
5836 Register CoordinateReg =
I.getOperand(4).getReg();
5839 "ImageReg is not an image type.");
5844 ComponentOrCompareReg =
I.getOperand(5).getReg();
5845 OffsetReg =
I.getOperand(6).getReg();
5848 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
5852 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
5853 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
5854 Dim != SPIRV::Dim::DIM_Rect) {
5856 "Gather operations are only supported for 2D, Cube, and Rect images.");
5863 if (!loadHandleBeforePosition(
5868 MachineIRBuilder MIRBuilder(
I);
5869 SPIRVTypeInst SampledImageType =
5874 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
5882 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
5884 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
5886 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5891 .
addUse(ComponentOrCompareReg);
5893 uint32_t ImageOperands = 0;
5894 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
5895 if (Dim == SPIRV::Dim::DIM_Cube) {
5897 "Gather operations with offset are not supported for Cube images.");
5901 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
5903 ImageOperands |= SPIRV::ImageOperand::Offset;
5907 if (ImageOperands != 0) {
5908 MIB.
addImm(ImageOperands);
5910 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
5918bool SPIRVInstructionSelector::generateImageReadOrFetch(
5921 const ImageOperands *ImOps)
const {
5924 "ImageReg is not an image type.");
5926 bool IsSignedInteger =
5931 bool IsFetch = (SampledOp.getImm() == 1);
5933 auto AddOperands = [&](MachineInstrBuilder &MIB) {
5934 uint32_t ImageOperandsMask = 0;
5935 if (IsSignedInteger)
5936 ImageOperandsMask |= 0x1000;
5938 if (IsFetch && ImOps) {
5940 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
5941 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
5943 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
5945 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
5949 if (ImageOperandsMask != 0) {
5950 MIB.
addImm(ImageOperandsMask);
5951 if (IsFetch && ImOps) {
5954 if (ImOps->Offset &&
5955 (ImageOperandsMask &
5956 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
5957 MIB.
addUse(*ImOps->Offset);
5963 if (ResultSize == 4) {
5966 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
5973 BMI.constrainAllUses(
TII,
TRI, RBI);
5977 SPIRVTypeInst ReadType = widenTypeToVec4(ResType, Pos);
5981 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
5987 BMI.constrainAllUses(
TII,
TRI, RBI);
5989 if (ResultSize == 1) {
5998 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
6001bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
6002 SPIRVTypeInst ResType,
6003 MachineInstr &
I)
const {
6004 Register ResourcePtr =
I.getOperand(2).getReg();
6006 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
6015 MachineIRBuilder MIRBuilder(
I);
6020 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6026 if (
I.getNumExplicitOperands() > 3) {
6027 Register IndexReg =
I.getOperand(3).getReg();
6034bool SPIRVInstructionSelector::selectPushConstantGetPointer(
6035 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6040bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
6041 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6042 Register ObjReg =
I.getOperand(2).getReg();
6043 if (!BuildCOPY(ResVReg, ObjReg,
I))
6053 decorateUsesAsNonUniform(ResVReg);
6057void SPIRVInstructionSelector::decorateUsesAsNonUniform(
6060 while (WorkList.
size() > 0) {
6064 bool IsDecorated =
false;
6066 if (
Use.getOpcode() == SPIRV::OpDecorate &&
6067 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
6073 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
6075 if (ResultReg == CurrentReg)
6083 SPIRV::Decoration::NonUniformEXT, {});
6088bool SPIRVInstructionSelector::extractSubvector(
6090 MachineInstr &InsertionPoint)
const {
6092 [[maybe_unused]] uint64_t InputSize =
6095 assert(InputSize > 1 &&
"The input must be a vector.");
6096 assert(ResultSize > 1 &&
"The result must be a vector.");
6097 assert(ResultSize < InputSize &&
6098 "Cannot extract more element than there are in the input.");
6101 const TargetRegisterClass *ScalarRegClass = GR.
getRegClass(ScalarType);
6102 for (uint64_t
I = 0;
I < ResultSize;
I++) {
6105 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6114 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
6116 TII.get(SPIRV::OpCompositeConstruct))
6120 for (
Register ComponentReg : ComponentRegisters)
6121 MIB.
addUse(ComponentReg);
6126bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
6127 MachineInstr &
I)
const {
6134 Register ImageReg =
I.getOperand(1).getReg();
6142 Register CoordinateReg =
I.getOperand(2).getReg();
6143 Register DataReg =
I.getOperand(3).getReg();
6146 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
6154Register SPIRVInstructionSelector::buildPointerToResource(
6155 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
6156 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
6157 StringRef Name, MachineIRBuilder MIRBuilder)
const {
6159 if (ArraySize == 1) {
6160 SPIRVTypeInst PtrType =
6163 "SpirvResType did not have an explicit layout.");
6168 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
6169 SPIRVTypeInst VarPointerType =
6172 VarPointerType, Set,
Binding, Name, MIRBuilder);
6174 SPIRVTypeInst ResPointerType =
6187bool SPIRVInstructionSelector::selectFirstBitSet16(
6188 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
6189 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
6191 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
6195 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
6198bool SPIRVInstructionSelector::selectFirstBitSet32(
6200 unsigned BitSetOpcode)
const {
6201 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6204 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6211bool SPIRVInstructionSelector::selectFirstBitSet64(
6213 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
6226 if (ComponentCount > 2) {
6227 auto Func = [
this, SwapPrimarySide](
Register ResVReg, SPIRVTypeInst ResType,
6229 unsigned Opcode) ->
bool {
6230 return this->selectFirstBitSet64(ResVReg, ResType,
I, SrcReg, Opcode,
6234 return handle64BitOverflow(ResVReg, ResType,
I, SrcReg, BitSetOpcode, Func);
6238 MachineIRBuilder MIRBuilder(
I);
6240 BaseType, 2 * ComponentCount, MIRBuilder,
false);
6244 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
6250 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
6257 bool IsScalarRes = ResType->
getOpcode() != SPIRV::OpTypeVector;
6260 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
6261 SPIRV::OpVectorExtractDynamic))
6263 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
6264 SPIRV::OpVectorExtractDynamic))
6268 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6269 TII.get(SPIRV::OpVectorShuffle))
6277 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
6283 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6284 TII.get(SPIRV::OpVectorShuffle))
6292 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
6312 SelectOp = SPIRV::OpSelectSISCond;
6313 AddOp = SPIRV::OpIAddS;
6321 SelectOp = SPIRV::OpSelectVIVCond;
6322 AddOp = SPIRV::OpIAddV;
6328 Register RegSecondaryOffset = Reg0;
6332 if (SwapPrimarySide) {
6333 PrimaryReg = LowReg;
6334 SecondaryReg = HighReg;
6335 RegPrimaryOffset = Reg0;
6336 RegSecondaryOffset = Reg32;
6341 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
6342 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
6347 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
6348 SPIRV::OpINotEqual))
6355 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
6356 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
6361 if (SwapPrimarySide) {
6363 if (!selectOpWithSrcs(RegAdd, ResType,
I,
6364 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
6375 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
6376 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
6381 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
6382 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
6385 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
6389bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
6390 SPIRVTypeInst ResType,
6392 bool IsSigned)
const {
6394 Register OpReg =
I.getOperand(2).getReg();
6397 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
6398 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
6402 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6404 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6406 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6409 return diagnoseUnsupported(
6411 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
6415bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
6416 SPIRVTypeInst ResType,
6417 MachineInstr &
I)
const {
6419 Register OpReg =
I.getOperand(2).getReg();
6424 unsigned ExtendOpcode = SPIRV::OpUConvert;
6425 unsigned BitSetOpcode = GL::FindILsb;
6429 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6431 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6433 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6436 return diagnoseUnsupported(
I,
6437 "spv_firstbitlow only supports 16,32,64 bits.");
6441bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
6442 SPIRVTypeInst ResType,
6443 MachineInstr &
I)
const {
6447 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
6450 .
addUse(
I.getOperand(2).getReg())
6453 unsigned Alignment =
I.getOperand(3).getImm();
6467 while (!Worklist.
empty()) {
6469 switch (
T->getOpcode()) {
6470 case SPIRV::OpTypeInt:
6471 case SPIRV::OpTypeFloat:
6472 case SPIRV::OpTypePointer:
6474 case SPIRV::OpTypeVector:
6475 case SPIRV::OpTypeMatrix:
6476 case SPIRV::OpTypeArray: {
6477 Register OperandReg =
T->getOperand(1).getReg();
6481 case SPIRV::OpTypeStruct:
6482 for (
unsigned Idx = 1,
E =
T->getNumOperands(); Idx <
E; ++Idx) {
6483 Register OperandReg =
T->getOperand(Idx).getReg();
6495bool SPIRVInstructionSelector::selectAbort(MachineInstr &
I)
const {
6496 assert(
I.getNumExplicitOperands() == 2);
6498 Register MsgReg =
I.getOperand(1).getReg();
6500 assert(MsgType &&
"Message argument of llvm.spv.abort has no SPIR-V type");
6503 return diagnoseUnsupported(
6505 "llvm.spv.abort message type must be a concrete SPIR-V type (numerical "
6506 "scalar, pointer, vector, matrix, or aggregate of such types)");
6509 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
6516bool SPIRVInstructionSelector::selectTrap(MachineInstr &
I)
const {
6525 uint32_t MsgVal = ~0
u;
6526 if (
I.getOpcode() == TargetOpcode::G_UBSANTRAP)
6527 MsgVal =
static_cast<uint32_t
>(
I.getOperand(0).
getImm());
6530 Register MsgReg = buildI32ConstantInEntryBlock(MsgVal,
I, MsgType);
6533 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
6540bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
6541 SPIRVTypeInst ResType,
6542 MachineInstr &
I)
const {
6546 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
6549 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
6552 unsigned Alignment =
I.getOperand(2).getImm();
6559bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
6564 const MachineInstr *PrevI =
I.getPrevNode();
6566 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
6570 .
addMBB(
I.getOperand(0).getMBB())
6575 .
addMBB(
I.getOperand(0).getMBB())
6580bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
6591 const MachineInstr *NextI =
I.getNextNode();
6593 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
6599 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
6601 .
addUse(
I.getOperand(0).getReg())
6602 .
addMBB(
I.getOperand(1).getMBB())
6608bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
6609 MachineInstr &
I)
const {
6611 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
6613 const unsigned NumOps =
I.getNumOperands();
6614 for (
unsigned i = 1; i <
NumOps; i += 2) {
6615 MIB.
addUse(
I.getOperand(i + 0).getReg());
6616 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
6622bool SPIRVInstructionSelector::selectGlobalValue(
6623 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
6625 MachineIRBuilder MIRBuilder(
I);
6626 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
6629 std::string GlobalIdent;
6631 unsigned &
ID = UnnamedGlobalIDs[GV];
6633 ID = UnnamedGlobalIDs.
size();
6634 GlobalIdent =
"__unnamed_" + Twine(
ID).str();
6660 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
6667 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
6672 MachineInstrBuilder MIB1 =
6673 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
6676 MachineInstrBuilder MIB2 =
6678 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
6682 GR.
add(ConstVal, MIB2);
6690 MachineInstrBuilder MIB3 =
6691 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
6694 GR.
add(ConstVal, MIB3);
6698 assert(NewReg != ResVReg);
6699 return BuildCOPY(ResVReg, NewReg,
I);
6709 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
6712 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
6718 SPIRVTypeInst ResType =
6722 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
6727 if (
GlobalVar->isExternallyInitialized() &&
6728 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
6729 constexpr unsigned ReadWriteINTEL = 3u;
6732 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
6738bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
6739 SPIRVTypeInst ResType,
6740 MachineInstr &
I)
const {
6742 return selectExtInst(ResVReg, ResType,
I, CL::log10);
6750 MachineIRBuilder MIRBuilder(
I);
6755 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6758 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6760 .
add(
I.getOperand(1))
6765 ResType->
getOpcode() == SPIRV::OpTypeFloat);
6775 APFloat::rmNearestTiesToEven, &LosesInfo);
6779 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
6780 ? SPIRV::OpVectorTimesScalar
6791bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
6792 SPIRVTypeInst ResType,
6793 MachineInstr &
I)
const {
6796 return selectExtInst(ResVReg, ResType,
I, CL::pown);
6802 Register ExpReg =
I.getOperand(2).getReg();
6804 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
6805 SPIRV::OpConvertSToF))
6807 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
6814bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
6815 SPIRVTypeInst ResType,
6816 MachineInstr &
I)
const {
6832 MachineIRBuilder MIRBuilder(
I);
6835 ResType, MIRBuilder, SPIRV::StorageClass::Function);
6848 MachineBasicBlock &EntryBB =
I.getMF()->
front();
6850 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
6853 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
6859 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6862 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
6865 .
add(
I.getOperand(
I.getNumExplicitDefs()))
6869 Register IntegralPartReg =
I.getOperand(1).getReg();
6870 if (IntegralPartReg.
isValid()) {
6872 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6883 assert(
false &&
"GLSL::Modf is deprecated.");
6894bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
6895 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
6896 SPIRVTypeInst ResType, MachineInstr &
I)
const {
6897 MachineIRBuilder MIRBuilder(
I);
6898 const SPIRVTypeInst Vec3Ty =
6901 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
6913 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
6917 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
6923 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6930 assert(
I.getOperand(2).isReg());
6931 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
6935 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6946bool SPIRVInstructionSelector::loadBuiltinInputID(
6947 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
6948 SPIRVTypeInst ResType, MachineInstr &
I)
const {
6949 MachineIRBuilder MIRBuilder(
I);
6951 ResType, MIRBuilder, SPIRV::StorageClass::Input);
6966 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
6970 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
6979SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
6980 MachineInstr &
I)
const {
6981 MachineIRBuilder MIRBuilder(
I);
6982 if (
Type->getOpcode() != SPIRV::OpTypeVector)
6992bool SPIRVInstructionSelector::loadHandleBeforePosition(
6993 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
6994 MachineInstr &Pos)
const {
6997 Intrinsic::spv_resource_handlefrombinding);
7005 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
7006 MachineIRBuilder MIRBuilder(HandleDef);
7007 SPIRVTypeInst VarType = ResType;
7008 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
7010 if (IsStructuredBuffer) {
7015 if (ResType->
getOpcode() == SPIRV::OpTypeImage && ArraySize == 0)
7017 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
7020 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
7021 ArraySize, IndexReg, Name, MIRBuilder);
7025 uint32_t LoadOpcode =
7026 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
7036bool SPIRVInstructionSelector::errorIfInstrOutsideShader(
7037 MachineInstr &
I)
const {
7039 return diagnoseUnsupported(
7040 I,
"this instruction is only supported in shaders.");
7045InstructionSelector *
7049 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...