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;
335 bool IsSigned)
const;
337 bool IsSigned,
unsigned Opcode)
const;
339 bool IsSigned)
const;
345 bool IsSigned)
const;
386 GL::GLSLExtInst GLInst,
bool setMIFlags =
true,
387 bool useMISrc =
true,
389 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
390 CL::OpenCLExtInst CLInst,
bool setMIFlags =
true,
391 bool useMISrc =
true,
393 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
394 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
395 bool setMIFlags =
true,
bool useMISrc =
true,
397 bool selectExtInst(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
398 const ExtInstList &ExtInsts,
bool setMIFlags =
true,
399 bool useMISrc =
true,
402 bool selectLog10(
Register ResVReg, SPIRVTypeInst ResType,
403 MachineInstr &
I)
const;
405 bool selectFpowi(
Register ResVReg, SPIRVTypeInst ResType,
406 MachineInstr &
I)
const;
408 bool selectSaturate(
Register ResVReg, SPIRVTypeInst ResType,
409 MachineInstr &
I)
const;
411 bool selectWaveOpInst(
Register ResVReg, SPIRVTypeInst ResType,
412 MachineInstr &
I,
unsigned Opcode)
const;
414 bool selectBarrierInst(MachineInstr &
I,
unsigned Scope,
unsigned MemSem,
415 bool WithGroupSync)
const;
417 bool selectWaveActiveCountBits(
Register ResVReg, SPIRVTypeInst ResType,
418 MachineInstr &
I)
const;
420 bool selectWaveActiveAllEqual(
Register ResVReg, SPIRVTypeInst ResType,
421 MachineInstr &
I)
const;
425 bool selectHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
426 MachineInstr &
I)
const;
428 bool selectCounterHandleFromBinding(
Register &ResVReg, SPIRVTypeInst ResType,
429 MachineInstr &
I)
const;
431 bool selectReadImageIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
432 MachineInstr &
I)
const;
433 bool selectGetDimensionsIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
434 MachineInstr &
I)
const;
435 bool selectGetDimensionsLevelsIntrinsic(
Register &ResVReg,
436 SPIRVTypeInst ResType,
437 MachineInstr &
I)
const;
438 bool selectGetDimensionsMSIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
439 MachineInstr &
I)
const;
442 std::optional<Register> LodReg = std::nullopt)
const;
443 bool selectSampleBasicIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
444 MachineInstr &
I)
const;
445 bool selectCalculateLodIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
446 MachineInstr &
I)
const;
447 bool selectSampleBiasIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
448 MachineInstr &
I)
const;
449 bool selectSampleGradIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
450 MachineInstr &
I)
const;
451 bool selectSampleLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
452 MachineInstr &
I)
const;
453 bool selectLoadLevelIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
454 MachineInstr &
I)
const;
455 bool selectSampleCmpIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
456 MachineInstr &
I)
const;
457 bool selectSampleCmpLevelZeroIntrinsic(
Register &ResVReg,
458 SPIRVTypeInst ResType,
459 MachineInstr &
I)
const;
460 bool selectGatherIntrinsic(
Register &ResVReg, SPIRVTypeInst ResType,
461 MachineInstr &
I)
const;
462 bool selectImageWriteIntrinsic(MachineInstr &
I)
const;
463 bool selectResourceGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
464 MachineInstr &
I)
const;
465 bool selectPushConstantGetPointer(
Register &ResVReg, SPIRVTypeInst ResType,
466 MachineInstr &
I)
const;
467 bool selectResourceNonUniformIndex(
Register &ResVReg, SPIRVTypeInst ResType,
468 MachineInstr &
I)
const;
469 bool selectModf(
Register ResVReg, SPIRVTypeInst ResType,
470 MachineInstr &
I)
const;
471 bool selectUpdateCounter(
Register &ResVReg, SPIRVTypeInst ResType,
472 MachineInstr &
I)
const;
473 bool selectFrexp(
Register ResVReg, SPIRVTypeInst ResType,
474 MachineInstr &
I)
const;
476 bool selectCopySign(
Register ResVReg, SPIRVTypeInst ResType,
477 MachineInstr &
I)
const;
479 bool selectLdexp(
Register ResVReg, SPIRVTypeInst ResType,
480 MachineInstr &
I)
const;
481 bool selectSincos(
Register ResVReg, SPIRVTypeInst ResType,
482 MachineInstr &
I)
const;
483 bool selectExp10(
Register ResVReg, SPIRVTypeInst ResType,
484 MachineInstr &
I)
const;
485 bool selectDerivativeInst(
Register ResVReg, SPIRVTypeInst ResType,
486 MachineInstr &
I,
const unsigned DPdOpCode)
const;
488 Register buildI32ConstantInEntryBlock(uint32_t Val, MachineInstr &
I,
489 SPIRVTypeInst ResType =
nullptr)
const;
491 Register buildZerosVal(SPIRVTypeInst ResType, MachineInstr &
I)
const;
492 bool isScalarOrVectorIntConstantZero(
Register Reg)
const;
493 Register buildZerosValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
495 MachineInstr &
I)
const;
496 Register buildOnesValF(SPIRVTypeInst ResType, MachineInstr &
I)
const;
498 MachineInstr &
I)
const;
500 bool wrapIntoSpecConstantOp(MachineInstr &
I,
503 SPIRVTypeInst getUcharPtrType(MachineInstr &
I,
504 SPIRV::StorageClass::StorageClass SC)
const;
505 Register getUcharPtrTypeReg(MachineInstr &
I,
506 SPIRV::StorageClass::StorageClass SC)
const;
507 MachineInstrBuilder buildSpecConstantOp(MachineInstr &
I,
Register Dest,
509 uint32_t Opcode)
const;
510 MachineInstrBuilder buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
511 SPIRVTypeInst SrcPtrTy)
const;
512 Register buildPointerToResource(SPIRVTypeInst ResType,
513 SPIRV::StorageClass::StorageClass SC,
514 uint32_t Set, uint32_t
Binding,
515 uint32_t ArraySize,
Register IndexReg,
517 MachineIRBuilder MIRBuilder)
const;
518 SPIRVTypeInst widenTypeToVec4(SPIRVTypeInst
Type, MachineInstr &
I)
const;
519 bool extractSubvector(
Register &ResVReg, SPIRVTypeInst ResType,
520 Register &ReadReg, MachineInstr &InsertionPoint)
const;
521 bool generateImageReadOrFetch(
Register &ResVReg, SPIRVTypeInst ResType,
524 const ImageOperands *ImOps =
nullptr)
const;
525 bool generateSampleImage(
Register ResVReg, SPIRVTypeInst ResType,
527 Register CoordinateReg,
const ImageOperands &ImOps,
530 bool loadVec3BuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
531 Register ResVReg, SPIRVTypeInst ResType,
532 MachineInstr &
I)
const;
533 bool loadBuiltinInputID(SPIRV::BuiltIn::BuiltIn BuiltInValue,
534 Register ResVReg, SPIRVTypeInst ResType,
535 MachineInstr &
I)
const;
536 bool loadHandleBeforePosition(
Register &HandleReg, SPIRVTypeInst ResType,
537 GIntrinsic &HandleDef, MachineInstr &Pos)
const;
538 void decorateUsesAsNonUniform(
Register &NonUniformReg)
const;
539 bool errorIfInstrOutsideShader(MachineInstr &
I)
const;
542 handle64BitOverflow(
Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
543 Register SrcReg,
unsigned int Opcode,
544 std::function<
bool(
Register, SPIRVTypeInst,
545 MachineInstr &,
Register,
unsigned)>
549bool sampledTypeIsSignedInteger(
const llvm::Type *HandleType) {
551 if (
TET->getTargetExtName() ==
"spirv.Image") {
554 assert(
TET->getTargetExtName() ==
"spirv.SignedImage");
555 return TET->getTypeParameter(0)->isIntegerTy();
559#define GET_GLOBALISEL_IMPL
560#include "SPIRVGenGlobalISel.inc"
561#undef GET_GLOBALISEL_IMPL
567 TRI(*ST.getRegisterInfo()), RBI(RBI), GR(*ST.getSPIRVGlobalRegistry()),
570#include
"SPIRVGenGlobalISel.inc"
573#include
"SPIRVGenGlobalISel.inc"
585 InstructionSelector::setupMF(MF, VT, CoverageInfo, PSI, BFI);
590 if (HasVRegsReset == &MF)
605 for (
const auto &
MBB : MF) {
606 for (
const auto &
MI :
MBB) {
609 if (
MI.getOpcode() != SPIRV::ASSIGN_TYPE)
616 if (DstType != SrcType)
621 if (DstRC != SrcRC && SrcRC)
633 while (!Stack.empty()) {
638 switch (
MI->getOpcode()) {
639 case TargetOpcode::G_INTRINSIC:
640 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
641 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS: {
644 if (IntrID != Intrinsic::spv_const_composite &&
645 IntrID != Intrinsic::spv_undef && IntrID != Intrinsic::spv_poison)
649 case TargetOpcode::G_BUILD_VECTOR:
650 case TargetOpcode::G_SPLAT_VECTOR:
652 i < OpDef->getNumOperands(); i++) {
657 Stack.push_back(OpNestedDef);
660 case TargetOpcode::G_CONSTANT:
661 case TargetOpcode::G_FCONSTANT:
662 case TargetOpcode::G_IMPLICIT_DEF:
663 case SPIRV::OpConstantTrue:
664 case SPIRV::OpConstantFalse:
665 case SPIRV::OpConstantI:
666 case SPIRV::OpConstantF:
667 case SPIRV::OpConstantComposite:
668 case SPIRV::OpConstantCompositeContinuedINTEL:
669 case SPIRV::OpConstantSampler:
670 case SPIRV::OpConstantNull:
672 case SPIRV::OpPoisonKHR:
673 case SPIRV::OpConstantFunctionPointerINTEL:
700 case Intrinsic::spv_all:
701 case Intrinsic::spv_alloca:
702 case Intrinsic::spv_any:
703 case Intrinsic::spv_bitcast:
704 case Intrinsic::spv_const_composite:
705 case Intrinsic::spv_degrees:
706 case Intrinsic::spv_distance:
707 case Intrinsic::spv_extractelt:
708 case Intrinsic::spv_extractv:
709 case Intrinsic::spv_faceforward:
710 case Intrinsic::spv_fdot:
711 case Intrinsic::spv_firstbitlow:
712 case Intrinsic::spv_firstbitshigh:
713 case Intrinsic::spv_firstbituhigh:
714 case Intrinsic::spv_frac:
715 case Intrinsic::spv_gep:
716 case Intrinsic::spv_global_offset:
717 case Intrinsic::spv_global_size:
718 case Intrinsic::spv_group_id:
719 case Intrinsic::spv_insertelt:
720 case Intrinsic::spv_insertv:
721 case Intrinsic::spv_isinf:
722 case Intrinsic::spv_isnan:
723 case Intrinsic::spv_isfinite:
724 case Intrinsic::spv_isnormal:
725 case Intrinsic::spv_lerp:
726 case Intrinsic::spv_length:
727 case Intrinsic::spv_normalize:
728 case Intrinsic::spv_num_subgroups:
729 case Intrinsic::spv_num_workgroups:
730 case Intrinsic::spv_ptrcast:
731 case Intrinsic::spv_radians:
732 case Intrinsic::spv_reflect:
733 case Intrinsic::spv_refract:
734 case Intrinsic::spv_resource_getbasepointer:
735 case Intrinsic::spv_resource_getpointer:
736 case Intrinsic::spv_resource_handlefrombinding:
737 case Intrinsic::spv_resource_handlefromimplicitbinding:
738 case Intrinsic::spv_resource_nonuniformindex:
739 case Intrinsic::spv_resource_sample:
740 case Intrinsic::spv_rsqrt:
741 case Intrinsic::spv_saturate:
742 case Intrinsic::spv_sdot:
743 case Intrinsic::spv_sign:
744 case Intrinsic::spv_smoothstep:
745 case Intrinsic::spv_subgroup_id:
746 case Intrinsic::spv_subgroup_local_invocation_id:
747 case Intrinsic::spv_subgroup_max_size:
748 case Intrinsic::spv_subgroup_size:
749 case Intrinsic::spv_thread_id:
750 case Intrinsic::spv_thread_id_in_group:
751 case Intrinsic::spv_udot:
752 case Intrinsic::spv_undef:
753 case Intrinsic::spv_value_md:
754 case Intrinsic::spv_workgroup_size:
766 case SPIRV::OpTypeVoid:
767 case SPIRV::OpTypeBool:
768 case SPIRV::OpTypeInt:
769 case SPIRV::OpTypeFloat:
770 case SPIRV::OpTypeVector:
771 case SPIRV::OpTypeVectorIdEXT:
772 case SPIRV::OpTypeMatrix:
773 case SPIRV::OpTypeImage:
774 case SPIRV::OpTypeSampler:
775 case SPIRV::OpTypeSampledImage:
776 case SPIRV::OpTypeArray:
777 case SPIRV::OpTypeRuntimeArray:
778 case SPIRV::OpTypeStruct:
779 case SPIRV::OpTypeOpaque:
780 case SPIRV::OpTypePointer:
781 case SPIRV::OpTypeFunction:
782 case SPIRV::OpTypeEvent:
783 case SPIRV::OpTypeDeviceEvent:
784 case SPIRV::OpTypeReserveId:
785 case SPIRV::OpTypeQueue:
786 case SPIRV::OpTypePipe:
787 case SPIRV::OpTypeForwardPointer:
788 case SPIRV::OpTypePipeStorage:
789 case SPIRV::OpTypeNamedBarrier:
790 case SPIRV::OpTypeAccelerationStructureNV:
791 case SPIRV::OpTypeCooperativeMatrixNV:
792 case SPIRV::OpTypeCooperativeMatrixKHR:
802 if (
MI.getNumDefs() == 0)
805 for (
const auto &MO :
MI.all_defs()) {
807 if (
Reg.isPhysical()) {
811 for (
const auto &
UseMI :
MRI.use_nodbg_instructions(
Reg)) {
812 if (
UseMI.getOpcode() != SPIRV::OpName) {
819 if (
MI.getOpcode() == TargetOpcode::LOCAL_ESCAPE ||
MI.isFakeUse() ||
820 MI.isLifetimeMarker()) {
823 <<
"Not dead: Opcode is LOCAL_ESCAPE, fake use, or lifetime marker.\n");
834 if (
MI.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
835 MI.getOpcode() == TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS) {
838 LLVM_DEBUG(
dbgs() <<
"Dead: Intrinsic with no real side effects.\n");
843 if (
MI.mayStore() ||
MI.isCall() ||
844 (
MI.mayLoad() &&
MI.hasOrderedMemoryRef()) ||
MI.isPosition() ||
845 MI.isDebugInstr() ||
MI.isTerminator() ||
MI.isJumpTableDebugInfo()) {
846 LLVM_DEBUG(
dbgs() <<
"Not dead: instruction has side effects.\n");
857 LLVM_DEBUG(
dbgs() <<
"Dead: known opcode with no side effects\n");
864void SPIRVInstructionSelector::removeOpNamesForDeadMI(MachineInstr &
MI)
const {
866 for (
const auto &MO :
MI.all_defs()) {
870 SmallVector<MachineInstr *, 4> UselessOpNames;
873 "There is still a use of the dead function.");
876 for (MachineInstr *OpNameMI : UselessOpNames) {
878 OpNameMI->eraseFromParent();
883void SPIRVInstructionSelector::removeDeadInstruction(MachineInstr &
MI)
const {
886 removeOpNamesForDeadMI(
MI);
887 MI.eraseFromParent();
890bool SPIRVInstructionSelector::select(MachineInstr &
I) {
891 resetVRegsType(*
I.getParent()->getParent());
893 assert(
I.getParent() &&
"Instruction should be in a basic block!");
894 assert(
I.getParent()->getParent() &&
"Instruction should be in a function!");
899 removeDeadInstruction(
I);
906 if (Opcode == SPIRV::ASSIGN_TYPE) {
907 Register DstReg =
I.getOperand(0).getReg();
908 Register SrcReg =
I.getOperand(1).getReg();
911 Def->getOpcode() != TargetOpcode::G_CONSTANT &&
912 Def->getOpcode() != TargetOpcode::G_FCONSTANT) {
913 if (
Def->getOpcode() == TargetOpcode::G_SELECT) {
914 Register SelectDstReg =
Def->getOperand(0).getReg();
915 bool SuccessToSelectSelect [[maybe_unused]] = selectSelect(
917 assert(SuccessToSelectSelect);
919 Def->eraseFromParent();
926 bool Res = selectImpl(
I, *CoverageInfo);
928 if (!Res &&
Def->getOpcode() != TargetOpcode::G_CONSTANT) {
929 dbgs() <<
"Unexpected pattern in ASSIGN_TYPE.\nInstruction: ";
933 assert(Res ||
Def->getOpcode() == TargetOpcode::G_CONSTANT);
945 }
else if (
I.getNumDefs() == 1) {
957 removeDeadInstruction(
I);
962 if (
I.getNumOperands() !=
I.getNumExplicitOperands()) {
963 LLVM_DEBUG(
errs() <<
"Generic instr has unexpected implicit operands\n");
969 bool HasDefs =
I.getNumDefs() > 0;
972 assert(!HasDefs || ResType ||
I.getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
973 I.getOpcode() == TargetOpcode::G_IMPLICIT_DEF);
974 if (spvSelect(ResVReg, ResType,
I)) {
976 for (
unsigned i = 0; i <
I.getNumDefs(); ++i)
987 case TargetOpcode::G_CONSTANT:
988 case TargetOpcode::G_FCONSTANT:
995 MachineInstr &
I)
const {
998 if (DstRC != SrcRC && SrcRC)
1000 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::COPY))
1007bool SPIRVInstructionSelector::spvSelect(
Register ResVReg,
1008 SPIRVTypeInst ResType,
1009 MachineInstr &
I)
const {
1010 const unsigned Opcode =
I.getOpcode();
1012 return selectImpl(
I, *CoverageInfo);
1014 case TargetOpcode::G_CONSTANT:
1015 case TargetOpcode::G_FCONSTANT:
1016 return selectConst(ResVReg, ResType,
I);
1017 case TargetOpcode::G_GLOBAL_VALUE:
1018 return selectGlobalValue(ResVReg,
I);
1019 case TargetOpcode::G_IMPLICIT_DEF:
1020 return selectOpUndef(ResVReg, ResType,
I);
1021 case TargetOpcode::G_FREEZE:
1022 return selectFreeze(ResVReg, ResType,
I);
1024 case TargetOpcode::G_INTRINSIC:
1025 case TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS:
1026 case TargetOpcode::G_INTRINSIC_CONVERGENT:
1027 case TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS:
1028 return selectIntrinsic(ResVReg, ResType,
I);
1029 case TargetOpcode::G_BITREVERSE:
1030 return selectBitreverse(ResVReg, ResType,
I);
1032 case TargetOpcode::G_BUILD_VECTOR:
1033 return selectBuildVector(ResVReg, ResType,
I);
1034 case TargetOpcode::G_SPLAT_VECTOR:
1035 return selectSplatVector(ResVReg, ResType,
I);
1036 case TargetOpcode::G_CONCAT_VECTORS:
1037 return selectConcatVectors(ResVReg, ResType,
I);
1039 case TargetOpcode::G_SHUFFLE_VECTOR: {
1040 MachineBasicBlock &BB = *
I.getParent();
1041 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
1044 .
addUse(
I.getOperand(1).getReg())
1045 .
addUse(
I.getOperand(2).getReg());
1046 for (
auto V :
I.getOperand(3).getShuffleMask())
1051 case TargetOpcode::G_MEMMOVE:
1052 case TargetOpcode::G_MEMCPY:
1053 case TargetOpcode::G_MEMCPY_INLINE:
1054 case TargetOpcode::G_MEMSET:
1055 case TargetOpcode::G_MEMSET_INLINE:
1056 return selectMemOperation(ResVReg,
I);
1058 case TargetOpcode::G_ICMP:
1059 return selectICmp(ResVReg, ResType,
I);
1060 case TargetOpcode::G_FCMP:
1061 return selectFCmp(ResVReg, ResType,
I);
1063 case TargetOpcode::G_FRAME_INDEX:
1064 return selectFrameIndex(ResVReg, ResType,
I);
1066 case TargetOpcode::G_LOAD:
1067 return selectLoad(ResVReg, ResType,
I);
1068 case TargetOpcode::G_STORE:
1069 return selectStore(
I);
1071 case TargetOpcode::G_BR:
1072 return selectBranch(
I);
1073 case TargetOpcode::G_BRCOND:
1074 return selectBranchCond(
I);
1076 case TargetOpcode::G_PHI:
1077 return selectPhi(ResVReg,
I);
1079 case TargetOpcode::G_FPTOSI:
1080 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1081 case TargetOpcode::G_FPTOUI:
1082 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1084 case TargetOpcode::G_FPTOSI_SAT:
1085 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToS);
1086 case TargetOpcode::G_FPTOUI_SAT:
1087 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertFToU);
1089 case TargetOpcode::G_SITOFP:
1090 return selectIToF(ResVReg, ResType,
I,
true, SPIRV::OpConvertSToF);
1091 case TargetOpcode::G_UITOFP:
1092 return selectIToF(ResVReg, ResType,
I,
false, SPIRV::OpConvertUToF);
1094 case TargetOpcode::G_CTPOP:
1095 return selectPopCount(ResVReg, ResType,
I, SPIRV::OpBitCount);
1096 case TargetOpcode::G_SMIN:
1097 return selectExtInst(ResVReg, ResType,
I, CL::s_min, GL::SMin);
1098 case TargetOpcode::G_UMIN:
1099 return selectExtInst(ResVReg, ResType,
I, CL::u_min, GL::UMin);
1101 case TargetOpcode::G_SMAX:
1102 return selectExtInst(ResVReg, ResType,
I, CL::s_max, GL::SMax);
1103 case TargetOpcode::G_UMAX:
1104 return selectExtInst(ResVReg, ResType,
I, CL::u_max, GL::UMax);
1106 case TargetOpcode::G_SCMP:
1107 return selectSUCmp(ResVReg, ResType,
I,
true);
1108 case TargetOpcode::G_UCMP:
1109 return selectSUCmp(ResVReg, ResType,
I,
false);
1110 case TargetOpcode::G_LROUND:
1111 case TargetOpcode::G_LLROUND: {
1116 regForLround, *(
I.getParent()->getParent()));
1118 CL::round, GL::Round,
false);
1120 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConvertFToS))
1127 case TargetOpcode::G_STRICT_FMA:
1128 case TargetOpcode::G_FMA: {
1131 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFmaKHR))
1134 .
addUse(
I.getOperand(1).getReg())
1135 .
addUse(
I.getOperand(2).getReg())
1136 .
addUse(
I.getOperand(3).getReg())
1141 return selectExtInst(ResVReg, ResType,
I, CL::fma, GL::Fma);
1144 case TargetOpcode::G_FLDEXP:
1145 case TargetOpcode::G_STRICT_FLDEXP:
1146 return selectLdexp(ResVReg, ResType,
I);
1148 case TargetOpcode::G_FPOW:
1149 return selectExtInst(ResVReg, ResType,
I, CL::pow, GL::Pow);
1150 case TargetOpcode::G_FPOWI:
1151 return selectFpowi(ResVReg, ResType,
I);
1153 case TargetOpcode::G_FEXP:
1154 return selectExtInst(ResVReg, ResType,
I, CL::exp, GL::Exp);
1155 case TargetOpcode::G_FEXP2:
1156 return selectExtInst(ResVReg, ResType,
I, CL::exp2, GL::Exp2);
1157 case TargetOpcode::G_FEXP10:
1158 return selectExp10(ResVReg, ResType,
I);
1160 case TargetOpcode::G_FMODF:
1161 return selectModf(ResVReg, ResType,
I);
1162 case TargetOpcode::G_FSINCOS:
1163 return selectSincos(ResVReg, ResType,
I);
1165 case TargetOpcode::G_FLOG:
1166 return selectExtInst(ResVReg, ResType,
I, CL::log, GL::Log);
1167 case TargetOpcode::G_FLOG2:
1168 return selectExtInst(ResVReg, ResType,
I, CL::log2, GL::Log2);
1169 case TargetOpcode::G_FLOG10:
1170 return selectLog10(ResVReg, ResType,
I);
1172 case TargetOpcode::G_FABS:
1173 return selectExtInst(ResVReg, ResType,
I, CL::fabs, GL::FAbs);
1174 case TargetOpcode::G_ABS:
1175 return selectExtInst(ResVReg, ResType,
I, CL::s_abs, GL::SAbs);
1177 case TargetOpcode::G_FMINNUM:
1178 return selectExtInst(ResVReg, ResType,
I, CL::fmin, GL::NMin);
1179 case TargetOpcode::G_FMAXNUM:
1180 return selectExtInst(ResVReg, ResType,
I, CL::fmax, GL::NMax);
1182 case TargetOpcode::G_FCOPYSIGN:
1183 return selectCopySign(ResVReg, ResType,
I);
1185 case TargetOpcode::G_FCEIL:
1186 return selectExtInst(ResVReg, ResType,
I, CL::ceil, GL::Ceil);
1187 case TargetOpcode::G_FFLOOR:
1188 return selectExtInst(ResVReg, ResType,
I, CL::floor, GL::Floor);
1190 case TargetOpcode::G_FCOS:
1191 return selectExtInst(ResVReg, ResType,
I, CL::cos, GL::Cos);
1192 case TargetOpcode::G_FSIN:
1193 return selectExtInst(ResVReg, ResType,
I, CL::sin, GL::Sin);
1194 case TargetOpcode::G_FTAN:
1195 return selectExtInst(ResVReg, ResType,
I, CL::tan, GL::Tan);
1196 case TargetOpcode::G_FACOS:
1197 return selectExtInst(ResVReg, ResType,
I, CL::acos, GL::Acos);
1198 case TargetOpcode::G_FASIN:
1199 return selectExtInst(ResVReg, ResType,
I, CL::asin, GL::Asin);
1200 case TargetOpcode::G_FATAN:
1201 return selectExtInst(ResVReg, ResType,
I, CL::atan, GL::Atan);
1202 case TargetOpcode::G_FATAN2:
1203 return selectExtInst(ResVReg, ResType,
I, CL::atan2, GL::Atan2);
1204 case TargetOpcode::G_FCOSH:
1205 return selectExtInst(ResVReg, ResType,
I, CL::cosh, GL::Cosh);
1206 case TargetOpcode::G_FSINH:
1207 return selectExtInst(ResVReg, ResType,
I, CL::sinh, GL::Sinh);
1208 case TargetOpcode::G_FTANH:
1209 return selectExtInst(ResVReg, ResType,
I, CL::tanh, GL::Tanh);
1211 case TargetOpcode::G_STRICT_FSQRT:
1212 case TargetOpcode::G_FSQRT:
1213 return selectExtInst(ResVReg, ResType,
I, CL::sqrt, GL::Sqrt);
1215 case TargetOpcode::G_CTTZ:
1216 case TargetOpcode::G_CTTZ_ZERO_POISON:
1217 return selectExtInst(ResVReg, ResType,
I, CL::ctz);
1218 case TargetOpcode::G_CTLZ:
1219 case TargetOpcode::G_CTLZ_ZERO_POISON:
1220 return selectExtInst(ResVReg, ResType,
I, CL::clz);
1222 case TargetOpcode::G_INTRINSIC_ROUND:
1223 return selectExtInst(ResVReg, ResType,
I, CL::round, GL::Round);
1224 case TargetOpcode::G_INTRINSIC_ROUNDEVEN:
1225 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1226 case TargetOpcode::G_INTRINSIC_TRUNC:
1227 return selectExtInst(ResVReg, ResType,
I, CL::trunc, GL::Trunc);
1228 case TargetOpcode::G_FRINT:
1229 case TargetOpcode::G_FNEARBYINT:
1230 return selectExtInst(ResVReg, ResType,
I, CL::rint, GL::RoundEven);
1232 case TargetOpcode::G_SMULH:
1233 return selectExtInst(ResVReg, ResType,
I, CL::s_mul_hi);
1234 case TargetOpcode::G_UMULH:
1235 return selectExtInst(ResVReg, ResType,
I, CL::u_mul_hi);
1237 case TargetOpcode::G_SADDSAT:
1238 return selectExtInst(ResVReg, ResType,
I, CL::s_add_sat);
1239 case TargetOpcode::G_UADDSAT:
1240 return selectExtInst(ResVReg, ResType,
I, CL::u_add_sat);
1241 case TargetOpcode::G_SSUBSAT:
1242 return selectExtInst(ResVReg, ResType,
I, CL::s_sub_sat);
1243 case TargetOpcode::G_USUBSAT:
1244 return selectExtInst(ResVReg, ResType,
I, CL::u_sub_sat);
1246 case TargetOpcode::G_FFREXP:
1247 return selectFrexp(ResVReg, ResType,
I);
1249 case TargetOpcode::G_UADDO:
1250 return selectOverflowArith(ResVReg, ResType,
I,
1252 : SPIRV::OpIAddCarryS);
1253 case TargetOpcode::G_USUBO:
1254 return selectOverflowArith(ResVReg, ResType,
I,
1256 : SPIRV::OpISubBorrowS);
1257 case TargetOpcode::G_UMULO:
1258 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpUMulExtended);
1259 case TargetOpcode::G_SMULO:
1260 return selectOverflowArith(ResVReg, ResType,
I, SPIRV::OpSMulExtended);
1262 case TargetOpcode::G_SEXT:
1263 return selectExt(ResVReg, ResType,
I,
true);
1264 case TargetOpcode::G_ANYEXT:
1265 case TargetOpcode::G_ZEXT:
1266 return selectExt(ResVReg, ResType,
I,
false);
1267 case TargetOpcode::G_TRUNC:
1268 return selectTrunc(ResVReg, ResType,
I);
1269 case TargetOpcode::G_FPTRUNC:
1270 case TargetOpcode::G_FPEXT:
1271 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpFConvert);
1273 case TargetOpcode::G_PTRTOINT:
1274 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertPtrToU);
1275 case TargetOpcode::G_INTTOPTR:
1276 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpConvertUToPtr);
1277 case TargetOpcode::G_BITCAST:
1278 return selectBitcast(ResVReg, ResType,
I);
1279 case TargetOpcode::G_ADDRSPACE_CAST:
1280 return selectAddrSpaceCast(ResVReg, ResType,
I);
1281 case TargetOpcode::G_PTRMASK:
1282 return selectPtrMask(ResVReg, ResType,
I);
1283 case TargetOpcode::G_PTR_ADD: {
1285 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
1289 assert(((*II).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1290 (*II).getOpcode() == TargetOpcode::COPY ||
1291 (*II).getOpcode() == SPIRV::OpVariable ||
1292 (*II).getOpcode() == SPIRV::OpUntypedVariableKHR) &&
1295 bool IsGVInit =
false;
1299 UseIt != UseEnd; UseIt = std::next(UseIt)) {
1300 if ((*UseIt).getOpcode() == TargetOpcode::G_GLOBAL_VALUE ||
1301 (*UseIt).getOpcode() == SPIRV::OpSpecConstantOp ||
1302 (*UseIt).getOpcode() == SPIRV::OpVariable ||
1303 (*UseIt).getOpcode() == SPIRV::OpUntypedVariableKHR) {
1315 const bool UseUntypedPointers =
1316 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1317 if (UseUntypedPointers) {
1318 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1321 .
addImm(
static_cast<uint32_t
>(
1322 SPIRV::Opcode::UntypedInBoundsPtrAccessChainKHR))
1325 .
addUse(
I.getOperand(2).getReg())
1332 if (GVPointeeType && ResPointeeType && GVPointeeType != ResPointeeType) {
1344 return diagnoseUnsupported(
1345 I,
"incompatible result and operand types in a bitcast");
1347 MachineInstrBuilder MIB =
1348 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
1355 : SPIRV::OpInBoundsPtrAccessChain))
1359 .
addUse(
I.getOperand(2).getReg())
1362 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1366 static_cast<uint32_t
>(SPIRV::Opcode::InBoundsPtrAccessChain))
1368 .
addUse(
I.getOperand(2).getReg())
1377 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSpecConstantOp))
1380 .
addImm(
static_cast<uint32_t
>(
1381 SPIRV::Opcode::InBoundsPtrAccessChain))
1384 .
addUse(
I.getOperand(2).getReg());
1389 case TargetOpcode::G_ATOMICRMW_OR:
1390 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicOr);
1391 case TargetOpcode::G_ATOMICRMW_ADD:
1392 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicIAdd);
1393 case TargetOpcode::G_ATOMICRMW_AND:
1394 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicAnd);
1395 case TargetOpcode::G_ATOMICRMW_MAX:
1396 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMax);
1397 case TargetOpcode::G_ATOMICRMW_MIN:
1398 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicSMin);
1399 case TargetOpcode::G_ATOMICRMW_SUB:
1400 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicISub);
1401 case TargetOpcode::G_ATOMICRMW_XOR:
1402 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicXor);
1403 case TargetOpcode::G_ATOMICRMW_UMAX:
1404 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMax);
1405 case TargetOpcode::G_ATOMICRMW_UMIN:
1406 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicUMin);
1407 case TargetOpcode::G_ATOMICRMW_XCHG:
1408 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicExchange);
1410 case TargetOpcode::G_ATOMICRMW_FADD:
1411 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT);
1412 case TargetOpcode::G_ATOMICRMW_FSUB:
1414 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFAddEXT,
1416 : SPIRV::OpFNegate);
1417 case TargetOpcode::G_ATOMICRMW_FMIN:
1418 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMinEXT);
1419 case TargetOpcode::G_ATOMICRMW_FMAX:
1420 return selectAtomicRMW(ResVReg, ResType,
I, SPIRV::OpAtomicFMaxEXT);
1422 case TargetOpcode::G_FENCE:
1423 return selectFence(
I);
1425 case TargetOpcode::G_STACKSAVE:
1426 return selectStackSave(ResVReg, ResType,
I);
1427 case TargetOpcode::G_STACKRESTORE:
1428 return selectStackRestore(
I);
1430 case TargetOpcode::G_UNMERGE_VALUES:
1433 case TargetOpcode::G_TRAP:
1434 case TargetOpcode::G_UBSANTRAP:
1435 return selectTrap(
I);
1440 case TargetOpcode::DBG_LABEL:
1442 case TargetOpcode::G_DEBUGTRAP:
1443 return selectDebugTrap(ResVReg, ResType,
I);
1444 case TargetOpcode::G_PREFETCH:
1445 return selectPrefetch(
I);
1452bool SPIRVInstructionSelector::selectDebugTrap(
Register ResVReg,
1453 SPIRVTypeInst ResType,
1454 MachineInstr &
I)
const {
1455 unsigned Opcode = SPIRV::OpNop;
1462bool SPIRVInstructionSelector::selectPrefetch(MachineInstr &
I)
const {
1471 MachineIRBuilder MIRBuilder(
I);
1473 const SPIRVTypeInst PointerSizeType =
1481 Register AddrVal =
I.getOperand(0).getReg();
1484 return selectExtInst(ExtReg, GR.
getOpTypeVoid(MIRBuilder),
I, CL::prefetch,
1486 {AddrVal, ConstIntOne});
1491bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1492 SPIRVTypeInst ResType,
1494 GL::GLSLExtInst GLInst,
1495 bool setMIFlags,
bool useMISrc,
1498 SPIRV::InstructionSet::InstructionSet::GLSL_std_450))
1499 return diagnoseUnsupported(
1501 "this instruction is only supported with the GLSL extended instruction "
1503 return selectExtInst(ResVReg, ResType,
I,
1504 {{SPIRV::InstructionSet::GLSL_std_450, GLInst}},
1505 setMIFlags, useMISrc, SrcRegs);
1508bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1509 SPIRVTypeInst ResType,
1511 CL::OpenCLExtInst CLInst,
1512 bool setMIFlags,
bool useMISrc,
1514 return selectExtInst(ResVReg, ResType,
I,
1515 {{SPIRV::InstructionSet::OpenCL_std, CLInst}},
1516 setMIFlags, useMISrc, SrcRegs);
1519bool SPIRVInstructionSelector::selectExtInst(
1520 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
1521 CL::OpenCLExtInst CLInst, GL::GLSLExtInst GLInst,
bool setMIFlags,
1523 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CLInst},
1524 {SPIRV::InstructionSet::GLSL_std_450, GLInst}};
1525 return selectExtInst(ResVReg, ResType,
I, ExtInsts, setMIFlags, useMISrc,
1529bool SPIRVInstructionSelector::selectExtInst(
Register ResVReg,
1530 SPIRVTypeInst ResType,
1533 bool setMIFlags,
bool useMISrc,
1536 for (
const auto &[InstructionSet, Opcode] : Insts) {
1540 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1543 .
addImm(
static_cast<uint32_t
>(InstructionSet))
1548 const unsigned NumOps =
I.getNumOperands();
1551 I.getOperand(Index).getType() ==
1552 MachineOperand::MachineOperandType::MO_IntrinsicID)
1555 MIB.
add(
I.getOperand(Index));
1567bool SPIRVInstructionSelector::selectCopySign(
Register ResVReg,
1568 SPIRVTypeInst ResType,
1569 MachineInstr &
I)
const {
1571 return selectExtInst(ResVReg, ResType,
I, CL::copysign);
1576 Register MagnitudeReg =
I.getOperand(1).getReg();
1577 Register SignReg =
I.getOperand(2).getReg();
1585 unsigned AndOpcode, OrOpcode;
1586 if (ComponentCount > 1) {
1590 AndOpcode = SPIRV::OpBitwiseAndV;
1591 OrOpcode = SPIRV::OpBitwiseOrV;
1595 AndOpcode = SPIRV::OpBitwiseAndS;
1596 OrOpcode = SPIRV::OpBitwiseOrS;
1602 return selectOpWithSrcs(ResReg, IntType,
I, SrcRegs, Opcode);
1605 Register MagnitudeInt, SignInt, MagnitudeBits, SignBits, CombinedInt;
1606 if (!EmitBitOp(MagnitudeInt, {MagnitudeReg}, SPIRV::OpBitcast) ||
1607 !EmitBitOp(SignInt, {SignReg}, SPIRV::OpBitcast) ||
1608 !EmitBitOp(MagnitudeBits, {MagnitudeInt, NotSignMask}, AndOpcode) ||
1609 !EmitBitOp(SignBits, {SignInt, SignMask}, AndOpcode) ||
1610 !EmitBitOp(CombinedInt, {MagnitudeBits, SignBits}, OrOpcode))
1613 return selectOpWithSrcs(ResVReg, ResType,
I, {CombinedInt}, SPIRV::OpBitcast);
1616bool SPIRVInstructionSelector::selectFrexp(
Register ResVReg,
1617 SPIRVTypeInst ResType,
1618 MachineInstr &
I)
const {
1619 ExtInstList ExtInsts = {{SPIRV::InstructionSet::OpenCL_std, CL::frexp},
1620 {SPIRV::InstructionSet::GLSL_std_450, GL::Frexp}};
1621 for (
const auto &Ex : ExtInsts) {
1622 SPIRV::InstructionSet::InstructionSet
Set = Ex.first;
1623 uint32_t Opcode = Ex.second;
1627 MachineIRBuilder MIRBuilder(
I);
1634 PointeeTy, MIRBuilder, SPIRV::StorageClass::Function);
1641 const bool IsUntyped =
1642 PointerType->getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1644 BuildMI(*It->getParent(), It, It->getDebugLoc(),
1645 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
1646 : SPIRV::OpVariable))
1649 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
1654 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1657 .
addImm(
static_cast<uint32_t
>(Ex.first))
1659 .
add(
I.getOperand(2))
1663 Register ExpResReg =
I.getOperand(1).getReg();
1665 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1675bool SPIRVInstructionSelector::selectLdexp(
Register ResVReg,
1676 SPIRVTypeInst ResType,
1677 MachineInstr &
I)
const {
1678 Register XReg =
I.getOperand(1).getReg();
1679 Register ExpReg =
I.getOperand(2).getReg();
1685 if (ResType->
getOpcode() == SPIRV::OpTypeVector &&
1686 ExpType->
getOpcode() != SPIRV::OpTypeVector) {
1688 ExpReg = buildVectorSplat(ExpReg, NumElts,
I);
1691 return selectExtInst(ResVReg, ResType,
I, CL::ldexp, GL::Ldexp,
1692 true,
false, {XReg, ExpReg});
1695bool SPIRVInstructionSelector::selectSincos(
Register ResVReg,
1696 SPIRVTypeInst ResType,
1697 MachineInstr &
I)
const {
1698 Register CosResVReg =
I.getOperand(1).getReg();
1699 unsigned SrcIdx =
I.getNumExplicitDefs();
1704 MachineIRBuilder MIRBuilder(
I);
1706 ResType, MIRBuilder, SPIRV::StorageClass::Function);
1713 const bool IsUntyped =
1714 PointerType->getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
1716 BuildMI(*It->getParent(), It, It->getDebugLoc(),
1717 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
1718 : SPIRV::OpVariable))
1721 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
1725 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1728 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
1730 .
add(
I.getOperand(SrcIdx))
1734 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
1742 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1745 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1747 .
add(
I.getOperand(SrcIdx))
1749 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
1752 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
1754 .
add(
I.getOperand(SrcIdx))
1761bool SPIRVInstructionSelector::selectOpWithSrcs(
Register ResVReg,
1762 SPIRVTypeInst ResType,
1765 unsigned Opcode)
const {
1766 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
1776bool SPIRVInstructionSelector::selectPopCount16(
Register ResVReg,
1777 SPIRVTypeInst ResType,
1780 unsigned Opcode)
const {
1781 MachineIRBuilder MIRBuilder(
I);
1783 Register OpReg =
I.getOperand(1).getReg();
1789 SPIRVTypeInst ExtType =
1796 if (!selectOpWithSrcs(ExtReg, ExtType,
I, {OpReg}, SPIRV::OpUConvert))
1800 if (!selectPopCount32(PopCountReg, ExtType,
I, ExtReg, Opcode))
1803 return selectOpWithSrcs(ResVReg, ResType,
I, {PopCountReg}, ExtOpcode);
1806bool SPIRVInstructionSelector::selectPopCount32(
Register ResVReg,
1807 SPIRVTypeInst ResType,
1810 unsigned Opcode)
const {
1811 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
1814bool SPIRVInstructionSelector::selectPopCount64(
Register ResVReg,
1815 SPIRVTypeInst ResType,
1818 unsigned Opcode)
const {
1819 MachineIRBuilder MIRBuilder(
I);
1828 SPIRVTypeInst WorkingType =
1835 if (!selectOpWithSrcs(Trunc, WorkingType,
I, {SrcReg}, SPIRV::OpUConvert))
1839 if (!selectOpWithSrcs(LowCount, WorkingType,
I, {Trunc}, SPIRV::OpBitCount))
1847 IsScalar ? SPIRV::OpShiftRightLogicalS : SPIRV::OpShiftRightLogicalV;
1849 if (!selectOpWithSrcs(Shift, SrcType,
I, {SrcReg, ShiftAmount}, ShiftOp))
1853 if (!selectOpWithSrcs(Trunc, WorkingType,
I, {Shift}, SPIRV::OpUConvert))
1857 if (!selectOpWithSrcs(HighCount, WorkingType,
I, {Trunc}, SPIRV::OpBitCount))
1862 if (!selectOpWithSrcs(Sum, WorkingType,
I, {HighCount, LowCount},
1863 IsScalar ? SPIRV::OpIAddS : SPIRV::OpIAddV))
1867 unsigned ConvOp = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
1868 return selectOpWithSrcs(ResVReg, ResType,
I, {Sum}, ConvOp);
1871bool SPIRVInstructionSelector::selectPopCount(
Register ResVReg,
1872 SPIRVTypeInst ResType,
1874 unsigned Opcode)
const {
1879 if (!STI.getTargetTriple().isVulkanOS())
1880 return selectUnOp(ResVReg, ResType,
I, Opcode);
1882 Register OpReg =
I.getOperand(1).getReg();
1885 : SPIRV::OpUConvert;
1889 return selectPopCount16(ResVReg, ResType,
I, ExtOpcode, Opcode);
1891 return selectPopCount32(ResVReg, ResType,
I, OpReg, Opcode);
1893 return selectPopCount64(ResVReg, ResType,
I, OpReg, Opcode);
1895 return diagnoseUnsupported(
I,
"unsupported operand bit width for popcount");
1899bool SPIRVInstructionSelector::selectUnOp(
Register ResVReg,
1900 SPIRVTypeInst ResType,
1902 unsigned Opcode)
const {
1904 Register SrcReg =
I.getOperand(1).getReg();
1909 unsigned DefOpCode = DefIt->getOpcode();
1910 if (DefOpCode == SPIRV::ASSIGN_TYPE || DefOpCode == TargetOpcode::COPY) {
1913 if (
auto *VRD =
getVRegDef(*
MRI, DefIt->getOperand(1).getReg()))
1914 DefOpCode = VRD->getOpcode();
1916 if (DefOpCode == TargetOpcode::G_GLOBAL_VALUE ||
1917 DefOpCode == TargetOpcode::G_CONSTANT ||
1918 DefOpCode == SPIRV::OpVariable ||
1919 DefOpCode == SPIRV::OpUntypedVariableKHR ||
1920 DefOpCode == SPIRV::OpConstantI) {
1926 uint32_t SpecOpcode = 0;
1928 case SPIRV::OpConvertPtrToU:
1929 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertPtrToU);
1931 case SPIRV::OpConvertUToPtr:
1932 SpecOpcode =
static_cast<uint32_t
>(SPIRV::Opcode::ConvertUToPtr);
1937 TII.get(SPIRV::OpSpecConstantOp))
1947 return selectOpWithSrcs(ResVReg, ResType,
I, {
I.getOperand(1).getReg()},
1951bool SPIRVInstructionSelector::selectBitcast(
Register ResVReg,
1952 SPIRVTypeInst ResType,
1953 MachineInstr &
I)
const {
1954 Register OpReg =
I.getOperand(1).getReg();
1955 SPIRVTypeInst OpType =
1958 return diagnoseUnsupported(
1959 I,
"incompatible result and operand types in a bitcast");
1960 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpBitcast);
1966 std::optional<Align> AlignOverride = std::nullopt) {
1971 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
1972 if (
MemOp->isNonTemporal())
1973 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
1975 if (!ST->isShader() &&
MemOp->getAlign().value())
1976 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned);
1980 if (ST->canUseExtension(SPIRV::Extension::SPV_INTEL_memory_access_aliasing)) {
1981 if (
auto *MD =
MemOp->getAAInfo().Scope) {
1985 static_cast<uint32_t>(SPIRV::MemoryOperand::AliasScopeINTELMask);
1987 if (
auto *MD =
MemOp->getAAInfo().NoAlias) {
1991 static_cast<uint32_t>(SPIRV::MemoryOperand::NoAliasINTELMask);
1995 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None)) {
1997 if (SpvMemOp &
static_cast<uint32_t>(SPIRV::MemoryOperand::Aligned))
1998 MIB.
addImm(AlignOverride.value_or(
MemOp->getAlign()).value());
2009 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Volatile);
2011 SpvMemOp |=
static_cast<uint32_t>(SPIRV::MemoryOperand::Nontemporal);
2013 if (SpvMemOp !=
static_cast<uint32_t>(SPIRV::MemoryOperand::None))
2017bool SPIRVInstructionSelector::selectLoad(
Register ResVReg,
2018 SPIRVTypeInst ResType,
2019 MachineInstr &
I)
const {
2021 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2026 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2027 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2029 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2031 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2035 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2039 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2040 return generateImageReadOrFetch(ResVReg, ResType, NewHandleReg, IdxReg,
2041 I.getDebugLoc(),
I);
2045 MachineIRBuilder MIRBuilder(
I);
2047 if (
I.getNumMemOperands()) {
2048 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2049 if (MemOp->isAtomic())
2050 return selectAtomicLoad(ResVReg, ResType,
I);
2053 auto MIB = MIRBuilder.
buildInstr(SPIRV::OpLoad)
2057 if (!
I.getNumMemOperands()) {
2058 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2060 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2069Register SPIRVInstructionSelector::createPtrSizedIntReg(
2070 MachineIRBuilder &MIRBuilder)
const {
2071 SPIRVTypeInst IntType =
2081SPIRVInstructionSelector::convertPtrToInt(
Register PtrVal,
2082 MachineIRBuilder &MIRBuilder)
const {
2083 SPIRVTypeInst IntType =
2085 Register IntReg = createPtrSizedIntReg(MIRBuilder);
2086 MIRBuilder.
buildInstr(SPIRV::OpConvertPtrToU)
2094Register SPIRVInstructionSelector::castPtrToPtrToInt(
2095 Register Ptr, SPIRV::StorageClass::StorageClass SC,
2096 MachineIRBuilder &MIRBuilder)
const {
2097 SPIRVTypeInst IntType =
2099 SPIRVTypeInst PtrType =
2113bool SPIRVInstructionSelector::selectAtomicPtrValue(
2114 Register ResVReg, SPIRVTypeInst ResType, MachineIRBuilder &MIRBuilder,
2115 function_ref<
Register(SPIRVTypeInst IntType)> EmitAtomic)
const {
2123 Register IntResult = EmitAtomic(IntType);
2125 MIRBuilder.
buildInstr(SPIRV::OpConvertUToPtr)
2133bool SPIRVInstructionSelector::selectAtomicLoad(
Register ResVReg,
2134 SPIRVTypeInst ResType,
2135 MachineInstr &
I)
const {
2136 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2139 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2142 return diagnoseUnsupported(
2143 I,
"Lowering to SPIR-V of atomic load is only "
2144 "allowed for integer, floating point or pointer types");
2146 assert(
I.getNumMemOperands());
2147 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2148 assert(MemOp.isAtomic());
2150 uint32_t
Scope =
static_cast<uint32_t
>(
2152 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
2158 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2159 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2162 Register MemSemReg = buildI32ConstantInEntryBlock(Sem,
I);
2164 MachineIRBuilder MIRBuilder(
I);
2168 return diagnoseUnsupported(
2169 I,
"Lowering to SPIR-V of atomic load is only "
2170 "allowed for pointer types for physical addressing model");
2175 SPIRV::StorageClass::StorageClass SC =
2177 return selectAtomicPtrValue(
2178 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2179 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2180 Register IntResult = createPtrSizedIntReg(MIRBuilder);
2191 auto AtomicLoad = MIRBuilder.
buildInstr(SPIRV::OpAtomicLoad)
2202bool SPIRVInstructionSelector::selectStore(MachineInstr &
I)
const {
2204 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2205 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2210 (IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getbasepointer ||
2211 IntPtrDef->getIntrinsicID() == Intrinsic::spv_resource_getpointer)) {
2213 Register HandleReg = IntPtrDef->getOperand(2).getReg();
2218 if (!loadHandleBeforePosition(NewHandleReg, HandleType, *HandleDef,
I)) {
2222 Register IdxReg = IntPtrDef->getOperand(3).getReg();
2223 if (HandleType->
getOpcode() == SPIRV::OpTypeImage) {
2224 SPIRVTypeInst SampledType =
2226 SPIRVTypeInst StoreValCompType =
2228 if (StoreValCompType && StoreValCompType != SampledType) {
2231 SPIRVTypeInst PackedType = widenTypeToVec4(SampledType,
I);
2234 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
2239 StoreVal = PackedReg;
2242 auto BMI =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2243 TII.get(SPIRV::OpImageWrite))
2249 if (sampledTypeIsSignedInteger(LLVMHandleType))
2252 BMI.constrainAllUses(
TII,
TRI, RBI);
2259 if (PointeeTy && PointeeTy->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
2260 StoreTy->
getOpcode() != SPIRV::OpTypeVectorIdEXT &&
2274 if (
I.getNumMemOperands()) {
2275 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2276 if (MemOp->isAtomic())
2277 return selectAtomicStore(
I);
2284 if (PtrSC == SPIRV::StorageClass::UniformConstant ||
2285 PtrSC == SPIRV::StorageClass::Input ||
2286 PtrSC == SPIRV::StorageClass::PushConstant)
2287 return diagnoseUnsupported(
2288 I,
"store into a read-only SPIR-V storage class is not allowed");
2290 MachineIRBuilder MIRBuilder(
I);
2292 if (!
I.getNumMemOperands()) {
2293 assert(
I.getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS ||
2295 TargetOpcode::G_INTRINSIC_CONVERGENT_W_SIDE_EFFECTS);
2304bool SPIRVInstructionSelector::selectAtomicStore(MachineInstr &
I)
const {
2305 LLVMContext &
Context =
I.getMF()->getFunction().getContext();
2308 Register StoreVal =
I.getOperand(0 + OpOffset).getReg();
2309 Register Ptr =
I.getOperand(1 + OpOffset).getReg();
2314 if (!PointeeType && PtrType &&
2315 PtrType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR)
2318 return diagnoseUnsupported(
I,
2319 "Lowering to SPIR-V of atomic store is only "
2320 "allowed for integer or floating point types");
2322 assert(
I.getNumMemOperands());
2323 const MachineMemOperand &MemOp = **
I.memoperands_begin();
2324 assert(MemOp.isAtomic());
2326 uint32_t
Scope =
static_cast<uint32_t
>(
2328 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
2334 if (MemOp.isVolatile() && STI.getTargetTriple().isVulkanOS())
2335 MemSem |=
static_cast<uint32_t
>(SPIRV::MemorySemantics::Volatile);
2338 Register MemSemReg = buildI32ConstantInEntryBlock(Sem,
I);
2339 MachineIRBuilder MIRBuilder(
I);
2343 return diagnoseUnsupported(
2344 I,
"Lowering to SPIR-V of atomic store is only "
2345 "allowed for pointer types for physical addressing model");
2350 SPIRV::StorageClass::StorageClass SC =
2352 return selectAtomicPtrValue(
2353 Register(), SPIRVTypeInst(), MIRBuilder, [&](SPIRVTypeInst IntType) {
2355 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2368 return diagnoseUnsupported(
I,
2369 "Lowering to SPIR-V of atomic store is only "
2370 "allowed for integer or floating point types");
2372 auto AtomicStore = MIRBuilder.
buildInstr(SPIRV::OpAtomicStore)
2382bool SPIRVInstructionSelector::selectMaskedGather(
Register ResVReg,
2383 SPIRVTypeInst ResType,
2384 MachineInstr &
I)
const {
2385 assert(
I.getNumExplicitDefs() == 1 &&
"Expected single def for gather");
2393 const Register PtrsReg =
I.getOperand(2).getReg();
2394 const uint32_t
Alignment =
I.getOperand(3).getImm();
2395 const Register MaskReg =
I.getOperand(4).getReg();
2396 const Register PassthruReg =
I.getOperand(5).getReg();
2397 const Register AlignmentReg = buildI32ConstantInEntryBlock(Alignment,
I);
2401 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedGatherINTEL))
2412bool SPIRVInstructionSelector::selectMaskedScatter(MachineInstr &
I)
const {
2413 assert(
I.getNumExplicitDefs() == 0 &&
"Expected no defs for scatter");
2420 const Register ValuesReg =
I.getOperand(1).getReg();
2421 const Register PtrsReg =
I.getOperand(2).getReg();
2422 const uint32_t
Alignment =
I.getOperand(3).getImm();
2423 const Register MaskReg =
I.getOperand(4).getReg();
2424 const Register AlignmentReg = buildI32ConstantInEntryBlock(Alignment,
I);
2428 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMaskedScatterINTEL))
2437bool SPIRVInstructionSelector::diagnoseUnsupported(
const MachineInstr &
I,
2438 const Twine &
Msg)
const {
2439 const Function &
F =
I.getMF()->getFunction();
2440 F.getContext().diagnose(
2441 DiagnosticInfoUnsupported(
F,
Msg,
I.getDebugLoc(),
DS_Error));
2445bool SPIRVInstructionSelector::selectStackSave(
Register ResVReg,
2446 SPIRVTypeInst ResType,
2447 MachineInstr &
I)
const {
2448 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2449 return diagnoseUnsupported(
2450 I,
"llvm.stacksave intrinsic: this instruction requires the following "
2451 "SPIR-V extension: SPV_INTEL_variable_length_array");
2453 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSaveMemoryINTEL))
2460bool SPIRVInstructionSelector::selectStackRestore(MachineInstr &
I)
const {
2461 if (!STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_variable_length_array))
2462 return diagnoseUnsupported(
2464 "llvm.stackrestore intrinsic: this instruction requires the following "
2465 "SPIR-V extension: SPV_INTEL_variable_length_array");
2466 if (!
I.getOperand(0).isReg())
2469 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpRestoreMemoryINTEL))
2470 .
addUse(
I.getOperand(0).getReg())
2476SPIRVInstructionSelector::getOrCreateMemSetGlobal(MachineInstr &
I)
const {
2477 MachineIRBuilder MIRBuilder(
I);
2478 assert(
I.getOperand(1).isReg() &&
I.getOperand(2).isReg());
2485 GlobalVariable *GV =
new GlobalVariable(*CurFunction.
getParent(), LLVMArrTy,
2489 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2490 Type *ArrTy = ArrayType::get(ValTy, Num);
2492 ArrTy, MIRBuilder, SPIRV::StorageClass::UniformConstant);
2495 ArrTy, MIRBuilder, SPIRV::AccessQualifier::None,
false);
2506 const bool IsUntyped = VarTy->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
2507 auto MIBVar =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2508 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
2509 : SPIRV::OpVariable))
2512 .
addImm(SPIRV::StorageClass::UniformConstant);
2525bool SPIRVInstructionSelector::selectCopyMemory(MachineInstr &
I,
2528 Register DstReg =
I.getOperand(0).getReg();
2534 return diagnoseUnsupported(
2535 I,
"OpCopyMemory requires operands to have the same type");
2540 return diagnoseUnsupported(
2541 I,
"Unable to determine pointee type size for OpCopyMemory");
2542 const DataLayout &
DL =
I.getMF()->getFunction().getDataLayout();
2543 if (CopySize !=
DL.getTypeStoreSize(
const_cast<Type *
>(LLVMPointeeTy)))
2544 return diagnoseUnsupported(
2545 I,
"OpCopyMemory requires the size to match the pointee type size");
2548 const unsigned Opcode =
2549 IsLogical ? SPIRV::OpCopyMemory : SPIRV::OpCopyMemorySized;
2550 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
2554 MIB.
addUse(
I.getOperand(2).getReg());
2556 if (
I.getNumMemOperands()) {
2557 MachineIRBuilder MIRBuilder(
I);
2558 const MachineMemOperand *DstMemOp = *
I.memoperands_begin();
2560 Align SrcAlign = DstAlign;
2563 if (
I.getNumMemOperands() > 1 && !STI.
isShader())
2564 SrcAlign = (*std::next(
I.memoperands_begin()))->
getAlign();
2574 std::min(DstAlign, SrcAlign));
2581bool SPIRVInstructionSelector::selectMemOperation(
Register ResVReg,
2582 MachineInstr &
I)
const {
2584 Register SizeReg =
I.getOperand(2).getReg();
2586 SizeDef && SizeDef->
getOpcode() == TargetOpcode::G_CONSTANT &&
2590 Register SrcReg =
I.getOperand(1).getReg();
2591 if (
I.getOpcode() == TargetOpcode::G_MEMSET ||
2592 I.getOpcode() == TargetOpcode::G_MEMSET_INLINE) {
2593 Register VarReg = getOrCreateMemSetGlobal(
I);
2596 Type *ValTy = Type::getInt8Ty(
I.getMF()->getFunction().getContext());
2598 ValTy,
I, SPIRV::StorageClass::UniformConstant);
2600 if (!selectOpWithSrcs(SrcReg, SourceTy,
I, {VarReg}, SPIRV::OpBitcast))
2603 if (!selectCopyMemory(
I, SrcReg))
2605 if (ResVReg.
isValid() && ResVReg !=
I.getOperand(0).getReg())
2606 if (!BuildCOPY(ResVReg,
I.getOperand(0).getReg(),
I))
2611bool SPIRVInstructionSelector::selectAtomicRMW(
Register ResVReg,
2612 SPIRVTypeInst ResType,
2615 unsigned NegateOpcode)
const {
2617 const MachineMemOperand *MemOp = *
I.memoperands_begin();
2618 uint32_t
Scope =
static_cast<uint32_t
>(
2620 MemOp->getSyncScopeID()));
2621 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
2623 Register Ptr =
I.getOperand(1).getReg();
2624 uint32_t ScSem =
static_cast<uint32_t
>(
2628 Register MemSemReg = buildI32ConstantInEntryBlock(
2632 Register ValueReg =
I.getOperand(2).getReg();
2633 if (NegateOpcode != 0) {
2636 if (!selectOpWithSrcs(TmpReg, ResType,
I, {ValueReg}, NegateOpcode))
2642 if (NewOpcode != SPIRV::OpAtomicExchange)
2643 return diagnoseUnsupported(
2644 I,
"Lowering to SPIR-V of this atomic operation is not "
2645 "allowed for pointer types");
2647 return diagnoseUnsupported(
2648 I,
"Lowering to SPIR-V of atomic exchange is only "
2649 "allowed for pointer types for physical addressing model");
2656 MachineIRBuilder MIRBuilder(
I);
2658 return selectAtomicPtrValue(
2659 ResVReg, ResType, MIRBuilder, [&](SPIRVTypeInst IntType) {
2661 Register CastedPtr = castPtrToPtrToInt(Ptr, SC, MIRBuilder);
2662 Register ExchangeResReg = createPtrSizedIntReg(MIRBuilder);
2663 MIRBuilder.
buildInstr(SPIRV::OpAtomicExchange)
2671 return ExchangeResReg;
2675 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(NewOpcode))
2686bool SPIRVInstructionSelector::selectUnmergeValues(MachineInstr &
I)
const {
2687 unsigned ArgI =
I.getNumOperands() - 1;
2689 I.getOperand(ArgI).isReg() ?
I.getOperand(ArgI).getReg() :
Register(0);
2690 SPIRVTypeInst SrcType =
2694 "cannot select G_UNMERGE_VALUES with a non-vector argument");
2698 unsigned CurrentIndex = 0;
2699 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2700 Register ResVReg =
I.getOperand(i).getReg();
2709 ResType = ScalarType;
2718 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorShuffle))
2724 for (
unsigned j = 0;
j < NumElements; ++
j) {
2725 MIB.
addImm(CurrentIndex + j);
2727 CurrentIndex += NumElements;
2731 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2743bool SPIRVInstructionSelector::selectFence(MachineInstr &
I)
const {
2746 ? SPIRV::MemorySemantics::UniformMemory |
2747 SPIRV::MemorySemantics::WorkgroupMemory |
2748 SPIRV::MemorySemantics::ImageMemory
2749 : SPIRV::MemorySemantics::WorkgroupMemory |
2750 SPIRV::MemorySemantics::CrossWorkgroupMemory |
2751 SPIRV::MemorySemantics::ImageMemory;
2753 STI.getTargetTriple(),
static_cast<uint32_t
>(
getMemSemantics(AO)), ScSem);
2754 Register MemSemReg = buildI32ConstantInEntryBlock(MemSem,
I);
2758 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
2760 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpMemoryBarrier))
2767bool SPIRVInstructionSelector::selectOverflowArith(
Register ResVReg,
2768 SPIRVTypeInst ResType,
2770 unsigned Opcode)
const {
2771 Type *ResTy =
nullptr;
2774 return diagnoseUnsupported(
2776 "Not enough info to select the arithmetic with overflow instruction");
2778 return diagnoseUnsupported(
I,
2779 "Expect struct type result for the arithmetic "
2780 "with overflow instruction");
2786 MachineIRBuilder MIRBuilder(
I);
2788 ResTy, MIRBuilder, SPIRV::AccessQualifier::ReadWrite,
false);
2789 assert(
I.getNumDefs() > 1 &&
"Not enought operands");
2796 Register ZeroReg = buildZerosVal(ResType,
I);
2801 if (ResName.
size() > 0)
2809 for (
unsigned i =
I.getNumDefs(); i <
I.getNumOperands(); ++i)
2810 MIB.
addUse(
I.getOperand(i).getReg());
2816 for (
unsigned i = 0; i <
I.getNumDefs(); ++i) {
2818 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
2819 .
addDef(i == 1 ? HigherVReg :
I.getOperand(i).getReg())
2826 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
2827 .
addDef(
I.getOperand(1).getReg())
2835bool SPIRVInstructionSelector::selectAtomicCmpXchg(
Register ResVReg,
2836 SPIRVTypeInst ResType,
2837 MachineInstr &
I)
const {
2839 "selectAtomicCmpXchg only handles the spv_cmpxchg intrinsic");
2840 Register Ptr =
I.getOperand(2).getReg();
2841 Register ScopeReg =
I.getOperand(5).getReg();
2842 Register MemSemEqReg =
I.getOperand(6).getReg();
2843 Register MemSemNeqReg =
I.getOperand(7).getReg();
2845 Register Val =
I.getOperand(4).getReg();
2849 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpAtomicCompareExchange))
2868 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2875 BuildMI(*
I.getParent(),
I,
DL,
TII.get(SPIRV::OpCompositeInsert))
2887 case SPIRV::StorageClass::DeviceOnlyINTEL:
2888 case SPIRV::StorageClass::HostOnlyINTEL:
2897 bool IsGRef =
false;
2898 bool IsAllowedRefs =
2899 llvm::all_of(
MRI->use_instructions(ResVReg), [&IsGRef](
auto const &It) {
2900 unsigned Opcode = It.getOpcode();
2901 if (Opcode == SPIRV::OpConstantComposite ||
2902 Opcode == SPIRV::OpSpecConstantComposite ||
2903 Opcode == SPIRV::OpVariable ||
2904 Opcode == SPIRV::OpUntypedVariableKHR ||
2905 isSpvIntrinsic(It, Intrinsic::spv_init_global))
2906 return IsGRef = true;
2907 return Opcode == SPIRV::OpName;
2909 return IsAllowedRefs && IsGRef;
2912SPIRVTypeInst SPIRVInstructionSelector::getUcharPtrType(
2913 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2915 Type::getInt8Ty(
I.getMF()->getFunction().getContext()),
I, SC);
2918Register SPIRVInstructionSelector::getUcharPtrTypeReg(
2919 MachineInstr &
I, SPIRV::StorageClass::StorageClass SC)
const {
2924SPIRVInstructionSelector::buildSpecConstantOp(MachineInstr &
I,
Register Dest,
2926 uint32_t Opcode)
const {
2927 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
2928 TII.get(SPIRV::OpSpecConstantOp))
2936SPIRVInstructionSelector::buildConstGenericPtr(MachineInstr &
I,
Register SrcPtr,
2937 SPIRVTypeInst SrcPtrTy)
const {
2938 SPIRVTypeInst GenericPtrTy =
2942 SPIRV::StorageClass::Generic),
2946 MachineInstrBuilder MIB = buildSpecConstantOp(
2948 static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric));
2958bool SPIRVInstructionSelector::selectAddrSpaceCast(
Register ResVReg,
2959 SPIRVTypeInst ResType,
2960 MachineInstr &
I)
const {
2964 Register SrcPtr =
I.getOperand(1).getReg();
2969 return BuildCOPY(ResVReg, SrcPtr,
I);
2979 unsigned SpecOpcode = [&]() ->
unsigned {
2980 if (SrcSC == SPIRV::StorageClass::CodeSectionINTEL)
2981 return static_cast<uint32_t
>(SPIRV::Opcode::Bitcast);
2983 return static_cast<uint32_t
>(SPIRV::Opcode::PtrCastToGeneric);
2985 return static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr);
2996 if (SrcSC != SPIRV::StorageClass::CodeSectionINTEL &&
2997 (SpecOpcode || ViaGeneric)) {
2998 SPIRVTypeInst UcharSrcPtrTy = getUcharPtrType(
I, SrcSC);
2999 if (SrcPtrTy != UcharSrcPtrTy) {
3002 MachineInstrBuilder MIB = buildSpecConstantOp(
3004 static_cast<uint32_t
>(SPIRV::Opcode::Bitcast));
3008 SrcPtrTy = UcharSrcPtrTy;
3012 buildSpecConstantOp(
I, ResVReg, SrcPtr, getUcharPtrTypeReg(
I, DstSC),
3014 .constrainAllUses(
TII,
TRI, RBI);
3015 }
else if (ViaGeneric) {
3016 MachineInstrBuilder MIB = buildConstGenericPtr(
I, SrcPtr, SrcPtrTy);
3018 buildSpecConstantOp(
3020 static_cast<uint32_t
>(SPIRV::Opcode::GenericCastToPtr))
3021 .constrainAllUses(
TII,
TRI, RBI);
3028 return BuildCOPY(ResVReg, SrcPtr,
I);
3030 if ((SrcSC == SPIRV::StorageClass::Function &&
3031 DstSC == SPIRV::StorageClass::Private) ||
3032 (DstSC == SPIRV::StorageClass::Function &&
3033 SrcSC == SPIRV::StorageClass::Private))
3034 return BuildCOPY(ResVReg, SrcPtr,
I);
3038 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
3041 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
3044 SPIRVTypeInst GenericPtrTy =
3063 return selectUnOp(ResVReg, ResType,
I,
3064 SPIRV::OpPtrCastToCrossWorkgroupINTEL);
3066 return selectUnOp(ResVReg, ResType,
I,
3067 SPIRV::OpCrossWorkgroupCastToPtrINTEL);
3069 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpPtrCastToGeneric);
3071 return selectUnOp(ResVReg, ResType,
I, SPIRV::OpGenericCastToPtr);
3081bool SPIRVInstructionSelector::selectPtrMask(
Register ResVReg,
3082 SPIRVTypeInst ResType,
3083 MachineInstr &
I)
const {
3085 return diagnoseUnsupported(
3086 I,
"G_PTRMASK is not supported with logical SPIR-V");
3091 Register PtrReg =
I.getOperand(1).getReg();
3092 Register MaskReg =
I.getOperand(2).getReg();
3111 ? SPIRV::OpBitwiseAndV
3112 : SPIRV::OpBitwiseAndS;
3135 return SPIRV::OpFOrdEqual;
3137 return SPIRV::OpFOrdGreaterThanEqual;
3139 return SPIRV::OpFOrdGreaterThan;
3141 return SPIRV::OpFOrdLessThanEqual;
3143 return SPIRV::OpFOrdLessThan;
3145 return SPIRV::OpFOrdNotEqual;
3147 return SPIRV::OpOrdered;
3149 return SPIRV::OpFUnordEqual;
3151 return SPIRV::OpFUnordGreaterThanEqual;
3153 return SPIRV::OpFUnordGreaterThan;
3155 return SPIRV::OpFUnordLessThanEqual;
3157 return SPIRV::OpFUnordLessThan;
3159 return SPIRV::OpFUnordNotEqual;
3161 return SPIRV::OpUnordered;
3171 return SPIRV::OpIEqual;
3173 return SPIRV::OpINotEqual;
3175 return SPIRV::OpSGreaterThanEqual;
3177 return SPIRV::OpSGreaterThan;
3179 return SPIRV::OpSLessThanEqual;
3181 return SPIRV::OpSLessThan;
3183 return SPIRV::OpUGreaterThanEqual;
3185 return SPIRV::OpUGreaterThan;
3187 return SPIRV::OpULessThanEqual;
3189 return SPIRV::OpULessThan;
3198 return SPIRV::OpPtrEqual;
3200 return SPIRV::OpPtrNotEqual;
3211 return SPIRV::OpLogicalEqual;
3213 return SPIRV::OpLogicalNotEqual;
3251bool SPIRVInstructionSelector::selectAnyOrAll(
Register ResVReg,
3252 SPIRVTypeInst ResType,
3254 unsigned OpAnyOrAll)
const {
3255 assert(
I.getNumOperands() == 3);
3256 assert(
I.getOperand(2).isReg());
3258 Register InputRegister =
I.getOperand(2).getReg();
3261 assert(InputType &&
"VReg has no type assigned");
3265 assert(ResVReg ==
I.getOperand(0).getReg());
3266 return BuildCOPY(ResVReg, InputRegister,
I);
3270 unsigned SpirvNotEqualId =
3271 IsFloatTy ? SPIRV::OpFOrdNotEqual : SPIRV::OpINotEqual;
3273 SPIRVTypeInst SpvBoolTy = SpvBoolScalarTy;
3278 IsBoolTy ? InputRegister
3286 IsFloatTy ? buildZerosValF(InputType,
I) : buildZerosVal(InputType,
I);
3288 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SpirvNotEqualId))
3305bool SPIRVInstructionSelector::selectAll(
Register ResVReg,
3306 SPIRVTypeInst ResType,
3307 MachineInstr &
I)
const {
3308 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAll);
3311bool SPIRVInstructionSelector::selectAny(
Register ResVReg,
3312 SPIRVTypeInst ResType,
3313 MachineInstr &
I)
const {
3314 return selectAnyOrAll(ResVReg, ResType,
I, SPIRV::OpAny);
3318bool SPIRVInstructionSelector::selectFloatDot(
Register ResVReg,
3319 SPIRVTypeInst ResType,
3320 MachineInstr &
I)
const {
3321 assert(
I.getNumOperands() == 4);
3322 assert(
I.getOperand(2).isReg());
3323 assert(
I.getOperand(3).isReg());
3325 [[maybe_unused]] SPIRVTypeInst VecType =
3330 VecType->
getOpcode() == SPIRV::OpTypeVectorIdEXT) &&
3331 "dot product requires either a vector of at least 2 components or"
3332 " the SPV_EXT_long vector extension.");
3334 [[maybe_unused]] SPIRVTypeInst EltType =
3343 .
addUse(
I.getOperand(2).getReg())
3344 .
addUse(
I.getOperand(3).getReg())
3349bool SPIRVInstructionSelector::selectIntegerDot(
Register ResVReg,
3350 SPIRVTypeInst ResType,
3353 assert(
I.getNumOperands() == 4);
3354 assert(
I.getOperand(2).isReg());
3355 assert(
I.getOperand(3).isReg());
3358 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3362 .
addUse(
I.getOperand(2).getReg())
3363 .
addUse(
I.getOperand(3).getReg())
3370bool SPIRVInstructionSelector::selectIntegerDotExpansion(
3371 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3372 assert(
I.getNumOperands() == 4);
3373 assert(
I.getOperand(2).isReg());
3374 assert(
I.getOperand(3).isReg());
3378 Register Vec0 =
I.getOperand(2).getReg();
3379 Register Vec1 =
I.getOperand(3).getReg();
3383 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulV))
3392 VecType->
getOpcode() == SPIRV::OpTypeVectorIdEXT) &&
3393 "dot product requires either a vector of at least 2 components "
3394 "or the SPV_EXT_long_vector extension.");
3397 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3407 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
3418 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3430bool SPIRVInstructionSelector::selectOpIsInf(
Register ResVReg,
3431 SPIRVTypeInst ResType,
3432 MachineInstr &
I)
const {
3434 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsInf))
3437 .
addUse(
I.getOperand(2).getReg())
3442bool SPIRVInstructionSelector::selectOpIsNan(
Register ResVReg,
3443 SPIRVTypeInst ResType,
3444 MachineInstr &
I)
const {
3446 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
3449 .
addUse(
I.getOperand(2).getReg())
3454bool SPIRVInstructionSelector::selectOpIsFinite(
Register ResVReg,
3455 SPIRVTypeInst ResType,
3456 MachineInstr &
I)
const {
3459 Register Src =
I.getOperand(2).getReg();
3502bool SPIRVInstructionSelector::selectOpIsNormal(
Register ResVReg,
3503 SPIRVTypeInst ResType,
3504 MachineInstr &
I)
const {
3506 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNormal))
3509 .
addUse(
I.getOperand(2).getReg())
3514template <
bool Signed>
3515bool SPIRVInstructionSelector::selectDot4AddPacked(
Register ResVReg,
3516 SPIRVTypeInst ResType,
3517 MachineInstr &
I)
const {
3518 assert(
I.getNumOperands() == 5);
3519 assert(
I.getOperand(2).isReg());
3520 assert(
I.getOperand(3).isReg());
3521 assert(
I.getOperand(4).isReg());
3524 Register Acc =
I.getOperand(2).getReg();
3528 auto DotOp =
Signed ? SPIRV::OpSDot : SPIRV::OpUDot;
3530 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(DotOp))
3535 MIB.
addImm(SPIRV::BuiltIn::PackedVectorFormat4x8Bit);
3538 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3550template <
bool Signed>
3551bool SPIRVInstructionSelector::selectDot4AddPackedExpansion(
3552 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3553 assert(
I.getNumOperands() == 5);
3554 assert(
I.getOperand(2).isReg());
3555 assert(
I.getOperand(3).isReg());
3556 assert(
I.getOperand(4).isReg());
3559 Register Acc =
I.getOperand(2).getReg();
3565 Signed ? SPIRV::OpBitFieldSExtract : SPIRV::OpBitFieldUExtract;
3569 for (
unsigned i = 0; i < 4; i++) {
3592 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIMulS))
3612 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
3627bool SPIRVInstructionSelector::selectSaturate(
Register ResVReg,
3628 SPIRVTypeInst ResType,
3629 MachineInstr &
I)
const {
3630 assert(
I.getNumOperands() == 3);
3631 assert(
I.getOperand(2).isReg());
3633 Register VZero = buildZerosValF(ResType,
I);
3634 Register VOne = buildOnesValF(ResType,
I);
3636 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
3639 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3641 .
addUse(
I.getOperand(2).getReg())
3648bool SPIRVInstructionSelector::selectSign(
Register ResVReg,
3649 SPIRVTypeInst ResType,
3650 MachineInstr &
I)
const {
3651 assert(
I.getNumOperands() == 3);
3652 assert(
I.getOperand(2).isReg());
3654 Register InputRegister =
I.getOperand(2).getReg();
3656 auto &
DL =
I.getDebugLoc();
3659 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3666 bool NeedsConversion = IsFloatTy || SignBitWidth != ResBitWidth;
3668 auto SignOpcode = IsFloatTy ? GL::FSign : GL::SSign;
3676 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
3681 if (NeedsConversion) {
3682 auto ConvertOpcode = IsFloatTy ? SPIRV::OpConvertFToS : SPIRV::OpSConvert;
3693bool SPIRVInstructionSelector::selectWaveOpInst(
Register ResVReg,
3694 SPIRVTypeInst ResType,
3696 unsigned Opcode)
const {
3700 auto BMI =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
3706 for (
unsigned J = 2; J <
I.getNumOperands(); J++) {
3707 BMI.addUse(
I.getOperand(J).getReg());
3714bool SPIRVInstructionSelector::selectBarrierInst(MachineInstr &
I,
3717 bool WithGroupSync)
const {
3719 WithGroupSync ? SPIRV::OpControlBarrier : SPIRV::OpMemoryBarrier;
3721 MemSem |= SPIRV::MemorySemantics::AcquireRelease;
3723 assert(((Scope != SPIRV::Scope::Workgroup) ||
3724 ((MemSem & SPIRV::MemorySemantics::WorkgroupMemory) > 0)) &&
3725 "Workgroup Scope must set WorkGroupMemory semantic "
3726 "in Barrier instruction");
3728 assert(((Scope != SPIRV::Scope::Device) ||
3729 ((MemSem & SPIRV::MemorySemantics::UniformMemory) > 0 &&
3730 (MemSem & SPIRV::MemorySemantics::ImageMemory) > 0)) &&
3731 "Device Scope must set UniformMemory and ImageMemory semantic "
3732 "in Barrier instruction");
3738 if (WithGroupSync) {
3739 Register ExecReg = buildI32ConstantInEntryBlock(SPIRV::Scope::Workgroup,
I);
3743 Register ScopeReg = buildI32ConstantInEntryBlock(Scope,
I);
3744 Register MemSemReg = buildI32ConstantInEntryBlock(MemSem,
I);
3746 MI.addUse(ScopeReg).addUse(MemSemReg).constrainAllUses(
TII,
TRI, RBI);
3750bool SPIRVInstructionSelector::selectWaveActiveCountBits(
3751 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3756 if (!selectWaveOpInst(BallotReg, BallotType,
I,
3757 SPIRV::OpGroupNonUniformBallot))
3762 TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3767 .
addImm(SPIRV::GroupOperation::Reduce)
3774bool SPIRVInstructionSelector::selectWaveActiveAllEqual(
Register ResVReg,
3775 SPIRVTypeInst ResType,
3776 MachineInstr &
I)
const {
3781 Register InputReg =
I.getOperand(2).getReg();
3786 bool IsVector = NumElems > 1 ||
3787 (InputType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
3801 return selectWaveOpInst(ResVReg, ElemBoolType,
I,
3802 SPIRV::OpGroupNonUniformAllEqual);
3807 ElementResults.
reserve(NumElems);
3809 for (
unsigned Idx = 0;
Idx < NumElems; ++
Idx) {
3822 ElemInput = Extracted;
3828 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformAllEqual))
3839 auto MIB =
BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpCompositeConstruct))
3850bool SPIRVInstructionSelector::selectWavePrefixBitCount(
Register ResVReg,
3851 SPIRVTypeInst ResType,
3852 MachineInstr &
I)
const {
3854 assert(
I.getNumOperands() == 3);
3856 auto Op =
I.getOperand(2);
3866 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3868 if (InputType->
getOpcode() != SPIRV::OpTypeBool)
3869 return diagnoseUnsupported(
I,
"WavePrefixBitCount requires boolean input");
3890 BuildMI(BB,
I,
DL,
TII.get(SPIRV::OpGroupNonUniformBallotBitCount))
3894 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
3901bool SPIRVInstructionSelector::selectWaveReduceMax(
Register ResVReg,
3902 SPIRVTypeInst ResType,
3904 bool IsUnsigned)
const {
3905 return selectWaveReduce(
3906 ResVReg, ResType,
I, IsUnsigned,
3907 [&](
Register InputRegister,
bool IsUnsigned) {
3908 const bool IsFloatTy =
3910 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMax
3911 : SPIRV::OpGroupNonUniformSMax;
3912 return IsFloatTy ? SPIRV::OpGroupNonUniformFMax : IntOp;
3916bool SPIRVInstructionSelector::selectWaveReduceMin(
Register ResVReg,
3917 SPIRVTypeInst ResType,
3919 bool IsUnsigned)
const {
3920 return selectWaveReduce(
3921 ResVReg, ResType,
I, IsUnsigned,
3922 [&](
Register InputRegister,
bool IsUnsigned) {
3923 const bool IsFloatTy =
3925 const auto IntOp = IsUnsigned ? SPIRV::OpGroupNonUniformUMin
3926 : SPIRV::OpGroupNonUniformSMin;
3927 return IsFloatTy ? SPIRV::OpGroupNonUniformFMin : IntOp;
3931bool SPIRVInstructionSelector::selectWaveReduceSum(
Register ResVReg,
3932 SPIRVTypeInst ResType,
3933 MachineInstr &
I)
const {
3934 return selectWaveReduce(ResVReg, ResType,
I,
false,
3935 [&](
Register InputRegister,
bool IsUnsigned) {
3937 InputRegister, SPIRV::OpTypeFloat);
3938 return IsFloatTy ? SPIRV::OpGroupNonUniformFAdd
3939 : SPIRV::OpGroupNonUniformIAdd;
3943bool SPIRVInstructionSelector::selectWaveReduceProduct(
Register ResVReg,
3944 SPIRVTypeInst ResType,
3945 MachineInstr &
I)
const {
3946 return selectWaveReduce(ResVReg, ResType,
I,
false,
3947 [&](
Register InputRegister,
bool IsUnsigned) {
3949 InputRegister, SPIRV::OpTypeFloat);
3950 return IsFloatTy ? SPIRV::OpGroupNonUniformFMul
3951 : SPIRV::OpGroupNonUniformIMul;
3955template <
typename PickOpcodeFn>
3956bool SPIRVInstructionSelector::selectWaveReduce(
3957 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
3958 PickOpcodeFn &&PickOpcode)
const {
3959 assert(
I.getNumOperands() == 3);
3960 assert(
I.getOperand(2).isReg());
3962 Register InputRegister =
I.getOperand(2).getReg();
3966 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
3969 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
3975 .
addImm(SPIRV::GroupOperation::Reduce)
3976 .
addUse(
I.getOperand(2).getReg())
3981bool SPIRVInstructionSelector::selectWaveReduceOp(
Register ResVReg,
3982 SPIRVTypeInst ResType,
3984 unsigned Opcode)
const {
3985 return selectWaveReduce(
3986 ResVReg, ResType,
I,
false,
3987 [&](
Register InputRegister,
bool IsUnsigned) {
return Opcode; });
3990bool SPIRVInstructionSelector::selectWaveExclusiveScanSum(
3991 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
3992 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
3993 [&](
Register InputRegister,
bool IsUnsigned) {
3995 InputRegister, SPIRV::OpTypeFloat);
3997 ? SPIRV::OpGroupNonUniformFAdd
3998 : SPIRV::OpGroupNonUniformIAdd;
4002bool SPIRVInstructionSelector::selectWaveExclusiveScanProduct(
4003 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
4004 return selectWaveExclusiveScan(ResVReg, ResType,
I,
false,
4005 [&](
Register InputRegister,
bool IsUnsigned) {
4007 InputRegister, SPIRV::OpTypeFloat);
4009 ? SPIRV::OpGroupNonUniformFMul
4010 : SPIRV::OpGroupNonUniformIMul;
4014template <
typename PickOpcodeFn>
4015bool SPIRVInstructionSelector::selectWaveExclusiveScan(
4016 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
bool IsUnsigned,
4017 PickOpcodeFn &&PickOpcode)
const {
4018 assert(
I.getNumOperands() == 3);
4019 assert(
I.getOperand(2).isReg());
4021 Register InputRegister =
I.getOperand(2).getReg();
4025 return diagnoseUnsupported(
I,
"Input Type could not be determined.");
4028 const unsigned Opcode = PickOpcode(InputRegister, IsUnsigned);
4034 .
addImm(SPIRV::GroupOperation::ExclusiveScan)
4035 .
addUse(
I.getOperand(2).getReg())
4040bool SPIRVInstructionSelector::selectQuadSwap(
Register ResVReg,
4041 SPIRVTypeInst ResType,
4044 assert(
I.getNumOperands() == 3);
4045 assert(
I.getOperand(2).isReg());
4047 Register InputRegister =
I.getOperand(2).getReg();
4053 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGroupNonUniformQuadSwap))
4064bool SPIRVInstructionSelector::selectBitreverseViaI32(
Register ResVReg,
4065 SPIRVTypeInst ResType,
4072 unsigned ShiftOp = SPIRV::OpShiftRightLogicalS;
4077 : SPIRV::OpUConvert;
4079 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4082 ShiftOp = SPIRV::OpShiftRightLogicalV;
4087 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4088 TII.get(SPIRV::OpConstantComposite))
4091 for (
unsigned It = 0; It <
N; ++It)
4095 ShiftConst = CompositeReg;
4100 if (!selectOpWithSrcs(ExtReg, Int32Type,
I, {
Op}, ExtendOpcode))
4105 if (!selectBitreverseNative(BitrevReg, Int32Type,
I, ExtReg))
4110 if (!selectOpWithSrcs(ShiftReg, Int32Type,
I, {BitrevReg, ShiftConst},
4115 return selectOpWithSrcs(ResVReg, ResType,
I, {ShiftReg}, ExtendOpcode);
4118bool SPIRVInstructionSelector::handle64BitOverflow(
4120 unsigned int Opcode,
4127 "handle64BitOverflow should only be used for integer types");
4129 assert(ComponentCount < 5 &&
"Vec 5+ will generate invalid SPIR-V ops");
4131 MachineIRBuilder MIRBuilder(
I);
4133 SPIRVTypeInst I64x2Type =
4135 SPIRVTypeInst Vec2ResType =
4138 std::vector<Register> PartialRegs;
4140 unsigned CurrentComponent = 0;
4141 for (; CurrentComponent + 1 < ComponentCount; CurrentComponent += 2) {
4145 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4146 TII.get(SPIRV::OpVectorShuffle))
4151 .
addImm(CurrentComponent)
4152 .
addImm(CurrentComponent + 1);
4162 PartialRegs.push_back(SubVecReg);
4165 if (CurrentComponent != ComponentCount) {
4171 if (!selectOpWithSrcs(FinalElemReg, I64Type,
I, {SrcReg, ConstIntLastIdx},
4172 SPIRV::OpVectorExtractDynamic))
4181 PartialRegs.push_back(FinalElemResReg);
4185 return selectOpWithSrcs(ResVReg, ResType,
I, PartialRegs,
4186 SPIRV::OpCompositeConstruct);
4189bool SPIRVInstructionSelector::selectBitreverse64(
Register ResVReg,
4190 SPIRVTypeInst ResType,
4194 if (ComponentCount > 2)
4195 return handle64BitOverflow(
4196 ResVReg, ResType,
I, SrcReg, SPIRV::OpBitReverse,
4198 unsigned O) {
return this->selectBitreverse64(R,
T,
I, S); });
4200 MachineIRBuilder MIRBuilder(
I);
4204 I32Type, 2 * ComponentCount, MIRBuilder,
false);
4208 if (!selectOpWithSrcs(Vec32, VecI32Type,
I, {SrcReg}, SPIRV::OpBitcast))
4213 if (!selectBitreverseNative(Reverse32, VecI32Type,
I, Vec32))
4220 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4221 TII.get(SPIRV::OpVectorShuffle))
4226 for (
unsigned J = 0; J < ComponentCount; ++J) {
4233 return selectOpWithSrcs(ResVReg, ResType,
I, {SwappedVec}, SPIRV::OpBitcast);
4236bool SPIRVInstructionSelector::selectBitreverseNative(
Register ResVReg,
4237 SPIRVTypeInst ResType,
4241 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitReverse))
4249bool SPIRVInstructionSelector::selectBitreverse(
Register ResVReg,
4250 SPIRVTypeInst ResType,
4251 MachineInstr &
I)
const {
4252 Register OpReg =
I.getOperand(1).getReg();
4261 return selectBitreverseViaI32(ResVReg, ResType,
I, OpReg);
4263 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4265 return selectBitreverse64(ResVReg, ResType,
I, OpReg);
4267 return SPIRVInstructionSelector::diagnoseUnsupported(
4268 I,
"G_BITREVERSE only support 16,32,64 bits.");
4272 return selectBitreverseNative(ResVReg, ResType,
I, OpReg);
4283 unsigned AndOp = SPIRV::OpBitwiseAndS;
4284 unsigned OrOp = SPIRV::OpBitwiseOrS;
4285 unsigned ShlOp = SPIRV::OpShiftLeftLogicalS;
4286 unsigned ShrOp = SPIRV::OpShiftRightLogicalS;
4287 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4289 AndOp = SPIRV::OpBitwiseAndV;
4290 OrOp = SPIRV::OpBitwiseOrV;
4291 ShlOp = SPIRV::OpShiftLeftLogicalV;
4292 ShrOp = SPIRV::OpShiftRightLogicalV;
4298 const unsigned Shift) ->
Register {
4301 (ResType->
getOpcode() != SPIRV::OpTypeVectorIdEXT ||
4308 Register MaskReg = CreateConst(Mask);
4309 Register ShiftReg = CreateConst(Shift);
4316 if (!selectOpWithSrcs(
T1, ResType,
I, {
Input, ShiftReg}, ShrOp) ||
4317 !selectOpWithSrcs(T2, ResType,
I, {
T1, MaskReg}, AndOp) ||
4318 !selectOpWithSrcs(T3, ResType,
I, {
Input, MaskReg}, AndOp) ||
4319 !selectOpWithSrcs(T4, ResType,
I, {T3, ShiftReg}, ShlOp) ||
4320 !selectOpWithSrcs(Result, ResType,
I, {T2, T4}, OrOp))
4329 while ((Shift >>= 1) > 0) {
4336 return BuildCOPY(ResVReg, Result,
I);
4339bool SPIRVInstructionSelector::selectFreeze(
Register ResVReg,
4340 SPIRVTypeInst ResType,
4341 MachineInstr &
I)
const {
4342 assert(
I.getOperand(0).isReg() &&
I.getOperand(1).isReg() &&
4343 "G_FREEZE must define and use a register");
4344 Register OpReg =
I.getOperand(1).getReg();
4348 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
4362 if (
Def->getOpcode() == TargetOpcode::COPY)
4365 switch (
Def->getOpcode()) {
4366 case SPIRV::ASSIGN_TYPE:
4367 if (MachineInstr *AssignToDef =
4369 if (AssignToDef->getOpcode() == TargetOpcode::G_IMPLICIT_DEF)
4370 Reg =
Def->getOperand(2).getReg();
4373 case SPIRV::OpUndef:
4374 Reg =
Def->getOperand(1).getReg();
4377 unsigned DestOpCode;
4379 DestOpCode = SPIRV::OpConstantNull;
4380 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze of a "
4381 "static undef/poison lowered to OpConstantNull\n");
4383 DestOpCode = TargetOpcode::COPY;
4385 LLVM_DEBUG(
dbgs() <<
"SPV_KHR_poison_freeze is not enabled. freeze "
4386 "skipped, lowered as a copy of the operand\n");
4388 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DestOpCode))
4389 .
addDef(
I.getOperand(0).getReg())
4397bool SPIRVInstructionSelector::selectBuildVector(
Register ResVReg,
4398 SPIRVTypeInst ResType,
4399 MachineInstr &
I)
const {
4403 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4407 if (
I.getNumExplicitOperands() -
I.getNumExplicitDefs() !=
N)
4412 for (
unsigned i =
I.getNumExplicitDefs();
4413 i <
I.getNumExplicitOperands() && IsConst; ++i)
4418 return diagnoseUnsupported(
4419 I,
"There must be at least two constituent operands in a vector");
4424 for (
unsigned i =
I.getNumExplicitDefs();
4425 i <
I.getNumExplicitOperands() && IsNullVector; ++i) {
4431 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
4438 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4439 TII.get(IsConst ? SPIRV::OpConstantComposite
4440 : SPIRV::OpCompositeConstruct))
4443 for (
unsigned i =
I.getNumExplicitDefs(); i <
I.getNumExplicitOperands(); ++i)
4444 MIB.
addUse(
I.getOperand(i).getReg());
4449bool SPIRVInstructionSelector::selectSplatVector(
Register ResVReg,
4450 SPIRVTypeInst ResType,
4451 MachineInstr &
I)
const {
4455 else if (ResType->
getOpcode() == SPIRV::OpTypeArray)
4460 unsigned OpIdx =
I.getNumExplicitDefs();
4461 if (!
I.getOperand(OpIdx).isReg())
4465 Register OpReg =
I.getOperand(OpIdx).getReg();
4469 return diagnoseUnsupported(
4470 I,
"There must be at least two constituent operands in a vector");
4473 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4474 TII.get(IsConst ? SPIRV::OpConstantComposite
4475 : SPIRV::OpCompositeConstruct))
4478 for (
unsigned i = 0; i <
N; ++i)
4484bool SPIRVInstructionSelector::selectConcatVectors(
Register ResVReg,
4485 SPIRVTypeInst ResType,
4486 MachineInstr &
I)
const {
4492 "Cannot select G_CONCAT_VECTORS with a non-vector result");
4494 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4495 TII.get(SPIRV::OpCompositeConstruct))
4498 for (
unsigned OpIdx =
I.getNumExplicitDefs();
4499 OpIdx <
I.getNumExplicitOperands(); ++OpIdx)
4500 MIB.
addUse(
I.getOperand(OpIdx).getReg());
4505bool SPIRVInstructionSelector::selectDiscard(
Register ResVReg,
4506 SPIRVTypeInst ResType,
4507 MachineInstr &
I)
const {
4513 SPIRV::Extension::SPV_EXT_demote_to_helper_invocation) ||
4515 Opcode = SPIRV::OpDemoteToHelperInvocation;
4517 Opcode = SPIRV::OpKill;
4522 ToErase.eraseFromParent();
4531bool SPIRVInstructionSelector::selectCmp(
Register ResVReg,
4532 SPIRVTypeInst ResType,
unsigned CmpOpc,
4533 MachineInstr &
I)
const {
4534 Register Cmp0 =
I.getOperand(2).getReg();
4535 Register Cmp1 =
I.getOperand(3).getReg();
4538 "CMP operands should have the same type");
4539 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(CmpOpc))
4549bool SPIRVInstructionSelector::selectICmp(
Register ResVReg,
4550 SPIRVTypeInst ResType,
4551 MachineInstr &
I)
const {
4555 Register CmpOperand =
I.getOperand(2).getReg();
4557 bool IsPtrCmp = CmpOperandType && CmpOperandType.
isPointer();
4562 Register Op1 =
I.getOperand(3).getReg();
4566 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpBitcast))
4571 I.getOperand(3).setReg(NewOp1);
4577 return selectCmp(ResVReg, ResType, CmpOpc,
I);
4582Register SPIRVInstructionSelector::buildI32ConstantInEntryBlock(
4583 uint32_t Val, MachineInstr &
I, SPIRVTypeInst ResType)
const {
4585 SPIRVTypeInst SpvI32Ty =
4587 auto *ConstInt = ConstantInt::get(LLVMTy, Val);
4592 MachineBasicBlock &EntryBB = *InsertIt->getParent();
4593 MachineInstr *
MI =
nullptr;
4597 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantNull))
4601 uint64_t ImmVal = APInt(32, Val).getZExtValue();
4602 MI =
BuildMI(EntryBB, InsertIt, DbgLoc,
TII.get(SPIRV::OpConstantI))
4608 GR.
add(ConstInt,
MI);
4613bool SPIRVInstructionSelector::selectFCmp(
Register ResVReg,
4614 SPIRVTypeInst ResType,
4615 MachineInstr &
I)
const {
4624 for (
unsigned Idx = 0;
Idx < 2; ++
Idx) {
4626 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIsNan))
4629 .
addUse(
I.getOperand(2 + Idx).getReg())
4636 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLogicalOr))
4643 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLogicalNot))
4651 return selectCmp(ResVReg, ResType, CmpOp,
I);
4654bool SPIRVInstructionSelector::selectExp10(
Register ResVReg,
4655 SPIRVTypeInst ResType,
4656 MachineInstr &
I)
const {
4658 return selectExtInst(ResVReg, ResType,
I, CL::exp10);
4668 MachineIRBuilder MIRBuilder(
I);
4675 APFloat ConstVal(3.3219280948873623);
4679 APFloat::rmNearestTiesToEven, &LosesInfo);
4684 isVectorType(ResType) ? SPIRV::OpVectorTimesScalar : SPIRV::OpFMulS;
4686 if (!selectOpWithSrcs(ArgReg, ResType,
I,
4687 {
I.getOperand(1).getReg(), ConstReg}, Opcode))
4689 if (!selectExtInst(ResVReg, ResType,
I,
4690 {{SPIRV::InstructionSet::GLSL_std_450, GL::Exp2}},
false,
4700Register SPIRVInstructionSelector::buildZerosVal(SPIRVTypeInst ResType,
4701 MachineInstr &
I)
const {
4709bool SPIRVInstructionSelector::isScalarOrVectorIntConstantZero(
4715 if (!CompType || CompType->
getOpcode() != SPIRV::OpTypeInt)
4723 if (
Def->getOpcode() == SPIRV::OpConstantNull)
4726 if (
Def->getOpcode() == TargetOpcode::G_CONSTANT ||
4727 Def->getOpcode() == SPIRV::OpConstantI)
4740 if (
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ||
4741 (
Def->getOpcode() == TargetOpcode::G_INTRINSIC_W_SIDE_EFFECTS &&
4743 Intrinsic::spv_const_composite)) {
4744 unsigned StartOp =
Def->getOpcode() == TargetOpcode::G_BUILD_VECTOR ? 1 : 2;
4745 for (
unsigned i = StartOp; i <
Def->getNumOperands(); ++i) {
4746 if (!IsZero(
Def->getOperand(i).getReg()))
4755Register SPIRVInstructionSelector::buildZerosValF(SPIRVTypeInst ResType,
4756 MachineInstr &
I)
const {
4765Register SPIRVInstructionSelector::buildOnesValF(SPIRVTypeInst ResType,
4766 MachineInstr &
I)
const {
4776 SPIRVTypeInst ResType,
4777 MachineInstr &
I)
const {
4788 MachineInstr &
I)
const {
4793 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
4794 TII.get(SPIRV::OpCompositeConstruct))
4797 for (
unsigned J = 0; J < NumElts; ++J)
4803bool SPIRVInstructionSelector::selectSelect(
Register ResVReg,
4804 SPIRVTypeInst ResType,
4805 MachineInstr &
I)
const {
4806 Register SelectFirstArg =
I.getOperand(2).getReg();
4807 Register SelectSecondArg =
I.getOperand(3).getReg();
4816 Register CondReg =
I.getOperand(1).getReg();
4817 bool IsScalarBool = GR.
isScalarOfType(CondReg, SPIRV::OpTypeBool);
4825 CondReg = buildVectorSplat(CondReg, NumElts,
I);
4826 IsScalarBool =
false;
4829 Opcode = IsScalarBool ? SPIRV::OpSelectVFSCond : SPIRV::OpSelectVFVCond;
4830 }
else if (IsPtrTy) {
4831 Opcode = IsScalarBool ? SPIRV::OpSelectVPSCond : SPIRV::OpSelectVPVCond;
4833 Opcode = IsScalarBool ? SPIRV::OpSelectVISCond : SPIRV::OpSelectVIVCond;
4836 assert(IsScalarBool &&
"OpSelect with a scalar result requires a scalar "
4837 "boolean condition");
4839 Opcode = SPIRV::OpSelectSFSCond;
4840 }
else if (IsPtrTy) {
4841 Opcode = SPIRV::OpSelectSPSCond;
4843 Opcode = SPIRV::OpSelectSISCond;
4846 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
4858bool SPIRVInstructionSelector::selectBoolToInt(
Register ResVReg,
4859 SPIRVTypeInst ResType,
4861 MachineInstr &InsertAt,
4862 bool IsSigned)
const {
4864 Register ZeroReg = buildZerosVal(ResType, InsertAt);
4865 Register OneReg = buildOnesVal(IsSigned, ResType, InsertAt);
4866 bool IsScalarBool = GR.
isScalarOfType(BooleanVReg, SPIRV::OpTypeBool);
4868 IsScalarBool ? SPIRV::OpSelectSISCond : SPIRV::OpSelectVIVCond;
4880bool SPIRVInstructionSelector::selectIToF(
Register ResVReg,
4881 SPIRVTypeInst ResType,
4882 MachineInstr &
I,
bool IsSigned,
4883 unsigned Opcode)
const {
4884 Register SrcReg =
I.getOperand(1).getReg();
4895 selectBoolToInt(SrcReg, TmpType,
I.getOperand(1).getReg(),
I, IsSigned);
4897 return selectOpWithSrcs(ResVReg, ResType,
I, {SrcReg}, Opcode);
4900bool SPIRVInstructionSelector::selectExt(
Register ResVReg,
4901 SPIRVTypeInst ResType, MachineInstr &
I,
4902 bool IsSigned)
const {
4903 Register SrcReg =
I.getOperand(1).getReg();
4905 return selectBoolToInt(ResVReg, ResType,
I.getOperand(1).getReg(),
I,
4909 if (ResType == SrcType)
4910 return BuildCOPY(ResVReg, SrcReg,
I);
4912 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
4913 return selectUnOp(ResVReg, ResType,
I, Opcode);
4916bool SPIRVInstructionSelector::selectSUCmp(
Register ResVReg,
4917 SPIRVTypeInst ResType,
4919 bool IsSigned)
const {
4920 MachineIRBuilder MIRBuilder(
I);
4921 MachineRegisterInfo *
MRI = MIRBuilder.
getMRI();
4926 if (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4934 TII.get(IsSigned ? SPIRV::OpSLessThanEqual : SPIRV::OpULessThanEqual))
4937 .
addUse(
I.getOperand(1).getReg())
4938 .
addUse(
I.getOperand(2).getReg())
4943 TII.get(IsSigned ? SPIRV::OpSLessThan : SPIRV::OpULessThan))
4946 .
addUse(
I.getOperand(1).getReg())
4947 .
addUse(
I.getOperand(2).getReg())
4955 unsigned SelectOpcode =
4956 (
N > 1 || (ResType->
getOpcode() == SPIRV::OpTypeVectorIdEXT &&
4958 ? SPIRV::OpSelectVIVCond
4959 : SPIRV::OpSelectSISCond;
4964 .
addUse(buildOnesVal(
true, ResType,
I))
4965 .
addUse(buildZerosVal(ResType,
I))
4972 .
addUse(buildOnesVal(
false, ResType,
I))
4977bool SPIRVInstructionSelector::selectIntToBool(
Register IntReg,
4980 SPIRVTypeInst IntTy,
4981 SPIRVTypeInst BoolTy)
const {
4985 isVectorType(IntTy) ? SPIRV::OpBitwiseAndV : SPIRV::OpBitwiseAndS;
4987 Register One = buildOnesVal(
false, IntTy,
I);
4995 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpINotEqual))
5004bool SPIRVInstructionSelector::selectTrunc(
Register ResVReg,
5005 SPIRVTypeInst ResType,
5006 MachineInstr &
I)
const {
5007 Register IntReg =
I.getOperand(1).getReg();
5010 return selectIntToBool(IntReg, ResVReg,
I, ArgType, ResType);
5011 if (ArgType == ResType)
5012 return BuildCOPY(ResVReg, IntReg,
I);
5014 unsigned Opcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
5015 return selectUnOp(ResVReg, ResType,
I, Opcode);
5018bool SPIRVInstructionSelector::selectConst(
Register ResVReg,
5019 SPIRVTypeInst ResType,
5020 MachineInstr &
I)
const {
5021 unsigned Opcode =
I.getOpcode();
5022 unsigned TpOpcode = ResType->
getOpcode();
5024 if (ResType.
isPointer() || TpOpcode == SPIRV::OpTypeEvent) {
5025 assert(Opcode == TargetOpcode::G_CONSTANT &&
5026 I.getOperand(1).getCImm()->isZero());
5027 MachineBasicBlock &DepMBB =
I.getMF()->front();
5030 }
else if (TpOpcode == SPIRV::OpTypeVectorIdEXT) {
5035 "Expected <1 x T> Vector!");
5036 if (Opcode == TargetOpcode::G_FCONSTANT)
5042 }
else if (Opcode == TargetOpcode::G_FCONSTANT) {
5050 return Reg == ResVReg ?
true : BuildCOPY(ResVReg,
Reg,
I);
5053bool SPIRVInstructionSelector::selectOpUndef(
Register ResVReg,
5054 SPIRVTypeInst ResType,
5055 MachineInstr &
I)
const {
5056 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5063bool SPIRVInstructionSelector::selectInsertVal(
Register ResVReg,
5064 SPIRVTypeInst ResType,
5065 MachineInstr &
I)
const {
5067 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeInsert))
5071 .
addUse(
I.getOperand(3).getReg())
5073 .
addUse(
I.getOperand(2).getReg());
5074 for (
unsigned i = 4; i <
I.getNumOperands(); i++)
5080bool SPIRVInstructionSelector::selectExtractVal(
Register ResVReg,
5081 SPIRVTypeInst ResType,
5082 MachineInstr &
I)
const {
5083 Type *MaybeResTy =
nullptr;
5088 "Expected aggregate type for extractv instruction");
5090 SPIRV::AccessQualifier::ReadWrite,
false);
5094 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
5097 .
addUse(
I.getOperand(2).getReg());
5098 for (
unsigned i = 3; i <
I.getNumOperands(); i++)
5104bool SPIRVInstructionSelector::selectInsertElt(
Register ResVReg,
5105 SPIRVTypeInst ResType,
5106 MachineInstr &
I)
const {
5108 return selectInsertVal(ResVReg, ResType,
I);
5110 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorInsertDynamic))
5113 .
addUse(
I.getOperand(2).getReg())
5114 .
addUse(
I.getOperand(3).getReg())
5115 .
addUse(
I.getOperand(4).getReg())
5120bool SPIRVInstructionSelector::selectExtractElt(
Register ResVReg,
5121 SPIRVTypeInst ResType,
5122 MachineInstr &
I)
const {
5124 return selectExtractVal(ResVReg, ResType,
I);
5126 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpVectorExtractDynamic))
5129 .
addUse(
I.getOperand(2).getReg())
5130 .
addUse(
I.getOperand(3).getReg())
5135bool SPIRVInstructionSelector::selectGEP(
Register ResVReg,
5136 SPIRVTypeInst ResType,
5137 MachineInstr &
I)
const {
5138 const bool IsGEPInBounds =
I.getOperand(2).getImm();
5141 const bool UseUntypedPointers =
5142 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
5147 if (UseUntypedPointers) {
5149 Opcode = IsGEPInBounds ? SPIRV::OpUntypedInBoundsAccessChainKHR
5150 : SPIRV::OpUntypedAccessChainKHR;
5152 Opcode = IsGEPInBounds ? SPIRV::OpUntypedInBoundsPtrAccessChainKHR
5153 : SPIRV::OpUntypedPtrAccessChainKHR;
5162 IsGEPInBounds ? SPIRV::OpInBoundsAccessChain : SPIRV::OpAccessChain;
5164 Opcode = IsGEPInBounds ? SPIRV::OpInBoundsPtrAccessChain
5165 : SPIRV::OpPtrAccessChain;
5170 auto Res =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
5175 if (UseUntypedPointers) {
5190 while (Def &&
Def->getOpcode() == TargetOpcode::COPY &&
5191 Def->getOperand(1).isReg())
5193 if (Def &&
Def->getOpcode() == TargetOpcode::G_GLOBAL_VALUE)
5194 if (
const auto *GVar =
5197 SPIRV::AccessQualifier::ReadWrite,
5201 return diagnoseUnsupported(
5202 I,
"could not deduce the base type of an untyped access chain");
5207 Res.addUse(BaseReg);
5209 const bool IsAccessChainOpcode =
5210 (Opcode == SPIRV::OpAccessChain ||
5211 Opcode == SPIRV::OpInBoundsAccessChain ||
5212 Opcode == SPIRV::OpUntypedAccessChainKHR ||
5213 Opcode == SPIRV::OpUntypedInBoundsAccessChainKHR);
5217 "Cannot translate GEP to OpAccessChain.");
5220 const unsigned StartingIndex = IsAccessChainOpcode ? 5 : 4;
5221 for (
unsigned i = StartingIndex; i <
I.getNumExplicitOperands(); ++i)
5222 Res.addUse(
I.getOperand(i).getReg());
5223 Res.constrainAllUses(
TII,
TRI, RBI);
5232 if (Extract.
getOpcode() == SPIRV::OpCompositeExtract) {
5238 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5244 TII.get(SPIRV::OpCompositeInsert))
5258bool SPIRVInstructionSelector::wrapIntoSpecConstantOp(
5260 unsigned Lim =
I.getNumExplicitOperands();
5261 for (
unsigned i =
I.getNumExplicitDefs() + 1; i < Lim; ++i) {
5262 Register OpReg =
I.getOperand(i).getReg();
5266 OpDefine->
getOpcode() == TargetOpcode::G_ADDRSPACE_CAST ||
5267 OpDefine->
getOpcode() == TargetOpcode::G_INTTOPTR ||
5280 SPIRVTypeInst WrapType = OpType;
5281 if (OpType->
getOpcode() == SPIRV::OpTypePointer &&
5283 SPIRV::StorageClass::CodeSectionINTEL) {
5285 SPIRV::StorageClass::Function,
I);
5292 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
5293 TII.get(SPIRV::OpSpecConstantOp))
5296 .
addImm(
static_cast<uint32_t
>(SPIRV::Opcode::Bitcast))
5298 GR.
add(OpDefine, MIB);
5304bool SPIRVInstructionSelector::selectDerivativeInst(
5305 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
5306 const unsigned DPdOpCode)
const {
5309 if (!errorIfInstrOutsideShader(
I))
5315 Register SrcReg =
I.getOperand(2).getReg();
5320 return BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5323 .
addUse(
I.getOperand(2).getReg());
5325 MachineIRBuilder MIRBuilder(
I);
5328 if (componentCount != 1)
5336 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5341 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(DPdOpCode))
5346 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFConvert))
5354bool SPIRVInstructionSelector::selectIntrinsic(
Register ResVReg,
5355 SPIRVTypeInst ResType,
5356 MachineInstr &
I)
const {
5360 case Intrinsic::spv_load:
5361 return selectLoad(ResVReg, ResType,
I);
5362 case Intrinsic::spv_atomic_load:
5363 return selectAtomicLoad(ResVReg, ResType,
I);
5364 case Intrinsic::spv_store:
5365 return selectStore(
I);
5366 case Intrinsic::spv_atomic_store:
5367 return selectAtomicStore(
I);
5368 case Intrinsic::spv_extractv:
5369 return selectExtractVal(ResVReg, ResType,
I);
5370 case Intrinsic::spv_insertv:
5371 return selectInsertVal(ResVReg, ResType,
I);
5372 case Intrinsic::spv_extractelt:
5373 return selectExtractElt(ResVReg, ResType,
I);
5374 case Intrinsic::spv_insertelt:
5375 return selectInsertElt(ResVReg, ResType,
I);
5376 case Intrinsic::spv_gep:
5377 return selectGEP(ResVReg, ResType,
I);
5378 case Intrinsic::spv_bitcast: {
5379 Register OpReg =
I.getOperand(2).getReg();
5380 SPIRVTypeInst OpType =
5384 return selectOpWithSrcs(ResVReg, ResType,
I, {OpReg}, SPIRV::OpBitcast);
5386 case Intrinsic::spv_unref_global:
5387 case Intrinsic::spv_init_global: {
5393 Register GVarVReg =
MI->getOperand(0).getReg();
5394 if (!selectGlobalValue(GVarVReg, *
MI, Init))
5399 if (
MI->getOpcode() == TargetOpcode::G_GLOBAL_VALUE) {
5401 MI->eraseFromParent();
5405 case Intrinsic::spv_undef: {
5406 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
5412 case Intrinsic::spv_poison:
5413 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpPoisonKHR))
5418 case Intrinsic::spv_freeze:
5419 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpFreezeKHR))
5422 .
addUse(
I.getOperand(2).getReg())
5425 case Intrinsic::spv_named_boolean_spec_constant: {
5426 auto Opcode =
I.getOperand(3).getImm() ? SPIRV::OpSpecConstantTrue
5427 : SPIRV::OpSpecConstantFalse;
5429 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
5430 .
addDef(
I.getOperand(0).getReg())
5433 unsigned SpecId =
I.getOperand(2).getImm();
5435 SPIRV::Decoration::SpecId, {SpecId});
5439 case Intrinsic::spv_const_composite: {
5441 bool IsNull =
I.getNumExplicitDefs() + 1 ==
I.getNumExplicitOperands();
5447 if (!wrapIntoSpecConstantOp(
I, CompositeArgs))
5449 std::function<bool(
Register)> HasSpecConstOperand =
5459 for (
unsigned J =
Def->getNumExplicitDefs() + 1;
5460 J < Def->getNumExplicitOperands(); ++J) {
5461 if (
Def->getOperand(J).isReg() &&
5462 HasSpecConstOperand(
Def->getOperand(J).getReg()))
5468 bool HasSpecConst =
llvm::any_of(CompositeArgs, HasSpecConstOperand);
5469 unsigned CompositeOpc = HasSpecConst ? SPIRV::OpSpecConstantComposite
5470 : SPIRV::OpConstantComposite;
5471 unsigned ContinuedOpc = HasSpecConst
5472 ? SPIRV::OpSpecConstantCompositeContinuedINTEL
5473 : SPIRV::OpConstantCompositeContinuedINTEL;
5474 MachineIRBuilder MIR(
I);
5476 MIR, CompositeOpc, 3, ContinuedOpc, CompositeArgs, ResVReg,
5478 for (
auto *Instr : Instructions) {
5479 Instr->setDebugLoc(
I.getDebugLoc());
5484 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5491 case Intrinsic::spv_assign_name: {
5492 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpName));
5493 MIB.
addUse(
I.getOperand(
I.getNumExplicitDefs() + 1).getReg());
5494 for (
unsigned i =
I.getNumExplicitDefs() + 2;
5495 i <
I.getNumExplicitOperands(); ++i) {
5496 MIB.
addImm(
I.getOperand(i).getImm());
5501 case Intrinsic::spv_switch: {
5502 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSwitch));
5503 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5504 if (
I.getOperand(i).isReg())
5505 MIB.
addReg(
I.getOperand(i).getReg());
5506 else if (
I.getOperand(i).isCImm())
5507 addNumImm(
I.getOperand(i).getCImm()->getValue(), MIB);
5508 else if (
I.getOperand(i).isMBB())
5509 MIB.
addMBB(
I.getOperand(i).getMBB());
5516 case Intrinsic::spv_loop_merge: {
5517 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopMerge));
5518 for (
unsigned i = 1; i <
I.getNumExplicitOperands(); ++i) {
5519 if (
I.getOperand(i).isMBB())
5520 MIB.
addMBB(
I.getOperand(i).getMBB());
5527 case Intrinsic::spv_loop_control_intel: {
5529 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoopControlINTEL));
5530 for (
unsigned J = 1; J <
I.getNumExplicitOperands(); ++J)
5535 case Intrinsic::spv_selection_merge: {
5537 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSelectionMerge));
5538 assert(
I.getOperand(1).isMBB() &&
5539 "operand 1 to spv_selection_merge must be a basic block");
5540 MIB.
addMBB(
I.getOperand(1).getMBB());
5541 MIB.
addImm(getSelectionOperandForImm(
I.getOperand(2).getImm()));
5545 case Intrinsic::spv_cmpxchg:
5546 return selectAtomicCmpXchg(ResVReg, ResType,
I);
5547 case Intrinsic::spv_unreachable:
5548 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUnreachable))
5551 case Intrinsic::spv_abort:
5552 return selectAbort(
I);
5553 case Intrinsic::spv_alloca:
5554 return selectFrameIndex(ResVReg, ResType,
I);
5555 case Intrinsic::spv_alloca_array:
5556 return selectAllocaArray(ResVReg, ResType,
I);
5557 case Intrinsic::spv_assume:
5559 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAssumeTrueKHR))
5560 .
addUse(
I.getOperand(1).getReg())
5565 case Intrinsic::spv_expect:
5567 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExpectKHR))
5570 .
addUse(
I.getOperand(2).getReg())
5571 .
addUse(
I.getOperand(3).getReg())
5576 case Intrinsic::arithmetic_fence:
5577 if (STI.
canUseExtension(SPIRV::Extension::SPV_EXT_arithmetic_fence)) {
5578 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpArithmeticFenceEXT))
5581 .
addUse(
I.getOperand(2).getReg())
5585 return BuildCOPY(ResVReg,
I.getOperand(2).getReg(),
I);
5587 case Intrinsic::spv_thread_id:
5593 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalInvocationId, ResVReg,
5595 case Intrinsic::spv_thread_id_in_group:
5601 return loadVec3BuiltinInputID(SPIRV::BuiltIn::LocalInvocationId, ResVReg,
5603 case Intrinsic::spv_group_id:
5609 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupId, ResVReg, ResType,
5611 case Intrinsic::spv_flattened_thread_id_in_group:
5618 return loadBuiltinInputID(SPIRV::BuiltIn::LocalInvocationIndex, ResVReg,
5620 case Intrinsic::spv_workgroup_size:
5621 return loadVec3BuiltinInputID(SPIRV::BuiltIn::WorkgroupSize, ResVReg,
5623 case Intrinsic::spv_global_size:
5624 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalSize, ResVReg, ResType,
5626 case Intrinsic::spv_global_offset:
5627 return loadVec3BuiltinInputID(SPIRV::BuiltIn::GlobalOffset, ResVReg,
5629 case Intrinsic::spv_num_workgroups:
5630 return loadVec3BuiltinInputID(SPIRV::BuiltIn::NumWorkgroups, ResVReg,
5632 case Intrinsic::spv_subgroup_size:
5633 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupSize, ResVReg, ResType,
5635 case Intrinsic::spv_num_subgroups:
5636 return loadBuiltinInputID(SPIRV::BuiltIn::NumSubgroups, ResVReg, ResType,
5638 case Intrinsic::spv_subgroup_id:
5639 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupId, ResVReg, ResType,
I);
5640 case Intrinsic::spv_subgroup_local_invocation_id:
5641 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupLocalInvocationId,
5642 ResVReg, ResType,
I);
5643 case Intrinsic::spv_subgroup_max_size:
5644 return loadBuiltinInputID(SPIRV::BuiltIn::SubgroupMaxSize, ResVReg, ResType,
5646 case Intrinsic::spv_fdot:
5647 return selectFloatDot(ResVReg, ResType,
I);
5648 case Intrinsic::spv_udot:
5649 case Intrinsic::spv_sdot:
5650 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5652 return selectIntegerDot(ResVReg, ResType,
I,
5653 IID == Intrinsic::spv_sdot);
5654 return selectIntegerDotExpansion(ResVReg, ResType,
I);
5655 case Intrinsic::spv_dot4add_i8packed:
5656 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5658 return selectDot4AddPacked<true>(ResVReg, ResType,
I);
5659 return selectDot4AddPackedExpansion<true>(ResVReg, ResType,
I);
5660 case Intrinsic::spv_dot4add_u8packed:
5661 if (STI.
canUseExtension(SPIRV::Extension::SPV_KHR_integer_dot_product) ||
5663 return selectDot4AddPacked<false>(ResVReg, ResType,
I);
5664 return selectDot4AddPackedExpansion<false>(ResVReg, ResType,
I);
5665 case Intrinsic::spv_all:
5666 return selectAll(ResVReg, ResType,
I);
5667 case Intrinsic::spv_any:
5668 return selectAny(ResVReg, ResType,
I);
5669 case Intrinsic::spv_distance:
5670 return selectExtInst(ResVReg, ResType,
I, CL::distance, GL::Distance);
5671 case Intrinsic::spv_lerp:
5672 return selectExtInst(ResVReg, ResType,
I, CL::mix, GL::FMix);
5673 case Intrinsic::spv_length:
5674 return selectExtInst(ResVReg, ResType,
I, CL::length, GL::Length);
5675 case Intrinsic::spv_degrees:
5676 return selectExtInst(ResVReg, ResType,
I, CL::degrees, GL::Degrees);
5677 case Intrinsic::spv_faceforward:
5678 return selectExtInst(ResVReg, ResType,
I, GL::FaceForward);
5679 case Intrinsic::spv_frac:
5680 return selectExtInst(ResVReg, ResType,
I, CL::fract, GL::Fract);
5681 case Intrinsic::spv_isinf:
5682 return selectOpIsInf(ResVReg, ResType,
I);
5683 case Intrinsic::spv_isnan:
5684 return selectOpIsNan(ResVReg, ResType,
I);
5685 case Intrinsic::spv_isfinite:
5686 return selectOpIsFinite(ResVReg, ResType,
I);
5687 case Intrinsic::spv_isnormal:
5688 return selectOpIsNormal(ResVReg, ResType,
I);
5689 case Intrinsic::spv_normalize:
5690 return selectExtInst(ResVReg, ResType,
I, CL::normalize, GL::Normalize);
5691 case Intrinsic::spv_refract:
5692 return selectExtInst(ResVReg, ResType,
I, GL::Refract);
5693 case Intrinsic::spv_reflect:
5694 return selectExtInst(ResVReg, ResType,
I, GL::Reflect);
5695 case Intrinsic::spv_rsqrt:
5696 return selectExtInst(ResVReg, ResType,
I, CL::rsqrt, GL::InverseSqrt);
5697 case Intrinsic::spv_sign:
5698 return selectSign(ResVReg, ResType,
I);
5699 case Intrinsic::spv_smoothstep:
5700 return selectExtInst(ResVReg, ResType,
I, CL::smoothstep, GL::SmoothStep);
5701 case Intrinsic::spv_firstbituhigh:
5702 return selectFirstBitHigh(ResVReg, ResType,
I,
false);
5703 case Intrinsic::spv_firstbitshigh:
5704 return selectFirstBitHigh(ResVReg, ResType,
I,
true);
5705 case Intrinsic::spv_firstbitlow:
5706 return selectFirstBitLow(ResVReg, ResType,
I);
5707 case Intrinsic::spv_all_memory_barrier:
5708 return selectBarrierInst(
I, SPIRV::Scope::Device,
5709 SPIRV::MemorySemantics::UniformMemory |
5710 SPIRV::MemorySemantics::ImageMemory |
5711 SPIRV::MemorySemantics::WorkgroupMemory,
5713 case Intrinsic::spv_all_memory_barrier_with_group_sync:
5714 return selectBarrierInst(
I, SPIRV::Scope::Device,
5715 SPIRV::MemorySemantics::UniformMemory |
5716 SPIRV::MemorySemantics::ImageMemory |
5717 SPIRV::MemorySemantics::WorkgroupMemory,
5719 case Intrinsic::spv_device_memory_barrier:
5720 return selectBarrierInst(
I, SPIRV::Scope::Device,
5721 SPIRV::MemorySemantics::UniformMemory |
5722 SPIRV::MemorySemantics::ImageMemory,
5724 case Intrinsic::spv_device_memory_barrier_with_group_sync:
5725 return selectBarrierInst(
I, SPIRV::Scope::Device,
5726 SPIRV::MemorySemantics::UniformMemory |
5727 SPIRV::MemorySemantics::ImageMemory,
5729 case Intrinsic::spv_group_memory_barrier:
5730 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5731 SPIRV::MemorySemantics::WorkgroupMemory,
5733 case Intrinsic::spv_group_memory_barrier_with_group_sync:
5734 return selectBarrierInst(
I, SPIRV::Scope::Workgroup,
5735 SPIRV::MemorySemantics::WorkgroupMemory,
5737 case Intrinsic::spv_generic_cast_to_ptr_explicit: {
5738 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 1).getReg();
5739 SPIRV::StorageClass::StorageClass ResSC =
5742 return diagnoseUnsupported(
I,
"The target storage class is not castable "
5743 "from the Generic storage class");
5744 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpGenericCastToPtrExplicit))
5752 case Intrinsic::spv_lifetime_start:
5753 case Intrinsic::spv_lifetime_end: {
5754 unsigned Op = IID == Intrinsic::spv_lifetime_start ? SPIRV::OpLifetimeStart
5755 : SPIRV::OpLifetimeStop;
5756 int64_t
Size =
I.getOperand(
I.getNumExplicitDefs() + 1).getImm();
5757 Register PtrReg =
I.getOperand(
I.getNumExplicitDefs() + 2).getReg();
5766 case Intrinsic::spv_saturate:
5767 return selectSaturate(ResVReg, ResType,
I);
5768 case Intrinsic::spv_nclamp:
5769 return selectExtInst(ResVReg, ResType,
I, CL::fclamp, GL::NClamp);
5770 case Intrinsic::spv_uclamp:
5771 return selectExtInst(ResVReg, ResType,
I, CL::u_clamp, GL::UClamp);
5772 case Intrinsic::spv_sclamp:
5773 return selectExtInst(ResVReg, ResType,
I, CL::s_clamp, GL::SClamp);
5774 case Intrinsic::spv_subgroup_prefix_bit_count:
5775 return selectWavePrefixBitCount(ResVReg, ResType,
I);
5776 case Intrinsic::spv_wave_active_countbits:
5777 return selectWaveActiveCountBits(ResVReg, ResType,
I);
5778 case Intrinsic::spv_wave_all_equal:
5779 return selectWaveActiveAllEqual(ResVReg, ResType,
I);
5780 case Intrinsic::spv_wave_all:
5781 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAll);
5782 case Intrinsic::spv_wave_any:
5783 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformAny);
5784 case Intrinsic::spv_subgroup_ballot:
5785 return selectWaveOpInst(ResVReg, ResType,
I,
5786 SPIRV::OpGroupNonUniformBallot);
5787 case Intrinsic::spv_wave_is_first_lane:
5788 return selectWaveOpInst(ResVReg, ResType,
I, SPIRV::OpGroupNonUniformElect);
5789 case Intrinsic::spv_wave_reduce_or:
5790 return selectWaveReduceOp(ResVReg, ResType,
I,
5791 SPIRV::OpGroupNonUniformBitwiseOr);
5792 case Intrinsic::spv_wave_reduce_xor:
5793 return selectWaveReduceOp(ResVReg, ResType,
I,
5794 SPIRV::OpGroupNonUniformBitwiseXor);
5795 case Intrinsic::spv_wave_reduce_and:
5796 return selectWaveReduceOp(ResVReg, ResType,
I,
5797 SPIRV::OpGroupNonUniformBitwiseAnd);
5798 case Intrinsic::spv_wave_reduce_umax:
5799 return selectWaveReduceMax(ResVReg, ResType,
I,
true);
5800 case Intrinsic::spv_wave_reduce_max:
5801 return selectWaveReduceMax(ResVReg, ResType,
I,
false);
5802 case Intrinsic::spv_wave_reduce_umin:
5803 return selectWaveReduceMin(ResVReg, ResType,
I,
true);
5804 case Intrinsic::spv_wave_reduce_min:
5805 return selectWaveReduceMin(ResVReg, ResType,
I,
false);
5806 case Intrinsic::spv_wave_reduce_sum:
5807 return selectWaveReduceSum(ResVReg, ResType,
I);
5808 case Intrinsic::spv_wave_product:
5809 return selectWaveReduceProduct(ResVReg, ResType,
I);
5810 case Intrinsic::spv_wave_readlane:
5811 return selectWaveOpInst(ResVReg, ResType,
I,
5812 SPIRV::OpGroupNonUniformShuffle);
5813 case Intrinsic::spv_wave_readlane_first:
5814 return selectWaveOpInst(ResVReg, ResType,
I,
5815 SPIRV::OpGroupNonUniformBroadcastFirst);
5816 case Intrinsic::spv_wave_prefix_sum:
5817 return selectWaveExclusiveScanSum(ResVReg, ResType,
I);
5818 case Intrinsic::spv_wave_prefix_product:
5819 return selectWaveExclusiveScanProduct(ResVReg, ResType,
I);
5820 case Intrinsic::spv_quad_read_across_x: {
5821 return selectQuadSwap(ResVReg, ResType,
I, 0);
5823 case Intrinsic::spv_quad_read_across_y: {
5824 return selectQuadSwap(ResVReg, ResType,
I, 1);
5826 case Intrinsic::spv_quad_read_across_diagonal: {
5827 return selectQuadSwap(ResVReg, ResType,
I, 2);
5829 case Intrinsic::spv_radians:
5830 return selectExtInst(ResVReg, ResType,
I, CL::radians, GL::Radians);
5834 case Intrinsic::instrprof_increment:
5835 case Intrinsic::instrprof_increment_step:
5836 case Intrinsic::instrprof_value_profile:
5839 case Intrinsic::spv_value_md:
5841 case Intrinsic::spv_resource_handlefrombinding: {
5842 return selectHandleFromBinding(ResVReg, ResType,
I);
5844 case Intrinsic::spv_resource_counterhandlefrombinding:
5845 return selectCounterHandleFromBinding(ResVReg, ResType,
I);
5846 case Intrinsic::spv_resource_updatecounter:
5847 return selectUpdateCounter(ResVReg, ResType,
I);
5848 case Intrinsic::spv_resource_store_typedbuffer: {
5849 return selectImageWriteIntrinsic(
I);
5851 case Intrinsic::spv_resource_load_typedbuffer: {
5852 return selectReadImageIntrinsic(ResVReg, ResType,
I);
5854 case Intrinsic::spv_resource_load_level: {
5855 return selectLoadLevelIntrinsic(ResVReg, ResType,
I);
5857 case Intrinsic::spv_resource_getdimensions_x:
5858 case Intrinsic::spv_resource_getdimensions_xy:
5859 case Intrinsic::spv_resource_getdimensions_xyz: {
5860 return selectGetDimensionsIntrinsic(ResVReg, ResType,
I);
5862 case Intrinsic::spv_resource_getdimensions_levels_x:
5863 case Intrinsic::spv_resource_getdimensions_levels_xy:
5864 case Intrinsic::spv_resource_getdimensions_levels_xyz: {
5865 return selectGetDimensionsLevelsIntrinsic(ResVReg, ResType,
I);
5867 case Intrinsic::spv_resource_getdimensions_ms_xy:
5868 case Intrinsic::spv_resource_getdimensions_ms_xyz: {
5869 return selectGetDimensionsMSIntrinsic(ResVReg, ResType,
I);
5871 case Intrinsic::spv_resource_calculate_lod:
5872 case Intrinsic::spv_resource_calculate_lod_unclamped:
5873 return selectCalculateLodIntrinsic(ResVReg, ResType,
I);
5874 case Intrinsic::spv_resource_sample:
5875 case Intrinsic::spv_resource_sample_clamp:
5876 return selectSampleBasicIntrinsic(ResVReg, ResType,
I);
5877 case Intrinsic::spv_resource_samplebias:
5878 case Intrinsic::spv_resource_samplebias_clamp:
5879 return selectSampleBiasIntrinsic(ResVReg, ResType,
I);
5880 case Intrinsic::spv_resource_samplegrad:
5881 case Intrinsic::spv_resource_samplegrad_clamp:
5882 return selectSampleGradIntrinsic(ResVReg, ResType,
I);
5883 case Intrinsic::spv_resource_samplelevel:
5884 return selectSampleLevelIntrinsic(ResVReg, ResType,
I);
5885 case Intrinsic::spv_resource_samplecmp:
5886 case Intrinsic::spv_resource_samplecmp_clamp:
5887 return selectSampleCmpIntrinsic(ResVReg, ResType,
I);
5888 case Intrinsic::spv_resource_samplecmplevelzero:
5889 return selectSampleCmpLevelZeroIntrinsic(ResVReg, ResType,
I);
5890 case Intrinsic::spv_resource_gather:
5891 case Intrinsic::spv_resource_gather_cmp:
5892 return selectGatherIntrinsic(ResVReg, ResType,
I);
5893 case Intrinsic::spv_resource_getbasepointer:
5894 case Intrinsic::spv_resource_getpointer: {
5895 return selectResourceGetPointer(ResVReg, ResType,
I);
5897 case Intrinsic::spv_pushconstant_getpointer: {
5898 return selectPushConstantGetPointer(ResVReg, ResType,
I);
5900 case Intrinsic::spv_discard: {
5901 return selectDiscard(ResVReg, ResType,
I);
5903 case Intrinsic::spv_resource_nonuniformindex: {
5904 return selectResourceNonUniformIndex(ResVReg, ResType,
I);
5906 case Intrinsic::spv_unpackhalf2x16: {
5907 return selectExtInst(ResVReg, ResType,
I, GL::UnpackHalf2x16);
5909 case Intrinsic::spv_packhalf2x16: {
5910 return selectExtInst(ResVReg, ResType,
I, GL::PackHalf2x16);
5912 case Intrinsic::spv_ddx:
5913 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdx);
5914 case Intrinsic::spv_ddy:
5915 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdy);
5916 case Intrinsic::spv_ddx_coarse:
5917 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxCoarse);
5918 case Intrinsic::spv_ddy_coarse:
5919 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyCoarse);
5920 case Intrinsic::spv_ddx_fine:
5921 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdxFine);
5922 case Intrinsic::spv_ddy_fine:
5923 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpDPdyFine);
5924 case Intrinsic::spv_fwidth:
5925 return selectDerivativeInst(ResVReg, ResType,
I, SPIRV::OpFwidth);
5926 case Intrinsic::spv_masked_gather:
5927 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5928 return selectMaskedGather(ResVReg, ResType,
I);
5929 return diagnoseUnsupported(
5930 I,
"llvm.masked.gather requires SPV_INTEL_masked_gather_scatter");
5931 case Intrinsic::spv_masked_scatter:
5932 if (STI.
canUseExtension(SPIRV::Extension::SPV_INTEL_masked_gather_scatter))
5933 return selectMaskedScatter(
I);
5934 return diagnoseUnsupported(
5935 I,
"llvm.masked.scatter requires SPV_INTEL_masked_gather_scatter");
5936 case Intrinsic::returnaddress:
5937 case Intrinsic::frameaddress: {
5939 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpConstantNull))
5946 return diagnoseUnsupported(
I,
"intrinsic selection not implemented.");
5951bool SPIRVInstructionSelector::selectHandleFromBinding(
Register &ResVReg,
5952 SPIRVTypeInst ResType,
5953 MachineInstr &
I)
const {
5958bool SPIRVInstructionSelector::selectCounterHandleFromBinding(
5959 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
5961 assert(Intr.getIntrinsicID() ==
5962 Intrinsic::spv_resource_counterhandlefrombinding);
5965 Register MainHandleReg = Intr.getOperand(2).getReg();
5967 assert(MainHandleDef->getIntrinsicID() ==
5968 Intrinsic::spv_resource_handlefrombinding);
5972 uint32_t ArraySize =
getIConstVal(MainHandleDef->getOperand(4).getReg(),
MRI);
5973 Register IndexReg = MainHandleDef->getOperand(5).getReg();
5974 std::string CounterName =
5979 MachineIRBuilder MIRBuilder(
I);
5981 buildPointerToResource(SPIRVTypeInst(GR.
getPointeeType(ResType)),
5983 ArraySize, IndexReg, CounterName, MIRBuilder);
5985 return BuildCOPY(ResVReg, CounterVarReg,
I);
5988bool SPIRVInstructionSelector::selectUpdateCounter(
Register &ResVReg,
5989 SPIRVTypeInst ResType,
5990 MachineInstr &
I)
const {
5992 assert(Intr.getIntrinsicID() == Intrinsic::spv_resource_updatecounter);
5994 Register CounterHandleReg = Intr.getOperand(2).getReg();
5995 int64_t IncrVal = Intr.getOperand(3).getImm();
6002 SPIRVTypeInst CounterVarPointeeType = GR.
getPointeeType(CounterVarType);
6003 assert(CounterVarPointeeType &&
6004 CounterVarPointeeType->
getOpcode() == SPIRV::OpTypeStruct &&
6005 "Counter variable must be a struct");
6007 SPIRV::StorageClass::StorageBuffer &&
6008 "Counter variable must be in the storage buffer storage class");
6010 "Counter variable must have exactly 1 member in the struct");
6011 const SPIRVTypeInst MemberType =
6014 "Counter variable struct must have a single i32 member");
6018 MachineIRBuilder MIRBuilder(
I);
6020 Type::getInt32Ty(
I.getMF()->getFunction().getContext());
6023 LLVMIntType, MIRBuilder, SPIRV::StorageClass::StorageBuffer);
6029 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6032 .
addUse(CounterHandleReg)
6038 Register Scope = buildI32ConstantInEntryBlock(SPIRV::Scope::Device,
I);
6040 buildI32ConstantInEntryBlock(SPIRV::MemorySemantics::None,
I);
6043 buildI32ConstantInEntryBlock(
static_cast<uint32_t
>(IncrVal),
I);
6046 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAtomicIAdd))
6055 return BuildCOPY(ResVReg, AtomicRes,
I);
6063 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpIAddS))
6071bool SPIRVInstructionSelector::selectReadImageIntrinsic(
Register &ResVReg,
6072 SPIRVTypeInst ResType,
6073 MachineInstr &
I)
const {
6081 Register ImageReg =
I.getOperand(2).getReg();
6089 Register IdxReg =
I.getOperand(3).getReg();
6091 MachineInstr &Pos =
I;
6093 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, IdxReg, Loc,
6097bool SPIRVInstructionSelector::generateSampleImage(
6100 DebugLoc Loc, MachineInstr &Pos)
const {
6111 if (!loadHandleBeforePosition(NewSamplerReg,
6117 MachineIRBuilder MIRBuilder(Pos);
6130 bool IsExplicitLod = ImOps.GradX.has_value() || ImOps.GradY.has_value() ||
6131 ImOps.Lod.has_value();
6132 unsigned Opcode = IsExplicitLod ? SPIRV::OpImageSampleExplicitLod
6133 : SPIRV::OpImageSampleImplicitLod;
6135 Opcode = IsExplicitLod ? SPIRV::OpImageSampleDrefExplicitLod
6136 : SPIRV::OpImageSampleDrefImplicitLod;
6145 MIB.
addUse(*ImOps.Compare);
6147 uint32_t ImageOperands = 0;
6149 ImageOperands |= SPIRV::ImageOperand::Bias;
6151 ImageOperands |= SPIRV::ImageOperand::Lod;
6152 if (ImOps.GradX && ImOps.GradY)
6153 ImageOperands |= SPIRV::ImageOperand::Grad;
6154 if (ImOps.Offset && !isScalarOrVectorIntConstantZero(*ImOps.Offset)) {
6156 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6159 "Non-constant offsets are not supported in sample instructions.");
6164 ImageOperands |= SPIRV::ImageOperand::MinLod;
6166 if (ImageOperands != 0) {
6167 MIB.
addImm(ImageOperands);
6168 if (ImageOperands & SPIRV::ImageOperand::Bias)
6170 if (ImageOperands & SPIRV::ImageOperand::Lod)
6172 if (ImageOperands & SPIRV::ImageOperand::Grad) {
6173 MIB.
addUse(*ImOps.GradX);
6174 MIB.
addUse(*ImOps.GradY);
6177 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6178 MIB.
addUse(*ImOps.Offset);
6179 if (ImageOperands & SPIRV::ImageOperand::MinLod)
6180 MIB.
addUse(*ImOps.MinLod);
6187bool SPIRVInstructionSelector::selectImageQuerySize(
6189 std::optional<Register> LodReg)
const {
6192 "ImageReg is not an image type.");
6194 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6196 unsigned NumComponents = 0;
6198 case SPIRV::Dim::DIM_1D:
6199 case SPIRV::Dim::DIM_Buffer:
6200 NumComponents =
IsArray ? 2 : 1;
6202 case SPIRV::Dim::DIM_2D:
6203 case SPIRV::Dim::DIM_Cube:
6204 case SPIRV::Dim::DIM_Rect:
6205 NumComponents =
IsArray ? 3 : 2;
6207 case SPIRV::Dim::DIM_3D:
6211 I.emitGenericError(
"Unsupported image dimension for OpImageQuerySize.");
6221 if (!LodReg && IsSampled && !IsMultisampled &&
6222 Dim != SPIRV::Dim::DIM_Buffer && Dim != SPIRV::Dim::DIM_Rect)
6225 LodReg ? SPIRV::OpImageQuerySizeLod : SPIRV::OpImageQuerySize;
6226 SPIRVTypeInst ResType =
6231 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6241bool SPIRVInstructionSelector::selectGetDimensionsIntrinsic(
6242 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6243 Register ImageReg =
I.getOperand(2).getReg();
6250 return selectImageQuerySize(NewImageReg, ResVReg,
I);
6253bool SPIRVInstructionSelector::selectGetDimensionsLevelsIntrinsic(
6254 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6255 Register ImageReg =
I.getOperand(2).getReg();
6264 Register LodReg =
I.getOperand(3).getReg();
6267 "OpImageQuerySizeLod and OpImageQueryLevels require a sampled image");
6269 if (!selectImageQuerySize(NewImageReg, SizeReg,
I, LodReg)) {
6276 TII.get(SPIRV::OpImageQueryLevels))
6283 TII.get(SPIRV::OpCompositeConstruct))
6293bool SPIRVInstructionSelector::selectGetDimensionsMSIntrinsic(
6294 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6295 Register ImageReg =
I.getOperand(2).getReg();
6306 "OpImageQuerySamples requires a multisampled image");
6308 if (!selectImageQuerySize(NewImageReg, SizeReg,
I)) {
6316 TII.get(SPIRV::OpImageQuerySamples))
6323 TII.get(SPIRV::OpCompositeConstruct))
6333bool SPIRVInstructionSelector::selectCalculateLodIntrinsic(
6334 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6335 Register ImageReg =
I.getOperand(2).getReg();
6336 Register SamplerReg =
I.getOperand(3).getReg();
6337 Register CoordinateReg =
I.getOperand(4).getReg();
6353 if (!loadHandleBeforePosition(
6358 MachineIRBuilder MIRBuilder(
I);
6364 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6374 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageQueryLod))
6381 unsigned ExtractedIndex =
6383 Intrinsic::spv_resource_calculate_lod_unclamped
6387 MachineInstrBuilder MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6388 TII.get(SPIRV::OpCompositeExtract))
6398bool SPIRVInstructionSelector::selectSampleBasicIntrinsic(
6399 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6400 Register ImageReg =
I.getOperand(2).getReg();
6401 Register SamplerReg =
I.getOperand(3).getReg();
6402 Register CoordinateReg =
I.getOperand(4).getReg();
6403 ImageOperands ImOps;
6404 if (
I.getNumOperands() > 5)
6405 ImOps.Offset =
I.getOperand(5).getReg();
6406 if (
I.getNumOperands() > 6)
6407 ImOps.MinLod =
I.getOperand(6).getReg();
6408 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6409 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6412bool SPIRVInstructionSelector::selectSampleBiasIntrinsic(
6413 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6414 Register ImageReg =
I.getOperand(2).getReg();
6415 Register SamplerReg =
I.getOperand(3).getReg();
6416 Register CoordinateReg =
I.getOperand(4).getReg();
6417 ImageOperands ImOps;
6418 ImOps.Bias =
I.getOperand(5).getReg();
6419 if (
I.getNumOperands() > 6)
6420 ImOps.Offset =
I.getOperand(6).getReg();
6421 if (
I.getNumOperands() > 7)
6422 ImOps.MinLod =
I.getOperand(7).getReg();
6423 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6424 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6427bool SPIRVInstructionSelector::selectSampleGradIntrinsic(
6428 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6429 Register ImageReg =
I.getOperand(2).getReg();
6430 Register SamplerReg =
I.getOperand(3).getReg();
6431 Register CoordinateReg =
I.getOperand(4).getReg();
6432 ImageOperands ImOps;
6433 ImOps.GradX =
I.getOperand(5).getReg();
6434 ImOps.GradY =
I.getOperand(6).getReg();
6435 if (
I.getNumOperands() > 7)
6436 ImOps.Offset =
I.getOperand(7).getReg();
6437 if (
I.getNumOperands() > 8)
6438 ImOps.MinLod =
I.getOperand(8).getReg();
6439 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6440 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6443bool SPIRVInstructionSelector::selectSampleLevelIntrinsic(
6444 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6445 Register ImageReg =
I.getOperand(2).getReg();
6446 Register SamplerReg =
I.getOperand(3).getReg();
6447 Register CoordinateReg =
I.getOperand(4).getReg();
6448 ImageOperands ImOps;
6449 ImOps.Lod =
I.getOperand(5).getReg();
6450 if (
I.getNumOperands() > 6)
6451 ImOps.Offset =
I.getOperand(6).getReg();
6452 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6453 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6456bool SPIRVInstructionSelector::selectSampleCmpIntrinsic(
Register &ResVReg,
6457 SPIRVTypeInst ResType,
6458 MachineInstr &
I)
const {
6459 Register ImageReg =
I.getOperand(2).getReg();
6460 Register SamplerReg =
I.getOperand(3).getReg();
6461 Register CoordinateReg =
I.getOperand(4).getReg();
6462 ImageOperands ImOps;
6463 ImOps.Compare =
I.getOperand(5).getReg();
6464 if (
I.getNumOperands() > 6)
6465 ImOps.Offset =
I.getOperand(6).getReg();
6466 if (
I.getNumOperands() > 7)
6467 ImOps.MinLod =
I.getOperand(7).getReg();
6468 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6469 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6472bool SPIRVInstructionSelector::selectLoadLevelIntrinsic(
Register &ResVReg,
6473 SPIRVTypeInst ResType,
6474 MachineInstr &
I)
const {
6475 Register ImageReg =
I.getOperand(2).getReg();
6476 Register CoordinateReg =
I.getOperand(3).getReg();
6477 Register LodReg =
I.getOperand(4).getReg();
6479 ImageOperands ImOps;
6481 if (
I.getNumOperands() > 5)
6482 ImOps.Offset =
I.getOperand(5).getReg();
6494 return generateImageReadOrFetch(ResVReg, ResType, NewImageReg, CoordinateReg,
6495 I.getDebugLoc(),
I, &ImOps);
6498bool SPIRVInstructionSelector::selectSampleCmpLevelZeroIntrinsic(
6499 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6500 Register ImageReg =
I.getOperand(2).getReg();
6501 Register SamplerReg =
I.getOperand(3).getReg();
6502 Register CoordinateReg =
I.getOperand(4).getReg();
6503 ImageOperands ImOps;
6504 ImOps.Compare =
I.getOperand(5).getReg();
6505 if (
I.getNumOperands() > 6)
6506 ImOps.Offset =
I.getOperand(6).getReg();
6509 return generateSampleImage(ResVReg, ResType, ImageReg, SamplerReg,
6510 CoordinateReg, ImOps,
I.getDebugLoc(),
I);
6513bool SPIRVInstructionSelector::selectGatherIntrinsic(
Register &ResVReg,
6514 SPIRVTypeInst ResType,
6515 MachineInstr &
I)
const {
6516 Register ImageReg =
I.getOperand(2).getReg();
6517 Register SamplerReg =
I.getOperand(3).getReg();
6518 Register CoordinateReg =
I.getOperand(4).getReg();
6521 "ImageReg is not an image type.");
6526 ComponentOrCompareReg =
I.getOperand(5).getReg();
6527 OffsetReg =
I.getOperand(6).getReg();
6530 if (!loadHandleBeforePosition(NewImageReg, ImageType, *ImageDef,
I)) {
6534 auto Dim =
static_cast<SPIRV::Dim::Dim
>(ImageType->
getOperand(2).
getImm());
6535 if (Dim != SPIRV::Dim::DIM_2D && Dim != SPIRV::Dim::DIM_Cube &&
6536 Dim != SPIRV::Dim::DIM_Rect) {
6538 "Gather operations are only supported for 2D, Cube, and Rect images.");
6545 if (!loadHandleBeforePosition(
6550 MachineIRBuilder MIRBuilder(
I);
6551 SPIRVTypeInst SampledImageType =
6556 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpSampledImage))
6564 bool IsGatherCmp =
IntrId == Intrinsic::spv_resource_gather_cmp;
6566 IsGatherCmp ? SPIRV::OpImageDrefGather : SPIRV::OpImageGather;
6568 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(Opcode))
6573 .
addUse(ComponentOrCompareReg);
6575 uint32_t ImageOperands = 0;
6576 if (OffsetReg && !isScalarOrVectorIntConstantZero(OffsetReg)) {
6577 if (Dim == SPIRV::Dim::DIM_Cube) {
6579 "Gather operations with offset are not supported for Cube images.");
6583 ImageOperands |= SPIRV::ImageOperand::ConstOffset;
6585 ImageOperands |= SPIRV::ImageOperand::Offset;
6589 if (ImageOperands != 0) {
6590 MIB.
addImm(ImageOperands);
6592 (SPIRV::ImageOperand::ConstOffset | SPIRV::ImageOperand::Offset))
6600bool SPIRVInstructionSelector::generateImageReadOrFetch(
6603 const ImageOperands *ImOps)
const {
6606 "ImageReg is not an image type.");
6608 bool IsSignedInteger =
6613 bool IsFetch = (SampledOp.getImm() == 1);
6615 auto AddOperands = [&](MachineInstrBuilder &MIB) {
6616 uint32_t ImageOperandsMask = 0;
6617 if (IsSignedInteger)
6618 ImageOperandsMask |= 0x1000;
6620 if (IsFetch && ImOps) {
6622 ImageOperandsMask |= SPIRV::ImageOperand::Lod;
6623 if (ImOps->Offset && !isScalarOrVectorIntConstantZero(*ImOps->Offset)) {
6625 ImageOperandsMask |= SPIRV::ImageOperand::ConstOffset;
6627 ImageOperandsMask |= SPIRV::ImageOperand::Offset;
6631 if (ImageOperandsMask != 0) {
6632 MIB.
addImm(ImageOperandsMask);
6633 if (IsFetch && ImOps) {
6636 if (ImOps->Offset &&
6637 (ImageOperandsMask &
6638 (SPIRV::ImageOperand::Offset | SPIRV::ImageOperand::ConstOffset)))
6639 MIB.
addUse(*ImOps->Offset);
6648 SPIRVTypeInst SampledType =
6651 SPIRVTypeInst ReadType =
6652 widenTypeToVec4(IsPacked ? SampledType : ResType, Pos);
6653 bool ReadTypeMatchesResult = ReadType == ResType;
6655 Register ReadReg = ReadTypeMatchesResult
6661 TII.get(IsFetch ? SPIRV::OpImageFetch : SPIRV::OpImageRead))
6667 BMI.constrainAllUses(
TII,
TRI, RBI);
6669 if (ReadTypeMatchesResult)
6682 if (ResultSize == 1) {
6691 return extractSubvector(ResVReg, ResType, ReadReg, Pos);
6694bool SPIRVInstructionSelector::selectResourceGetPointer(
Register &ResVReg,
6695 SPIRVTypeInst ResType,
6696 MachineInstr &
I)
const {
6697 Register ResourcePtr =
I.getOperand(2).getReg();
6699 if (
RegType->getOpcode() == SPIRV::OpTypeImage) {
6708 MachineIRBuilder MIRBuilder(
I);
6713 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAccessChain))
6719 if (
I.getNumExplicitOperands() > 3) {
6720 Register IndexReg =
I.getOperand(3).getReg();
6727bool SPIRVInstructionSelector::selectPushConstantGetPointer(
6728 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6733bool SPIRVInstructionSelector::selectResourceNonUniformIndex(
6734 Register &ResVReg, SPIRVTypeInst ResType, MachineInstr &
I)
const {
6735 Register ObjReg =
I.getOperand(2).getReg();
6736 if (!BuildCOPY(ResVReg, ObjReg,
I))
6746 decorateUsesAsNonUniform(ResVReg);
6750void SPIRVInstructionSelector::decorateUsesAsNonUniform(
6753 {NonUniformReg,
nullptr}};
6754 llvm::SmallSet<Register, 8> Visited;
6755 while (WorkList.
size() > 0) {
6758 if (!Visited.
insert(CurrentReg).second)
6761 bool IsDecorated =
false;
6763 if (
Use.getOpcode() == SPIRV::OpDecorate &&
6764 Use.getOperand(1).getImm() == SPIRV::Decoration::NonUniformEXT) {
6770 if (
Use.getOperand(0).isReg() &&
Use.getOperand(0).isDef()) {
6772 if (ResultReg == CurrentReg)
6780 MachineInstr &InsertPt =
6783 SPIRV::Decoration::NonUniformEXT, {});
6788bool SPIRVInstructionSelector::extractSubvector(
6790 MachineInstr &InsertionPoint)
const {
6792 [[maybe_unused]]
uint64_t InputSize =
6795 [[maybe_unused]]
bool IsLongVectorEXT =
6797 assert((InputSize > 1 || IsLongVectorEXT) &&
"The input must be a vector.");
6798 assert((ResultSize > 1 || IsLongVectorEXT) &&
"The result must be a vector.");
6799 assert(ResultSize < InputSize &&
6800 "Cannot extract more element than there are in the input.");
6807 InsertionPoint.
getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
6816 MachineInstrBuilder MIB =
BuildMI(*InsertionPoint.
getParent(), InsertionPoint,
6818 TII.get(SPIRV::OpCompositeConstruct))
6822 for (
Register ComponentReg : ComponentRegisters)
6823 MIB.
addUse(ComponentReg);
6828bool SPIRVInstructionSelector::selectImageWriteIntrinsic(
6829 MachineInstr &
I)
const {
6836 Register ImageReg =
I.getOperand(1).getReg();
6844 Register CoordinateReg =
I.getOperand(2).getReg();
6845 Register DataReg =
I.getOperand(3).getReg();
6848 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpImageWrite))
6856Register SPIRVInstructionSelector::buildPointerToResource(
6857 SPIRVTypeInst SpirvResType, SPIRV::StorageClass::StorageClass SC,
6858 uint32_t Set, uint32_t
Binding, uint32_t ArraySize,
Register IndexReg,
6859 StringRef Name, MachineIRBuilder MIRBuilder)
const {
6861 if (ArraySize == 1) {
6862 SPIRVTypeInst PtrType =
6865 "SpirvResType did not have an explicit layout.");
6870 const Type *VarType = ArrayType::get(
const_cast<Type *
>(ResType), ArraySize);
6871 SPIRVTypeInst VarPointerType =
6874 VarPointerType, Set,
Binding, Name, MIRBuilder);
6876 SPIRVTypeInst ResPointerType =
6889bool SPIRVInstructionSelector::selectFirstBitSet16(
6890 Register ResVReg, SPIRVTypeInst ResType, MachineInstr &
I,
6891 unsigned ExtendOpcode,
unsigned BitSetOpcode)
const {
6893 if (!selectOpWithSrcs(ExtReg, ResType,
I, {
I.getOperand(2).getReg()},
6897 return selectFirstBitSet32(ResVReg, ResType,
I, ExtReg, BitSetOpcode);
6900bool SPIRVInstructionSelector::selectFirstBitSet32(
6902 unsigned BitSetOpcode)
const {
6903 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
6906 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
6913bool SPIRVInstructionSelector::selectFirstBitSet64(
6915 unsigned BitSetOpcode,
bool SwapPrimarySide)
const {
6929 if (ComponentCount > 2) {
6930 auto Func = [
this, SwapPrimarySide](
Register ResVReg, SPIRVTypeInst ResType,
6932 unsigned Opcode) ->
bool {
6933 return this->selectFirstBitSet64(ResVReg, ResType,
I, SrcReg, Opcode,
6937 return handle64BitOverflow(ResVReg, ResType,
I, SrcReg, BitSetOpcode, Func);
6941 MachineIRBuilder MIRBuilder(
I);
6943 BaseType, 2 * ComponentCount, MIRBuilder,
false);
6947 if (!selectOpWithSrcs(BitcastReg, PostCastType,
I, {SrcReg},
6953 if (!selectFirstBitSet32(FBSReg, PostCastType,
I, BitcastReg, BitSetOpcode))
6963 if (!selectOpWithSrcs(HighReg, ResType,
I, {FBSReg, ConstIntOne},
6964 SPIRV::OpVectorExtractDynamic))
6966 if (!selectOpWithSrcs(LowReg, ResType,
I, {FBSReg, ConstIntZero},
6967 SPIRV::OpVectorExtractDynamic))
6971 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6972 TII.get(SPIRV::OpVectorShuffle))
6980 for (
unsigned J = 1; J < ComponentCount * 2; J += 2) {
6986 MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
6987 TII.get(SPIRV::OpVectorShuffle))
6995 for (
unsigned J = 0; J < ComponentCount * 2; J += 2) {
7015 SelectOp = SPIRV::OpSelectSISCond;
7016 AddOp = SPIRV::OpIAddS;
7024 SelectOp = SPIRV::OpSelectVIVCond;
7025 AddOp = SPIRV::OpIAddV;
7031 Register RegSecondaryOffset = Reg0;
7035 if (SwapPrimarySide) {
7036 PrimaryReg = LowReg;
7037 SecondaryReg = HighReg;
7038 RegPrimaryOffset = Reg0;
7039 RegSecondaryOffset = Reg32;
7044 if (!selectOpWithSrcs(RegSecondaryHasVal, BoolType,
I,
7045 {SecondaryReg, NegOneReg}, SPIRV::OpINotEqual))
7050 if (!selectOpWithSrcs(RegPrimaryHasVal, BoolType,
I, {PrimaryReg, NegOneReg},
7051 SPIRV::OpINotEqual))
7058 if (!selectOpWithSrcs(RegReturnBits, ResType,
I,
7059 {RegSecondaryHasVal, SecondaryReg, NegOneReg},
7064 if (SwapPrimarySide) {
7066 if (!selectOpWithSrcs(RegAdd, ResType,
I,
7067 {RegSecondaryHasVal, RegSecondaryOffset, Reg0},
7078 if (!selectOpWithSrcs(RegReturnBits2, ResType,
I,
7079 {RegPrimaryHasVal, PrimaryReg, RegReturnBits},
7084 if (!selectOpWithSrcs(RegAdd2, ResType,
I,
7085 {RegPrimaryHasVal, RegPrimaryOffset, RegAdd}, SelectOp))
7088 return selectOpWithSrcs(ResVReg, ResType,
I, {RegReturnBits2, RegAdd2},
7092bool SPIRVInstructionSelector::selectFirstBitHigh(
Register ResVReg,
7093 SPIRVTypeInst ResType,
7095 bool IsSigned)
const {
7097 Register OpReg =
I.getOperand(2).getReg();
7100 unsigned ExtendOpcode = IsSigned ? SPIRV::OpSConvert : SPIRV::OpUConvert;
7101 unsigned BitSetOpcode = IsSigned ? GL::FindSMsb : GL::FindUMsb;
7105 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
7107 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
7109 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
7112 return diagnoseUnsupported(
7114 "spv_firstbituhigh and spv_firstbitshigh only support 16,32,64 bits.");
7118bool SPIRVInstructionSelector::selectFirstBitLow(
Register ResVReg,
7119 SPIRVTypeInst ResType,
7120 MachineInstr &
I)
const {
7122 Register OpReg =
I.getOperand(2).getReg();
7127 unsigned ExtendOpcode = SPIRV::OpUConvert;
7128 unsigned BitSetOpcode = GL::FindILsb;
7132 return selectFirstBitSet16(ResVReg, ResType,
I, ExtendOpcode, BitSetOpcode);
7134 return selectFirstBitSet32(ResVReg, ResType,
I, OpReg, BitSetOpcode);
7136 return selectFirstBitSet64(ResVReg, ResType,
I, OpReg, BitSetOpcode,
7139 return diagnoseUnsupported(
I,
7140 "spv_firstbitlow only supports 16,32,64 bits.");
7144bool SPIRVInstructionSelector::selectAllocaArray(
Register ResVReg,
7145 SPIRVTypeInst ResType,
7146 MachineInstr &
I)
const {
7151 bool UseUntypedPointers =
7152 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7153 unsigned Opcode = UseUntypedPointers
7154 ? SPIRV::OpUntypedVariableLengthArrayINTEL
7155 : SPIRV::OpVariableLengthArrayINTEL;
7157 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(Opcode))
7163 if (UseUntypedPointers) {
7166 "untyped variable length array result must have a recorded element "
7171 MIB.
addUse(
I.getOperand(2).getReg());
7175 unsigned Alignment =
I.getOperand(3).getImm();
7189 while (!Worklist.
empty()) {
7191 switch (
T->getOpcode()) {
7192 case SPIRV::OpTypeInt:
7193 case SPIRV::OpTypeFloat:
7194 case SPIRV::OpTypePointer:
7196 case SPIRV::OpTypeVector:
7197 case SPIRV::OpTypeVectorIdEXT:
7198 case SPIRV::OpTypeMatrix:
7199 case SPIRV::OpTypeArray: {
7200 Register OperandReg =
T->getOperand(1).getReg();
7204 case SPIRV::OpTypeStruct:
7205 for (
unsigned Idx = 1,
E =
T->getNumOperands(); Idx <
E; ++Idx) {
7206 Register OperandReg =
T->getOperand(Idx).getReg();
7221 Register TypeReg = Ty->getOperand(0).getReg();
7222 if (!Visited.
insert(TypeReg).second)
7225 switch (Ty->getOpcode()) {
7226 case SPIRV::OpTypePointer:
7227 if (Ty->getOperand(1).getImm() == SPIRV::StorageClass::StorageBuffer)
7231 case SPIRV::OpTypeArray:
7232 case SPIRV::OpTypeRuntimeArray:
7235 case SPIRV::OpTypeStruct:
7236 for (
unsigned I = 1;
I < Ty->getNumOperands(); ++
I)
7252bool SPIRVInstructionSelector::selectAbort(MachineInstr &
I)
const {
7253 assert(
I.getNumExplicitOperands() == 2);
7255 Register MsgReg =
I.getOperand(1).getReg();
7257 assert(MsgType &&
"Message argument of llvm.spv.abort has no SPIR-V type");
7260 return diagnoseUnsupported(
7262 "llvm.spv.abort message type must be a concrete SPIR-V type (numerical "
7263 "scalar, pointer, vector, matrix, or aggregate of such types)");
7266 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7273bool SPIRVInstructionSelector::selectTrap(MachineInstr &
I)
const {
7284 uint32_t MsgVal = ~0
u;
7285 if (
I.getOpcode() == TargetOpcode::G_UBSANTRAP)
7286 MsgVal =
static_cast<uint32_t
>(
I.getOperand(0).
getImm());
7289 Register MsgReg = buildI32ConstantInEntryBlock(MsgVal,
I, MsgType);
7292 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpAbortKHR))
7299bool SPIRVInstructionSelector::selectFrameIndex(
Register ResVReg,
7300 SPIRVTypeInst ResType,
7301 MachineInstr &
I)
const {
7308 bool UseUntypedPointers =
7309 ResType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7311 UseUntypedPointers ? SPIRV::OpUntypedVariableKHR : SPIRV::OpVariable;
7314 MachineIRBuilder MIRBuilder(
I);
7317 .
addImm(SPIRV::Extension::SPV_KHR_variable_pointers);
7319 .
addImm(SPIRV::Capability::VariablePointersStorageBuffer);
7322 auto MIB =
BuildMI(*It->getParent(), It, It->getDebugLoc(),
TII.get(Opcode))
7325 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7329 if (UseUntypedPointers) {
7333 return diagnoseUnsupported(
7334 I,
"could not deduce the data type of an untyped variable");
7340 unsigned Alignment =
I.getOperand(2).getImm();
7347bool SPIRVInstructionSelector::selectBranch(MachineInstr &
I)
const {
7352 const MachineInstr *PrevI =
I.getPrevNode();
7354 if (PrevI !=
nullptr && PrevI->
getOpcode() == TargetOpcode::G_BRCOND) {
7358 .
addMBB(
I.getOperand(0).getMBB())
7363 .
addMBB(
I.getOperand(0).getMBB())
7368bool SPIRVInstructionSelector::selectBranchCond(MachineInstr &
I)
const {
7379 const MachineInstr *NextI =
I.getNextNode();
7381 if (NextI !=
nullptr && NextI->
getOpcode() == SPIRV::OpBranchConditional)
7387 MachineBasicBlock *NextMBB =
I.getMF()->getBlockNumbered(NextMBBNum);
7389 .
addUse(
I.getOperand(0).getReg())
7390 .
addMBB(
I.getOperand(1).getMBB())
7396bool SPIRVInstructionSelector::selectPhi(
Register ResVReg,
7397 MachineInstr &
I)
const {
7399 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(TargetOpcode::PHI))
7401 const unsigned NumOps =
I.getNumOperands();
7402 for (
unsigned i = 1; i <
NumOps; i += 2) {
7403 MIB.
addUse(
I.getOperand(i + 0).getReg());
7404 MIB.
addMBB(
I.getOperand(i + 1).getMBB());
7410bool SPIRVInstructionSelector::selectGlobalValue(
7411 Register ResVReg, MachineInstr &
I,
const MachineInstr *Init)
const {
7413 MachineIRBuilder MIRBuilder(
I);
7414 const GlobalValue *GV =
I.
getOperand(1).getGlobal();
7417 std::string GlobalIdent;
7419 unsigned &
ID = UnnamedGlobalIDs[GV];
7421 ID = UnnamedGlobalIDs.
size();
7422 GlobalIdent =
"__unnamed_" + Twine(ID).str();
7448 GVFun ? SPIRV::StorageClass::CodeSectionINTEL
7455 MachineRegisterInfo *
MRI = MIRBuilder.
getMRI();
7460 MachineInstrBuilder MIB1 =
7461 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7464 MachineInstrBuilder MIB2 =
7466 TII.get(SPIRV::OpConstantFunctionPointerINTEL))
7470 GR.
add(ConstVal, MIB2);
7478 MachineInstrBuilder MIB3 =
7479 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpUndef))
7482 GR.
add(ConstVal, MIB3);
7488 assert(NewReg != ResVReg);
7489 return BuildCOPY(ResVReg, NewReg,
I);
7499 const std::optional<SPIRV::LinkageType::LinkageType> LnkType =
7502 if (LnkType && *LnkType == SPIRV::LinkageType::Import)
7508 SPIRVTypeInst ResType =
7512 GlobalVar->isConstant(), LnkType, MIRBuilder,
true);
7517 if (
GlobalVar->isExternallyInitialized() &&
7518 STI.getTargetTriple().getVendor() ==
Triple::AMD) {
7519 constexpr unsigned ReadWriteINTEL = 3u;
7522 MachineInstrBuilder MIB(*MF, --MIRBuilder.
getInsertPt());
7528bool SPIRVInstructionSelector::selectLog10(
Register ResVReg,
7529 SPIRVTypeInst ResType,
7530 MachineInstr &
I)
const {
7532 return selectExtInst(ResVReg, ResType,
I, CL::log10);
7540 MachineIRBuilder MIRBuilder(
I);
7545 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7548 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::GLSL_std_450))
7550 .
add(
I.getOperand(1))
7564 APFloat::rmNearestTiesToEven, &LosesInfo);
7569 isVectorType(ResType) ? SPIRV::OpVectorTimesScalar : SPIRV::OpFMulS;
7579bool SPIRVInstructionSelector::selectFpowi(
Register ResVReg,
7580 SPIRVTypeInst ResType,
7581 MachineInstr &
I)
const {
7584 return selectExtInst(ResVReg, ResType,
I, CL::pown);
7590 Register ExpReg =
I.getOperand(2).getReg();
7592 if (!selectOpWithSrcs(FloatExpReg, ResType,
I, {ExpReg},
7593 SPIRV::OpConvertSToF))
7595 return selectExtInst(ResVReg, ResType,
I, GL::Pow,
7602bool SPIRVInstructionSelector::selectModf(
Register ResVReg,
7603 SPIRVTypeInst ResType,
7604 MachineInstr &
I)
const {
7620 MachineIRBuilder MIRBuilder(
I);
7621 SPIRVTypeInst FloatType =
7625 FloatType, MIRBuilder, SPIRV::StorageClass::Function);
7638 MachineBasicBlock &EntryBB =
I.getMF()->
front();
7639 const bool IsUntyped =
7640 PtrType->
getOpcode() == SPIRV::OpTypeUntypedPointerKHR;
7642 BuildMI(EntryBB, VarPos,
I.getDebugLoc(),
7643 TII.get(IsUntyped ? SPIRV::OpUntypedVariableKHR
7644 : SPIRV::OpVariable))
7647 .
addImm(
static_cast<uint32_t
>(SPIRV::StorageClass::Function));
7655 BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpExtInst))
7658 .
addImm(
static_cast<uint32_t
>(SPIRV::InstructionSet::OpenCL_std))
7661 .
add(
I.getOperand(
I.getNumExplicitDefs()))
7665 Register IntegralPartReg =
I.getOperand(1).getReg();
7668 auto LoadMIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7678 assert(
false &&
"GLSL::Modf is deprecated.");
7689bool SPIRVInstructionSelector::loadVec3BuiltinInputID(
7690 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7691 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7692 MachineIRBuilder MIRBuilder(
I);
7693 const SPIRVTypeInst Vec3Ty =
7696 Vec3Ty, MIRBuilder, SPIRV::StorageClass::Input);
7708 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7712 MachineRegisterInfo *
MRI = MIRBuilder.
getMRI();
7718 BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7725 assert(
I.getOperand(2).isReg());
7726 const uint32_t ThreadId =
foldImm(
I.getOperand(2),
MRI);
7730 auto MIB =
BuildMI(BB,
I,
I.getDebugLoc(),
TII.get(SPIRV::OpCompositeExtract))
7741bool SPIRVInstructionSelector::loadBuiltinInputID(
7742 SPIRV::BuiltIn::BuiltIn BuiltInValue,
Register ResVReg,
7743 SPIRVTypeInst ResType, MachineInstr &
I)
const {
7744 MachineIRBuilder MIRBuilder(
I);
7746 ResType, MIRBuilder, SPIRV::StorageClass::Input);
7761 SPIRV::StorageClass::Input,
nullptr,
true, std::nullopt, MIRBuilder,
7765 auto MIB =
BuildMI(*
I.getParent(),
I,
I.getDebugLoc(),
TII.get(SPIRV::OpLoad))
7774SPIRVTypeInst SPIRVInstructionSelector::widenTypeToVec4(SPIRVTypeInst
Type,
7775 MachineInstr &
I)
const {
7776 MachineIRBuilder MIRBuilder(
I);
7787bool SPIRVInstructionSelector::loadHandleBeforePosition(
7788 Register &HandleReg, SPIRVTypeInst ResType, GIntrinsic &HandleDef,
7789 MachineInstr &Pos)
const {
7792 Intrinsic::spv_resource_handlefrombinding);
7800 bool IsStructuredBuffer = ResType->
getOpcode() == SPIRV::OpTypePointer;
7801 MachineIRBuilder MIRBuilder(HandleDef);
7802 SPIRVTypeInst VarType = ResType;
7803 SPIRV::StorageClass::StorageClass SC = SPIRV::StorageClass::UniformConstant;
7805 if (IsStructuredBuffer) {
7814 .
addImm(SPIRV::Capability::RuntimeDescriptorArrayEXT);
7817 buildPointerToResource(SPIRVTypeInst(VarType), SC, Set,
Binding,
7818 ArraySize, IndexReg, Name, MIRBuilder);
7822 uint32_t LoadOpcode =
7823 IsStructuredBuffer ? SPIRV::OpCopyObject : SPIRV::OpLoad;
7833bool SPIRVInstructionSelector::errorIfInstrOutsideShader(
7834 MachineInstr &
I)
const {
7836 return diagnoseUnsupported(
7837 I,
"this instruction is only supported in shaders.");
7842InstructionSelector *
7846 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.
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