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::selectInterlockedOp(
Register ResVReg,
2449 SPIRVTypeInst ResType,
2451 unsigned Opcode)
const {
2452 Register Ptr =
I.getOperand(2).getReg();
2456 assert((SC == SPIRV::StorageClass::Workgroup ||
2457 SC == SPIRV::StorageClass::StorageBuffer) &&
2458 "InterlockedAdd requires Workgroup or StorageBuffer storage class");
2459 uint32_t
Scope =
static_cast<uint32_t
>(SC == SPIRV::StorageClass::Workgroup
2460 ? SPIRV::Scope::Workgroup
2461 : SPIRV::Scope::Device);
2462 Register ScopeReg = buildI32Constant(Scope,
I);
2465 Register MemSemReg = buildI32Constant(MemSem,
I);
2467 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
2478bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2479 unsigned ArgI =
I.getNumOperands() - 1;
2481 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2482 SPIRVTypeInst SrcType =
2484 if (!SrcType || SrcType->
getOpcode() != SPIRV::OpTypeVector)
2486 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2490 unsigned CurrentIndex = 0;
2491 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2492 Register ResVReg =
I.getOperand(i).getReg();
2495 LLT ResLLT = MRI->
getType(ResVReg);
2501 ResType = ScalarType;
2507 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
2510 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2516 for (
unsigned j = 0;
j < NumElements; ++
j) {
2517 MIB.
addImm(CurrentIndex + j);
2519 CurrentIndex += NumElements;
2523 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2535bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2538 Register MemSemReg = buildI32Constant(MemSem,
I);
2540 uint32_t
Scope =
static_cast<uint32_t
>(
2542 Register ScopeReg = buildI32Constant(Scope,
I);
2544 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2551bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2552 SPIRVTypeInst ResType,
2554 unsigned Opcode)
const {
2555 Type *ResTy =
nullptr;
2558 return diagnoseUnsupported(
2560 "Not enough info to select the arithmetic with overflow instruction");
2562 return diagnoseUnsupported(
I,
2563 "Expect struct type result for the arithmetic "
2564 "with overflow instruction");
2570 MachineIRBuilder MIRBuilder(
I);
2572 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2573 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2579 Register ZeroReg = buildZerosVal(ResType,
I);
2584 if (ResName.
size() > 0)
2589 BuildMI(BB, MIRBuilder.getInsertPt(),
I.getDebugLoc(),
TII.get(Opcode))
2592 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2593 MIB.
addUse(
I.getOperand(i).getReg());
2598 MRI->
setRegClass(HigherVReg, &SPIRV::iIDRegClass);
2599 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2601 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2602 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2609 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2610 .
addDef(
I.getOperand(1).getReg())
2618bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2619 SPIRVTypeInst ResType,
2620 MachineInstr &
I)
const {
2622 "selectAtomicCmpXchg only handles the spv_cmpxchg intrinsic");
2623 Register Ptr =
I.getOperand(2).getReg();
2624 Register ScopeReg =
I.getOperand(5).getReg();
2625 Register MemSemEqReg =
I.getOperand(6).getReg();
2626 Register MemSemNeqReg =
I.getOperand(7).getReg();
2628 Register Val =
I.getOperand(4).getReg();
2632 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2651 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2658 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2670 case SPIRV::StorageClass::DeviceOnlyINTEL:
2671 case SPIRV::StorageClass::HostOnlyINTEL:
2680 bool IsGRef =
false;
2681 bool IsAllowedRefs =
2683 unsigned Opcode = It.getOpcode();
2684 if (Opcode == SPIRV::OpConstantComposite ||
2685 Opcode == SPIRV::OpSpecConstantComposite ||
2686 Opcode == SPIRV::OpVariable ||
2687 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2688 return IsGRef = true;
2689 return Opcode == SPIRV::OpName;
2691 return IsAllowedRefs && IsGRef;
2694Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2695 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2697 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC));
2701SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2703 uint32_t Opcode)
const {
2704 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2705 TII.get(SPIRV::OpSpecConstantOp))
2713SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2714 SPIRVTypeInst SrcPtrTy)
const {
2715 SPIRVTypeInst GenericPtrTy =
2719 SPIRV::StorageClass::Generic),
2721 MachineFunction *MF =
I.getParent()->getParent();
2723 MachineInstrBuilder MIB = buildSpecConstantOp(
2725 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2735bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2736 SPIRVTypeInst ResType,
2737 MachineInstr &
I)
const {
2741 Register SrcPtr =
I.getOperand(1).getReg();
2745 if (SrcPtrTy->
getOpcode() != SPIRV::OpTypePointer ||
2746 ResType->
getOpcode() != SPIRV::OpTypePointer)
2747 return BuildCOPY(ResVReg, SrcPtr,
I);
2757 unsigned SpecOpcode =
2759 ?
static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric)
2762 ? static_cast<uint32_t>(
SPIRV::Opcode::GenericCastToPtr)
2769 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
2771 .constrainAllUses(
TII,
TRI, RBI);
2773 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
2775 buildSpecConstantOp(
2777 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
2778 .constrainAllUses(
TII,
TRI, RBI);
2785 return BuildCOPY(ResVReg, SrcPtr,
I);
2787 if ((SrcSC == SPIRV::StorageClass::Function &&
2788 DstSC == SPIRV::StorageClass::Private) ||
2789 (DstSC == SPIRV::StorageClass::Function &&
2790 SrcSC == SPIRV::StorageClass::Private))
2791 return BuildCOPY(ResVReg, SrcPtr,
I);
2795 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2798 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2801 SPIRVTypeInst GenericPtrTy =
2820 return selectUnOp(ResVReg, ResType,
I,
2821 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
2823 return selectUnOp(ResVReg, ResType,
I,
2824 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
2826 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
2828 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
2838bool SPIRVInstructionSelector::selectPtrMask(
Register ResVReg,
2839 SPIRVTypeInst ResType,
2840 MachineInstr &
I)
const {
2842 return diagnoseUnsupported(
2843 I,
"G_PTRMASK is not supported with logical SPIR-V");
2848 Register PtrReg =
I.getOperand(1).getReg();
2849 Register MaskReg =
I.getOperand(2).getReg();
2868 ? SPIRV::OpBitwiseAndV
2869 : SPIRV::OpBitwiseAndS;
2892 return SPIRV::OpFOrdEqual;
2894 return SPIRV::OpFOrdGreaterThanEqual;
2896 return SPIRV::OpFOrdGreaterThan;
2898 return SPIRV::OpFOrdLessThanEqual;
2900 return SPIRV::OpFOrdLessThan;
2902 return SPIRV::OpFOrdNotEqual;
2904 return SPIRV::OpOrdered;
2906 return SPIRV::OpFUnordEqual;
2908 return SPIRV::OpFUnordGreaterThanEqual;
2910 return SPIRV::OpFUnordGreaterThan;
2912 return SPIRV::OpFUnordLessThanEqual;
2914 return SPIRV::OpFUnordLessThan;
2916 return SPIRV::OpFUnordNotEqual;
2918 return SPIRV::OpUnordered;
2928 return SPIRV::OpIEqual;
2930 return SPIRV::OpINotEqual;
2932 return SPIRV::OpSGreaterThanEqual;
2934 return SPIRV::OpSGreaterThan;
2936 return SPIRV::OpSLessThanEqual;
2938 return SPIRV::OpSLessThan;
2940 return SPIRV::OpUGreaterThanEqual;
2942 return SPIRV::OpUGreaterThan;
2944 return SPIRV::OpULessThanEqual;
2946 return SPIRV::OpULessThan;
2955 return SPIRV::OpPtrEqual;
2957 return SPIRV::OpPtrNotEqual;
2968 return SPIRV::OpLogicalEqual;
2970 return SPIRV::OpLogicalNotEqual;
3004bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
3005 SPIRVTypeInst ResType,
3007 unsigned OpAnyOrAll)
const {
3008 assert(
I.getNumOperands() == 3);
3009 assert(
I.getOperand(2).isReg());
3011 Register InputRegister =
I.getOperand(2).getReg();
3014 assert(InputType &&
"VReg has no type assigned");
3017 bool IsVectorTy = InputType->
getOpcode() == SPIRV::OpTypeVector;
3018 if (IsBoolTy && !IsVectorTy) {
3019 assert(ResVReg ==
I.getOperand(0).getReg());
3020 return BuildCOPY(ResVReg, InputRegister,
I);
3024 unsigned SpirvNotEqualId =
3025 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
3027 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
3032 IsBoolTy ? InputRegister
3040 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
3042 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
3059bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
3060 SPIRVTypeInst ResType,
3061 MachineInstr &
I)
const {
3062 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
3065bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
3066 SPIRVTypeInst ResType,
3067 MachineInstr &
I)
const {
3068 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
3072bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
3073 SPIRVTypeInst ResType,
3074 MachineInstr &
I)
const {
3075 assert(
I.getNumOperands() == 4);
3076 assert(
I.getOperand(2).isReg());
3077 assert(
I.getOperand(3).isReg());
3079 [[maybe_unused]] SPIRVTypeInst VecType =
3084 "dot product requires a vector of at least 2 components");
3086 [[maybe_unused]] SPIRVTypeInst EltType =
3095 .
addUse(
I.getOperand(2).getReg())
3096 .
addUse(
I.getOperand(3).getReg())
3101bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
3102 SPIRVTypeInst ResType,
3105 assert(
I.getNumOperands() == 4);
3106 assert(
I.getOperand(2).isReg());
3107 assert(
I.getOperand(3).isReg());
3110 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3114 .
addUse(
I.getOperand(2).getReg())
3115 .
addUse(
I.getOperand(3).getReg())
3122bool SPIRVInstructionSelector::selectIntegerDotExpansion(
3123 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3124 assert(
I.getNumOperands() == 4);
3125 assert(
I.getOperand(2).isReg());
3126 assert(
I.getOperand(3).isReg());
3130 Register Vec0 =
I.getOperand(2).getReg();
3131 Register Vec1 =
I.getOperand(3).getReg();
3135 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
3144 "dot product requires a vector of at least 2 components");
3147 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3157 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3168 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3180bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
3181 SPIRVTypeInst ResType,
3182 MachineInstr &
I)
const {
3184 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
3187 .
addUse(
I.getOperand(2).getReg())
3192bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
3193 SPIRVTypeInst ResType,
3194 MachineInstr &
I)
const {
3196 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
3199 .
addUse(
I.getOperand(2).getReg())
3204bool SPIRVInstructionSelector::selectOpIsFinite(
Register ResVReg,
3205 SPIRVTypeInst ResType,
3206 MachineInstr &
I)
const {
3208 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsFinite))
3211 .
addUse(
I.getOperand(2).getReg())
3216bool SPIRVInstructionSelector::selectOpIsNormal(
Register ResVReg,
3217 SPIRVTypeInst ResType,
3218 MachineInstr &
I)
const {
3220 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNormal))
3223 .
addUse(
I.getOperand(2).getReg())
3228template <
bool Signed>
3229bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
3230 SPIRVTypeInst ResType,
3231 MachineInstr &
I)
const {
3232 assert(
I.getNumOperands() == 5);
3233 assert(
I.getOperand(2).isReg());
3234 assert(
I.getOperand(3).isReg());
3235 assert(
I.getOperand(4).isReg());
3238 Register Acc =
I.getOperand(2).getReg();
3242 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3244 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
3249 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
3252 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3264template <
bool Signed>
3265bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
3266 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3267 assert(
I.getNumOperands() == 5);
3268 assert(
I.getOperand(2).isReg());
3269 assert(
I.getOperand(3).isReg());
3270 assert(
I.getOperand(4).isReg());
3273 Register Acc =
I.getOperand(2).getReg();
3279 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
3283 for (
unsigned i = 0; i < 4; i++) {
3306 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
3326 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3341bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
3342 SPIRVTypeInst ResType,
3343 MachineInstr &
I)
const {
3344 assert(
I.getNumOperands() == 3);
3345 assert(
I.getOperand(2).isReg());
3347 Register VZero = buildZerosValF(ResType,
I);
3348 Register VOne = buildOnesValF(ResType,
I);
3350 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
3353 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3355 .
addUse(
I.getOperand(2).getReg())
3362bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
3363 SPIRVTypeInst ResType,
3364 MachineInstr &
I)
const {
3365 assert(
I.getNumOperands() == 3);
3366 assert(
I.getOperand(2).isReg());
3368 Register InputRegister =
I.getOperand(2).getReg();
3370 auto &
DL =
I.getDebugLoc();
3373 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3380 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
3382 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
3390 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3395 if (NeedsConversion) {
3396 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
3407bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
3408 SPIRVTypeInst ResType,
3410 unsigned Opcode)
const {
3414 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
3420 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
3421 BMI.addUse(
I.getOperand(J).getReg());
3428bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
3431 bool WithGroupSync)
const {
3433 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
3435 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
3437 assert(((Scope != SPIRV::Scope::Workgroup) ||
3438 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
3439 "Workgroup Scope must set WorkGroupMemory semantic "
3440 "in Barrier instruction");
3442 assert(((Scope != SPIRV::Scope::Device) ||
3443 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
3444 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
3445 "Device Scope must set UniformMemory and ImageMemory semantic "
3446 "in Barrier instruction");
3452 if (WithGroupSync) {
3453 Register ExecReg = buildI32Constant(SPIRV::Scope::Workgroup,
I);
3457 Register ScopeReg = buildI32Constant(Scope,
I);
3458 Register MemSemReg = buildI32Constant(MemSem,
I);
3460 MI.addUse(ScopeReg).addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
3464bool SPIRVInstructionSelector::selectWaveActiveCountBits(
3465 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3470 if (!selectWaveOpInst(BallotReg, BallotType,
I,
3471 SPIRV::OpGroupNonUniformBallot))
3476 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3481 .
addImm(SPIRV::GroupOperation::Reduce)
3488bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
3489 SPIRVTypeInst ResType,
3490 MachineInstr &
I)
const {
3495 Register InputReg =
I.getOperand(2).getReg();
3500 bool IsVector = NumElems > 1;
3513 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
3514 SPIRV::OpGroupNonUniformAllEqual);
3519 ElementResults.
reserve(NumElems);
3521 for (
unsigned Idx = 0; Idx < NumElems; ++Idx) {
3534 ElemInput = Extracted;
3540 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3551 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3562bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3563 SPIRVTypeInst ResType,
3564 MachineInstr &
I)
const {
3566 assert(
I.getNumOperands() == 3);
3568 auto Op =
I.getOperand(2);
3578 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3580 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3581 return diagnoseUnsupported(
I,
"WavePrefixBitCount requires boolean input");
3602 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3606 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3613bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3614 SPIRVTypeInst ResType,
3616 bool IsUnsigned)
const {
3617 return selectWaveReduce(
3618 ResVReg, ResType,
I, IsUnsigned,
3619 [&](
Register InputRegister,
bool IsUnsigned) {
3620 const bool IsFloatTy =
3622 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3623 : SPIRV::OpGroupNonUniformSMax;
3624 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3628bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3629 SPIRVTypeInst ResType,
3631 bool IsUnsigned)
const {
3632 return selectWaveReduce(
3633 ResVReg, ResType,
I, IsUnsigned,
3634 [&](
Register InputRegister,
bool IsUnsigned) {
3635 const bool IsFloatTy =
3637 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3638 : SPIRV::OpGroupNonUniformSMin;
3639 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3643bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3644 SPIRVTypeInst ResType,
3645 MachineInstr &
I)
const {
3646 return selectWaveReduce(ResVReg, ResType,
I,
false,
3647 [&](
Register InputRegister,
bool IsUnsigned) {
3649 InputRegister, SPIRV::OpTypeFloat);
3650 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3651 : SPIRV::OpGroupNonUniformIAdd;
3655bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3656 SPIRVTypeInst ResType,
3657 MachineInstr &
I)
const {
3658 return selectWaveReduce(ResVReg, ResType,
I,
false,
3659 [&](
Register InputRegister,
bool IsUnsigned) {
3661 InputRegister, SPIRV::OpTypeFloat);
3662 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3663 : SPIRV::OpGroupNonUniformIMul;
3667template <
typename PickOpcodeFn>
3668bool SPIRVInstructionSelector::selectWaveReduce(
3669 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3670 PickOpcodeFn &&PickOpcode)
const {
3671 assert(
I.getNumOperands() == 3);
3672 assert(
I.getOperand(2).isReg());
3674 Register InputRegister =
I.getOperand(2).getReg();
3678 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3681 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3687 .
addImm(SPIRV::GroupOperation::Reduce)
3688 .
addUse(
I.getOperand(2).getReg())
3693bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3694 SPIRVTypeInst ResType,
3696 unsigned Opcode)
const {
3697 return selectWaveReduce(
3698 ResVReg, ResType,
I,
false,
3699 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3702bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3703 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3704 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3705 [&](
Register InputRegister,
bool IsUnsigned) {
3707 InputRegister, SPIRV::OpTypeFloat);
3709 ? SPIRV::OpGroupNonUniformFAdd
3710 : SPIRV::OpGroupNonUniformIAdd;
3714bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
3715 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3716 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3717 [&](
Register InputRegister,
bool IsUnsigned) {
3719 InputRegister, SPIRV::OpTypeFloat);
3721 ? SPIRV::OpGroupNonUniformFMul
3722 : SPIRV::OpGroupNonUniformIMul;
3726template <
typename PickOpcodeFn>
3727bool SPIRVInstructionSelector::selectWaveExclusiveScan(
3728 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3729 PickOpcodeFn &&PickOpcode)
const {
3730 assert(
I.getNumOperands() == 3);
3731 assert(
I.getOperand(2).isReg());
3733 Register InputRegister =
I.getOperand(2).getReg();
3737 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3740 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3746 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3747 .
addUse(
I.getOperand(2).getReg())
3752bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
3753 SPIRVTypeInst ResType,
3756 assert(
I.getNumOperands() == 3);
3757 assert(
I.getOperand(2).isReg());
3759 Register InputRegister =
I.getOperand(2).getReg();
3765 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
3776bool SPIRVInstructionSelector::selectBitreverseViaI32(
Register ResVReg,
3777 SPIRVTypeInst ResType,
3784 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
3789 : SPIRV::OpUConvert;
3793 ShiftOp = SPIRV::OpShiftRightLogicalV;
3798 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3799 TII.get(SPIRV::OpConstantComposite))
3802 for (
unsigned It = 0; It <
N; ++It)
3806 ShiftConst = CompositeReg;
3811 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
3816 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
3821 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
3826 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
3829bool SPIRVInstructionSelector::handle64BitOverflow(
3831 unsigned int Opcode,
3838 "handle64BitOverflow should only be used for integer types");
3840 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
3842 MachineIRBuilder MIRBuilder(
I);
3844 SPIRVTypeInst I64x2Type =
3846 SPIRVTypeInst Vec2ResType =
3849 std::vector<Register> PartialRegs;
3851 unsigned CurrentComponent = 0;
3852 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
3856 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3857 TII.get(SPIRV::OpVectorShuffle))
3862 .
addImm(CurrentComponent)
3863 .
addImm(CurrentComponent + 1);
3873 PartialRegs.push_back(SubVecReg);
3876 if (CurrentComponent != ComponentCount) {
3882 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
3883 SPIRV::OpVectorExtractDynamic))
3892 PartialRegs.push_back(FinalElemResReg);
3896 return selectOpWithSrcs(ResVReg, ResType,
I, PartialRegs,
3897 SPIRV::OpCompositeConstruct);
3900bool SPIRVInstructionSelector::selectBitreverse64(
Register ResVReg,
3901 SPIRVTypeInst ResType,
3905 if (ComponentCount > 2)
3906 return handle64BitOverflow(
3907 ResVReg, ResType,
I, SrcReg, SPIRV::OpBitReverse,
3909 unsigned O) {
return this->selectBitreverse64(R,
T,
I, S); });
3911 MachineIRBuilder MIRBuilder(
I);
3915 I32Type, 2 * ComponentCount, MIRBuilder,
false);
3919 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
3924 if (!selectBitreverseNative(Reverse32, VecI32Type,
I, Vec32))
3931 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
3932 TII.get(SPIRV::OpVectorShuffle))
3937 for (
unsigned J = 0; J < ComponentCount; ++J) {
3944 return selectOpWithSrcs(ResVReg, ResType,
I, {SwappedVec}, SPIRV::OpBitcast);
3947bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
3948 SPIRVTypeInst ResType,
3952 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
3960bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
3961 SPIRVTypeInst ResType,
3962 MachineInstr &
I)
const {
3963 Register OpReg =
I.getOperand(1).getReg();
3972 return selectBitreverseViaI32(ResVReg, ResType,
I, OpReg);
3974 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
3976 return selectBitreverse64(ResVReg, ResType,
I, OpReg);
3978 return SPIRVInstructionSelector::diagnoseUnsupported(
3979 I,
"G_BITREVERSE only support 16,32,64 bits.");
3983 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
3994 unsigned AndOp = SPIRV::OpBitwiseAndS;
3995 unsigned OrOp = SPIRV::OpBitwiseOrS;
3996 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
3997 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
3999 AndOp = SPIRV::OpBitwiseAndV;
4000 OrOp = SPIRV::OpBitwiseOrV;
4001 ShlOp = SPIRV::OpShiftLeftLogicalV;
4002 ShrOp = SPIRV::OpShiftRightLogicalV;
4008 const unsigned Shift) ->
Register {
4016 Register MaskReg = CreateConst(Mask);
4017 Register ShiftReg = CreateConst(Shift);
4024 if (!selectOpWithSrcs(
T1, ResType,
I, {Input, ShiftReg}, ShrOp) ||
4025 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
4026 !selectOpWithSrcs(T3, ResType,
I, {Input, MaskReg}, AndOp) ||
4027 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
4028 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
4036 uint64_t
Mask = ~0ull;
4037 while ((Shift >>= 1) > 0) {
4044 return BuildCOPY(ResVReg, Result,
I);
4047bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
4048 SPIRVTypeInst ResType,
4049 MachineInstr &
I)
const {
4050 assert(
I.getOperand(0).isReg() &&
I.getOperand(1).isReg() &&
4051 "G_FREEZE must define and use a register");
4052 Register OpReg =
I.getOperand(1).getReg();
4056 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
4069 if (MachineInstr *Def = MRI->
getVRegDef(OpReg)) {
4070 if (
Def->getOpcode() == TargetOpcode::COPY)
4073 switch (
Def->getOpcode()) {
4074 case SPIRV::ASSIGN_TYPE:
4075 if (MachineInstr *AssignToDef =
4077 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
4078 Reg =
Def->getOperand(2).getReg();
4081 case SPIRV::OpUndef:
4082 Reg =
Def->getOperand(1).getReg();
4085 unsigned DestOpCode;
4087 DestOpCode = SPIRV::OpConstantNull;
4088 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze of a "
4089 "static undef/poison lowered to OpConstantNull\n");
4091 DestOpCode = TargetOpcode::COPY;
4093 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze "
4094 "skipped, lowered as a copy of the operand\n");
4096 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
4097 .
addDef(
I.getOperand(0).getReg())
4105bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
4106 SPIRVTypeInst ResType,
4107 MachineInstr &
I)
const {
4109 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4111 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4115 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
4120 for (
unsigned i =
I.getNumExplicitDefs();
4121 i <
I.getNumExplicitOperands() && IsConst; ++i)
4125 if (!IsConst &&
N < 2)
4126 return diagnoseUnsupported(
4127 I,
"There must be at least two constituent operands in a vector");
4132 for (
unsigned i =
I.getNumExplicitDefs();
4133 i <
I.getNumExplicitOperands() && IsNullVector; ++i) {
4134 MachineInstr *
Def =
getDef(
I.getOperand(i), MRI);
4139 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4146 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4147 TII.get(IsConst ? SPIRV::OpConstantComposite
4148 : SPIRV::OpCompositeConstruct))
4151 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
4152 MIB.
addUse(
I.getOperand(i).getReg());
4157bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
4158 SPIRVTypeInst ResType,
4159 MachineInstr &
I)
const {
4161 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4163 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4169 if (!
I.getOperand(
OpIdx).isReg())
4176 if (!IsConst &&
N < 2)
4177 return diagnoseUnsupported(
4178 I,
"There must be at least two constituent operands in a vector");
4181 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4182 TII.get(IsConst ? SPIRV::OpConstantComposite
4183 : SPIRV::OpCompositeConstruct))
4186 for (
unsigned i = 0; i <
N; ++i)
4192bool SPIRVInstructionSelector::selectConcatVectors(
Register ResVReg,
4193 SPIRVTypeInst ResType,
4194 MachineInstr &
I)
const {
4198 if (ResType->
getOpcode() != SPIRV::OpTypeVector)
4200 "Cannot select G_CONCAT_VECTORS with a non-vector result");
4202 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4203 TII.get(SPIRV::OpCompositeConstruct))
4213bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
4214 SPIRVTypeInst ResType,
4215 MachineInstr &
I)
const {
4220 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
4222 Opcode = SPIRV::OpDemoteToHelperInvocation;
4224 Opcode = SPIRV::OpKill;
4226 if (MachineInstr *NextI =
I.getNextNode()) {
4228 NextI->eraseFromParent();
4238bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
4239 SPIRVTypeInst ResType,
unsigned CmpOpc,
4240 MachineInstr &
I)
const {
4241 Register Cmp0 =
I.getOperand(2).getReg();
4242 Register Cmp1 =
I.getOperand(3).getReg();
4245 "CMP operands should have the same type");
4246 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
4256bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
4257 SPIRVTypeInst ResType,
4258 MachineInstr &
I)
const {
4259 auto Pred =
I.getOperand(1).getPredicate();
4262 Register CmpOperand =
I.getOperand(2).getReg();
4267 Register Op1 =
I.getOperand(3).getReg();
4271 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
4276 I.getOperand(3).setReg(NewOp1);
4282 return selectCmp(ResVReg, ResType, CmpOpc,
I);
4286SPIRVInstructionSelector::buildI32Constant(uint32_t Val, MachineInstr &
I,
4287 SPIRVTypeInst ResType)
const {
4289 SPIRVTypeInst SpvI32Ty =
4292 auto ConstInt = ConstantInt::get(LLVMTy, Val);
4299 ?
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4302 :
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantI))
4305 .
addImm(APInt(32, Val).getZExtValue());
4307 GR.
add(ConstInt,
MI);
4314Register SPIRVInstructionSelector::buildI32ConstantInEntryBlock(
4315 uint32_t Val, MachineInstr &
I, SPIRVTypeInst ResType)
const {
4317 SPIRVTypeInst SpvI32Ty =
4319 auto *ConstInt = ConstantInt::get(LLVMTy, Val);
4324 MachineBasicBlock &EntryBB = *InsertIt->getParent();
4325 MachineInstr *
MI =
nullptr;
4329 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantNull))
4333 uint64_t ImmVal = APInt(32, Val).getZExtValue();
4334 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantI))
4340 GR.
add(ConstInt,
MI);
4345bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
4346 SPIRVTypeInst ResType,
4347 MachineInstr &
I)
const {
4349 return selectCmp(ResVReg, ResType, CmpOp,
I);
4352bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
4353 SPIRVTypeInst ResType,
4354 MachineInstr &
I)
const {
4356 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
4363 if (ResType->
getOpcode() != SPIRV::OpTypeVector &&
4364 ResType->
getOpcode() != SPIRV::OpTypeFloat)
4367 MachineIRBuilder MIRBuilder(
I);
4374 APFloat ConstVal(3.3219280948873623);
4378 APFloat::rmNearestTiesToEven, &LosesInfo);
4382 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
4383 ? SPIRV::OpVectorTimesScalar
4386 if (!selectOpWithSrcs(ArgReg, ResType,
I,
4387 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
4389 if (!selectExtInst(ResVReg, ResType,
I,
4390 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
4400Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
4401 MachineInstr &
I)
const {
4404 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4409bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
4415 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
4423 if (
Def->getOpcode() == SPIRV::OpConstantNull)
4426 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
4427 Def->getOpcode() == SPIRV::OpConstantI)
4440 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
4441 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
4443 Intrinsic::spv_const_composite)) {
4444 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
4445 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
4446 if (!IsZero(
Def->getOperand(i).getReg()))
4455Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
4456 MachineInstr &
I)
const {
4460 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4465Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
4466 MachineInstr &
I)
const {
4470 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4476 SPIRVTypeInst ResType,
4477 MachineInstr &
I)
const {
4481 if (ResType->
getOpcode() == SPIRV::OpTypeVector)
4486bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
4487 SPIRVTypeInst ResType,
4488 MachineInstr &
I)
const {
4489 Register SelectFirstArg =
I.getOperand(2).getReg();
4490 Register SelectSecondArg =
I.getOperand(3).getReg();
4499 SPIRV::OpTypeVector;
4506 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
4507 }
else if (IsPtrTy) {
4508 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
4510 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
4513 assert(IsScalarBool &&
"OpSelect with a scalar result requires a scalar "
4514 "boolean condition");
4516 Opcode = SPIRV::OpSelectSFSCond;
4517 }
else if (IsPtrTy) {
4518 Opcode = SPIRV::OpSelectSPSCond;
4520 Opcode = SPIRV::OpSelectSISCond;
4523 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4526 .
addUse(
I.getOperand(1).getReg())
4535bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
4536 SPIRVTypeInst ResType,
4538 MachineInstr &InsertAt,
4539 bool IsSigned)
const {
4541 Register ZeroReg = buildZerosVal(ResType, InsertAt);
4542 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
4543 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
4545 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4557bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
4558 SPIRVTypeInst ResType,
4559 MachineInstr &
I,
bool IsSigned,
4560 unsigned Opcode)
const {
4561 Register SrcReg =
I.getOperand(1).getReg();
4567 if (ResType->
getOpcode() == SPIRV::OpTypeVector) {
4572 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I,
false);
4574 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
4577bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
4578 SPIRVTypeInst ResType, MachineInstr &
I,
4579 bool IsSigned)
const {
4580 Register SrcReg =
I.getOperand(1).getReg();
4582 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
4586 if (ResType == SrcType)
4587 return BuildCOPY(ResVReg, SrcReg,
I);
4589 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4590 return selectUnOp(ResVReg, ResType,
I, Opcode);
4593bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
4594 SPIRVTypeInst ResType,
4596 bool IsSigned)
const {
4597 MachineIRBuilder MIRBuilder(
I);
4598 MachineRegisterInfo *MRI = MIRBuilder.getMRI();
4610 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
4613 .
addUse(
I.getOperand(1).getReg())
4614 .
addUse(
I.getOperand(2).getReg())
4619 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
4622 .
addUse(
I.getOperand(1).getReg())
4623 .
addUse(
I.getOperand(2).getReg())
4631 unsigned SelectOpcode =
4632 N > 1 ? SPIRV::OpSelectVIVCond : SPIRV::OpSelectSISCond;
4637 .
addUse(buildOnesVal(
true, ResType,
I))
4638 .
addUse(buildZerosVal(ResType,
I))
4645 .
addUse(buildOnesVal(
false, ResType,
I))
4650bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
4653 SPIRVTypeInst IntTy,
4654 SPIRVTypeInst BoolTy)
const {
4657 bool IsVectorTy = IntTy->
getOpcode() == SPIRV::OpTypeVector;
4658 unsigned Opcode = IsVectorTy ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
4660 Register One = buildOnesVal(
false, IntTy,
I);
4668 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
4677bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
4678 SPIRVTypeInst ResType,
4679 MachineInstr &
I)
const {
4680 Register IntReg =
I.getOperand(1).getReg();
4683 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
4684 if (ArgType == ResType)
4685 return BuildCOPY(ResVReg, IntReg,
I);
4687 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4688 return selectUnOp(ResVReg, ResType,
I, Opcode);
4691bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
4692 SPIRVTypeInst ResType,
4693 MachineInstr &
I)
const {
4694 unsigned Opcode =
I.getOpcode();
4695 unsigned TpOpcode = ResType->
getOpcode();
4697 if (TpOpcode == SPIRV::OpTypePointer || TpOpcode == SPIRV::OpTypeEvent) {
4698 assert(Opcode == TargetOpcode::G_CONSTANT &&
4699 I.getOperand(1).getCImm()->isZero());
4700 MachineBasicBlock &DepMBB =
I.getMF()->front();
4703 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
4710 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
4713bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
4714 SPIRVTypeInst ResType,
4715 MachineInstr &
I)
const {
4716 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4723bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
4724 SPIRVTypeInst ResType,
4725 MachineInstr &
I)
const {
4727 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
4731 .
addUse(
I.getOperand(3).getReg())
4733 .
addUse(
I.getOperand(2).getReg());
4734 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
4740bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
4741 SPIRVTypeInst ResType,
4742 MachineInstr &
I)
const {
4743 Type *MaybeResTy =
nullptr;
4748 "Expected aggregate type for extractv instruction");
4750 SPIRV::AccessQualifier::ReadWrite,
false);
4754 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
4757 .
addUse(
I.getOperand(2).getReg());
4758 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
4764bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
4765 SPIRVTypeInst ResType,
4766 MachineInstr &
I)
const {
4767 if (
getImm(
I.getOperand(4), MRI))
4768 return selectInsertVal(ResVReg, ResType,
I);
4770 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
4773 .
addUse(
I.getOperand(2).getReg())
4774 .
addUse(
I.getOperand(3).getReg())
4775 .
addUse(
I.getOperand(4).getReg())
4780bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
4781 SPIRVTypeInst ResType,
4782 MachineInstr &
I)
const {
4783 if (
getImm(
I.getOperand(3), MRI))
4784 return selectExtractVal(ResVReg, ResType,
I);
4786 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
4789 .
addUse(
I.getOperand(2).getReg())
4790 .
addUse(
I.getOperand(3).getReg())
4795bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
4796 SPIRVTypeInst ResType,
4797 MachineInstr &
I)
const {
4798 const bool IsGEPInBounds =
I.getOperand(2).getImm();
4804 ? (IsGEPInBounds ? SPIRV::OpInBoundsAccessChain
4805 : SPIRV::OpAccessChain)
4806 : (IsGEPInBounds ?
SPIRV::OpInBoundsPtrAccessChain
4807 :
SPIRV::OpPtrAccessChain);
4809 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4813 .
addUse(
I.getOperand(3).getReg());
4815 (Opcode == SPIRV::OpPtrAccessChain ||
4816 Opcode == SPIRV::OpInBoundsPtrAccessChain ||
4817 (
getImm(
I.getOperand(4), MRI) &&
foldImm(
I.getOperand(4), MRI) == 0)) &&
4818 "Cannot translate GEP to OpAccessChain. First index must be 0.");
4821 const unsigned StartingIndex =
4822 (Opcode == SPIRV::OpAccessChain || Opcode == SPIRV::OpInBoundsAccessChain)
4825 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
4826 Res.addUse(
I.getOperand(i).getReg());
4827 Res.constrainAllUses(
TII,
TRI, RBI);
4832bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
4834 unsigned Lim =
I.getNumExplicitOperands();
4835 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
4836 Register OpReg =
I.getOperand(i).getReg();
4837 MachineInstr *OpDefine = MRI->
getVRegDef(OpReg);
4839 if (!OpDefine || !OpType ||
isConstReg(MRI, OpDefine) ||
4840 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
4841 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
4848 MachineFunction *MF =
I.getMF();
4860 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4861 TII.get(SPIRV::OpSpecConstantOp))
4864 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
4866 GR.
add(OpDefine, MIB);
4872bool SPIRVInstructionSelector::selectDerivativeInst(
4873 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
4874 const unsigned DPdOpCode)
const {
4877 if (!errorIfInstrOutsideShader(
I))
4883 Register SrcReg =
I.getOperand(2).getReg();
4888 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4891 .
addUse(
I.getOperand(2).getReg());
4893 MachineIRBuilder MIRBuilder(
I);
4896 if (componentCount != 1)
4900 const TargetRegisterClass *RegClass = GR.
getRegClass(SrcType);
4904 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4909 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
4914 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
4922bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
4923 SPIRVTypeInst ResType,
4924 MachineInstr &
I)
const {
4928 case Intrinsic::spv_load:
4929 return selectLoad(ResVReg, ResType,
I);
4930 case Intrinsic::spv_atomic_load:
4931 return selectAtomicLoad(ResVReg, ResType,
I);
4932 case Intrinsic::spv_store:
4933 return selectStore(
I);
4934 case Intrinsic::spv_atomic_store:
4935 return selectAtomicStore(
I);
4936 case Intrinsic::spv_extractv:
4937 return selectExtractVal(ResVReg, ResType,
I);
4938 case Intrinsic::spv_insertv:
4939 return selectInsertVal(ResVReg, ResType,
I);
4940 case Intrinsic::spv_extractelt:
4941 return selectExtractElt(ResVReg, ResType,
I);
4942 case Intrinsic::spv_insertelt:
4943 return selectInsertElt(ResVReg, ResType,
I);
4944 case Intrinsic::spv_gep:
4945 return selectGEP(ResVReg, ResType,
I);
4946 case Intrinsic::spv_bitcast: {
4947 Register OpReg =
I.getOperand(2).getReg();
4948 SPIRVTypeInst OpType =
4952 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
4954 case Intrinsic::spv_unref_global:
4955 case Intrinsic::spv_init_global: {
4956 MachineInstr *
MI = MRI->
getVRegDef(
I.getOperand(1).getReg());
4961 Register GVarVReg =
MI->getOperand(0).getReg();
4962 if (!selectGlobalValue(GVarVReg, *
MI, Init))
4967 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
4969 MI->eraseFromParent();
4973 case Intrinsic::spv_undef: {
4974 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
4980 case Intrinsic::spv_poison:
4981 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpPoisonKHR))
4986 case Intrinsic::spv_freeze:
4987 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
4990 .
addUse(
I.getOperand(2).getReg())
4993 case Intrinsic::spv_named_boolean_spec_constant: {
4994 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
4995 : SPIRV::OpSpecConstantFalse;
4997 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
4998 .
addDef(
I.getOperand(0).getReg())
5001 unsigned SpecId =
I.getOperand(2).getImm();
5003 SPIRV::Decoration::SpecId, {SpecId});
5007 case Intrinsic::spv_const_composite: {
5009 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
5015 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
5017 std::function<bool(
Register)> HasSpecConstOperand =
5027 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
5028 J < Def->getNumExplicitOperands(); ++J) {
5029 if (
Def->getOperand(J).isReg() &&
5030 HasSpecConstOperand(
Def->getOperand(J).getReg()))
5036 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
5037 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
5038 : SPIRV::OpConstantComposite;
5039 unsigned ContinuedOpc = HasSpecConst
5040 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
5041 : SPIRV::OpConstantCompositeContinuedINTEL;
5042 MachineIRBuilder MIR(
I);
5044 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
5046 for (
auto *Instr : Instructions) {
5047 Instr->setDebugLoc(
I.getDebugLoc());
5052 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5059 case Intrinsic::spv_assign_name: {
5060 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
5061 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
5062 for (
unsigned i =
I.getNumExplicitDefs() + 2;
5063 i <
I.getNumExplicitOperands(); ++i) {
5064 MIB.
addImm(
I.getOperand(i).getImm());
5069 case Intrinsic::spv_switch: {
5070 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
5071 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5072 if (
I.getOperand(i).isReg())
5073 MIB.
addReg(
I.getOperand(i).getReg());
5074 else if (
I.getOperand(i).isCImm())
5075 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
5076 else if (
I.getOperand(i).isMBB())
5077 MIB.
addMBB(
I.getOperand(i).getMBB());
5084 case Intrinsic::spv_loop_merge: {
5085 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
5086 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5087 if (
I.getOperand(i).isMBB())
5088 MIB.
addMBB(
I.getOperand(i).getMBB());
5095 case Intrinsic::spv_loop_control_intel: {
5097 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
5098 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
5103 case Intrinsic::spv_selection_merge: {
5105 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
5106 assert(
I.getOperand(1).isMBB() &&
5107 "operand 1 to spv_selection_merge must be a basic block");
5108 MIB.
addMBB(
I.getOperand(1).getMBB());
5109 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
5113 case Intrinsic::spv_cmpxchg:
5114 return selectAtomicCmpXchg(ResVReg, ResType,
I);
5115 case Intrinsic::spv_unreachable:
5116 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
5119 case Intrinsic::spv_abort:
5120 return selectAbort(
I);
5121 case Intrinsic::spv_alloca:
5122 return selectFrameIndex(ResVReg, ResType,
I);
5123 case Intrinsic::spv_alloca_array:
5124 return selectAllocaArray(ResVReg, ResType,
I);
5125 case Intrinsic::spv_assume:
5127 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
5128 .
addUse(
I.getOperand(1).getReg())
5133 case Intrinsic::spv_expect:
5135 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
5138 .
addUse(
I.getOperand(2).getReg())
5139 .
addUse(
I.getOperand(3).getReg())
5144 case Intrinsic::arithmetic_fence:
5145 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
5146 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
5149 .
addUse(
I.getOperand(2).getReg())
5153 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
5155 case Intrinsic::spv_thread_id:
5161 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
5163 case Intrinsic::spv_thread_id_in_group:
5169 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
5171 case Intrinsic::spv_group_id:
5177 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
5179 case Intrinsic::spv_flattened_thread_id_in_group:
5186 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
5188 case Intrinsic::spv_workgroup_size:
5189 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
5191 case Intrinsic::spv_global_size:
5192 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
5194 case Intrinsic::spv_global_offset:
5195 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
5197 case Intrinsic::spv_num_workgroups:
5198 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
5200 case Intrinsic::spv_subgroup_size:
5201 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
5203 case Intrinsic::spv_num_subgroups:
5204 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
5206 case Intrinsic::spv_subgroup_id:
5207 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
5208 case Intrinsic::spv_subgroup_local_invocation_id:
5209 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
5210 ResVReg, ResType,
I);
5211 case Intrinsic::spv_subgroup_max_size:
5212 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
5214 case Intrinsic::spv_fdot:
5215 return selectFloatDot(ResVReg, ResType,
I);
5216 case Intrinsic::spv_udot:
5217 case Intrinsic::spv_sdot:
5218 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5220 return selectIntegerDot(ResVReg, ResType,
I,
5221 IID == Intrinsic::spv_sdot);
5222 return selectIntegerDotExpansion(ResVReg, ResType,
I);
5223 case Intrinsic::spv_dot4add_i8packed:
5224 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5226 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
5227 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
5228 case Intrinsic::spv_dot4add_u8packed:
5229 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5231 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
5232 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
5233 case Intrinsic::spv_all:
5234 return selectAll(ResVReg, ResType,
I);
5235 case Intrinsic::spv_any:
5236 return selectAny(ResVReg, ResType,
I);
5237 case Intrinsic::spv_cross:
5238 return selectExtInst(ResVReg, ResType,
I, CL::cross, GL::Cross);
5239 case Intrinsic::spv_distance:
5240 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
5241 case Intrinsic::spv_lerp:
5242 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
5243 case Intrinsic::spv_length:
5244 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
5245 case Intrinsic::spv_degrees:
5246 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
5247 case Intrinsic::spv_faceforward:
5248 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
5249 case Intrinsic::spv_frac:
5250 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
5251 case Intrinsic::spv_isinf:
5252 return selectOpIsInf(ResVReg, ResType,
I);
5253 case Intrinsic::spv_isnan:
5254 return selectOpIsNan(ResVReg, ResType,
I);
5255 case Intrinsic::spv_isfinite:
5256 return selectOpIsFinite(ResVReg, ResType,
I);
5257 case Intrinsic::spv_isnormal:
5258 return selectOpIsNormal(ResVReg, ResType,
I);
5259 case Intrinsic::spv_normalize:
5260 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
5261 case Intrinsic::spv_refract:
5262 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
5263 case Intrinsic::spv_reflect:
5264 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
5265 case Intrinsic::spv_rsqrt:
5266 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
5267 case Intrinsic::spv_sign:
5268 return selectSign(ResVReg, ResType,
I);
5269 case Intrinsic::spv_smoothstep:
5270 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
5271 case Intrinsic::spv_firstbituhigh:
5272 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
5273 case Intrinsic::spv_firstbitshigh:
5274 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
5275 case Intrinsic::spv_firstbitlow:
5276 return selectFirstBitLow(ResVReg, ResType,
I);
5277 case Intrinsic::spv_all_memory_barrier:
5278 return selectBarrierInst(
I, SPIRV::Scope::Device,
5279 SPIRV::MemorySemantics::UniformMemory |
5280 SPIRV::MemorySemantics::ImageMemory |
5281 SPIRV::MemorySemantics::WorkgroupMemory,
5283 case Intrinsic::spv_all_memory_barrier_with_group_sync:
5284 return selectBarrierInst(
I, SPIRV::Scope::Device,
5285 SPIRV::MemorySemantics::UniformMemory |
5286 SPIRV::MemorySemantics::ImageMemory |
5287 SPIRV::MemorySemantics::WorkgroupMemory,
5289 case Intrinsic::spv_device_memory_barrier:
5290 return selectBarrierInst(
I, SPIRV::Scope::Device,
5291 SPIRV::MemorySemantics::UniformMemory |
5292 SPIRV::MemorySemantics::ImageMemory,
5294 case Intrinsic::spv_device_memory_barrier_with_group_sync:
5295 return selectBarrierInst(
I, SPIRV::Scope::Device,
5296 SPIRV::MemorySemantics::UniformMemory |
5297 SPIRV::MemorySemantics::ImageMemory,
5299 case Intrinsic::spv_group_memory_barrier:
5300 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5301 SPIRV::MemorySemantics::WorkgroupMemory,
5303 case Intrinsic::spv_group_memory_barrier_with_group_sync:
5304 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5305 SPIRV::MemorySemantics::WorkgroupMemory,
5307 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
5308 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
5309 SPIRV::StorageClass::StorageClass ResSC =
5312 return diagnoseUnsupported(
I,
"The target storage class is not castable "
5313 "from the Generic storage class");
5314 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
5322 case Intrinsic::spv_lifetime_start:
5323 case Intrinsic::spv_lifetime_end: {
5324 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
5325 : SPIRV::OpLifetimeStop;
5326 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
5327 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
5336 case Intrinsic::spv_saturate:
5337 return selectSaturate(ResVReg, ResType,
I);
5338 case Intrinsic::spv_nclamp:
5339 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
5340 case Intrinsic::spv_uclamp:
5341 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
5342 case Intrinsic::spv_sclamp:
5343 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
5344 case Intrinsic::spv_subgroup_prefix_bit_count:
5345 return selectWavePrefixBitCount(ResVReg, ResType,
I);
5346 case Intrinsic::spv_wave_active_countbits:
5347 return selectWaveActiveCountBits(ResVReg, ResType,
I);
5348 case Intrinsic::spv_wave_all_equal:
5349 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
5350 case Intrinsic::spv_wave_all:
5351 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
5352 case Intrinsic::spv_wave_any:
5353 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
5354 case Intrinsic::spv_subgroup_ballot:
5355 return selectWaveOpInst(ResVReg, ResType,
I,
5356 SPIRV::OpGroupNonUniformBallot);
5357 case Intrinsic::spv_wave_is_first_lane:
5358 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
5359 case Intrinsic::spv_wave_reduce_or:
5360 return selectWaveReduceOp(ResVReg, ResType,
I,
5361 SPIRV::OpGroupNonUniformBitwiseOr);
5362 case Intrinsic::spv_wave_reduce_xor:
5363 return selectWaveReduceOp(ResVReg, ResType,
I,
5364 SPIRV::OpGroupNonUniformBitwiseXor);
5365 case Intrinsic::spv_wave_reduce_and:
5366 return selectWaveReduceOp(ResVReg, ResType,
I,
5367 SPIRV::OpGroupNonUniformBitwiseAnd);
5368 case Intrinsic::spv_interlocked_add:
5369 return selectInterlockedOp(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
5370 case Intrinsic::spv_interlocked_or:
5371 return selectInterlockedOp(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
5372 case Intrinsic::spv_wave_reduce_umax:
5373 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
5374 case Intrinsic::spv_wave_reduce_max:
5375 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
5376 case Intrinsic::spv_wave_reduce_umin:
5377 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
5378 case Intrinsic::spv_wave_reduce_min:
5379 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
5380 case Intrinsic::spv_wave_reduce_sum:
5381 return selectWaveReduceSum(ResVReg, ResType,
I);
5382 case Intrinsic::spv_wave_product:
5383 return selectWaveReduceProduct(ResVReg, ResType,
I);
5384 case Intrinsic::spv_wave_readlane:
5385 return selectWaveOpInst(ResVReg, ResType,
I,
5386 SPIRV::OpGroupNonUniformShuffle);
5387 case Intrinsic::spv_wave_prefix_sum:
5388 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
5389 case Intrinsic::spv_wave_prefix_product:
5390 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
5391 case Intrinsic::spv_quad_read_across_x: {
5392 return selectQuadSwap(ResVReg, ResType,
I, 0);
5394 case Intrinsic::spv_quad_read_across_y: {
5395 return selectQuadSwap(ResVReg, ResType,
I, 1);
5397 case Intrinsic::spv_quad_read_across_diagonal: {
5398 return selectQuadSwap(ResVReg, ResType,
I, 2);
5400 case Intrinsic::spv_step:
5401 return selectExtInst(ResVReg, ResType,
I, CL::step, GL::Step);
5402 case Intrinsic::spv_radians:
5403 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
5407 case Intrinsic::instrprof_increment:
5408 case Intrinsic::instrprof_increment_step:
5409 case Intrinsic::instrprof_value_profile:
5412 case Intrinsic::spv_value_md:
5414 case Intrinsic::spv_resource_handlefrombinding: {
5415 return selectHandleFromBinding(ResVReg, ResType,
I);
5417 case Intrinsic::spv_resource_counterhandlefrombinding:
5418 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
5419 case Intrinsic::spv_resource_updatecounter:
5420 return selectUpdateCounter(ResVReg, ResType,
I);
5421 case Intrinsic::spv_resource_store_typedbuffer: {
5422 return selectImageWriteIntrinsic(
I);
5424 case Intrinsic::spv_resource_load_typedbuffer: {
5425 return selectReadImageIntrinsic(ResVReg, ResType,
I);
5427 case Intrinsic::spv_resource_load_level: {
5428 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
5430 case Intrinsic::spv_resource_getdimensions_x:
5431 case Intrinsic::spv_resource_getdimensions_xy:
5432 case Intrinsic::spv_resource_getdimensions_xyz: {
5433 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
5435 case Intrinsic::spv_resource_getdimensions_levels_x:
5436 case Intrinsic::spv_resource_getdimensions_levels_xy:
5437 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
5438 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
5440 case Intrinsic::spv_resource_getdimensions_ms_xy:
5441 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
5442 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
5444 case Intrinsic::spv_resource_calculate_lod:
5445 case Intrinsic::spv_resource_calculate_lod_unclamped:
5446 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
5447 case Intrinsic::spv_resource_sample:
5448 case Intrinsic::spv_resource_sample_clamp:
5449 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
5450 case Intrinsic::spv_resource_samplebias:
5451 case Intrinsic::spv_resource_samplebias_clamp:
5452 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
5453 case Intrinsic::spv_resource_samplegrad:
5454 case Intrinsic::spv_resource_samplegrad_clamp:
5455 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
5456 case Intrinsic::spv_resource_samplelevel:
5457 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
5458 case Intrinsic::spv_resource_samplecmp:
5459 case Intrinsic::spv_resource_samplecmp_clamp:
5460 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
5461 case Intrinsic::spv_resource_samplecmplevelzero:
5462 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
5463 case Intrinsic::spv_resource_gather:
5464 case Intrinsic::spv_resource_gather_cmp:
5465 return selectGatherIntrinsic(ResVReg, ResType,
I);
5466 case Intrinsic::spv_resource_getbasepointer:
5467 case Intrinsic::spv_resource_getpointer: {
5468 return selectResourceGetPointer(ResVReg, ResType,
I);
5470 case Intrinsic::spv_pushconstant_getpointer: {
5471 return selectPushConstantGetPointer(ResVReg, ResType,
I);
5473 case Intrinsic::spv_discard: {
5474 return selectDiscard(ResVReg, ResType,
I);
5476 case Intrinsic::spv_resource_nonuniformindex: {
5477 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
5479 case Intrinsic::spv_unpackhalf2x16: {
5480 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
5482 case Intrinsic::spv_packhalf2x16: {
5483 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
5485 case Intrinsic::spv_ddx:
5486 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
5487 case Intrinsic::spv_ddy:
5488 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
5489 case Intrinsic::spv_ddx_coarse:
5490 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
5491 case Intrinsic::spv_ddy_coarse:
5492 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
5493 case Intrinsic::spv_ddx_fine:
5494 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
5495 case Intrinsic::spv_ddy_fine:
5496 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
5497 case Intrinsic::spv_fwidth:
5498 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
5499 case Intrinsic::spv_masked_gather:
5500 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5501 return selectMaskedGather(ResVReg, ResType,
I);
5502 return diagnoseUnsupported(
5503 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
5504 case Intrinsic::spv_masked_scatter:
5505 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5506 return selectMaskedScatter(
I);
5507 return diagnoseUnsupported(
5508 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
5509 case Intrinsic::returnaddress:
5510 case Intrinsic::frameaddress: {
5512 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5519 return diagnoseUnsupported(
I,
"intrinsic selection not implemented.");
5524bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
5525 SPIRVTypeInst ResType,
5526 MachineInstr &
I)
const {
5529 if (ResType->
getOpcode() == SPIRV::OpTypeImage)
5536bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
5537 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5539 assert(Intr.getIntrinsicID() ==
5540 Intrinsic::spv_resource_counterhandlefrombinding);
5543 Register MainHandleReg = Intr.getOperand(2).getReg();
5545 assert(MainHandleDef->getIntrinsicID() ==
5546 Intrinsic::spv_resource_handlefrombinding);
5550 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(), MRI);
5551 Register IndexReg = MainHandleDef->getOperand(5).getReg();
5552 std::string CounterName =
5557 MachineIRBuilder MIRBuilder(
I);
5559 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
5561 ArraySize, IndexReg, CounterName, MIRBuilder);
5563 return BuildCOPY(ResVReg, CounterVarReg,
I);
5566bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
5567 SPIRVTypeInst ResType,
5568 MachineInstr &
I)
const {
5570 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
5572 Register CounterHandleReg = Intr.getOperand(2).getReg();
5573 Register IncrReg = Intr.getOperand(3).getReg();
5580 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
5581 assert(CounterVarPointeeType &&
5582 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
5583 "Counter variable must be a struct");
5585 SPIRV::StorageClass::StorageBuffer &&
5586 "Counter variable must be in the storage buffer storage class");
5588 "Counter variable must have exactly 1 member in the struct");
5589 const SPIRVTypeInst MemberType =
5592 "Counter variable struct must have a single i32 member");
5596 MachineIRBuilder MIRBuilder(
I);
5598 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
5601 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
5607 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
5610 .
addUse(CounterHandleReg)
5617 Register Semantics = buildI32Constant(SPIRV::MemorySemantics::None,
I);
5620 Register Incr = buildI32Constant(
static_cast<uint32_t
>(IncrVal),
I);
5623 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
5632 return BuildCOPY(ResVReg, AtomicRes,
I);
5640 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
5648bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
5649 SPIRVTypeInst ResType,
5650 MachineInstr &
I)
const {
5658 Register ImageReg =
I.getOperand(2).getReg();
5666 Register IdxReg =
I.getOperand(3).getReg();
5668 MachineInstr &Pos =
I;
5670 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
5674bool SPIRVInstructionSelector::generateSampleImage(
5677 DebugLoc Loc, MachineInstr &Pos)
const {
5688 if (!loadHandleBeforePosition(NewSamplerReg,
5694 MachineIRBuilder MIRBuilder(Pos);
5707 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
5708 ImOps.Lod.has_value();
5709 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
5710 : SPIRV::OpImageSampleImplicitLod;
5712 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
5713 : SPIRV::OpImageSampleDrefImplicitLod;
5722 MIB.
addUse(*ImOps.Compare);
5724 uint32_t ImageOperands = 0;
5726 ImageOperands |= SPIRV::ImageOperand::Bias;
5728 ImageOperands |= SPIRV::ImageOperand::Lod;
5729 if (ImOps.GradX && ImOps.GradY)
5730 ImageOperands |= SPIRV::ImageOperand::Grad;
5731 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
5733 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
5736 "Non-constant offsets are not supported in sample instructions.");
5741 ImageOperands |= SPIRV::ImageOperand::MinLod;
5743 if (ImageOperands != 0) {
5744 MIB.
addImm(ImageOperands);
5745 if (ImageOperands & SPIRV::ImageOperand::Bias)
5747 if (ImageOperands & SPIRV::ImageOperand::Lod)
5749 if (ImageOperands & SPIRV::ImageOperand::Grad) {
5750 MIB.
addUse(*ImOps.GradX);
5751 MIB.
addUse(*ImOps.GradY);
5754 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
5755 MIB.
addUse(*ImOps.Offset);
5756 if (ImageOperands & SPIRV::ImageOperand::MinLod)
5757 MIB.
addUse(*ImOps.MinLod);
5764bool SPIRVInstructionSelector::selectImageQuerySize(
5766 std::optional<Register> LodReg)
const {
5768 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
5771 "ImageReg is not an image type.");
5773 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
5775 unsigned NumComponents = 0;
5777 case SPIRV::Dim::DIM_1D:
5778 case SPIRV::Dim::DIM_Buffer:
5779 NumComponents =
IsArray ? 2 : 1;
5781 case SPIRV::Dim::DIM_2D:
5782 case SPIRV::Dim::DIM_Cube:
5783 case SPIRV::Dim::DIM_Rect:
5784 NumComponents =
IsArray ? 3 : 2;
5786 case SPIRV::Dim::DIM_3D:
5790 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
5795 SPIRVTypeInst ResType =
5800 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5810bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
5811 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5812 Register ImageReg =
I.getOperand(2).getReg();
5819 return selectImageQuerySize(NewImageReg, ResVReg,
I);
5822bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
5823 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5824 Register ImageReg =
I.getOperand(2).getReg();
5833 Register LodReg =
I.getOperand(3).getReg();
5836 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
5838 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
5845 TII.get(SPIRV::OpImageQueryLevels))
5852 TII.get(SPIRV::OpCompositeConstruct))
5862bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
5863 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5864 Register ImageReg =
I.getOperand(2).getReg();
5875 "OpImageQuerySamples requires a multisampled image");
5877 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
5885 TII.get(SPIRV::OpImageQuerySamples))
5892 TII.get(SPIRV::OpCompositeConstruct))
5902bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
5903 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5904 Register ImageReg =
I.getOperand(2).getReg();
5905 Register SamplerReg =
I.getOperand(3).getReg();
5906 Register CoordinateReg =
I.getOperand(4).getReg();
5922 if (!loadHandleBeforePosition(
5927 MachineIRBuilder MIRBuilder(
I);
5933 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
5943 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
5950 unsigned ExtractedIndex =
5952 Intrinsic::spv_resource_calculate_lod_unclamped
5956 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5957 TII.get(SPIRV::OpCompositeExtract))
5967bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
5968 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5969 Register ImageReg =
I.getOperand(2).getReg();
5970 Register SamplerReg =
I.getOperand(3).getReg();
5971 Register CoordinateReg =
I.getOperand(4).getReg();
5972 ImageOperands ImOps;
5973 if (
I.getNumOperands() > 5)
5974 ImOps.Offset =
I.getOperand(5).getReg();
5975 if (
I.getNumOperands() > 6)
5976 ImOps.MinLod =
I.getOperand(6).getReg();
5977 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5978 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5981bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
5982 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5983 Register ImageReg =
I.getOperand(2).getReg();
5984 Register SamplerReg =
I.getOperand(3).getReg();
5985 Register CoordinateReg =
I.getOperand(4).getReg();
5986 ImageOperands ImOps;
5987 ImOps.Bias =
I.getOperand(5).getReg();
5988 if (
I.getNumOperands() > 6)
5989 ImOps.Offset =
I.getOperand(6).getReg();
5990 if (
I.getNumOperands() > 7)
5991 ImOps.MinLod =
I.getOperand(7).getReg();
5992 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
5993 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
5996bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
5997 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5998 Register ImageReg =
I.getOperand(2).getReg();
5999 Register SamplerReg =
I.getOperand(3).getReg();
6000 Register CoordinateReg =
I.getOperand(4).getReg();
6001 ImageOperands ImOps;
6002 ImOps.GradX =
I.getOperand(5).getReg();
6003 ImOps.GradY =
I.getOperand(6).getReg();
6004 if (
I.getNumOperands() > 7)
6005 ImOps.Offset =
I.getOperand(7).getReg();
6006 if (
I.getNumOperands() > 8)
6007 ImOps.MinLod =
I.getOperand(8).getReg();
6008 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6009 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6012bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
6013 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6014 Register ImageReg =
I.getOperand(2).getReg();
6015 Register SamplerReg =
I.getOperand(3).getReg();
6016 Register CoordinateReg =
I.getOperand(4).getReg();
6017 ImageOperands ImOps;
6018 ImOps.Lod =
I.getOperand(5).getReg();
6019 if (
I.getNumOperands() > 6)
6020 ImOps.Offset =
I.getOperand(6).getReg();
6021 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6022 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6025bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
6026 SPIRVTypeInst ResType,
6027 MachineInstr &
I)
const {
6028 Register ImageReg =
I.getOperand(2).getReg();
6029 Register SamplerReg =
I.getOperand(3).getReg();
6030 Register CoordinateReg =
I.getOperand(4).getReg();
6031 ImageOperands ImOps;
6032 ImOps.Compare =
I.getOperand(5).getReg();
6033 if (
I.getNumOperands() > 6)
6034 ImOps.Offset =
I.getOperand(6).getReg();
6035 if (
I.getNumOperands() > 7)
6036 ImOps.MinLod =
I.getOperand(7).getReg();
6037 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6038 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6041bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
6042 SPIRVTypeInst ResType,
6043 MachineInstr &
I)
const {
6044 Register ImageReg =
I.getOperand(2).getReg();
6045 Register CoordinateReg =
I.getOperand(3).getReg();
6046 Register LodReg =
I.getOperand(4).getReg();
6048 ImageOperands ImOps;
6050 if (
I.getNumOperands() > 5)
6051 ImOps.Offset =
I.getOperand(5).getReg();
6063 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
6064 I.getDebugLoc(),
I, &ImOps);
6067bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
6068 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6069 Register ImageReg =
I.getOperand(2).getReg();
6070 Register SamplerReg =
I.getOperand(3).getReg();
6071 Register CoordinateReg =
I.getOperand(4).getReg();
6072 ImageOperands ImOps;
6073 ImOps.Compare =
I.getOperand(5).getReg();
6074 if (
I.getNumOperands() > 6)
6075 ImOps.Offset =
I.getOperand(6).getReg();
6078 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6079 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6082bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
6083 SPIRVTypeInst ResType,
6084 MachineInstr &
I)
const {
6085 Register ImageReg =
I.getOperand(2).getReg();
6086 Register SamplerReg =
I.getOperand(3).getReg();
6087 Register CoordinateReg =
I.getOperand(4).getReg();
6090 "ImageReg is not an image type.");
6095 ComponentOrCompareReg =
I.getOperand(5).getReg();
6096 OffsetReg =
I.getOperand(6).getReg();
6099 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
6103 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6104 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
6105 Dim != SPIRV::Dim::DIM_Rect) {
6107 "Gather operations are only supported for 2D, Cube, and Rect images.");
6114 if (!loadHandleBeforePosition(
6119 MachineIRBuilder MIRBuilder(
I);
6120 SPIRVTypeInst SampledImageType =
6125 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6133 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
6135 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
6137 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6142 .
addUse(ComponentOrCompareReg);
6144 uint32_t ImageOperands = 0;
6145 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
6146 if (Dim == SPIRV::Dim::DIM_Cube) {
6148 "Gather operations with offset are not supported for Cube images.");
6152 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6154 ImageOperands |= SPIRV::ImageOperand::Offset;
6158 if (ImageOperands != 0) {
6159 MIB.
addImm(ImageOperands);
6161 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6169bool SPIRVInstructionSelector::generateImageReadOrFetch(
6172 const ImageOperands *ImOps)
const {
6175 "ImageReg is not an image type.");
6177 bool IsSignedInteger =
6182 bool IsFetch = (SampledOp.getImm() == 1);
6184 auto AddOperands = [&](MachineInstrBuilder &MIB) {
6185 uint32_t ImageOperandsMask = 0;
6186 if (IsSignedInteger)
6187 ImageOperandsMask |= 0x1000;
6189 if (IsFetch && ImOps) {
6191 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
6192 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
6194 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
6196 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
6200 if (ImageOperandsMask != 0) {
6201 MIB.
addImm(ImageOperandsMask);
6202 if (IsFetch && ImOps) {
6205 if (ImOps->Offset &&
6206 (ImageOperandsMask &
6207 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
6208 MIB.
addUse(*ImOps->Offset);
6214 if (ResultSize == 4) {
6217 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
6224 BMI.constrainAllUses(
TII,
TRI, RBI);
6228 SPIRVTypeInst ReadType = widenTypeToVec4(ResType, Pos);
6232 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
6238 BMI.constrainAllUses(
TII,
TRI, RBI);
6240 if (ResultSize == 1) {
6249 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
6252bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
6253 SPIRVTypeInst ResType,
6254 MachineInstr &
I)
const {
6255 Register ResourcePtr =
I.getOperand(2).getReg();
6257 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
6266 MachineIRBuilder MIRBuilder(
I);
6271 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6277 if (
I.getNumExplicitOperands() > 3) {
6278 Register IndexReg =
I.getOperand(3).getReg();
6285bool SPIRVInstructionSelector::selectPushConstantGetPointer(
6286 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6291bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
6292 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6293 Register ObjReg =
I.getOperand(2).getReg();
6294 if (!BuildCOPY(ResVReg, ObjReg,
I))
6304 decorateUsesAsNonUniform(ResVReg);
6308void SPIRVInstructionSelector::decorateUsesAsNonUniform(
6311 while (WorkList.
size() > 0) {
6315 bool IsDecorated =
false;
6317 if (
Use.getOpcode() == SPIRV::OpDecorate &&
6318 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
6324 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
6326 if (ResultReg == CurrentReg)
6334 SPIRV::Decoration::NonUniformEXT, {});
6339bool SPIRVInstructionSelector::extractSubvector(
6341 MachineInstr &InsertionPoint)
const {
6343 [[maybe_unused]] uint64_t InputSize =
6346 assert(InputSize > 1 &&
"The input must be a vector.");
6347 assert(ResultSize > 1 &&
"The result must be a vector.");
6348 assert(ResultSize < InputSize &&
6349 "Cannot extract more element than there are in the input.");
6352 const TargetRegisterClass *ScalarRegClass = GR.
getRegClass(ScalarType);
6353 for (uint64_t
I = 0;
I < ResultSize;
I++) {
6356 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6365 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
6367 TII.get(SPIRV::OpCompositeConstruct))
6371 for (
Register ComponentReg : ComponentRegisters)
6372 MIB.
addUse(ComponentReg);
6377bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
6378 MachineInstr &
I)
const {
6385 Register ImageReg =
I.getOperand(1).getReg();
6393 Register CoordinateReg =
I.getOperand(2).getReg();
6394 Register DataReg =
I.getOperand(3).getReg();
6397 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
6405Register SPIRVInstructionSelector::buildPointerToResource(
6406 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
6407 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
6408 StringRef Name, MachineIRBuilder MIRBuilder)
const {
6410 if (ArraySize == 1) {
6411 SPIRVTypeInst PtrType =
6414 "SpirvResType did not have an explicit layout.");
6419 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
6420 SPIRVTypeInst VarPointerType =
6423 VarPointerType, Set,
Binding, Name, MIRBuilder);
6425 SPIRVTypeInst ResPointerType =
6438bool SPIRVInstructionSelector::selectFirstBitSet16(
6439 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
6440 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
6442 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
6446 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
6449bool SPIRVInstructionSelector::selectFirstBitSet32(
6451 unsigned BitSetOpcode)
const {
6452 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6455 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6462bool SPIRVInstructionSelector::selectFirstBitSet64(
6464 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
6477 if (ComponentCount > 2) {
6478 auto Func = [
this, SwapPrimarySide](
Register ResVReg, SPIRVTypeInst ResType,
6480 unsigned Opcode) ->
bool {
6481 return this->selectFirstBitSet64(ResVReg, ResType,
I, SrcReg, Opcode,
6485 return handle64BitOverflow(ResVReg, ResType,
I, SrcReg, BitSetOpcode, Func);
6489 MachineIRBuilder MIRBuilder(
I);
6491 BaseType, 2 * ComponentCount, MIRBuilder,
false);
6495 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
6501 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
6508 bool IsScalarRes = ResType->
getOpcode() != SPIRV::OpTypeVector;
6511 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
6512 SPIRV::OpVectorExtractDynamic))
6514 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
6515 SPIRV::OpVectorExtractDynamic))
6519 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6520 TII.get(SPIRV::OpVectorShuffle))
6528 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
6534 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6535 TII.get(SPIRV::OpVectorShuffle))
6543 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
6563 SelectOp = SPIRV::OpSelectSISCond;
6564 AddOp = SPIRV::OpIAddS;
6572 SelectOp = SPIRV::OpSelectVIVCond;
6573 AddOp = SPIRV::OpIAddV;
6579 Register RegSecondaryOffset = Reg0;
6583 if (SwapPrimarySide) {
6584 PrimaryReg = LowReg;
6585 SecondaryReg = HighReg;
6586 RegPrimaryOffset = Reg0;
6587 RegSecondaryOffset = Reg32;
6592 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
6593 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
6598 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
6599 SPIRV::OpINotEqual))
6606 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
6607 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
6612 if (SwapPrimarySide) {
6614 if (!selectOpWithSrcs(RegAdd, ResType,
I,
6615 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
6626 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
6627 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
6632 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
6633 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
6636 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
6640bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
6641 SPIRVTypeInst ResType,
6643 bool IsSigned)
const {
6645 Register OpReg =
I.getOperand(2).getReg();
6648 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
6649 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
6653 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6655 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6657 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6660 return diagnoseUnsupported(
6662 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
6666bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
6667 SPIRVTypeInst ResType,
6668 MachineInstr &
I)
const {
6670 Register OpReg =
I.getOperand(2).getReg();
6675 unsigned ExtendOpcode = SPIRV::OpUConvert;
6676 unsigned BitSetOpcode = GL::FindILsb;
6680 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
6682 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
6684 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
6687 return diagnoseUnsupported(
I,
6688 "spv_firstbitlow only supports 16,32,64 bits.");
6692bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
6693 SPIRVTypeInst ResType,
6694 MachineInstr &
I)
const {
6698 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVariableLengthArrayINTEL))
6701 .
addUse(
I.getOperand(2).getReg())
6704 unsigned Alignment =
I.getOperand(3).getImm();
6718 while (!Worklist.
empty()) {
6720 switch (
T->getOpcode()) {
6721 case SPIRV::OpTypeInt:
6722 case SPIRV::OpTypeFloat:
6723 case SPIRV::OpTypePointer:
6725 case SPIRV::OpTypeVector:
6726 case SPIRV::OpTypeMatrix:
6727 case SPIRV::OpTypeArray: {
6728 Register OperandReg =
T->getOperand(1).getReg();
6732 case SPIRV::OpTypeStruct:
6733 for (
unsigned Idx = 1,
E =
T->getNumOperands(); Idx <
E; ++Idx) {
6734 Register OperandReg =
T->getOperand(Idx).getReg();
6746bool SPIRVInstructionSelector::selectAbort(MachineInstr &
I)
const {
6747 assert(
I.getNumExplicitOperands() == 2);
6749 Register MsgReg =
I.getOperand(1).getReg();
6751 assert(MsgType &&
"Message argument of llvm.spv.abort has no SPIR-V type");
6754 return diagnoseUnsupported(
6756 "llvm.spv.abort message type must be a concrete SPIR-V type (numerical "
6757 "scalar, pointer, vector, matrix, or aggregate of such types)");
6760 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
6767bool SPIRVInstructionSelector::selectTrap(MachineInstr &
I)
const {
6776 uint32_t MsgVal = ~0
u;
6777 if (
I.getOpcode() == TargetOpcode::G_UBSANTRAP)
6778 MsgVal =
static_cast<uint32_t
>(
I.getOperand(0).
getImm());
6781 Register MsgReg = buildI32ConstantInEntryBlock(MsgVal,
I, MsgType);
6784 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
6791bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
6792 SPIRVTypeInst ResType,
6793 MachineInstr &
I)
const {
6797 BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(SPIRV::OpVariable))
6800 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function))
6803 unsigned Alignment =
I.getOperand(2).getImm();
6810bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
6815 const MachineInstr *PrevI =
I.getPrevNode();
6817 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
6821 .
addMBB(
I.getOperand(0).getMBB())
6826 .
addMBB(
I.getOperand(0).getMBB())
6831bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
6842 const MachineInstr *NextI =
I.getNextNode();
6844 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
6850 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
6852 .
addUse(
I.getOperand(0).getReg())
6853 .
addMBB(
I.getOperand(1).getMBB())
6859bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
6860 MachineInstr &
I)
const {
6862 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
6864 const unsigned NumOps =
I.getNumOperands();
6865 for (
unsigned i = 1; i <
NumOps; i += 2) {
6866 MIB.
addUse(
I.getOperand(i + 0).getReg());
6867 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
6873bool SPIRVInstructionSelector::selectGlobalValue(
6874 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
6876 MachineIRBuilder MIRBuilder(
I);
6877 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
6880 std::string GlobalIdent;
6882 unsigned &
ID = UnnamedGlobalIDs[GV];
6884 ID = UnnamedGlobalIDs.
size();
6885 GlobalIdent =
"__unnamed_" + Twine(
ID).str();
6911 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
6918 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
6923 MachineInstrBuilder MIB1 =
6924 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
6927 MachineInstrBuilder MIB2 =
6929 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
6933 GR.
add(ConstVal, MIB2);
6941 MachineInstrBuilder MIB3 =
6942 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
6945 GR.
add(ConstVal, MIB3);
6951 assert(NewReg != ResVReg);
6952 return BuildCOPY(ResVReg, NewReg,
I);
6962 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
6965 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
6971 SPIRVTypeInst ResType =
6975 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
6980 if (
GlobalVar->isExternallyInitialized() &&
6981 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
6982 constexpr unsigned ReadWriteINTEL = 3u;
6985 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
6991bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
6992 SPIRVTypeInst ResType,
6993 MachineInstr &
I)
const {
6995 return selectExtInst(ResVReg, ResType,
I, CL::log10);
7003 MachineIRBuilder MIRBuilder(
I);
7008 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7011 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
7013 .
add(
I.getOperand(1))
7018 ResType->
getOpcode() == SPIRV::OpTypeFloat);
7028 APFloat::rmNearestTiesToEven, &LosesInfo);
7032 auto Opcode = ResType->
getOpcode() == SPIRV::OpTypeVector
7033 ? SPIRV::OpVectorTimesScalar
7044bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
7045 SPIRVTypeInst ResType,
7046 MachineInstr &
I)
const {
7049 return selectExtInst(ResVReg, ResType,
I, CL::pown);
7055 Register ExpReg =
I.getOperand(2).getReg();
7057 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
7058 SPIRV::OpConvertSToF))
7060 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
7067bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
7068 SPIRVTypeInst ResType,
7069 MachineInstr &
I)
const {
7085 MachineIRBuilder MIRBuilder(
I);
7086 SPIRVTypeInst FloatType =
7090 FloatType, MIRBuilder, SPIRV::StorageClass::Function);
7103 MachineBasicBlock &EntryBB =
I.getMF()->
front();
7105 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
TII.get(SPIRV::OpVariable))
7108 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7114 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7117 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
7120 .
add(
I.getOperand(
I.getNumExplicitDefs()))
7124 Register IntegralPartReg =
I.getOperand(1).getReg();
7127 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7137 assert(
false &&
"GLSL::Modf is deprecated.");
7148bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
7149 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7150 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7151 MachineIRBuilder MIRBuilder(
I);
7152 const SPIRVTypeInst Vec3Ty =
7155 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
7167 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7171 MachineRegisterInfo *MRI = MIRBuilder.
getMRI();
7177 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7184 assert(
I.getOperand(2).isReg());
7185 const uint32_t ThreadId =
foldImm(
I.getOperand(2), MRI);
7189 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
7200bool SPIRVInstructionSelector::loadBuiltinInputID(
7201 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7202 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7203 MachineIRBuilder MIRBuilder(
I);
7205 ResType, MIRBuilder, SPIRV::StorageClass::Input);
7220 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7224 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7233SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
7234 MachineInstr &
I)
const {
7235 MachineIRBuilder MIRBuilder(
I);
7236 if (
Type->getOpcode() != SPIRV::OpTypeVector)
7246bool SPIRVInstructionSelector::loadHandleBeforePosition(
7247 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
7248 MachineInstr &Pos)
const {
7251 Intrinsic::spv_resource_handlefrombinding);
7259 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
7260 MachineIRBuilder MIRBuilder(HandleDef);
7261 SPIRVTypeInst VarType = ResType;
7262 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
7264 if (IsStructuredBuffer) {
7269 if (ResType->
getOpcode() == SPIRV::OpTypeImage && ArraySize == 0)
7271 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
7274 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
7275 ArraySize, IndexReg, Name, MIRBuilder);
7279 uint32_t LoadOpcode =
7280 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
7290bool SPIRVInstructionSelector::errorIfInstrOutsideShader(
7291 MachineInstr &
I)
const {
7293 return diagnoseUnsupported(
7294 I,
"this instruction is only supported in shaders.");
7299InstructionSelector *
7303 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.
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)
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)
RelativeUniformCounterPtr ValuesPtrExpr VTableAddr Value
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...