37#include "llvm/IR/IntrinsicsSPIRV.h"
43#define DEBUG_TYPE "spirv-isel"
50 std::vector<std::pair<SPIRV::InstructionSet::InstructionSet, uint32_t>>;
55 std::optional<Register> Bias;
56 std::optional<Register>
Offset;
57 std::optional<Register> MinLod;
58 std::optional<Register> GradX;
59 std::optional<Register> GradY;
60 std::optional<Register> Lod;
61 std::optional<Register> Compare;
64llvm::SPIRV::SelectionControl::SelectionControl
65getSelectionOperandForImm(
int Imm) {
67 return SPIRV::SelectionControl::Flatten;
69 return SPIRV::SelectionControl::DontFlatten;
71 return SPIRV::SelectionControl::None;
75#define GET_GLOBALISEL_PREDICATE_BITSET
76#include "SPIRVGenGlobalISel.inc"
77#undef GET_GLOBALISEL_PREDICATE_BITSET
104#define GET_GLOBALISEL_PREDICATES_DECL
105#include "SPIRVGenGlobalISel.inc"
106#undef GET_GLOBALISEL_PREDICATES_DECL
108#define GET_GLOBALISEL_TEMPORARIES_DECL
109#include "SPIRVGenGlobalISel.inc"
110#undef GET_GLOBALISEL_TEMPORARIES_DECL
134 unsigned BitSetOpcode)
const;
138 unsigned BitSetOpcode)
const;
142 unsigned BitSetOpcode,
bool SwapPrimarySide)
const;
149 unsigned Opcode)
const;
152 unsigned Opcode)
const;
173 unsigned NewOpcode,
unsigned NegateOpcode = 0)
const;
182 Register castPtrToPtrToInt(
Register Ptr, SPIRV::StorageClass::StorageClass SC,
186 bool selectAtomicPtrValue(
202 unsigned OpType)
const;
270 unsigned Opcode)
const;
274 unsigned Opcode)
const;
278 unsigned Opcode)
const;
282 unsigned Opcode)
const;
284 template <
bool Signed>
287 template <
bool Signed>
294 template <
typename PickOpcodeFn>
297 PickOpcodeFn &&PickOpcode)
const;
314 template <
typename PickOpcodeFn>
317 PickOpcodeFn &&PickOpcode)
const;
339 bool IsSigned)
const;
341 bool IsSigned,
unsigned Opcode)
const;
343 bool IsSigned)
const;
349 bool IsSigned)
const;
390 GL::GLSLExtInst GLInst,
bool setMIFlags =
true,
391 bool useMISrc =
true,
393 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
394 CL::OpenCLExtInst CLInst,
bool setMIFlags =
true,
395 bool useMISrc =
true,
397 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
398 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
399 bool setMIFlags =
true,
bool useMISrc =
true,
401 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
402 const ExtInstList &ExtInsts,
bool setMIFlags =
true,
403 bool useMISrc =
true,
406 bool selectLog10(
Register ResVReg, SPIRVTypeInst ResType,
407 MachineInstr &
I)
const;
409 bool selectFpowi(
Register ResVReg, SPIRVTypeInst ResType,
410 MachineInstr &
I)
const;
412 bool selectSaturate(
Register ResVReg, SPIRVTypeInst ResType,
413 MachineInstr &
I)
const;
415 bool selectWaveOpInst(
Register ResVReg, SPIRVTypeInst ResType,
416 MachineInstr &
I,
unsigned Opcode)
const;
418 bool selectBarrierInst(MachineInstr &
I,
unsigned Scope,
unsigned MemSem,
419 bool WithGroupSync)
const;
421 bool selectWaveActiveCountBits(
Register ResVReg, SPIRVTypeInst ResType,
422 MachineInstr &
I)
const;
424 bool selectWaveActiveAllEqual(
Register ResVReg, SPIRVTypeInst ResType,
425 MachineInstr &
I)
const;
429 bool selectHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
430 MachineInstr &
I)
const;
432 bool selectCounterHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
433 MachineInstr &
I)
const;
435 bool selectReadImageIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
436 MachineInstr &
I)
const;
437 bool selectGetDimensionsIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
438 MachineInstr &
I)
const;
439 bool selectGetDimensionsLevelsIntrinsic(
Register &ResVReg,
440 SPIRVTypeInst ResType,
441 MachineInstr &
I)
const;
442 bool selectGetDimensionsMSIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
443 MachineInstr &
I)
const;
446 std::optional<Register> LodReg = std::nullopt)
const;
447 bool selectSampleBasicIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
448 MachineInstr &
I)
const;
449 bool selectCalculateLodIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
450 MachineInstr &
I)
const;
451 bool selectSampleBiasIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
452 MachineInstr &
I)
const;
453 bool selectSampleGradIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
454 MachineInstr &
I)
const;
455 bool selectSampleLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
456 MachineInstr &
I)
const;
457 bool selectLoadLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
458 MachineInstr &
I)
const;
459 bool selectSampleCmpIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
460 MachineInstr &
I)
const;
461 bool selectSampleCmpLevelZeroIntrinsic(
Register &ResVReg,
462 SPIRVTypeInst ResType,
463 MachineInstr &
I)
const;
464 bool selectGatherIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
465 MachineInstr &
I)
const;
466 bool selectImageWriteIntrinsic(MachineInstr &
I)
const;
467 bool selectResourceGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
468 MachineInstr &
I)
const;
469 bool selectPushConstantGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
470 MachineInstr &
I)
const;
471 bool selectResourceNonUniformIndex(
Register &ResVReg, SPIRVTypeInst ResType,
472 MachineInstr &
I)
const;
473 bool selectModf(
Register ResVReg, SPIRVTypeInst ResType,
474 MachineInstr &
I)
const;
475 bool selectUpdateCounter(
Register &ResVReg, SPIRVTypeInst ResType,
476 MachineInstr &
I)
const;
477 bool selectFrexp(
Register ResVReg, SPIRVTypeInst ResType,
478 MachineInstr &
I)
const;
480 bool selectCopySign(
Register ResVReg, SPIRVTypeInst ResType,
481 MachineInstr &
I)
const;
483 bool selectLdexp(
Register ResVReg, SPIRVTypeInst ResType,
484 MachineInstr &
I)
const;
485 bool selectSincos(
Register ResVReg, SPIRVTypeInst ResType,
486 MachineInstr &
I)
const;
487 bool selectExp10(
Register ResVReg, SPIRVTypeInst ResType,
488 MachineInstr &
I)
const;
489 bool selectDerivativeInst(
Register ResVReg, SPIRVTypeInst ResType,
490 MachineInstr &
I,
const unsigned DPdOpCode)
const;
492 Register buildI32ConstantInEntryBlock(uint32_t Val, MachineInstr &
I,
493 SPIRVTypeInst ResType =
nullptr)
const;
495 Register buildZerosVal(SPIRVTypeInst ResType, MachineInstr &
I)
const;
496 bool isScalarOrVectorIntConstantZero(
Register Reg)
const;
497 Register buildZerosValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
499 MachineInstr &
I)
const;
500 Register buildOnesValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
502 MachineInstr &
I)
const;
504 bool wrapIntoSpecConstantOp(MachineInstr &
I,
507 SPIRVTypeInst getUcharPtrType(MachineInstr &
I,
508 SPIRV::StorageClass::StorageClass SC)
const;
509 Register getUcharPtrTypeReg(MachineInstr &
I,
510 SPIRV::StorageClass::StorageClass SC)
const;
511 MachineInstrBuilder buildSpecConstantOp(MachineInstr &
I,
Register Dest,
513 uint32_t Opcode)
const;
515 MachineInstrBuilder buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
516 SPIRVTypeInst SrcPtrTy)
const;
517 Register buildPointerToResource(SPIRVTypeInst ResType,
518 SPIRV::StorageClass::StorageClass SC,
519 uint32_t Set, uint32_t
Binding,
520 uint32_t ArraySize,
Register IndexReg,
522 MachineIRBuilder MIRBuilder)
const;
523 SPIRVTypeInst widenTypeToVec4(SPIRVTypeInst
Type, MachineInstr &
I)
const;
524 bool extractSubvector(
Register &ResVReg, SPIRVTypeInst ResType,
525 Register &ReadReg, MachineInstr &InsertionPoint)
const;
526 bool generateImageReadOrFetch(
Register &ResVReg, SPIRVTypeInst ResType,
529 const ImageOperands *ImOps =
nullptr)
const;
530 bool generateSampleImage(
Register ResVReg, SPIRVTypeInst ResType,
532 Register CoordinateReg,
const ImageOperands &ImOps,
535 bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
536 Register ResVReg, SPIRVTypeInst ResType,
537 MachineInstr &
I)
const;
538 bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
539 Register ResVReg, SPIRVTypeInst ResType,
540 MachineInstr &
I)
const;
541 bool loadHandleBeforePosition(
Register &HandleReg, SPIRVTypeInst ResType,
542 GIntrinsic &HandleDef, MachineInstr &Pos)
const;
543 void decorateUsesAsNonUniform(
Register &NonUniformReg)
const;
544 bool errorIfInstrOutsideShader(MachineInstr &
I)
const;
547 handle64BitOverflow(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
548 Register SrcReg,
unsigned int Opcode,
549 std::function<
bool(
Register, SPIRVTypeInst,
550 MachineInstr &,
Register,
unsigned)>
554bool sampledTypeIsSignedInteger(
const llvm::Type *HandleType) {
556 if (
TET->getTargetExtName() ==
"spirv.Image") {
559 assert(
TET->getTargetExtName() ==
"spirv.SignedImage");
560 return TET->getTypeParameter(0)->isIntegerTy();
564#define GET_GLOBALISEL_IMPL
565#include "SPIRVGenGlobalISel.inc"
566#undef GET_GLOBALISEL_IMPL
572 TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
575#include
"SPIRVGenGlobalISel.inc"
578#include
"SPIRVGenGlobalISel.inc"
590 InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
595 if (HasVRegsReset == &MF)
610 for (
const auto &
MBB : MF) {
611 for (
const auto &
MI :
MBB) {
614 if (
MI.getOpcode() != SPIRV::ASSIGN_TYPE)
621 if (DstType != SrcType)
626 if (DstRC != SrcRC && SrcRC)
638 while (!Stack.empty()) {
643 switch (
MI->getOpcode()) {
644 case TargetOpcode::G_INTRINSIC:
645 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
646 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS: {
649 if (IntrID != Intrinsic::spv_const_composite &&
650 IntrID != Intrinsic::spv_undef && IntrID != Intrinsic::spv_poison)
654 case TargetOpcode::G_BUILD_VECTOR:
655 case TargetOpcode::G_SPLAT_VECTOR:
657 i < OpDef->getNumOperands(); i++) {
662 Stack.push_back(OpNestedDef);
665 case TargetOpcode::G_CONSTANT:
666 case TargetOpcode::G_FCONSTANT:
667 case TargetOpcode::G_IMPLICIT_DEF:
668 case SPIRV::OpConstantTrue:
669 case SPIRV::OpConstantFalse:
670 case SPIRV::OpConstantI:
671 case SPIRV::OpConstantF:
672 case SPIRV::OpConstantComposite:
673 case SPIRV::OpConstantCompositeContinuedINTEL:
674 case SPIRV::OpConstantSampler:
675 case SPIRV::OpConstantNull:
677 case SPIRV::OpPoisonKHR:
678 case SPIRV::OpConstantFunctionPointerINTEL:
705 case Intrinsic::spv_all:
706 case Intrinsic::spv_alloca:
707 case Intrinsic::spv_any:
708 case Intrinsic::spv_bitcast:
709 case Intrinsic::spv_const_composite:
710 case Intrinsic::spv_degrees:
711 case Intrinsic::spv_distance:
712 case Intrinsic::spv_extractelt:
713 case Intrinsic::spv_extractv:
714 case Intrinsic::spv_faceforward:
715 case Intrinsic::spv_fdot:
716 case Intrinsic::spv_firstbitlow:
717 case Intrinsic::spv_firstbitshigh:
718 case Intrinsic::spv_firstbituhigh:
719 case Intrinsic::spv_frac:
720 case Intrinsic::spv_gep:
721 case Intrinsic::spv_global_offset:
722 case Intrinsic::spv_global_size:
723 case Intrinsic::spv_group_id:
724 case Intrinsic::spv_insertelt:
725 case Intrinsic::spv_insertv:
726 case Intrinsic::spv_isinf:
727 case Intrinsic::spv_isnan:
728 case Intrinsic::spv_isfinite:
729 case Intrinsic::spv_isnormal:
730 case Intrinsic::spv_lerp:
731 case Intrinsic::spv_length:
732 case Intrinsic::spv_normalize:
733 case Intrinsic::spv_num_subgroups:
734 case Intrinsic::spv_num_workgroups:
735 case Intrinsic::spv_ptrcast:
736 case Intrinsic::spv_radians:
737 case Intrinsic::spv_reflect:
738 case Intrinsic::spv_refract:
739 case Intrinsic::spv_resource_getbasepointer:
740 case Intrinsic::spv_resource_getpointer:
741 case Intrinsic::spv_resource_handlefrombinding:
742 case Intrinsic::spv_resource_handlefromimplicitbinding:
743 case Intrinsic::spv_resource_nonuniformindex:
744 case Intrinsic::spv_resource_sample:
745 case Intrinsic::spv_rsqrt:
746 case Intrinsic::spv_saturate:
747 case Intrinsic::spv_sdot:
748 case Intrinsic::spv_sign:
749 case Intrinsic::spv_smoothstep:
750 case Intrinsic::spv_subgroup_id:
751 case Intrinsic::spv_subgroup_local_invocation_id:
752 case Intrinsic::spv_subgroup_max_size:
753 case Intrinsic::spv_subgroup_size:
754 case Intrinsic::spv_thread_id:
755 case Intrinsic::spv_thread_id_in_group:
756 case Intrinsic::spv_udot:
757 case Intrinsic::spv_undef:
758 case Intrinsic::spv_value_md:
759 case Intrinsic::spv_workgroup_size:
771 case SPIRV::OpTypeVoid:
772 case SPIRV::OpTypeBool:
773 case SPIRV::OpTypeInt:
774 case SPIRV::OpTypeFloat:
775 case SPIRV::OpTypeVector:
776 case SPIRV::OpTypeVectorIdEXT:
777 case SPIRV::OpTypeMatrix:
778 case SPIRV::OpTypeImage:
779 case SPIRV::OpTypeSampler:
780 case SPIRV::OpTypeSampledImage:
781 case SPIRV::OpTypeArray:
782 case SPIRV::OpTypeRuntimeArray:
783 case SPIRV::OpTypeStruct:
784 case SPIRV::OpTypeOpaque:
785 case SPIRV::OpTypePointer:
786 case SPIRV::OpTypeFunction:
787 case SPIRV::OpTypeEvent:
788 case SPIRV::OpTypeDeviceEvent:
789 case SPIRV::OpTypeReserveId:
790 case SPIRV::OpTypeQueue:
791 case SPIRV::OpTypePipe:
792 case SPIRV::OpTypeForwardPointer:
793 case SPIRV::OpTypePipeStorage:
794 case SPIRV::OpTypeNamedBarrier:
795 case SPIRV::OpTypeAccelerationStructureNV:
796 case SPIRV::OpTypeCooperativeMatrixNV:
797 case SPIRV::OpTypeCooperativeMatrixKHR:
807 if (
MI.getNumDefs() == 0)
810 for (
const auto &MO :
MI.all_defs()) {
812 if (
Reg.isPhysical()) {
816 for (
const auto &
UseMI :
MRI.use_nodbg_instructions(
Reg)) {
817 if (
UseMI.getOpcode() != SPIRV::OpName) {
824 if (
MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE ||
MI.isFakeUse() ||
825 MI.isLifetimeMarker()) {
828 <<
"Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
839 if (
MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
840 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
843 LLVM_DEBUG(
dbgs() <<
"Dead: Intrinsic with no real side effects.\n");
848 if (
MI.mayStore() ||
MI.isCall() ||
849 (
MI.mayLoad() &&
MI.hasOrderedMemoryRef()) ||
MI.isPosition() ||
850 MI.isDebugInstr() ||
MI.isTerminator() ||
MI.isJumpTableDebugInfo()) {
851 LLVM_DEBUG(
dbgs() <<
"Not dead: instruction has side effects.\n");
862 LLVM_DEBUG(
dbgs() <<
"Dead: known opcode with no side effects\n");
869void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &
MI)
const {
871 for (
const auto &MO :
MI.all_defs()) {
875 SmallVector<MachineInstr *, 4> UselessOpNames;
878 "There is still a use of the dead function.");
881 for (MachineInstr *OpNameMI : UselessOpNames) {
883 OpNameMI->eraseFromParent();
888void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &
MI)
const {
891 removeOpNamesForDeadMI(
MI);
892 MI.eraseFromParent();
895bool SPIRVInstructionSelector::select(MachineInstr &
I) {
896 resetVRegsType(*
I.getParent()->getParent());
898 assert(
I.getParent() &&
"Instruction should be in a basic block!");
899 assert(
I.getParent()->getParent() &&
"Instruction should be in a function!");
904 removeDeadInstruction(
I);
911 if (Opcode == SPIRV::ASSIGN_TYPE) {
912 Register DstReg =
I.getOperand(0).getReg();
913 Register SrcReg =
I.getOperand(1).getReg();
916 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
917 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
918 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
919 Register SelectDstReg =
Def->getOperand(0).getReg();
920 bool SuccessToSelectSelect [[maybe_unused]] = selectSelect(
922 assert(SuccessToSelectSelect);
924 Def->eraseFromParent();
931 bool Res = selectImpl(
I, *CoverageInfo);
933 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
934 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
938 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
950 }
else if (
I.getNumDefs() == 1) {
962 removeDeadInstruction(
I);
967 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
968 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
974 bool HasDefs =
I.getNumDefs() > 0;
977 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
978 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
979 if (spvSelect(ResVReg, ResType,
I)) {
981 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
992 case TargetOpcode::G_CONSTANT:
993 case TargetOpcode::G_FCONSTANT:
1000 MachineInstr &
I)
const {
1003 if (DstRC != SrcRC && SrcRC)
1005 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
1012bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
1013 SPIRVTypeInst ResType,
1014 MachineInstr &
I)
const {
1015 const unsigned Opcode =
I.getOpcode();
1017 return selectImpl(
I, *CoverageInfo);
1019 case TargetOpcode::G_CONSTANT:
1020 case TargetOpcode::G_FCONSTANT:
1021 return selectConst(ResVReg, ResType,
I);
1022 case TargetOpcode::G_GLOBAL_VALUE:
1023 return selectGlobalValue(ResVReg,
I);
1024 case TargetOpcode::G_IMPLICIT_DEF:
1025 return selectOpUndef(ResVReg, ResType,
I);
1026 case TargetOpcode::G_FREEZE:
1027 return selectFreeze(ResVReg, ResType,
I);
1029 case TargetOpcode::G_INTRINSIC:
1030 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
1031 case TargetOpcode::G_INTRINSIC_CONVERGENT:
1032 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
1033 return selectIntrinsic(ResVReg, ResType,
I);
1034 case TargetOpcode::G_BITREVERSE:
1035 return selectBitreverse(ResVReg, ResType,
I);
1037 case TargetOpcode::G_BUILD_VECTOR:
1038 return selectBuildVector(ResVReg, ResType,
I);
1039 case TargetOpcode::G_SPLAT_VECTOR:
1040 return selectSplatVector(ResVReg, ResType,
I);
1041 case TargetOpcode::G_CONCAT_VECTORS:
1042 return selectConcatVectors(ResVReg, ResType,
I);
1044 case TargetOpcode::G_SHUFFLE_VECTOR: {
1045 MachineBasicBlock &BB = *
I.getParent();
1046 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
1049 .
addUse(
I.getOperand(1).getReg())
1050 .
addUse(
I.getOperand(2).getReg());
1051 for (
auto V :
I.getOperand(3).getShuffleMask())
1056 case TargetOpcode::G_MEMMOVE:
1057 case TargetOpcode::G_MEMCPY:
1058 case TargetOpcode::G_MEMCPY_INLINE:
1059 case TargetOpcode::G_MEMSET:
1060 case TargetOpcode::G_MEMSET_INLINE:
1061 return selectMemOperation(ResVReg,
I);
1063 case TargetOpcode::G_ICMP:
1064 return selectICmp(ResVReg, ResType,
I);
1065 case TargetOpcode::G_FCMP:
1066 return selectFCmp(ResVReg, ResType,
I);
1068 case TargetOpcode::G_FRAME_INDEX:
1069 return selectFrameIndex(ResVReg, ResType,
I);
1071 case TargetOpcode::G_LOAD:
1072 return selectLoad(ResVReg, ResType,
I);
1073 case TargetOpcode::G_STORE:
1074 return selectStore(
I);
1076 case TargetOpcode::G_BR:
1077 return selectBranch(
I);
1078 case TargetOpcode::G_BRCOND:
1079 return selectBranchCond(
I);
1081 case TargetOpcode::G_PHI:
1082 return selectPhi(ResVReg,
I);
1084 case TargetOpcode::G_FPTOSI:
1085 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1086 case TargetOpcode::G_FPTOUI:
1087 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1089 case TargetOpcode::G_FPTOSI_SAT:
1090 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1091 case TargetOpcode::G_FPTOUI_SAT:
1092 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1094 case TargetOpcode::G_SITOFP:
1095 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
1096 case TargetOpcode::G_UITOFP:
1097 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
1099 case TargetOpcode::G_CTPOP:
1100 return selectPopCount(ResVReg, ResType,
I, SPIRV::OpBitCount);
1101 case TargetOpcode::G_SMIN:
1102 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
1103 case TargetOpcode::G_UMIN:
1104 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
1106 case TargetOpcode::G_SMAX:
1107 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
1108 case TargetOpcode::G_UMAX:
1109 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
1111 case TargetOpcode::G_SCMP:
1112 return selectSUCmp(ResVReg, ResType,
I,
true);
1113 case TargetOpcode::G_UCMP:
1114 return selectSUCmp(ResVReg, ResType,
I,
false);
1115 case TargetOpcode::G_LROUND:
1116 case TargetOpcode::G_LLROUND: {
1121 regForLround, *(
I.getParent()->getParent()));
1123 CL::round, GL::Round,
false);
1125 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1132 case TargetOpcode::G_STRICT_FMA:
1133 case TargetOpcode::G_FMA: {
1136 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1139 .
addUse(
I.getOperand(1).getReg())
1140 .
addUse(
I.getOperand(2).getReg())
1141 .
addUse(
I.getOperand(3).getReg())
1146 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1149 case TargetOpcode::G_FLDEXP:
1150 case TargetOpcode::G_STRICT_FLDEXP:
1151 return selectLdexp(ResVReg, ResType,
I);
1153 case TargetOpcode::G_FPOW:
1154 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1155 case TargetOpcode::G_FPOWI:
1156 return selectFpowi(ResVReg, ResType,
I);
1158 case TargetOpcode::G_FEXP:
1159 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1160 case TargetOpcode::G_FEXP2:
1161 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1162 case TargetOpcode::G_FEXP10:
1163 return selectExp10(ResVReg, ResType,
I);
1165 case TargetOpcode::G_FMODF:
1166 return selectModf(ResVReg, ResType,
I);
1167 case TargetOpcode::G_FSINCOS:
1168 return selectSincos(ResVReg, ResType,
I);
1170 case TargetOpcode::G_FLOG:
1171 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1172 case TargetOpcode::G_FLOG2:
1173 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1174 case TargetOpcode::G_FLOG10:
1175 return selectLog10(ResVReg, ResType,
I);
1177 case TargetOpcode::G_FABS:
1178 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1179 case TargetOpcode::G_ABS:
1180 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1182 case TargetOpcode::G_FMINNUM:
1183 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1184 case TargetOpcode::G_FMAXNUM:
1185 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1187 case TargetOpcode::G_FCOPYSIGN:
1188 return selectCopySign(ResVReg, ResType,
I);
1190 case TargetOpcode::G_FCEIL:
1191 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1192 case TargetOpcode::G_FFLOOR:
1193 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1195 case TargetOpcode::G_FCOS:
1196 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1197 case TargetOpcode::G_FSIN:
1198 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1199 case TargetOpcode::G_FTAN:
1200 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1201 case TargetOpcode::G_FACOS:
1202 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1203 case TargetOpcode::G_FASIN:
1204 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1205 case TargetOpcode::G_FATAN:
1206 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1207 case TargetOpcode::G_FATAN2:
1208 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1209 case TargetOpcode::G_FCOSH:
1210 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1211 case TargetOpcode::G_FSINH:
1212 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1213 case TargetOpcode::G_FTANH:
1214 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1216 case TargetOpcode::G_STRICT_FSQRT:
1217 case TargetOpcode::G_FSQRT:
1218 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1220 case TargetOpcode::G_CTTZ:
1221 case TargetOpcode::G_CTTZ_ZERO_POISON:
1222 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1223 case TargetOpcode::G_CTLZ:
1224 case TargetOpcode::G_CTLZ_ZERO_POISON:
1225 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1227 case TargetOpcode::G_INTRINSIC_ROUND:
1228 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1229 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1230 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1231 case TargetOpcode::G_INTRINSIC_TRUNC:
1232 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1233 case TargetOpcode::G_FRINT:
1234 case TargetOpcode::G_FNEARBYINT:
1235 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1237 case TargetOpcode::G_SMULH:
1238 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1239 case TargetOpcode::G_UMULH:
1240 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1242 case TargetOpcode::G_SADDSAT:
1243 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1244 case TargetOpcode::G_UADDSAT:
1245 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1246 case TargetOpcode::G_SSUBSAT:
1247 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1248 case TargetOpcode::G_USUBSAT:
1249 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1251 case TargetOpcode::G_FFREXP:
1252 return selectFrexp(ResVReg, ResType,
I);
1254 case TargetOpcode::G_UADDO:
1255 return selectOverflowArith(ResVReg, ResType,
I,
1257 : SPIRV::OpIAddCarryS);
1258 case TargetOpcode::G_USUBO:
1259 return selectOverflowArith(ResVReg, ResType,
I,
1261 : SPIRV::OpISubBorrowS);
1262 case TargetOpcode::G_UMULO:
1263 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1264 case TargetOpcode::G_SMULO:
1265 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1267 case TargetOpcode::G_SEXT:
1268 return selectExt(ResVReg, ResType,
I,
true);
1269 case TargetOpcode::G_ANYEXT:
1270 case TargetOpcode::G_ZEXT:
1271 return selectExt(ResVReg, ResType,
I,
false);
1272 case TargetOpcode::G_TRUNC:
1273 return selectTrunc(ResVReg, ResType,
I);
1274 case TargetOpcode::G_FPTRUNC:
1275 case TargetOpcode::G_FPEXT:
1276 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1278 case TargetOpcode::G_PTRTOINT:
1279 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1280 case TargetOpcode::G_INTTOPTR:
1281 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1282 case TargetOpcode::G_BITCAST:
1283 return selectBitcast(ResVReg, ResType,
I);
1284 case TargetOpcode::G_ADDRSPACE_CAST:
1285 return selectAddrSpaceCast(ResVReg, ResType,
I);
1286 case TargetOpcode::G_PTRMASK:
1287 return selectPtrMask(ResVReg, ResType,
I);
1288 case TargetOpcode::G_PTR_ADD: {
1291 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1294 const bool IsAddrSpaceCast =
1295 (*II).getOpcode() == TargetOpcode::G_ADDRSPACE_CAST;
1296 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1297 (*II).getOpcode() == TargetOpcode::COPY || IsAddrSpaceCast ||
1298 (*II).getOpcode() == SPIRV::OpVariable ||
1299 (*II).getOpcode() == SPIRV::OpUntypedVariableKHR) &&
1302 bool IsGVInit =
false;
1306 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1307 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1308 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1309 (*UseIt).getOpcode() == SPIRV::OpVariable ||
1310 (*UseIt).getOpcode() == SPIRV::OpUntypedVariableKHR) {
1320 if (!
getImm((*II).getOperand(1),
MRI) ||
1322 return diagnoseUnsupported(
1323 I,
"address space cast of a non-null integer used as a G_PTR_ADD "
1325 SPIRVTypeInst NullPtrTy =
1327 MachineBasicBlock &DepMBB =
I.getMF()->front();
1338 const bool UseUntypedPointers =
1339 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1340 if (UseUntypedPointers) {
1341 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1344 .
addImm(
static_cast<uint32_t
>(
1345 SPIRV::Opcode::UntypedInBoundsPtrAccessChainKHR))
1348 .
addUse(
I.getOperand(2).getReg())
1355 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1367 return diagnoseUnsupported(
1368 I,
"incompatible result and operand types in a bitcast");
1370 MachineInstrBuilder MIB =
1371 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1378 : SPIRV::OpInBoundsPtrAccessChain))
1382 .
addUse(
I.getOperand(2).getReg())
1385 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1389 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1391 .
addUse(
I.getOperand(2).getReg())
1400 Register ChainBase = IsAddrSpaceCast ? GV : bitcastConstPtrToI8(
I, GV);
1405 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1408 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1410 .
addUse(
I.getOperand(2).getReg())
1415 case TargetOpcode::G_ATOMICRMW_OR:
1416 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1417 case TargetOpcode::G_ATOMICRMW_ADD:
1418 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1419 case TargetOpcode::G_ATOMICRMW_AND:
1420 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1421 case TargetOpcode::G_ATOMICRMW_MAX:
1422 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1423 case TargetOpcode::G_ATOMICRMW_MIN:
1424 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1425 case TargetOpcode::G_ATOMICRMW_SUB:
1426 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1427 case TargetOpcode::G_ATOMICRMW_XOR:
1428 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1429 case TargetOpcode::G_ATOMICRMW_UMAX:
1430 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1431 case TargetOpcode::G_ATOMICRMW_UMIN:
1432 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1433 case TargetOpcode::G_ATOMICRMW_XCHG:
1434 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1436 case TargetOpcode::G_ATOMICRMW_FADD:
1437 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1438 case TargetOpcode::G_ATOMICRMW_FSUB:
1440 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1442 : SPIRV::OpFNegate);
1443 case TargetOpcode::G_ATOMICRMW_FMIN:
1444 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1445 case TargetOpcode::G_ATOMICRMW_FMAX:
1446 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1448 case TargetOpcode::G_FENCE:
1449 return selectFence(
I);
1451 case TargetOpcode::G_STACKSAVE:
1452 return selectStackSave(ResVReg, ResType,
I);
1453 case TargetOpcode::G_STACKRESTORE:
1454 return selectStackRestore(
I);
1456 case TargetOpcode::G_UNMERGE_VALUES:
1459 case TargetOpcode::G_TRAP:
1460 case TargetOpcode::G_UBSANTRAP:
1461 return selectTrap(
I);
1466 case TargetOpcode::DBG_LABEL:
1468 case TargetOpcode::G_DEBUGTRAP:
1469 return selectDebugTrap(ResVReg, ResType,
I);
1470 case TargetOpcode::G_PREFETCH:
1471 return selectPrefetch(
I);
1478bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1479 SPIRVTypeInst ResType,
1480 MachineInstr &
I)
const {
1481 if (!STI.
canUseExtInstSet(SPIRV::InstructionSet::NonSemantic_DebugBreak)) {
1487 MachineIRBuilder MIRBuilder(
I);
1490 return selectExtInst(ExtReg, VoidTy,
I,
1491 {{SPIRV::InstructionSet::NonSemantic_DebugBreak,
1492 SPIRV::NonSemanticDebugBreakExtInst::DebugBreak}},
1496bool SPIRVInstructionSelector::selectPrefetch(MachineInstr &
I)
const {
1505 MachineIRBuilder MIRBuilder(
I);
1507 const SPIRVTypeInst PointerSizeType =
1515 Register AddrVal =
I.getOperand(0).getReg();
1518 return selectExtInst(ExtReg, GR.
getOpTypeVoid(MIRBuilder),
I, CL::prefetch,
1520 {AddrVal, ConstIntOne});
1525bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1526 SPIRVTypeInst ResType,
1528 GL::GLSLExtInst GLInst,
1529 bool setMIFlags,
bool useMISrc,
1532 SPIRV::InstructionSet::InstructionSet::GLSL_std_450))
1533 return diagnoseUnsupported(
1535 "this instruction is only supported with the GLSL extended instruction "
1537 return selectExtInst(ResVReg, ResType,
I,
1538 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1539 setMIFlags, useMISrc, SrcRegs);
1542bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1543 SPIRVTypeInst ResType,
1545 CL::OpenCLExtInst CLInst,
1546 bool setMIFlags,
bool useMISrc,
1548 return selectExtInst(ResVReg, ResType,
I,
1549 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1550 setMIFlags, useMISrc, SrcRegs);
1553bool SPIRVInstructionSelector::selectExtInst(
1554 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1555 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1557 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1558 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1559 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1563bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1564 SPIRVTypeInst ResType,
1567 bool setMIFlags,
bool useMISrc,
1570 for (
const auto &[InstructionSet, Opcode] : Insts) {
1574 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1577 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1582 const unsigned NumOps =
I.getNumOperands();
1585 I.getOperand(Index).getType() ==
1586 MachineOperand::MachineOperandType::MO_IntrinsicID)
1589 MIB.
add(
I.getOperand(Index));
1601bool SPIRVInstructionSelector::selectCopySign(
Register ResVReg,
1602 SPIRVTypeInst ResType,
1603 MachineInstr &
I)
const {
1605 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1610 Register MagnitudeReg =
I.getOperand(1).getReg();
1611 Register SignReg =
I.getOperand(2).getReg();
1619 unsigned AndOpcode, OrOpcode;
1620 if (ComponentCount > 1) {
1624 AndOpcode = SPIRV::OpBitwiseAndV;
1625 OrOpcode = SPIRV::OpBitwiseOrV;
1629 AndOpcode = SPIRV::OpBitwiseAndS;
1630 OrOpcode = SPIRV::OpBitwiseOrS;
1636 return selectOpWithSrcs(ResReg, IntType,
I, SrcRegs, Opcode);
1639 Register MagnitudeInt, SignInt, MagnitudeBits, SignBits, CombinedInt;
1640 if (!EmitBitOp(MagnitudeInt, {MagnitudeReg}, SPIRV::OpBitcast) ||
1641 !EmitBitOp(SignInt, {SignReg}, SPIRV::OpBitcast) ||
1642 !EmitBitOp(MagnitudeBits, {MagnitudeInt, NotSignMask}, AndOpcode) ||
1643 !EmitBitOp(SignBits, {SignInt, SignMask}, AndOpcode) ||
1644 !EmitBitOp(CombinedInt, {MagnitudeBits, SignBits}, OrOpcode))
1647 return selectOpWithSrcs(ResVReg, ResType,
I, {CombinedInt}, SPIRV::OpBitcast);
1650bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1651 SPIRVTypeInst ResType,
1652 MachineInstr &
I)
const {
1653 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1654 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1655 for (
const auto &Ex : ExtInsts) {
1656 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1657 uint32_t Opcode = Ex.second;
1661 MachineIRBuilder MIRBuilder(
I);
1664 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1671 const bool IsUntyped =
1672 PointerType->getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1674 BuildMI(*It->getParent(), It, It->getDebugLoc(),
1675 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
1676 : SPIRV::OpVariable))
1679 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
1684 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1687 .
addImm(
static_cast<uint32_t
>(Ex.first))
1689 .
add(
I.getOperand(2))
1693 Register ExpResReg =
I.getOperand(1).getReg();
1695 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1705bool SPIRVInstructionSelector::selectLdexp(
Register ResVReg,
1706 SPIRVTypeInst ResType,
1707 MachineInstr &
I)
const {
1708 Register XReg =
I.getOperand(1).getReg();
1709 Register ExpReg =
I.getOperand(2).getReg();
1715 if (ResType->
getOpcode() == SPIRV::OpTypeVector &&
1716 ExpType->
getOpcode() != SPIRV::OpTypeVector) {
1718 ExpReg = buildVectorSplat(ExpReg, NumElts,
I);
1721 return selectExtInst(ResVReg, ResType,
I, CL::ldexp, GL::Ldexp,
1722 true,
false, {XReg, ExpReg});
1725bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1726 SPIRVTypeInst ResType,
1727 MachineInstr &
I)
const {
1728 Register CosResVReg =
I.getOperand(1).getReg();
1729 unsigned SrcIdx =
I.getNumExplicitDefs();
1734 MachineIRBuilder MIRBuilder(
I);
1736 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1743 const bool IsUntyped =
1744 PointerType->getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1746 BuildMI(*It->getParent(), It, It->getDebugLoc(),
1747 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
1748 : SPIRV::OpVariable))
1751 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
1755 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1758 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1760 .
add(
I.getOperand(SrcIdx))
1764 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1772 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1775 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1777 .
add(
I.getOperand(SrcIdx))
1779 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1782 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1784 .
add(
I.getOperand(SrcIdx))
1791bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1792 SPIRVTypeInst ResType,
1795 unsigned Opcode)
const {
1796 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1806bool SPIRVInstructionSelector::selectPopCount16(
Register ResVReg,
1807 SPIRVTypeInst ResType,
1810 unsigned Opcode)
const {
1811 MachineIRBuilder MIRBuilder(
I);
1813 Register OpReg =
I.getOperand(1).getReg();
1819 SPIRVTypeInst ExtType =
1826 if (!selectOpWithSrcs(ExtReg, ExtType,
I, {OpReg}, SPIRV::OpUConvert))
1830 if (!selectPopCount32(PopCountReg, ExtType,
I, ExtReg, Opcode))
1833 return selectOpWithSrcs(ResVReg, ResType,
I, {PopCountReg}, ExtOpcode);
1836bool SPIRVInstructionSelector::selectPopCount32(
Register ResVReg,
1837 SPIRVTypeInst ResType,
1840 unsigned Opcode)
const {
1841 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
1844bool SPIRVInstructionSelector::selectPopCount64(
Register ResVReg,
1845 SPIRVTypeInst ResType,
1848 unsigned Opcode)
const {
1849 MachineIRBuilder MIRBuilder(
I);
1858 SPIRVTypeInst WorkingType =
1865 if (!selectOpWithSrcs(Trunc, WorkingType,
I, {SrcReg}, SPIRV::OpUConvert))
1869 if (!selectOpWithSrcs(LowCount, WorkingType,
I, {Trunc}, SPIRV::OpBitCount))
1877 IsScalar ? SPIRV::OpShiftRightLogicalS : SPIRV::OpShiftRightLogicalV;
1879 if (!selectOpWithSrcs(Shift, SrcType,
I, {SrcReg, ShiftAmount}, ShiftOp))
1883 if (!selectOpWithSrcs(Trunc, WorkingType,
I, {Shift}, SPIRV::OpUConvert))
1887 if (!selectOpWithSrcs(HighCount, WorkingType,
I, {Trunc}, SPIRV::OpBitCount))
1892 if (!selectOpWithSrcs(Sum, WorkingType,
I, {HighCount, LowCount},
1893 IsScalar ? SPIRV::OpIAddS : SPIRV::OpIAddV))
1897 unsigned ConvOp = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
1898 return selectOpWithSrcs(ResVReg, ResType,
I, {Sum}, ConvOp);
1901bool SPIRVInstructionSelector::selectPopCount(
Register ResVReg,
1902 SPIRVTypeInst ResType,
1904 unsigned Opcode)
const {
1909 if (!STI.getTargetTriple().isVulkanOS())
1910 return selectUnOp(ResVReg, ResType,
I, Opcode);
1912 Register OpReg =
I.getOperand(1).getReg();
1915 : SPIRV::OpUConvert;
1919 return selectPopCount16(ResVReg, ResType,
I, ExtOpcode, Opcode);
1921 return selectPopCount32(ResVReg, ResType,
I, OpReg, Opcode);
1923 return selectPopCount64(ResVReg, ResType,
I, OpReg, Opcode);
1925 return diagnoseUnsupported(
I,
"unsupported operand bit width for popcount");
1929bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1930 SPIRVTypeInst ResType,
1932 unsigned Opcode)
const {
1934 Register SrcReg =
I.getOperand(1).getReg();
1939 unsigned DefOpCode = DefIt->getOpcode();
1940 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1943 if (
auto *VRD =
getVRegDef(*
MRI, DefIt->getOperand(1).getReg()))
1944 DefOpCode = VRD->getOpcode();
1946 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1947 DefOpCode == TargetOpcode::G_CONSTANT ||
1948 DefOpCode == SPIRV::OpVariable ||
1949 DefOpCode == SPIRV::OpUntypedVariableKHR ||
1950 DefOpCode == SPIRV::OpConstantI) {
1956 uint32_t SpecOpcode = 0;
1958 case SPIRV::OpConvertPtrToU:
1959 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1961 case SPIRV::OpConvertUToPtr:
1962 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1967 TII.get(SPIRV::OpSpecConstantOp))
1977 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1981bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1982 SPIRVTypeInst ResType,
1983 MachineInstr &
I)
const {
1984 Register OpReg =
I.getOperand(1).getReg();
1985 SPIRVTypeInst OpType =
1988 return diagnoseUnsupported(
1989 I,
"incompatible result and operand types in a bitcast");
1990 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1996 std::optional<Align> AlignOverride = std::nullopt) {
2001 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
2002 if (
MemOp->isNonTemporal())
2003 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
2005 if (!ST->isShader() &&
MemOp->getAlign().value())
2006 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
2010 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
2011 if (
auto *MD =
MemOp->getAAInfo().Scope) {
2015 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
2017 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
2021 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
2025 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
2027 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
2028 MIB.
addImm(AlignOverride.value_or(
MemOp->getAlign()).value());
2039 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
2041 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
2043 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
2047bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
2048 SPIRVTypeInst ResType,
2049 MachineInstr &
I)
const {
2051 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2056 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2057 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2059 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2061 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2065 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2069 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2070 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
2071 I.getDebugLoc(),
I);
2075 MachineIRBuilder MIRBuilder(
I);
2077 if (
I.getNumMemOperands()) {
2078 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2079 if (MemOp->isAtomic())
2080 return selectAtomicLoad(ResVReg, ResType,
I);
2083 auto MIB = MIRBuilder.
buildInstr(SPIRV::OpLoad)
2087 if (!
I.getNumMemOperands()) {
2088 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2090 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2099Register SPIRVInstructionSelector::createPtrSizedIntReg(
2100 MachineIRBuilder &MIRBuilder)
const {
2101 SPIRVTypeInst IntType =
2111SPIRVInstructionSelector::convertPtrToInt(
Register PtrVal,
2112 MachineIRBuilder &MIRBuilder)
const {
2113 SPIRVTypeInst IntType =
2115 Register IntReg = createPtrSizedIntReg(MIRBuilder);
2116 MIRBuilder.
buildInstr(SPIRV::OpConvertPtrToU)
2124Register SPIRVInstructionSelector::castPtrToPtrToInt(
2125 Register Ptr, SPIRV::StorageClass::StorageClass SC,
2126 MachineIRBuilder &MIRBuilder)
const {
2127 SPIRVTypeInst IntType =
2129 SPIRVTypeInst PtrType =
2143bool SPIRVInstructionSelector::selectAtomicPtrValue(
2144 Register ResVReg, SPIRVTypeInst ResType, MachineIRBuilder &MIRBuilder,
2145 function_ref<
Register(SPIRVTypeInst IntType)> EmitAtomic)
const {
2153 Register IntResult = EmitAtomic(IntType);
2155 MIRBuilder.
buildInstr(SPIRV::OpConvertUToPtr)
2163bool SPIRVInstructionSelector::selectAtomicLoad(
Register ResVReg,
2164 SPIRVTypeInst ResType,
2165 MachineInstr &
I)
const {
2166 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2169 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2172 return diagnoseUnsupported(
2173 I,
"Lowering to SPIR-V of atomic load is only "
2174 "allowed for integer, floating point or pointer types");
2176 assert(
I.getNumMemOperands());
2177 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2178 assert(MemOp.isAtomic());
2180 uint32_t
Scope =
static_cast<uint32_t
>(
2182 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
2188 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2189 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2192 Register MemSemReg = buildI32ConstantInEntryBlock(Sem,
I);
2194 MachineIRBuilder MIRBuilder(
I);
2198 return diagnoseUnsupported(
2199 I,
"Lowering to SPIR-V of atomic load is only "
2200 "allowed for pointer types for physical addressing model");
2205 SPIRV::StorageClass::StorageClass SC =
2207 return selectAtomicPtrValue(
2208 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2209 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2210 Register IntResult = createPtrSizedIntReg(MIRBuilder);
2221 auto AtomicLoad = MIRBuilder.
buildInstr(SPIRV::OpAtomicLoad)
2232bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
2234 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2235 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2240 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2241 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2243 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2248 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2252 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2253 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2254 SPIRVTypeInst SampledType =
2256 SPIRVTypeInst StoreValCompType =
2258 if (StoreValCompType && StoreValCompType != SampledType) {
2261 SPIRVTypeInst PackedType = widenTypeToVec4(SampledType,
I);
2264 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
2269 StoreVal = PackedReg;
2272 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2273 TII.get(SPIRV::OpImageWrite))
2279 if (sampledTypeIsSignedInteger(LLVMHandleType))
2282 BMI.constrainAllUses(
TII,
TRI, RBI);
2289 if (PointeeTy && PointeeTy->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
2290 StoreTy->
getOpcode() != SPIRV::OpTypeVectorIdEXT &&
2304 if (
I.getNumMemOperands()) {
2305 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2306 if (MemOp->isAtomic())
2307 return selectAtomicStore(
I);
2314 if (PtrSC == SPIRV::StorageClass::UniformConstant ||
2315 PtrSC == SPIRV::StorageClass::Input ||
2316 PtrSC == SPIRV::StorageClass::PushConstant)
2317 return diagnoseUnsupported(
2318 I,
"store into a read-only SPIR-V storage class is not allowed");
2320 MachineIRBuilder MIRBuilder(
I);
2322 if (!
I.getNumMemOperands()) {
2323 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2325 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2334bool SPIRVInstructionSelector::selectAtomicStore(MachineInstr &
I)
const {
2335 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2338 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2339 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2344 if (!PointeeType && PtrType &&
2345 PtrType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR)
2348 return diagnoseUnsupported(
I,
2349 "Lowering to SPIR-V of atomic store is only "
2350 "allowed for integer or floating point types");
2352 assert(
I.getNumMemOperands());
2353 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2354 assert(MemOp.isAtomic());
2356 uint32_t
Scope =
static_cast<uint32_t
>(
2358 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
2364 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2365 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2368 Register MemSemReg = buildI32ConstantInEntryBlock(Sem,
I);
2369 MachineIRBuilder MIRBuilder(
I);
2373 return diagnoseUnsupported(
2374 I,
"Lowering to SPIR-V of atomic store is only "
2375 "allowed for pointer types for physical addressing model");
2380 SPIRV::StorageClass::StorageClass SC =
2382 return selectAtomicPtrValue(
2383 Register(), SPIRVTypeInst(), MIRBuilder, [&](SPIRVTypeInst IntType) {
2385 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2398 return diagnoseUnsupported(
I,
2399 "Lowering to SPIR-V of atomic store is only "
2400 "allowed for integer or floating point types");
2402 auto AtomicStore = MIRBuilder.
buildInstr(SPIRV::OpAtomicStore)
2412bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
2413 SPIRVTypeInst ResType,
2414 MachineInstr &
I)
const {
2415 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
2423 const Register PtrsReg =
I.getOperand(2).getReg();
2424 const uint32_t
Alignment =
I.getOperand(3).getImm();
2425 const Register MaskReg =
I.getOperand(4).getReg();
2426 const Register PassthruReg =
I.getOperand(5).getReg();
2427 const Register AlignmentReg = buildI32ConstantInEntryBlock(Alignment,
I);
2431 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
2442bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
2443 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
2450 const Register ValuesReg =
I.getOperand(1).getReg();
2451 const Register PtrsReg =
I.getOperand(2).getReg();
2452 const uint32_t
Alignment =
I.getOperand(3).getImm();
2453 const Register MaskReg =
I.getOperand(4).getReg();
2454 const Register AlignmentReg = buildI32ConstantInEntryBlock(Alignment,
I);
2458 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
2467bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
2468 const Twine &
Msg)
const {
2469 const Function &
F =
I.getMF()->getFunction();
2470 F.getContext().diagnose(
2471 DiagnosticInfoUnsupported(
F,
Msg,
I.getDebugLoc(),
DS_Error));
2475bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
2476 SPIRVTypeInst ResType,
2477 MachineInstr &
I)
const {
2478 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2479 return diagnoseUnsupported(
2480 I,
"llvm.stacksave intrinsic: this instruction requires the following "
2481 "SPIR-V extension: SPV_INTEL_variable_length_array");
2483 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
2490bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
2491 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2492 return diagnoseUnsupported(
2494 "llvm.stackrestore intrinsic: this instruction requires the following "
2495 "SPIR-V extension: SPV_INTEL_variable_length_array");
2496 if (!
I.getOperand(0).isReg())
2499 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
2500 .
addUse(
I.getOperand(0).getReg())
2506SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
2507 MachineIRBuilder MIRBuilder(
I);
2508 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
2515 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
2519 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2520 Type *ArrTy = ArrayType::get(ValTy, Num);
2522 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
2525 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
2536 const bool IsUntyped = VarTy->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
2537 auto MIBVar =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2538 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
2539 : SPIRV::OpVariable))
2542 .
addImm(SPIRV::StorageClass::UniformConstant);
2555bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
2558 Register DstReg =
I.getOperand(0).getReg();
2564 return diagnoseUnsupported(
2565 I,
"OpCopyMemory requires operands to have the same type");
2570 return diagnoseUnsupported(
2571 I,
"Unable to determine pointee type size for OpCopyMemory");
2572 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
2573 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
2574 return diagnoseUnsupported(
2575 I,
"OpCopyMemory requires the size to match the pointee type size");
2578 const unsigned Opcode =
2579 IsLogical ? SPIRV::OpCopyMemory : SPIRV::OpCopyMemorySized;
2580 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
2584 MIB.
addUse(
I.getOperand(2).getReg());
2586 if (
I.getNumMemOperands()) {
2587 MachineIRBuilder MIRBuilder(
I);
2588 const MachineMemOperand *DstMemOp = *
I.memoperands_begin();
2590 Align SrcAlign = DstAlign;
2593 if (
I.getNumMemOperands() > 1 && !STI.
isShader())
2594 SrcAlign = (*std::next(
I.memoperands_begin()))->
getAlign();
2604 std::min(DstAlign, SrcAlign));
2611bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
2612 MachineInstr &
I)
const {
2614 Register SizeReg =
I.getOperand(2).getReg();
2616 SizeDef && SizeDef->
getOpcode() == TargetOpcode::G_CONSTANT &&
2620 Register SrcReg =
I.getOperand(1).getReg();
2621 if (
I.getOpcode() == TargetOpcode::G_MEMSET ||
2622 I.getOpcode() == TargetOpcode::G_MEMSET_INLINE) {
2623 Register VarReg = getOrCreateMemSetGlobal(
I);
2626 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2628 ValTy,
I, SPIRV::StorageClass::UniformConstant);
2630 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
2633 if (!selectCopyMemory(
I, SrcReg))
2635 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
2636 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
2641bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
2642 SPIRVTypeInst ResType,
2645 unsigned NegateOpcode)
const {
2647 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2648 uint32_t
Scope =
static_cast<uint32_t
>(
2650 MemOp->getSyncScopeID()));
2651 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
2653 Register Ptr =
I.getOperand(1).getReg();
2654 uint32_t ScSem =
static_cast<uint32_t
>(
2658 Register MemSemReg = buildI32ConstantInEntryBlock(
2662 Register ValueReg =
I.getOperand(2).getReg();
2663 if (NegateOpcode != 0) {
2666 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
2672 if (NewOpcode != SPIRV::OpAtomicExchange)
2673 return diagnoseUnsupported(
2674 I,
"Lowering to SPIR-V of this atomic operation is not "
2675 "allowed for pointer types");
2677 return diagnoseUnsupported(
2678 I,
"Lowering to SPIR-V of atomic exchange is only "
2679 "allowed for pointer types for physical addressing model");
2686 MachineIRBuilder MIRBuilder(
I);
2688 return selectAtomicPtrValue(
2689 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2691 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2692 Register ExchangeResReg = createPtrSizedIntReg(MIRBuilder);
2693 MIRBuilder.
buildInstr(SPIRV::OpAtomicExchange)
2701 return ExchangeResReg;
2705 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2716bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2717 unsigned ArgI =
I.getNumOperands() - 1;
2719 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2720 SPIRVTypeInst SrcType =
2724 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2728 unsigned CurrentIndex = 0;
2729 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2730 Register ResVReg =
I.getOperand(i).getReg();
2739 ResType = ScalarType;
2748 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2754 for (
unsigned j = 0;
j < NumElements; ++
j) {
2755 MIB.
addImm(CurrentIndex + j);
2757 CurrentIndex += NumElements;
2761 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2773bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2778 ? SPIRV::MemorySemantics::UniformMemory |
2779 SPIRV::MemorySemantics::WorkgroupMemory |
2780 SPIRV::MemorySemantics::ImageMemory
2781 : SPIRV::MemorySemantics::WorkgroupMemory |
2782 SPIRV::MemorySemantics::CrossWorkgroupMemory;
2784 STI.getTargetTriple(),
static_cast<uint32_t
>(
getMemSemantics(AO)), ScSem);
2785 Register MemSemReg = buildI32ConstantInEntryBlock(MemSem,
I);
2789 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
2791 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2798bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2799 SPIRVTypeInst ResType,
2801 unsigned Opcode)
const {
2802 Type *ResTy =
nullptr;
2805 return diagnoseUnsupported(
2807 "Not enough info to select the arithmetic with overflow instruction");
2809 return diagnoseUnsupported(
I,
2810 "Expect struct type result for the arithmetic "
2811 "with overflow instruction");
2817 MachineIRBuilder MIRBuilder(
I);
2819 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2820 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2827 Register ZeroReg = buildZerosVal(ResType,
I);
2832 if (ResName.
size() > 0)
2840 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2841 MIB.
addUse(
I.getOperand(i).getReg());
2847 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2849 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2850 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2857 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2858 .
addDef(
I.getOperand(1).getReg())
2866bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2867 SPIRVTypeInst ResType,
2868 MachineInstr &
I)
const {
2870 "selectAtomicCmpXchg only handles the spv_cmpxchg intrinsic");
2871 Register Ptr =
I.getOperand(2).getReg();
2872 Register ScopeReg =
I.getOperand(5).getReg();
2873 Register MemSemEqReg =
I.getOperand(6).getReg();
2874 Register MemSemNeqReg =
I.getOperand(7).getReg();
2876 Register Val =
I.getOperand(4).getReg();
2880 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2899 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2906 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2918 case SPIRV::StorageClass::DeviceOnlyINTEL:
2919 case SPIRV::StorageClass::HostOnlyINTEL:
2929 bool IsGRef =
false;
2930 bool IsAllowedRefs =
2931 llvm::all_of(
MRI->use_instructions(ResVReg), [&IsGRef](
auto const &It) {
2932 unsigned Opcode = It.getOpcode();
2933 if (Opcode == SPIRV::OpConstantComposite ||
2934 Opcode == SPIRV::OpSpecConstantComposite ||
2935 Opcode == SPIRV::OpSpecConstantOp || Opcode == SPIRV::OpVariable ||
2936 Opcode == SPIRV::OpUntypedVariableKHR ||
2937 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2938 return IsGRef = true;
2939 return Opcode == SPIRV::OpName;
2941 return IsAllowedRefs && IsGRef;
2944SPIRVTypeInst SPIRVInstructionSelector::getUcharPtrType(
2945 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2947 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC);
2950Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2951 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2956SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2958 uint32_t Opcode)
const {
2959 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2960 TII.get(SPIRV::OpSpecConstantOp))
2969Register SPIRVInstructionSelector::bitcastConstPtrToI8(MachineInstr &
I,
2978 MachineInstrBuilder Cast =
2980 static_cast<uint32_t
>(SPIRV::Opcode::Bitcast));
2987SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2988 SPIRVTypeInst SrcPtrTy)
const {
2989 SPIRVTypeInst GenericPtrTy =
2993 SPIRV::StorageClass::Generic),
2997 MachineInstrBuilder MIB = buildSpecConstantOp(
2999 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
3009bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
3010 SPIRVTypeInst ResType,
3011 MachineInstr &
I)
const {
3015 Register SrcPtr =
I.getOperand(1).getReg();
3020 return BuildCOPY(ResVReg, SrcPtr,
I);
3030 unsigned SpecOpcode = [&]() ->
unsigned {
3031 if (SrcSC == SPIRV::StorageClass::CodeSectionINTEL)
3032 return static_cast<uint32_t
>(SPIRV::Opcode::Bitcast);
3034 return static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric);
3036 return static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr);
3047 if (SrcSC != SPIRV::StorageClass::CodeSectionINTEL &&
3048 (SpecOpcode || ViaGeneric)) {
3049 SrcPtr = bitcastConstPtrToI8(
I, SrcPtr);
3053 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
3055 .constrainAllUses(
TII,
TRI, RBI);
3056 }
else if (ViaGeneric) {
3057 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
3059 buildSpecConstantOp(
3061 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
3062 .constrainAllUses(
TII,
TRI, RBI);
3069 return BuildCOPY(ResVReg, SrcPtr,
I);
3071 if ((SrcSC == SPIRV::StorageClass::Function &&
3072 DstSC == SPIRV::StorageClass::Private) ||
3073 (DstSC == SPIRV::StorageClass::Function &&
3074 SrcSC == SPIRV::StorageClass::Private))
3075 return BuildCOPY(ResVReg, SrcPtr,
I);
3079 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
3082 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
3085 SPIRVTypeInst GenericPtrTy =
3104 return selectUnOp(ResVReg, ResType,
I,
3105 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
3107 return selectUnOp(ResVReg, ResType,
I,
3108 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
3110 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
3112 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
3122bool SPIRVInstructionSelector::selectPtrMask(
Register ResVReg,
3123 SPIRVTypeInst ResType,
3124 MachineInstr &
I)
const {
3126 return diagnoseUnsupported(
3127 I,
"G_PTRMASK is not supported with logical SPIR-V");
3132 Register PtrReg =
I.getOperand(1).getReg();
3133 Register MaskReg =
I.getOperand(2).getReg();
3152 ? SPIRV::OpBitwiseAndV
3153 : SPIRV::OpBitwiseAndS;
3176 return SPIRV::OpFOrdEqual;
3178 return SPIRV::OpFOrdGreaterThanEqual;
3180 return SPIRV::OpFOrdGreaterThan;
3182 return SPIRV::OpFOrdLessThanEqual;
3184 return SPIRV::OpFOrdLessThan;
3186 return SPIRV::OpFOrdNotEqual;
3188 return SPIRV::OpOrdered;
3190 return SPIRV::OpFUnordEqual;
3192 return SPIRV::OpFUnordGreaterThanEqual;
3194 return SPIRV::OpFUnordGreaterThan;
3196 return SPIRV::OpFUnordLessThanEqual;
3198 return SPIRV::OpFUnordLessThan;
3200 return SPIRV::OpFUnordNotEqual;
3202 return SPIRV::OpUnordered;
3212 return SPIRV::OpIEqual;
3214 return SPIRV::OpINotEqual;
3216 return SPIRV::OpSGreaterThanEqual;
3218 return SPIRV::OpSGreaterThan;
3220 return SPIRV::OpSLessThanEqual;
3222 return SPIRV::OpSLessThan;
3224 return SPIRV::OpUGreaterThanEqual;
3226 return SPIRV::OpUGreaterThan;
3228 return SPIRV::OpULessThanEqual;
3230 return SPIRV::OpULessThan;
3239 return SPIRV::OpPtrEqual;
3241 return SPIRV::OpPtrNotEqual;
3252 return SPIRV::OpLogicalEqual;
3254 return SPIRV::OpLogicalNotEqual;
3292bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
3293 SPIRVTypeInst ResType,
3295 unsigned OpAnyOrAll)
const {
3296 assert(
I.getNumOperands() == 3);
3297 assert(
I.getOperand(2).isReg());
3299 Register InputRegister =
I.getOperand(2).getReg();
3302 assert(InputType &&
"VReg has no type assigned");
3306 assert(ResVReg ==
I.getOperand(0).getReg());
3307 return BuildCOPY(ResVReg, InputRegister,
I);
3311 unsigned SpirvNotEqualId =
3312 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
3314 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
3319 IsBoolTy ? InputRegister
3327 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
3329 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
3346bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
3347 SPIRVTypeInst ResType,
3348 MachineInstr &
I)
const {
3349 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
3352bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
3353 SPIRVTypeInst ResType,
3354 MachineInstr &
I)
const {
3355 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
3359bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
3360 SPIRVTypeInst ResType,
3361 MachineInstr &
I)
const {
3362 assert(
I.getNumOperands() == 4);
3363 assert(
I.getOperand(2).isReg());
3364 assert(
I.getOperand(3).isReg());
3366 [[maybe_unused]] SPIRVTypeInst VecType =
3371 VecType->
getOpcode() == SPIRV::OpTypeVectorIdEXT) &&
3372 "dot product requires either a vector of at least 2 components or"
3373 " the SPV_EXT_long vector extension.");
3375 [[maybe_unused]] SPIRVTypeInst EltType =
3384 .
addUse(
I.getOperand(2).getReg())
3385 .
addUse(
I.getOperand(3).getReg())
3390bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
3391 SPIRVTypeInst ResType,
3394 assert(
I.getNumOperands() == 4);
3395 assert(
I.getOperand(2).isReg());
3396 assert(
I.getOperand(3).isReg());
3399 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3403 .
addUse(
I.getOperand(2).getReg())
3404 .
addUse(
I.getOperand(3).getReg())
3411bool SPIRVInstructionSelector::selectIntegerDotExpansion(
3412 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3413 assert(
I.getNumOperands() == 4);
3414 assert(
I.getOperand(2).isReg());
3415 assert(
I.getOperand(3).isReg());
3419 Register Vec0 =
I.getOperand(2).getReg();
3420 Register Vec1 =
I.getOperand(3).getReg();
3424 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
3433 VecType->
getOpcode() == SPIRV::OpTypeVectorIdEXT) &&
3434 "dot product requires either a vector of at least 2 components "
3435 "or the SPV_EXT_long_vector extension.");
3438 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3448 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3459 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3471bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
3472 SPIRVTypeInst ResType,
3473 MachineInstr &
I)
const {
3475 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
3478 .
addUse(
I.getOperand(2).getReg())
3483bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
3484 SPIRVTypeInst ResType,
3485 MachineInstr &
I)
const {
3487 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
3490 .
addUse(
I.getOperand(2).getReg())
3495bool SPIRVInstructionSelector::selectOpIsFinite(
Register ResVReg,
3496 SPIRVTypeInst ResType,
3497 MachineInstr &
I)
const {
3500 Register Src =
I.getOperand(2).getReg();
3543bool SPIRVInstructionSelector::selectOpIsNormal(
Register ResVReg,
3544 SPIRVTypeInst ResType,
3545 MachineInstr &
I)
const {
3547 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNormal))
3550 .
addUse(
I.getOperand(2).getReg())
3555template <
bool Signed>
3556bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
3557 SPIRVTypeInst ResType,
3558 MachineInstr &
I)
const {
3559 assert(
I.getNumOperands() == 5);
3560 assert(
I.getOperand(2).isReg());
3561 assert(
I.getOperand(3).isReg());
3562 assert(
I.getOperand(4).isReg());
3565 Register Acc =
I.getOperand(2).getReg();
3569 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3571 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
3576 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
3579 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3591template <
bool Signed>
3592bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
3593 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3594 assert(
I.getNumOperands() == 5);
3595 assert(
I.getOperand(2).isReg());
3596 assert(
I.getOperand(3).isReg());
3597 assert(
I.getOperand(4).isReg());
3600 Register Acc =
I.getOperand(2).getReg();
3606 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
3610 for (
unsigned i = 0; i < 4; i++) {
3633 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
3653 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3668bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
3669 SPIRVTypeInst ResType,
3670 MachineInstr &
I)
const {
3671 assert(
I.getNumOperands() == 3);
3672 assert(
I.getOperand(2).isReg());
3674 Register VZero = buildZerosValF(ResType,
I);
3675 Register VOne = buildOnesValF(ResType,
I);
3677 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
3680 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3682 .
addUse(
I.getOperand(2).getReg())
3689bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
3690 SPIRVTypeInst ResType,
3691 MachineInstr &
I)
const {
3692 assert(
I.getNumOperands() == 3);
3693 assert(
I.getOperand(2).isReg());
3695 Register InputRegister =
I.getOperand(2).getReg();
3697 auto &
DL =
I.getDebugLoc();
3700 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3707 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
3709 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
3717 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3722 if (NeedsConversion) {
3723 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
3734bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
3735 SPIRVTypeInst ResType,
3737 unsigned Opcode)
const {
3741 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
3747 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
3748 BMI.addUse(
I.getOperand(J).getReg());
3755bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
3758 bool WithGroupSync)
const {
3760 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
3762 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
3764 assert(((Scope != SPIRV::Scope::Workgroup) ||
3765 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
3766 "Workgroup Scope must set WorkGroupMemory semantic "
3767 "in Barrier instruction");
3769 assert(((Scope != SPIRV::Scope::Device) ||
3770 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
3771 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
3772 "Device Scope must set UniformMemory and ImageMemory semantic "
3773 "in Barrier instruction");
3779 if (WithGroupSync) {
3780 Register ExecReg = buildI32ConstantInEntryBlock(SPIRV::Scope::Workgroup,
I);
3784 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
3785 Register MemSemReg = buildI32ConstantInEntryBlock(MemSem,
I);
3787 MI.addUse(ScopeReg).addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
3791bool SPIRVInstructionSelector::selectWaveActiveCountBits(
3792 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3797 if (!selectWaveOpInst(BallotReg, BallotType,
I,
3798 SPIRV::OpGroupNonUniformBallot))
3803 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3808 .
addImm(SPIRV::GroupOperation::Reduce)
3815bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
3816 SPIRVTypeInst ResType,
3817 MachineInstr &
I)
const {
3822 Register InputReg =
I.getOperand(2).getReg();
3827 bool IsVector = NumElems > 1 ||
3828 (InputType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
3842 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
3843 SPIRV::OpGroupNonUniformAllEqual);
3848 ElementResults.
reserve(NumElems);
3850 for (
unsigned Idx = 0;
Idx < NumElems; ++
Idx) {
3863 ElemInput = Extracted;
3869 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3880 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3891bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3892 SPIRVTypeInst ResType,
3893 MachineInstr &
I)
const {
3895 assert(
I.getNumOperands() == 3);
3897 auto Op =
I.getOperand(2);
3907 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3909 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3910 return diagnoseUnsupported(
I,
"WavePrefixBitCount requires boolean input");
3931 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3935 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3942bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3943 SPIRVTypeInst ResType,
3945 bool IsUnsigned)
const {
3946 return selectWaveReduce(
3947 ResVReg, ResType,
I, IsUnsigned,
3948 [&](
Register InputRegister,
bool IsUnsigned) {
3949 const bool IsFloatTy =
3951 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3952 : SPIRV::OpGroupNonUniformSMax;
3953 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3957bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3958 SPIRVTypeInst ResType,
3960 bool IsUnsigned)
const {
3961 return selectWaveReduce(
3962 ResVReg, ResType,
I, IsUnsigned,
3963 [&](
Register InputRegister,
bool IsUnsigned) {
3964 const bool IsFloatTy =
3966 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3967 : SPIRV::OpGroupNonUniformSMin;
3968 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3972bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3973 SPIRVTypeInst ResType,
3974 MachineInstr &
I)
const {
3975 return selectWaveReduce(ResVReg, ResType,
I,
false,
3976 [&](
Register InputRegister,
bool IsUnsigned) {
3978 InputRegister, SPIRV::OpTypeFloat);
3979 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3980 : SPIRV::OpGroupNonUniformIAdd;
3984bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3985 SPIRVTypeInst ResType,
3986 MachineInstr &
I)
const {
3987 return selectWaveReduce(ResVReg, ResType,
I,
false,
3988 [&](
Register InputRegister,
bool IsUnsigned) {
3990 InputRegister, SPIRV::OpTypeFloat);
3991 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3992 : SPIRV::OpGroupNonUniformIMul;
3996template <
typename PickOpcodeFn>
3997bool SPIRVInstructionSelector::selectWaveReduce(
3998 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3999 PickOpcodeFn &&PickOpcode)
const {
4000 assert(
I.getNumOperands() == 3);
4001 assert(
I.getOperand(2).isReg());
4003 Register InputRegister =
I.getOperand(2).getReg();
4007 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
4010 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
4016 .
addImm(SPIRV::GroupOperation::Reduce)
4017 .
addUse(
I.getOperand(2).getReg())
4022bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
4023 SPIRVTypeInst ResType,
4025 unsigned Opcode)
const {
4026 return selectWaveReduce(
4027 ResVReg, ResType,
I,
false,
4028 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
4031bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
4032 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4033 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
4034 [&](
Register InputRegister,
bool IsUnsigned) {
4036 InputRegister, SPIRV::OpTypeFloat);
4038 ? SPIRV::OpGroupNonUniformFAdd
4039 : SPIRV::OpGroupNonUniformIAdd;
4043bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
4044 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4045 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
4046 [&](
Register InputRegister,
bool IsUnsigned) {
4048 InputRegister, SPIRV::OpTypeFloat);
4050 ? SPIRV::OpGroupNonUniformFMul
4051 : SPIRV::OpGroupNonUniformIMul;
4055template <
typename PickOpcodeFn>
4056bool SPIRVInstructionSelector::selectWaveExclusiveScan(
4057 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
4058 PickOpcodeFn &&PickOpcode)
const {
4059 assert(
I.getNumOperands() == 3);
4060 assert(
I.getOperand(2).isReg());
4062 Register InputRegister =
I.getOperand(2).getReg();
4066 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
4069 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
4075 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
4076 .
addUse(
I.getOperand(2).getReg())
4081bool SPIRVInstructionSelector::selectWaveMultiPrefixOp(
4082 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
unsigned Opcode,
4083 StringRef BuiltinName)
const {
4085 SPIRV::Extension::SPV_EXT_shader_subgroup_partitioned) &&
4087 SPIRV::Extension::SPV_NV_shader_subgroup_partitioned))
4088 return diagnoseUnsupported(
4089 I, Twine(BuiltinName) +
4090 " requires SPV_EXT_shader_subgroup_partitioned or "
4091 "SPV_NV_shader_subgroup_partitioned");
4093 assert(
I.getNumOperands() == 7);
4098 auto MaskMIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4099 TII.get(SPIRV::OpCompositeConstruct))
4102 for (
unsigned Op = 3;
Op != 7; ++
Op)
4103 MaskMIB.
addUse(
I.getOperand(
Op).getReg());
4106 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4111 .
addImm(SPIRV::GroupOperation::PartitionedExclusiveScanNV)
4112 .
addUse(
I.getOperand(2).getReg())
4118bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
4119 SPIRVTypeInst ResType,
4122 assert(
I.getNumOperands() == 3);
4123 assert(
I.getOperand(2).isReg());
4125 Register InputRegister =
I.getOperand(2).getReg();
4131 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
4142bool SPIRVInstructionSelector::selectBitreverseViaI32(
Register ResVReg,
4143 SPIRVTypeInst ResType,
4150 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
4155 : SPIRV::OpUConvert;
4157 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4160 ShiftOp = SPIRV::OpShiftRightLogicalV;
4165 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4166 TII.get(SPIRV::OpConstantComposite))
4169 for (
unsigned It = 0; It <
N; ++It)
4173 ShiftConst = CompositeReg;
4178 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
4183 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
4188 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
4193 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
4196bool SPIRVInstructionSelector::handle64BitOverflow(
4198 unsigned int Opcode,
4205 "handle64BitOverflow should only be used for integer types");
4207 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
4209 MachineIRBuilder MIRBuilder(
I);
4211 SPIRVTypeInst I64x2Type =
4213 SPIRVTypeInst Vec2ResType =
4216 std::vector<Register> PartialRegs;
4218 unsigned CurrentComponent = 0;
4219 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
4223 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4224 TII.get(SPIRV::OpVectorShuffle))
4229 .
addImm(CurrentComponent)
4230 .
addImm(CurrentComponent + 1);
4240 PartialRegs.push_back(SubVecReg);
4243 if (CurrentComponent != ComponentCount) {
4249 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
4250 SPIRV::OpVectorExtractDynamic))
4259 PartialRegs.push_back(FinalElemResReg);
4263 return selectOpWithSrcs(ResVReg, ResType,
I, PartialRegs,
4264 SPIRV::OpCompositeConstruct);
4267bool SPIRVInstructionSelector::selectBitreverse64(
Register ResVReg,
4268 SPIRVTypeInst ResType,
4272 if (ComponentCount > 2)
4273 return handle64BitOverflow(
4274 ResVReg, ResType,
I, SrcReg, SPIRV::OpBitReverse,
4276 unsigned O) {
return this->selectBitreverse64(R,
T,
I, S); });
4278 MachineIRBuilder MIRBuilder(
I);
4282 I32Type, 2 * ComponentCount, MIRBuilder,
false);
4286 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
4291 if (!selectBitreverseNative(Reverse32, VecI32Type,
I, Vec32))
4298 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4299 TII.get(SPIRV::OpVectorShuffle))
4304 for (
unsigned J = 0; J < ComponentCount; ++J) {
4311 return selectOpWithSrcs(ResVReg, ResType,
I, {SwappedVec}, SPIRV::OpBitcast);
4314bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
4315 SPIRVTypeInst ResType,
4319 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
4327bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
4328 SPIRVTypeInst ResType,
4329 MachineInstr &
I)
const {
4330 Register OpReg =
I.getOperand(1).getReg();
4339 return selectBitreverseViaI32(ResVReg, ResType,
I, OpReg);
4341 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4343 return selectBitreverse64(ResVReg, ResType,
I, OpReg);
4345 return SPIRVInstructionSelector::diagnoseUnsupported(
4346 I,
"G_BITREVERSE only support 16,32,64 bits.");
4350 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4361 unsigned AndOp = SPIRV::OpBitwiseAndS;
4362 unsigned OrOp = SPIRV::OpBitwiseOrS;
4363 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
4364 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
4365 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4367 AndOp = SPIRV::OpBitwiseAndV;
4368 OrOp = SPIRV::OpBitwiseOrV;
4369 ShlOp = SPIRV::OpShiftLeftLogicalV;
4370 ShrOp = SPIRV::OpShiftRightLogicalV;
4376 const unsigned Shift) ->
Register {
4379 (ResType->
getOpcode() != SPIRV::OpTypeVectorIdEXT ||
4386 Register MaskReg = CreateConst(Mask);
4387 Register ShiftReg = CreateConst(Shift);
4394 if (!selectOpWithSrcs(
T1, ResType,
I, {
Input, ShiftReg}, ShrOp) ||
4395 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
4396 !selectOpWithSrcs(T3, ResType,
I, {
Input, MaskReg}, AndOp) ||
4397 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
4398 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
4407 while ((Shift >>= 1) > 0) {
4414 return BuildCOPY(ResVReg, Result,
I);
4417bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
4418 SPIRVTypeInst ResType,
4419 MachineInstr &
I)
const {
4420 assert(
I.getOperand(0).isReg() &&
I.getOperand(1).isReg() &&
4421 "G_FREEZE must define and use a register");
4422 Register OpReg =
I.getOperand(1).getReg();
4426 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
4440 if (
Def->getOpcode() == TargetOpcode::COPY)
4443 switch (
Def->getOpcode()) {
4444 case SPIRV::ASSIGN_TYPE:
4445 if (MachineInstr *AssignToDef =
4447 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
4448 Reg =
Def->getOperand(2).getReg();
4451 case SPIRV::OpUndef:
4452 Reg =
Def->getOperand(1).getReg();
4455 unsigned DestOpCode;
4457 DestOpCode = SPIRV::OpConstantNull;
4458 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze of a "
4459 "static undef/poison lowered to OpConstantNull\n");
4461 DestOpCode = TargetOpcode::COPY;
4463 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze "
4464 "skipped, lowered as a copy of the operand\n");
4466 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
4467 .
addDef(
I.getOperand(0).getReg())
4475bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
4476 SPIRVTypeInst ResType,
4477 MachineInstr &
I)
const {
4481 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4485 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
4490 for (
unsigned i =
I.getNumExplicitDefs();
4491 i <
I.getNumExplicitOperands() && IsConst; ++i)
4496 return diagnoseUnsupported(
4497 I,
"There must be at least two constituent operands in a vector");
4502 for (
unsigned i =
I.getNumExplicitDefs();
4503 i <
I.getNumExplicitOperands() && IsNullVector; ++i) {
4509 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4516 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4517 TII.get(IsConst ? SPIRV::OpConstantComposite
4518 : SPIRV::OpCompositeConstruct))
4521 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
4522 MIB.
addUse(
I.getOperand(i).getReg());
4527bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
4528 SPIRVTypeInst ResType,
4529 MachineInstr &
I)
const {
4533 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4538 unsigned OpIdx =
I.getNumExplicitDefs();
4539 if (!
I.getOperand(OpIdx).isReg())
4543 Register OpReg =
I.getOperand(OpIdx).getReg();
4547 return diagnoseUnsupported(
4548 I,
"There must be at least two constituent operands in a vector");
4551 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4552 TII.get(IsConst ? SPIRV::OpConstantComposite
4553 : SPIRV::OpCompositeConstruct))
4556 for (
unsigned i = 0; i <
N; ++i)
4562bool SPIRVInstructionSelector::selectConcatVectors(
Register ResVReg,
4563 SPIRVTypeInst ResType,
4564 MachineInstr &
I)
const {
4570 "Cannot select G_CONCAT_VECTORS with a non-vector result");
4572 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4573 TII.get(SPIRV::OpCompositeConstruct))
4576 for (
unsigned OpIdx =
I.getNumExplicitDefs();
4577 OpIdx <
I.getNumExplicitOperands(); ++OpIdx)
4578 MIB.
addUse(
I.getOperand(OpIdx).getReg());
4583bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
4584 SPIRVTypeInst ResType,
4585 MachineInstr &
I)
const {
4591 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
4593 Opcode = SPIRV::OpDemoteToHelperInvocation;
4595 Opcode = SPIRV::OpKill;
4600 ToErase.eraseFromParent();
4609bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
4610 SPIRVTypeInst ResType,
unsigned CmpOpc,
4611 MachineInstr &
I)
const {
4612 Register Cmp0 =
I.getOperand(2).getReg();
4613 Register Cmp1 =
I.getOperand(3).getReg();
4616 "CMP operands should have the same type");
4617 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
4627bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
4628 SPIRVTypeInst ResType,
4629 MachineInstr &
I)
const {
4633 Register CmpOperand =
I.getOperand(2).getReg();
4635 bool IsPtrCmp = CmpOperandType && CmpOperandType.
isPointer();
4640 Register Op1 =
I.getOperand(3).getReg();
4644 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
4649 I.getOperand(3).setReg(NewOp1);
4655 return selectCmp(ResVReg, ResType, CmpOpc,
I);
4660Register SPIRVInstructionSelector::buildI32ConstantInEntryBlock(
4661 uint32_t Val, MachineInstr &
I, SPIRVTypeInst ResType)
const {
4663 SPIRVTypeInst SpvI32Ty =
4665 auto *ConstInt = ConstantInt::get(LLVMTy, Val);
4670 MachineBasicBlock &EntryBB = *InsertIt->getParent();
4671 MachineInstr *
MI =
nullptr;
4675 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantNull))
4679 uint64_t ImmVal = APInt(32, Val).getZExtValue();
4680 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantI))
4686 GR.
add(ConstInt,
MI);
4691bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
4692 SPIRVTypeInst ResType,
4693 MachineInstr &
I)
const {
4702 for (
unsigned Idx = 0;
Idx < 2; ++
Idx) {
4704 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
4707 .
addUse(
I.getOperand(2 + Idx).getReg())
4714 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLogicalOr))
4721 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLogicalNot))
4729 return selectCmp(ResVReg, ResType, CmpOp,
I);
4732bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
4733 SPIRVTypeInst ResType,
4734 MachineInstr &
I)
const {
4736 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
4746 MachineIRBuilder MIRBuilder(
I);
4753 APFloat ConstVal(3.3219280948873623);
4757 APFloat::rmNearestTiesToEven, &LosesInfo);
4762 isVectorType(ResType) ? SPIRV::OpVectorTimesScalar : SPIRV::OpFMulS;
4764 if (!selectOpWithSrcs(ArgReg, ResType,
I,
4765 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
4767 if (!selectExtInst(ResVReg, ResType,
I,
4768 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
4778Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
4779 MachineInstr &
I)
const {
4787bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
4793 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
4801 if (
Def->getOpcode() == SPIRV::OpConstantNull)
4804 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
4805 Def->getOpcode() == SPIRV::OpConstantI)
4818 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
4819 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
4821 Intrinsic::spv_const_composite)) {
4822 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
4823 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
4824 if (!IsZero(
Def->getOperand(i).getReg()))
4833Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
4834 MachineInstr &
I)
const {
4843Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
4844 MachineInstr &
I)
const {
4854 SPIRVTypeInst ResType,
4855 MachineInstr &
I)
const {
4866 MachineInstr &
I)
const {
4871 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4872 TII.get(SPIRV::OpCompositeConstruct))
4875 for (
unsigned J = 0; J < NumElts; ++J)
4881bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
4882 SPIRVTypeInst ResType,
4883 MachineInstr &
I)
const {
4884 Register SelectFirstArg =
I.getOperand(2).getReg();
4885 Register SelectSecondArg =
I.getOperand(3).getReg();
4893 Register CondReg =
I.getOperand(1).getReg();
4894 bool IsScalarBool = GR.
isScalarOfType(CondReg, SPIRV::OpTypeBool);
4902 CondReg = buildVectorSplat(CondReg, NumElts,
I);
4903 IsScalarBool =
false;
4906 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
4907 }
else if (IsPtrTy) {
4908 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
4910 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
4913 assert(IsScalarBool &&
"OpSelect with a scalar result requires a scalar "
4914 "boolean condition");
4916 Opcode = SPIRV::OpSelectSFSCond;
4917 }
else if (IsPtrTy) {
4918 Opcode = SPIRV::OpSelectSPSCond;
4920 Opcode = SPIRV::OpSelectSISCond;
4923 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4935bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
4936 SPIRVTypeInst ResType,
4938 MachineInstr &InsertAt,
4939 bool IsSigned)
const {
4941 Register ZeroReg = buildZerosVal(ResType, InsertAt);
4942 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
4943 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
4945 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4957bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
4958 SPIRVTypeInst ResType,
4959 MachineInstr &
I,
bool IsSigned,
4960 unsigned Opcode)
const {
4961 Register SrcReg =
I.getOperand(1).getReg();
4972 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I, IsSigned);
4974 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
4977bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
4978 SPIRVTypeInst ResType, MachineInstr &
I,
4979 bool IsSigned)
const {
4980 Register SrcReg =
I.getOperand(1).getReg();
4982 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
4986 if (ResType == SrcType)
4987 return BuildCOPY(ResVReg, SrcReg,
I);
4989 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4990 return selectUnOp(ResVReg, ResType,
I, Opcode);
4993bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
4994 SPIRVTypeInst ResType,
4996 bool IsSigned)
const {
4997 MachineIRBuilder MIRBuilder(
I);
4998 MachineRegisterInfo *
MRI = MIRBuilder.
getMRI();
5003 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
5011 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
5014 .
addUse(
I.getOperand(1).getReg())
5015 .
addUse(
I.getOperand(2).getReg())
5020 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
5023 .
addUse(
I.getOperand(1).getReg())
5024 .
addUse(
I.getOperand(2).getReg())
5032 unsigned SelectOpcode =
5033 (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
5035 ? SPIRV::OpSelectVIVCond
5036 : SPIRV::OpSelectSISCond;
5041 .
addUse(buildOnesVal(
true, ResType,
I))
5042 .
addUse(buildZerosVal(ResType,
I))
5049 .
addUse(buildOnesVal(
false, ResType,
I))
5054bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
5057 SPIRVTypeInst IntTy,
5058 SPIRVTypeInst BoolTy)
const {
5062 isVectorType(IntTy) ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
5064 Register One = buildOnesVal(
false, IntTy,
I);
5072 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
5081bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
5082 SPIRVTypeInst ResType,
5083 MachineInstr &
I)
const {
5084 Register IntReg =
I.getOperand(1).getReg();
5087 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
5088 if (ArgType == ResType)
5089 return BuildCOPY(ResVReg, IntReg,
I);
5091 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
5092 return selectUnOp(ResVReg, ResType,
I, Opcode);
5095bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
5096 SPIRVTypeInst ResType,
5097 MachineInstr &
I)
const {
5098 unsigned Opcode =
I.getOpcode();
5099 unsigned TpOpcode = ResType->
getOpcode();
5101 if (ResType.
isPointer() || TpOpcode == SPIRV::OpTypeEvent) {
5102 assert(Opcode == TargetOpcode::G_CONSTANT &&
5103 I.getOperand(1).getCImm()->isZero());
5104 MachineBasicBlock &DepMBB =
I.getMF()->
front();
5107 }
else if (TpOpcode == SPIRV::OpTypeVectorIdEXT) {
5112 "Expected <1 x T> Vector!");
5113 if (Opcode == TargetOpcode::G_FCONSTANT)
5119 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
5127 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
5130bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
5131 SPIRVTypeInst ResType,
5132 MachineInstr &
I)
const {
5133 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5140bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
5141 SPIRVTypeInst ResType,
5142 MachineInstr &
I)
const {
5144 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
5148 .
addUse(
I.getOperand(3).getReg())
5150 .
addUse(
I.getOperand(2).getReg());
5151 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
5157bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
5158 SPIRVTypeInst ResType,
5159 MachineInstr &
I)
const {
5160 Type *MaybeResTy =
nullptr;
5165 "Expected aggregate type for extractv instruction");
5167 SPIRV::AccessQualifier::ReadWrite,
false);
5171 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
5174 .
addUse(
I.getOperand(2).getReg());
5175 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
5181bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
5182 SPIRVTypeInst ResType,
5183 MachineInstr &
I)
const {
5185 return selectInsertVal(ResVReg, ResType,
I);
5187 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
5190 .
addUse(
I.getOperand(2).getReg())
5191 .
addUse(
I.getOperand(3).getReg())
5192 .
addUse(
I.getOperand(4).getReg())
5197bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
5198 SPIRVTypeInst ResType,
5199 MachineInstr &
I)
const {
5201 return selectExtractVal(ResVReg, ResType,
I);
5203 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
5206 .
addUse(
I.getOperand(2).getReg())
5207 .
addUse(
I.getOperand(3).getReg())
5212bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
5213 SPIRVTypeInst ResType,
5214 MachineInstr &
I)
const {
5215 const bool IsGEPInBounds =
I.getOperand(2).getImm();
5218 const bool UseUntypedPointers =
5219 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
5224 if (UseUntypedPointers) {
5226 Opcode = IsGEPInBounds ? SPIRV::OpUntypedInBoundsAccessChainKHR
5227 : SPIRV::OpUntypedAccessChainKHR;
5229 Opcode = IsGEPInBounds ? SPIRV::OpUntypedInBoundsPtrAccessChainKHR
5230 : SPIRV::OpUntypedPtrAccessChainKHR;
5239 IsGEPInBounds ? SPIRV::OpInBoundsAccessChain : SPIRV::OpAccessChain;
5241 Opcode = IsGEPInBounds ? SPIRV::OpInBoundsPtrAccessChain
5242 : SPIRV::OpPtrAccessChain;
5247 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5252 if (UseUntypedPointers) {
5267 while (Def &&
Def->getOpcode() == TargetOpcode::COPY &&
5268 Def->getOperand(1).isReg())
5270 if (Def &&
Def->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)
5271 if (
const auto *GVar =
5274 SPIRV::AccessQualifier::ReadWrite,
5278 return diagnoseUnsupported(
5279 I,
"could not deduce the base type of an untyped access chain");
5284 Res.addUse(BaseReg);
5286 const bool IsAccessChainOpcode =
5287 (Opcode == SPIRV::OpAccessChain ||
5288 Opcode == SPIRV::OpInBoundsAccessChain ||
5289 Opcode == SPIRV::OpUntypedAccessChainKHR ||
5290 Opcode == SPIRV::OpUntypedInBoundsAccessChainKHR);
5294 "Cannot translate GEP to OpAccessChain.");
5297 const unsigned StartingIndex = IsAccessChainOpcode ? 5 : 4;
5298 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
5299 Res.addUse(
I.getOperand(i).getReg());
5300 Res.constrainAllUses(
TII,
TRI, RBI);
5309 if (Extract.
getOpcode() == SPIRV::OpCompositeExtract) {
5315 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5321 TII.get(SPIRV::OpCompositeInsert))
5335bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
5337 unsigned Lim =
I.getNumExplicitOperands();
5338 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
5339 Register OpReg =
I.getOperand(i).getReg();
5343 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
5344 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
5357 SPIRVTypeInst WrapType = OpType;
5358 if (OpType->
getOpcode() == SPIRV::OpTypePointer &&
5360 SPIRV::StorageClass::CodeSectionINTEL) {
5362 SPIRV::StorageClass::Function,
I);
5369 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5370 TII.get(SPIRV::OpSpecConstantOp))
5373 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
5375 GR.
add(OpDefine, MIB);
5381bool SPIRVInstructionSelector::selectDerivativeInst(
5382 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
5383 const unsigned DPdOpCode)
const {
5386 if (!errorIfInstrOutsideShader(
I))
5392 Register SrcReg =
I.getOperand(2).getReg();
5397 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5400 .
addUse(
I.getOperand(2).getReg());
5402 MachineIRBuilder MIRBuilder(
I);
5405 if (componentCount != 1)
5413 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5418 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5423 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5431bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
5432 SPIRVTypeInst ResType,
5433 MachineInstr &
I)
const {
5437 case Intrinsic::spv_load:
5438 return selectLoad(ResVReg, ResType,
I);
5439 case Intrinsic::spv_atomic_load:
5440 return selectAtomicLoad(ResVReg, ResType,
I);
5441 case Intrinsic::spv_store:
5442 return selectStore(
I);
5443 case Intrinsic::spv_atomic_store:
5444 return selectAtomicStore(
I);
5445 case Intrinsic::spv_extractv:
5446 return selectExtractVal(ResVReg, ResType,
I);
5447 case Intrinsic::spv_insertv:
5448 return selectInsertVal(ResVReg, ResType,
I);
5449 case Intrinsic::spv_extractelt:
5450 return selectExtractElt(ResVReg, ResType,
I);
5451 case Intrinsic::spv_insertelt:
5452 return selectInsertElt(ResVReg, ResType,
I);
5453 case Intrinsic::spv_gep:
5454 return selectGEP(ResVReg, ResType,
I);
5455 case Intrinsic::spv_bitcast: {
5456 Register OpReg =
I.getOperand(2).getReg();
5457 SPIRVTypeInst OpType =
5461 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
5463 case Intrinsic::spv_unref_global:
5464 case Intrinsic::spv_init_global: {
5470 Register GVarVReg =
MI->getOperand(0).getReg();
5471 if (!selectGlobalValue(GVarVReg, *
MI, Init))
5476 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
5478 MI->eraseFromParent();
5482 case Intrinsic::spv_undef: {
5483 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5489 case Intrinsic::spv_poison:
5490 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpPoisonKHR))
5495 case Intrinsic::spv_freeze:
5496 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
5499 .
addUse(
I.getOperand(2).getReg())
5502 case Intrinsic::spv_named_boolean_spec_constant: {
5503 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
5504 : SPIRV::OpSpecConstantFalse;
5506 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
5507 .
addDef(
I.getOperand(0).getReg())
5510 unsigned SpecId =
I.getOperand(2).getImm();
5512 SPIRV::Decoration::SpecId, {SpecId});
5516 case Intrinsic::spv_const_composite: {
5518 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
5524 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
5526 std::function<bool(
Register)> HasSpecConstOperand =
5536 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
5537 J < Def->getNumExplicitOperands(); ++J) {
5538 if (
Def->getOperand(J).isReg() &&
5539 HasSpecConstOperand(
Def->getOperand(J).getReg()))
5545 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
5546 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
5547 : SPIRV::OpConstantComposite;
5548 unsigned ContinuedOpc = HasSpecConst
5549 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
5550 : SPIRV::OpConstantCompositeContinuedINTEL;
5551 MachineIRBuilder MIR(
I);
5553 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
5555 for (
auto *Instr : Instructions) {
5556 Instr->setDebugLoc(
I.getDebugLoc());
5561 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5568 case Intrinsic::spv_assign_name: {
5569 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
5570 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
5571 for (
unsigned i =
I.getNumExplicitDefs() + 2;
5572 i <
I.getNumExplicitOperands(); ++i) {
5573 MIB.
addImm(
I.getOperand(i).getImm());
5578 case Intrinsic::spv_switch: {
5579 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
5580 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5581 if (
I.getOperand(i).isReg())
5582 MIB.
addReg(
I.getOperand(i).getReg());
5583 else if (
I.getOperand(i).isCImm())
5584 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
5585 else if (
I.getOperand(i).isMBB())
5586 MIB.
addMBB(
I.getOperand(i).getMBB());
5593 case Intrinsic::spv_loop_merge: {
5594 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
5595 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5596 if (
I.getOperand(i).isMBB())
5597 MIB.
addMBB(
I.getOperand(i).getMBB());
5604 case Intrinsic::spv_loop_control_intel: {
5606 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
5607 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
5612 case Intrinsic::spv_selection_merge: {
5614 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
5615 assert(
I.getOperand(1).isMBB() &&
5616 "operand 1 to spv_selection_merge must be a basic block");
5617 MIB.
addMBB(
I.getOperand(1).getMBB());
5618 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
5622 case Intrinsic::spv_cmpxchg:
5623 return selectAtomicCmpXchg(ResVReg, ResType,
I);
5624 case Intrinsic::spv_unreachable:
5625 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
5628 case Intrinsic::spv_abort:
5629 return selectAbort(
I);
5630 case Intrinsic::spv_alloca:
5631 return selectFrameIndex(ResVReg, ResType,
I);
5632 case Intrinsic::spv_alloca_array:
5633 return selectAllocaArray(ResVReg, ResType,
I);
5634 case Intrinsic::spv_assume:
5636 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
5637 .
addUse(
I.getOperand(1).getReg())
5642 case Intrinsic::spv_expect:
5644 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
5647 .
addUse(
I.getOperand(2).getReg())
5648 .
addUse(
I.getOperand(3).getReg())
5653 case Intrinsic::arithmetic_fence:
5654 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
5655 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
5658 .
addUse(
I.getOperand(2).getReg())
5662 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
5664 case Intrinsic::spv_thread_id:
5670 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
5672 case Intrinsic::spv_thread_id_in_group:
5678 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
5680 case Intrinsic::spv_group_id:
5686 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
5688 case Intrinsic::spv_flattened_thread_id_in_group:
5695 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
5697 case Intrinsic::spv_workgroup_size:
5698 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
5700 case Intrinsic::spv_global_size:
5701 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
5703 case Intrinsic::spv_global_offset:
5704 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
5706 case Intrinsic::spv_num_workgroups:
5707 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
5709 case Intrinsic::spv_subgroup_size:
5710 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
5712 case Intrinsic::spv_num_subgroups:
5713 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
5715 case Intrinsic::spv_subgroup_id:
5716 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
5717 case Intrinsic::spv_subgroup_local_invocation_id:
5718 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
5719 ResVReg, ResType,
I);
5720 case Intrinsic::spv_subgroup_max_size:
5721 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
5723 case Intrinsic::spv_fdot:
5724 return selectFloatDot(ResVReg, ResType,
I);
5725 case Intrinsic::spv_udot:
5726 case Intrinsic::spv_sdot:
5727 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5729 return selectIntegerDot(ResVReg, ResType,
I,
5730 IID == Intrinsic::spv_sdot);
5731 return selectIntegerDotExpansion(ResVReg, ResType,
I);
5732 case Intrinsic::spv_dot4add_i8packed:
5733 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5735 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
5736 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
5737 case Intrinsic::spv_dot4add_u8packed:
5738 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5740 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
5741 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
5742 case Intrinsic::spv_all:
5743 return selectAll(ResVReg, ResType,
I);
5744 case Intrinsic::spv_any:
5745 return selectAny(ResVReg, ResType,
I);
5746 case Intrinsic::spv_distance:
5747 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
5748 case Intrinsic::spv_lerp:
5749 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
5750 case Intrinsic::spv_length:
5751 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
5752 case Intrinsic::spv_degrees:
5753 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
5754 case Intrinsic::spv_faceforward:
5755 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
5756 case Intrinsic::spv_frac:
5757 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
5758 case Intrinsic::spv_isinf:
5759 return selectOpIsInf(ResVReg, ResType,
I);
5760 case Intrinsic::spv_isnan:
5761 return selectOpIsNan(ResVReg, ResType,
I);
5762 case Intrinsic::spv_isfinite:
5763 return selectOpIsFinite(ResVReg, ResType,
I);
5764 case Intrinsic::spv_isnormal:
5765 return selectOpIsNormal(ResVReg, ResType,
I);
5766 case Intrinsic::spv_normalize:
5767 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
5768 case Intrinsic::spv_refract:
5769 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
5770 case Intrinsic::spv_reflect:
5771 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
5772 case Intrinsic::spv_rsqrt:
5773 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
5774 case Intrinsic::spv_sign:
5775 return selectSign(ResVReg, ResType,
I);
5776 case Intrinsic::spv_smoothstep:
5777 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
5778 case Intrinsic::spv_firstbituhigh:
5779 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
5780 case Intrinsic::spv_firstbitshigh:
5781 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
5782 case Intrinsic::spv_firstbitlow:
5783 return selectFirstBitLow(ResVReg, ResType,
I);
5784 case Intrinsic::spv_all_memory_barrier:
5785 return selectBarrierInst(
I, SPIRV::Scope::Device,
5786 SPIRV::MemorySemantics::UniformMemory |
5787 SPIRV::MemorySemantics::ImageMemory |
5788 SPIRV::MemorySemantics::WorkgroupMemory,
5790 case Intrinsic::spv_all_memory_barrier_with_group_sync:
5791 return selectBarrierInst(
I, SPIRV::Scope::Device,
5792 SPIRV::MemorySemantics::UniformMemory |
5793 SPIRV::MemorySemantics::ImageMemory |
5794 SPIRV::MemorySemantics::WorkgroupMemory,
5796 case Intrinsic::spv_device_memory_barrier:
5797 return selectBarrierInst(
I, SPIRV::Scope::Device,
5798 SPIRV::MemorySemantics::UniformMemory |
5799 SPIRV::MemorySemantics::ImageMemory,
5801 case Intrinsic::spv_device_memory_barrier_with_group_sync:
5802 return selectBarrierInst(
I, SPIRV::Scope::Device,
5803 SPIRV::MemorySemantics::UniformMemory |
5804 SPIRV::MemorySemantics::ImageMemory,
5806 case Intrinsic::spv_group_memory_barrier:
5807 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5808 SPIRV::MemorySemantics::WorkgroupMemory,
5810 case Intrinsic::spv_group_memory_barrier_with_group_sync:
5811 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5812 SPIRV::MemorySemantics::WorkgroupMemory,
5814 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
5815 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
5816 SPIRV::StorageClass::StorageClass ResSC =
5819 return diagnoseUnsupported(
I,
"The target storage class is not castable "
5820 "from the Generic storage class");
5821 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
5829 case Intrinsic::spv_lifetime_start:
5830 case Intrinsic::spv_lifetime_end: {
5831 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
5832 : SPIRV::OpLifetimeStop;
5833 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
5834 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
5843 case Intrinsic::spv_saturate:
5844 return selectSaturate(ResVReg, ResType,
I);
5845 case Intrinsic::spv_nclamp:
5846 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
5847 case Intrinsic::spv_uclamp:
5848 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
5849 case Intrinsic::spv_sclamp:
5850 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
5851 case Intrinsic::spv_subgroup_prefix_bit_count:
5852 return selectWavePrefixBitCount(ResVReg, ResType,
I);
5853 case Intrinsic::spv_wave_active_countbits:
5854 return selectWaveActiveCountBits(ResVReg, ResType,
I);
5855 case Intrinsic::spv_wave_all_equal:
5856 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
5857 case Intrinsic::spv_wave_all:
5858 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
5859 case Intrinsic::spv_wave_any:
5860 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
5861 case Intrinsic::spv_subgroup_ballot:
5862 return selectWaveOpInst(ResVReg, ResType,
I,
5863 SPIRV::OpGroupNonUniformBallot);
5864 case Intrinsic::spv_wave_is_first_lane:
5865 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
5866 case Intrinsic::spv_wave_reduce_or:
5867 return selectWaveReduceOp(ResVReg, ResType,
I,
5868 SPIRV::OpGroupNonUniformBitwiseOr);
5869 case Intrinsic::spv_wave_reduce_xor:
5870 return selectWaveReduceOp(ResVReg, ResType,
I,
5871 SPIRV::OpGroupNonUniformBitwiseXor);
5872 case Intrinsic::spv_wave_reduce_and:
5873 return selectWaveReduceOp(ResVReg, ResType,
I,
5874 SPIRV::OpGroupNonUniformBitwiseAnd);
5875 case Intrinsic::spv_wave_reduce_umax:
5876 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
5877 case Intrinsic::spv_wave_reduce_max:
5878 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
5879 case Intrinsic::spv_wave_reduce_umin:
5880 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
5881 case Intrinsic::spv_wave_reduce_min:
5882 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
5883 case Intrinsic::spv_wave_reduce_sum:
5884 return selectWaveReduceSum(ResVReg, ResType,
I);
5885 case Intrinsic::spv_wave_product:
5886 return selectWaveReduceProduct(ResVReg, ResType,
I);
5887 case Intrinsic::spv_wave_readlane:
5888 return selectWaveOpInst(ResVReg, ResType,
I,
5889 SPIRV::OpGroupNonUniformShuffle);
5890 case Intrinsic::spv_wave_readlane_first:
5891 return selectWaveOpInst(ResVReg, ResType,
I,
5892 SPIRV::OpGroupNonUniformBroadcastFirst);
5893 case Intrinsic::spv_wave_prefix_sum:
5894 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
5895 case Intrinsic::spv_subgroup_multi_prefix_bit_and:
5896 return selectWaveMultiPrefixOp(ResVReg, ResType,
I,
5897 SPIRV::OpGroupNonUniformBitwiseAnd,
5898 "WaveMultiPrefixBitAnd");
5899 case Intrinsic::spv_subgroup_multi_prefix_bit_or:
5900 return selectWaveMultiPrefixOp(ResVReg, ResType,
I,
5901 SPIRV::OpGroupNonUniformBitwiseOr,
5902 "WaveMultiPrefixBitOr");
5903 case Intrinsic::spv_subgroup_multi_prefix_bit_xor:
5904 return selectWaveMultiPrefixOp(ResVReg, ResType,
I,
5905 SPIRV::OpGroupNonUniformBitwiseXor,
5906 "WaveMultiPrefixBitXor");
5907 case Intrinsic::spv_subgroup_multi_prefix_product: {
5910 return selectWaveMultiPrefixOp(ResVReg, ResType,
I,
5911 IsFloatTy ? SPIRV::OpGroupNonUniformFMul
5912 : SPIRV::OpGroupNonUniformIMul,
5913 "WaveMultiPrefixProduct");
5915 case Intrinsic::spv_subgroup_multi_prefix_sum: {
5918 return selectWaveMultiPrefixOp(ResVReg, ResType,
I,
5919 IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
5920 : SPIRV::OpGroupNonUniformIAdd,
5921 "WaveMultiPrefixSum");
5923 case Intrinsic::spv_wave_prefix_product:
5924 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
5925 case Intrinsic::spv_quad_all:
5926 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadAllKHR))
5929 .
addUse(
I.getOperand(2).getReg())
5932 case Intrinsic::spv_quad_read_across_x: {
5933 return selectQuadSwap(ResVReg, ResType,
I, 0);
5935 case Intrinsic::spv_quad_read_across_y: {
5936 return selectQuadSwap(ResVReg, ResType,
I, 1);
5938 case Intrinsic::spv_quad_read_across_diagonal: {
5939 return selectQuadSwap(ResVReg, ResType,
I, 2);
5941 case Intrinsic::spv_radians:
5942 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
5946 case Intrinsic::instrprof_increment:
5947 case Intrinsic::instrprof_increment_step:
5948 case Intrinsic::instrprof_value_profile:
5951 case Intrinsic::spv_value_md:
5953 case Intrinsic::spv_resource_handlefrombinding: {
5954 return selectHandleFromBinding(ResVReg, ResType,
I);
5956 case Intrinsic::spv_resource_counterhandlefrombinding:
5957 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
5958 case Intrinsic::spv_resource_updatecounter:
5959 return selectUpdateCounter(ResVReg, ResType,
I);
5960 case Intrinsic::spv_resource_store_typedbuffer: {
5961 return selectImageWriteIntrinsic(
I);
5963 case Intrinsic::spv_resource_load_typedbuffer: {
5964 return selectReadImageIntrinsic(ResVReg, ResType,
I);
5966 case Intrinsic::spv_resource_load_level: {
5967 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
5969 case Intrinsic::spv_resource_getdimensions_x:
5970 case Intrinsic::spv_resource_getdimensions_xy:
5971 case Intrinsic::spv_resource_getdimensions_xyz: {
5972 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
5974 case Intrinsic::spv_resource_getdimensions_levels_x:
5975 case Intrinsic::spv_resource_getdimensions_levels_xy:
5976 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
5977 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
5979 case Intrinsic::spv_resource_getdimensions_ms_xy:
5980 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
5981 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
5983 case Intrinsic::spv_resource_calculate_lod:
5984 case Intrinsic::spv_resource_calculate_lod_unclamped:
5985 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
5986 case Intrinsic::spv_resource_sample:
5987 case Intrinsic::spv_resource_sample_clamp:
5988 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
5989 case Intrinsic::spv_resource_samplebias:
5990 case Intrinsic::spv_resource_samplebias_clamp:
5991 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
5992 case Intrinsic::spv_resource_samplegrad:
5993 case Intrinsic::spv_resource_samplegrad_clamp:
5994 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
5995 case Intrinsic::spv_resource_samplelevel:
5996 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
5997 case Intrinsic::spv_resource_samplecmp:
5998 case Intrinsic::spv_resource_samplecmp_clamp:
5999 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
6000 case Intrinsic::spv_resource_samplecmplevelzero:
6001 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
6002 case Intrinsic::spv_resource_gather:
6003 case Intrinsic::spv_resource_gather_cmp:
6004 return selectGatherIntrinsic(ResVReg, ResType,
I);
6005 case Intrinsic::spv_resource_getbasepointer:
6006 case Intrinsic::spv_resource_getpointer: {
6007 return selectResourceGetPointer(ResVReg, ResType,
I);
6009 case Intrinsic::spv_pushconstant_getpointer: {
6010 return selectPushConstantGetPointer(ResVReg, ResType,
I);
6012 case Intrinsic::spv_discard: {
6013 return selectDiscard(ResVReg, ResType,
I);
6015 case Intrinsic::spv_resource_nonuniformindex: {
6016 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
6018 case Intrinsic::spv_unpackhalf2x16: {
6019 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
6021 case Intrinsic::spv_packhalf2x16: {
6022 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
6024 case Intrinsic::spv_ddx:
6025 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
6026 case Intrinsic::spv_ddy:
6027 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
6028 case Intrinsic::spv_ddx_coarse:
6029 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
6030 case Intrinsic::spv_ddy_coarse:
6031 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
6032 case Intrinsic::spv_ddx_fine:
6033 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
6034 case Intrinsic::spv_ddy_fine:
6035 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
6036 case Intrinsic::spv_fwidth:
6037 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
6038 case Intrinsic::spv_masked_gather:
6039 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
6040 return selectMaskedGather(ResVReg, ResType,
I);
6041 return diagnoseUnsupported(
6042 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
6043 case Intrinsic::spv_masked_scatter:
6044 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
6045 return selectMaskedScatter(
I);
6046 return diagnoseUnsupported(
6047 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
6048 case Intrinsic::returnaddress:
6049 case Intrinsic::frameaddress: {
6051 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
6058 return diagnoseUnsupported(
I,
"intrinsic selection not implemented.");
6063bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
6064 SPIRVTypeInst ResType,
6065 MachineInstr &
I)
const {
6070bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
6071 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6073 assert(Intr.getIntrinsicID() ==
6074 Intrinsic::spv_resource_counterhandlefrombinding);
6077 Register MainHandleReg = Intr.getOperand(2).getReg();
6079 assert(MainHandleDef->getIntrinsicID() ==
6080 Intrinsic::spv_resource_handlefrombinding);
6084 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(),
MRI);
6085 Register IndexReg = MainHandleDef->getOperand(5).getReg();
6086 std::string CounterName =
6091 MachineIRBuilder MIRBuilder(
I);
6093 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
6095 ArraySize, IndexReg, CounterName, MIRBuilder);
6097 return BuildCOPY(ResVReg, CounterVarReg,
I);
6100bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
6101 SPIRVTypeInst ResType,
6102 MachineInstr &
I)
const {
6104 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
6106 Register CounterHandleReg = Intr.getOperand(2).getReg();
6107 int64_t IncrVal = Intr.getOperand(3).getImm();
6114 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
6115 assert(CounterVarPointeeType &&
6116 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
6117 "Counter variable must be a struct");
6119 SPIRV::StorageClass::StorageBuffer &&
6120 "Counter variable must be in the storage buffer storage class");
6122 "Counter variable must have exactly 1 member in the struct");
6123 const SPIRVTypeInst MemberType =
6126 "Counter variable struct must have a single i32 member");
6130 MachineIRBuilder MIRBuilder(
I);
6132 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
6135 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
6141 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6144 .
addUse(CounterHandleReg)
6150 Register Scope = buildI32ConstantInEntryBlock(SPIRV::Scope::Device,
I);
6152 buildI32ConstantInEntryBlock(SPIRV::MemorySemantics::None,
I);
6155 buildI32ConstantInEntryBlock(
static_cast<uint32_t
>(IncrVal),
I);
6158 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
6167 return BuildCOPY(ResVReg, AtomicRes,
I);
6175 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
6183bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
6184 SPIRVTypeInst ResType,
6185 MachineInstr &
I)
const {
6193 Register ImageReg =
I.getOperand(2).getReg();
6201 Register IdxReg =
I.getOperand(3).getReg();
6203 MachineInstr &Pos =
I;
6205 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
6209bool SPIRVInstructionSelector::generateSampleImage(
6212 DebugLoc Loc, MachineInstr &Pos)
const {
6223 if (!loadHandleBeforePosition(NewSamplerReg,
6229 MachineIRBuilder MIRBuilder(Pos);
6242 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
6243 ImOps.Lod.has_value();
6244 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
6245 : SPIRV::OpImageSampleImplicitLod;
6247 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
6248 : SPIRV::OpImageSampleDrefImplicitLod;
6257 MIB.
addUse(*ImOps.Compare);
6259 uint32_t ImageOperands = 0;
6261 ImageOperands |= SPIRV::ImageOperand::Bias;
6263 ImageOperands |= SPIRV::ImageOperand::Lod;
6264 if (ImOps.GradX && ImOps.GradY)
6265 ImageOperands |= SPIRV::ImageOperand::Grad;
6266 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
6268 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6271 "Non-constant offsets are not supported in sample instructions.");
6276 ImageOperands |= SPIRV::ImageOperand::MinLod;
6278 if (ImageOperands != 0) {
6279 MIB.
addImm(ImageOperands);
6280 if (ImageOperands & SPIRV::ImageOperand::Bias)
6282 if (ImageOperands & SPIRV::ImageOperand::Lod)
6284 if (ImageOperands & SPIRV::ImageOperand::Grad) {
6285 MIB.
addUse(*ImOps.GradX);
6286 MIB.
addUse(*ImOps.GradY);
6289 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6290 MIB.
addUse(*ImOps.Offset);
6291 if (ImageOperands & SPIRV::ImageOperand::MinLod)
6292 MIB.
addUse(*ImOps.MinLod);
6299bool SPIRVInstructionSelector::selectImageQuerySize(
6301 std::optional<Register> LodReg)
const {
6304 "ImageReg is not an image type.");
6306 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6308 unsigned NumComponents = 0;
6310 case SPIRV::Dim::DIM_1D:
6311 case SPIRV::Dim::DIM_Buffer:
6312 NumComponents =
IsArray ? 2 : 1;
6314 case SPIRV::Dim::DIM_2D:
6315 case SPIRV::Dim::DIM_Cube:
6316 case SPIRV::Dim::DIM_Rect:
6317 NumComponents =
IsArray ? 3 : 2;
6319 case SPIRV::Dim::DIM_3D:
6323 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
6333 if (!LodReg && IsSampled && !IsMultisampled &&
6334 Dim != SPIRV::Dim::DIM_Buffer && Dim != SPIRV::Dim::DIM_Rect)
6337 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
6338 SPIRVTypeInst ResType =
6343 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6353bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
6354 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6355 Register ImageReg =
I.getOperand(2).getReg();
6362 return selectImageQuerySize(NewImageReg, ResVReg,
I);
6365bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
6366 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6367 Register ImageReg =
I.getOperand(2).getReg();
6376 Register LodReg =
I.getOperand(3).getReg();
6379 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
6381 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
6388 TII.get(SPIRV::OpImageQueryLevels))
6395 TII.get(SPIRV::OpCompositeConstruct))
6405bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
6406 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6407 Register ImageReg =
I.getOperand(2).getReg();
6418 "OpImageQuerySamples requires a multisampled image");
6420 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
6428 TII.get(SPIRV::OpImageQuerySamples))
6435 TII.get(SPIRV::OpCompositeConstruct))
6445bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
6446 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6447 Register ImageReg =
I.getOperand(2).getReg();
6448 Register SamplerReg =
I.getOperand(3).getReg();
6449 Register CoordinateReg =
I.getOperand(4).getReg();
6465 if (!loadHandleBeforePosition(
6470 MachineIRBuilder MIRBuilder(
I);
6476 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6486 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
6493 unsigned ExtractedIndex =
6495 Intrinsic::spv_resource_calculate_lod_unclamped
6499 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6500 TII.get(SPIRV::OpCompositeExtract))
6510bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
6511 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6512 Register ImageReg =
I.getOperand(2).getReg();
6513 Register SamplerReg =
I.getOperand(3).getReg();
6514 Register CoordinateReg =
I.getOperand(4).getReg();
6515 ImageOperands ImOps;
6516 if (
I.getNumOperands() > 5)
6517 ImOps.Offset =
I.getOperand(5).getReg();
6518 if (
I.getNumOperands() > 6)
6519 ImOps.MinLod =
I.getOperand(6).getReg();
6520 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6521 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6524bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
6525 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6526 Register ImageReg =
I.getOperand(2).getReg();
6527 Register SamplerReg =
I.getOperand(3).getReg();
6528 Register CoordinateReg =
I.getOperand(4).getReg();
6529 ImageOperands ImOps;
6530 ImOps.Bias =
I.getOperand(5).getReg();
6531 if (
I.getNumOperands() > 6)
6532 ImOps.Offset =
I.getOperand(6).getReg();
6533 if (
I.getNumOperands() > 7)
6534 ImOps.MinLod =
I.getOperand(7).getReg();
6535 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6536 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6539bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
6540 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6541 Register ImageReg =
I.getOperand(2).getReg();
6542 Register SamplerReg =
I.getOperand(3).getReg();
6543 Register CoordinateReg =
I.getOperand(4).getReg();
6544 ImageOperands ImOps;
6545 ImOps.GradX =
I.getOperand(5).getReg();
6546 ImOps.GradY =
I.getOperand(6).getReg();
6547 if (
I.getNumOperands() > 7)
6548 ImOps.Offset =
I.getOperand(7).getReg();
6549 if (
I.getNumOperands() > 8)
6550 ImOps.MinLod =
I.getOperand(8).getReg();
6551 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6552 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6555bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
6556 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6557 Register ImageReg =
I.getOperand(2).getReg();
6558 Register SamplerReg =
I.getOperand(3).getReg();
6559 Register CoordinateReg =
I.getOperand(4).getReg();
6560 ImageOperands ImOps;
6561 ImOps.Lod =
I.getOperand(5).getReg();
6562 if (
I.getNumOperands() > 6)
6563 ImOps.Offset =
I.getOperand(6).getReg();
6564 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6565 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6568bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
6569 SPIRVTypeInst ResType,
6570 MachineInstr &
I)
const {
6571 Register ImageReg =
I.getOperand(2).getReg();
6572 Register SamplerReg =
I.getOperand(3).getReg();
6573 Register CoordinateReg =
I.getOperand(4).getReg();
6574 ImageOperands ImOps;
6575 ImOps.Compare =
I.getOperand(5).getReg();
6576 if (
I.getNumOperands() > 6)
6577 ImOps.Offset =
I.getOperand(6).getReg();
6578 if (
I.getNumOperands() > 7)
6579 ImOps.MinLod =
I.getOperand(7).getReg();
6580 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6581 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6584bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
6585 SPIRVTypeInst ResType,
6586 MachineInstr &
I)
const {
6587 Register ImageReg =
I.getOperand(2).getReg();
6588 Register CoordinateReg =
I.getOperand(3).getReg();
6589 Register LodReg =
I.getOperand(4).getReg();
6591 ImageOperands ImOps;
6593 if (
I.getNumOperands() > 5)
6594 ImOps.Offset =
I.getOperand(5).getReg();
6606 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
6607 I.getDebugLoc(),
I, &ImOps);
6610bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
6611 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6612 Register ImageReg =
I.getOperand(2).getReg();
6613 Register SamplerReg =
I.getOperand(3).getReg();
6614 Register CoordinateReg =
I.getOperand(4).getReg();
6615 ImageOperands ImOps;
6616 ImOps.Compare =
I.getOperand(5).getReg();
6617 if (
I.getNumOperands() > 6)
6618 ImOps.Offset =
I.getOperand(6).getReg();
6621 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6622 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6625bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
6626 SPIRVTypeInst ResType,
6627 MachineInstr &
I)
const {
6628 Register ImageReg =
I.getOperand(2).getReg();
6629 Register SamplerReg =
I.getOperand(3).getReg();
6630 Register CoordinateReg =
I.getOperand(4).getReg();
6633 "ImageReg is not an image type.");
6638 ComponentOrCompareReg =
I.getOperand(5).getReg();
6639 OffsetReg =
I.getOperand(6).getReg();
6642 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
6646 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6647 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
6648 Dim != SPIRV::Dim::DIM_Rect) {
6650 "Gather operations are only supported for 2D, Cube, and Rect images.");
6657 if (!loadHandleBeforePosition(
6662 MachineIRBuilder MIRBuilder(
I);
6663 SPIRVTypeInst SampledImageType =
6668 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6676 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
6678 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
6680 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6685 .
addUse(ComponentOrCompareReg);
6687 uint32_t ImageOperands = 0;
6688 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
6689 if (Dim == SPIRV::Dim::DIM_Cube) {
6691 "Gather operations with offset are not supported for Cube images.");
6695 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6697 ImageOperands |= SPIRV::ImageOperand::Offset;
6701 if (ImageOperands != 0) {
6702 MIB.
addImm(ImageOperands);
6704 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6712bool SPIRVInstructionSelector::generateImageReadOrFetch(
6715 const ImageOperands *ImOps)
const {
6718 "ImageReg is not an image type.");
6720 bool IsSignedInteger =
6725 bool IsFetch = (SampledOp.getImm() == 1);
6727 auto AddOperands = [&](MachineInstrBuilder &MIB) {
6728 uint32_t ImageOperandsMask = 0;
6729 if (IsSignedInteger)
6730 ImageOperandsMask |= 0x1000;
6732 if (IsFetch && ImOps) {
6734 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
6735 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
6737 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
6739 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
6743 if (ImageOperandsMask != 0) {
6744 MIB.
addImm(ImageOperandsMask);
6745 if (IsFetch && ImOps) {
6748 if (ImOps->Offset &&
6749 (ImageOperandsMask &
6750 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
6751 MIB.
addUse(*ImOps->Offset);
6760 SPIRVTypeInst SampledType =
6763 SPIRVTypeInst ReadType =
6764 widenTypeToVec4(IsPacked ? SampledType : ResType, Pos);
6765 bool ReadTypeMatchesResult = ReadType == ResType;
6767 Register ReadReg = ReadTypeMatchesResult
6773 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
6779 BMI.constrainAllUses(
TII,
TRI, RBI);
6781 if (ReadTypeMatchesResult)
6794 if (ResultSize == 1) {
6803 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
6806bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
6807 SPIRVTypeInst ResType,
6808 MachineInstr &
I)
const {
6809 Register ResourcePtr =
I.getOperand(2).getReg();
6811 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
6824 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6830 if (
I.getNumExplicitOperands() > 3) {
6831 Register IndexReg =
I.getOperand(3).getReg();
6838bool SPIRVInstructionSelector::selectPushConstantGetPointer(
6839 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6844bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
6845 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6846 Register ObjReg =
I.getOperand(2).getReg();
6847 if (!BuildCOPY(ResVReg, ObjReg,
I))
6857 decorateUsesAsNonUniform(ResVReg);
6861void SPIRVInstructionSelector::decorateUsesAsNonUniform(
6864 {NonUniformReg,
nullptr}};
6865 llvm::SmallSet<Register, 8> Visited;
6866 while (WorkList.
size() > 0) {
6869 if (!Visited.
insert(CurrentReg).second)
6872 bool IsDecorated =
false;
6874 if (
Use.getOpcode() == SPIRV::OpDecorate &&
6875 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
6881 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
6883 if (ResultReg == CurrentReg)
6891 MachineInstr &InsertPt =
6894 SPIRV::Decoration::NonUniformEXT, {});
6899bool SPIRVInstructionSelector::extractSubvector(
6901 MachineInstr &InsertionPoint)
const {
6903 [[maybe_unused]]
uint64_t InputSize =
6906 [[maybe_unused]]
bool IsLongVectorEXT =
6908 assert((InputSize > 1 || IsLongVectorEXT) &&
"The input must be a vector.");
6909 assert((ResultSize > 1 || IsLongVectorEXT) &&
"The result must be a vector.");
6910 assert(ResultSize < InputSize &&
6911 "Cannot extract more element than there are in the input.");
6918 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6927 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
6929 TII.get(SPIRV::OpCompositeConstruct))
6933 for (
Register ComponentReg : ComponentRegisters)
6934 MIB.
addUse(ComponentReg);
6939bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
6940 MachineInstr &
I)
const {
6947 Register ImageReg =
I.getOperand(1).getReg();
6955 Register CoordinateReg =
I.getOperand(2).getReg();
6956 Register DataReg =
I.getOperand(3).getReg();
6959 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
6967Register SPIRVInstructionSelector::buildPointerToResource(
6968 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
6969 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
6970 StringRef Name, MachineIRBuilder MIRBuilder)
const {
6972 if (ArraySize == 1) {
6973 SPIRVTypeInst PtrType =
6976 "SpirvResType did not have an explicit layout.");
6981 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
6982 SPIRVTypeInst VarPointerType =
6985 VarPointerType, Set,
Binding, Name, MIRBuilder);
6987 SPIRVTypeInst ResPointerType =
7000bool SPIRVInstructionSelector::selectFirstBitSet16(
7001 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
7002 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
7004 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
7008 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
7011bool SPIRVInstructionSelector::selectFirstBitSet32(
7013 unsigned BitSetOpcode)
const {
7014 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7017 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
7024bool SPIRVInstructionSelector::selectFirstBitSet64(
7026 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
7040 if (ComponentCount > 2) {
7041 auto Func = [
this, SwapPrimarySide](
Register ResVReg, SPIRVTypeInst ResType,
7043 unsigned Opcode) ->
bool {
7044 return this->selectFirstBitSet64(ResVReg, ResType,
I, SrcReg, Opcode,
7048 return handle64BitOverflow(ResVReg, ResType,
I, SrcReg, BitSetOpcode, Func);
7052 MachineIRBuilder MIRBuilder(
I);
7054 BaseType, 2 * ComponentCount, MIRBuilder,
false);
7058 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
7064 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
7074 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
7075 SPIRV::OpVectorExtractDynamic))
7077 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
7078 SPIRV::OpVectorExtractDynamic))
7082 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
7083 TII.get(SPIRV::OpVectorShuffle))
7091 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
7097 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
7098 TII.get(SPIRV::OpVectorShuffle))
7106 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
7126 SelectOp = SPIRV::OpSelectSISCond;
7127 AddOp = SPIRV::OpIAddS;
7135 SelectOp = SPIRV::OpSelectVIVCond;
7136 AddOp = SPIRV::OpIAddV;
7142 Register RegSecondaryOffset = Reg0;
7146 if (SwapPrimarySide) {
7147 PrimaryReg = LowReg;
7148 SecondaryReg = HighReg;
7149 RegPrimaryOffset = Reg0;
7150 RegSecondaryOffset = Reg32;
7155 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
7156 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
7161 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
7162 SPIRV::OpINotEqual))
7169 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
7170 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
7175 if (SwapPrimarySide) {
7177 if (!selectOpWithSrcs(RegAdd, ResType,
I,
7178 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
7189 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
7190 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
7195 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
7196 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
7199 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
7203bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
7204 SPIRVTypeInst ResType,
7206 bool IsSigned)
const {
7208 Register OpReg =
I.getOperand(2).getReg();
7211 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
7212 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
7216 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
7218 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
7220 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
7223 return diagnoseUnsupported(
7225 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
7229bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
7230 SPIRVTypeInst ResType,
7231 MachineInstr &
I)
const {
7233 Register OpReg =
I.getOperand(2).getReg();
7238 unsigned ExtendOpcode = SPIRV::OpUConvert;
7239 unsigned BitSetOpcode = GL::FindILsb;
7243 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
7245 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
7247 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
7250 return diagnoseUnsupported(
I,
7251 "spv_firstbitlow only supports 16,32,64 bits.");
7255bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
7256 SPIRVTypeInst ResType,
7257 MachineInstr &
I)
const {
7262 bool UseUntypedPointers =
7263 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7264 unsigned Opcode = UseUntypedPointers
7265 ? SPIRV::OpUntypedVariableLengthArrayINTEL
7266 : SPIRV::OpVariableLengthArrayINTEL;
7268 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
7274 if (UseUntypedPointers) {
7277 "untyped variable length array result must have a recorded element "
7282 MIB.
addUse(
I.getOperand(2).getReg());
7286 unsigned Alignment =
I.getOperand(3).getImm();
7300 while (!Worklist.
empty()) {
7302 switch (
T->getOpcode()) {
7303 case SPIRV::OpTypeInt:
7304 case SPIRV::OpTypeFloat:
7305 case SPIRV::OpTypePointer:
7307 case SPIRV::OpTypeVector:
7308 case SPIRV::OpTypeVectorIdEXT:
7309 case SPIRV::OpTypeMatrix:
7310 case SPIRV::OpTypeArray: {
7311 Register OperandReg =
T->getOperand(1).getReg();
7315 case SPIRV::OpTypeStruct:
7316 for (
unsigned Idx = 1,
E =
T->getNumOperands(); Idx <
E; ++Idx) {
7317 Register OperandReg =
T->getOperand(Idx).getReg();
7332 Register TypeReg = Ty->getOperand(0).getReg();
7333 if (!Visited.
insert(TypeReg).second)
7336 switch (Ty->getOpcode()) {
7337 case SPIRV::OpTypePointer:
7338 if (Ty->getOperand(1).getImm() == SPIRV::StorageClass::StorageBuffer)
7342 case SPIRV::OpTypeArray:
7343 case SPIRV::OpTypeRuntimeArray:
7346 case SPIRV::OpTypeStruct:
7347 for (
unsigned I = 1;
I < Ty->getNumOperands(); ++
I)
7363bool SPIRVInstructionSelector::selectAbort(MachineInstr &
I)
const {
7364 assert(
I.getNumExplicitOperands() == 2);
7366 Register MsgReg =
I.getOperand(1).getReg();
7368 assert(MsgType &&
"Message argument of llvm.spv.abort has no SPIR-V type");
7371 return diagnoseUnsupported(
7373 "llvm.spv.abort message type must be a concrete SPIR-V type (numerical "
7374 "scalar, pointer, vector, matrix, or aggregate of such types)");
7377 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7384bool SPIRVInstructionSelector::selectTrap(MachineInstr &
I)
const {
7395 uint32_t MsgVal = ~0
u;
7396 if (
I.getOpcode() == TargetOpcode::G_UBSANTRAP)
7397 MsgVal =
static_cast<uint32_t
>(
I.getOperand(0).
getImm());
7400 Register MsgReg = buildI32ConstantInEntryBlock(MsgVal,
I, MsgType);
7403 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7410bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
7411 SPIRVTypeInst ResType,
7412 MachineInstr &
I)
const {
7419 bool UseUntypedPointers =
7420 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7422 UseUntypedPointers ? SPIRV::OpUntypedVariableKHR : SPIRV::OpVariable;
7425 MachineIRBuilder MIRBuilder(
I);
7428 .
addImm(SPIRV::Extension::SPV_KHR_variable_pointers);
7430 .
addImm(SPIRV::Capability::VariablePointersStorageBuffer);
7433 auto MIB =
BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(Opcode))
7436 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7440 if (UseUntypedPointers) {
7444 return diagnoseUnsupported(
7445 I,
"could not deduce the data type of an untyped variable");
7451 unsigned Alignment =
I.getOperand(2).getImm();
7458bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
7463 const MachineInstr *PrevI =
I.getPrevNode();
7465 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
7469 .
addMBB(
I.getOperand(0).getMBB())
7474 .
addMBB(
I.getOperand(0).getMBB())
7479bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
7490 const MachineInstr *NextI =
I.getNextNode();
7492 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
7498 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
7500 .
addUse(
I.getOperand(0).getReg())
7501 .
addMBB(
I.getOperand(1).getMBB())
7507bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
7508 MachineInstr &
I)
const {
7510 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
7512 const unsigned NumOps =
I.getNumOperands();
7513 for (
unsigned i = 1; i <
NumOps; i += 2) {
7514 MIB.
addUse(
I.getOperand(i + 0).getReg());
7515 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
7521bool SPIRVInstructionSelector::selectGlobalValue(
7522 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
7524 MachineIRBuilder MIRBuilder(
I);
7525 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
7528 std::string GlobalIdent;
7530 unsigned &
ID = UnnamedGlobalIDs[GV];
7532 ID = UnnamedGlobalIDs.
size();
7533 GlobalIdent =
"__unnamed_" + Twine(ID).str();
7559 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
7566 MachineRegisterInfo *
MRI = MIRBuilder.
getMRI();
7571 MachineInstrBuilder MIB1 =
7572 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7575 MachineInstrBuilder MIB2 =
7577 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
7581 GR.
add(ConstVal, MIB2);
7589 MachineInstrBuilder MIB3 =
7590 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7593 GR.
add(ConstVal, MIB3);
7599 assert(NewReg != ResVReg);
7600 return BuildCOPY(ResVReg, NewReg,
I);
7610 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
7613 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
7619 SPIRVTypeInst ResType =
7623 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
7628 if (
GlobalVar->isExternallyInitialized() &&
7629 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
7630 constexpr unsigned ReadWriteINTEL = 3u;
7633 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
7639bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
7640 SPIRVTypeInst ResType,
7641 MachineInstr &
I)
const {
7643 return selectExtInst(ResVReg, ResType,
I, CL::log10);
7651 MachineIRBuilder MIRBuilder(
I);
7656 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7659 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
7661 .
add(
I.getOperand(1))
7675 APFloat::rmNearestTiesToEven, &LosesInfo);
7680 isVectorType(ResType) ? SPIRV::OpVectorTimesScalar : SPIRV::OpFMulS;
7690bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
7691 SPIRVTypeInst ResType,
7692 MachineInstr &
I)
const {
7695 return selectExtInst(ResVReg, ResType,
I, CL::pown);
7701 Register ExpReg =
I.getOperand(2).getReg();
7703 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
7704 SPIRV::OpConvertSToF))
7706 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
7713bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
7714 SPIRVTypeInst ResType,
7715 MachineInstr &
I)
const {
7731 MachineIRBuilder MIRBuilder(
I);
7732 SPIRVTypeInst FloatType =
7736 FloatType, MIRBuilder, SPIRV::StorageClass::Function);
7749 MachineBasicBlock &EntryBB =
I.getMF()->
front();
7750 const bool IsUntyped =
7751 PtrType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7753 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
7754 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
7755 : SPIRV::OpVariable))
7758 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7766 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7769 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
7772 .
add(
I.getOperand(
I.getNumExplicitDefs()))
7776 Register IntegralPartReg =
I.getOperand(1).getReg();
7779 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7789 assert(
false &&
"GLSL::Modf is deprecated.");
7800bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
7801 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7802 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7803 MachineIRBuilder MIRBuilder(
I);
7804 const SPIRVTypeInst Vec3Ty =
7807 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
7819 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7823 MachineRegisterInfo *
MRI = MIRBuilder.
getMRI();
7829 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7836 assert(
I.getOperand(2).isReg());
7837 const uint32_t ThreadId =
foldImm(
I.getOperand(2),
MRI);
7841 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
7852bool SPIRVInstructionSelector::loadBuiltinInputID(
7853 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7854 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7855 MachineIRBuilder MIRBuilder(
I);
7857 ResType, MIRBuilder, SPIRV::StorageClass::Input);
7872 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7876 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7885SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
7886 MachineInstr &
I)
const {
7887 MachineIRBuilder MIRBuilder(
I);
7898bool SPIRVInstructionSelector::loadHandleBeforePosition(
7899 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
7900 MachineInstr &Pos)
const {
7903 Intrinsic::spv_resource_handlefrombinding);
7911 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
7912 MachineIRBuilder MIRBuilder(HandleDef);
7913 SPIRVTypeInst VarType = ResType;
7914 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
7916 if (IsStructuredBuffer) {
7925 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
7928 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
7929 ArraySize, IndexReg, Name, MIRBuilder);
7933 uint32_t LoadOpcode =
7934 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
7944bool SPIRVInstructionSelector::errorIfInstrOutsideShader(
7945 MachineInstr &
I)
const {
7947 return diagnoseUnsupported(
7948 I,
"this instruction is only supported in shaders.");
7953InstructionSelector *
7957 return new SPIRVInstructionSelector(TM,
Subtarget, RBI);
MachineInstrBuilder & UseMI
MachineInstrBuilder MachineInstrBuilder & DefMI
const AArch64Subtarget & Subtarget
#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
static MaybeAlign getAlign(Value *Ptr)
LLVMTypeRef LLVMIntType(unsigned NumBits)
const size_t AbstractManglingParser< Derived, Alloc >::NumOps
Loop::LoopBounds::Direction Direction
Register const TargetRegisterInfo * TRI
Promote Memory to Register
uint64_t IntrinsicInst * II
static StringRef getName(Value *V)
static unsigned getFCmpOpcode(CmpInst::Predicate Pred, unsigned Size)
MachineRegisterInfo & MRI
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 bool containsStorageBufferPointer(SPIRVTypeInst Ty, const SPIRVGlobalRegistry &GR, SmallSet< Register, 8 > &Visited)
static bool isConstReg(MachineRegisterInfo *MRI, MachineInstr *OpDef)
static unsigned getPtrCmpOpcode(unsigned Pred)
static void addMemoryOperands(const MachineMemOperand *MemOp, MachineInstrBuilder &MIB, MachineIRBuilder &MIRBuilder, SPIRVGlobalRegistry &GR, std::optional< Align > AlignOverride=std::nullopt)
bool isDead(const MachineInstr &MI, const MachineRegisterInfo &MRI)
static Register convertPtrToInt(Register Reg, LLT ConvTy, SPIRVTypeInst SpvType, LegalizerHelper &Helper, MachineRegisterInfo &MRI, SPIRVGlobalRegistry *GR)
#define SPIRV_BACKEND_SERVICE_FUN_NAME
BaseType
A given derived pointer can have multiple base pointers through phi/selects.
This file defines the SmallSet class.
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 & BFloat()
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.
static APInt getSignMask(unsigned BitWidth)
Get the SignMask for a specific bit width.
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.
static LLVM_ABI std::optional< GFConstant > getConstant(Register Const, const MachineRegisterInfo &MRI)
static LLVM_ABI std::optional< GIConstant > getConstant(Register Const, const MachineRegisterInfo &MRI)
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.
instr_iterator instr_end()
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 substituteRegister(Register FromReg, Register ToReg, unsigned SubIdx, const TargetRegisterInfo &RegInfo)
Replace all occurrences of FromReg with ToReg:SubIdx, properly composing subreg indices where necessa...
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.
LLVM_ABI Align getAlign() const
Return the minimum known alignment in bytes of the actual memory reference.
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 LLVM_READONLY 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.
bool hasOneUse(Register RegNo) const
hasOneUse - Return true if there is exactly one instruction using the specified 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.
const SDValue & getOperand(unsigned i) const
bool isScalarOrVectorSigned(SPIRVTypeInst Type) const
SPIRVTypeInst getOrCreateOpTypeSampledImage(SPIRVTypeInst ImageType, MachineIRBuilder &MIRBuilder)
void assignSPIRVTypeToVReg(SPIRVTypeInst Type, Register VReg, const MachineFunction &MF)
SPIRVTypeInst getOrCreateSPIRVPointerType(const Type *BaseType, MachineIRBuilder &MIRBuilder, SPIRV::StorageClass::StorageClass SC, bool ForceTyped=false)
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
SPIRVTypeInst getOpTypeVoid(MachineIRBuilder &MIRBuilder)
void invalidateMachineInstr(MachineInstr *MI)
SPIRVTypeInst getOrCreateSPIRVBoolType(MachineIRBuilder &MIRBuilder, bool EmitIR)
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
SPIRVTypeInst getUntypedPtrElementType(Register Reg) 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 isAnyTypeFloat() 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.
SmallSet - This maintains a set of unique values, optimizing for the case when the set is small (less...
std::pair< const_iterator, bool > insert(const T &V)
insert - Insert an element into the set if it isn't already there.
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.
Represent a constant reference to a string, i.e.
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
@ BFloatTyID
16-bit floating point type (7-bit significand)
@ 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.
An efficient, type-erasing, non-owning reference to a callable.
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 Align[]
Key for Kernel::Arg::Metadata::mAlign.
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.
ElementType
The element type of an SRV or UAV resource.
constexpr uint64_t PointerSize
aarch64 pointer size.
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.
unsigned getOpcode(const VPValue *V)
Return the instruction opcode for the recipe defining V or 0 for unsupported recipes and VPValues not...
This is an optimization pass for GlobalISel generic memory operations.
bool all_of(R &&range, UnaryPredicate P)
Provide wrappers to std::all_of which take ranges instead of having to pass begin/end explicitly.
void addStringImm(StringRef Str, MCInst &Inst)
MachineBasicBlock::iterator getOpVariableMBBIt(MachineFunction &MF)
MachineInstrBuilder BuildMI(MachineFunction &MF, const MIMetadata &MIMD, const MCInstrDesc &MCID)
Builder interface. Specify how to create the initial instruction itself.
bool isTypeFoldingSupported(unsigned Opcode)
uint32_t getMemSemanticsWithStorageClass(const Triple &TT, uint32_t OrderSem, uint32_t StorageClassSem)
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.
iterator_range< T > make_range(T x, T y)
Convenience function for iterating over sub-ranges.
Register createVirtualRegister(SPIRVTypeInst SpvType, SPIRVGlobalRegistry *GR, MachineRegisterInfo *MRI, const MachineFunction &MF)
unsigned getArrayComponentCount(const MachineRegisterInfo *MRI, const MachineInstr *ResType)
iterator_range< early_inc_iterator_impl< detail::IterOfRange< RangeT > > > make_early_inc_range(RangeT &&Range)
Make a range that does early increment to allow mutation of the underlying range without disrupting i...
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...
SPIRV::Scope::Scope getMemScope(const Triple &TT, LLVMContext &Ctx, SyncScope::ID Id)
uint64_t getIConstVal(Register ConstReg, const MachineRegisterInfo *MRI)
SmallVector< MachineInstr *, 4 > createContinuedInstructions(MachineIRBuilder &MIRBuilder, unsigned Opcode, unsigned MinWC, unsigned ContinuedOpcode, ArrayRef< Register > Args, Register ReturnRegister, Register TypeID)
SPIRV::MemorySemantics::MemorySemantics getMemSemanticsForStorageClass(SPIRV::StorageClass::StorageClass SC)
constexpr unsigned storageClassToAddressSpace(SPIRV::StorageClass::StorageClass SC)
bool isVectorType(SPIRVTypeInst SPVTy)
RelativeUniformCounterPtr ValuesPtrExpr VTableAddr Value
void buildOpName(Register Target, StringRef Name, MachineIRBuilder &MIRBuilder)
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.
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)
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...
MCRegisterClass TargetRegisterClass