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;
195 unsigned OpType)
const;
262 unsigned Opcode)
const;
266 unsigned Opcode)
const;
270 unsigned Opcode)
const;
274 unsigned Opcode)
const;
276 template <
bool Signed>
279 template <
bool Signed>
286 template <
typename PickOpcodeFn>
289 PickOpcodeFn &&PickOpcode)
const;
306 template <
typename PickOpcodeFn>
309 PickOpcodeFn &&PickOpcode)
const;
327 bool IsSigned)
const;
329 bool IsSigned,
unsigned Opcode)
const;
331 bool IsSigned)
const;
337 bool IsSigned)
const;
378 GL::GLSLExtInst GLInst,
bool setMIFlags =
true,
379 bool useMISrc =
true,
381 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
382 CL::OpenCLExtInst CLInst,
bool setMIFlags =
true,
383 bool useMISrc =
true,
385 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
386 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
387 bool setMIFlags =
true,
bool useMISrc =
true,
389 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
390 const ExtInstList &ExtInsts,
bool setMIFlags =
true,
391 bool useMISrc =
true,
394 bool selectLog10(
Register ResVReg, SPIRVTypeInst ResType,
395 MachineInstr &
I)
const;
397 bool selectFpowi(
Register ResVReg, SPIRVTypeInst ResType,
398 MachineInstr &
I)
const;
400 bool selectSaturate(
Register ResVReg, SPIRVTypeInst ResType,
401 MachineInstr &
I)
const;
403 bool selectWaveOpInst(
Register ResVReg, SPIRVTypeInst ResType,
404 MachineInstr &
I,
unsigned Opcode)
const;
406 bool selectBarrierInst(MachineInstr &
I,
unsigned Scope,
unsigned MemSem,
407 bool WithGroupSync)
const;
409 bool selectWaveActiveCountBits(
Register ResVReg, SPIRVTypeInst ResType,
410 MachineInstr &
I)
const;
412 bool selectWaveActiveAllEqual(
Register ResVReg, SPIRVTypeInst ResType,
413 MachineInstr &
I)
const;
417 bool selectHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
418 MachineInstr &
I)
const;
420 bool selectCounterHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
421 MachineInstr &
I)
const;
423 bool selectReadImageIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
424 MachineInstr &
I)
const;
425 bool selectGetDimensionsIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
426 MachineInstr &
I)
const;
427 bool selectGetDimensionsLevelsIntrinsic(
Register &ResVReg,
428 SPIRVTypeInst ResType,
429 MachineInstr &
I)
const;
430 bool selectGetDimensionsMSIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
431 MachineInstr &
I)
const;
434 std::optional<Register> LodReg = std::nullopt)
const;
435 bool selectSampleBasicIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
436 MachineInstr &
I)
const;
437 bool selectCalculateLodIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
438 MachineInstr &
I)
const;
439 bool selectSampleBiasIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
440 MachineInstr &
I)
const;
441 bool selectSampleGradIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
442 MachineInstr &
I)
const;
443 bool selectSampleLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
444 MachineInstr &
I)
const;
445 bool selectLoadLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
446 MachineInstr &
I)
const;
447 bool selectSampleCmpIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
448 MachineInstr &
I)
const;
449 bool selectSampleCmpLevelZeroIntrinsic(
Register &ResVReg,
450 SPIRVTypeInst ResType,
451 MachineInstr &
I)
const;
452 bool selectGatherIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
453 MachineInstr &
I)
const;
454 bool selectImageWriteIntrinsic(MachineInstr &
I)
const;
455 bool selectResourceGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
456 MachineInstr &
I)
const;
457 bool selectPushConstantGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
458 MachineInstr &
I)
const;
459 bool selectResourceNonUniformIndex(
Register &ResVReg, SPIRVTypeInst ResType,
460 MachineInstr &
I)
const;
461 bool selectModf(
Register ResVReg, SPIRVTypeInst ResType,
462 MachineInstr &
I)
const;
463 bool selectUpdateCounter(
Register &ResVReg, SPIRVTypeInst ResType,
464 MachineInstr &
I)
const;
465 bool selectFrexp(
Register ResVReg, SPIRVTypeInst ResType,
466 MachineInstr &
I)
const;
467 bool selectSincos(
Register ResVReg, SPIRVTypeInst ResType,
468 MachineInstr &
I)
const;
469 bool selectExp10(
Register ResVReg, SPIRVTypeInst ResType,
470 MachineInstr &
I)
const;
471 bool selectDerivativeInst(
Register ResVReg, SPIRVTypeInst ResType,
472 MachineInstr &
I,
const unsigned DPdOpCode)
const;
474 Register buildI32Constant(uint32_t Val, MachineInstr &
I,
475 SPIRVTypeInst ResType =
nullptr)
const;
476 Register buildI32ConstantInEntryBlock(uint32_t Val, MachineInstr &
I,
477 SPIRVTypeInst ResType =
nullptr)
const;
479 Register buildZerosVal(SPIRVTypeInst ResType, MachineInstr &
I)
const;
480 bool isScalarOrVectorIntConstantZero(
Register Reg)
const;
481 Register buildZerosValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
483 MachineInstr &
I)
const;
484 Register buildOnesValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
486 bool wrapIntoSpecConstantOp(MachineInstr &
I,
489 Register getUcharPtrTypeReg(MachineInstr &
I,
490 SPIRV::StorageClass::StorageClass SC)
const;
491 MachineInstrBuilder buildSpecConstantOp(MachineInstr &
I,
Register Dest,
493 uint32_t Opcode)
const;
494 MachineInstrBuilder buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
495 SPIRVTypeInst SrcPtrTy)
const;
496 Register buildPointerToResource(SPIRVTypeInst ResType,
497 SPIRV::StorageClass::StorageClass SC,
498 uint32_t Set, uint32_t
Binding,
499 uint32_t ArraySize,
Register IndexReg,
501 MachineIRBuilder MIRBuilder)
const;
502 SPIRVTypeInst widenTypeToVec4(SPIRVTypeInst
Type, MachineInstr &
I)
const;
503 bool extractSubvector(
Register &ResVReg, SPIRVTypeInst ResType,
504 Register &ReadReg, MachineInstr &InsertionPoint)
const;
505 bool generateImageReadOrFetch(
Register &ResVReg, SPIRVTypeInst ResType,
508 const ImageOperands *ImOps =
nullptr)
const;
509 bool generateSampleImage(
Register ResVReg, SPIRVTypeInst ResType,
511 Register CoordinateReg,
const ImageOperands &ImOps,
514 bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
515 Register ResVReg, SPIRVTypeInst ResType,
516 MachineInstr &
I)
const;
517 bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
518 Register ResVReg, SPIRVTypeInst ResType,
519 MachineInstr &
I)
const;
520 bool loadHandleBeforePosition(
Register &HandleReg, SPIRVTypeInst ResType,
521 GIntrinsic &HandleDef, MachineInstr &Pos)
const;
522 void decorateUsesAsNonUniform(
Register &NonUniformReg)
const;
523 bool errorIfInstrOutsideShader(MachineInstr &
I)
const;
525 std::optional<SplitParts> splitEvenOddLanes(
Register PopCountReg,
526 unsigned ComponentCount,
528 SPIRVTypeInst I32Type)
const;
531 handle64BitOverflow(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
532 Register SrcReg,
unsigned int Opcode,
533 std::function<
bool(
Register, SPIRVTypeInst,
534 MachineInstr &,
Register,
unsigned)>
538bool sampledTypeIsSignedInteger(
const llvm::Type *HandleType) {
540 if (
TET->getTargetExtName() ==
"spirv.Image") {
543 assert(
TET->getTargetExtName() ==
"spirv.SignedImage");
544 return TET->getTypeParameter(0)->isIntegerTy();
548#define GET_GLOBALISEL_IMPL
549#include "SPIRVGenGlobalISel.inc"
550#undef GET_GLOBALISEL_IMPL
556 TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
559#include
"SPIRVGenGlobalISel.inc"
562#include
"SPIRVGenGlobalISel.inc"
574 InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
578void SPIRVInstructionSelector::resetVRegsType(MachineFunction &MF) {
579 if (HasVRegsReset == &MF)
594 for (
const auto &
MBB : MF) {
595 for (
const auto &
MI :
MBB) {
598 if (
MI.getOpcode() != SPIRV::ASSIGN_TYPE)
602 LLT DstType = MRI.
getType(DstReg);
604 LLT SrcType = MRI.
getType(SrcReg);
605 if (DstType != SrcType)
610 if (DstRC != SrcRC && SrcRC)
622 while (!Stack.empty()) {
627 switch (
MI->getOpcode()) {
628 case TargetOpcode::G_INTRINSIC:
629 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
630 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS: {
633 if (IntrID != Intrinsic::spv_const_composite &&
634 IntrID != Intrinsic::spv_undef && IntrID != Intrinsic::spv_poison)
638 case TargetOpcode::G_BUILD_VECTOR:
639 case TargetOpcode::G_SPLAT_VECTOR:
641 i < OpDef->getNumOperands(); i++) {
646 Stack.push_back(OpNestedDef);
649 case TargetOpcode::G_CONSTANT:
650 case TargetOpcode::G_FCONSTANT:
651 case TargetOpcode::G_IMPLICIT_DEF:
652 case SPIRV::OpConstantTrue:
653 case SPIRV::OpConstantFalse:
654 case SPIRV::OpConstantI:
655 case SPIRV::OpConstantF:
656 case SPIRV::OpConstantComposite:
657 case SPIRV::OpConstantCompositeContinuedINTEL:
658 case SPIRV::OpConstantSampler:
659 case SPIRV::OpConstantNull:
661 case SPIRV::OpPoisonKHR:
662 case SPIRV::OpConstantFunctionPointerINTEL:
689 case Intrinsic::spv_all:
690 case Intrinsic::spv_alloca:
691 case Intrinsic::spv_any:
692 case Intrinsic::spv_bitcast:
693 case Intrinsic::spv_const_composite:
694 case Intrinsic::spv_cross:
695 case Intrinsic::spv_degrees:
696 case Intrinsic::spv_distance:
697 case Intrinsic::spv_extractelt:
698 case Intrinsic::spv_extractv:
699 case Intrinsic::spv_faceforward:
700 case Intrinsic::spv_fdot:
701 case Intrinsic::spv_firstbitlow:
702 case Intrinsic::spv_firstbitshigh:
703 case Intrinsic::spv_firstbituhigh:
704 case Intrinsic::spv_frac:
705 case Intrinsic::spv_gep:
706 case Intrinsic::spv_global_offset:
707 case Intrinsic::spv_global_size:
708 case Intrinsic::spv_group_id:
709 case Intrinsic::spv_insertelt:
710 case Intrinsic::spv_insertv:
711 case Intrinsic::spv_isinf:
712 case Intrinsic::spv_isnan:
713 case Intrinsic::spv_isfinite:
714 case Intrinsic::spv_isnormal:
715 case Intrinsic::spv_lerp:
716 case Intrinsic::spv_length:
717 case Intrinsic::spv_normalize:
718 case Intrinsic::spv_num_subgroups:
719 case Intrinsic::spv_num_workgroups:
720 case Intrinsic::spv_ptrcast:
721 case Intrinsic::spv_radians:
722 case Intrinsic::spv_reflect:
723 case Intrinsic::spv_refract:
724 case Intrinsic::spv_resource_getbasepointer:
725 case Intrinsic::spv_resource_getpointer:
726 case Intrinsic::spv_resource_handlefrombinding:
727 case Intrinsic::spv_resource_handlefromimplicitbinding:
728 case Intrinsic::spv_resource_nonuniformindex:
729 case Intrinsic::spv_resource_sample:
730 case Intrinsic::spv_rsqrt:
731 case Intrinsic::spv_saturate:
732 case Intrinsic::spv_sdot:
733 case Intrinsic::spv_sign:
734 case Intrinsic::spv_smoothstep:
735 case Intrinsic::spv_step:
736 case Intrinsic::spv_subgroup_id:
737 case Intrinsic::spv_subgroup_local_invocation_id:
738 case Intrinsic::spv_subgroup_max_size:
739 case Intrinsic::spv_subgroup_size:
740 case Intrinsic::spv_thread_id:
741 case Intrinsic::spv_thread_id_in_group:
742 case Intrinsic::spv_udot:
743 case Intrinsic::spv_undef:
744 case Intrinsic::spv_value_md:
745 case Intrinsic::spv_workgroup_size:
757 case SPIRV::OpTypeVoid:
758 case SPIRV::OpTypeBool:
759 case SPIRV::OpTypeInt:
760 case SPIRV::OpTypeFloat:
761 case SPIRV::OpTypeVector:
762 case SPIRV::OpTypeMatrix:
763 case SPIRV::OpTypeImage:
764 case SPIRV::OpTypeSampler:
765 case SPIRV::OpTypeSampledImage:
766 case SPIRV::OpTypeArray:
767 case SPIRV::OpTypeRuntimeArray:
768 case SPIRV::OpTypeStruct:
769 case SPIRV::OpTypeOpaque:
770 case SPIRV::OpTypePointer:
771 case SPIRV::OpTypeFunction:
772 case SPIRV::OpTypeEvent:
773 case SPIRV::OpTypeDeviceEvent:
774 case SPIRV::OpTypeReserveId:
775 case SPIRV::OpTypeQueue:
776 case SPIRV::OpTypePipe:
777 case SPIRV::OpTypeForwardPointer:
778 case SPIRV::OpTypePipeStorage:
779 case SPIRV::OpTypeNamedBarrier:
780 case SPIRV::OpTypeAccelerationStructureNV:
781 case SPIRV::OpTypeCooperativeMatrixNV:
782 case SPIRV::OpTypeCooperativeMatrixKHR:
792 if (
MI.getNumDefs() == 0)
795 for (
const auto &MO :
MI.all_defs()) {
797 if (
Reg.isPhysical()) {
802 if (
UseMI.getOpcode() != SPIRV::OpName) {
809 if (
MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE ||
MI.isFakeUse() ||
810 MI.isLifetimeMarker()) {
813 <<
"Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
824 if (
MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
825 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
828 LLVM_DEBUG(
dbgs() <<
"Dead: Intrinsic with no real side effects.\n");
833 if (
MI.mayStore() ||
MI.isCall() ||
834 (
MI.mayLoad() &&
MI.hasOrderedMemoryRef()) ||
MI.isPosition() ||
835 MI.isDebugInstr() ||
MI.isTerminator() ||
MI.isJumpTableDebugInfo()) {
836 LLVM_DEBUG(
dbgs() <<
"Not dead: instruction has side effects.\n");
847 LLVM_DEBUG(
dbgs() <<
"Dead: known opcode with no side effects\n");
854void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &
MI)
const {
856 for (
const auto &MO :
MI.all_defs()) {
860 SmallVector<MachineInstr *, 4> UselessOpNames;
863 "There is still a use of the dead function.");
866 for (MachineInstr *OpNameMI : UselessOpNames) {
868 OpNameMI->eraseFromParent();
873void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &
MI)
const {
876 removeOpNamesForDeadMI(
MI);
877 MI.eraseFromParent();
880bool SPIRVInstructionSelector::select(MachineInstr &
I) {
881 resetVRegsType(*
I.getParent()->getParent());
883 assert(
I.getParent() &&
"Instruction should be in a basic block!");
884 assert(
I.getParent()->getParent() &&
"Instruction should be in a function!");
889 removeDeadInstruction(
I);
896 if (Opcode == SPIRV::ASSIGN_TYPE) {
897 Register DstReg =
I.getOperand(0).getReg();
898 Register SrcReg =
I.getOperand(1).getReg();
901 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
902 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
903 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
904 Register SelectDstReg =
Def->getOperand(0).getReg();
905 bool SuccessToSelectSelect [[maybe_unused]] = selectSelect(
907 assert(SuccessToSelectSelect);
909 Def->eraseFromParent();
916 bool Res = selectImpl(
I, *CoverageInfo);
918 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
919 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
923 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
935 }
else if (
I.getNumDefs() == 1) {
947 removeDeadInstruction(
I);
952 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
953 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
959 bool HasDefs =
I.getNumDefs() > 0;
962 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
963 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
964 if (spvSelect(ResVReg, ResType,
I)) {
966 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
977 case TargetOpcode::G_CONSTANT:
978 case TargetOpcode::G_FCONSTANT:
985 MachineInstr &
I)
const {
988 if (DstRC != SrcRC && SrcRC)
990 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
997bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
998 SPIRVTypeInst ResType,
999 MachineInstr &
I)
const {
1000 const unsigned Opcode =
I.getOpcode();
1002 return selectImpl(
I, *CoverageInfo);
1004 case TargetOpcode::G_CONSTANT:
1005 case TargetOpcode::G_FCONSTANT:
1006 return selectConst(ResVReg, ResType,
I);
1007 case TargetOpcode::G_GLOBAL_VALUE:
1008 return selectGlobalValue(ResVReg,
I);
1009 case TargetOpcode::G_IMPLICIT_DEF:
1010 return selectOpUndef(ResVReg, ResType,
I);
1011 case TargetOpcode::G_FREEZE:
1012 return selectFreeze(ResVReg, ResType,
I);
1014 case TargetOpcode::G_INTRINSIC:
1015 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
1016 case TargetOpcode::G_INTRINSIC_CONVERGENT:
1017 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
1018 return selectIntrinsic(ResVReg, ResType,
I);
1019 case TargetOpcode::G_BITREVERSE:
1020 return selectBitreverse(ResVReg, ResType,
I);
1022 case TargetOpcode::G_BUILD_VECTOR:
1023 return selectBuildVector(ResVReg, ResType,
I);
1024 case TargetOpcode::G_SPLAT_VECTOR:
1025 return selectSplatVector(ResVReg, ResType,
I);
1026 case TargetOpcode::G_CONCAT_VECTORS:
1027 return selectConcatVectors(ResVReg, ResType,
I);
1029 case TargetOpcode::G_SHUFFLE_VECTOR: {
1030 MachineBasicBlock &BB = *
I.getParent();
1031 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
1034 .
addUse(
I.getOperand(1).getReg())
1035 .
addUse(
I.getOperand(2).getReg());
1036 for (
auto V :
I.getOperand(3).getShuffleMask())
1041 case TargetOpcode::G_MEMMOVE:
1042 case TargetOpcode::G_MEMCPY:
1043 case TargetOpcode::G_MEMCPY_INLINE:
1044 case TargetOpcode::G_MEMSET:
1045 case TargetOpcode::G_MEMSET_INLINE:
1046 return selectMemOperation(ResVReg,
I);
1048 case TargetOpcode::G_ICMP:
1049 return selectICmp(ResVReg, ResType,
I);
1050 case TargetOpcode::G_FCMP:
1051 return selectFCmp(ResVReg, ResType,
I);
1053 case TargetOpcode::G_FRAME_INDEX:
1054 return selectFrameIndex(ResVReg, ResType,
I);
1056 case TargetOpcode::G_LOAD:
1057 return selectLoad(ResVReg, ResType,
I);
1058 case TargetOpcode::G_STORE:
1059 return selectStore(
I);
1061 case TargetOpcode::G_BR:
1062 return selectBranch(
I);
1063 case TargetOpcode::G_BRCOND:
1064 return selectBranchCond(
I);
1066 case TargetOpcode::G_PHI:
1067 return selectPhi(ResVReg,
I);
1069 case TargetOpcode::G_FPTOSI:
1070 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1071 case TargetOpcode::G_FPTOUI:
1072 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1074 case TargetOpcode::G_FPTOSI_SAT:
1075 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1076 case TargetOpcode::G_FPTOUI_SAT:
1077 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1079 case TargetOpcode::G_SITOFP:
1080 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
1081 case TargetOpcode::G_UITOFP:
1082 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
1084 case TargetOpcode::G_CTPOP:
1085 return selectPopCount(ResVReg, ResType,
I, SPIRV::OpBitCount);
1086 case TargetOpcode::G_SMIN:
1087 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
1088 case TargetOpcode::G_UMIN:
1089 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
1091 case TargetOpcode::G_SMAX:
1092 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
1093 case TargetOpcode::G_UMAX:
1094 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
1096 case TargetOpcode::G_SCMP:
1097 return selectSUCmp(ResVReg, ResType,
I,
true);
1098 case TargetOpcode::G_UCMP:
1099 return selectSUCmp(ResVReg, ResType,
I,
false);
1100 case TargetOpcode::G_LROUND:
1101 case TargetOpcode::G_LLROUND: {
1104 MRI->
setRegClass(regForLround, &SPIRV::iIDRegClass);
1106 regForLround, *(
I.getParent()->getParent()));
1108 CL::round, GL::Round,
false);
1110 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1117 case TargetOpcode::G_STRICT_FMA:
1118 case TargetOpcode::G_FMA: {
1121 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1124 .
addUse(
I.getOperand(1).getReg())
1125 .
addUse(
I.getOperand(2).getReg())
1126 .
addUse(
I.getOperand(3).getReg())
1131 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1134 case TargetOpcode::G_STRICT_FLDEXP:
1135 return selectExtInst(ResVReg, ResType,
I, CL::ldexp);
1137 case TargetOpcode::G_FPOW:
1138 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1139 case TargetOpcode::G_FPOWI:
1140 return selectFpowi(ResVReg, ResType,
I);
1142 case TargetOpcode::G_FEXP:
1143 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1144 case TargetOpcode::G_FEXP2:
1145 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1146 case TargetOpcode::G_FEXP10:
1147 return selectExp10(ResVReg, ResType,
I);
1149 case TargetOpcode::G_FMODF:
1150 return selectModf(ResVReg, ResType,
I);
1151 case TargetOpcode::G_FSINCOS:
1152 return selectSincos(ResVReg, ResType,
I);
1154 case TargetOpcode::G_FLOG:
1155 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1156 case TargetOpcode::G_FLOG2:
1157 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1158 case TargetOpcode::G_FLOG10:
1159 return selectLog10(ResVReg, ResType,
I);
1161 case TargetOpcode::G_FABS:
1162 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1163 case TargetOpcode::G_ABS:
1164 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1166 case TargetOpcode::G_FMINNUM:
1167 case TargetOpcode::G_FMINIMUM:
1168 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1169 case TargetOpcode::G_FMAXNUM:
1170 case TargetOpcode::G_FMAXIMUM:
1171 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1173 case TargetOpcode::G_FCOPYSIGN:
1174 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1176 case TargetOpcode::G_FCEIL:
1177 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1178 case TargetOpcode::G_FFLOOR:
1179 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1181 case TargetOpcode::G_FCOS:
1182 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1183 case TargetOpcode::G_FSIN:
1184 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1185 case TargetOpcode::G_FTAN:
1186 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1187 case TargetOpcode::G_FACOS:
1188 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1189 case TargetOpcode::G_FASIN:
1190 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1191 case TargetOpcode::G_FATAN:
1192 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1193 case TargetOpcode::G_FATAN2:
1194 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1195 case TargetOpcode::G_FCOSH:
1196 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1197 case TargetOpcode::G_FSINH:
1198 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1199 case TargetOpcode::G_FTANH:
1200 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1202 case TargetOpcode::G_STRICT_FSQRT:
1203 case TargetOpcode::G_FSQRT:
1204 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1206 case TargetOpcode::G_CTTZ:
1207 case TargetOpcode::G_CTTZ_ZERO_POISON:
1208 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1209 case TargetOpcode::G_CTLZ:
1210 case TargetOpcode::G_CTLZ_ZERO_POISON:
1211 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1213 case TargetOpcode::G_INTRINSIC_ROUND:
1214 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1215 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1216 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1217 case TargetOpcode::G_INTRINSIC_TRUNC:
1218 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1219 case TargetOpcode::G_FRINT:
1220 case TargetOpcode::G_FNEARBYINT:
1221 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1223 case TargetOpcode::G_SMULH:
1224 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1225 case TargetOpcode::G_UMULH:
1226 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1228 case TargetOpcode::G_SADDSAT:
1229 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1230 case TargetOpcode::G_UADDSAT:
1231 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1232 case TargetOpcode::G_SSUBSAT:
1233 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1234 case TargetOpcode::G_USUBSAT:
1235 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1237 case TargetOpcode::G_FFREXP:
1238 return selectFrexp(ResVReg, ResType,
I);
1240 case TargetOpcode::G_UADDO:
1241 return selectOverflowArith(ResVReg, ResType,
I,
1242 ResType->
getOpcode() == SPIRV::OpTypeVector
1243 ? SPIRV::OpIAddCarryV
1244 : SPIRV::OpIAddCarryS);
1245 case TargetOpcode::G_USUBO:
1246 return selectOverflowArith(ResVReg, ResType,
I,
1247 ResType->
getOpcode() == SPIRV::OpTypeVector
1248 ? SPIRV::OpISubBorrowV
1249 : SPIRV::OpISubBorrowS);
1250 case TargetOpcode::G_UMULO:
1251 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1252 case TargetOpcode::G_SMULO:
1253 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1255 case TargetOpcode::G_SEXT:
1256 return selectExt(ResVReg, ResType,
I,
true);
1257 case TargetOpcode::G_ANYEXT:
1258 case TargetOpcode::G_ZEXT:
1259 return selectExt(ResVReg, ResType,
I,
false);
1260 case TargetOpcode::G_TRUNC:
1261 return selectTrunc(ResVReg, ResType,
I);
1262 case TargetOpcode::G_FPTRUNC:
1263 case TargetOpcode::G_FPEXT:
1264 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1266 case TargetOpcode::G_PTRTOINT:
1267 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1268 case TargetOpcode::G_INTTOPTR:
1269 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1270 case TargetOpcode::G_BITCAST:
1271 return selectBitcast(ResVReg, ResType,
I);
1272 case TargetOpcode::G_ADDRSPACE_CAST:
1273 return selectAddrSpaceCast(ResVReg, ResType,
I);
1274 case TargetOpcode::G_PTRMASK:
1275 return selectPtrMask(ResVReg, ResType,
I);
1276 case TargetOpcode::G_PTR_ADD: {
1278 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1282 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1283 (*II).getOpcode() == TargetOpcode::COPY ||
1284 (*II).getOpcode() == SPIRV::OpVariable) &&
1285 getImm(
I.getOperand(2), MRI));
1287 bool IsGVInit =
false;
1291 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1292 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1293 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1294 (*UseIt).getOpcode() == SPIRV::OpVariable) {
1304 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1316 return diagnoseUnsupported(
1317 I,
"incompatible result and operand types in a bitcast");
1319 MachineInstrBuilder MIB =
1320 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1327 : SPIRV::OpInBoundsPtrAccessChain))
1331 .
addUse(
I.getOperand(2).getReg())
1334 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1338 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1340 .
addUse(
I.getOperand(2).getReg())
1349 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1352 .
addImm(
static_cast<uint32_t
>(
1353 SPIRV::Opcode::InBoundsPtrAccessChain))
1356 .
addUse(
I.getOperand(2).getReg());
1361 case TargetOpcode::G_ATOMICRMW_OR:
1362 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1363 case TargetOpcode::G_ATOMICRMW_ADD:
1364 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1365 case TargetOpcode::G_ATOMICRMW_AND:
1366 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1367 case TargetOpcode::G_ATOMICRMW_MAX:
1368 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1369 case TargetOpcode::G_ATOMICRMW_MIN:
1370 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1371 case TargetOpcode::G_ATOMICRMW_SUB:
1372 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1373 case TargetOpcode::G_ATOMICRMW_XOR:
1374 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1375 case TargetOpcode::G_ATOMICRMW_UMAX:
1376 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1377 case TargetOpcode::G_ATOMICRMW_UMIN:
1378 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1379 case TargetOpcode::G_ATOMICRMW_XCHG:
1380 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1382 case TargetOpcode::G_ATOMICRMW_FADD:
1383 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1384 case TargetOpcode::G_ATOMICRMW_FSUB:
1386 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1387 ResType->
getOpcode() == SPIRV::OpTypeVector
1389 : SPIRV::OpFNegate);
1390 case TargetOpcode::G_ATOMICRMW_FMIN:
1391 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1392 case TargetOpcode::G_ATOMICRMW_FMAX:
1393 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1395 case TargetOpcode::G_FENCE:
1396 return selectFence(
I);
1398 case TargetOpcode::G_STACKSAVE:
1399 return selectStackSave(ResVReg, ResType,
I);
1400 case TargetOpcode::G_STACKRESTORE:
1401 return selectStackRestore(
I);
1403 case TargetOpcode::G_UNMERGE_VALUES:
1406 case TargetOpcode::G_TRAP:
1407 case TargetOpcode::G_UBSANTRAP:
1408 return selectTrap(
I);
1413 case TargetOpcode::DBG_LABEL:
1415 case TargetOpcode::G_DEBUGTRAP:
1416 return selectDebugTrap(ResVReg, ResType,
I);
1423bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1424 SPIRVTypeInst ResType,
1425 MachineInstr &
I)
const {
1426 unsigned Opcode = SPIRV::OpNop;
1433bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1434 SPIRVTypeInst ResType,
1436 GL::GLSLExtInst GLInst,
1437 bool setMIFlags,
bool useMISrc,
1440 SPIRV::InstructionSet::InstructionSet::GLSL_std_450))
1441 return diagnoseUnsupported(
1443 "this instruction is only supported with the GLSL extended instruction "
1445 return selectExtInst(ResVReg, ResType,
I,
1446 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1447 setMIFlags, useMISrc, SrcRegs);
1450bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1451 SPIRVTypeInst ResType,
1453 CL::OpenCLExtInst CLInst,
1454 bool setMIFlags,
bool useMISrc,
1456 return selectExtInst(ResVReg, ResType,
I,
1457 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1458 setMIFlags, useMISrc, SrcRegs);
1461bool SPIRVInstructionSelector::selectExtInst(
1462 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1463 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1465 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1466 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1467 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1471bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1472 SPIRVTypeInst ResType,
1475 bool setMIFlags,
bool useMISrc,
1478 for (
const auto &[InstructionSet, Opcode] : Insts) {
1482 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1485 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1490 const unsigned NumOps =
I.getNumOperands();
1493 I.getOperand(Index).getType() ==
1494 MachineOperand::MachineOperandType::MO_IntrinsicID)
1497 MIB.
add(
I.getOperand(Index));
1509bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1510 SPIRVTypeInst ResType,
1511 MachineInstr &
I)
const {
1512 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1513 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1514 for (
const auto &Ex : ExtInsts) {
1515 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1516 uint32_t Opcode = Ex.second;
1520 MachineIRBuilder MIRBuilder(
I);
1523 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1528 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1531 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1535 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1538 .
addImm(
static_cast<uint32_t
>(Ex.first))
1540 .
add(
I.getOperand(2))
1544 Register ExpResReg =
I.getOperand(1).getReg();
1546 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1556bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1557 SPIRVTypeInst ResType,
1558 MachineInstr &
I)
const {
1559 Register CosResVReg =
I.getOperand(1).getReg();
1560 unsigned SrcIdx =
I.getNumExplicitDefs();
1565 MachineIRBuilder MIRBuilder(
I);
1567 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1572 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
1575 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
1577 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1580 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1582 .
add(
I.getOperand(SrcIdx))
1585 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1593 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1596 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1598 .
add(
I.getOperand(SrcIdx))
1600 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1603 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1605 .
add(
I.getOperand(SrcIdx))
1612bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1613 SPIRVTypeInst ResType,
1616 unsigned Opcode)
const {
1617 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1627std::optional<SplitParts> SPIRVInstructionSelector::splitEvenOddLanes(
1628 Register PopCountReg,
unsigned ComponentCount, MachineInstr &
I,
1629 SPIRVTypeInst I32Type)
const {
1632 if (ComponentCount == 1) {
1635 Parts.IsScalar =
true;
1636 Parts.Type = I32Type;
1644 if (!selectOpWithSrcs(Parts.High, I32Type,
I, {PopCountReg, IdxOne},
1645 SPIRV::OpVectorExtractDynamic))
1646 return std::nullopt;
1648 if (!selectOpWithSrcs(Parts.Low, I32Type,
I, {PopCountReg, IdxZero},
1649 SPIRV::OpVectorExtractDynamic))
1650 return std::nullopt;
1654 MachineIRBuilder MIRBuilder(
I);
1655 Parts.IsScalar =
false;
1662 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1663 TII.get(SPIRV::OpVectorShuffle))
1668 for (
unsigned J = 1; J < ComponentCount * 2; J += 2)
1673 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
1674 TII.get(SPIRV::OpVectorShuffle))
1679 for (
unsigned J = 0; J < ComponentCount * 2; J += 2)
1687bool SPIRVInstructionSelector::selectPopCount16(
Register ResVReg,
1688 SPIRVTypeInst ResType,
1691 unsigned Opcode)
const {
1692 Register OpReg =
I.getOperand(1).getReg();
1695 MachineIRBuilder MIRBuilder(
I);
1697 SPIRVTypeInst I32VectorType =
1700 bool IsVector = NumElems > 1;
1701 SPIRVTypeInst ExtType = IsVector ? I32VectorType : I32Type;
1704 if (!selectOpWithSrcs(ExtReg, ExtType,
I, {OpReg}, SPIRV::OpUConvert))
1708 if (!selectPopCount32(PopCountReg, ExtType,
I, ExtReg, Opcode))
1711 return selectOpWithSrcs(ResVReg, ResType,
I, {PopCountReg}, ExtOpcode);
1714bool SPIRVInstructionSelector::selectPopCount32(
Register ResVReg,
1715 SPIRVTypeInst ResType,
1718 unsigned Opcode)
const {
1719 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
1722bool SPIRVInstructionSelector::selectPopCount64(
Register ResVReg,
1723 SPIRVTypeInst ResType,
1726 unsigned Opcode)
const {
1728 if (ComponentCount > 2)
1729 return handle64BitOverflow(
1730 ResVReg, ResType,
I, SrcReg, Opcode,
1732 unsigned O) {
return this->selectPopCount64(R,
T,
I, S, O); });
1734 MachineIRBuilder MIRBuilder(
I);
1739 I32Type, 2 * ComponentCount, MIRBuilder,
false);
1743 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
1748 if (!selectPopCount32(Pop32, VecI32Type,
I, Vec32, Opcode))
1752 auto MaybeParts = splitEvenOddLanes(Pop32, ComponentCount,
I, I32Type);
1755 SplitParts &Parts = *MaybeParts;
1758 unsigned OpAdd = Parts.IsScalar ? SPIRV::OpIAddS : SPIRV::OpIAddV;
1760 if (!selectOpWithSrcs(Sum, Parts.Type,
I, {Parts.High, Parts.Low}, OpAdd))
1765 unsigned ConvOp = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
1766 return selectOpWithSrcs(ResVReg, ResType,
I, {Sum}, ConvOp);
1769bool SPIRVInstructionSelector::selectPopCount(
Register ResVReg,
1770 SPIRVTypeInst ResType,
1772 unsigned Opcode)
const {
1777 if (!STI.getTargetTriple().isVulkanOS())
1778 return selectUnOp(ResVReg, ResType,
I, Opcode);
1780 Register OpReg =
I.getOperand(1).getReg();
1783 : SPIRV::OpUConvert;
1787 return selectPopCount16(ResVReg, ResType,
I, ExtOpcode, Opcode);
1789 return selectPopCount32(ResVReg, ResType,
I, OpReg, Opcode);
1791 return selectPopCount64(ResVReg, ResType,
I, OpReg, Opcode);
1793 return diagnoseUnsupported(
I,
"unsupported operand bit width for popcount");
1797bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1798 SPIRVTypeInst ResType,
1800 unsigned Opcode)
const {
1802 Register SrcReg =
I.getOperand(1).getReg();
1807 unsigned DefOpCode = DefIt->getOpcode();
1808 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1811 if (
auto *VRD =
getVRegDef(*MRI, DefIt->getOperand(1).getReg()))
1812 DefOpCode = VRD->getOpcode();
1814 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1815 DefOpCode == TargetOpcode::G_CONSTANT ||
1816 DefOpCode == SPIRV::OpVariable || DefOpCode == SPIRV::OpConstantI) {
1822 uint32_t SpecOpcode = 0;
1824 case SPIRV::OpConvertPtrToU:
1825 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1827 case SPIRV::OpConvertUToPtr:
1828 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1833 TII.get(SPIRV::OpSpecConstantOp))
1843 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1847bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1848 SPIRVTypeInst ResType,
1849 MachineInstr &
I)
const {
1850 Register OpReg =
I.getOperand(1).getReg();
1851 SPIRVTypeInst OpType =
1854 return diagnoseUnsupported(
1855 I,
"incompatible result and operand types in a bitcast");
1856 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1866 if (
MemOp->isVolatile())
1867 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1868 if (
MemOp->isNonTemporal())
1869 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1871 if (!ST->isShader() &&
MemOp->getAlign().value())
1872 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1876 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1877 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1881 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1883 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1887 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1891 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1893 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1905 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1907 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1909 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
1913bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
1914 SPIRVTypeInst ResType,
1915 MachineInstr &
I)
const {
1917 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1922 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
1923 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
1925 Register HandleReg = IntPtrDef->getOperand(2).getReg();
1927 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
1931 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
1935 Register IdxReg = IntPtrDef->getOperand(3).getReg();
1936 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
1937 I.getDebugLoc(),
I);
1941 MachineIRBuilder MIRBuilder(
I);
1943 if (
I.getNumMemOperands()) {
1944 const MachineMemOperand *MemOp = *
I.memoperands_begin();
1945 if (MemOp->isAtomic())
1946 return selectAtomicLoad(ResVReg, ResType,
I);
1949 auto MIB = MIRBuilder.buildInstr(SPIRV::OpLoad)
1953 if (!
I.getNumMemOperands()) {
1954 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
1956 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
1965bool SPIRVInstructionSelector::selectAtomicLoad(
Register ResVReg,
1966 SPIRVTypeInst ResType,
1967 MachineInstr &
I)
const {
1968 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
1971 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
1974 return diagnoseUnsupported(
1975 I,
"Lowering to SPIR-V of atomic load is only "
1976 "allowed for integer, floating point or pointer types");
1978 assert(
I.getNumMemOperands());
1979 const MachineMemOperand &MemOp = **
I.memoperands_begin();
1980 assert(MemOp.isAtomic());
1984 Register ScopeReg = buildI32Constant(Scope,
I);
1990 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
1991 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
1994 MachineIRBuilder MIRBuilder(
I);
1998 return diagnoseUnsupported(
1999 I,
"Lowering to SPIR-V of atomic load is only "
2000 "allowed for pointer types for physical addressing model");
2007 SPIRVTypeInst PtrAsIntSpirvType =
2018 PtrAsIntSpirvType, MIRBuilder,
2021 MIRBuilder.getMF());
2023 MIRBuilder.buildInstr(SPIRV::OpBitcast)
2024 .addDef(PtrCastedToMatchValReg)
2027 .constrainAllUses(
TII,
TRI, RBI);
2029 MIRBuilder.buildInstr(SPIRV::OpAtomicLoad)
2032 .addUse(PtrCastedToMatchValReg)
2035 .constrainAllUses(
TII,
TRI, RBI);
2036 MIRBuilder.buildInstr(SPIRV::OpConvertUToPtr)
2040 .constrainAllUses(
TII,
TRI, RBI);
2043 auto AtomicLoad = MIRBuilder.buildInstr(SPIRV::OpAtomicLoad)
2049 AtomicLoad.constrainAllUses(
TII,
TRI, RBI);
2054bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
2056 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2057 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2062 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2063 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2065 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2070 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2074 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2075 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2076 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2077 TII.get(SPIRV::OpImageWrite))
2083 if (sampledTypeIsSignedInteger(LLVMHandleType))
2086 BMI.constrainAllUses(
TII,
TRI, RBI);
2091 if (
I.getNumMemOperands()) {
2092 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2093 if (MemOp->isAtomic())
2094 return selectAtomicStore(
I);
2097 MachineIRBuilder MIRBuilder(
I);
2098 auto MIB = MIRBuilder.buildInstr(SPIRV::OpStore).
addUse(Ptr).
addUse(StoreVal);
2099 if (!
I.getNumMemOperands()) {
2100 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2102 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2111bool SPIRVInstructionSelector::selectAtomicStore(MachineInstr &
I)
const {
2112 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2115 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2116 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2121 assert(
I.getNumMemOperands());
2122 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2123 assert(MemOp.isAtomic());
2127 Register ScopeReg = buildI32Constant(Scope,
I);
2133 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2134 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2136 MachineIRBuilder MIRBuilder(
I);
2140 return diagnoseUnsupported(
2141 I,
"Lowering to SPIR-V of atomic store is only "
2142 "allowed for pointer types for physical addressing model");
2148 SPIRVTypeInst PtrAsIntSpirvType =
2155 MIRBuilder.buildInstr(SPIRV::OpConvertPtrToU)
2159 .constrainAllUses(
TII,
TRI, RBI);
2165 PtrAsIntSpirvType, MIRBuilder,
2168 MIRBuilder.getMF());
2170 MIRBuilder.buildInstr(SPIRV::OpBitcast)
2171 .addDef(PtrCastedToMatchValReg)
2174 .constrainAllUses(
TII,
TRI, RBI);
2176 StoreVal = PtrToUVal;
2177 Ptr = PtrCastedToMatchValReg;
2178 PointeeType = PtrAsIntSpirvType;
2182 return diagnoseUnsupported(
I,
2183 "Lowering to SPIR-V of atomic store is only "
2184 "allowed for integer or floating point types");
2186 auto AtomicStore = MIRBuilder.buildInstr(SPIRV::OpAtomicStore)
2191 AtomicStore.constrainAllUses(
TII,
TRI, RBI);
2196bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
2197 SPIRVTypeInst ResType,
2198 MachineInstr &
I)
const {
2199 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
2207 const Register PtrsReg =
I.getOperand(2).getReg();
2208 const uint32_t Alignment =
I.getOperand(3).getImm();
2209 const Register MaskReg =
I.getOperand(4).getReg();
2210 const Register PassthruReg =
I.getOperand(5).getReg();
2211 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2215 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
2226bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
2227 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
2234 const Register ValuesReg =
I.getOperand(1).getReg();
2235 const Register PtrsReg =
I.getOperand(2).getReg();
2236 const uint32_t Alignment =
I.getOperand(3).getImm();
2237 const Register MaskReg =
I.getOperand(4).getReg();
2238 const Register AlignmentReg = buildI32Constant(Alignment,
I);
2242 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
2251bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
2252 const Twine &Msg)
const {
2253 const Function &
F =
I.getMF()->getFunction();
2254 F.getContext().diagnose(
2255 DiagnosticInfoUnsupported(
F, Msg,
I.getDebugLoc(),
DS_Error));
2259bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
2260 SPIRVTypeInst ResType,
2261 MachineInstr &
I)
const {
2262 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2263 return diagnoseUnsupported(
2264 I,
"llvm.stacksave intrinsic: this instruction requires the following "
2265 "SPIR-V extension: SPV_INTEL_variable_length_array");
2267 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
2274bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
2275 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2276 return diagnoseUnsupported(
2278 "llvm.stackrestore intrinsic: this instruction requires the following "
2279 "SPIR-V extension: SPV_INTEL_variable_length_array");
2280 if (!
I.getOperand(0).isReg())
2283 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
2284 .
addUse(
I.getOperand(0).getReg())
2290SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
2291 MachineIRBuilder MIRBuilder(
I);
2292 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
2299 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
2303 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2304 Type *ArrTy = ArrayType::get(ValTy, Num);
2306 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
2309 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
2316 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
2319 .
addImm(SPIRV::StorageClass::UniformConstant)
2330bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
2333 Register DstReg =
I.getOperand(0).getReg();
2337 return diagnoseUnsupported(
2338 I,
"OpCopyMemory requires operands to have the same type");
2339 uint64_t CopySize =
getIConstVal(
I.getOperand(2).getReg(), MRI);
2343 return diagnoseUnsupported(
2344 I,
"Unable to determine pointee type size for OpCopyMemory");
2345 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
2346 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
2347 return diagnoseUnsupported(
2348 I,
"OpCopyMemory requires the size to match the pointee type size");
2349 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemory))
2352 if (
I.getNumMemOperands()) {
2353 MachineIRBuilder MIRBuilder(
I);
2360bool SPIRVInstructionSelector::selectCopyMemorySized(MachineInstr &
I,
2363 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCopyMemorySized))
2364 .
addUse(
I.getOperand(0).getReg())
2366 .
addUse(
I.getOperand(2).getReg());
2367 if (
I.getNumMemOperands()) {
2368 MachineIRBuilder MIRBuilder(
I);
2375bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
2376 MachineInstr &
I)
const {
2378 Register SizeReg =
I.getOperand(2).getReg();
2380 SizeDef && SizeDef->
getOpcode() == TargetOpcode::G_CONSTANT &&
2384 Register SrcReg =
I.getOperand(1).getReg();
2385 if (
I.getOpcode() == TargetOpcode::G_MEMSET ||
2386 I.getOpcode() == TargetOpcode::G_MEMSET_INLINE) {
2387 Register VarReg = getOrCreateMemSetGlobal(
I);
2390 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2392 ValTy,
I, SPIRV::StorageClass::UniformConstant);
2394 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
2398 if (!selectCopyMemory(
I, SrcReg))
2401 if (!selectCopyMemorySized(
I, SrcReg))
2404 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
2405 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
2410bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
2411 SPIRVTypeInst ResType,
2414 unsigned NegateOpcode)
const {
2416 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2419 Register ScopeReg = buildI32Constant(Scope,
I);
2421 Register Ptr =
I.getOperand(1).getReg();
2422 uint32_t ScSem =
static_cast<uint32_t
>(
2426 Register MemSemReg = buildI32Constant(MemSem,
I);
2428 Register ValueReg =
I.getOperand(2).getReg();
2429 if (NegateOpcode != 0) {
2432 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
2437 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2448bool SPIRVInstructionSelector::selectInterlockedAdd(
Register ResVReg,
2449 SPIRVTypeInst ResType,
2450 MachineInstr &
I)
const {
2451 Register Ptr =
I.getOperand(2).getReg();
2455 assert((SC == SPIRV::StorageClass::Workgroup ||
2456 SC == SPIRV::StorageClass::StorageBuffer) &&
2457 "InterlockedAdd requires Workgroup or StorageBuffer storage class");
2458 uint32_t
Scope =
static_cast<uint32_t
>(SC == SPIRV::StorageClass::Workgroup
2459 ? SPIRV::Scope::Workgroup
2460 : SPIRV::Scope::Device);
2461 Register ScopeReg = buildI32Constant(Scope,
I);
2464 Register MemSemReg = buildI32Constant(MemSem,
I);
2466 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
2477bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2478 unsigned ArgI =
I.getNumOperands() - 1;
2480 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2481 SPIRVTypeInst SrcType =
2483 if (!SrcType || SrcType->
getOpcode() != SPIRV::OpTypeVector)
2485 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2489 unsigned CurrentIndex = 0;
2490 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2491 Register ResVReg =
I.getOperand(i).getReg();
2494 LLT ResLLT = MRI->
getType(ResVReg);
2500 ResType = ScalarType;
2506 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
2509 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2515 for (
unsigned j = 0;
j < NumElements; ++
j) {
2516 MIB.
addImm(CurrentIndex + j);
2518 CurrentIndex += NumElements;
2522 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2534bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2537 Register MemSemReg = buildI32Constant(MemSem,
I);
2539 uint32_t
Scope =
static_cast<uint32_t
>(
2541 Register ScopeReg = buildI32Constant(Scope,
I);
2543 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2550bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2551 SPIRVTypeInst ResType,
2553 unsigned Opcode)
const {
2554 Type *ResTy =
nullptr;
2557 return diagnoseUnsupported(
2559 "Not enough info to select the arithmetic with overflow instruction");
2561 return diagnoseUnsupported(
I,
2562 "Expect struct type result for the arithmetic "
2563 "with overflow instruction");
2569 MachineIRBuilder MIRBuilder(
I);
2571 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2572 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2578 Register ZeroReg = buildZerosVal(ResType,
I);
2583 if (ResName.
size() > 0)
2588 BuildMI(BB, MIRBuilder.getInsertPt(),
I.getDebugLoc(),
TII.get(Opcode))
2591 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2592 MIB.
addUse(
I.getOperand(i).getReg());
2597 MRI->
setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2598 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2600 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2601 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2608 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2609 .
addDef(
I.getOperand(1).getReg())
2617bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2618 SPIRVTypeInst ResType,
2619 MachineInstr &
I)
const {
2621 "selectAtomicCmpXchg only handles the spv_cmpxchg intrinsic");
2622 Register Ptr =
I.getOperand(2).getReg();
2623 Register ScopeReg =
I.getOperand(5).getReg();
2624 Register MemSemEqReg =
I.getOperand(6).getReg();
2625 Register MemSemNeqReg =
I.getOperand(7).getReg();
2627 Register Val =
I.getOperand(4).getReg();
2631 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2650 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2657 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2669 case SPIRV::StorageClass::DeviceOnlyINTEL:
2670 case SPIRV::StorageClass::HostOnlyINTEL:
2679 bool IsGRef =
false;
2680 bool IsAllowedRefs =
2682 unsigned Opcode = It.getOpcode();
2683 if (Opcode == SPIRV::OpConstantComposite ||
2684 Opcode == SPIRV::OpSpecConstantComposite ||
2685 Opcode == SPIRV::OpVariable ||
2686 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2687 return IsGRef = true;
2688 return Opcode == SPIRV::OpName;
2690 return IsAllowedRefs && IsGRef;
2693Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2694 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2696 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2700SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2702 uint32_t Opcode)
const {
2703 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2704 TII.get(SPIRV::OpSpecConstantOp))
2712SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2713 SPIRVTypeInst SrcPtrTy)
const {
2714 SPIRVTypeInst GenericPtrTy =
2718 SPIRV::StorageClass::Generic),
2720 MachineFunction *MF =
I.getParent()->getParent();
2722 MachineInstrBuilder MIB = buildSpecConstantOp(
2724 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2734bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2735 SPIRVTypeInst ResType,
2736 MachineInstr &
I)
const {
2740 Register SrcPtr =
I.getOperand(1).getReg();
2744 if (SrcPtrTy->
getOpcode() != SPIRV::OpTypePointer ||
2745 ResType->
getOpcode() != SPIRV::OpTypePointer)
2746 return BuildCOPY(ResVReg, SrcPtr,
I);
2756 unsigned SpecOpcode =
2758 ?
static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric)
2761 ? static_cast<uint32_t>(
SPIRV::Opcode::GenericCastToPtr)
2768 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2770 .constrainAllUses(
TII,
TRI, RBI);
2772 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2774 buildSpecConstantOp(
2776 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2777 .constrainAllUses(
TII,
TRI, RBI);
2784 return BuildCOPY(ResVReg, SrcPtr,
I);
2786 if ((SrcSC == SPIRV::StorageClass::Function &&
2787 DstSC == SPIRV::StorageClass::Private) ||
2788 (DstSC == SPIRV::StorageClass::Function &&
2789 SrcSC == SPIRV::StorageClass::Private))
2790 return BuildCOPY(ResVReg, SrcPtr,
I);
2794 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2797 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2800 SPIRVTypeInst GenericPtrTy =
2819 return selectUnOp(ResVReg, ResType,
I,
2820 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
2822 return selectUnOp(ResVReg, ResType,
I,
2823 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
2825 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2827 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2837bool SPIRVInstructionSelector::selectPtrMask(
Register ResVReg,
2838 SPIRVTypeInst ResType,
2839 MachineInstr &
I)
const {
2841 return diagnoseUnsupported(
2842 I,
"G_PTRMASK is not supported with logical SPIR-V");
2847 Register PtrReg =
I.getOperand(1).getReg();
2848 Register MaskReg =
I.getOperand(2).getReg();
2867 ? SPIRV::OpBitwiseAndV
2868 : SPIRV::OpBitwiseAndS;
2891 return SPIRV::OpFOrdEqual;
2893 return SPIRV::OpFOrdGreaterThanEqual;
2895 return SPIRV::OpFOrdGreaterThan;
2897 return SPIRV::OpFOrdLessThanEqual;
2899 return SPIRV::OpFOrdLessThan;
2901 return SPIRV::OpFOrdNotEqual;
2903 return SPIRV::OpOrdered;
2905 return SPIRV::OpFUnordEqual;
2907 return SPIRV::OpFUnordGreaterThanEqual;
2909 return SPIRV::OpFUnordGreaterThan;
2911 return SPIRV::OpFUnordLessThanEqual;
2913 return SPIRV::OpFUnordLessThan;
2915 return SPIRV::OpFUnordNotEqual;
2917 return SPIRV::OpUnordered;
2927 return SPIRV::OpIEqual;
2929 return SPIRV::OpINotEqual;
2931 return SPIRV::OpSGreaterThanEqual;
2933 return SPIRV::OpSGreaterThan;
2935 return SPIRV::OpSLessThanEqual;
2937 return SPIRV::OpSLessThan;
2939 return SPIRV::OpUGreaterThanEqual;
2941 return SPIRV::OpUGreaterThan;
2943 return SPIRV::OpULessThanEqual;
2945 return SPIRV::OpULessThan;
2954 return SPIRV::OpPtrEqual;
2956 return SPIRV::OpPtrNotEqual;
2967 return SPIRV::OpLogicalEqual;
2969 return SPIRV::OpLogicalNotEqual;
3003bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
3004 SPIRVTypeInst ResType,
3006 unsigned OpAnyOrAll)
const {
3007 assert(
I.getNumOperands() == 3);
3008 assert(
I.getOperand(2).isReg());
3010 Register InputRegister =
I.getOperand(2).getReg();
3013 assert(InputType &&
"VReg has no type assigned");
3016 bool IsVectorTy = InputType->
getOpcode() == SPIRV::OpTypeVector;
3017 if (IsBoolTy && !IsVectorTy) {
3018 assert(ResVReg ==
I.getOperand(0).getReg());
3019 return BuildCOPY(ResVReg, InputRegister,
I);
3023 unsigned SpirvNotEqualId =
3024 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
3026 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
3031 IsBoolTy ? InputRegister
3039 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
3041 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
3058bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
3059 SPIRVTypeInst ResType,
3060 MachineInstr &
I)
const {
3061 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
3064bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
3065 SPIRVTypeInst ResType,
3066 MachineInstr &
I)
const {
3067 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
3071bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
3072 SPIRVTypeInst ResType,
3073 MachineInstr &
I)
const {
3074 assert(
I.getNumOperands() == 4);
3075 assert(
I.getOperand(2).isReg());
3076 assert(
I.getOperand(3).isReg());
3078 [[maybe_unused]] SPIRVTypeInst VecType =
3083 "dot product requires a vector of at least 2 components");
3085 [[maybe_unused]] SPIRVTypeInst EltType =
3094 .
addUse(
I.getOperand(2).getReg())
3095 .
addUse(
I.getOperand(3).getReg())
3100bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
3101 SPIRVTypeInst ResType,
3104 assert(
I.getNumOperands() == 4);
3105 assert(
I.getOperand(2).isReg());
3106 assert(
I.getOperand(3).isReg());
3109 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3113 .
addUse(
I.getOperand(2).getReg())
3114 .
addUse(
I.getOperand(3).getReg())
3121bool SPIRVInstructionSelector::selectIntegerDotExpansion(
3122 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3123 assert(
I.getNumOperands() == 4);
3124 assert(
I.getOperand(2).isReg());
3125 assert(
I.getOperand(3).isReg());
3129 Register Vec0 =
I.getOperand(2).getReg();
3130 Register Vec1 =
I.getOperand(3).getReg();
3134 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
3143 "dot product requires a vector of at least 2 components");
3146 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3156 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3167 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3179bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
3180 SPIRVTypeInst ResType,
3181 MachineInstr &
I)
const {
3183 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
3186 .
addUse(
I.getOperand(2).getReg())
3191bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
3192 SPIRVTypeInst ResType,
3193 MachineInstr &
I)
const {
3195 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
3198 .
addUse(
I.getOperand(2).getReg())
3203bool SPIRVInstructionSelector::selectOpIsFinite(
Register ResVReg,
3204 SPIRVTypeInst ResType,
3205 MachineInstr &
I)
const {
3207 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsFinite))
3210 .
addUse(
I.getOperand(2).getReg())
3215bool SPIRVInstructionSelector::selectOpIsNormal(
Register ResVReg,
3216 SPIRVTypeInst ResType,
3217 MachineInstr &
I)
const {
3219 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNormal))
3222 .
addUse(
I.getOperand(2).getReg())
3227template <
bool Signed>
3228bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
3229 SPIRVTypeInst ResType,
3230 MachineInstr &
I)
const {
3231 assert(
I.getNumOperands() == 5);
3232 assert(
I.getOperand(2).isReg());
3233 assert(
I.getOperand(3).isReg());
3234 assert(
I.getOperand(4).isReg());
3237 Register Acc =
I.getOperand(2).getReg();
3241 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3243 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
3248 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
3251 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3263template <
bool Signed>
3264bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
3265 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3266 assert(
I.getNumOperands() == 5);
3267 assert(
I.getOperand(2).isReg());
3268 assert(
I.getOperand(3).isReg());
3269 assert(
I.getOperand(4).isReg());
3272 Register Acc =
I.getOperand(2).getReg();
3278 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
3282 for (
unsigned i = 0; i < 4; i++) {
3305 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
3325 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3340bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
3341 SPIRVTypeInst ResType,
3342 MachineInstr &
I)
const {
3343 assert(
I.getNumOperands() == 3);
3344 assert(
I.getOperand(2).isReg());
3346 Register VZero = buildZerosValF(ResType,
I);
3347 Register VOne = buildOnesValF(ResType,
I);
3349 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
3352 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3354 .
addUse(
I.getOperand(2).getReg())
3361bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
3362 SPIRVTypeInst ResType,
3363 MachineInstr &
I)
const {
3364 assert(
I.getNumOperands() == 3);
3365 assert(
I.getOperand(2).isReg());
3367 Register InputRegister =
I.getOperand(2).getReg();
3369 auto &
DL =
I.getDebugLoc();
3372 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3379 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
3381 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
3389 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3394 if (NeedsConversion) {
3395 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
3406bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
3407 SPIRVTypeInst ResType,
3409 unsigned Opcode)
const {
3413 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
3419 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
3420 BMI.addUse(
I.getOperand(J).getReg());
3427bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
3430 bool WithGroupSync)
const {
3432 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
3434 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
3436 assert(((Scope != SPIRV::Scope::Workgroup) ||
3437 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
3438 "Workgroup Scope must set WorkGroupMemory semantic "
3439 "in Barrier instruction");
3441 assert(((Scope != SPIRV::Scope::Device) ||
3442 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
3443 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
3444 "Device Scope must set UniformMemory and ImageMemory semantic "
3445 "in Barrier instruction");
3451 if (WithGroupSync) {
3452 Register ExecReg = buildI32Constant(SPIRV::Scope::Workgroup,
I);
3456 Register ScopeReg = buildI32Constant(Scope,
I);
3457 Register MemSemReg = buildI32Constant(MemSem,
I);
3459 MI.addUse(ScopeReg).addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
3463bool SPIRVInstructionSelector::selectWaveActiveCountBits(
3464 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3469 if (!selectWaveOpInst(BallotReg, BallotType,
I,
3470 SPIRV::OpGroupNonUniformBallot))
3475 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3480 .
addImm(SPIRV::GroupOperation::Reduce)
3487bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
3488 SPIRVTypeInst ResType,
3489 MachineInstr &
I)
const {
3494 Register InputReg =
I.getOperand(2).getReg();
3499 bool IsVector = NumElems > 1;
3512 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
3513 SPIRV::OpGroupNonUniformAllEqual);
3518 ElementResults.
reserve(NumElems);
3520 for (
unsigned Idx = 0; Idx < NumElems; ++Idx) {
3533 ElemInput = Extracted;
3539 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3550 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3561bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3562 SPIRVTypeInst ResType,
3563 MachineInstr &
I)
const {
3565 assert(
I.getNumOperands() == 3);
3567 auto Op =
I.getOperand(2);
3577 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3579 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3580 return diagnoseUnsupported(
I,
"WavePrefixBitCount requires boolean input");
3601 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3605 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3612bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3613 SPIRVTypeInst ResType,
3615 bool IsUnsigned)
const {
3616 return selectWaveReduce(
3617 ResVReg, ResType,
I, IsUnsigned,
3618 [&](
Register InputRegister,
bool IsUnsigned) {
3619 const bool IsFloatTy =
3621 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3622 : SPIRV::OpGroupNonUniformSMax;
3623 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3627bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3628 SPIRVTypeInst ResType,
3630 bool IsUnsigned)
const {
3631 return selectWaveReduce(
3632 ResVReg, ResType,
I, IsUnsigned,
3633 [&](
Register InputRegister,
bool IsUnsigned) {
3634 const bool IsFloatTy =
3636 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3637 : SPIRV::OpGroupNonUniformSMin;
3638 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3642bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3643 SPIRVTypeInst ResType,
3644 MachineInstr &
I)
const {
3645 return selectWaveReduce(ResVReg, ResType,
I,
false,
3646 [&](
Register InputRegister,
bool IsUnsigned) {
3648 InputRegister, SPIRV::OpTypeFloat);
3649 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3650 : SPIRV::OpGroupNonUniformIAdd;
3654bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3655 SPIRVTypeInst ResType,
3656 MachineInstr &
I)
const {
3657 return selectWaveReduce(ResVReg, ResType,
I,
false,
3658 [&](
Register InputRegister,
bool IsUnsigned) {
3660 InputRegister, SPIRV::OpTypeFloat);
3661 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3662 : SPIRV::OpGroupNonUniformIMul;
3666template <
typename PickOpcodeFn>
3667bool SPIRVInstructionSelector::selectWaveReduce(
3668 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3669 PickOpcodeFn &&PickOpcode)
const {
3670 assert(
I.getNumOperands() == 3);
3671 assert(
I.getOperand(2).isReg());
3673 Register InputRegister =
I.getOperand(2).getReg();
3677 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3680 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3686 .
addImm(SPIRV::GroupOperation::Reduce)
3687 .
addUse(
I.getOperand(2).getReg())
3692bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3693 SPIRVTypeInst ResType,
3695 unsigned Opcode)
const {
3696 return selectWaveReduce(
3697 ResVReg, ResType,
I,
false,
3698 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3701bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3702 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3703 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3704 [&](
Register InputRegister,
bool IsUnsigned) {
3706 InputRegister, SPIRV::OpTypeFloat);
3708 ? SPIRV::OpGroupNonUniformFAdd
3709 : SPIRV::OpGroupNonUniformIAdd;
3713bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3714 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3715 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3716 [&](
Register InputRegister,
bool IsUnsigned) {
3718 InputRegister, SPIRV::OpTypeFloat);
3720 ? SPIRV::OpGroupNonUniformFMul
3721 : SPIRV::OpGroupNonUniformIMul;
3725template <
typename PickOpcodeFn>
3726bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3727 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3728 PickOpcodeFn &&PickOpcode)
const {
3729 assert(
I.getNumOperands() == 3);
3730 assert(
I.getOperand(2).isReg());
3732 Register InputRegister =
I.getOperand(2).getReg();
3736 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3739 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3745 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3746 .
addUse(
I.getOperand(2).getReg())
3751bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
3752 SPIRVTypeInst ResType,
3755 assert(
I.getNumOperands() == 3);
3756 assert(
I.getOperand(2).isReg());
3758 Register InputRegister =
I.getOperand(2).getReg();
3764 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
3775bool SPIRVInstructionSelector::selectBitreverse16(
Register ResVReg,
3776 SPIRVTypeInst ResType,
3781 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
3786 : SPIRV::OpUConvert;
3790 ShiftOp = SPIRV::OpShiftRightLogicalV;
3795 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3796 TII.get(SPIRV::OpConstantComposite))
3799 for (
unsigned It = 0; It <
N; ++It)
3803 ShiftConst = CompositeReg;
3808 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
3813 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
3818 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
3823 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
3826bool SPIRVInstructionSelector::handle64BitOverflow(
3828 unsigned int Opcode,
3835 "handle64BitOverflow should only be used for integer types");
3837 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
3839 MachineIRBuilder MIRBuilder(
I);
3841 SPIRVTypeInst I64x2Type =
3843 SPIRVTypeInst Vec2ResType =
3846 std::vector<Register> PartialRegs;
3848 unsigned CurrentComponent = 0;
3849 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
3853 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3854 TII.get(SPIRV::OpVectorShuffle))
3859 .
addImm(CurrentComponent)
3860 .
addImm(CurrentComponent + 1);
3870 PartialRegs.push_back(SubVecReg);
3873 if (CurrentComponent != ComponentCount) {
3879 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
3880 SPIRV::OpVectorExtractDynamic))
3889 PartialRegs.push_back(FinalElemResReg);
3893 return selectOpWithSrcs(ResVReg, ResType,
I, PartialRegs,
3894 SPIRV::OpCompositeConstruct);
3897bool SPIRVInstructionSelector::selectBitreverse64(
Register ResVReg,
3898 SPIRVTypeInst ResType,
3902 if (ComponentCount > 2)
3903 return handle64BitOverflow(
3904 ResVReg, ResType,
I, SrcReg, SPIRV::OpBitReverse,
3906 unsigned O) {
return this->selectBitreverse64(R,
T,
I, S); });
3908 MachineIRBuilder MIRBuilder(
I);
3912 I32Type, 2 * ComponentCount, MIRBuilder,
false);
3916 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
3921 if (!selectBitreverseNative(Reverse32, VecI32Type,
I, Vec32))
3928 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3929 TII.get(SPIRV::OpVectorShuffle))
3934 for (
unsigned J = 0; J < ComponentCount; ++J) {
3941 return selectOpWithSrcs(ResVReg, ResType,
I, {SwappedVec}, SPIRV::OpBitcast);
3944bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
3945 SPIRVTypeInst ResType,
3949 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
3957bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
3958 SPIRVTypeInst ResType,
3959 MachineInstr &
I)
const {
3960 Register OpReg =
I.getOperand(1).getReg();
3968 return selectBitreverse16(ResVReg, ResType,
I, OpReg);
3970 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
3972 return selectBitreverse64(ResVReg, ResType,
I, OpReg);
3974 return SPIRVInstructionSelector::diagnoseUnsupported(
3975 I,
"G_BITREVERSE only support 16,32,64 bits.");
3979 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
3990 unsigned AndOp = SPIRV::OpBitwiseAndS;
3991 unsigned OrOp = SPIRV::OpBitwiseOrS;
3992 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
3993 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
3995 AndOp = SPIRV::OpBitwiseAndV;
3996 OrOp = SPIRV::OpBitwiseOrV;
3997 ShlOp = SPIRV::OpShiftLeftLogicalV;
3998 ShrOp = SPIRV::OpShiftRightLogicalV;
4004 const unsigned Shift) ->
Register {
4012 Register MaskReg = CreateConst(Mask);
4013 Register ShiftReg = CreateConst(Shift);
4020 if (!selectOpWithSrcs(
T1, ResType,
I, {Input, ShiftReg}, ShrOp) ||
4021 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
4022 !selectOpWithSrcs(T3, ResType,
I, {Input, MaskReg}, AndOp) ||
4023 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
4024 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
4032 uint64_t
Mask = ~0ull;
4033 while ((Shift >>= 1) > 0) {
4040 return BuildCOPY(ResVReg, Result,
I);
4043bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
4044 SPIRVTypeInst ResType,
4045 MachineInstr &
I)
const {
4046 assert(
I.getOperand(0).isReg() &&
I.getOperand(1).isReg() &&
4047 "G_FREEZE must define and use a register");
4048 Register OpReg =
I.getOperand(1).getReg();
4052 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
4065 if (MachineInstr *Def = MRI->
getVRegDef(OpReg)) {
4066 if (
Def->getOpcode() == TargetOpcode::COPY)
4069 switch (
Def->getOpcode()) {
4070 case SPIRV::ASSIGN_TYPE:
4071 if (MachineInstr *AssignToDef =
4073 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
4074 Reg =
Def->getOperand(2).getReg();
4077 case SPIRV::OpUndef:
4078 Reg =
Def->getOperand(1).getReg();
4081 unsigned DestOpCode;
4083 DestOpCode = SPIRV::OpConstantNull;
4084 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze of a "
4085 "static undef/poison lowered to OpConstantNull\n");
4087 DestOpCode = TargetOpcode::COPY;
4089 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze "
4090 "skipped, lowered as a copy of the operand\n");
4092 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
4093 .
addDef(
I.getOperand(0).getReg())
4101bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
4102 SPIRVTypeInst ResType,
4103 MachineInstr &
I)
const {
4105 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4107 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4111 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
4116 for (
unsigned i =
I.getNumExplicitDefs();
4117 i <
I.getNumExplicitOperands() && IsConst; ++i)
4121 if (!IsConst &&
N < 2)
4122 return diagnoseUnsupported(
4123 I,
"There must be at least two constituent operands in a vector");
4128 for (
unsigned i =
I.getNumExplicitDefs();
4129 i <
I.getNumExplicitOperands() && IsNullVector; ++i) {
4130 MachineInstr *
Def =
getDef(
I.getOperand(i), MRI);
4135 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4142 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4143 TII.get(IsConst ? SPIRV::OpConstantComposite
4144 : SPIRV::OpCompositeConstruct))
4147 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
4148 MIB.
addUse(
I.getOperand(i).getReg());
4153bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
4154 SPIRVTypeInst ResType,
4155 MachineInstr &
I)
const {
4157 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4159 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4165 if (!
I.getOperand(
OpIdx).isReg())
4172 if (!IsConst &&
N < 2)
4173 return diagnoseUnsupported(
4174 I,
"There must be at least two constituent operands in a vector");
4177 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4178 TII.get(IsConst ? SPIRV::OpConstantComposite
4179 : SPIRV::OpCompositeConstruct))
4182 for (
unsigned i = 0; i <
N; ++i)
4188bool SPIRVInstructionSelector::selectConcatVectors(
Register ResVReg,
4189 SPIRVTypeInst ResType,
4190 MachineInstr &
I)
const {
4194 if (ResType->
getOpcode() != SPIRV::OpTypeVector)
4196 "Cannot select G_CONCAT_VECTORS with a non-vector result");
4198 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4199 TII.get(SPIRV::OpCompositeConstruct))
4209bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
4210 SPIRVTypeInst ResType,
4211 MachineInstr &
I)
const {
4216 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
4218 Opcode = SPIRV::OpDemoteToHelperInvocation;
4220 Opcode = SPIRV::OpKill;
4222 if (MachineInstr *NextI =
I.getNextNode()) {
4224 NextI->eraseFromParent();
4234bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
4235 SPIRVTypeInst ResType,
unsigned CmpOpc,
4236 MachineInstr &
I)
const {
4237 Register Cmp0 =
I.getOperand(2).getReg();
4238 Register Cmp1 =
I.getOperand(3).getReg();
4241 "CMP operands should have the same type");
4242 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
4252bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
4253 SPIRVTypeInst ResType,
4254 MachineInstr &
I)
const {
4255 auto Pred =
I.getOperand(1).getPredicate();
4258 Register CmpOperand =
I.getOperand(2).getReg();
4263 Register Op1 =
I.getOperand(3).getReg();
4267 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
4272 I.getOperand(3).setReg(NewOp1);
4278 return selectCmp(ResVReg, ResType, CmpOpc,
I);
4282SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
4283 SPIRVTypeInst ResType)
const {
4285 SPIRVTypeInst SpvI32Ty =
4288 auto ConstInt = ConstantInt::get(LLVMTy, Val);
4295 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4298 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
4301 .
addImm(APInt(32, Val).getZExtValue());
4303 GR.
add(ConstInt,
MI);
4310Register SPIRVInstructionSelector::buildI32ConstantInEntryBlock(
4311 uint32_t Val, MachineInstr &
I, SPIRVTypeInst ResType)
const {
4313 SPIRVTypeInst SpvI32Ty =
4315 auto *ConstInt = ConstantInt::get(LLVMTy, Val);
4320 MachineBasicBlock &EntryBB = *InsertIt->getParent();
4321 MachineInstr *
MI =
nullptr;
4325 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantNull))
4329 uint64_t ImmVal = APInt(32, Val).getZExtValue();
4330 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantI))
4336 GR.
add(ConstInt,
MI);
4341bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
4342 SPIRVTypeInst ResType,
4343 MachineInstr &
I)
const {
4345 return selectCmp(ResVReg, ResType, CmpOp,
I);
4348bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
4349 SPIRVTypeInst ResType,
4350 MachineInstr &
I)
const {
4352 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
4359 if (ResType->
getOpcode() != SPIRV::OpTypeVector &&
4360 ResType->
getOpcode() != SPIRV::OpTypeFloat)
4363 MachineIRBuilder MIRBuilder(
I);
4370 APFloat ConstVal(3.3219280948873623);
4374 APFloat::rmNearestTiesToEven, &LosesInfo);
4378 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
4379 ? SPIRV::OpVectorTimesScalar
4382 if (!selectOpWithSrcs(ArgReg, ResType,
I,
4383 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
4385 if (!selectExtInst(ResVReg, ResType,
I,
4386 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
4396Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
4397 MachineInstr &
I)
const {
4400 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4405bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
4411 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
4419 if (
Def->getOpcode() == SPIRV::OpConstantNull)
4422 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
4423 Def->getOpcode() == SPIRV::OpConstantI)
4436 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
4437 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
4439 Intrinsic::spv_const_composite)) {
4440 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
4441 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
4442 if (!IsZero(
Def->getOperand(i).getReg()))
4451Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
4452 MachineInstr &
I)
const {
4456 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4461Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
4462 MachineInstr &
I)
const {
4466 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4472 SPIRVTypeInst ResType,
4473 MachineInstr &
I)
const {
4477 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4482bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
4483 SPIRVTypeInst ResType,
4484 MachineInstr &
I)
const {
4485 Register SelectFirstArg =
I.getOperand(2).getReg();
4486 Register SelectSecondArg =
I.getOperand(3).getReg();
4495 SPIRV::OpTypeVector;
4502 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
4503 }
else if (IsPtrTy) {
4504 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
4506 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
4509 assert(IsScalarBool &&
"OpSelect with a scalar result requires a scalar "
4510 "boolean condition");
4512 Opcode = SPIRV::OpSelectSFSCond;
4513 }
else if (IsPtrTy) {
4514 Opcode = SPIRV::OpSelectSPSCond;
4516 Opcode = SPIRV::OpSelectSISCond;
4519 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4522 .
addUse(
I.getOperand(1).getReg())
4531bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
4532 SPIRVTypeInst ResType,
4534 MachineInstr &InsertAt,
4535 bool IsSigned)
const {
4537 Register ZeroReg = buildZerosVal(ResType, InsertAt);
4538 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
4539 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
4541 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4553bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
4554 SPIRVTypeInst ResType,
4555 MachineInstr &
I,
bool IsSigned,
4556 unsigned Opcode)
const {
4557 Register SrcReg =
I.getOperand(1).getReg();
4563 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
4568 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I,
false);
4570 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
4573bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
4574 SPIRVTypeInst ResType, MachineInstr &
I,
4575 bool IsSigned)
const {
4576 Register SrcReg =
I.getOperand(1).getReg();
4578 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
4582 if (ResType == SrcType)
4583 return BuildCOPY(ResVReg, SrcReg,
I);
4585 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4586 return selectUnOp(ResVReg, ResType,
I, Opcode);
4589bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
4590 SPIRVTypeInst ResType,
4592 bool IsSigned)
const {
4593 MachineIRBuilder MIRBuilder(
I);
4594 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
4606 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
4609 .
addUse(
I.getOperand(1).getReg())
4610 .
addUse(
I.getOperand(2).getReg())
4615 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
4618 .
addUse(
I.getOperand(1).getReg())
4619 .
addUse(
I.getOperand(2).getReg())
4627 unsigned SelectOpcode =
4628 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
4633 .
addUse(buildOnesVal(
true, ResType,
I))
4634 .
addUse(buildZerosVal(ResType,
I))
4641 .
addUse(buildOnesVal(
false, ResType,
I))
4646bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
4649 SPIRVTypeInst IntTy,
4650 SPIRVTypeInst BoolTy)
const {
4653 bool IsVectorTy = IntTy->
getOpcode() == SPIRV::OpTypeVector;
4654 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
4656 Register One = buildOnesVal(
false, IntTy,
I);
4664 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
4673bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
4674 SPIRVTypeInst ResType,
4675 MachineInstr &
I)
const {
4676 Register IntReg =
I.getOperand(1).getReg();
4679 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
4680 if (ArgType == ResType)
4681 return BuildCOPY(ResVReg, IntReg,
I);
4683 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4684 return selectUnOp(ResVReg, ResType,
I, Opcode);
4687bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
4688 SPIRVTypeInst ResType,
4689 MachineInstr &
I)
const {
4690 unsigned Opcode =
I.getOpcode();
4691 unsigned TpOpcode = ResType->
getOpcode();
4693 if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {
4694 assert(Opcode == TargetOpcode::G_CONSTANT &&
4695 I.getOperand(1).getCImm()->isZero());
4696 MachineBasicBlock &DepMBB =
I.getMF()->front();
4699 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
4706 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
4709bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
4710 SPIRVTypeInst ResType,
4711 MachineInstr &
I)
const {
4712 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4719bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
4720 SPIRVTypeInst ResType,
4721 MachineInstr &
I)
const {
4723 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
4727 .
addUse(
I.getOperand(3).getReg())
4729 .
addUse(
I.getOperand(2).getReg());
4730 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
4736bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
4737 SPIRVTypeInst ResType,
4738 MachineInstr &
I)
const {
4739 Type *MaybeResTy =
nullptr;
4744 "Expected aggregate type for extractv instruction");
4746 SPIRV::AccessQualifier::ReadWrite,
false);
4750 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
4753 .
addUse(
I.getOperand(2).getReg());
4754 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
4760bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
4761 SPIRVTypeInst ResType,
4762 MachineInstr &
I)
const {
4763 if (
getImm(
I.getOperand(4), MRI))
4764 return selectInsertVal(ResVReg, ResType,
I);
4766 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
4769 .
addUse(
I.getOperand(2).getReg())
4770 .
addUse(
I.getOperand(3).getReg())
4771 .
addUse(
I.getOperand(4).getReg())
4776bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
4777 SPIRVTypeInst ResType,
4778 MachineInstr &
I)
const {
4779 if (
getImm(
I.getOperand(3), MRI))
4780 return selectExtractVal(ResVReg, ResType,
I);
4782 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
4785 .
addUse(
I.getOperand(2).getReg())
4786 .
addUse(
I.getOperand(3).getReg())
4791bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
4792 SPIRVTypeInst ResType,
4793 MachineInstr &
I)
const {
4794 const bool IsGEPInBounds =
I.getOperand(2).getImm();
4800 ? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
4801 : SPIRV::OpAccessChain)
4802 : (IsGEPInBounds ?
SPIRV::OpInBoundsPtrAccessChain
4803 :
SPIRV::OpPtrAccessChain);
4805 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4809 .
addUse(
I.getOperand(3).getReg());
4811 (Opcode == SPIRV::OpPtrAccessChain ||
4812 Opcode == SPIRV::OpInBoundsPtrAccessChain ||
4813 (
getImm(
I.getOperand(4), MRI) &&
foldImm(
I.getOperand(4), MRI) == 0)) &&
4814 "Cannot translate GEP to OpAccessChain. First index must be 0.");
4817 const unsigned StartingIndex =
4818 (Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)
4821 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
4822 Res.addUse(
I.getOperand(i).getReg());
4823 Res.constrainAllUses(
TII,
TRI, RBI);
4828bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
4830 unsigned Lim =
I.getNumExplicitOperands();
4831 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
4832 Register OpReg =
I.getOperand(i).getReg();
4833 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
4835 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
4836 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
4837 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
4844 MachineFunction *MF =
I.getMF();
4856 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4857 TII.get(SPIRV::OpSpecConstantOp))
4860 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
4862 GR.
add(OpDefine, MIB);
4868bool SPIRVInstructionSelector::selectDerivativeInst(
4869 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
4870 const unsigned DPdOpCode)
const {
4873 if (!errorIfInstrOutsideShader(
I))
4879 Register SrcReg =
I.getOperand(2).getReg();
4884 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4887 .
addUse(
I.getOperand(2).getReg());
4889 MachineIRBuilder MIRBuilder(
I);
4892 if (componentCount != 1)
4896 const TargetRegisterClass *RegClass = GR.
getRegClass(SrcType);
4900 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4905 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4910 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4918bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
4919 SPIRVTypeInst ResType,
4920 MachineInstr &
I)
const {
4924 case Intrinsic::spv_load:
4925 return selectLoad(ResVReg, ResType,
I);
4926 case Intrinsic::spv_atomic_load:
4927 return selectAtomicLoad(ResVReg, ResType,
I);
4928 case Intrinsic::spv_store:
4929 return selectStore(
I);
4930 case Intrinsic::spv_atomic_store:
4931 return selectAtomicStore(
I);
4932 case Intrinsic::spv_extractv:
4933 return selectExtractVal(ResVReg, ResType,
I);
4934 case Intrinsic::spv_insertv:
4935 return selectInsertVal(ResVReg, ResType,
I);
4936 case Intrinsic::spv_extractelt:
4937 return selectExtractElt(ResVReg, ResType,
I);
4938 case Intrinsic::spv_insertelt:
4939 return selectInsertElt(ResVReg, ResType,
I);
4940 case Intrinsic::spv_gep:
4941 return selectGEP(ResVReg, ResType,
I);
4942 case Intrinsic::spv_bitcast: {
4943 Register OpReg =
I.getOperand(2).getReg();
4944 SPIRVTypeInst OpType =
4948 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
4950 case Intrinsic::spv_unref_global:
4951 case Intrinsic::spv_init_global: {
4952 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
4957 Register GVarVReg =
MI->getOperand(0).getReg();
4958 if (!selectGlobalValue(GVarVReg, *
MI, Init))
4963 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
4965 MI->eraseFromParent();
4969 case Intrinsic::spv_undef: {
4970 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4976 case Intrinsic::spv_poison:
4977 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpPoisonKHR))
4982 case Intrinsic::spv_freeze:
4983 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
4986 .
addUse(
I.getOperand(2).getReg())
4989 case Intrinsic::spv_named_boolean_spec_constant: {
4990 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
4991 : SPIRV::OpSpecConstantFalse;
4993 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
4994 .
addDef(
I.getOperand(0).getReg())
4997 unsigned SpecId =
I.getOperand(2).getImm();
4999 SPIRV::Decoration::SpecId, {SpecId});
5003 case Intrinsic::spv_const_composite: {
5005 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
5011 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
5013 std::function<bool(
Register)> HasSpecConstOperand =
5023 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
5024 J < Def->getNumExplicitOperands(); ++J) {
5025 if (
Def->getOperand(J).isReg() &&
5026 HasSpecConstOperand(
Def->getOperand(J).getReg()))
5032 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
5033 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
5034 : SPIRV::OpConstantComposite;
5035 unsigned ContinuedOpc = HasSpecConst
5036 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
5037 : SPIRV::OpConstantCompositeContinuedINTEL;
5038 MachineIRBuilder MIR(
I);
5040 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
5042 for (
auto *Instr : Instructions) {
5043 Instr->setDebugLoc(
I.getDebugLoc());
5048 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5055 case Intrinsic::spv_assign_name: {
5056 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
5057 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
5058 for (
unsigned i =
I.getNumExplicitDefs() + 2;
5059 i <
I.getNumExplicitOperands(); ++i) {
5060 MIB.
addImm(
I.getOperand(i).getImm());
5065 case Intrinsic::spv_switch: {
5066 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
5067 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5068 if (
I.getOperand(i).isReg())
5069 MIB.
addReg(
I.getOperand(i).getReg());
5070 else if (
I.getOperand(i).isCImm())
5071 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
5072 else if (
I.getOperand(i).isMBB())
5073 MIB.
addMBB(
I.getOperand(i).getMBB());
5080 case Intrinsic::spv_loop_merge: {
5081 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
5082 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5083 if (
I.getOperand(i).isMBB())
5084 MIB.
addMBB(
I.getOperand(i).getMBB());
5091 case Intrinsic::spv_loop_control_intel: {
5093 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
5094 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
5099 case Intrinsic::spv_selection_merge: {
5101 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
5102 assert(
I.getOperand(1).isMBB() &&
5103 "operand 1 to spv_selection_merge must be a basic block");
5104 MIB.
addMBB(
I.getOperand(1).getMBB());
5105 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
5109 case Intrinsic::spv_cmpxchg:
5110 return selectAtomicCmpXchg(ResVReg, ResType,
I);
5111 case Intrinsic::spv_unreachable:
5112 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
5115 case Intrinsic::spv_abort:
5116 return selectAbort(
I);
5117 case Intrinsic::spv_alloca:
5118 return selectFrameIndex(ResVReg, ResType,
I);
5119 case Intrinsic::spv_alloca_array:
5120 return selectAllocaArray(ResVReg, ResType,
I);
5121 case Intrinsic::spv_assume:
5123 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
5124 .
addUse(
I.getOperand(1).getReg())
5129 case Intrinsic::spv_expect:
5131 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
5134 .
addUse(
I.getOperand(2).getReg())
5135 .
addUse(
I.getOperand(3).getReg())
5140 case Intrinsic::arithmetic_fence:
5141 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
5142 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
5145 .
addUse(
I.getOperand(2).getReg())
5149 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
5151 case Intrinsic::spv_thread_id:
5157 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
5159 case Intrinsic::spv_thread_id_in_group:
5165 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
5167 case Intrinsic::spv_group_id:
5173 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
5175 case Intrinsic::spv_flattened_thread_id_in_group:
5182 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
5184 case Intrinsic::spv_workgroup_size:
5185 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
5187 case Intrinsic::spv_global_size:
5188 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
5190 case Intrinsic::spv_global_offset:
5191 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
5193 case Intrinsic::spv_num_workgroups:
5194 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
5196 case Intrinsic::spv_subgroup_size:
5197 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
5199 case Intrinsic::spv_num_subgroups:
5200 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
5202 case Intrinsic::spv_subgroup_id:
5203 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
5204 case Intrinsic::spv_subgroup_local_invocation_id:
5205 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
5206 ResVReg, ResType,
I);
5207 case Intrinsic::spv_subgroup_max_size:
5208 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
5210 case Intrinsic::spv_fdot:
5211 return selectFloatDot(ResVReg, ResType,
I);
5212 case Intrinsic::spv_udot:
5213 case Intrinsic::spv_sdot:
5214 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5216 return selectIntegerDot(ResVReg, ResType,
I,
5217 IID == Intrinsic::spv_sdot);
5218 return selectIntegerDotExpansion(ResVReg, ResType,
I);
5219 case Intrinsic::spv_dot4add_i8packed:
5220 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5222 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
5223 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
5224 case Intrinsic::spv_dot4add_u8packed:
5225 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5227 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
5228 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
5229 case Intrinsic::spv_all:
5230 return selectAll(ResVReg, ResType,
I);
5231 case Intrinsic::spv_any:
5232 return selectAny(ResVReg, ResType,
I);
5233 case Intrinsic::spv_cross:
5234 return selectExtInst(ResVReg, ResType,
I, CL::cross, GL::Cross);
5235 case Intrinsic::spv_distance:
5236 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
5237 case Intrinsic::spv_lerp:
5238 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
5239 case Intrinsic::spv_length:
5240 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
5241 case Intrinsic::spv_degrees:
5242 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
5243 case Intrinsic::spv_faceforward:
5244 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
5245 case Intrinsic::spv_frac:
5246 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
5247 case Intrinsic::spv_isinf:
5248 return selectOpIsInf(ResVReg, ResType,
I);
5249 case Intrinsic::spv_isnan:
5250 return selectOpIsNan(ResVReg, ResType,
I);
5251 case Intrinsic::spv_isfinite:
5252 return selectOpIsFinite(ResVReg, ResType,
I);
5253 case Intrinsic::spv_isnormal:
5254 return selectOpIsNormal(ResVReg, ResType,
I);
5255 case Intrinsic::spv_normalize:
5256 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
5257 case Intrinsic::spv_refract:
5258 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
5259 case Intrinsic::spv_reflect:
5260 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
5261 case Intrinsic::spv_rsqrt:
5262 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
5263 case Intrinsic::spv_sign:
5264 return selectSign(ResVReg, ResType,
I);
5265 case Intrinsic::spv_smoothstep:
5266 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
5267 case Intrinsic::spv_firstbituhigh:
5268 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
5269 case Intrinsic::spv_firstbitshigh:
5270 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
5271 case Intrinsic::spv_firstbitlow:
5272 return selectFirstBitLow(ResVReg, ResType,
I);
5273 case Intrinsic::spv_all_memory_barrier:
5274 return selectBarrierInst(
I, SPIRV::Scope::Device,
5275 SPIRV::MemorySemantics::UniformMemory |
5276 SPIRV::MemorySemantics::ImageMemory |
5277 SPIRV::MemorySemantics::WorkgroupMemory,
5279 case Intrinsic::spv_all_memory_barrier_with_group_sync:
5280 return selectBarrierInst(
I, SPIRV::Scope::Device,
5281 SPIRV::MemorySemantics::UniformMemory |
5282 SPIRV::MemorySemantics::ImageMemory |
5283 SPIRV::MemorySemantics::WorkgroupMemory,
5285 case Intrinsic::spv_device_memory_barrier:
5286 return selectBarrierInst(
I, SPIRV::Scope::Device,
5287 SPIRV::MemorySemantics::UniformMemory |
5288 SPIRV::MemorySemantics::ImageMemory,
5290 case Intrinsic::spv_device_memory_barrier_with_group_sync:
5291 return selectBarrierInst(
I, SPIRV::Scope::Device,
5292 SPIRV::MemorySemantics::UniformMemory |
5293 SPIRV::MemorySemantics::ImageMemory,
5295 case Intrinsic::spv_group_memory_barrier:
5296 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5297 SPIRV::MemorySemantics::WorkgroupMemory,
5299 case Intrinsic::spv_group_memory_barrier_with_group_sync:
5300 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5301 SPIRV::MemorySemantics::WorkgroupMemory,
5303 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
5304 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
5305 SPIRV::StorageClass::StorageClass ResSC =
5308 return diagnoseUnsupported(
I,
"The target storage class is not castable "
5309 "from the Generic storage class");
5310 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
5318 case Intrinsic::spv_lifetime_start:
5319 case Intrinsic::spv_lifetime_end: {
5320 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
5321 : SPIRV::OpLifetimeStop;
5322 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
5323 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
5332 case Intrinsic::spv_saturate:
5333 return selectSaturate(ResVReg, ResType,
I);
5334 case Intrinsic::spv_nclamp:
5335 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
5336 case Intrinsic::spv_uclamp:
5337 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
5338 case Intrinsic::spv_sclamp:
5339 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
5340 case Intrinsic::spv_subgroup_prefix_bit_count:
5341 return selectWavePrefixBitCount(ResVReg, ResType,
I);
5342 case Intrinsic::spv_wave_active_countbits:
5343 return selectWaveActiveCountBits(ResVReg, ResType,
I);
5344 case Intrinsic::spv_wave_all_equal:
5345 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
5346 case Intrinsic::spv_wave_all:
5347 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
5348 case Intrinsic::spv_wave_any:
5349 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
5350 case Intrinsic::spv_subgroup_ballot:
5351 return selectWaveOpInst(ResVReg, ResType,
I,
5352 SPIRV::OpGroupNonUniformBallot);
5353 case Intrinsic::spv_wave_is_first_lane:
5354 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
5355 case Intrinsic::spv_wave_reduce_or:
5356 return selectWaveReduceOp(ResVReg, ResType,
I,
5357 SPIRV::OpGroupNonUniformBitwiseOr);
5358 case Intrinsic::spv_wave_reduce_xor:
5359 return selectWaveReduceOp(ResVReg, ResType,
I,
5360 SPIRV::OpGroupNonUniformBitwiseXor);
5361 case Intrinsic::spv_wave_reduce_and:
5362 return selectWaveReduceOp(ResVReg, ResType,
I,
5363 SPIRV::OpGroupNonUniformBitwiseAnd);
5364 case Intrinsic::spv_interlocked_add:
5365 return selectInterlockedAdd(ResVReg, ResType,
I);
5366 case Intrinsic::spv_wave_reduce_umax:
5367 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
5368 case Intrinsic::spv_wave_reduce_max:
5369 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
5370 case Intrinsic::spv_wave_reduce_umin:
5371 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
5372 case Intrinsic::spv_wave_reduce_min:
5373 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
5374 case Intrinsic::spv_wave_reduce_sum:
5375 return selectWaveReduceSum(ResVReg, ResType,
I);
5376 case Intrinsic::spv_wave_product:
5377 return selectWaveReduceProduct(ResVReg, ResType,
I);
5378 case Intrinsic::spv_wave_readlane:
5379 return selectWaveOpInst(ResVReg, ResType,
I,
5380 SPIRV::OpGroupNonUniformShuffle);
5381 case Intrinsic::spv_wave_prefix_sum:
5382 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
5383 case Intrinsic::spv_wave_prefix_product:
5384 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
5385 case Intrinsic::spv_quad_read_across_x: {
5386 return selectQuadSwap(ResVReg, ResType,
I, 0);
5388 case Intrinsic::spv_quad_read_across_y: {
5389 return selectQuadSwap(ResVReg, ResType,
I, 1);
5391 case Intrinsic::spv_quad_read_across_diagonal: {
5392 return selectQuadSwap(ResVReg, ResType,
I, 2);
5394 case Intrinsic::spv_step:
5395 return selectExtInst(ResVReg, ResType,
I, CL::step, GL::Step);
5396 case Intrinsic::spv_radians:
5397 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
5401 case Intrinsic::instrprof_increment:
5402 case Intrinsic::instrprof_increment_step:
5403 case Intrinsic::instrprof_value_profile:
5406 case Intrinsic::spv_value_md:
5408 case Intrinsic::spv_resource_handlefrombinding: {
5409 return selectHandleFromBinding(ResVReg, ResType,
I);
5411 case Intrinsic::spv_resource_counterhandlefrombinding:
5412 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
5413 case Intrinsic::spv_resource_updatecounter:
5414 return selectUpdateCounter(ResVReg, ResType,
I);
5415 case Intrinsic::spv_resource_store_typedbuffer: {
5416 return selectImageWriteIntrinsic(
I);
5418 case Intrinsic::spv_resource_load_typedbuffer: {
5419 return selectReadImageIntrinsic(ResVReg, ResType,
I);
5421 case Intrinsic::spv_resource_load_level: {
5422 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
5424 case Intrinsic::spv_resource_getdimensions_x:
5425 case Intrinsic::spv_resource_getdimensions_xy:
5426 case Intrinsic::spv_resource_getdimensions_xyz: {
5427 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
5429 case Intrinsic::spv_resource_getdimensions_levels_x:
5430 case Intrinsic::spv_resource_getdimensions_levels_xy:
5431 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
5432 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
5434 case Intrinsic::spv_resource_getdimensions_ms_xy:
5435 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
5436 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
5438 case Intrinsic::spv_resource_calculate_lod:
5439 case Intrinsic::spv_resource_calculate_lod_unclamped:
5440 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
5441 case Intrinsic::spv_resource_sample:
5442 case Intrinsic::spv_resource_sample_clamp:
5443 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
5444 case Intrinsic::spv_resource_samplebias:
5445 case Intrinsic::spv_resource_samplebias_clamp:
5446 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
5447 case Intrinsic::spv_resource_samplegrad:
5448 case Intrinsic::spv_resource_samplegrad_clamp:
5449 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
5450 case Intrinsic::spv_resource_samplelevel:
5451 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
5452 case Intrinsic::spv_resource_samplecmp:
5453 case Intrinsic::spv_resource_samplecmp_clamp:
5454 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
5455 case Intrinsic::spv_resource_samplecmplevelzero:
5456 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
5457 case Intrinsic::spv_resource_gather:
5458 case Intrinsic::spv_resource_gather_cmp:
5459 return selectGatherIntrinsic(ResVReg, ResType,
I);
5460 case Intrinsic::spv_resource_getbasepointer:
5461 case Intrinsic::spv_resource_getpointer: {
5462 return selectResourceGetPointer(ResVReg, ResType,
I);
5464 case Intrinsic::spv_pushconstant_getpointer: {
5465 return selectPushConstantGetPointer(ResVReg, ResType,
I);
5467 case Intrinsic::spv_discard: {
5468 return selectDiscard(ResVReg, ResType,
I);
5470 case Intrinsic::spv_resource_nonuniformindex: {
5471 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
5473 case Intrinsic::spv_unpackhalf2x16: {
5474 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
5476 case Intrinsic::spv_packhalf2x16: {
5477 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
5479 case Intrinsic::spv_ddx:
5480 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
5481 case Intrinsic::spv_ddy:
5482 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
5483 case Intrinsic::spv_ddx_coarse:
5484 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
5485 case Intrinsic::spv_ddy_coarse:
5486 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
5487 case Intrinsic::spv_ddx_fine:
5488 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
5489 case Intrinsic::spv_ddy_fine:
5490 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
5491 case Intrinsic::spv_fwidth:
5492 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
5493 case Intrinsic::spv_masked_gather:
5494 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5495 return selectMaskedGather(ResVReg, ResType,
I);
5496 return diagnoseUnsupported(
5497 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
5498 case Intrinsic::spv_masked_scatter:
5499 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5500 return selectMaskedScatter(
I);
5501 return diagnoseUnsupported(
5502 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
5503 case Intrinsic::returnaddress:
5504 case Intrinsic::frameaddress: {
5506 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5513 return diagnoseUnsupported(
I,
"intrinsic selection not implemented.");
5518bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
5519 SPIRVTypeInst ResType,
5520 MachineInstr &
I)
const {
5523 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
5530bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
5531 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5533 assert(Intr.getIntrinsicID() ==
5534 Intrinsic::spv_resource_counterhandlefrombinding);
5537 Register MainHandleReg = Intr.getOperand(2).getReg();
5539 assert(MainHandleDef->getIntrinsicID() ==
5540 Intrinsic::spv_resource_handlefrombinding);
5544 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
5545 Register IndexReg = MainHandleDef->getOperand(5).getReg();
5546 std::string CounterName =
5551 MachineIRBuilder MIRBuilder(
I);
5553 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
5555 ArraySize, IndexReg, CounterName, MIRBuilder);
5557 return BuildCOPY(ResVReg, CounterVarReg,
I);
5560bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
5561 SPIRVTypeInst ResType,
5562 MachineInstr &
I)
const {
5564 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
5566 Register CounterHandleReg = Intr.getOperand(2).getReg();
5567 Register IncrReg = Intr.getOperand(3).getReg();
5574 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
5575 assert(CounterVarPointeeType &&
5576 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
5577 "Counter variable must be a struct");
5579 SPIRV::StorageClass::StorageBuffer &&
5580 "Counter variable must be in the storage buffer storage class");
5582 "Counter variable must have exactly 1 member in the struct");
5583 const SPIRVTypeInst MemberType =
5586 "Counter variable struct must have a single i32 member");
5590 MachineIRBuilder MIRBuilder(
I);
5592 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
5595 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
5601 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5604 .
addUse(CounterHandleReg)
5611 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
5614 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
5617 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
5626 return BuildCOPY(ResVReg, AtomicRes,
I);
5634 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
5642bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
5643 SPIRVTypeInst ResType,
5644 MachineInstr &
I)
const {
5652 Register ImageReg =
I.getOperand(2).getReg();
5660 Register IdxReg =
I.getOperand(3).getReg();
5662 MachineInstr &Pos =
I;
5664 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
5668bool SPIRVInstructionSelector::generateSampleImage(
5671 DebugLoc Loc, MachineInstr &Pos)
const {
5682 if (!loadHandleBeforePosition(NewSamplerReg,
5688 MachineIRBuilder MIRBuilder(Pos);
5701 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
5702 ImOps.Lod.has_value();
5703 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
5704 : SPIRV::OpImageSampleImplicitLod;
5706 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
5707 : SPIRV::OpImageSampleDrefImplicitLod;
5716 MIB.
addUse(*ImOps.Compare);
5718 uint32_t ImageOperands = 0;
5720 ImageOperands |= SPIRV::ImageOperand::Bias;
5722 ImageOperands |= SPIRV::ImageOperand::Lod;
5723 if (ImOps.GradX && ImOps.GradY)
5724 ImageOperands |= SPIRV::ImageOperand::Grad;
5725 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
5727 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
5730 "Non-constant offsets are not supported in sample instructions.");
5735 ImageOperands |= SPIRV::ImageOperand::MinLod;
5737 if (ImageOperands != 0) {
5738 MIB.
addImm(ImageOperands);
5739 if (ImageOperands & SPIRV::ImageOperand::Bias)
5741 if (ImageOperands & SPIRV::ImageOperand::Lod)
5743 if (ImageOperands & SPIRV::ImageOperand::Grad) {
5744 MIB.
addUse(*ImOps.GradX);
5745 MIB.
addUse(*ImOps.GradY);
5748 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
5749 MIB.
addUse(*ImOps.Offset);
5750 if (ImageOperands & SPIRV::ImageOperand::MinLod)
5751 MIB.
addUse(*ImOps.MinLod);
5758bool SPIRVInstructionSelector::selectImageQuerySize(
5760 std::optional<Register> LodReg)
const {
5762 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
5765 "ImageReg is not an image type.");
5767 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
5769 unsigned NumComponents = 0;
5771 case SPIRV::Dim::DIM_1D:
5772 case SPIRV::Dim::DIM_Buffer:
5773 NumComponents =
IsArray ? 2 : 1;
5775 case SPIRV::Dim::DIM_2D:
5776 case SPIRV::Dim::DIM_Cube:
5777 case SPIRV::Dim::DIM_Rect:
5778 NumComponents =
IsArray ? 3 : 2;
5780 case SPIRV::Dim::DIM_3D:
5784 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
5789 SPIRVTypeInst ResType =
5794 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5804bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
5805 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5806 Register ImageReg =
I.getOperand(2).getReg();
5813 return selectImageQuerySize(NewImageReg, ResVReg,
I);
5816bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
5817 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5818 Register ImageReg =
I.getOperand(2).getReg();
5827 Register LodReg =
I.getOperand(3).getReg();
5830 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
5832 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
5839 TII.get(SPIRV::OpImageQueryLevels))
5846 TII.get(SPIRV::OpCompositeConstruct))
5856bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
5857 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5858 Register ImageReg =
I.getOperand(2).getReg();
5869 "OpImageQuerySamples requires a multisampled image");
5871 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
5879 TII.get(SPIRV::OpImageQuerySamples))
5886 TII.get(SPIRV::OpCompositeConstruct))
5896bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
5897 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5898 Register ImageReg =
I.getOperand(2).getReg();
5899 Register SamplerReg =
I.getOperand(3).getReg();
5900 Register CoordinateReg =
I.getOperand(4).getReg();
5916 if (!loadHandleBeforePosition(
5921 MachineIRBuilder MIRBuilder(
I);
5927 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
5937 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
5944 unsigned ExtractedIndex =
5946 Intrinsic::spv_resource_calculate_lod_unclamped
5950 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5951 TII.get(SPIRV::OpCompositeExtract))
5961bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
5962 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5963 Register ImageReg =
I.getOperand(2).getReg();
5964 Register SamplerReg =
I.getOperand(3).getReg();
5965 Register CoordinateReg =
I.getOperand(4).getReg();
5966 ImageOperands ImOps;
5967 if (
I.getNumOperands() > 5)
5968 ImOps.Offset =
I.getOperand(5).getReg();
5969 if (
I.getNumOperands() > 6)
5970 ImOps.MinLod =
I.getOperand(6).getReg();
5971 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5972 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5975bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
5976 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5977 Register ImageReg =
I.getOperand(2).getReg();
5978 Register SamplerReg =
I.getOperand(3).getReg();
5979 Register CoordinateReg =
I.getOperand(4).getReg();
5980 ImageOperands ImOps;
5981 ImOps.Bias =
I.getOperand(5).getReg();
5982 if (
I.getNumOperands() > 6)
5983 ImOps.Offset =
I.getOperand(6).getReg();
5984 if (
I.getNumOperands() > 7)
5985 ImOps.MinLod =
I.getOperand(7).getReg();
5986 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5987 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5990bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
5991 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5992 Register ImageReg =
I.getOperand(2).getReg();
5993 Register SamplerReg =
I.getOperand(3).getReg();
5994 Register CoordinateReg =
I.getOperand(4).getReg();
5995 ImageOperands ImOps;
5996 ImOps.GradX =
I.getOperand(5).getReg();
5997 ImOps.GradY =
I.getOperand(6).getReg();
5998 if (
I.getNumOperands() > 7)
5999 ImOps.Offset =
I.getOperand(7).getReg();
6000 if (
I.getNumOperands() > 8)
6001 ImOps.MinLod =
I.getOperand(8).getReg();
6002 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6003 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6006bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
6007 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6008 Register ImageReg =
I.getOperand(2).getReg();
6009 Register SamplerReg =
I.getOperand(3).getReg();
6010 Register CoordinateReg =
I.getOperand(4).getReg();
6011 ImageOperands ImOps;
6012 ImOps.Lod =
I.getOperand(5).getReg();
6013 if (
I.getNumOperands() > 6)
6014 ImOps.Offset =
I.getOperand(6).getReg();
6015 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6016 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6019bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
6020 SPIRVTypeInst ResType,
6021 MachineInstr &
I)
const {
6022 Register ImageReg =
I.getOperand(2).getReg();
6023 Register SamplerReg =
I.getOperand(3).getReg();
6024 Register CoordinateReg =
I.getOperand(4).getReg();
6025 ImageOperands ImOps;
6026 ImOps.Compare =
I.getOperand(5).getReg();
6027 if (
I.getNumOperands() > 6)
6028 ImOps.Offset =
I.getOperand(6).getReg();
6029 if (
I.getNumOperands() > 7)
6030 ImOps.MinLod =
I.getOperand(7).getReg();
6031 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6032 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6035bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
6036 SPIRVTypeInst ResType,
6037 MachineInstr &
I)
const {
6038 Register ImageReg =
I.getOperand(2).getReg();
6039 Register CoordinateReg =
I.getOperand(3).getReg();
6040 Register LodReg =
I.getOperand(4).getReg();
6042 ImageOperands ImOps;
6044 if (
I.getNumOperands() > 5)
6045 ImOps.Offset =
I.getOperand(5).getReg();
6057 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
6058 I.getDebugLoc(),
I, &ImOps);
6061bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
6062 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6063 Register ImageReg =
I.getOperand(2).getReg();
6064 Register SamplerReg =
I.getOperand(3).getReg();
6065 Register CoordinateReg =
I.getOperand(4).getReg();
6066 ImageOperands ImOps;
6067 ImOps.Compare =
I.getOperand(5).getReg();
6068 if (
I.getNumOperands() > 6)
6069 ImOps.Offset =
I.getOperand(6).getReg();
6072 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6073 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6076bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
6077 SPIRVTypeInst ResType,
6078 MachineInstr &
I)
const {
6079 Register ImageReg =
I.getOperand(2).getReg();
6080 Register SamplerReg =
I.getOperand(3).getReg();
6081 Register CoordinateReg =
I.getOperand(4).getReg();
6084 "ImageReg is not an image type.");
6089 ComponentOrCompareReg =
I.getOperand(5).getReg();
6090 OffsetReg =
I.getOperand(6).getReg();
6093 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
6097 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6098 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
6099 Dim != SPIRV::Dim::DIM_Rect) {
6101 "Gather operations are only supported for 2D, Cube, and Rect images.");
6108 if (!loadHandleBeforePosition(
6113 MachineIRBuilder MIRBuilder(
I);
6114 SPIRVTypeInst SampledImageType =
6119 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6127 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
6129 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
6131 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6136 .
addUse(ComponentOrCompareReg);
6138 uint32_t ImageOperands = 0;
6139 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
6140 if (Dim == SPIRV::Dim::DIM_Cube) {
6142 "Gather operations with offset are not supported for Cube images.");
6146 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6148 ImageOperands |= SPIRV::ImageOperand::Offset;
6152 if (ImageOperands != 0) {
6153 MIB.
addImm(ImageOperands);
6155 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6163bool SPIRVInstructionSelector::generateImageReadOrFetch(
6166 const ImageOperands *ImOps)
const {
6169 "ImageReg is not an image type.");
6171 bool IsSignedInteger =
6176 bool IsFetch = (SampledOp.getImm() == 1);
6178 auto AddOperands = [&](MachineInstrBuilder &MIB) {
6179 uint32_t ImageOperandsMask = 0;
6180 if (IsSignedInteger)
6181 ImageOperandsMask |= 0x1000;
6183 if (IsFetch && ImOps) {
6185 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
6186 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
6188 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
6190 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
6194 if (ImageOperandsMask != 0) {
6195 MIB.
addImm(ImageOperandsMask);
6196 if (IsFetch && ImOps) {
6199 if (ImOps->Offset &&
6200 (ImageOperandsMask &
6201 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
6202 MIB.
addUse(*ImOps->Offset);
6208 if (ResultSize == 4) {
6211 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
6218 BMI.constrainAllUses(
TII,
TRI, RBI);
6222 SPIRVTypeInst ReadType = widenTypeToVec4(ResType, Pos);
6226 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
6232 BMI.constrainAllUses(
TII,
TRI, RBI);
6234 if (ResultSize == 1) {
6243 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
6246bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
6247 SPIRVTypeInst ResType,
6248 MachineInstr &
I)
const {
6249 Register ResourcePtr =
I.getOperand(2).getReg();
6251 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
6260 MachineIRBuilder MIRBuilder(
I);
6265 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6271 if (
I.getNumExplicitOperands() > 3) {
6272 Register IndexReg =
I.getOperand(3).getReg();
6279bool SPIRVInstructionSelector::selectPushConstantGetPointer(
6280 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6285bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
6286 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6287 Register ObjReg =
I.getOperand(2).getReg();
6288 if (!BuildCOPY(ResVReg, ObjReg,
I))
6298 decorateUsesAsNonUniform(ResVReg);
6302void SPIRVInstructionSelector::decorateUsesAsNonUniform(
6305 while (WorkList.
size() > 0) {
6309 bool IsDecorated =
false;
6311 if (
Use.getOpcode() == SPIRV::OpDecorate &&
6312 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
6318 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
6320 if (ResultReg == CurrentReg)
6328 SPIRV::Decoration::NonUniformEXT, {});
6333bool SPIRVInstructionSelector::extractSubvector(
6335 MachineInstr &InsertionPoint)
const {
6337 [[maybe_unused]] uint64_t InputSize =
6340 assert(InputSize > 1 &&
"The input must be a vector.");
6341 assert(ResultSize > 1 &&
"The result must be a vector.");
6342 assert(ResultSize < InputSize &&
6343 "Cannot extract more element than there are in the input.");
6346 const TargetRegisterClass *ScalarRegClass = GR.
getRegClass(ScalarType);
6347 for (uint64_t
I = 0;
I < ResultSize;
I++) {
6350 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6359 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
6361 TII.get(SPIRV::OpCompositeConstruct))
6365 for (
Register ComponentReg : ComponentRegisters)
6366 MIB.
addUse(ComponentReg);
6371bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
6372 MachineInstr &
I)
const {
6379 Register ImageReg =
I.getOperand(1).getReg();
6387 Register CoordinateReg =
I.getOperand(2).getReg();
6388 Register DataReg =
I.getOperand(3).getReg();
6391 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
6399Register SPIRVInstructionSelector::buildPointerToResource(
6400 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
6401 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
6402 StringRef Name, MachineIRBuilder MIRBuilder)
const {
6404 if (ArraySize == 1) {
6405 SPIRVTypeInst PtrType =
6408 "SpirvResType did not have an explicit layout.");
6413 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
6414 SPIRVTypeInst VarPointerType =
6417 VarPointerType, Set,
Binding, Name, MIRBuilder);
6419 SPIRVTypeInst ResPointerType =
6432bool SPIRVInstructionSelector::selectFirstBitSet16(
6433 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
6434 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
6436 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
6440 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
6443bool SPIRVInstructionSelector::selectFirstBitSet32(
6445 unsigned BitSetOpcode)
const {
6446 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6449 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6456bool SPIRVInstructionSelector::selectFirstBitSet64(
6458 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
6471 if (ComponentCount > 2) {
6472 auto Func = [
this, SwapPrimarySide](
Register ResVReg, SPIRVTypeInst ResType,
6474 unsigned Opcode) ->
bool {
6475 return this->selectFirstBitSet64(ResVReg, ResType,
I, SrcReg, Opcode,
6479 return handle64BitOverflow(ResVReg, ResType,
I, SrcReg, BitSetOpcode, Func);
6483 MachineIRBuilder MIRBuilder(
I);
6485 BaseType, 2 * ComponentCount, MIRBuilder,
false);
6489 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
6495 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
6502 bool IsScalarRes = ResType->
getOpcode() != SPIRV::OpTypeVector;
6505 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
6506 SPIRV::OpVectorExtractDynamic))
6508 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
6509 SPIRV::OpVectorExtractDynamic))
6513 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6514 TII.get(SPIRV::OpVectorShuffle))
6522 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
6528 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6529 TII.get(SPIRV::OpVectorShuffle))
6537 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
6557 SelectOp = SPIRV::OpSelectSISCond;
6558 AddOp = SPIRV::OpIAddS;
6566 SelectOp = SPIRV::OpSelectVIVCond;
6567 AddOp = SPIRV::OpIAddV;
6573 Register RegSecondaryOffset = Reg0;
6577 if (SwapPrimarySide) {
6578 PrimaryReg = LowReg;
6579 SecondaryReg = HighReg;
6580 RegPrimaryOffset = Reg0;
6581 RegSecondaryOffset = Reg32;
6586 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
6587 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
6592 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
6593 SPIRV::OpINotEqual))
6600 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
6601 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
6606 if (SwapPrimarySide) {
6608 if (!selectOpWithSrcs(RegAdd, ResType,
I,
6609 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
6620 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
6621 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
6626 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
6627 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
6630 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
6634bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
6635 SPIRVTypeInst ResType,
6637 bool IsSigned)
const {
6639 Register OpReg =
I.getOperand(2).getReg();
6642 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
6643 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
6647 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6649 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6651 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6654 return diagnoseUnsupported(
6656 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
6660bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
6661 SPIRVTypeInst ResType,
6662 MachineInstr &
I)
const {
6664 Register OpReg =
I.getOperand(2).getReg();
6669 unsigned ExtendOpcode = SPIRV::OpUConvert;
6670 unsigned BitSetOpcode = GL::FindILsb;
6674 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6676 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6678 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6681 return diagnoseUnsupported(
I,
6682 "spv_firstbitlow only supports 16,32,64 bits.");
6686bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
6687 SPIRVTypeInst ResType,
6688 MachineInstr &
I)
const {
6692 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
6695 .
addUse(
I.getOperand(2).getReg())
6698 unsigned Alignment =
I.getOperand(3).getImm();
6712 while (!Worklist.
empty()) {
6714 switch (
T->getOpcode()) {
6715 case SPIRV::OpTypeInt:
6716 case SPIRV::OpTypeFloat:
6717 case SPIRV::OpTypePointer:
6719 case SPIRV::OpTypeVector:
6720 case SPIRV::OpTypeMatrix:
6721 case SPIRV::OpTypeArray: {
6722 Register OperandReg =
T->getOperand(1).getReg();
6726 case SPIRV::OpTypeStruct:
6727 for (
unsigned Idx = 1,
E =
T->getNumOperands(); Idx <
E; ++Idx) {
6728 Register OperandReg =
T->getOperand(Idx).getReg();
6740bool SPIRVInstructionSelector::selectAbort(MachineInstr &
I)
const {
6741 assert(
I.getNumExplicitOperands() == 2);
6743 Register MsgReg =
I.getOperand(1).getReg();
6745 assert(MsgType &&
"Message argument of llvm.spv.abort has no SPIR-V type");
6748 return diagnoseUnsupported(
6750 "llvm.spv.abort message type must be a concrete SPIR-V type (numerical "
6751 "scalar, pointer, vector, matrix, or aggregate of such types)");
6754 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
6761bool SPIRVInstructionSelector::selectTrap(MachineInstr &
I)
const {
6770 uint32_t MsgVal = ~0
u;
6771 if (
I.getOpcode() == TargetOpcode::G_UBSANTRAP)
6772 MsgVal =
static_cast<uint32_t
>(
I.getOperand(0).
getImm());
6775 Register MsgReg = buildI32ConstantInEntryBlock(MsgVal,
I, MsgType);
6778 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
6785bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
6786 SPIRVTypeInst ResType,
6787 MachineInstr &
I)
const {
6791 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
6794 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
6797 unsigned Alignment =
I.getOperand(2).getImm();
6804bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
6809 const MachineInstr *PrevI =
I.getPrevNode();
6811 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
6815 .
addMBB(
I.getOperand(0).getMBB())
6820 .
addMBB(
I.getOperand(0).getMBB())
6825bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
6836 const MachineInstr *NextI =
I.getNextNode();
6838 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
6844 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
6846 .
addUse(
I.getOperand(0).getReg())
6847 .
addMBB(
I.getOperand(1).getMBB())
6853bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
6854 MachineInstr &
I)
const {
6856 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
6858 const unsigned NumOps =
I.getNumOperands();
6859 for (
unsigned i = 1; i <
NumOps; i += 2) {
6860 MIB.
addUse(
I.getOperand(i + 0).getReg());
6861 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
6867bool SPIRVInstructionSelector::selectGlobalValue(
6868 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
6870 MachineIRBuilder MIRBuilder(
I);
6871 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
6874 std::string GlobalIdent;
6876 unsigned &
ID = UnnamedGlobalIDs[GV];
6878 ID = UnnamedGlobalIDs.
size();
6879 GlobalIdent =
"__unnamed_" + Twine(
ID).str();
6905 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
6912 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
6917 MachineInstrBuilder MIB1 =
6918 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
6921 MachineInstrBuilder MIB2 =
6923 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
6927 GR.
add(ConstVal, MIB2);
6935 MachineInstrBuilder MIB3 =
6936 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
6939 GR.
add(ConstVal, MIB3);
6945 assert(NewReg != ResVReg);
6946 return BuildCOPY(ResVReg, NewReg,
I);
6956 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
6959 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
6965 SPIRVTypeInst ResType =
6969 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
6974 if (
GlobalVar->isExternallyInitialized() &&
6975 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
6976 constexpr unsigned ReadWriteINTEL = 3u;
6979 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
6985bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
6986 SPIRVTypeInst ResType,
6987 MachineInstr &
I)
const {
6989 return selectExtInst(ResVReg, ResType,
I, CL::log10);
6997 MachineIRBuilder MIRBuilder(
I);
7002 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7005 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
7007 .
add(
I.getOperand(1))
7012 ResType->
getOpcode() == SPIRV::OpTypeFloat);
7022 APFloat::rmNearestTiesToEven, &LosesInfo);
7026 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
7027 ? SPIRV::OpVectorTimesScalar
7038bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
7039 SPIRVTypeInst ResType,
7040 MachineInstr &
I)
const {
7043 return selectExtInst(ResVReg, ResType,
I, CL::pown);
7049 Register ExpReg =
I.getOperand(2).getReg();
7051 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
7052 SPIRV::OpConvertSToF))
7054 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
7061bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
7062 SPIRVTypeInst ResType,
7063 MachineInstr &
I)
const {
7079 MachineIRBuilder MIRBuilder(
I);
7080 SPIRVTypeInst FloatType =
7084 FloatType, MIRBuilder, SPIRV::StorageClass::Function);
7097 MachineBasicBlock &EntryBB =
I.getMF()->
front();
7099 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
7102 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7108 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7111 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
7114 .
add(
I.getOperand(
I.getNumExplicitDefs()))
7118 Register IntegralPartReg =
I.getOperand(1).getReg();
7121 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7131 assert(
false &&
"GLSL::Modf is deprecated.");
7142bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
7143 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7144 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7145 MachineIRBuilder MIRBuilder(
I);
7146 const SPIRVTypeInst Vec3Ty =
7149 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
7161 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7165 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7171 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7178 assert(
I.getOperand(2).isReg());
7179 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
7183 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
7194bool SPIRVInstructionSelector::loadBuiltinInputID(
7195 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7196 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7197 MachineIRBuilder MIRBuilder(
I);
7199 ResType, MIRBuilder, SPIRV::StorageClass::Input);
7214 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7218 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7227SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
7228 MachineInstr &
I)
const {
7229 MachineIRBuilder MIRBuilder(
I);
7230 if (
Type->getOpcode() != SPIRV::OpTypeVector)
7240bool SPIRVInstructionSelector::loadHandleBeforePosition(
7241 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
7242 MachineInstr &Pos)
const {
7245 Intrinsic::spv_resource_handlefrombinding);
7253 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
7254 MachineIRBuilder MIRBuilder(HandleDef);
7255 SPIRVTypeInst VarType = ResType;
7256 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
7258 if (IsStructuredBuffer) {
7263 if (ResType->
getOpcode() == SPIRV::OpTypeImage && ArraySize == 0)
7265 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
7268 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
7269 ArraySize, IndexReg, Name, MIRBuilder);
7273 uint32_t LoadOpcode =
7274 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
7284bool SPIRVInstructionSelector::errorIfInstrOutsideShader(
7285 MachineInstr &
I)
const {
7287 return diagnoseUnsupported(
7288 I,
"this instruction is only supported in shaders.");
7293InstructionSelector *
7297 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
bool use_nodbg_empty(Register RegNo) const
use_nodbg_empty - Return true if there are no non-Debug instructions using the specified register.
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.
MachineInstr * getDef(const MachineOperand &MO, const MachineRegisterInfo *MRI)
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)
void buildOpDecorate(Register Reg, MachineIRBuilder &MIRBuilder, SPIRV::Decoration::Decoration Dec, ArrayRef< uint32_t > DecArgs, StringRef StrImm)
LLVM_ABI bool isNullOrNullSplat(const MachineInstr &MI, const MachineRegisterInfo &MRI, bool AllowUndefs=false)
Return true if the value is a constant 0 integer or a splatted vector of a constant 0 integer (with n...
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.
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...